NVIDIA OptiX ray tracing engine is an software framework for reaching optimum ray tracing efficiency on the GPU. Functions utilizing OptiX can fail in methods which can be troublesome to diagnose: an invalid API argument, a black body, or a GPU-side bug buried underneath hundreds of concurrent threads.
Debugging services within the NVIDIA OptiX Toolkit (OTK) might help. OTK is a GitHub repository containing a set of utilities supporting workflows widespread in GPU ray tracing functions. OTK is licensed with a BSD 3-clause model license, so you might be free to repeat and modify any of the code.
This submit covers two OTK debugging services: constant checking of OptiX and CUDA API error codes and focused device-side debug printing. OTK consists of an instance program, DemandPbrtScene, that reveals device-side debug printing utilized in context.
How do OptiX logging and validation work?
Earlier than attending to the help supplied by OTK, some background on the OptiX log could be helpful. Whenever you create an OptiX gadget context, you provide an choices construction that permits you to configure a log callback and a validation mode. OptiX can help in validating inputs to API capabilities when the validation mode is ready to OPTIX_DEVICE_CONTEXT_VALIDATION_MODE_ALL.
OptiX writes human-readable messages for validation errors to the log. The log output needs to be the primary place you search for API errors like OPTIX_ERROR_INVALID_VALUE. Validation does impose some extra value within the API. It’s really helpful to allow all validation in debug and testing builds and omitting validation for launch builds. For extra particulars, see the OptiX SDK samples.
How one can constantly test return codes for errors
It’s greatest to detect errors as early as doable as detailed on this part, earlier than subsequent failures masks the unique downside.
API mechanisms
Most capabilities in OptiX return an OptixResult error code that, when non-zero, signifies an error. The CUDA runtime API and the CUDA driver API comply with the same sample. All three APIs help the next:
A definite enumerated sort for the error code; for instance OptixResult
A operate to return a symbolic identify for an error code as a string; for instance OPTIX_ERROR_INVALID_VALUE
A operate to return a human-readable error message for an error code; for instance, Invalid worth
The signatures of those capabilities are barely completely different for the three APIs, however the mechanisms are the identical.
Error checking coverage
Dealing with these error codes manually at each API name web site is tedious and error-prone. It’s higher to make use of macros or operate calls to implement a constant coverage for dealing with errors. OTK gives macros that implement two insurance policies when an error is detected:
Different insurance policies are simply applied by reusing among the supplied equipment and creating an acceptable macro.
Minimal use of macro equipment
This error checking mechanism makes use of macros to a minimal extent and delegates to inline capabilities to do the precise work. You may set breakpoints within the inline operate definitions to have the debugger cease execution when the operate detects an error.
The macros exist to supply diagnostic details about the code that brought about the error:
expr: A string type of the provided argument to the macro. That is the expression that evaluates to the error code.
__FILE__: The identify of the supply file the place the macro was invoked.
__LINE__: The road quantity throughout the supply file the place the macro was invoked.
This info from the macro name web site is handed to the inline operate that does the precise error checking.
Unified error checking throughout APIs
The inline template operate checkError checks the standing code for an error and creates a diagnostic message if it detects a failure. Within the case of those three APIs, a easy solid of the error code to bool suffices to point an error. All three APIs use a standing code of zero to point success and non-zero to point failure.
Error messages are formatted as follows:
The place expr is the evaluated expression, nnn is the results of casting the standing code to an int, identify is the symbolic identify of the standing code, and message is the human-readable error message. If both the identify or message are empty, the operate omits them.
The inline template operate makeErrorString is accountable for constructing this message. It calls the inline template capabilities getErrorName and getErrorMessage to construct the mixed message.
Every API has a definite sort for API standing codes, so you possibly can specialize the template capabilities to carry out the suitable API name to get the prolonged error info.
Utilization
OTK gives one header per API that gives the mandatory specializations of the template capabilities described (Desk 1).
Merely embrace the headers for the APIs you might be utilizing and use the only macro OTK_ERROR_CHECK round all name websites. The next instance makes use of all three APIs.
OTK_ERROR_CHECK( cuCtxGetCurrent( &m_cudaContext ) );
OTK_ERROR_CHECK( cuStreamCreate( &m_stream, CU_STREAM_DEFAULT ) );
OTK_ERROR_CHECK( optixInit() );
How one can carry out focused device-side debug printing
The issue with graphics functions is that there are too some ways to code up a black display.
To debug issues in your OptiX gadget code, you possibly can take a number of approaches. A few apparent selections spring to thoughts:
Use the CUDA debugger on a debug construct of the gadget code
Get info from a launch construct of the gadget code with printf
Many functions run too slowly when compiled in debug mode, impairing the usage of interactive debuggers.
The primary problem with printf-style debugging is that there are such a lot of threads working concurrently on the GPU that you would be able to find yourself drowning in a firehose of output. As well as, it could possibly be that the difficulty solely arises after a certain quantity of interplay with the appliance. Debug output previous to the issue manifesting visually is simply noise that will get in the way in which of discovering the specified info.
DebugLocation
The header gives a reusable mechanism for debug output. The construction DebugLocation controls the conduct:
{
bool enabled;
bool dumpSuppressed;
bool debugIndexSet;
uint3 debugIndex;
};
The enabled member turns your complete mechanism on or off. The dumpSuppressed member turns off the debug output even when the mechanism is enabled. The debugIndexSet member signifies {that a} legitimate launch index has been saved in debugIndex.
The mechanism will output debug info when the next circumstances are true:
enabled is true
dumpSuppressed is fake
debugIndexSet is true, and
The present launch index matches debugIndex
Embody an occasion of the DebugLocation construction within the launch parameters of the OptiX pipeline for interactive management of debug output.
The debugInfoDump operate
The template operate debugInfoDump gives the interface for emitting debug info:
static __forceinline__ __device__
bool debugInfoDump( const DebugLocation& debug,
const Callback &callback )
The Callback template parameter needs to be a struct or class matching the next:
{
void setColor( float pink, float inexperienced, float blue );
void dump( const uint3& index );
};
The setColor methodology is used to attract a visible field across the debug location for simple identification of the purpose on the display for which info is dumped. The everyday utilization is to set the colour for the output pixel equivalent to the present launch index. If no visible indication of the debug location is desired, the strategy can merely be empty.
The dump methodology is used to print no matter info the appliance considers related on the provided launch index.
Displaying the debug location
When enabled, the setColor methodology on the callback construction attracts a field on display to point the present debug location, even when dump output is suppressed. The field is hidden when disabled.
A pink pixel on the debug location sits inside a one-pixel-wide black field, itself inside a one-pixel-wide white field. This gives a high-contrast indicator of the situation the place dump messages are occurring. If the output buffer is just not a conventional colour buffer, you might be free to map the provided pink, inexperienced, and blue values to some distinctive worth for visualization.
One-shot mode
To keep away from drowning in debug output, it’s helpful to have the debug output in a one-shot mode, the place the output is dumped as soon as in response to consumer management. You may sequence this mode as follows:
Allow the DebugLocation mechanism
Launch as standard
When the consumer interactively selects the debug location, set dumpSuppressed to true, debugIndexSet to true, and debugIndex to the chosen location
Subsequent launches show the debug location, however the mechanism gives no dumps
The consumer interacts with the appliance to govern it into the suitable state, probably transferring the debug location alongside the way in which
When the consumer signifies that they need debug info for the present location, set dumpSuppressed to false
Launch to get the debug output
Set dumpSuppressed again to true after launch
DemandPbrtScene instance
The DemandPbrtScene instance in OTK demonstrates demand loaded geometry for pbrt model 3 scenes. It makes use of the DebugLocation mechanism, together with one-shot conduct and interactive toggling of debug output and interactive number of the debug place. It makes use of ImGui because the UI framework.
To run this instance from OTK, you have to a scene file for pbrt-v3.
The OptiX Toolkit gives reusable debugging and testing utilities for widespread OptiX improvement issues: constant API error checking and focused device-side debug output. The code is out there via the NVIDIA/optix-toolkit GitHub repo.
Able to get began? Obtain the OptiX Toolkit from GitHub and begin by enabling OptiX validation in your debug builds, wrapping your CUDA and OptiX calls with OTK_ERROR_CHECK, and utilizing DebugLocation to isolate GPU-side bugs earlier in improvement. OTK is out there underneath a permissive BSD 3-clause-style license, so you possibly can copy, adapt, and combine these utilities straight into your individual OptiX functions.
