Compute Sanitizer API 参考手册

简介

概述

Compute Sanitizer API 支持创建针对 CUDA 应用程序的清理和跟踪工具。此类工具的示例包括内存和竞争条件检查器。Compute Sanitizer API 由三个 API 组成：回调 API、补丁 API 和内存 API。它以动态库的形式在受支持的平台上交付。

用法

兼容性和要求

Compute Sanitizer 工具需要 CUDA 11.0 或更高版本。

Compute Sanitizer API 需要 CUDA 10.1 或更高版本。如果 CUDA 驱动程序版本与 Compute Sanitizer 版本不兼容，Compute Sanitizer API 调用将失败，并显示 SANITIZER_ERROR_NOT_INITIALIZED。

回调 API

Compute Sanitizer 回调 API 允许您在用户代码中注册回调。当应用程序调用 CUDA 运行时或驱动程序函数，或者当 CUDA 驱动程序中发生某些事件时，将调用该回调。回调 API 使用以下术语。

  • 回调域： 回调被分组到域中，以便更轻松地将回调函数与相关的 CUDA 函数或事件组关联。 Sanitizer_CallbackDomain 定义了以下回调域。

    1. CUDA 驱动程序函数

    2. CUDA 运行时函数

    3. CUDA 资源跟踪

    4. CUDA 同步通知

    5. CUDA 网格启动

    6. CUDA memcpy 操作

    7. CUDA memset 操作

    8. CUDA 批量内存操作

  • 回调 ID： 每个回调在相应的回调域内都给定一个唯一的 ID，以便在回调函数中标识它。CUDA 驱动程序 API ID 在 sanitizer_driver_cbid.h 中定义，CUDA 运行时 API ID 在 sanitizer_runtime_cbid.h 中定义。其他回调 ID 在 sanitizer_callbacks.h 中定义。所有这些头文件都作为 sanitizer.h 的一部分包含在内。

  • 回调函数： 回调函数的类型必须为 Sanitizer_CallbackFunc。此函数类型有两个参数来指定回调：域和标识回调发生原因的 ID。该类型还有一个 cbdata 参数，用于传递特定于回调的数据。

  • 订阅者： 订阅者用于将每个回调函数与一个或多个 CUDA API 函数关联。在任何时候，最多只能有一个使用 sanitizerSubscribe 初始化的订阅者。在初始化新的订阅者之前，必须使用 sanitizerUnsubscribe 最终确定现有的订阅者。

订阅者应在进行任何 CUDA API 调用之前初始化，以确保报告数据的正确性。

每个回调域在下面详细描述。除非明确说明，否则不支持从回调函数中调用任何 CUDA 运行时或驱动程序 API。这样做可能会导致应用程序挂起。但是，支持从回调函数中调用 Compute Sanitizer 内存 API。

驱动程序和运行时 API 回调

将回调 API 与 SANITIZER_CB_DOMAIN_DRIVER_API 或 SANITIZER_CB_DOMAIN_RUNTIME_API 域一起使用，可以将回调函数与一个或多个 CUDA API 函数关联。当在应用程序中调用这些 CUDA 函数时，也会调用回调函数。对于这些域，回调函数的 cbdata 参数类型将为 Sanitizer_CallbackData。

您可以从驱动程序或运行时 API 回调函数中调用 cudaDeviceSynchronize、cudaStreamSynchronize、cuCtxSynchronize 和 cuStreamSynchronize。

以下代码显示了用于将回调函数与一个或多个 CUDA API 函数关联的典型序列。为了简单起见，删除了错误检查代码。

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

首先，sanitizerSubscribe 用于使用 my_callback 回调函数初始化订阅者。接下来，sanitizerEnableDomain 用于将该回调与所有 CUDA 运行时函数关联。使用此代码序列将导致每次调用任何 CUDA 运行时 API 函数时都调用 my_callback 两次，一次在进入 CUDA 函数时，一次在 CUDA 函数退出之前。Compute Sanitizer 回调 API 函数 sanitizerEnableCallback 和 sanitizerEnableAllDomains 也可用于将 CUDA API 函数与回调关联。

以下代码显示了一个典型的回调函数。

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
        ...
    }
    ...
}

在回调函数中，可以使用 Sanitizer_CallbackDomain 和 Sanitizer_CallbackId 参数来确定哪个 CUDA API 函数调用正在触发此回调。在上面的示例中，我们正在检查 CUDA 运行时 cudaMemcpy 函数。cbdata 参数保存了一个有用的信息结构，可以在回调中使用。在本例中，我们使用结构的 callbackSite 成员来检测回调是否发生在进入 cudaMemcpy 时，我们使用 functionParams 成员来访问 cudaMemcpy 的参数。要访问参数，我们首先将 functionParams 强制转换为与 cudaMemcpy 函数对应的结构类型。这些参数结构包含在 generated_cuda_runtime_api_meta.h、generated_cuda_meta.h 和许多其他文件中。

资源回调

将回调 API 与 SANITIZER_CB_DOMAIN_RESOURCE 域一起使用，可以将回调函数与某些 CUDA 资源创建和销毁事件关联。例如，当创建 CUDA 上下文时，将使用回调 ID 等于 SANITIZER_CBID_RESOURCE_CONTEXT_CREATED 调用回调函数。对于此域，cbdata 参数是以下类型之一

  • Sanitizer_ResourceContextData 用于 CUDA 上下文创建和销毁

  • Sanitizer_ResourceStreamData 用于 CUDA 流创建和销毁

  • Sanitizer_ResourceModuleData 用于 CUDA 模块加载和卸载

  • Sanitizer_ResourceMemoryData 用于 CUDA 内存分配和释放

同步回调

将回调 API 与 SANITIZER_CB_DOMAIN_SYNCHRONIZE 域一起使用，可以将回调函数与 CUDA 上下文和流同步关联。例如，当 CUDA 上下文同步时，将使用回调 ID 等于 SANITIZER_CBID_SYNCHRONIZE_CONTEXT_SYNCHRONIZED 调用回调函数。对于此域，cbdata 参数的类型为 Sanitizer_SynchronizeData。

启动回调

将回调 API 与 SANITIZER_CB_DOMAIN_LAUNCH 域一起使用，可以将回调函数与 CUDA 内核启动关联。例如，当 CUDA 内核启动开始时，将使用回调 ID 等于 SANITIZER_CBID_LAUNCH_BEGIN 调用回调函数。对于此域，cbdata 参数的类型为 Sanitizer_LaunchData。

Memcpy 回调

将回调 API 与 SANITIZER_CB_DOMAIN_MEMCPY 域一起使用，可以将回调函数与 CUDA memcpy 操作关联。例如，当调用 cudaMemcpy API 函数时，将使用回调 ID 等于 SANITIZER_CBID_MEMCPY_STARTING 调用回调函数。对于此域，cbdata 参数的类型为 Sanitizer_MemcpyData。

Memset 回调

将回调 API 与 SANITIZER_CB_DOMAIN_MEMSET 域一起使用，可以将回调函数与 CUDA memset 操作关联。例如，当调用 cudaMemset API 函数时，将使用回调 ID 等于 SANITIZER_CBID_MEMSET_STARTING 调用回调函数。对于此域，cbdata 参数的类型为 Sanitizer_MemsetData。

批量内存操作回调

将回调 API 与 SANITIZER_CB_DOMAIN_BATCH_MEMOP 域一起使用，可以将回调函数与 CUDA 批量内存操作关联。例如，当调用 cuStreamWriteValue API 函数时，将使用回调 ID 等于 SANITIZER_CBID_BATCH_MEMOP_WRITE 调用回调函数。对于此域，cbdata 参数的类型为 Sanitizer_BatchMemopData。

补丁 API

Compute Sanitizer 补丁 API 允许您加载补丁函数并将它们插入到用户代码中。当应用程序的 CUDA 代码执行某些指令或调用某些 CUDA 设备函数时，将调用补丁函数。补丁 API 使用以下术语

  • 指令 ID：每个可修补的事件都给定一个唯一的 ID，该 ID 可以传递给补丁 API 函数，以指定应修补这些事件。指令 ID 由 Sanitizer_InstructionId 定义。

  • Instrumentation point: Compute Sanitizer API 正在检测原始 CUDA 代码中的位置。执行时，用户代码路径将被修改，以便在修补事件之前或之后执行补丁。所有补丁都在事件之前执行，设备端 malloc 除外。

  • 补丁：Compute Sanitizer 将插入到另一个现有 CUDA 代码中的 CUDA __device__ 函数。补丁函数签名必须与 API 期望的签名匹配（有关预期的签名类型，请参见下文）。

编写补丁

补丁必须遵循 Compute Sanitizer API 针对给定指令 ID 所需的函数签名。指令 ID 到函数签名的映射记录在 sanitizer_patching.h 中 Sanitizer_InstructionId 的注释中。例如，如果我们希望使用指令 ID SANITIZER_INSTRUCTION_MEMORY_ACCESS 修补内存访问，则需要使用 SanitizerCallbackMemoryAccess 类型。

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;
}

在此补丁中，我们记录对我们先前分配的结构的写入和读取访问。extern "C" 确保补丁名称不会被破坏，从而允许我们直接在调用 sanitizerPatchInstructions 中使用其名称作为字符串（请参阅下面）。

可以在单个 CUDA 文件中定义多个补丁。然后必须使用以下 nvcc 选项编译此文件

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

如果首选 fatbin 而不是 cubin 作为输出文件，则 --cubin 选项可以替换为 --fatbin。

插入补丁

生成补丁后，可以使用以下过程将其插入到用户代码中

  1. 加载补丁。有两个 API 用于加载补丁：sanitizerAddPatchesFromFile 和 sanitizerAddPatches。它们使用与 cuModuleLoad 和 cuModuleLoadData 相同的输入格式。

  2. 通过使用 sanitizerPatchInstructions API 选择要修补的指令。

  3. 通过使用 sanitizerPatchModule API 修补用户代码。

  4. 可选地，通过使用 sanitizerSetCallbackData API 设置补丁的回调数据。

以下代码显示了使用这些 API 的典型序列。为了简单起见，删除了错误检查。

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);

所有后续使用来自此 CUDA 模块的代码的启动都将被检测，并且将在每次内存访问之前调用 my_memory_access_callback。但是，回调数据仅针对给定内核的所有后续启动设置。获得内核 CUfunction 的一种简单方法是通过 Sanitizer 启动回调。可以使用 sanitizerUnpatchModule API 删除检测。

内存 API

Compute Sanitizer 内存 API 为 CUDA 内存 API 提供了替换函数，可以从 Compute Sanitizer 回调 中安全地调用这些函数。

  • sanitizerAlloc 是 cudaMalloc 的替代品。

  • sanitizerFree 是 cudaFree 的替代品。

  • sanitizerMemcpyHostToDeviceAsync 是用于主机到设备复制的 cudaMemcpyAsync 的替代品。

  • sanitizerMemcpyDeviceToHost 是用于设备到主机复制的 cudaMemcpy 的替代品。

  • sanitizerMemset 是 cudaMemset 的替代品。

这些函数也可以在普通用户代码中调用，在普通用户代码中，它们可以与 CUDA API 混合使用。例如，使用 sanitizerAlloc 分配的内存可以使用 cudaFree 释放。但是，由于只有 CUDA API 调用才会导致 回调 被调用，这可能会导致不连贯的跟踪状态，应避免这样做。

特殊情况

在某些特定情况下，Compute Sanitizer API 的行为可能与一般情况不同。本节列出了这些情况。

设备图启动

当 Compute Sanitizer API 检测的应用程序使用设备启动的 CUDA 图时，应考虑以下行为

  • 当从设备启动 CUDA 图时，不会从主机调用回调。但是，当从主机启动或上传设备可启动图时，仍然会调用与图启动相关的回调。

  • 如果在每个节点设置了不同的用户数据指针，则在设备启动的 CUDA 图内部时，设备上补丁 API 发出的回调可能具有来自不同节点的用户数据。

为了补偿这些限制，可以使用 sanitizerSetDeviceGraphData 设置设备启动的图特定数据。可以在主机上启动设备可启动图或包含设备图启动的图期间调用它。然后可以从此主机启动的图启动的任何设备图中检索使用 sanitizerSetDeviceGraphData 设置的数据。要从补丁 API 设备回调中检索此数据，可以使用以下代码

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

但是，建议在每次解引用时执行 NULL 检查。

限制

目前没有已知问题。

声明

声明

所有 NVIDIA 设计规范、参考板、文件、图纸、诊断程序、列表和其他文档（统称为“资料”）均按“原样”提供。NVIDIA 不对这些资料作任何明示、暗示、法定或其他形式的保证，并且明确声明不承担所有关于不侵权、适销性和特定用途适用性的暗示保证。

本文档提供的信息据信是准确可靠的。但是，对于因使用此类信息或因使用此类信息而可能导致的侵犯专利或第三方的其他权利的后果，NVIDIA 公司不承担任何责任。在 NVIDIA 公司的任何专利权项下，未通过暗示或其他方式授予任何许可。本出版物中提及的规格如有更改，恕不另行通知。本出版物取代并替换之前提供的所有其他信息。未经 NVIDIA 公司明确书面批准，NVIDIA 公司产品未被授权用作生命维持设备或系统中的关键组件。

商标

NVIDIA 和 NVIDIA 徽标是 NVIDIA 公司在美国和其他国家/地区的商标和/或注册商标。其他公司和产品名称可能是与其关联的各自公司的商标。