Sanitizer Patching API

实现 Sanitizer Patching API 的函数、类型和枚举。

枚举

SanitizerPatchResult: Sanitizer 补丁结果代码。
Sanitizer_BarrierFlags: 描述屏障的标志。
Sanitizer_CacheControlInstructionKind: 缓存控制操作。
Sanitizer_CallFlags: 描述函数调用的标志。
Sanitizer_CudaBarrierInstructionKind: CUDA 屏障操作类型。
Sanitizer_DeviceMemoryFlags: 描述内存访问的标志。
Sanitizer_FunctionLoadedStatus
Sanitizer_InstructionId: 检测工具。
Sanitizer_WarpgroupMMAAsyncFlags: 描述 warpgroup 对齐的 MMA 异步操作的标志。

函数

SanitizerResult sanitizerAddPatches(const void *image, CUcontext ctx): 加载包含可供补丁 API 使用的补丁的模块。
SanitizerResult sanitizerAddPatchesFromFile(const char *filename, CUcontext ctx): 加载包含可供补丁 API 使用的补丁的模块。
SanitizerResult sanitizerGetCallbackPcAndSize(CUcontext ctx, const char *deviceCallbackName, uint64_t *pc, uint64_t *size): 获取设备回调的 PC 和大小。
SanitizerResult sanitizerGetFunctionLoadedStatus(CUfunction func, Sanitizer_FunctionLoadedStatus *loadingStatus): 获取函数的加载状态。
SanitizerResult sanitizerGetFunctionPcAndSize(CUmodule module, const char *functionName, uint64_t *pc, uint64_t *size): 获取 CUDA 函数的 PC 和大小。
SanitizerResult sanitizerPatchInstructions(const Sanitizer_InstructionId instructionId, CUmodule module, const char *deviceCallbackName): 设置模块中要应用的检测点和补丁。
SanitizerResult sanitizerPatchModule(CUmodule module): 执行模块的实际检测。
SanitizerResult sanitizerSetCallbackData(CUfunction kernel, const void *userdata): 指定回调的用户数据指针。
SanitizerResult sanitizerSetDeviceGraphData(CUgraphExec graphExec, Sanitizer_StreamHandle stream, const void *userdata): 指定从设备启动的图中回调可访问的用户数据指针，这些图由指定的主机启动的 graphExec 启动。
SanitizerResult sanitizerSetLaunchCallbackData(Sanitizer_LaunchHandle launch, CUfunction kernel, Sanitizer_StreamHandle stream, const void *userdata): 指定回调的用户数据指针。
SanitizerResult sanitizerUnpatchModule(CUmodule module): 移除模块的现有检测。

类型定义

SanitizerCallbackAsyncReduction: 共享内存上异步归约操作的函数类型。
SanitizerCallbackAsyncStore: 共享内存上异步存储操作的函数类型。
SanitizerCallbackBarrier: 屏障回调的函数类型。
SanitizerCallbackBlockEnter: CUDA 块进入回调的函数类型。
SanitizerCallbackBlockExit: CUDA 块退出回调的函数类型。
SanitizerCallbackBulkCopyGlobalToShared: 从全局内存到共享内存的异步批量复制的函数类型。
SanitizerCallbackCacheControl: 缓存控制指令回调的函数类型。
SanitizerCallbackCall: 函数调用回调的函数类型。
SanitizerCallbackClusterBarrierArrive: 集群屏障到达的函数类型。
SanitizerCallbackCudaBarrier: CUDA 屏障操作回调的函数类型。
SanitizerCallbackDeviceSideFree: 设备端 free 调用的函数类型。
SanitizerCallbackDeviceSideMalloc: 设备端 malloc 调用的函数类型。
SanitizerCallbackMatrixMemoryAccess: 矩阵共享内存访问回调的函数类型。
SanitizerCallbackMemcpyAsync: 从全局内存到共享内存的异步复制的函数类型。
SanitizerCallbackMemcpyAsyncBarrier: 用于 mbarrier 完成的 cuda 屏障的函数类型。
SanitizerCallbackMemoryAccess: 内存访问回调的函数类型。
SanitizerCallbackMemsetShared: 在共享内存上进行 memset 操作的函数类型。
SanitizerCallbackPipelineCommit: 流水线提交的函数类型。
SanitizerCallbackPipelineWait: 流水线等待的函数类型。
SanitizerCallbackRet: 函数返回回调的函数类型。
SanitizerCallbackSetSmemSize: 设置分配给块的共享内存大小的函数类型。
SanitizerCallbackShfl: shfl 回调的函数类型。
SanitizerCallbackSyncwarp: syncwarp 回调的函数类型。
SanitizerCallbackTensorCoreBarrier: Blackwell 张量核心屏障的函数类型。
SanitizerCallbackWarpgroupFence: warpgroup MMA fence 的函数类型。
SanitizerCallbackWarpgroupMMAAsync: warpgroup 对齐的异步 MMA 的函数类型。
SanitizerCallbackWarpgroupWaitGroup: warpgroup MMA 等待组的函数类型。
Sanitizer_LaunchHandle

枚举

enum SanitizerPatchResult

Sanitizer 补丁结果代码。

Sanitizer 补丁返回的错误和结果代码。如果补丁返回 SANITIZER_PATCH_ERROR，线程将退出。在 Volta 和更新的架构上，线程所属的完整 warp 将退出。

值

enumerator SANITIZER_PATCH_SUCCESS: 没有错误。

enumerator SANITIZER_PATCH_ERROR: 在补丁中检测到错误。

enumerator SANITIZER_PATCH_FORCE_INT

enum Sanitizer_BarrierFlags

描述屏障的标志。

描述屏障的标志。这些值将在 SanitizerCallbackBarrier 回调的 flags 值中进行或运算组合。

值

enumerator SANITIZER_BARRIER_FLAG_NONE: 空标志。

enumerator SANITIZER_BARRIER_FLAG_UNALIGNED_ALLOWED

指定可以非对齐方式调用屏障。

此标志仅在 SM 7.0 及以上版本上有效。

enumerator SANITIZER_BARRIER_FLAG_FORCE_INT

enum Sanitizer_CacheControlInstructionKind

缓存控制操作。

值

enumerator SANITIZER_CACHE_CONTROL_INVALID: 无效的操作 ID。

enumerator SANITIZER_CACHE_CONTROL_L1_PREFETCH: 预取到 L1 缓存。

enumerator SANITIZER_CACHE_CONTROL_L2_PREFETCH: 预取到 L2 缓存。

enumerator SANITIZER_CACHE_CONTROL_FORCE_INT

enum Sanitizer_CallFlags

描述函数调用的标志。

描述函数调用的标志。这些值将在 SanitizerCallbackCall 回调的 flags 值中进行或运算组合。

值

enumerator SANITIZER_CALL_FLAG_NONE: 空标志。

enumerator SANITIZER_CALL_FLAG_UNALIGNED_ALLOWED

指定此函数调用中的屏障可以非对齐方式调用。

此标志仅在 SM 7.0 及以上版本上有效。

enumerator SANITIZER_CALL_FLAG_FORCE_INT

enum Sanitizer_CudaBarrierInstructionKind

CUDA 屏障操作类型。

有关这些操作的更详细描述，请参阅 CUDA 工具包文档的 CUDA 屏障接口部分。

值

enumerator SANITIZER_CUDA_BARRIER_INVALID: 无效的操作 ID。

enumerator SANITIZER_CUDA_BARRIER_INIT: 屏障初始化。

enumerator SANITIZER_CUDA_BARRIER_ARRIVE

屏障到达操作。

在 Hopper 和更新的架构上，屏障数据是到达操作的 count 参数。

enumerator SANITIZER_CUDA_BARRIER_ARRIVE_DROP

屏障到达和丢弃操作。

在 Hopper 和更新的架构上，屏障数据是到达操作的 count 参数。

enumerator SANITIZER_CUDA_BARRIER_ARRIVE_NOCOMPLETE

不完成阶段的屏障到达操作。

屏障数据是到达操作的 count 参数。

enumerator SANITIZER_CUDA_BARRIER_ARRIVE_DROP_NOCOMPLETE

不完成阶段的屏障到达和丢弃操作。

屏障数据是到达操作的 count 参数。

enumerator SANITIZER_CUDA_BARRIER_WAIT: 屏障等待操作。

enumerator SANITIZER_CUDA_BARRIER_INVALIDATE: 屏障失效。

enumerator SANITIZER_CUDA_BARRIER_FORCE_INT

enum Sanitizer_DeviceMemoryFlags

描述内存访问的标志。

描述内存访问的标志。这些值将在 SanitizerCallbackMemoryAccess 回调的 flags 值中进行或运算组合。

值

enumerator SANITIZER_MEMORY_DEVICE_FLAG_NONE: 空标志。

enumerator SANITIZER_MEMORY_DEVICE_FLAG_READ: 指定访问是读取操作。

enumerator SANITIZER_MEMORY_DEVICE_FLAG_WRITE: 指定访问是写入操作。

enumerator SANITIZER_MEMORY_DEVICE_FLAG_ATOMSYS: 指定访问是系统范围的原子操作。

enumerator SANITIZER_MEMORY_DEVICE_FLAG_PREFETCH: 指定访问是缓存预取操作。

enumerator SANITIZER_MEMORY_DEVICE_FLAG_FORCE_INT

enum Sanitizer_FunctionLoadedStatus

值

enumerator SANITIZER_FUNCTION_NOT_LOADED: 函数未加载。

enumerator SANITIZER_FUNCTION_PARTIALLY_LOADED: 函数正在加载。

enumerator SANITIZER_FUNCTION_LOADED: 函数已完全加载。

enumerator SANITIZER_FUNCTION_LOADED_FORCE_INT

enum Sanitizer_InstructionId

检测工具。

检测工具。每个条目代表可以插入回调补丁的指令类型或函数调用。

值

enumerator SANITIZER_INSTRUCTION_INVALID: 无效的指令 ID。

enumerator SANITIZER_INSTRUCTION_BLOCK_ENTER

CUDA 块进入。

这在任何用户代码之前调用。回调的类型必须是 SanitizerCallbackBlockEnter。

enumerator SANITIZER_INSTRUCTION_BLOCK_EXIT

CUDA 块退出。

这在所有用户代码执行完毕后调用。回调的类型必须是 SanitizerCallbackBlockExit。

enumerator SANITIZER_INSTRUCTION_GLOBAL_MEMORY_ACCESS

全局内存访问。

这可以是存储、加载或原子操作。回调的类型必须是 SanitizerCallbackMemoryAccess。

enumerator SANITIZER_INSTRUCTION_SHARED_MEMORY_ACCESS

共享内存访问。

这可以是存储、加载或原子操作。回调的类型必须是 SanitizerCallbackMemoryAccess。

enumerator SANITIZER_INSTRUCTION_LOCAL_MEMORY_ACCESS

本地内存访问。

这可以是存储或加载操作。回调的类型必须是 SanitizerCallbackMemoryAccess。

enumerator SANITIZER_INSTRUCTION_BARRIER

屏障。

回调的类型必须是 SanitizerCallbackBarrier。

enumerator SANITIZER_INSTRUCTION_SYNCWARP

Syncwarp。

回调的类型必须是 SanitizerCallbackSyncwarp。

enumerator SANITIZER_INSTRUCTION_SHFL

Shfl。

回调的类型必须是 SanitizerCallbackShfl。

enumerator SANITIZER_INSTRUCTION_CALL

函数调用。

回调的类型必须是 SanitizerCallbackCall。

enumerator SANITIZER_INSTRUCTION_RET

函数返回。

回调的类型必须是 SanitizerCallbackRet。

enumerator SANITIZER_INSTRUCTION_DEVICE_SIDE_MALLOC

设备端 malloc。

回调的类型必须是 SanitizerCallbackDeviceSideMalloc。

enumerator SANITIZER_INSTRUCTION_DEVICE_SIDE_FREE

设备端 free。

回调的类型必须是 SanitizerCallbackDeviceSideFree。

enumerator SANITIZER_INSTRUCTION_CUDA_BARRIER

CUDA 屏障操作。

回调的类型必须是 SanitizerCallbackCudaBarrier。

enumerator SANITIZER_INSTRUCTION_MEMCPY_ASYNC

全局内存到共享内存的异步复制。

回调的类型必须是 SanitizerCallbackMemcpyAsync。

enumerator SANITIZER_INSTRUCTION_PIPELINE_COMMIT

流水线提交。

回调的类型必须是 SanitizerCallbackPipelineCommit。

enumerator SANITIZER_INSTRUCTION_PIPELINE_WAIT

流水线等待。

回调的类型必须是 SanitizerCallbackPipelineWait。

enumerator SANITIZER_INSTRUCTION_REMOTE_SHARED_MEMORY_ACCESS

远程共享内存访问。

这可以是存储或加载操作。回调的类型必须是 SanitizerCallbackMemoryAccess。

enumerator SANITIZER_INSTRUCTION_DEVICE_ALIGNED_MALLOC

设备端对齐的 malloc。

回调的类型必须是 SanitizerCallbackDeviceSideMalloc。

enumerator SANITIZER_INSTRUCTION_MATRIX_MEMORY_ACCESS

矩阵共享内存访问。

回调的类型必须是 SanitizerCallbackMatrixMemoryAccess。

enumerator SANITIZER_INSTRUCTION_CACHE_CONTROL

缓存控制指令。

回调的类型必须是 SanitizerCallbackCacheControl。

enumerator SANITIZER_INSTRUCTION_CLUSTER_BARRIER_ARRIVE

集群屏障到达指令。

回调的类型必须是 SanitizerCallbackClusterBarrierArrive。

enumerator SANITIZER_INSTRUCTION_CLUSTER_BARRIER_WAIT

集群屏障等待指令。

回调的类型必须是 SanitizerCallbackClusterBarrierWait。

enumerator SANITIZER_INSTRUCTION_WARPGROUP_MMA_ASYNC

Warpgroup 对齐的异步 MMA 指令。

回调的类型必须是 SanitizerCallbackWarpgroupMMAAsync。

enumerator SANITIZER_INSTRUCTION_WARPGROUP_WAIT_GROUP

Warpgroup 等待 MMA 组指令。

回调的类型必须是 SanitizerCallbackWarpgroupWaitGroup。

enumerator SANITIZER_INSTRUCTION_WARPGROUP_FENCE

Warpgroup 栅栏指令。

回调的类型必须是 SanitizerCallbackWarpgroupFence。

enumerator SANITIZER_INSTRUCTION_ASYNC_STORE

异步存储指令。

回调的类型必须是 SanitizerCallbackAsyncStore。

enumerator SANITIZER_INSTRUCTION_ASYNC_REDUCTION

异步归约指令。

回调的类型必须是 SanitizerCallbackAsyncReduction。

enumerator SANITIZER_INSTRUCTION_SET_SHARED_MEMORY_SIZE

设置分配给块指令的共享内存大小。

回调的类型必须是 SanitizerCallbackSetSmemSize。

enumerator SANITIZER_INSTRUCTION_BARRIER_RELEASE

屏障释放后。

回调的类型必须是 SanitizerCallbackBarrier。

enumerator SANITIZER_INSTRUCTION_BULK_COPY_GLOBAL_TO_SHARED

从全局内存到共享内存的批量复制指令。

回调的类型必须是 SanitizerCallbackBulkCopyGlobalToShared。

enumerator SANITIZER_INSTRUCTION_TENSOR_CORE_BARRIER

张量核心屏障。

回调的类型必须是 SanitizerCallbackTensorCoreBarrier。

enumerator SANITIZER_INSTRUCTION_MEMSET_SHARED

从全局内存到共享内存的批量复制指令。

回调的类型必须是 SanitizerCallbackMemsetShared。

enumerator SANITIZER_INSTRUCTION_SYNCWARP_RELEASE

Syncwarp 释放后。

回调的类型必须是 SanitizerCallbackSyncwarp。

enumerator SANITIZER_INSTRUCTION_MEMCPY_ASYNC_BARRIER

使用 Cuda 屏障进行 memcpy 异步完成。

回调的类型必须是 SanitizerCallbackMemcpyAsyncBarrier。

enumerator SANITIZER_INSTRUCTION_FORCE_INT

enum Sanitizer_WarpgroupMMAAsyncFlags

描述 warpgroup 对齐的 MMA 异步操作的标志。

描述 warpgroup 对齐的异步 MMA 的标志。这些值将在 SanitizerCallbackWarpgroupMMAAsync 回调的 flags 值中进行或运算组合。

值

enumerator SANITIZER_WARPGROUP_MMA_ASYNC_FLAG_NONE: 空标志。

enumerator SANITIZER_WARPGROUP_MMA_ASYNC_FLAG_COMMIT_GROUP

指定 MMA 异步指令划定了一个 MMA 异步组的界限，并且它是该组中的最后一条指令。

有关更多详细信息，请参阅 PTX 文档中关于 wgmma_async.commit_group 的说明。即使 warpMask 为零，此属性也有效。

enumerator SANITIZER_WARPGROUP_MMA_ASYNC_FLAG_FORCE_INT

函数

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

加载包含可供补丁 API 使用的补丁的模块。

注意

线程安全性：API 用户必须序列化对 sanitizerAddPatchesFromFile、sanitizerAddPatches、sanitizerPatchInstructions 和 sanitizerPatchModule 的访问。例如，如果并发调用 sanitizerAddPatches(image) 和 sanitizerPatchInstruction(*, *, cbName)，并且 cbName 旨在在加载的图像中找到，则结果是未定义的。

注意

加载的补丁仅对指定的 CUDA 上下文有效。

参数

image – 指向要加载的模块数据的指针。此 API 支持与 CUDA 驱动程序 API 中的 cuModuleLoadData 和 cuModuleLoadFatBinary 函数相同的模块格式。
ctx – 加载补丁的 CUDA 上下文。如果 ctx 为 NULL，将使用当前上下文。

返回值

SANITIZER_SUCCESS – 成功时
SANITIZER_ERROR_NOT_INITIALIZED – 如果无法初始化清理器
SANITIZER_ERROR_INVALID_PARAMETER – 如果 image 未指向有效的 CUDA 模块。

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

加载包含可供补丁 API 使用的补丁的模块。

注意

线程安全性：API 用户必须序列化对 sanitizerAddPatchesFromFile、sanitizerAddPatches、sanitizerPatchInstructions 和 sanitizerPatchModule 的访问。例如，如果并发调用 sanitizerAddPatchesFromFile(filename) 和 sanitizerPatchInstruction(*, *, cbName)，并且 cbName 旨在在加载的模块中找到，则结果是未定义的。

注意

加载的补丁仅对指定的 CUDA 上下文有效。

参数

filename – 模块文件的路径。此 API 支持与 CUDA 驱动程序 API 中的 cuModuleLoad 函数相同的模块格式。
ctx – 加载补丁的 CUDA 上下文。如果 ctx 为 NULL，将使用当前上下文。

返回值

SANITIZER_SUCCESS – 成功时
SANITIZER_ERROR_NOT_INITIALIZED – 如果无法初始化清理器
SANITIZER_ERROR_INVALID_PARAMETER – 如果 filename 不是有效 CUDA 模块的路径。

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

获取设备回调的 PC 和大小。

参数

ctx – [in] 加载补丁的 CUDA 上下文。如果 ctx 为 NULL，将使用当前上下文。
deviceCallbackName – [in] 设备函数回调名称
pc – [out] 返回的回调 PC
size – [out] 返回的回调大小

返回值

SANITIZER_SUCCESS – 成功时
SANITIZER_ERROR_INVALID_PARAMETER – 如果找不到 deviceCallbackName 函数，或者 pc 为 NULL 或 size 为 NULL。

SanitizerResult sanitizerGetFunctionLoadedStatus(CUfunction func, Sanitizer_FunctionLoadedStatus *loadingStatus)

获取函数的加载状态。

需要驱动程序版本 >= 515。

参数

func – [in] 查询加载状态的 CUDA 函数。
loadingStatus – [out] 返回的加载状态

返回值

SANITIZER_SUCCESS – 成功时
SANITIZER_ERROR_INVALID_PARAMETER – 如果 func 为 NULL 或 loadingStatus 为 NULL。
SANITIZER_ERROR_NOT_SUPPORTED – 如果此驱动程序版本无法查询加载状态。

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

获取 CUDA 函数的 PC 和大小。

参数

module – [in] 包含该函数的 CUDA 模块
deviceCallbackName – [in] CUDA 函数名称
pc – [out] 返回的函数起始程序计数器 (PC)
size – [out] 返回的函数大小（以字节为单位）

返回值

SANITIZER_SUCCESS – 成功时
SANITIZER_ERROR_INVALID_PARAMETER – 如果找不到 functionName 函数，或者 pc 为 NULL 或 size 为 NULL。

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

设置模块中要应用的检测点和补丁。

标记所有与 instructionId 匹配的检测点将被修补，以便调用由 deviceCallbackName 标识的设备函数。API 客户端有责任确保此设备回调存在，并为此检测点匹配正确的回调格式。

注意

线程安全性：API 用户必须序列化对 sanitizerAddPatchesFromFile、sanitizerAddPatches、sanitizerPatchInstructions 和 sanitizerPatchModule 的访问。例如，如果并发调用 sanitizerAddPatches(fileName) 和 sanitizerPatchInstruction(*, *, cbName)，并且 cbName 旨在在加载的模块中找到，则结果是未定义的。

参数

instructionId – 要插入补丁的检测点
module – 要检测的 CUDA 模块
deviceCallbackName – 插入的补丁将在检测点调用的设备函数回调的名称。预计此函数在先前由 sanitizerAddPatchesFromFile 或 sanitizerAddPatches 加载的代码中找到。

返回值

SANITIZER_SUCCESS – 成功时
SANITIZER_ERROR_NOT_INITIALIZED – 如果无法初始化清理器
SANITIZER_ERROR_INVALID_PARAMETER – 如果 module 不是 CUDA 模块，或者找不到 deviceCallbackName 函数。

SanitizerResult sanitizerPatchModule(CUmodule module)

执行模块的实际检测。

基于先前对 sanitizerPatchInstructions 的调用，执行 CUDA 模块的检测。此函数还指定要作为 userdata 传递到所有回调函数的设备内存缓冲区。

注意

线程安全性：API 用户必须序列化对 sanitizerAddPatchesFromFile、sanitizerAddPatches、sanitizerPatchInstructions 和 sanitizerPatchModule 的访问。例如，如果并发调用 sanitizerPatchModule(mod, *) 和 sanitizerPatchInstruction(*, mod, *)，则结果是未定义的。

参数

module – 要检测的 CUDA 模块

返回值

SANITIZER_SUCCESS – 成功时
SANITIZER_ERROR_INVALID_PARAMETER – 如果 module 不是 CUDA 模块

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

指定回调的用户数据指针。

标记 kernel 的所有后续启动都使用 userdata 指针作为要传递到回调函数的设备内存缓冲区。

参数

kernel – 要链接到用户数据的 CUDA 函数。在此内核的后续启动中的回调将使用 userdata 作为回调数据。
userdata – 设备内存缓冲区。此数据将通过 userdata 参数传递给回调函数。

返回值

SANITIZER_SUCCESS – 成功时

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

指定从设备启动的图中回调可访问的用户数据指针，这些图由指定的主机启动的 graphExec 启动。

标记 graphExec 的所有后续启动，以使 userdata 在从设备启动的图形的设备回调中可用。userdata 不会在回调 userdata 参数中设置，而是必须通过另一种方式访问。请参阅 Sanitizer API 参考手册。此功能仅在驱动程序版本为 535 或更高版本时可用。

参数

graphExec – 将从设备启动 CUDA 图形的 CUDA graphExec。
stream – 与流启动关联的 CUDA 流。
userdata – 设备内存缓冲区。

返回值

SANITIZER_SUCCESS – 成功时

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

指定回调的用户数据指针。

标记 launch 使用 userdata 指针作为要传递到回调函数的设备内存缓冲区。此功能仅在驱动程序版本为 455 或更高版本时可用。

参数

launch – 要链接到用户数据的内核启动。此内核启动中的回调将使用 userdata 作为回调数据。
kernel – 与内核启动关联的 CUDA 函数。
stream – 与流启动关联的 CUDA 流。
userdata – 设备内存缓冲区。此数据将通过 userdata 参数传递给回调函数。

返回值

SANITIZER_SUCCESS – 成功时

SanitizerResult sanitizerUnpatchModule(CUmodule module)

移除模块的现有检测。

删除先前调用 sanitizerPatchModule 对 CUDA 模块执行的任何检测。

注意

线程安全性：API 用户必须序列化对同一模块上的 sanitizerPatchModule 和 sanitizerUnpatchModule 的访问。例如，如果并发调用 sanitizerPatchModule(mod) 和 sanitizerUnpatchModule(mod)，则结果是未定义的。

参数: module – 要在其上删除检测的 CUDA 模块
返回值: SANITIZER_SUCCESS – 成功时

类型定义

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

共享内存上异步归约操作的函数类型。

这可以由 red.async PTX 指令生成

Param userdata: 指向用户数据的指针。请参阅 sanitizerPatchModule。
Param pc: 修补指令的程序计数器。
Param address: 共享内存中的目标地址。
Param mbarAddress: mbarrier 对象的地址。
Param accessSize: 访问大小（以字节为单位）。有效值为 4 和 8。

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

共享内存上异步存储操作的函数类型。

这可以通过 st.async PTX 指令生成

Param userdata: 指向用户数据的指针。请参阅 sanitizerPatchModule。
Param pc: 修补指令的程序计数器。
Param address: 共享内存中的目标地址。
Param mbarAddress: mbarrier 对象的地址。
Param pNewValue: 指向正在写入的新值的指针。
Param accessSize: 访问大小（以字节为单位）。有效值为 4 和 8。

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

屏障回调的函数类型。

Param userdata: 指向用户数据的指针。参见 sanitizerPatchModule
Param pc: 已修补指令的程序计数器
Param barIndex: 屏障索引。
Param threadCount: 预期线程数（必须是 warp 大小的倍数）。
Param flags: 包含有关屏障的信息。请参阅 Sanitizer_BarrierFlags 以解释此值。0 表示所有线程都参与屏障。

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

CUDA 块进入回调的函数类型。

Param userdata: 指向用户数据的指针。参见 sanitizerPatchModule
Param pc: 块入口点的程序计数器

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

CUDA 块退出回调的函数类型。

Param userdata: 指向用户数据的指针。参见 sanitizerPatchModule
Param pc: 已修补指令的程序计数器

typedef SanitizerPatchResult (*SanitizerCallbackBulkCopyGlobalToShared)(void *userdata, uint64_t pc, uint64_t src, uint32_t dst, uint32_t barrier, uint32_t data, uint32_t isMulticast)

从全局内存到共享内存的异步批量复制的函数类型。

这可以通过 cp.async.bulk.shared::cluster.global 指令生成

warp 中的所有活动线程都具有相同的参数值。

Param userdata: 指向用户数据的指针。请参阅 sanitizerPatchModule。
Param pc: 修补指令的程序计数器。
Param dst: 目标 DSMEM 地址。
Param barrier: 关联的 mbarrier 完成机制的 DSMEM 地址
Param src: 源全局内存地址
Param data: 在位 0:15 中包含请求的 16 字节块的数量，在位 16:31 中包含多播掩码
Param isMulticast: 布尔值，指示操作是否为多播

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

缓存控制指令回调的函数类型。

Param userdata: 指向用户数据的指针。参见 sanitizerPatchModule
Param pc: 已修补指令的程序计数器
Param address: 被控制的内存地址
Param kind: 缓存控制的类型。参见 Sanitizer_CacheControlInstructionKind

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

函数调用回调的函数类型。

Param userdata: 指向用户数据的指针。参见 sanitizerPatchModule
Param pc: 已修补指令的程序计数器
Param targetPc: 被调用函数所在的 PC。
Param flags: 包含有关函数调用的信息。

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

集群屏障到达的函数类型。

集群屏障等待的函数类型。

这可以通过 cg::this_cluster().sync() (C++ API) 或 barrier.cluster.arrive (PTX API) 生成。

这可以通过 cg::this_cluster().sync() (C++ API) 或 barrier.cluster.wait (PTX API) 生成。

Param userdata: 指向用户数据的指针。参见 sanitizerPatchModule
Param pc: 已修补指令的程序计数器
Param userdata: 指向用户数据的指针。参见 sanitizerPatchModule
Param pc: 已修补指令的程序计数器
Retval SANITIZER_PATCH_SUCCESS: Warp 执行继续
Retval SANITIZER_PATCH_ERROR: Warp 应该退出

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

CUDA 屏障操作回调的函数类型。

Param userdata: 指向用户数据的指针。参见 sanitizerPatchModule
Param pc: 已修补指令的程序计数器
Param barrier: 可以用作唯一标识符的屏障地址
Param kind: 屏障动作类型。参见 Sanitizer_CudaBarrierInstructionKind
Param data: 屏障数据。这特定于每个动作类型，请参阅 Sanitizer_CudaBarrierInstructionKind

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

设备端 free 调用的函数类型。

注意

这在实际调用之前被调用。

Param userdata: 指向用户数据的指针。参见 sanitizerPatchModule
Param pc: 已修补指令的程序计数器
Param ptr: 传递给设备端 free 的指针。

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

设备端 malloc 调用的函数类型。

注意

这在调用完成后被调用。

Param userdata: 指向用户数据的指针。参见 sanitizerPatchModule
Param pc: 已修补指令的程序计数器
Param allocatedPtr: 设备端 malloc 返回的指针
Param allocatedSize: 用户向设备端 malloc 请求的大小。

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

矩阵共享内存访问回调的函数类型。

Param userdata: 指向用户数据的指针。参见 sanitizerPatchModule
Param pc: 已修补指令的程序计数器
Param address: 正在读取或写入的共享内存的地址。这是共享内存窗口内的偏移量
Param accessSize: 以字节为单位的访问大小。有效值为 16。
Param flags: 包含有关访问类型的信息。请参阅 Sanitizer_DeviceMemoryFlags 以解释此值。
Param count: 访问的矩阵数。
Param pNewValue: 如果访问是写入，则指向正在写入的新值的指针。如果访问是读取或原子操作，则指针将为 NULL。

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

从全局内存到共享内存的异步复制的函数类型。

Param userdata: 指向用户数据的指针。参见 sanitizerPatchModule
Param pc: 已修补指令的程序计数器
Param src: 正在读取的全局内存的地址。如果 src-size 为 0，则可以为 NULL。
Param dst: 正在写入的共享内存的地址。这是共享内存窗口内的偏移量
Param accessSize: 以字节为单位的访问大小。有效值为 4、8 和 16。

typedef SanitizerPatchResult (*SanitizerCallbackMemcpyAsyncBarrier)(void *userdata, uint64_t pc, uint32_t barrier)

用于 mbarrier 完成的 cuda 屏障的函数类型。

这可以通过 cp.async.mbarrier.arrive 生成

Param userdata: 指向用户数据的指针。请参阅 sanitizerPatchModule。
Param pc: 修补指令的程序计数器。
Param barrier: 屏障的共享内存地址。

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

内存访问回调的函数类型。

Param userdata

指向用户数据的指针。参见 sanitizerPatchModule

Param pc

已修补指令的程序计数器

Param ptr

正在访问的内存地址。对于本地或共享内存访问，这是本地或共享内存窗口内的偏移量。

Param accessSize

以字节为单位的访问大小。有效值为 1、2、4、8 和 16。

Param flags

包含有关访问类型的信息。请参阅 Sanitizer_DeviceMemoryFlags 以解释此值。

Param pData

指针，其值取决于访问类型

如果访问是写入，则 pData 指向正在写入的新值。
如果访问是读取且 pData 不为 NULL，则它指向正在使用的已加载字节的 32 位掩码（填充字节将不会出现）。
如果访问是原子操作，则指针将为 NULL。

typedef SanitizerPatchResult (*SanitizerCallbackMemsetShared)(void *userdata, uint64_t pc, uint32_t dst, uint32_t numBlocks)

在共享内存上进行 memset 操作的函数类型。

这可以通过 st.bulk 生成

warp 中的所有活动线程都具有相同的参数值。

Param userdata: 指向用户数据的指针。请参阅 sanitizerPatchModule。
Param pc: 修补指令的程序计数器。
Param dst: 目标共享内存地址。
Param numBlocks: 要设置为零的 8 字节块的数量。

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

流水线提交的函数类型。

这可以通过 pipeline::producer_commit (C++ API)、pipeline_commit (C API) 或 cp.async.commit_group (PTX API) 生成。

Param userdata: 指向用户数据的指针。参见 sanitizerPatchModule
Param pc: 已修补指令的程序计数器

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

流水线等待的函数类型。

这可以通过 pipeline::consumer_wait (C++ API)、pipeline_wait_prior (C API)、cp.async.wait_group 或 cp.async.wait_all (PTX API) 生成。

Param userdata: 指向用户数据的指针。参见 sanitizerPatchModule
Param pc: 已修补指令的程序计数器
Param groups: 管道将等待的组数。0 用于等待所有组。

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

函数返回回调的函数类型。

Param userdata: 指向用户数据的指针。参见 sanitizerPatchModule
Param pc: 已修补指令的程序计数器

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

设置分配给块的共享内存大小的函数类型。

这可以通过 setsmemsize.sync 指令生成

Param userdata: 指向用户数据的指针。请参阅 sanitizerPatchModule。
Param pc: 修补指令的程序计数器。
Param size: 请求的大小（以字节为单位）。

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

shfl 回调的函数类型。

Param userdata: 指向用户数据的指针。参见 sanitizerPatchModule
Param pc: 已修补指令的程序计数器

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

syncwarp 回调的函数类型。

Param userdata: 指向用户数据的指针。参见 sanitizerPatchModule
Param pc: 已修补指令的程序计数器
Param mask: 传递给 __syncwarp() 的线程掩码。

typedef SanitizerPatchResult (*SanitizerCallbackTensorCoreBarrier)(void *userdata, uint64_t pc, uint32_t barrier, uint32_t isMulticast, uint32_t multicastMask)

Blackwell 张量核心屏障的函数类型。

这可以通过 tcgen05.commit 指令生成

warp 中的所有活动线程都具有相同的参数值。

Param userdata: 指向用户数据的指针。请参阅 sanitizerPatchModule。
Param pc: 修补指令的程序计数器。
Param barrier: 关联的 mbarrier 完成机制的 DSMEM 地址
Param isMulticast: 布尔值，指示操作是否为多播
Param multicastMask: 多播掩码，如果 isMulticast 为 true

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

warpgroup MMA fence 的函数类型。

这可以通过 PTX 中的 wgmma.fence 指令生成。

Param userdata: 指向用户数据的指针。请参阅 sanitizerPatchModule。
Param pc: 修补指令的程序计数器。
参数 warpMask: 执行 fence 操作的线程掩码。预期值应为 0x0 或 0xffffffff (full)。预期该值在整个 warp group 中保持一致。报告其他值可能表示目标应用程序中存在编程错误。

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

warpgroup 对齐的异步 MMA 的函数类型。

这可以通过 PTX 中的 wgmma.mma_async 指令生成。

Param userdata: 指向用户数据的指针。参见 sanitizerPatchModule
Param pc: 已修补指令的程序计数器
参数 addressMatrixA: 正在读取的矩阵 A 在共享内存中的地址。仅当 sizeMatrixA 非零且 warpMask 为 full 时，此字段才有效。
参数 sizeMatrixA: 矩阵 A 在共享内存中的大小。值为 0 表示矩阵 A 从寄存器中读取。
参数 addressMatrixB: 正在读取的矩阵 B 在共享内存中的地址。仅当 warpMask 为 full 时，此字段才有效。
参数 sizeMatrixB: 矩阵 B 在共享内存中的大小。该值将始终为非零。
Param flags: 类型为 Sanitizer_WarpgroupMMAAsyncFlags。提供关于访问的信息。即使 warpMask 为零，也应考虑这些标志。
参数 warpMask: 将执行操作并读取操作数的线程掩码。预期值应为 0x0 或 0xffffffff (full)。预期该值在整个 warp group 中保持一致。报告其他值可能表示目标应用程序中存在编程错误。

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

warpgroup MMA 等待组的函数类型。

这可以通过 PTX 中的 wgmma.wait_group 指令生成。

Param userdata: 指向用户数据的指针。请参阅 sanitizerPatchModule。
Param pc: 修补指令的程序计数器。
参数 numGroups: 操作后将保持挂起状态的最大组数。值为零表示在操作之后，warp group 的所有 MMA async 操作都保证已完成。
参数 warpMask: 线程掩码，预期值应为 0x0 或 0xffffffff (full)。预期该值在整个 warp group 中保持一致。报告其他值可能表示目标应用程序中存在编程错误。如果该值有效，则该值对操作没有影响。

typedef struct Sanitizer_Launch_st *Sanitizer_LaunchHandle