CUDA虚拟内存管理分析

CUDA虚拟内存管理分析

10.1. 简介

虚拟内存管理 API 为应用程序提供了一种直接管理 CUDA 提供的统一虚拟地址空间的方法，用于将物理内存映射到 GPU 可访问的虚拟地址。在 CUDA 10.2 中引入的这些 API 还提供了一种与其他进程和图形 API （如 OpenGL 和 Vulkan）互操作的新方法，并提供用户可以调整以适应其应用程序的更新内存属性。

从历史上看，CUDA 编程模型中的内存分配调用（例如cudaMalloc()）返回了指向 GPU 内存的内存地址。这样获得的地址可以与任何 CUDA API 一起使用，也可以在设备内核中使用。但是，无法根据用户的内存需求调整分配的内存大小。为了增加分配的大小，用户必须显式分配更大的缓冲区，从初始分配中复制数据，释放它，然后继续跟踪较新的分配地址。这通常会导致应用程序的性能降低和峰值内存利用率提高。从本质上讲，用户有一个类似malloc的接口来分配GPU内存，但没有相应的realloc来补充它。虚拟内存管理 API 将地址和内存的概念分离，并允许应用程序单独处理它们。API 允许应用程序根据需要从虚拟地址范围映射和取消映射内存。

如果使用cudaEnablePeerAccess 启用对等设备对内存分配的访问，则所有过去和将来的用户分配都将映射到目标对等设备。这导致用户在不知不觉中支付将所有 cudaMalloc 分配映射到对等设备的运行时成本。但是，在大多数情况下，应用程序仅通过与另一个设备共享几个分配来进行通信，并且并非所有分配都需要映射到所有设备。使用虚拟内存管理，应用程序可以专门选择要从目标设备访问的某些分配。

CUDA 虚拟内存管理 API 向用户公开细粒度控制，用于管理应用程序中的 GPU 内存。它提供的 API 允许用户：

• 将分配给不同设备的内存放入连续的 VA 范围。
• 使用特定于平台的机制执行进程间通信以实现内存共享。
• 在支持它们的设备上选择使用较新的内存类型。

为了分配内存，虚拟内存管理编程模型公开了以下功能：

• 分配物理内存。
• 保留 VA 范围。
• 将分配的内存映射到 VA 范围。
• 控制对映射范围的访问权限。

请注意，本节中描述的 API 套件需要支持 UVA 的系统。

10.2. 查询支持

在尝试使用虚拟内存管理 API 之前，应用程序必须确保它们要使用的设备支持 CUDA 虚拟内存管理。以下代码示例演示如何查询虚拟内存管理支持：

int deviceSupportsVmm;

CUresult result = cuDeviceGetAttribute(&deviceSupportsVmm, CU_DEVICE_ATTRIBUTE_VIRTUAL_MEMORY_MANAGEMENT_SUPPORTED, device);

if (deviceSupportsVmm != 0) {

    // `device` supports Virtual Memory Management

}

10.3. 分配物理内存

使用虚拟内存管理 API 进行内存分配的第一步是创建一个物理内存块，该块将为分配提供支持。为了分配物理内存，应用程序必须使用cuMemCreate API。此函数创建的分配没有任何设备或主机映射。函数参数CUmemGenericAllocationHandle描述要分配的内存的属性，例如分配的位置、是否要将分配共享给另一个进程（或其他图形 API）或要分配的内存的物理属性。用户必须确保请求的分配大小必须与适当的粒度保持一致。可以使用 cuMemGetAllocationGranularity查询有关分配粒度要求的信息。以下代码片段显示了使用cuMemCreate 分配物理内存：

CUmemGenericAllocationHandle allocatePhysicalMemory(int device, size_t size) {

    CUmemAllocationProp prop = {};

    prop.type = CU_MEM_ALLOCATION_TYPE_PINNED;

    prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE;

    prop.location.id = device;

    cuMemGetAllocationGranularity(&granularity, &prop, CU_MEM_ALLOC_GRANULARITY_MINIMUM);

    // Ensure size matches granularity requirements for the allocation

    size_t padded_size = ROUND_UP(size, granularity);

    // Allocate physical memory

    CUmemGenericAllocationHandle allocHandle;

    cuMemCreate(&allocHandle, padded_size, &prop, 0);

    return allocHandle;

}

cuMemCreate分配的内存由CUmemGenericAllocationHandle返回的引用。这与 cudaMalloc 风格的分配不同，后者返回指向 GPU 内存的指针，该指针可由设备上执行的 CUDA 内核直接访问。分配的内存不能用于除使用cuMemGetAllocationPropertiesFromHandle 查询属性以外的任何操作。为了使此内存可访问，应用程序必须将此内存映射到cuMemAddressReserve 保留的 VA 范围，并提供对其的适当访问权限。应用程序必须使用cuMemRelease API 释放分配的内存。

10.3.1. 可共享内存分配

cuMemCreate用户现在可以在分配时向 CUDA 指示他们已指定特定分配用于进程间通信和图形互操作目的。应用程序可以通过设置CUmemAllocationProp::requestedHandleTypes为特定于平台的字段来执行此操作。在 Windows 上，CUmemAllocationProp::requestedHandleTypes设置为CU_MEM_HANDLE_TYPE_WIN32 应用程序时还必须在CUmemAllocationProp::win32HandleMetaData 中指定 LPSECURITYATTRIBUTES 属性。此安全属性定义导出的分配可以传输到其他进程的范围。

CUDA 虚拟内存管理 API 函数不支持与其内存的传统进程间通信函数。相反，它们公开了一种使用特定于操作系统的句柄的进程间通信的新机制。应用程序可以使用cuMemExportToShareableHandle 获取与分配对应的这些特定于操作系统的句柄。这样获得的句柄可以通过使用通常的操作系统本机机制进行进程间通信来传输。收件人进程应使用cuMemImportFromShareableHandle 导入分配。

用户必须确保在尝试导出分配有cuMemCreate 的内存之前查询对所请求的句柄类型的支持。以下代码片段以特定于平台的方式演示了句柄类型支持的查询。

int deviceSupportsIpcHandle;

#if defined(__linux__)

    cuDeviceGetAttribute(&deviceSupportsIpcHandle, CU_DEVICE_ATTRIBUTE_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR_SUPPORTED, device));

#else

    cuDeviceGetAttribute(&deviceSupportsIpcHandle, CU_DEVICE_ATTRIBUTE_HANDLE_TYPE_WIN32_HANDLE_SUPPORTED, device));

#endif

用户应按如下所示进行适当的设置CUmemAllocationProp::requestedHandleTypes：

#if defined(__linux__)

    prop.requestedHandleTypes = CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR;

#else

    prop.requestedHandleTypes = CU_MEM_HANDLE_TYPE_WIN32;

    prop.win32HandleMetaData = // Windows specific LPSECURITYATTRIBUTES attribute.

#endif

memMapIpcDrv 示例可用作将 IPC 与虚拟内存管理分配一起使用的示例。

10.3.2. 内存类型

在 CUDA 10.2 之前，应用程序没有用户控制的方式来分配某些设备可能支持的任何特殊类型的内存。使用cuMemCreateCUmemAllocationProp::allocFlags ，应用程序还可以使用 来指定内存类型要求，以选择加入任何特定的内存功能。应用程序还必须确保分配设备上支持请求的内存类型。

10.3.2.1. 可压缩内存

可压缩内存可用于加速对具有非结构化稀疏性和其他可压缩数据模式的数据的访问。压缩可以节省 DRAM 带宽、L2 读取带宽和 L2 容量，具体取决于正在操作的数据。想要在支持计算数据压缩的设备上分配可压缩内存的应用程序可以通过设置CUmemAllocationProp::allocFlags::compressionType为CU_MEM_ALLOCATION_COMP_GENERIC .用户必须使用CU_DEVICE_ATTRIBUTE_GENERIC_COMPRESSION_SUPPORTED 查询设备是否支持计算数据压缩。以下代码片段说明了查询可压缩内存支持 cuDeviceGetAttribute。

int compressionSupported = 0;

cuDeviceGetAttribute(&compressionSupported, CU_DEVICE_ATTRIBUTE_GENERIC_COMPRESSION_SUPPORTED, device);

在支持计算数据压缩的设备上，用户必须在分配时选择加入，如下所示：

prop.allocFlags.compressionType = CU_MEM_ALLOCATION_COMP_GENERIC;

由于硬件资源有限等各种原因，分配可能没有压缩属性，用户应使用cuMemGetAllocationPropertiesFromHandle查询分配内存的属性并检查压缩属性。

CUmemAllocationPropPrivate allocationProp = {};

cuMemGetAllocationPropertiesFromHandle(&allocationProp, allocationHandle);

if (allocationProp.allocFlags.compressionType == CU_MEM_ALLOCATION_COMP_GENERIC)

{

    // Obtained compressible memory allocation

}

10.4. 保留虚拟地址范围

由于使用虚拟内存管理时，地址和内存的概念是不同的，因此应用程序必须划出一个地址范围，该地址范围可以容纳cuMemCreate 所做的内存分配。保留的地址范围必须至少与用户计划放入其中的所有物理内存分配的大小之和一样大。

应用程序可以通过将适当的参数传递给cuMemAddressReserve 来保留虚拟地址范围。获取的地址范围将没有任何与之关联的设备或主机物理内存。保留的虚拟地址范围可以映射到属于系统中任何设备的内存块，从而为应用程序提供由属于不同设备的内存支持和映射的连续 VA 范围。应用程序应使用 cuMemAddressFree将虚拟地址范围返回到 CUDA 。用户必须确保在调用 cuMemAddressFree之前未映射整个 VA 范围。这些函数在概念上类似于mmap/munmap（在Linux上）或VirtualAlloc/VirtualFree（在Windows上）函数。以下代码片段说明了该函数的用法：

CUdeviceptr ptr;

// `ptr` holds the returned start of virtual address range reserved.

CUresult result = cuMemAddressReserve(&ptr, size, 0, 0, 0); // alignment = 0 for default alignment

10.5. 虚拟别名支持

虚拟内存管理 API 提供了一种创建多个虚拟内存映射或“代理”到同一分配的方法，使用具有不同虚拟地址的多个调用cuMemMap，即所谓的虚拟别名。除非 PTX ISA 中另有说明，否则在写入设备操作（网格启动、memcpy、memset 等）完成之前，对分配的一个代理的写入被视为与同一内存的任何其他代理不一致且不一致。在写入设备操作之前存在于 GPU 上的网格，但在写入设备操作完成后读取的网格也被视为具有不一致和不连贯的代理。

例如，假设设备指针 A 和 B 是相同内存分配的虚拟别名，以下代码片段被视为未定义：

__global__ void foo(char *A, char *B) {

  *A = 0x1;

  printf("%d\n", *B);    // Undefined behavior!  *B can take on either

// the previous value or some value in-between.

}

以下是定义的行为，假设这两个内核是单调排序的（按流或事件）。

__global__ void foo1(char *A) {

  *A = 0x1;

}

__global__ void foo2(char *B) {

  printf("%d\n", *B);    // *B == *A == 0x1 assuming foo2 waits for foo1

// to complete before launching

}

cudaMemcpyAsync(B, input, size, stream1);    // Aliases are allowed at

// operation boundaries

foo1<<<1,1,0,stream1>>>(A);                  // allowing foo1 to access A.

cudaEventRecord(event, stream1);

cudaStreamWaitEvent(stream2, event);

foo2<<<1,1,0,stream2>>>(B);

cudaStreamWaitEvent(stream3, event);

cudaMemcpyAsync(output, B, size, stream3);  // Both launches of foo2 and

                                            // cudaMemcpy (which both

                                            // read) wait for foo1 (which writes)

                                            // to complete before proceeding

10.6. 映射内存

前两节中分配的物理内存和划分的虚拟地址空间表示虚拟内存管理 API 引入的内存和地址区别。要使分配的内存可用，用户必须首先将内存放在地址空间中。从cuMemAddressReserve 获取的地址范围cuMemCreate和cuMemImportFromShareableHandle从中获取的物理分配，或者必须使用 cuMemMap相互关联。

用户可以关联来自多个设备的分配以驻留在连续的虚拟地址范围内，只要他们开辟了足够的地址空间。为了分离物理分配和地址范围，用户必须使用cuMemUnmap 取消映射映射的地址。用户可以根据需要多次映射和取消将内存映射到同一地址范围，只要他们确保不尝试在已映射的 VA 范围预留上创建映射即可。以下代码片段说明了该函数的用法：

CUdeviceptr ptr;

// `ptr`: address in the address range previously reserved by cuMemAddressReserve.

// `allocHandle`: CUmemGenericAllocationHandle obtained by a previous call to cuMemCreate.

CUresult result = cuMemMap(ptr, size, 0, allocHandle, 0);

10.7. 控制访问权限

虚拟内存管理 API 使应用程序能够通过访问控制机制显式保护其 VA 范围。使用cuMemMap 将分配映射到地址范围的区域不会使地址可访问，如果由 CUDA 内核访问，则会导致程序崩溃。用户必须使用cuMemSetAccess该功能专门选择访问控制，该功能允许或限制特定设备对映射地址范围的访问。以下代码片段说明了该函数的用法：

void setAccessOnDevice(int device, CUdeviceptr ptr, size_t size) {

    CUmemAccessDesc accessDesc = {};

    accessDesc.location.type = CU_MEM_LOCATION_TYPE_DEVICE;

    accessDesc.location.id = device;

    accessDesc.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE;

    // Make the address accessible

    cuMemSetAccess(ptr, size, &accessDesc, 1);

}

虚拟内存管理公开的访问控制机制允许用户明确表示要与系统上的其他对等设备共享哪些分配。如前所述，cudaEnablePeerAccess强制将所有先前和将来的 cudaMalloc'd 分配映射到目标对等设备。这在许多情况下可能很方便，因为用户不必担心跟踪系统中每个设备的每个分配的映射状态。但对于关注其应用程序性能的用户，此方法具有性能影响。通过分配粒度的访问控制，虚拟内存管理公开了一种机制，以最小的开销进行对等映射。

该示例可用作使用vectorAddMMAP虚拟内存管理 API 的示例。