Any allocation created with cudaMalloc can be registered as a memory pool using nvtxMemHeapRegister. The following code example allocates 64 bytes and registers the allocation as a memory pool.

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

Please note that Compute Sanitizer only supports nvtxMemHeapRegister with parameters usage = NVTX_MEM_HEAP_USAGE_TYPE_SUB_ALLOCATOR and type = NVTX_MEM_TYPE_VIRTUAL_ADDRESS. If you are using the CUDA runtime API, nvtxMemHeapRegister can be used with allocations created with cuMemAlloc.

An existing pool can be reset to its initial state using nvtxMemHeapReset. The following example resets the pool previously allocated.

// Requires <nvtx3/nvToolsExtMem.h>

nvtxMemHeapReset(nvtxDomain, nvtxPool);

In a similar fashion, a pool can be unregistered using nvtxMemHeapUnregister. An allocation cannot be used after it is unregistered, but the allocation must be freed using cudaFree to dispose of it.

// Requires <nvtx3/nvToolsExtMem.h>

nvtxMemHeapUnregister(nvtxDomain, nvtxPool);

For your convenience, calling cudaFree on a memory pool causes Compute Sanitizer to automatically unregister it.

Once a pool is created, users can create suballocations within this pool using nvtxMemRegionsRegister. For your convenience, you can register multiple regions at the same time. The following example creates a suballocation of 16 bytes at address ptr. Both ptr and ptr + 16 bytes must be part of the pool.

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

For your convenience, Initcheck assumes that a new suballocation is uninitialized, meaning failure to initialize it might result in error reports. Please note that Compute Sanitizer only supports nvtxMemRegionsRegister with parameter regionType = NVTX_MEM_TYPE_VIRTUAL_ADDRESS. Suballocations are considered as regular allocations for NVTX naming and permissions API, therefore it is possible to label them or change their permissions.

Existing suballocations can be resized using nvtxMemRegionsResize. The following example resizes our previous suballocation at address ptr from 16 bytes to 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);

In a similar fashion, existing allocations can be removed using nvtxMemRegionsUnregister. The following example removes our previous suballocation at address 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);

Omitting to unregister a suballocation is reported as a memory leak if Compute Sanitizer is used in combination with option --leak-check yes.

NVTX Permissions API allows any allocation permissions to be restricted using nvtxMemPermissionsAssign. For this example, we use the global program scope (by calling nvtxMemCudaGetProcessWidePermissions), meaning permissions are applied on all kernel launches. This example restricts the allocation at address ptr to read-only permissions.

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

Valid permissions are:

• Read: NVTX_MEM_PERMISSIONS_REGION_FLAGS_READ
• Write: NVTX_MEM_PERMISSIONS_REGION_FLAGS_WRITE
• Atomic: NVTX_MEM_PERMISSIONS_REGION_FLAGS_ATOMIC
• A combination of read, write and atomic (using XORs).
• Reset: NVTX_MEM_PERMISSIONS_REGION_FLAGS_RESET

Using special permission NVTX_MEM_PERMISSIONS_REGION_FLAGS_RESET resets assigned permissions for the specified allocation on the specified scope.

Allocations permissions can be restricted on a per-device basis, using nvtxMemCudaGetDeviceWidePermissions. The following example gets the permissions handle from device device, a handle that is used with nvtxMemPermissionsAssign to change permissions for the allocation at address ptr, previously restricted to read-only on the global scope, and now read-write for kernel launched on device (no atomic allowed).

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

Permissions can be assigned to a specific stream scope thanks to custom permissions objects. You can create one using nvtxMemPermissionsCreate, and bind it to a scope using nvtxMemPermissionsBind. The following example restricts the allocation at address ptr to read-only permissions.

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

On permissions object creation or binding, you can specify inheritance restriction flags. For example, excluding write permissions will block access for all allocations with unassigned permissions on that scope. These are applied:

• nvtxMemPermissionsCreate: applied for kernel launches on stream bound to the created object.
• nvtxMemPermissionsBind: applied for next kernel launch (on this CPU thread) and others using the same stream.

Please note that Compute Sanitizer only supports nvtxMemPermissionsBind with parameter scope = NVTX_MEM_PERMISSIONS_BIND_SCOPE_CUDA_STREAM.

Permissions objects currently bound can be unbound using nvtxMemPermissionsUnbind and destroyed using nvtxMemPermissionsDestroy. Permissions object destruction will result in an unbinding.

// Requires <nvtx3/nvToolsExtMem.h>

nvtxMemPermissionsUnbind(nvtxDomain, NVTX_MEM_PERMISSIONS_BIND_SCOPE_CUDA_STREAM)

nvtxMemPermissionsDestroy(nvtxDomain, permHandle);

Please note that Compute Sanitizer only supports nvtxMemPermissionsUnbind with parameter scope = NVTX_MEM_PERMISSIONS_BIND_SCOPE_CUDA_STREAM.

Peer devices access can be restricted for all allocations using nvtxMemCudaSetPeerAccess. If no permissions are set on an active scope for an allocation using nvtxMemPermissionsAssign, then default permissions set using nvtxMemCudaSetPeerAccess are applied. The following example restricts accesses to read-only on all devices except 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);