Sanitizer Patching API


Typedefs

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackAsyncReduction )(void *userdata, uint64_t pc, uint32_t address, uint32_t mbarAddress, uint32_t accessSize)
 Function type for an asynchronous reduction operation on shared memory.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackAsyncStore )(void *userdata, uint64_t pc, uint32_t address, uint32_t mbarAddress, void *pNewValue, uint32_t accessSize)
 Function type for an asynchronous store operation on shared memory.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackBarrier )(void *userdata, uint64_t pc, uint32_t barIndex, uint32_t threadCount, uint32_t flags)
 Function type for a barrier callback.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackBlockEnter )(void *userdata, uint64_t pc)
 Function type for a CUDA block enter callback.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackBlockExit )(void *userdata, uint64_t pc)
 Function type for a CUDA block exit callback.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackCacheControl )(void *userdata, uint64_t pc, void *address, Sanitizer_CacheControlInstructionKind kind)
 Function type for a cache control instruction callback.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackCall )(void *userdata, uint64_t pc, uint64_t targetPc, uint32_t flags)
 Function type for a function call callback.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackClusterBarrierArrive )(void *userdata, uint64_t pc)
 Function type for a cluster barrier arrive.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackCudaBarrier )(void *userdata, uint64_t pc, void *barrier, uint32_t kind, uint32_t data)
 Function type for a CUDA Barrier action callback.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackDeviceSideFree )(void *userdata, uint64_t pc, void *ptr)
 Function type for a device-side free call.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackDeviceSideMalloc )(void *userdata, uint64_t pc, void *allocatedPtr, uint64_t allocatedSize)
 Function type for a device-side malloc call.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackMatrixMemoryAccess )(void *userdata, uint64_t pc, uint32_t address, uint32_t accessSize, uint32_t flags, uint32_t count, const void *pNewValue)
 Function type for a matrix shared memory access callback.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackMemcpyAsync )(void *userdata, uint64_t pc, void *src, uint32_t dst, uint32_t accessSize)
 Function type for a global to shared memory asynchronous copy.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackMemoryAccess )(void *userdata, uint64_t pc, void *ptr, uint32_t accessSize, uint32_t flags, const void *pData)
 Function type for a memory access callback.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackPipelineCommit )(void *userdata, uint64_t pc)
 Function type for a pipeline commit.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackPipelineWait )(void *userdata, uint64_t pc, uint32_t groups)
 Function type for a pipeline wait.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackRet )(void *userdata, uint64_t pc)
 Function type for a function return callback.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackSetSmemSize )(void *userdata, uint64_t pc, uint32_t size)
 Function type for setting the shared memory size allocated to a block.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackShfl )(void *userdata, uint64_t pc)
 Function type for a shfl callback.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackSyncwarp )(void *userdata, uint64_t pc, uint32_t mask)
 Function type for a syncwarp callback.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackWarpgroupFence )(void *userdata, uint64_t pc, uint32_t warpMask)
 Function type for a warpgroup MMA fence.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackWarpgroupMMAAsync )(void *userdata, uint64_t pc, uint32_t addressMatrixA, uint32_t sizeMatrixA, uint32_t addressMatrixB, uint32_t sizeMatrixB, uint32_t flags, uint32_t warpMask)
 Function type for a warpgroup aligned async MMA.
typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackWarpgroupWaitGroup )(void *userdata, uint64_t pc, uint32_t numGroups, uint32_t warpMask)
 Function type for a warpgroup MMA wait group.

Enumerations

enum  Sanitizer_BarrierFlags {
  SANITIZER_BARRIER_FLAG_NONE = 0,
  SANITIZER_BARRIER_FLAG_UNALIGNED_ALLOWED = 0x1
}
 Flags describing a barrier. More...
enum  Sanitizer_CacheControlInstructionKind {
  SANITIZER_CACHE_CONTROL_INVALID = 0,
  SANITIZER_CACHE_CONTROL_L1_PREFETCH = 1,
  SANITIZER_CACHE_CONTROL_L2_PREFETCH = 2
}
 Cache control action. More...
enum  Sanitizer_CallFlags {
  SANITIZER_CALL_FLAG_NONE = 0,
  SANITIZER_CALL_FLAG_UNALIGNED_ALLOWED = 0x1
}
 Flags describing a function call. More...
enum  Sanitizer_CudaBarrierInstructionKind {
  SANITIZER_CUDA_BARRIER_INVALID = 0,
  SANITIZER_CUDA_BARRIER_INIT = 1,
  SANITIZER_CUDA_BARRIER_ARRIVE = 2,
  SANITIZER_CUDA_BARRIER_ARRIVE_DROP = 3,
  SANITIZER_CUDA_BARRIER_ARRIVE_NOCOMPLETE = 4,
  SANITIZER_CUDA_BARRIER_ARRIVE_DROP_NOCOMPLETE = 5,
  SANITIZER_CUDA_BARRIER_WAIT = 6,
  SANITIZER_CUDA_BARRIER_INVALIDATE = 7
}
 CUDA Barrier action kind. More...
enum  Sanitizer_DeviceMemoryFlags {
  SANITIZER_MEMORY_DEVICE_FLAG_NONE = 0,
  SANITIZER_MEMORY_DEVICE_FLAG_READ = 0x1,
  SANITIZER_MEMORY_DEVICE_FLAG_WRITE = 0x2,
  SANITIZER_MEMORY_DEVICE_FLAG_ATOMSYS = 0x4,
  SANITIZER_MEMORY_DEVICE_FLAG_PREFETCH = 0x8
}
 Flags describing a memory access. More...
enum  Sanitizer_FunctionLoadedStatus {
  SANITIZER_FUNCTION_NOT_LOADED = 0x0,
  SANITIZER_FUNCTION_PARTIALLY_LOADED = 0x1,
  SANITIZER_FUNCTION_LOADED = 0x2
}
enum  Sanitizer_InstructionId {
  SANITIZER_INSTRUCTION_INVALID = 0,
  SANITIZER_INSTRUCTION_BLOCK_ENTER = 1,
  SANITIZER_INSTRUCTION_BLOCK_EXIT = 2,
  SANITIZER_INSTRUCTION_GLOBAL_MEMORY_ACCESS = 3,
  SANITIZER_INSTRUCTION_SHARED_MEMORY_ACCESS = 4,
  SANITIZER_INSTRUCTION_LOCAL_MEMORY_ACCESS = 5,
  SANITIZER_INSTRUCTION_BARRIER = 6,
  SANITIZER_INSTRUCTION_SYNCWARP = 7,
  SANITIZER_INSTRUCTION_SHFL = 8,
  SANITIZER_INSTRUCTION_CALL = 9,
  SANITIZER_INSTRUCTION_RET = 10,
  SANITIZER_INSTRUCTION_DEVICE_SIDE_MALLOC = 11,
  SANITIZER_INSTRUCTION_DEVICE_SIDE_FREE = 12,
  SANITIZER_INSTRUCTION_CUDA_BARRIER = 13,
  SANITIZER_INSTRUCTION_MEMCPY_ASYNC = 14,
  SANITIZER_INSTRUCTION_PIPELINE_COMMIT = 15,
  SANITIZER_INSTRUCTION_PIPELINE_WAIT = 16,
  SANITIZER_INSTRUCTION_REMOTE_SHARED_MEMORY_ACCESS = 17,
  SANITIZER_INSTRUCTION_DEVICE_ALIGNED_MALLOC = 18,
  SANITIZER_INSTRUCTION_MATRIX_MEMORY_ACCESS = 19,
  SANITIZER_INSTRUCTION_CACHE_CONTROL = 20,
  SANITIZER_INSTRUCTION_CLUSTER_BARRIER_ARRIVE = 21,
  SANITIZER_INSTRUCTION_CLUSTER_BARRIER_WAIT = 22,
  SANITIZER_INSTRUCTION_WARPGROUP_MMA_ASYNC = 23,
  SANITIZER_INSTRUCTION_WARPGROUP_WAIT_GROUP = 24,
  SANITIZER_INSTRUCTION_WARPGROUP_FENCE = 25,
  SANITIZER_INSTRUCTION_ASYNC_STORE = 26,
  SANITIZER_INSTRUCTION_ASYNC_REDUCTION = 27,
  SANITIZER_INSTRUCTION_SET_SHARED_MEMORY_SIZE = 28,
  SANITIZER_INSTRUCTION_BARRIER_RELEASE = 29
}
 Instrumentation. More...
enum  Sanitizer_WarpgroupMMAAsyncFlags {
  SANITIZER_WARPGROUP_MMA_ASYNC_FLAG_NONE = 0,
  SANITIZER_WARPGROUP_MMA_ASYNC_FLAG_COMMIT_GROUP = 0x1
}
 Flags describing a warpgroup aligned MMA async. More...
enum  SanitizerPatchResult {
  SANITIZER_PATCH_SUCCESS = 0,
  SANITIZER_PATCH_ERROR = 1
}
 Sanitizer patch result codes. More...

Functions

SanitizerResult SANITIZERAPI sanitizerAddPatches (const void *image, CUcontext ctx)
 Load a module containing patches that can be used by the patching API.
SanitizerResult SANITIZERAPI sanitizerAddPatchesFromFile (const char *filename, CUcontext ctx)
 Load a module containing patches that can be used by the patching API.
SanitizerResult SANITIZERAPI sanitizerGetCallbackPcAndSize (CUcontext ctx, const char *deviceCallbackName, uint64_t *pc, uint64_t *size)
 Get PC and size of a device callback.
SanitizerResult SANITIZERAPI sanitizerGetFunctionLoadedStatus (CUfunction func, Sanitizer_FunctionLoadedStatus *loadingStatus)
 Get the loading status of a function. Requires a driver version >=515.
SanitizerResult SANITIZERAPI sanitizerGetFunctionPcAndSize (CUmodule module, const char *functionName, uint64_t *pc, uint64_t *size)
 Get PC and size of a CUDA function.
SanitizerResult SANITIZERAPI sanitizerPatchInstructions (const Sanitizer_InstructionId instructionId, CUmodule module, const char *deviceCallbackName)
 Set instrumentation points and patches to be applied in a module.
SanitizerResult SANITIZERAPI sanitizerPatchModule (CUmodule module)
 Perform the actual instrumentation of a module.
SanitizerResult SANITIZERAPI sanitizerSetCallbackData (CUfunction kernel, const void *userdata)
 Specifies the user data pointer for callbacks.
SanitizerResult SANITIZERAPI sanitizerSetDeviceGraphData (CUgraphExec graphExec, Sanitizer_StreamHandle stream, const void *userdata)
 Specifies the user data pointer accessible from callbacks in the device-launched graphs launched by the specified host-launched graphExec.
SanitizerResult SANITIZERAPI sanitizerSetLaunchCallbackData (Sanitizer_LaunchHandle launch, CUfunction kernel, Sanitizer_StreamHandle stream, const void *userdata)
 Specifies the user data pointer for callbacks.
SanitizerResult SANITIZERAPI sanitizerUnpatchModule (CUmodule module)
 Remove existing instrumentation of a module.

Detailed Description

Functions, types, and enums that implement the Sanitizer Patching API.

Typedef Documentation

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackAsyncReduction)(void *userdata, uint64_t pc, uint32_t address, uint32_t mbarAddress, uint32_t accessSize)

This can be generated by a red.async PTX instruction

userdata is a pointer to user data. See sanitizerPatchModule. pc is the program counter of the patched instruction. address is the destination address in shared memory. mbarAddress is the address of the mbarrier object. accessSize is the size of the access in bytes. Valid values are 4 and 8.

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackAsyncStore)(void *userdata, uint64_t pc, uint32_t address, uint32_t mbarAddress, void *pNewValue, uint32_t accessSize)

This can be generated by a st.async PTX instruction

userdata is a pointer to user data. See sanitizerPatchModule. pc is the program counter of the patched instruction. address is the destination address in shared memory. mbarAddress is the address of the mbarrier object. pNewValue is a pointer to the new value being written. accessSize is the size of the access in bytes. Valid values are 4 and 8.

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackBarrier)(void *userdata, uint64_t pc, uint32_t barIndex, uint32_t threadCount, uint32_t flags)

userdata is a pointer to user data. See sanitizerPatchModule pc is the program counter of the patched instruction barIndex is the barrier index. threadCount is the number of expected threads (must be a multiple of the warp size). flags contains information about the barrier. See Sanitizer_BarrierFlags to interpret this value. 0 means that all threads are participating in the barrier.

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackBlockEnter)(void *userdata, uint64_t pc)

userdata is a pointer to user data. See sanitizerPatchModule pc is the program counter of the entry point of the block

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackBlockExit)(void *userdata, uint64_t pc)

userdata is a pointer to user data. See sanitizerPatchModule pc is the program counter of the patched instruction

typedef SanitizerPatchResult(SANITIZERAPI* SanitizerCallbackCacheControl)(void *userdata, uint64_t pc, void *address, Sanitizer_CacheControlInstructionKind kind)

userdata is a pointer to user data. See sanitizerPatchModule pc is the program counter of the patched instruction address is the address of the memory being controlled kind is the type of cache control. See Sanitizer_CacheControlInstructionKind

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackCall)(void *userdata, uint64_t pc, uint64_t targetPc, uint32_t flags)

userdata is a pointer to user data. See sanitizerPatchModule pc is the program counter of the patched instruction targetPc is the PC where the called function is located. flags contains information about the function call.

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackClusterBarrierArrive)(void *userdata, uint64_t pc)

Function type for a cluster barrier wait.

This can be generated by a cg::this_cluster().sync() (C++ API), or a barrier.cluster.arrive (PTX API).

userdata is a pointer to user data. See sanitizerPatchModule pc is the program counter of the patched instruction

This can be generated by a cg::this_cluster().sync() (C++ API), or a barrier.cluster.wait (PTX API).

userdata is a pointer to user data. See sanitizerPatchModule pc is the program counter of the patched instruction

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackCudaBarrier)(void *userdata, uint64_t pc, void *barrier, uint32_t kind, uint32_t data)

userdata is a pointer to user data. See sanitizerPatchModule pc is the program counter of the patched instruction barrier Barrier address which can be used as a unique identifier kind Barrier action type. See Sanitizer_CudaBarrierInstructionKind data Barrier data. This is specific to each action type, refer to Sanitizer_CudaBarrierInstructionKind

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackDeviceSideFree)(void *userdata, uint64_t pc, void *ptr)

Note:
This is called prior to the actual call.
userdata is a pointer to user data. See sanitizerPatchModule pc is the program counter of the patched instruction ptr is the pointer passed to device-side free.

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackDeviceSideMalloc)(void *userdata, uint64_t pc, void *allocatedPtr, uint64_t allocatedSize)

Note:
This is called after the call has completed.
userdata is a pointer to user data. See sanitizerPatchModule pc is the program counter of the patched instruction allocatedPtr is the pointer returned by device-side malloc allocatedSize is the size requested by the user to device-side malloc.

typedef SanitizerPatchResult(SANITIZERAPI* SanitizerCallbackMatrixMemoryAccess)(void *userdata, uint64_t pc, uint32_t address, uint32_t accessSize, uint32_t flags, uint32_t count, const void *pNewValue)

userdata is a pointer to user data. See sanitizerPatchModule pc is the program counter of the patched instruction address is the address of the shared memory being read or written. This is an offset within the shared memory window accessSize is the size of the access in bytes. Valid value is 16. flags contains information about the type of access. See Sanitizer_DeviceMemoryFlags to interpret this value. count is the number of matrices accessed. pNewValue is a pointer to the new value being written if the access is a write. If the access is a read or an atomic, the pointer will be NULL.

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackMemcpyAsync)(void *userdata, uint64_t pc, void *src, uint32_t dst, uint32_t accessSize)

userdata is a pointer to user data. See sanitizerPatchModule pc is the program counter of the patched instruction src is the address of the global memory being read. This can be NULL if src-size is 0. dst is the address of the shared memory being written. This is an offset within the shared memory window accessSize is the size of the access in bytes. Valid values are 4, 8 and 16.

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackMemoryAccess)(void *userdata, uint64_t pc, void *ptr, uint32_t accessSize, uint32_t flags, const void *pData)

userdata is a pointer to user data. See sanitizerPatchModule pc is the program counter of the patched instruction ptr is the address of the memory being accessed. For local or shared memory access, this is the offset within the local or shared memory window. accessSize is the size of the access in bytes. Valid values are 1, 2, 4, 8, and 16. flags contains information about the type of access. See Sanitizer_DeviceMemoryFlags to interpret this value. pData is a pointer which value depends on the type of access:

  • If the access is a write, pData points to the new value being written.
  • If the access is a read and pData is not NULL, then it points to a 32-bit mask of loaded bytes being used (padding bytes will not appear).
  • If the access is an atomic, the pointer will be NULL.

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackPipelineCommit)(void *userdata, uint64_t pc)

This can be generated by a pipeline::producer_commit (C++ API), a pipeline_commit (C API) or a cp.async.commit_group (PTX API).

userdata is a pointer to user data. See sanitizerPatchModule pc is the program counter of the patched instruction

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackPipelineWait)(void *userdata, uint64_t pc, uint32_t groups)

This can be generated by a pipeline::consumer_wait (C++ API), a pipeline_wait_prior (C API), cp.async.wait_group or cp.async.wait_all (PTX API).

userdata is a pointer to user data. See sanitizerPatchModule pc is the program counter of the patched instruction groups is the number of groups the pipeline will wait for. 0 is used to wait for all groups.

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackRet)(void *userdata, uint64_t pc)

userdata is a pointer to user data. See sanitizerPatchModule pc is the program counter of the patched instruction

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackSetSmemSize)(void *userdata, uint64_t pc, uint32_t size)

This can be generated by a setsmemsize.sync instruction

userdata is a pointer to user data. See sanitizerPatchModule. pc is the program counter of the patched instruction. size is the requested size in bytes.

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackShfl)(void *userdata, uint64_t pc)

userdata is a pointer to user data. See sanitizerPatchModule pc is the program counter of the patched instruction

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackSyncwarp)(void *userdata, uint64_t pc, uint32_t mask)

userdata is a pointer to user data. See sanitizerPatchModule pc is the program counter of the patched instruction mask is the thread mask passed to __syncwarp().

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackWarpgroupFence)(void *userdata, uint64_t pc, uint32_t warpMask)

This can be generated by a wgmma.fence in PTX.

userdata is a pointer to user data. See sanitizerPatchModule. pc is the program counter of the patched instruction. warpMask is a mask of threads that will perform the fence operation. Expected values are either 0x0 or 0xffffffff (full). The value is expected to be the same across the warpgroup. Other values can be reported but signal a programming error in the target application.

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackWarpgroupMMAAsync)(void *userdata, uint64_t pc, uint32_t addressMatrixA, uint32_t sizeMatrixA, uint32_t addressMatrixB, uint32_t sizeMatrixB, uint32_t flags, uint32_t warpMask)

This can be generated by a wgmma.mma_async in PTX.

userdata is a pointer to user data. See sanitizerPatchModule pc is the program counter of the patched instruction addressMatrixA is the address in shared memory of the matrix A being read. This field is only valid if sizeMatrixA is non-zero and warpMask is full. sizeMatrixA is the size of the matrix A in shared memory. A value of 0 means that the matrix A is read from registers instead. addressMatrixB is the address in shared memory of the matrix B being read. This field is only valid if warpMask is full. sizeMatrixB is the size of the matrix B in shared memory. The value will always be non-zero. flags of type Sanitizer_WarpgroupMMAAsyncFlags provide information about the access. These flags are to be taken into account even if the warpMask is zero. warpMask is a mask of threads that will perform the operation and read the operands. Expected values are either 0x0 or 0xffffffff (full). The value is expected to be the same across the warpgroup. Other values can be reported but signal a programming error in the target application.

typedef SanitizerPatchResult(SANITIZERAPI * SanitizerCallbackWarpgroupWaitGroup)(void *userdata, uint64_t pc, uint32_t numGroups, uint32_t warpMask)

This can be generated by a wgmma.wait_group in PTX.

userdata is a pointer to user data. See sanitizerPatchModule. pc is the program counter of the patched instruction. numGroups is the maximum number of group that will be left pending after the operation. A value of zero means that all MMA async of the warpgroup are guaranteed to have completed after the operation. warpMask is a mask of threads for which the expected values are either 0x0 or 0xffffffff (full). The value is expected to be the same across the warpgroup. Other values can be reported but signal a programming error in the target application. If the value is valid, the value has no influence on the operation.


Enumeration Type Documentation

Flags describing a barrier. These values are to be or-combined in the value of flags for a SanitizerCallbackBarrier callback.

Enumerator:
SANITIZER_BARRIER_FLAG_NONE  Empty flag.
SANITIZER_BARRIER_FLAG_UNALIGNED_ALLOWED  Specifies that the barrier can be called unaligned. This flag is only valid on SM 7.0 and above.

Enumerator:
SANITIZER_CACHE_CONTROL_INVALID  Invalid action ID.
SANITIZER_CACHE_CONTROL_L1_PREFETCH  Prefetch to L1.
SANITIZER_CACHE_CONTROL_L2_PREFETCH  Prefetch to L2.

Flags describing a function call. These values are to be or-combined in the value of flags for a SanitizerCallbackCall callback.

Enumerator:
SANITIZER_CALL_FLAG_NONE  Empty flag.
SANITIZER_CALL_FLAG_UNALIGNED_ALLOWED  Specifies that barriers within this function call can be called unaligned. This flag is only valid on SM 7.0 and above.

Refer to the CUDA Barrier interface section of the CUDA toolkit documentation for a more extensive description of these actions.

Enumerator:
SANITIZER_CUDA_BARRIER_INVALID  Invalid action ID.
SANITIZER_CUDA_BARRIER_INIT  Barrier initialization.
SANITIZER_CUDA_BARRIER_ARRIVE  Barrier arrive operation. On Hopper and newer architectures, barrier data is the count argument to the arrive-on operation.
SANITIZER_CUDA_BARRIER_ARRIVE_DROP  Barrier arrive and drop operation. On Hopper and newer architectures, barrier data is the count argument to the arrive-on operation.
SANITIZER_CUDA_BARRIER_ARRIVE_NOCOMPLETE  Barrier arrive operation without phase completion. Barrier data is the count argument to the arrive-on operation.
SANITIZER_CUDA_BARRIER_ARRIVE_DROP_NOCOMPLETE  Barrier arrive and drop operation without phase completion. Barrier data is the count argument to the arrive-on operation.
SANITIZER_CUDA_BARRIER_WAIT  Barrier wait operation.
SANITIZER_CUDA_BARRIER_INVALIDATE  Barrier invalidation.

Flags describing a memory access. These values are to be or-combined in the value of flags for a SanitizerCallbackMemoryAccess callback.

Enumerator:
SANITIZER_MEMORY_DEVICE_FLAG_NONE  Empty flag.
SANITIZER_MEMORY_DEVICE_FLAG_READ  Specifies that the access is a read.
SANITIZER_MEMORY_DEVICE_FLAG_WRITE  Specifies that the access is a write.
SANITIZER_MEMORY_DEVICE_FLAG_ATOMSYS  Specifies that the access is a system-scoped atomic.
SANITIZER_MEMORY_DEVICE_FLAG_PREFETCH  Specifies that the access is a cache prefetch.

Enumerator:
SANITIZER_FUNCTION_NOT_LOADED  The function is not loaded.
SANITIZER_FUNCTION_PARTIALLY_LOADED  The function is being loaded.
SANITIZER_FUNCTION_LOADED  The function is fully loaded.

Instrumentation. Every entry represent an instruction type or a function call where a callback patch can be inserted.

Enumerator:
SANITIZER_INSTRUCTION_INVALID  Invalid instruction ID.
SANITIZER_INSTRUCTION_BLOCK_ENTER  CUDA block enter. This is called prior to any user code. The type of the callback must be SanitizerCallbackBlockEnter.
SANITIZER_INSTRUCTION_BLOCK_EXIT  CUDA block exit. This is called after all user code has executed. The type of the callback must be SanitizerCallbackBlockExit.
SANITIZER_INSTRUCTION_GLOBAL_MEMORY_ACCESS  Global Memory Access. This can be a store, load or atomic operation. The type of the callback must be SanitizerCallbackMemoryAccess.
SANITIZER_INSTRUCTION_SHARED_MEMORY_ACCESS  Shared Memory Access. This can be a store, load or atomic operation. The type of the callback must be SanitizerCallbackMemoryAccess.
SANITIZER_INSTRUCTION_LOCAL_MEMORY_ACCESS  Local Memory Access. This can be a store or load operation. The type of the callback must be SanitizerCallbackMemoryAccess.
SANITIZER_INSTRUCTION_BARRIER  Barrier. The type of the callback must be SanitizerCallbackBarrier.
SANITIZER_INSTRUCTION_SYNCWARP  Syncwarp. The type of the callback must be SanitizerCallbackSyncwarp.
SANITIZER_INSTRUCTION_SHFL  Shfl. The type of the callback must be SanitizerCallbackShfl.
SANITIZER_INSTRUCTION_CALL  Function call. The type of the callback must be SanitizerCallbackCall.
SANITIZER_INSTRUCTION_RET  Function return. The type of the callback must be SanitizerCallbackRet.
SANITIZER_INSTRUCTION_DEVICE_SIDE_MALLOC  Device-side malloc. The type of the callback must be SanitizerCallbackDeviceSideMalloc.
SANITIZER_INSTRUCTION_DEVICE_SIDE_FREE  Device-side free. The type of the callback must be SanitizerCallbackDeviceSideFree.
SANITIZER_INSTRUCTION_CUDA_BARRIER  CUDA Barrier operation. The type of the callback must be SanitizerCallbackCudaBarrier.
SANITIZER_INSTRUCTION_MEMCPY_ASYNC  Global to shared memory asynchronous copy. The type of the callback must be SanitizerCallbackMemcpyAsync.
SANITIZER_INSTRUCTION_PIPELINE_COMMIT  Pipeline commit. The type of the callback must be SanitizerCallbackPipelineCommit.
SANITIZER_INSTRUCTION_PIPELINE_WAIT  Pipeline wait. The type of the callback must be SanitizerCallbackPipelineWait.
SANITIZER_INSTRUCTION_REMOTE_SHARED_MEMORY_ACCESS  Remote Shared Memory Access. This can be a store or load operation. The type of the callback must be SanitizerCallbackMemoryAccess.
SANITIZER_INSTRUCTION_DEVICE_ALIGNED_MALLOC  Device-side aligned malloc. The type of the callback must be SanitizerCallbackDeviceSideMalloc.
SANITIZER_INSTRUCTION_MATRIX_MEMORY_ACCESS  Matrix shared memory access. The type of the callback must be SanitizerCallbackMatrixMemoryAccess.
SANITIZER_INSTRUCTION_CACHE_CONTROL  Cache control instruction. The type of the callback must be SanitizerCallbackCacheControl.
SANITIZER_INSTRUCTION_CLUSTER_BARRIER_ARRIVE  Cluster barrier arrive instruction. The type of the callback must be SanitizerCallbackClusterBarrierArrive.
SANITIZER_INSTRUCTION_CLUSTER_BARRIER_WAIT  Cluster barrier wait instruction. The type of the callback must be SanitizerCallbackClusterBarrierWait.
SANITIZER_INSTRUCTION_WARPGROUP_MMA_ASYNC  Warpgroup aligned async MMA instruction. The type of the callback must be SanitizerCallbackWarpgroupMMAAsync.
SANITIZER_INSTRUCTION_WARPGROUP_WAIT_GROUP  Warpgroup wait MMA group instruction. The type of the callback must be SanitizerCallbackWarpgroupWaitGroup.
SANITIZER_INSTRUCTION_WARPGROUP_FENCE  Warpgroup fence instruction. The type of the callback must be SanitizerCallbackWarpgroupFence.
SANITIZER_INSTRUCTION_ASYNC_STORE  Asynchronous store instruction. The type of the callback must be SanitizerCallbackAsyncStore.
SANITIZER_INSTRUCTION_ASYNC_REDUCTION  Asynchronous reduction instruction. The type of the callback must be SanitizerCallbackAsyncReduction.
SANITIZER_INSTRUCTION_SET_SHARED_MEMORY_SIZE  Set the shared memory size allocated to a block instruction. The type of the callback must SanitizerCallbackSetSmemSize
SANITIZER_INSTRUCTION_BARRIER_RELEASE  Barrier after it is released. The type of the callback must be SanitizerCallbackBarrier.

Flags describing a warpgroup aligned MMA async. These values are to be or-combined in the value of flags for a SanitizerCallbackWarpgroupMMAAsync callback.

Enumerator:
SANITIZER_WARPGROUP_MMA_ASYNC_FLAG_NONE  Empty flag.
SANITIZER_WARPGROUP_MMA_ASYNC_FLAG_COMMIT_GROUP  Specifies that the MMA async delimits a MMA async group of which it is the last instruction. Please refer to the PTX documentation for wgmma_async.commit_group for more details. This property is valid even if the warpMask is zero.

Error and result codes returned by Sanitizer patches. If a patch returns SANITIZER_PATCH_ERROR, the thread will be exited. On Volta and newer architectures, the full warp which the thread belongs to will be exited.

Enumerator:
SANITIZER_PATCH_SUCCESS  No error.
SANITIZER_PATCH_ERROR  An error was detected in the patch.


Function Documentation

SanitizerResult SANITIZERAPI sanitizerAddPatches ( const void *  image,
CUcontext  ctx 
)

Note:
Thread-safety: an API user must serialize access to sanitizerAddPatchesFromFile, sanitizerAddPatches, sanitizerPatchInstructions, and sanitizerPatchModule. For example if sanitizerAddPatches(image) and sanitizerPatchInstruction(*, *, cbName) are called concurrently and cbName is intended to be found in the loaded image, the results are undefined.

The patches loaded are only valid for the specified CUDA context.

Parameters:
image Pointer to module data to load. This API supports the same module formats as the cuModuleLoadData and cuModuleLoadFatBinary functions from the CUDA driver API.
ctx CUDA context in which to load the patches. If ctx is NULL, the current context will be used.
Return values:
SANITIZER_SUCCESS on success
SANITIZER_ERROR_NOT_INITIALIZED if unable to initialize the sanitizer
SANITIZER_ERROR_INVALID_PARAMETER if image does not point to a valid CUDA module.

SanitizerResult SANITIZERAPI sanitizerAddPatchesFromFile ( const char *  filename,
CUcontext  ctx 
)

Note:
Thread-safety: an API user must serialize access to sanitizerAddPatchesFromFile, sanitizerAddPatches, sanitizerPatchInstructions, and sanitizerPatchModule. For example if sanitizerAddPatchesFromFile(filename) and sanitizerPatchInstruction(*, *, cbName) are called concurrently and cbName is intended to be found in the loaded module, the results are undefined.

The patches loaded are only valid for the specified CUDA context.

Parameters:
filename Path to the module file. This API supports the same module formats as the cuModuleLoad function from the CUDA driver API.
ctx CUDA context in which to load the patches. If ctx is NULL, the current context will be used.
Return values:
SANITIZER_SUCCESS on success
SANITIZER_ERROR_NOT_INITIALIZED if unable to initialize the sanitizer
SANITIZER_ERROR_INVALID_PARAMETER if filename is not a path to a valid CUDA module.

SanitizerResult SANITIZERAPI sanitizerGetCallbackPcAndSize ( CUcontext  ctx,
const char *  deviceCallbackName,
uint64_t *  pc,
uint64_t *  size 
)

Parameters:
[in] ctx CUDA context in which the patches were loaded. If ctx is NULL, the current context will be used.
[in] deviceCallbackName device function callback name
[out] pc Callback PC returned
[out] size Callback size returned
Return values:
SANITIZER_SUCCESS on success
SANITIZER_ERROR_INVALID_PARAMETER if deviceCallbackName function cannot be located, if pc is NULL or if size is NULL.

SanitizerResult SANITIZERAPI sanitizerGetFunctionLoadedStatus ( CUfunction  func,
Sanitizer_FunctionLoadedStatus loadingStatus 
)

Parameters:
[in] func CUDA function for which the loading status is queried.
[out] loadingStatus Loading status returned
Return values:
SANITIZER_SUCCESS on success
SANITIZER_ERROR_INVALID_PARAMETER if func is NULL or if loadingStatus is NULL.
SANITIZER_ERROR_NOT_SUPPORTED if the loading status cannot be queried with this driver version.

SanitizerResult SANITIZERAPI sanitizerGetFunctionPcAndSize ( CUmodule  module,
const char *  functionName,
uint64_t *  pc,
uint64_t *  size 
)

Parameters:
[in] module CUDA module containing the function
[in] deviceCallbackName CUDA function name
[out] pc Function start program counter (PC) returned
[out] size Function size in bytes returned
Return values:
SANITIZER_SUCCESS on success
SANITIZER_ERROR_INVALID_PARAMETER if functionName function cannot be located, if pc is NULL or if size is NULL.

SanitizerResult SANITIZERAPI sanitizerPatchInstructions ( const Sanitizer_InstructionId  instructionId,
CUmodule  module,
const char *  deviceCallbackName 
)

Mark that all instrumentation points matching instructionId are to be patched in order to call the device function identified by deviceCallbackName. It is up to the API client to ensure that this device callback exists and match the correct callback format for this instrumentation point.

Note:
Thread-safety: an API user must serialize access to sanitizerAddPatchesFromFile, sanitizerAddPatches, sanitizerPatchInstructions, and sanitizerPatchModule. For example if sanitizerAddPatches(fileName) and sanitizerPatchInstruction(*, *, cbName) are called concurrently and cbName is intended to be found in the loaded module, the results are undefined.
Parameters:
instructionId Instrumentation point for which to insert patches
module CUDA module to instrument
deviceCallbackName Name of the device function callback that the inserted patch will call at the instrumented points. This function is expected to be found in code previously loaded by sanitizerAddPatchesFromFile or sanitizerAddPatches.
Return values:
SANITIZER_SUCCESS on success
SANITIZER_ERROR_NOT_INITIALIZED if unable to initialize the sanitizer
SANITIZER_ERROR_INVALID_PARAMETER if module is not a CUDA module or if deviceCallbackName function cannot be located.

SanitizerResult SANITIZERAPI sanitizerPatchModule ( CUmodule  module  ) 

Perform the instrumentation of a CUDA module based on previous calls to sanitizerPatchInstructions. This function also specifies the device memory buffer to be passed in as userdata to all callback functions.

Note:
Thread-safety: an API user must serialize access to sanitizerAddPatchesFromFile, sanitizerAddPatches, sanitizerPatchInstructions, and sanitizerPatchModule. For example if sanitizerPatchModule(mod, *) and sanitizerPatchInstruction(*, mod, *) are called concurrently, the results are undefined.
Parameters:
module CUDA module to instrument
Return values:
SANITIZER_SUCCESS on success
SANITIZER_ERROR_INVALID_PARAMETER if module is not a CUDA module

SanitizerResult SANITIZERAPI sanitizerSetCallbackData ( CUfunction  kernel,
const void *  userdata 
)

Mark all subsequent launches of kernel to use userdata pointer as the device memory buffer to pass in to callback functions.

Parameters:
kernel CUDA function to link to user data. Callbacks in subsequent launches on this kernel will use userdata as callback data.
userdata Device memory buffer. This data will be passed to callback functions via the userdata parameter.
Return values:
SANITIZER_SUCCESS on success

SanitizerResult SANITIZERAPI sanitizerSetDeviceGraphData ( CUgraphExec  graphExec,
Sanitizer_StreamHandle  stream,
const void *  userdata 
)

Mark all subsequent launch of graphExec to make available userdata in device callbacks from device-launched graphs. userdata will not be set in the callback userdata parameter but must be accessed through another mean instead. Please refer to the Sanitizer API reference manual. This function is only available if the driver version is 535 or newer.

Parameters:
graphExec CUDA graphExec that will launch CUDA graphs from the device.
stream CUDA stream associated with the stream launch.
userdata Device memory buffer.
Return values:
SANITIZER_SUCCESS on success

SanitizerResult SANITIZERAPI sanitizerSetLaunchCallbackData ( Sanitizer_LaunchHandle  launch,
CUfunction  kernel,
Sanitizer_StreamHandle  stream,
const void *  userdata 
)

Mark launch to use userdata pointer as the device memory buffer to pass in to callback functions. This function is only available if the driver version is 455 or newer.

Parameters:
launch Kernel launch to link to user data. Callbacks in this kernel launch will use userdata as callback data.
kernel CUDA function associated with the kernel launch.
stream CUDA stream associated with the stream launch.
userdata Device memory buffer. This data will be passed to callback functions via the userdata parameter.
Return values:
SANITIZER_SUCCESS on success

SanitizerResult SANITIZERAPI sanitizerUnpatchModule ( CUmodule  module  ) 

Remove any instrumentation of a CUDA module performed by previous calls to sanitizerPatchModule.

Note:
Thread-safety: an API user must serialize access to sanitizerPatchModule and sanitizerUnpatchModule on the same module. For example, if sanitizerPatchModule(mod) and sanitizerUnpatchModule(mod) are called concurrently, the results are undefined.
Parameters:
module CUDA module on which to remove instrumentation
Return values:
SANITIZER_SUCCESS on success


Generated on Tue Oct 31 02:32:23 2023 for SanitizerApi by  doxygen 1.5.8