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. |
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) |
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) |
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:
pData
points to the new value being written.pData
is not NULL
, then it points to a 32-bit mask of loaded bytes being used (padding bytes will not appear).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.
Flags describing a barrier. These values are to be or-combined in the value of flags for a SanitizerCallbackBarrier callback.
enum Sanitizer_CallFlags |
Flags describing a function call. These values are to be or-combined in the value of flags for a SanitizerCallbackCall callback.
Refer to the CUDA Barrier interface section of the CUDA toolkit documentation for a more extensive description of these actions.
Flags describing a memory access. These values are to be or-combined in the value of flags for a SanitizerCallbackMemoryAccess callback.
Instrumentation. Every entry represent an instruction type or a function call where a callback patch can be inserted.
Flags describing a warpgroup aligned MMA async. These values are to be or-combined in the value of flags for a SanitizerCallbackWarpgroupMMAAsync callback.
enum SanitizerPatchResult |
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.
SanitizerResult SANITIZERAPI sanitizerAddPatches | ( | const void * | image, | |
CUcontext | ctx | |||
) |
The patches loaded are only valid for the specified CUDA context.
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. |
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 | |||
) |
The patches loaded are only valid for the specified CUDA context.
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. |
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 | |||
) |
[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 |
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 | |||
) |
[in] | func | CUDA function for which the loading status is queried. |
[out] | loadingStatus | Loading status returned |
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 | |||
) |
[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 |
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.
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. |
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.
module | CUDA module to instrument |
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.
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. |
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.
graphExec | CUDA graphExec that will launch CUDA graphs from the device. | |
stream | CUDA stream associated with the stream launch. | |
userdata | Device memory buffer. |
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.
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. |
SANITIZER_SUCCESS | on success |
SanitizerResult SANITIZERAPI sanitizerUnpatchModule | ( | CUmodule | module | ) |
Remove any instrumentation of a CUDA module performed by previous calls to sanitizerPatchModule.
module | CUDA module on which to remove instrumentation |
SANITIZER_SUCCESS | on success |