虚拟内存管理API提供管理统一虚拟地址空间的方法

虚拟内存管理API提供管理统一虚拟地址空间的方法

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

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

将分配在不同设备上的内存放入一个连续的 VA 范围内。

使用平台特定机制执行内存共享的进程间通信。

在支持它们的设备上选择更新的内存类型。

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

分配物理内存。

保留 VA 范围。

将分配的内存映射到 VA 范围。

控制映射范围的访问权限。

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

E.2. Query for support

在尝试使用虚拟内存管理 API 之前，应用程序必须确保他们希望使用的设备支持 CUDA 虚拟内存管理。 以下代码示例显示了查询虚拟内存管理支持：

E.3. Allocating Physical Memory

通过虚拟内存管理 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;}

E.3.1. Shareable Memory Allocations

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

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

用户必须确保在尝试导出使用 cuMemCreate 分配的内存之前查询是否支持请求的句柄类型。以下代码片段说明了以特定平台方式查询句柄类型支持。

用户应适当设置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 与虚拟内存管理分配一起使用的示例。

E.3.2. Memory Type

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

E.3.2.1. Compressible Memory

int compressionSupported = 0;cuDeviceGetAttribute(&compressionSupported, CU_DEVICE_ATTRIBUTE_GENERIC_COMPRESSION_SUPPORTED, device);

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

prop.allocFlags.compressionType = CU_MEM_ALLOCATION_COMP_GENERIC;

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

E.4. Reserving a Virtual Address Range

由于使用虚拟内存管理，地址和内存的概念是不同的，因此应用程序必须划出一个地址范围，以容纳由 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

E.5. Virtual Aliasing Support

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

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

__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 boundariesfoo1<<<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

E.6. Mapping Memory

前两节分配的物理内存和挖出的虚拟地址空间代表了虚拟内存管理 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);

E.7. Control Access Rights

虚拟内存管理 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 分配映射到目标对等设备。 这在许多情况下很方便，因为用户不必担心跟踪每个分配到系统中每个设备的映射状态。 但是对于关心其应用程序性能的用户来说，这种方法具有性能影响。 通过分配粒度的访问控制，虚拟内存管理公开了一种机制，可以以最小的开销进行对等映射。

版权声明：本文内容由网络用户投稿，版权归原作者所有，本站不拥有其著作权，亦不承担相应法律责任。如果您发现本站中有涉嫌抄袭或描述失实的内容，请联系我们jiasou666@gmail.com 处理，核实后本网站将在24小时内删除侵权内容。