E.1. Introduction

虛擬內(nèi)存管理 API 為應(yīng)用程序提供了一種直接管理統(tǒng)一虛擬地址空間的方法，該空間由 CUDA 提供，用于將物理內(nèi)存映射到 GPU 可訪問(wèn)的虛擬地址。在 CUDA 10.2 中引入的這些 API 還提供了一種與其他進(jìn)程和圖形 API（如 OpenGL 和 Vulkan）進(jìn)行互操作的新方法，并提供了用戶可以調(diào)整以適應(yīng)其應(yīng)用程序的更新內(nèi)存屬性。

從歷史上看，CUDA 編程模型中的內(nèi)存分配調(diào)用（例如 cudaMalloc）返回了一個(gè)指向 GPU 內(nèi)存的內(nèi)存地址。這樣獲得的地址可以與任何 CUDA API 一起使用，也可以在設(shè)備內(nèi)核中使用。但是，分配的內(nèi)存無(wú)法根據(jù)用戶的內(nèi)存需求調(diào)整大小。為了增加分配的大小，用戶必須顯式分配更大的緩沖區(qū)，從初始分配中復(fù)制數(shù)據(jù)，釋放它，然后繼續(xù)跟蹤新分配的地址。這通常會(huì)導(dǎo)致應(yīng)用程序的性能降低和峰值內(nèi)存利用率更高。本質(zhì)上，用戶有一個(gè)類似 malloc 的接口來(lái)分配 GPU 內(nèi)存，但沒(méi)有相應(yīng)的 realloc 來(lái)補(bǔ)充它。虛擬內(nèi)存管理 API 將地址和內(nèi)存的概念解耦，并允許應(yīng)用程序分別處理它們。 API 允許應(yīng)用程序在他們認(rèn)為合適的時(shí)候從虛擬地址范圍映射和取消映射內(nèi)存。

在通過(guò) cudaEnablePeerAccess 啟用對(duì)等設(shè)備訪問(wèn)內(nèi)存分配的情況下，所有過(guò)去和未來(lái)的用戶分配都映射到目標(biāo)對(duì)等設(shè)備。這導(dǎo)致用戶無(wú)意中支付了將所有 cudaMalloc 分配映射到對(duì)等設(shè)備的運(yùn)行時(shí)成本。然而，在大多數(shù)情況下，應(yīng)用程序通過(guò)僅與另一個(gè)設(shè)備共享少量分配進(jìn)行通信，并且并非所有分配都需要映射到所有設(shè)備。使用虛擬內(nèi)存管理，應(yīng)用程序可以專門選擇某些分配可從目標(biāo)設(shè)備訪問(wèn)。

CUDA 虛擬內(nèi)存管理 API 向用戶提供細(xì)粒度控制，以管理應(yīng)用程序中的 GPU 內(nèi)存。它提供的 API 允許用戶：

將分配在不同設(shè)備上的內(nèi)存放入一個(gè)連續(xù)的 VA 范圍內(nèi)。

使用平臺(tái)特定機(jī)制執(zhí)行內(nèi)存共享的進(jìn)程間通信。

在支持它們的設(shè)備上選擇更新的內(nèi)存類型。

為了分配內(nèi)存，虛擬內(nèi)存管理編程模型公開(kāi)了以下功能：

分配物理內(nèi)存。

保留 VA 范圍。

將分配的內(nèi)存映射到 VA 范圍。

控制映射范圍的訪問(wèn)權(quán)限。

請(qǐng)注意，本節(jié)中描述的 API 套件需要支持 UVA 的系統(tǒng)。

E.2. Query for support

在嘗試使用虛擬內(nèi)存管理 API 之前，應(yīng)用程序必須確保他們希望使用的設(shè)備支持 CUDA 虛擬內(nèi)存管理。 以下代碼示例顯示了查詢虛擬內(nèi)存管理支持：

int deviceSupportsVmm;
CUresult result = cuDeviceGetAttribute(&deviceSupportsVmm, CU_DEVICE_ATTRIBUTE_VIRTUAL_MEMORY_MANAGEMENT_SUPPORTED, device);
if (deviceSupportsVmm != 0) {
    // `device` supports Virtual Memory Management 
}
   

E.3. Allocating Physical Memory

通過(guò)虛擬內(nèi)存管理 API 進(jìn)行內(nèi)存分配的第一步是創(chuàng)建一個(gè)物理內(nèi)存塊，為分配提供支持。 為了分配物理內(nèi)存，應(yīng)用程序必須使用 cuMemCreate API。 此函數(shù)創(chuàng)建的分配沒(méi)有任何設(shè)備或主機(jī)映射。 函數(shù)參數(shù) CUmemGenericAllocationHandle 描述了要分配的內(nèi)存的屬性，例如分配的位置、分配是否要共享給另一個(gè)進(jìn)程（或其他圖形 API），或者要分配的內(nèi)存的物理屬性。 用戶必須確保請(qǐng)求分配的大小必須與適當(dāng)?shù)牧６葘?duì)齊。 可以使用 cuMemGetAllocationGranularity 查詢有關(guān)分配粒度要求的信息。 以下代碼片段顯示了使用 cuMemCreate 分配物理內(nèi)存：

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 分配的內(nèi)存由它返回的 CUmemGenericAllocationHandle 引用。 這與 cudaMalloc風(fēng)格的分配不同，后者返回一個(gè)指向 GPU 內(nèi)存的指針，該指針可由在設(shè)備上執(zhí)行的 CUDA 內(nèi)核直接訪問(wèn)。 除了使用 cuMemGetAllocationPropertiesFromHandle 查詢屬性之外，分配的內(nèi)存不能用于任何操作。 為了使此內(nèi)存可訪問(wèn)，應(yīng)用程序必須將此內(nèi)存映射到由 cuMemAddressReserve 保留的 VA 范圍，并為其提供適當(dāng)?shù)脑L問(wèn)權(quán)限。 應(yīng)用程序必須使用 cuMemRelease API 釋放分配的內(nèi)存。

E.3.1. Shareable Memory Allocations

使用 cuMemCreate 用戶現(xiàn)在可以在分配時(shí)向 CUDA 指示他們已指定特定分配用于進(jìn)程間通信或圖形互操作目的。應(yīng)用程序可以通過(guò)將 CUmemAllocationProp：：requestedHandleTypes 設(shè)置為平臺(tái)特定字段來(lái)完成此操作。在 Windows 上，當(dāng) CUmemAllocationProp：：requestedHandleTypes 設(shè)置為 CU_MEM_HANDLE_TYPE_WIN32 時(shí)，應(yīng)用程序還必須在 CUmemAllocationProp：：win32HandleMetaData 中指定 LPSECURITYATTRIBUTES 屬性。該安全屬性定義了可以將導(dǎo)出的分配轉(zhuǎn)移到其他進(jìn)程的范圍。

CUDA 虛擬內(nèi)存管理 API 函數(shù)不支持傳統(tǒng)的進(jìn)程間通信函數(shù)及其內(nèi)存。相反，它們公開(kāi)了一種利用操作系統(tǒng)特定句柄的進(jìn)程間通信的新機(jī)制。應(yīng)用程序可以使用 cuMemExportToShareableHandle 獲取與分配相對(duì)應(yīng)的這些操作系統(tǒng)特定句柄。這樣獲得的句柄可以通過(guò)使用通常的 OS 本地機(jī)制進(jìn)行傳輸，以進(jìn)行進(jìn)程間通信。接收進(jìn)程應(yīng)使用 cuMemImportFromShareableHandle 導(dǎo)入分配。

用戶必須確保在嘗試導(dǎo)出使用 cuMemCreate 分配的內(nèi)存之前查詢是否支持請(qǐng)求的句柄類型。以下代碼片段說(shuō)明了以特定平臺(tái)方式查詢句柄類型支持。

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

用戶應(yīng)適當(dāng)設(shè)置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 與虛擬內(nèi)存管理分配一起使用的示例。

E.3.2. Memory Type

在 CUDA 10.2 之前，應(yīng)用程序沒(méi)有用戶控制的方式來(lái)分配某些設(shè)備可能支持的任何特殊類型的內(nèi)存。 使用 cuMemCreate 應(yīng)用程序還可以使用 CUmemAllocationProp：：allocFlags 指定內(nèi)存類型要求，以選擇任何特定的內(nèi)存功能。 應(yīng)用程序還必須確保分配設(shè)備支持請(qǐng)求的內(nèi)存類型。

E.3.2.1. Compressible Memory

可壓縮內(nèi)存可用于加速對(duì)具有非結(jié)構(gòu)化稀疏性和其他可壓縮數(shù)據(jù)模式的數(shù)據(jù)的訪問(wèn)。 壓縮可以節(jié)省 DRAM 帶寬、L2 讀取帶寬和 L2 容量，具體取決于正在操作的數(shù)據(jù)。 想要在支持計(jì)算數(shù)據(jù)壓縮的設(shè)備上分配可壓縮內(nèi)存的應(yīng)用程序可以通過(guò)將 CUmemAllocationProp：：allocFlags：：compressionType 設(shè)置為 CU_MEM_ALLOCATION_COMP_GENERIC 來(lái)實(shí)現(xiàn)。 用戶必須通過(guò) CU_DEVICE_ATTRIBUTE_GENERIC_COMPRESSION_SUPPORTED 查詢?cè)O(shè)備是否支持計(jì)算數(shù)據(jù)壓縮。 以下代碼片段說(shuō)明了查詢可壓縮內(nèi)存支持 cuDeviceGetAttribute。

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

在支持計(jì)算數(shù)據(jù)壓縮的設(shè)備上，用戶需要在分配時(shí)選擇加入，如下所示：

prop.allocFlags.compressionType = CU_MEM_ALLOCATION_COMP_GENERIC;

由于硬件資源有限等各種原因，分配的內(nèi)存可能沒(méi)有壓縮屬性，用戶需要使用cuMemGetAllocationPropertiesFromHandle查詢回分配內(nèi)存的屬性并檢查壓縮屬性。

CUmemAllocationPropPrivate allocationProp = {};
cuMemGetAllocationPropertiesFromHandle(&allocationProp, allocationHandle);

if (allocationProp.allocFlags.compressionType == CU_MEM_ALLOCATION_COMP_GENERIC)
{
    // Obtained compressible memory allocation
}

E.4. Reserving a Virtual Address Range

由于使用虛擬內(nèi)存管理，地址和內(nèi)存的概念是不同的，因此應(yīng)用程序必須劃出一個(gè)地址范圍，以容納由 cuMemCreate 進(jìn)行的內(nèi)存分配。保留的地址范圍必須至少與用戶計(jì)劃放入其中的所有物理內(nèi)存分配大小的總和一樣大。

應(yīng)用程序可以通過(guò)將適當(dāng)?shù)膮?shù)傳遞給 cuMemAddressReserve 來(lái)保留虛擬地址范圍。獲得的地址范圍不會(huì)有任何與之關(guān)聯(lián)的設(shè)備或主機(jī)物理內(nèi)存。保留的虛擬地址范圍可以映射到屬于系統(tǒng)中任何設(shè)備的內(nèi)存塊，從而為應(yīng)用程序提供由屬于不同設(shè)備的內(nèi)存支持和映射的連續(xù) VA 范圍。應(yīng)用程序應(yīng)使用 cuMemAddressFree 將虛擬地址范圍返回給 CUDA。用戶必須確保在調(diào)用 cuMemAddressFree 之前未映射整個(gè) VA 范圍。這些函數(shù)在概念上類似于 mmap/munmap（在 Linux 上）或 VirtualAlloc/VirtualFree（在 Windows 上）函數(shù)。以下代碼片段說(shuō)明了該函數(shù)的用法：

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

虛擬內(nèi)存管理 API 提供了一種創(chuàng)建多個(gè)虛擬內(nèi)存映射或“代理”到相同分配的方法，該方法使用對(duì)具有不同虛擬地址的 cuMemMap 的多次調(diào)用，即所謂的虛擬別名。 除非在 PTX ISA 中另有說(shuō)明，否則寫入分配的一個(gè)代理被認(rèn)為與同一內(nèi)存的任何其他代理不一致和不連貫，直到寫入設(shè)備操作（網(wǎng)格啟動(dòng)、memcpy、memset 等）完成。 在寫入設(shè)備操作之前出現(xiàn)在 GPU 上但在寫入設(shè)備操作完成后讀取的網(wǎng)格也被認(rèn)為具有不一致和不連貫的代理。

例如，下面的代碼片段被認(rèn)為是未定義的，假設(shè)設(shè)備指針 A 和 B 是相同內(nèi)存分配的虛擬別名：

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

以下是定義的行為，假設(shè)這兩個(gè)內(nèi)核是單調(diào)排序的（通過(guò)流或事件）。

__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

E.6. Mapping Memory

前兩節(jié)分配的物理內(nèi)存和挖出的虛擬地址空間代表了虛擬內(nèi)存管理 API 引入的內(nèi)存和地址區(qū)別。為了使分配的內(nèi)存可用，用戶必須首先將內(nèi)存放在地址空間中。從 cuMemAddressReserve 獲取的地址范圍和從 cuMemCreate 或 cuMemImportFromShareableHandle 獲取的物理分配必須通過(guò) cuMemMap 相互關(guān)聯(lián)。

用戶可以關(guān)聯(lián)來(lái)自多個(gè)設(shè)備的分配以駐留在連續(xù)的虛擬地址范圍內(nèi)，只要他們已經(jīng)劃分出足夠的地址空間。為了解耦物理分配和地址范圍，用戶必須通過(guò) cuMemUnmap 取消映射的地址。用戶可以根據(jù)需要多次將內(nèi)存映射和取消映射到同一地址范圍，只要他們確保不會(huì)嘗試在已映射的 VA 范圍保留上創(chuàng)建映射。以下代碼片段說(shuō)明了該函數(shù)的用法：

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

虛擬內(nèi)存管理 API 使應(yīng)用程序能夠通過(guò)訪問(wèn)控制機(jī)制顯式保護(hù)其 VA 范圍。 使用 cuMemMap 將分配映射到地址范圍的區(qū)域不會(huì)使地址可訪問(wèn)，并且如果被 CUDA 內(nèi)核訪問(wèn)會(huì)導(dǎo)致程序崩潰。 用戶必須使用 cuMemSetAccess 函數(shù)專門選擇訪問(wèn)控制，該函數(shù)允許或限制特定設(shè)備對(duì)映射地址范圍的訪問(wèn)。 以下代碼片段說(shuō)明了該函數(shù)的用法：

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

使用虛擬內(nèi)存管理公開(kāi)的訪問(wèn)控制機(jī)制允許用戶明確他們希望與系統(tǒng)上的其他對(duì)等設(shè)備共享哪些分配。 如前所述，cudaEnablePeerAccess 強(qiáng)制將所有先前和將來(lái)的 cudaMalloc 分配映射到目標(biāo)對(duì)等設(shè)備。 這在許多情況下很方便，因?yàn)橛脩舨槐負(fù)?dān)心跟蹤每個(gè)分配到系統(tǒng)中每個(gè)設(shè)備的映射狀態(tài)。 但是對(duì)于關(guān)心其應(yīng)用程序性能的用戶來(lái)說(shuō)，這種方法具有性能影響。 通過(guò)分配粒度的訪問(wèn)控制，虛擬內(nèi)存管理公開(kāi)了一種機(jī)制，可以以最小的開(kāi)銷進(jìn)行對(duì)等映射。

關(guān)于作者

Ken He 是 NVIDIA 企業(yè)級(jí)開(kāi)發(fā)者社區(qū)經(jīng)理 & 高級(jí)講師，擁有多年的 GPU 和人工智能開(kāi)發(fā)經(jīng)驗(yàn)。自 2017 年加入 NVIDIA 開(kāi)發(fā)者社區(qū)以來(lái)，完成過(guò)上百場(chǎng)培訓(xùn)，幫助上萬(wàn)個(gè)開(kāi)發(fā)者了解人工智能和 GPU 編程開(kāi)發(fā)。在計(jì)算機(jī)視覺(jué)，高性能計(jì)算領(lǐng)域完成過(guò)多個(gè)獨(dú)立項(xiàng)目。并且，在機(jī)器人和無(wú)人機(jī)領(lǐng)域，有過(guò)豐富的研發(fā)經(jīng)驗(yàn)。對(duì)于圖像識(shí)別，目標(biāo)的檢測(cè)與跟蹤完成過(guò)多種解決方案。曾經(jīng)參與 GPU 版氣象模式GRAPES，是其主要研發(fā)者。

審核編輯：郭婷