E.1. Introduction

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

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

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

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

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

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

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

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

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

保留 VA 范圍。

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

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

請注意，本節(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

通過虛擬內(nèi)存管理 API 進行內(nèi)存分配的第一步是創(chuàng)建一個物理內(nèi)存塊，為分配提供支持。 為了分配物理內(nèi)存，應(yīng)用程序必須使用 cuMemCreate API。 此函數(shù)創(chuàng)建的分配沒有任何設(shè)備或主機映射。 函數(shù)參數(shù) CUmemGenericAllocationHandle 描述了要分配的內(nèi)存的屬性，例如分配的位置、分配是否要共享給另一個進程（或其他圖形 API），或者要分配的內(nèi)存的物理屬性。 用戶必須確保請求分配的大小必須與適當(dāng)?shù)牧６葘R。 可以使用 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)格的分配不同，后者返回一個指向 GPU 內(nèi)存的指針，該指針可由在設(shè)備上執(zhí)行的 CUDA 內(nèi)核直接訪問。 除了使用 cuMemGetAllocationPropertiesFromHandle 查詢屬性之外，分配的內(nèi)存不能用于任何操作。 為了使此內(nèi)存可訪問，應(yīng)用程序必須將此內(nèi)存映射到由 cuMemAddressReserve 保留的 VA 范圍，并為其提供適當(dāng)?shù)脑L問權(quán)限。 應(yīng)用程序必須使用 cuMemRelease API 釋放分配的內(nèi)存。

E.3.1. Shareable Memory Allocations

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

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

用戶必須確保在嘗試導(dǎo)出使用 cuMemCreate 分配的內(nè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)用程序沒有用戶控制的方式來分配某些設(shè)備可能支持的任何特殊類型的內(nèi)存。 使用 cuMemCreate 應(yīng)用程序還可以使用 CUmemAllocationProp：：allocFlags 指定內(nèi)存類型要求，以選擇任何特定的內(nèi)存功能。 應(yīng)用程序還必須確保分配設(shè)備支持請求的內(nèi)存類型。

E.3.2.1. Compressible Memory

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

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

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

prop.allocFlags.compressionType = CU_MEM_ALLOCATION_COMP_GENERIC;

由于硬件資源有限等各種原因，分配的內(nè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)用程序必須劃出一個地址范圍，以容納由 cuMemCreate 進行的內(nèi)存分配。保留的地址范圍必須至少與用戶計劃放入其中的所有物理內(nèi)存分配大小的總和一樣大。

應(yīng)用程序可以通過將適當(dāng)?shù)膮?shù)傳遞給 cuMemAddressReserve 來保留虛擬地址范圍。獲得的地址范圍不會有任何與之關(guān)聯(lián)的設(shè)備或主機物理內(nèi)存。保留的虛擬地址范圍可以映射到屬于系統(tǒng)中任何設(shè)備的內(nèi)存塊，從而為應(yīng)用程序提供由屬于不同設(shè)備的內(nèi)存支持和映射的連續(xù) VA 范圍。應(yīng)用程序應(yīng)使用 cuMemAddressFree 將虛擬地址范圍返回給 CUDA。用戶必須確保在調(diào)用 cuMemAddressFree 之前未映射整個 VA 范圍。這些函數(shù)在概念上類似于 mmap/munmap（在 Linux 上）或 VirtualAlloc/VirtualFree（在 Windows 上）函數(shù)。以下代碼片段說明了該函數(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)建多個虛擬內(nèi)存映射或“代理”到相同分配的方法，該方法使用對具有不同虛擬地址的 cuMemMap 的多次調(diào)用，即所謂的虛擬別名。 除非在 PTX ISA 中另有說明，否則寫入分配的一個代理被認(rèn)為與同一內(nèi)存的任何其他代理不一致和不連貫，直到寫入設(shè)備操作（網(wǎ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è)這兩個內(nèi)核是單調(diào)排序的（通過流或事件）。

__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 獲取的物理分配必須通過 cuMemMap 相互關(guān)聯(lián)。

用戶可以關(guān)聯(lián)來自多個設(shè)備的分配以駐留在連續(xù)的虛擬地址范圍內(nèi)，只要他們已經(jīng)劃分出足夠的地址空間。為了解耦物理分配和地址范圍，用戶必須通過 cuMemUnmap 取消映射的地址。用戶可以根據(jù)需要多次將內(nèi)存映射和取消映射到同一地址范圍，只要他們確保不會嘗試在已映射的 VA 范圍保留上創(chuàng)建映射。以下代碼片段說明了該函數(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)用程序能夠通過訪問控制機制顯式保護其 VA 范圍。 使用 cuMemMap 將分配映射到地址范圍的區(qū)域不會使地址可訪問，并且如果被 CUDA 內(nèi)核訪問會導(dǎo)致程序崩潰。 用戶必須使用 cuMemSetAccess 函數(shù)專門選擇訪問控制，該函數(shù)允許或限制特定設(shè)備對映射地址范圍的訪問。 以下代碼片段說明了該函數(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)存管理公開的訪問控制機制允許用戶明確他們希望與系統(tǒng)上的其他對等設(shè)備共享哪些分配。 如前所述，cudaEnablePeerAccess 強制將所有先前和將來的 cudaMalloc 分配映射到目標(biāo)對等設(shè)備。 這在許多情況下很方便，因為用戶不必?fù)?dān)心跟蹤每個分配到系統(tǒng)中每個設(shè)備的映射狀態(tài)。 但是對于關(guān)心其應(yīng)用程序性能的用戶來說，這種方法具有性能影響。 通過分配粒度的訪問控制，虛擬內(nèi)存管理公開了一種機制，可以以最小的開銷進行對等映射。

關(guān)于作者

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

審核編輯：郭婷