NVTX API for Compute Sanitizer 参考手册

简介

概述

用于 Compute Sanitizer 的 NVTX 内存 API 允许 CUDA 程序向 Compute Sanitizer 通知内存限制：内存池管理或权限限制，以及内存标记。这些工具通过 NVTX (NVIDIA Tools Extension) 得到通知，NVTX 是一个仅包含头文件的 C 库，被各种 NVIDIA 工具使用。最新的 NVTX 头文件可以从 我们的 GitHub 仓库（实验性分支） 下载。

此 API 具有以下主要目标

  • 程序可以将分配标记为内存池，从而使 Compute Sanitizer 能够了解此特定分配的哪些部分实际被使用。当使用 Memcheck 工具时，如果程序访问了池的未注册部分，您将收到通知，否则可能会错过这些错误。当将 Initcheck 工具与选项 --track-unused-memory yes 结合使用时，您不会收到关于未注册区域中未使用内存的通知，从而避免误报。

  • 程序可以使用有意义的名称标记分配，从而使您可以通过名称（例如，泄漏或未使用的分配）来识别与特定错误关联的分配。

  • 程序可以将某些分配限制为一组特定的权限（例如，只读或只写），这些权限适用于特定的范围（例如，CUDA 流、设备或整个程序）。当使用 Memcheck 工具时，违反这些限制将导致错误。

用法

兼容性和要求

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

Compute Sanitizer 从 CUDA 11.3 开始支持 NVTX 内存 API，使用 --nvtx yes 选项。从 CUDA 12.0 开始，此选项默认启用。

Compute Sanitizer 需要在调用 NVTX 之前初始化 CUDA 运行时。

// NVTX calls are not allowed before CUDA runtime initialization.

// Forces CUDA runtime initialization.
cudaFree(0);

// NVTX calls are now allowed.

NVTX 结构必须进行零初始化。此页面上的示例使用 C++ 空初始化器 ({})。如果您使用的是 C，则可以使用 memset 或使用至少包含一个字段的初始化器语法（C 不支持空初始化器）。

NVTX 域

所有 NVTX 调用都需要您创建一个 NVTX 域。这可以使用 nvtxDomainCreateA 来实现。

// Requires <nvtx3/nvToolsExt.h>

auto nvtxDomain = nvtxDomainCreateA("my-domain");

目前，NVTX 域没有特定用途，但在未来的 Compute Sanitizer 版本中将会有用途。

子分配 API

池管理

使用 cudaMalloc 创建的任何分配都可以使用 nvtxMemHeapRegister 注册为内存池。以下代码示例分配 64 字节并将该分配注册为内存池。

// Requires <nvtx3/nvToolsExtMem.h>
// (see https://github.com/NVIDIA/NVTX/tree/dev-mem-api/c/include)

void *ptr;
cudaMalloc(&ptr, 64);

nvtxMemVirtualRangeDesc_t nvtxRangeDesc = {};
nvtxRangeDesc.size = 64;
nvtxRangeDesc.ptr = ptr;

nvtxMemHeapDesc_t nvtxHeapDesc = {};
nvtxHeapDesc.extCompatID = NVTX_EXT_COMPATID_MEM;
nvtxHeapDesc.structSize = sizeof(nvtxMemHeapDesc_t);
nvtxHeapDesc.usage = NVTX_MEM_HEAP_USAGE_TYPE_SUB_ALLOCATOR;
nvtxHeapDesc.type = NVTX_MEM_TYPE_VIRTUAL_ADDRESS;
nvtxHeapDesc.typeSpecificDescSize = sizeof(nvtxMemVirtualRangeDesc_t);
nvtxHeapDesc.typeSpecificDesc = &nvtxRangeDesc;

auto nvtxPool = nvtxMemHeapRegister(
    nvtxDomain,
    &nvtxHeapDesc);

请注意，Compute Sanitizer 仅支持参数为 usage = NVTX_MEM_HEAP_USAGE_TYPE_SUB_ALLOCATOR 和 type = NVTX_MEM_TYPE_VIRTUAL_ADDRESS 的 nvtxMemHeapRegister。如果您使用的是 CUDA 运行时 API，则可以将 nvtxMemHeapRegister 与使用 cuMemAlloc 创建的分配一起使用。

可以使用 nvtxMemHeapReset 将现有池重置为其初始状态。以下示例重置先前分配的池。

// Requires <nvtx3/nvToolsExtMem.h>

nvtxMemHeapReset(nvtxDomain, nvtxPool);

类似地，可以使用 nvtxMemHeapUnregister 注销池。注销后分配不能使用，但必须使用 cudaFree 释放分配以处理它。

// Requires <nvtx3/nvToolsExtMem.h>

nvtxMemHeapUnregister(nvtxDomain, nvtxPool);

为了您的方便，在内存池上调用 cudaFree 会导致 Compute Sanitizer 自动注销它。

子分配管理

创建池后，用户可以使用 nvtxMemRegionsRegister 在此池中创建子分配。为了您的方便，您可以同时注册多个区域。以下示例在地址 ptr 创建 16 字节的子分配。ptr 和 ptr + 16 bytes 都必须是池的一部分。

// Requires <nvtx3/nvToolsExtMem.h>

nvtxMemVirtualRangeDesc_t nvtxRangeDesc = {};
nvtxRangeDesc.size = 16;
nvtxRangeDesc.ptr = ptr;

nvtxMemRegionsRegisterBatch_t nvtxRegionsDesc = {};
nvtxRegionsDesc.extCompatID = NVTX_EXT_COMPATID_MEM;
nvtxRegionsDesc.structSize = sizeof(nvtxMemRegionsRegisterBatch_t);
nvtxRegionsDesc.regionType = NVTX_MEM_TYPE_VIRTUAL_ADDRESS;
nvtxRegionsDesc.heap = nvtxPool;
nvtxRegionsDesc.regionCount = 1;
nvtxRegionsDesc.regionDescElementSize = sizeof(nvtxMemVirtualRangeDesc_t);
nvtxRegionsDesc.regionDescElements = &nvtxRangeDesc;

nvtxMemRegionsRegister(nvtxDomain, &nvtxRegionsDesc);

为了您的方便，Initcheck 假设新的子分配未初始化，这意味着未能初始化它可能会导致错误报告。请注意，Compute Sanitizer 仅支持参数为 regionType = NVTX_MEM_TYPE_VIRTUAL_ADDRESS 的 nvtxMemRegionsRegister。子分配被视为 NVTX 命名 和 权限 API 的常规分配，因此可以标记它们或更改它们的权限。

可以使用 nvtxMemRegionsResize 调整现有子分配的大小。以下示例将地址 ptr 处的先前子分配从 16 字节调整为 32 字节。

// Requires <nvtx3/nvToolsExtMem.h>

nvtxMemVirtualRangeDesc_t nvtxRangeDesc = {};
nvtxRangeDesc.size = 32;
nvtxRangeDesc.ptr = ptr;

nvtxMemRegionsResizeBatch_t nvtxRegionsDesc = {};
nvtxRegionsDesc.extCompatID = NVTX_EXT_COMPATID_MEM;
nvtxRegionsDesc.structSize = sizeof(nvtxMemRegionsResizeBatch_t);
nvtxRegionsDesc.regionType = NVTX_MEM_TYPE_VIRTUAL_ADDRESS;
nvtxRegionsDesc.regionDescCount = 1;
nvtxRegionsDesc.regionDescElementSize = sizeof(nvtxMemVirtualRangeDesc_t);
nvtxRegionsDesc.regionDescElements = &nvtxRangeDesc;

nvtxMemRegionsResize(nvtxDomain, &nvtxRegionsDesc);

类似地，可以使用 nvtxMemRegionsUnregister 删除现有分配。以下示例删除地址 ptr 处的先前子分配。

nvtxMemRegionRef_t nvtxRegionRef;
nvtxRegionRef.pointer = ptr;

nvtxMemRegionsUnregisterBatch_t nvtxRegionsDesc = {};
nvtxRegionsDesc.extCompatID = NVTX_EXT_COMPATID_MEM;
nvtxRegionsDesc.structSize = sizeof(nvtxMemRegionsUnregisterBatch_t);
nvtxRegionsDesc.refType = NVTX_MEM_REGION_REF_TYPE_POINTER;
nvtxRegionsDesc.refCount = 1;
nvtxRegionsDesc.refElementSize = sizeof(nvtxMemRegionRef_t);
nvtxRegionsDesc.refElements = &nvtxRegionRef;

nvtxMemRegionsUnregister(nvtxDomain, &nvtxRegionsDesc);

如果 Compute Sanitizer 与选项 --leak-check yes 结合使用，则省略注销子分配会被报告为内存泄漏。

命名 API

可以为任何分配分配名称，以便未来的 Compute Sanitizer 错误报告可以通过其名称引用分配。此示例将地址 ptr 处的分配命名为：“My Allocation”。

// Requires <nvtx3/nvToolsExtMem.h>

nvtxMemRegionNameDesc_t nvtxLabelDesc;
nvtxLabelDesc.regionRefType = NVTX_MEM_REGION_REF_TYPE_POINTER;
nvtxLabelDesc.nameType = NVTX_MESSAGE_TYPE_ASCII;
nvtxLabelDesc.region.pointer = ptr;
nvtxLabelDesc.name.ascii = "My Allocation";

nvtxMemRegionsNameBatch_t nvtxRegionsDesc = {};
nvtxRegionsDesc.extCompatID = NVTX_EXT_COMPATID_MEM;
nvtxRegionsDesc.structSize = sizeof(nvtxMemRegionsNameBatch_t);
nvtxRegionsDesc.regionCount = 1;
nvtxRegionsDesc.regionElementSize = sizeof(nvtxMemRegionNameDesc_t);
nvtxRegionsDesc.regionElements = &nvtxLabelDesc;

nvtxMemRegionsName(nvtxDomain, &nvtxRegionsDesc);

请注意，Compute Sanitizer 仅支持 nameType = NVTX_MESSAGE_TYPE_ASCII 参数的 nvtxMemRegionsName，用于 regionElements 中的所有区域元素。到目前为止，只有泄漏和未使用的内存报告功能分配名称。

权限 API

基本权限管理

NVTX 权限 API 允许使用 nvtxMemPermissionsAssign 限制任何分配权限。在此示例中，我们使用全局程序范围（通过调用 nvtxMemCudaGetProcessWidePermissions），这意味着权限应用于所有内核启动。此示例将地址 ptr 处的分配限制为只读权限。

// Requires <nvtx3/nvToolsExtMem.h> and <nvtx3/nvToolsExtMemCudaRt.h>

auto processPermHandle = nvtxMemCudaGetProcessWidePermissions(nvtxDomain);

nvtxMemPermissionsAssignRegionDesc_t nvtxPermDesc;
nvtxPermDesc.flags = NVTX_MEM_PERMISSIONS_REGION_FLAGS_READ;
nvtxPermDesc.regionRefType = NVTX_MEM_REGION_REF_TYPE_POINTER;
nvtxPermDesc.region.pointer = ptr;

nvtxMemPermissionsAssignBatch_t nvtxRegionsDesc = {};
nvtxRegionsDesc.extCompatID = NVTX_EXT_COMPATID_MEM;
nvtxRegionsDesc.structSize = sizeof(nvtxMemPermissionsAssignBatch_t);
nvtxRegionsDesc.permissions = processPermHandle;
nvtxRegionsDesc.regionCount = 1;
nvtxRegionsDesc.regionElementSize = sizeof(nvtxMemPermissionsAssignRegionDesc_t);
nvtxRegionsDesc.regionElements = &nvtxPermDesc;

nvtxMemPermissionsAssign(nvtxDomain, &nvtxRegionsDesc);

有效权限为

  • 读取: NVTX_MEM_PERMISSIONS_REGION_FLAGS_READ

  • 写入: NVTX_MEM_PERMISSIONS_REGION_FLAGS_WRITE

  • 原子: NVTX_MEM_PERMISSIONS_REGION_FLAGS_ATOMIC

  • 读取、写入和原子的组合（使用异或）。

  • 重置: NVTX_MEM_PERMISSIONS_REGION_FLAGS_RESET

使用特殊权限 NVTX_MEM_PERMISSIONS_REGION_FLAGS_RESET 会重置指定范围内指定分配的已分配权限。

可以使用 nvtxMemCudaGetDeviceWidePermissions 在每个设备的基础上限制分配权限。以下示例从设备 device 获取权限句柄，该句柄与 nvtxMemPermissionsAssign 一起使用以更改地址 ptr 处的分配的权限，该分配先前在全局范围内被限制为只读，现在对于在 device 上启动的内核是读写（不允许原子操作）。

// Requires <nvtx3/nvToolsExtMem.h> and <nvtx3/nvToolsExtMemCudaRt.h>

auto devicePermHandle = nvtxMemCudaGetDeviceWidePermissions(nvtxDomain, device);

nvtxMemPermissionsAssignRegionDesc_t nvtxPermDesc;
nvtxPermDesc.flags = NVTX_MEM_PERMISSIONS_REGION_FLAGS_READ | NVTX_MEM_PERMISSIONS_REGION_FLAGS_WRITE;
nvtxPermDesc.regionRefType = NVTX_MEM_REGION_REF_TYPE_POINTER;
nvtxPermDesc.region.pointer = ptr;

nvtxMemPermissionsAssignBatch_t nvtxRegionsDesc = {};
nvtxRegionsDesc.extCompatID = NVTX_EXT_COMPATID_MEM;
nvtxRegionsDesc.structSize = sizeof(nvtxMemPermissionsAssignBatch_t);
nvtxRegionsDesc.permissions = devicePermHandle;
nvtxRegionsDesc.regionCount = 1;
nvtxRegionsDesc.regionElementSize = sizeof(nvtxMemPermissionsAssignRegionDesc_t);
nvtxRegionsDesc.regionElements = &nvtxPermDesc;

nvtxMemPermissionsAssign(nvtxDomain, &nvtxRegionsDesc);

高级权限管理

由于自定义权限对象，可以将权限分配给特定的流范围。您可以使用 nvtxMemPermissionsCreate 创建一个，并使用 nvtxMemPermissionsBind 将其绑定到范围。以下示例将地址 ptr 处的分配限制为只读权限。

// Requires <nvtx3/nvToolsExtMem.h> and <nvtx3/nvToolsExtMemCudaRt.h>

// Create new permissions object.
auto permHandle = nvtxMemPermissionsCreate(nvtxDomain, NVTX_MEM_PERMISSIONS_CREATE_FLAGS_NONE);

nvtxMemPermissionsAssignRegionDesc_t nvtxPermDesc;
nvtxPermDesc.flags = NVTX_MEM_PERMISSIONS_REGION_FLAGS_READ;
nvtxPermDesc.regionRefType = NVTX_MEM_REGION_REF_TYPE_POINTER;
nvtxPermDesc.region.pointer = ptr;

nvtxMemPermissionsAssignBatch_t nvtxRegionsDesc = {};
nvtxRegionsDesc.extCompatID = NVTX_EXT_COMPATID_MEM;
nvtxRegionsDesc.structSize = sizeof(nvtxMemPermissionsAssignBatch_t);
nvtxRegionsDesc.permissions = permHandle;
nvtxRegionsDesc.regionCount = 1;
nvtxRegionsDesc.regionElementSize = sizeof(nvtxMemPermissionsAssignRegionDesc_t);
nvtxRegionsDesc.regionElements = &nvtxPermDesc;

// Assign read-only permissions to allocation at address ptr.
// Permissions will be applied on scope bound to permHandle.
nvtxMemPermissionsAssign(nvtxDomain, &nvtxRegionsDesc);

// Binding will happen on next kernel launch on this CPU thread, meaning the
// stream for this launch will be the one bound to this permissions object.
nvtxMemPermissionsBind(
    nvtxDomain,
    permHandle,
    NVTX_MEM_PERMISSIONS_BIND_SCOPE_CUDA_STREAM,
    NVTX_MEM_PERMISSIONS_BIND_FLAGS_NONE);

// permHandle is now bound to stream.
MyKernel<<<BlocksNb, ThreadsNb, 0, stream>>>(ptr);

在权限对象创建或绑定时，您可以指定继承限制标志。例如，排除写入权限将阻止对该范围内所有具有未分配权限的分配进行访问。这些标志应用在

  • nvtxMemPermissionsCreate：应用于绑定到创建对象的流上的内核启动。

  • nvtxMemPermissionsBind：应用于下一个内核启动（在此 CPU 线程上）以及其他使用相同流的内核启动。

请注意，Compute Sanitizer 仅支持参数为 scope = NVTX_MEM_PERMISSIONS_BIND_SCOPE_CUDA_STREAM 的 nvtxMemPermissionsBind。

当前绑定的权限对象可以使用 nvtxMemPermissionsUnbind 解绑，并使用 nvtxMemPermissionsDestroy 销毁。权限对象销毁将导致解绑。

// Requires <nvtx3/nvToolsExtMem.h>

nvtxMemPermissionsUnbind(nvtxDomain, NVTX_MEM_PERMISSIONS_BIND_SCOPE_CUDA_STREAM)

nvtxMemPermissionsDestroy(nvtxDomain, permHandle);

请注意，Compute Sanitizer 仅支持参数为 scope = NVTX_MEM_PERMISSIONS_BIND_SCOPE_CUDA_STREAM 的 nvtxMemPermissionsUnbind。

可以使用 nvtxMemCudaSetPeerAccess 限制对所有分配的对等设备访问。如果未使用 nvtxMemPermissionsAssign 在活动范围内为分配设置任何权限，则应用使用 nvtxMemCudaSetPeerAccess 设置的默认权限。以下示例限制除 device 之外的所有设备的访问为只读。

// Requires <nvtx3/nvToolsExtMem.h>

auto permHandle = nvtxMemCudaGetDeviceWidePermissions(nvtxDomain, device);
nvtxMemCudaSetPeerAccess(
    nvtxDomain,
    permHandle,
    NVTX_MEM_CUDA_PEER_ALL_DEVICES,
    NVTX_MEM_PERMISSIONS_REGION_FLAGS_READ);
nvtxMemCudaSetPeerAccess(
    nvtxDomain,
    permHandle,
    device,
    NVTX_MEM_PERMISSIONS_REGION_FLAGS_READ | NVTX_MEM_PERMISSIONS_REGION_FLAGS_WRITE | NVTX_MEM_PERMISSIONS_REGION_FLAGS_ATOMIC);

局限性

请注意，Compute Sanitizer 对 NVTX 内存 API 的支持具有以下局限性

  • 分配名称在泄漏和未使用的内存报告中可见，但目前在其他错误报告中不可见。

  • 分配名称必须以 ASCII 编码，仅包含可打印字符，并且包含 1 到 49 个字符（必须符合以下正则表达式：^[:print:]{1,49}$）

  • 权限仅应用于内核启动。其他操作，如 cudaMemcpy 或 cudaMemset，目前不受支持。

声明

声明

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

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

商标

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