Using the Callback API with the SANITIZER_CB_DOMAIN_DRIVER_API or SANITIZER_CB_DOMAIN_RUNTIME_API domains, a callback function can be associated with one or more CUDA API functions. When those CUDA functions are called in the application, the callback function is invoked as well. For these domains, the cbdata argument to the callback function will be of the type Sanitizer_CallbackData.

You can call cudaDeviceSynchronize, cudaStreamSynchronize, cuCtxSynchronize and cuStreamSynchronize from within a driver or runtime API callback function.

The following code shows a typical sequence used to associate a callback function with one or more CUDA API functions. For simplicity, error checking code was removed.

Sanitizer_SubscriberHandle handle;
MyDataStruct *my_data = ...;
...
sanitizerSubscribe(&handle, my_callback, my_data);
sanitizerEnableDomain(1, handle, SANITIZER_CB_DOMAIN_RUNTIME_API);
        

First, sanitizerSubscribe is used to initialize a subscriber with the my_callback callback function. Next, sanitizerEnableDomain is used to associate that callback with all the CUDA runtime functions. Using this code sequence will cause my_callback to be called twice each time any of the CUDA runtime API functions are invoked, once on entry to the CUDA function and once just before the CUDA function exits. Compute Sanitizer callback API functions sanitizerEnableCallback and sanitizerEnableAllDomains can also be used to associate CUDA API functions with a callback.

The following code shows a typical callback function.

void SANITIZERAPI
my_callback(void *userdata,
            Sanitizer_CallbackDomain domain,
            Sanitizer_CallbackId cbid,
            const void *cbdata)
{
    const Sanitizer_CallbackData *cbInfo = (Sanitizer_CallbackData *)cbdata;
    MyDataStruct *my_data = (MyDataStruct *)userdata;

    if ((domain == SANITIZER_CB_DOMAIN_RUNTIME_API) &&
        (cbid == SANITIZER_RUNTIME_TRACE_CBID_cudaMemcpy_v3020) &&
        (cbInfo->callbackSite == SANITIZER_API_ENTER))
    {
        cudaMemcpy_v3020_params *funcParams = (cudaMemcpy_v3020_params *)(cbInfo->functionParams);
        size_t count = funcParams->count;
        enum cudaMemcpyKind kind = funcParams->kind
        ...
    }
    ...
}
        

In the callback function, Sanitizer_CallbackDomain and Sanitizer_CallbackId parameters can be used to determine which CUDA API function invocation is triggering this callback. In the example above, we are checking for the CUDA runtime cudaMemcpy function. The cbdata parameter holds a structure of useful information that can be used within the callback. In this case, we use the callbackSite member of the structure to detect that the callback is occurring on entry to cudaMemcpy, and we use the functionParams member to access the parameters to cudaMemcpy. To access the parameters, we first cast functionParams to a structure type corresponding to the cudaMemcpy function. These parameter structures are contained in generated_cuda_runtime_api_meta.h, generated_cuda_meta.h and a number of other files.

Using the Callback API with the SANITIZER_CB_DOMAIN_SYNCHRONIZE domain, a callback function can be associated with CUDA context and stream synchronizations. For example, when a CUDA context is synchronized, the callback function is invoked with a callback ID equal to SANITIZER_CBID_SYNCHRONIZE_CONTEXT_SYNCHRONIZED. For this domain, the cbdata argument is of the type Sanitizer_SynchronizeData.

Using the Callback API with the SANITIZER_CB_DOMAIN_LAUNCH domain, a callback function can be associated with CUDA kernel launches. For example, when a CUDA kernel launch has started, the callback function is invoked with a callback ID equal to SANITIZER_CBID_LAUNCH_BEGIN. For this domain, the cbdata argument is of the type Sanitizer_LaunchData.

Using the Callback API with the SANITIZER_CB_DOMAIN_MEMCPY domain, a callback function can be associated with CUDA memcpy operations. For example, when a cudaMemcpy API function is called, the callback function is invoked with a callback ID equal to SANITIZER_CBID_MEMCPY_STARTING. For this domain, the cbdata argument is of the type Sanitizer_MemcpyData.

Using the Callback API with the SANITIZER_CB_DOMAIN_MEMSET domain, a callback function can be associated with CUDA memset operations. For example, when a cudaMemset API function is called, the callback function is invoked with a callback ID equal to SANITIZER_CBID_MEMSET_STARTING. For this domain, the cbdata argument is of the type Sanitizer_MemsetData.

Using the Callback API with the SANITIZER_CB_DOMAIN_BATCH_MEMOP domain, a callback function can be associated with CUDA batch memory operations. For example, when a cuStreamWriteValue API function is called, the callback function is invoked with a callback ID equal to SANITIZER_CBID_BATCH_MEMOP_WRITE. For this domain, the cbdata argument is of the type Sanitizer_BatchMemopData.

The patch must follow the function signature required by the Compute Sanitizer API for a given instruction ID. The mapping of instruction ID to function signature is documented in the comments of Sanitizer_InstructionId in sanitizer_patching.h. For instance, if we wish to patch a memory access using the instruction ID SANITIZER_INSTRUCTION_MEMORY_ACCESS, we need to use the SanitizerCallbackMemoryAccess type.

extern "C" __device__
SanitizerPatchResult SANITIZERAPI my_memory_access_callback(
    void* userdata,
    uint64_t pc,
    void* ptr,
    uint32_t accessSize,
    uint32_t flags)
{
    MyDeviceDataStruct *my_data = (MyDeviceDataStruct *)userdata

    if ((flags & SANITIZER_MEMORY_DEVICE_FLAG_WRITE) != 0)
        // log write
    else
        // log read

    return SANITIZER_PATCH_SUCCESS;
}
        

In this patch, we log write and read accesses to a structure we allocated previously. extern "C" ensures that the patch name will not be mangled, allowing us to use its name as a string directly in calls to sanitizerPatchInstructions (see below).

There can be multiple patches defined in a single CUDA file. This file must then be compiled using the following nvcc options:

$ nvcc --cubin --compile-as-tools-patch MySanitizerPatches.cu -o MySanitizerPatches.cubin
        

The --cubin option can be replaced by --fatbin if a fatbin is preferred over a cubin as the output file.

Once the patch has been generated, it can be inserted into user code by using the following procedure:

1. Load the patch. There are two APIs used to load the patch: sanitizerAddPatchesFromFile and sanitizerAddPatches. They use the same input format as cuModuleLoad and cuModuleLoadData, respectively.
2. Select which instructions to patch by using the sanitizerPatchInstructions API.
3. Patch user code by using the sanitizerPatchModule API.
4. Optionally, set the callback data for patches by using the sanitizerSetCallbackData API.

The following code shows a typical sequence using these APIs. For simplicity, error checking was removed.

CUcontext ctx = ... // current CUDA context
sanitizerAddPatchesFromFile("MySanitizerPatches.cubin", ctx);

CUmodule module = ... // module containing the user code
sanitizerPatchInstructions(SANITIZER_INSTRUCTION_MEMORY_ACCESS, module, "my_memory_access_callback");

sanitizerPatchModule(module);

MyDeviceDataTracker *deviceDataTracker;
cudaMalloc(&deviceDataTracker, sizeof(*deviceDataTracker));

CUfunction function = ... // kernel to be launched for which we want to set the callbackdata for the patches
sanitizerSetCallbackData(function, deviceDataTracker);
        

All subsequent launches using code from this CUDA module will be instrumented and my_memory_access_callback will be invoked before every memory access. However, the callback data is only set for all subsequent launches of the given kernel. An easy way to have a kernel CUfunction, is through the Sanitizer launch callbacks. Instrumentation can be removed by using the sanitizerUnpatchModule API.

When the application instrumented by the Compute Sanitizer API makes use of device launched CUDA graphs, the following behaviors are to be taken into account:

• No callbacks will be called from the host when a CUDA graph is launched from the device. However, graph launch-related callbacks are still called when a device launchable graph is launched from the host, or uploaded.
• Patching API issued callbacks on the device when inside a device launched CUDA graph may have a userdata from a different node if different userdata pointers were set per node.

void* userdata = **((void***)cudaGetCurrentGraphExec());