F.1. Introduction

使用 cudaMalloc 和 cudaFree 管理內(nèi)存分配會導(dǎo)致 GPU 在所有正在執(zhí)行的 CUDA 流之間進(jìn)行同步。 Stream Order Memory Allocator 使應(yīng)用程序能夠通過啟動到 CUDA 流中的其他工作（例如內(nèi)核啟動和異步拷貝）來對內(nèi)存分配和釋放進(jìn)行排序。這通過利用流排序語義來重用內(nèi)存分配來改進(jìn)應(yīng)用程序內(nèi)存使用。分配器還允許應(yīng)用程序控制分配器的內(nèi)存緩存行為。當(dāng)設(shè)置了適當(dāng)?shù)尼尫砰撝禃r，緩存行為允許分配器在應(yīng)用程序表明它愿意接受更大的內(nèi)存占用時避免對操作系統(tǒng)進(jìn)行昂貴的調(diào)用。分配器還支持在進(jìn)程之間輕松安全地共享分配。

對于許多應(yīng)用程序，Stream Ordered Memory Allocator 減少了對自定義內(nèi)存管理抽象的需求，并使為需要它的應(yīng)用程序創(chuàng)建高性能自定義內(nèi)存管理變得更加容易。對于已經(jīng)具有自定義內(nèi)存分配器的應(yīng)用程序和庫，采用 Stream Ordered Memory Allocator 可以使多個庫共享由驅(qū)動程序管理的公共內(nèi)存池，從而減少過多的內(nèi)存消耗。此外，驅(qū)動程序可以根據(jù)其對分配器和其他流管理 API 的感知執(zhí)行優(yōu)化。最后，Nsight Compute 和 Next-Gen CUDA 調(diào)試器知道分配器是其 CUDA 11.3 工具包支持的一部分。

F.2. Query for Support

用戶可以通過使用設(shè)備屬性 cudaDevAttrMemoryPoolsSupported 調(diào)用 cudaDeviceGetAttribute（） 來確定設(shè)備是否支持流序內(nèi)存分配器。

從 CUDA 11.3 開始，可以使用 cudaDevAttrMemoryPoolSupportedHandleTypes 設(shè)備屬性查詢 IPC 內(nèi)存池支持。 以前的驅(qū)動程序?qū)⒎祷?cudaErrorInvalidValue，因為這些驅(qū)動程序不知道屬性枚舉。

int driverVersion = 0;
int deviceSupportsMemoryPools = 0;
int poolSupportedHandleTypes = 0;
cudaDriverGetVersion(&driverVersion);
if (driverVersion >= 11020) {
    cudaDeviceGetAttribute(&deviceSupportsMemoryPools,
                           cudaDevAttrMemoryPoolsSupported, device);
}
if (deviceSupportsMemoryPools != 0) {
    // `device` supports the Stream Ordered Memory Allocator
}

if (driverVersion >= 11030) {
    cudaDeviceGetAttribute(&poolSupportedHandleTypes,
              cudaDevAttrMemoryPoolSupportedHandleTypes, device);
}
if (poolSupportedHandleTypes & cudaMemHandleTypePosixFileDescriptor) {
   // Pools on the specified device can be created with posix file descriptor-based IPC
}

在查詢之前執(zhí)行驅(qū)動程序版本檢查可避免在尚未定義屬性的驅(qū)動程序上遇到 cudaErrorInvalidValue 錯誤。 可以使用 cudaGetLastError 來清除錯誤而不是避免它。

F.3. API Fundamentals （cudaMallocAsync and cudaFreeAsync）

API cudaMallocAsync 和 cudaFreeAsync 構(gòu)成了分配器的核心。 cudaMallocAsync 返回分配，cudaFreeAsync 釋放分配。 兩個 API 都接受流參數(shù)來定義分配何時變?yōu)榭捎煤屯Ｖ箍捎谩?cudaMallocAsync 返回的指針值是同步確定的，可用于構(gòu)建未來的工作。 重要的是要注意 cudaMallocAsync 在確定分配的位置時會忽略當(dāng)前設(shè)備/上下文。 相反，cudaMallocAsync 根據(jù)指定的內(nèi)存池或提供的流來確定常駐設(shè)備。 最簡單的使用模式是分配、使用和釋放內(nèi)存到同一個流中。

void *ptr;
size_t size = 512;
cudaMallocAsync(&ptr, size, cudaStreamPerThread);
// do work using the allocation
kernel<<<..., cudaStreamPerThread>>>(ptr, ...);
// An asynchronous free can be specified without synchronizing the CPU and GPU
cudaFreeAsync(ptr, cudaStreamPerThread);

用戶可以使用复制cudaFreeAsync()釋放使用复制cudaMalloc()分配的內(nèi)存。 在自由操作開始之前，用戶必須對訪問完成做出同樣的保證。

cudaMalloc(&ptr, size);
kernel<<<..., stream>>>(ptr, ...);
cudaFreeAsync(ptr, stream);

用戶可以使用复制cudaFree()釋放使用复制cudaMallocAsync分配的內(nèi)存。 通過复制cudaFree()API 釋放此類分配時，驅(qū)動程序假定對分配的所有訪問都已完成，并且不執(zhí)行進(jìn)一步的同步。 用戶可以使用复制cudaStreamQuery / cudaStreamSynchronize / cudaEventQuery / cudaEventSynchronize / cudaDeviceSynchronize來保證適當(dāng)?shù)漠惒焦ぷ魍瓿刹⑶褿PU不會嘗試訪問分配。

cudaMallocAsync(&ptr, size,stream);
kernel<<<..., stream>>>(ptr, ...);
// synchronize is needed to avoid prematurely freeing the memory
cudaStreamSynchronize(stream);
cudaFree(ptr);

F.4. Memory Pools and the cudaMemPool_t

內(nèi)存池封裝了虛擬地址和物理內(nèi)存資源，根據(jù)內(nèi)存池的屬性和屬性進(jìn)行分配和管理。內(nèi)存池的主要方面是它所管理的內(nèi)存的種類和位置。

所有對 cudaMallocAsync 的調(diào)用都使用內(nèi)存池的資源。在沒有指定內(nèi)存池的情況下，cudaMallocAsync API 使用提供的流設(shè)備的當(dāng)前內(nèi)存池。設(shè)備的當(dāng)前內(nèi)存池可以使用 cudaDeviceSetMempool 設(shè)置并使用 cudaDeviceGetMempool 查詢。默認(rèn)情況下（在沒有 cudaDeviceSetMempool 調(diào)用的情況下），當(dāng)前內(nèi)存池是設(shè)備的默認(rèn)內(nèi)存池。 cudaMallocFromPoolAsync 的 API cudaMallocFromPoolAsync 和 c++ 重載允許用戶指定要用于分配的池，而無需將其設(shè)置為當(dāng)前池。 API cudaDeviceGetDefaultMempool 和 cudaMemPoolCreate 為用戶提供內(nèi)存池的句柄。

注意：設(shè)備的內(nèi)存池當(dāng)前將是該設(shè)備的本地。因此，在不指定內(nèi)存池的情況下進(jìn)行分配將始終產(chǎn)生流設(shè)備本地的分配。

注意：cudaMemPoolSetAttribute 和 cudaMemPoolGetAttribute 控制內(nèi)存池的屬性。

F.5. Default/Impicit Pools

可以使用 cudaDeviceGetDefaultMempool API 檢索設(shè)備的默認(rèn)內(nèi)存池。 來自設(shè)備默認(rèn)內(nèi)存池的分配是位于該設(shè)備上的不可遷移設(shè)備分配。 這些分配將始終可以從該設(shè)備訪問。 默認(rèn)內(nèi)存池的可訪問性可以通過 cudaMemPoolSetAccess 進(jìn)行修改，并通過 cudaMemPoolGetAccess 進(jìn)行查詢。 由于不需要顯式創(chuàng)建默認(rèn)池，因此有時將它們稱為隱式池。 設(shè)備默認(rèn)內(nèi)存池不支持IPC。

F.6. Explicit Pools

API cudaMemPoolCreate 創(chuàng)建一個顯式池。 目前內(nèi)存池只能分配設(shè)備分配。 分配將駐留的設(shè)備必須在屬性結(jié)構(gòu)中指定。 顯式池的主要用例是 IPC 功能。

// create a pool similar to the implicit pool on device 0
int device = 0;
cudaMemPoolProps poolProps = { };
poolProps.allocType = cudaMemAllocationTypePinned;
poolProps.location.id = device;
poolProps.location.type = cudaMemLocationTypeDevice;

cudaMemPoolCreate(&memPool, &poolProps));

F.7. Physical Page Caching Behavior

默認(rèn)情況下，分配器嘗試最小化池?fù)碛械奈锢韮?nèi)存。 為了盡量減少分配和釋放物理內(nèi)存的操作系統(tǒng)調(diào)用，應(yīng)用程序必須為每個池配置內(nèi)存占用。 應(yīng)用程序可以使用釋放閾值屬性 （cudaMemPoolAttrReleaseThreshold） 執(zhí)行此操作。

釋放閾值是池在嘗試將內(nèi)存釋放回操作系統(tǒng)之前應(yīng)保留的內(nèi)存量（以字節(jié)為單位）。 當(dāng)內(nèi)存池持有超過釋放閾值字節(jié)的內(nèi)存時，分配器將嘗試在下一次調(diào)用流、事件或設(shè)備同步時將內(nèi)存釋放回操作系統(tǒng)。 將釋放閾值設(shè)置為 UINT64_MAX 將防止驅(qū)動程序在每次同步后嘗試收縮池。

Cuuint64_t setVal = UINT64_MAX;
cudaMemPoolSetAttribute(memPool, cudaMemPoolAttrReleaseThreshold, &setVal);

將复制cudaMemPoolAttrReleaseThreshold設(shè)置得足夠高以有效禁用內(nèi)存池收縮的應(yīng)用程序可能希望顯式收縮內(nèi)存池的內(nèi)存占用。复制cudaMemPoolTrimTo允許此類應(yīng)用程序這樣做。 在修剪內(nèi)存池的占用空間時，复制minBytesToKeep參數(shù)允許應(yīng)用程序保留它預(yù)期在后續(xù)執(zhí)行階段需要的內(nèi)存量。

Cuuint64_t setVal = UINT64_MAX;
cudaMemPoolSetAttribute(memPool, cudaMemPoolAttrReleaseThreshold, &setVal);

// application phase needing a lot of memory from the stream ordered allocator
for (i=0; i<10; i++) {
    for (j=0; j<10; j++) {
        cudaMallocAsync(&ptrs[j],size[j], stream);
    }
    kernel<<<...,stream>>>(ptrs,...);
    for (j=0; j<10; j++) {
        cudaFreeAsync(ptrs[j], stream);
    }
}

// Process does not need as much memory for the next phase.
// Synchronize so that the trim operation will know that the allocations are no 
// longer in use.
cudaStreamSynchronize(stream);
cudaMemPoolTrimTo(mempool, 0);

// Some other process/allocation mechanism can now use the physical memory 
// released by the trimming operation.

F.8. Resource Usage Statistics

在 CUDA 11.3 中，添加了池屬性 cudaMemPoolAttrReservedMemCurrent、cudaMemPoolAttrReservedMemHigh、cudaMemPoolAttrUsedMemCurrent 和 cudaMemPoolAttrUsedMemHigh 來查詢池的內(nèi)存使用情況。

查詢池的 cudaMemPoolAttrReservedMemCurrent 屬性會報告該池當(dāng)前消耗的總物理 GPU 內(nèi)存。 查詢池的 cudaMemPoolAttrUsedMemCurrent 會返回從池中分配且不可重用的所有內(nèi)存的總大小。

cudaMemPoolAttr*MemHigh 屬性是記錄自上次重置以來各個 cudaMemPoolAttr*MemCurrent 屬性達(dá)到的最大值的水印。 可以使用 cudaMemPoolSetAttribute API 將它們重置為當(dāng)前值。

// sample helper functions for getting the usage statistics in bulk
struct usageStatistics {
    cuuint64_t reserved;
    cuuint64_t reservedHigh;
    cuuint64_t used;
    cuuint64_t usedHigh;
};

void getUsageStatistics(cudaMemoryPool_t memPool, struct usageStatistics *statistics)
{
    cudaMemPoolGetAttribute(memPool, cudaMemPoolAttrReservedMemCurrent, statistics->reserved);
    cudaMemPoolGetAttribute(memPool, cudaMemPoolAttrReservedMemHigh, statistics->reservedHigh);
    cudaMemPoolGetAttribute(memPool, cudaMemPoolAttrUsedMemCurrent, statistics->used);
    cudaMemPoolGetAttribute(memPool, cudaMemPoolAttrUsedMemHigh, statistics->usedHigh);
}


// resetting the watermarks will make them take on the current value.
void resetStatistics(cudaMemoryPool_t memPool)
{
    cuuint64_t value = 0;
    cudaMemPoolSetAttribute(memPool, cudaMemPoolAttrReservedMemHigh, &value);
    cudaMemPoolSetAttribute(memPool, cudaMemPoolAttrUsedMemHigh, &value);
}

F.9. Memory Reuse Policies

為了服務(wù)分配請求，驅(qū)動程序在嘗試從操作系統(tǒng)分配更多內(nèi)存之前嘗試重用之前通過 cudaFreeAsync（） 釋放的內(nèi)存。 例如，流中釋放的內(nèi)存可以立即重新用于同一流中的后續(xù)分配請求。 類似地，當(dāng)一個流與 CPU 同步時，之前在該流中釋放的內(nèi)存可以重新用于任何流中的分配。

流序分配器有一些可控的分配策略。 池屬性 cudaMemPoolReuseFollowEventDependencies、cudaMemPoolReuseAllowOpportunistic 和 cudaMemPoolReuseAllowInternalDependencies 控制這些策略。 升級到更新的 CUDA 驅(qū)動程序可能會更改、增強(qiáng)、增加或重新排序重用策略。

F.9.1. cudaMemPoolReuseFollowEventDependencies

在分配更多物理 GPU 內(nèi)存之前，分配器會檢查由 CUDA 事件建立的依賴信息，并嘗試從另一個流中釋放的內(nèi)存中進(jìn)行分配。

cudaMallocAsync(&ptr, size, originalStream);
kernel<<<..., originalStream>>>(ptr, ...);
cudaFreeAsync(ptr, originalStream);
cudaEventRecord(event,originalStream);

// waiting on the event that captures the free in another stream 
// allows the allocator to reuse the memory to satisfy 
// a new allocation request in the other stream when
// cudaMemPoolReuseFollowEventDependencies is enabled.
cudaStreamWaitEvent(otherStream, event);
cudaMallocAsync(&ptr2, size, otherStream);

F.9.2. cudaMemPoolReuseAllowOpportunistic

根據(jù) cudaMemPoolReuseAllowOpportunistic 策略，分配器檢查釋放的分配以查看是否滿足釋放的流序語義（即流已通過釋放指示的執(zhí)行點）。 禁用此功能后，分配器仍將重用在流與 cpu 同步時可用的內(nèi)存。 禁用此策略不會阻止 cudaMemPoolReuseFollowEventDependencies 應(yīng)用。

cudaMallocAsync(&ptr, size, originalStream);
kernel<<<..., originalStream>>>(ptr, ...);
cudaFreeAsync(ptr, originalStream);


// after some time, the kernel finishes running
wait(10);

// When cudaMemPoolReuseAllowOpportunistic is enabled this allocation request
// can be fulfilled with the prior allocation based on the progress of originalStream.
cudaMallocAsync(&ptr2, size, otherStream);

F.9.3. cudaMemPoolReuseAllowInternalDependencies

如果無法從操作系統(tǒng)分配和映射更多物理內(nèi)存，驅(qū)動程序?qū)ふ移淇捎眯匀Q于另一個流的待處理進(jìn)度的內(nèi)存。 如果找到這樣的內(nèi)存，驅(qū)動程序會將所需的依賴項插入分配流并重用內(nèi)存。

cudaMallocAsync(&ptr, size, originalStream);
kernel<<<..., originalStream>>>(ptr, ...);
cudaFreeAsync(ptr, originalStream);

// When cudaMemPoolReuseAllowInternalDependencies is enabled
// and the driver fails to allocate more physical memory, the driver may
// effectively perform a cudaStreamWaitEvent in the allocating stream
// to make sure that future work in ‘otherStream’ happens after the work
// in the original stream that would be allowed to access the original allocation. 
cudaMallocAsync(&ptr2, size, otherStream);

F.9.4. Disabling Reuse Policies

雖然可控重用策略提高了內(nèi)存重用，但用戶可能希望禁用它們。 允許機(jī)會重用（即 cudaMemPoolReuseAllowOpportunistic）基于 CPU 和 GPU 執(zhí)行的交錯引入了運行到運行分配模式的差異。 當(dāng)用戶寧愿在分配失敗時顯式同步事件或流時，內(nèi)部依賴插入（即 cudaMemPoolReuseAllowInternalDependencies）可以以意想不到的和潛在的非確定性方式序列化工作。

F.10. Device Accessibility for Multi-GPU Support

就像通過虛擬內(nèi)存管理 API 控制的分配可訪問性一樣，內(nèi)存池分配可訪問性不遵循 cudaDeviceEnablePeerAccess 或 cuCtxEnablePeerAccess。相反，API cudaMemPoolSetAccess 修改了哪些設(shè)備可以訪問池中的分配。默認(rèn)情況下，可以從分配所在的設(shè)備訪問分配。無法撤銷此訪問權(quán)限。要啟用其他設(shè)備的訪問，訪問設(shè)備必須與內(nèi)存池的設(shè)備對等；檢查 cudaDeviceCanAccessPeer。如果未檢查對等功能，則設(shè)置訪問可能會失敗并顯示 cudaErrorInvalidDevice。如果沒有從池中進(jìn)行分配，即使設(shè)備不具備對等能力，cudaMemPoolSetAccess 調(diào)用也可能成功；在這種情況下，池中的下一次分配將失敗。

值得注意的是，cudaMemPoolSetAccess 會影響內(nèi)存池中的所有分配，而不僅僅是未來的分配。此外，cudaMemPoolGetAccess 報告的可訪問性適用于池中的所有分配，而不僅僅是未來的分配。建議不要頻繁更改給定 GPU 的池的可訪問性設(shè)置；一旦池可以從給定的 GPU 訪問，它應(yīng)該在池的整個生命周期內(nèi)都可以從該 GPU 訪問。

// snippet showing usage of cudaMemPoolSetAccess:
cudaError_t setAccessOnDevice(cudaMemPool_t memPool, int residentDevice,
              int accessingDevice) {
    cudaMemAccessDesc accessDesc = {};
    accessDesc.location.type = cudaMemLocationTypeDevice;
    accessDesc.location.id = accessingDevice;
    accessDesc.flags = cudaMemAccessFlagsProtReadWrite;

    int canAccess = 0;
    cudaError_t error = cudaDeviceCanAccessPeer(&canAccess, accessingDevice,
              residentDevice);
    if (error != cudaSuccess) {
        return error;
    } else if (canAccess == 0) {
        return cudaErrorPeerAccessUnsupported;
    }

    // Make the address accessible
    return cudaMemPoolSetAccess(memPool, &accessDesc, 1);
}

F.11. IPC Memory Pools

支持 IPC 的內(nèi)存池允許在進(jìn)程之間輕松、高效和安全地共享 GPU 內(nèi)存。 CUDA 的 IPC 內(nèi)存池提供與 CUDA 的虛擬內(nèi)存管理 API 相同的安全優(yōu)勢。

在具有內(nèi)存池的進(jìn)程之間共享內(nèi)存有兩個階段。 進(jìn)程首先需要共享對池的訪問權(quán)限，然后共享來自該池的特定分配。 第一階段建立并實施安全性。 第二階段協(xié)調(diào)每個進(jìn)程中使用的虛擬地址以及映射何時需要在導(dǎo)入過程中有效。

F.11.1. Creating and Sharing IPC Memory Pools

共享對池的訪問涉及檢索池的 OS 本機(jī)句柄（使用 cudaMemPoolExportToShareableHandle（） API），使用通常的 OS 本機(jī) IPC 機(jī)制將句柄轉(zhuǎn)移到導(dǎo)入進(jìn)程，并創(chuàng)建導(dǎo)入的內(nèi)存池（使用 cudaMemPoolImportFromShareableHandle（） API）。 要使 cudaMemPoolExportToShareableHandle 成功，必須使用池屬性結(jié)構(gòu)中指定的請求句柄類型創(chuàng)建內(nèi)存池。 請參考示例以了解在進(jìn)程之間傳輸操作系統(tǒng)本機(jī)句柄的適當(dāng) IPC 機(jī)制。 該過程的其余部分可以在以下代碼片段中找到。

// in exporting process
// create an exportable IPC capable pool on device 0
cudaMemPoolProps poolProps = { };
poolProps.allocType = cudaMemAllocationTypePinned;
poolProps.location.id = 0;
poolProps.location.type = cudaMemLocationTypeDevice;

// Setting handleTypes to a non zero value will make the pool exportable (IPC capable)
poolProps.handleTypes = CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR;

cudaMemPoolCreate(&memPool, &poolProps));

// FD based handles are integer types
int fdHandle = 0;


// Retrieve an OS native handle to the pool.
// Note that a pointer to the handle memory is passed in here.
cudaMemPoolExportToShareableHandle(&fdHandle,
             memPool,
             CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR,
             0);

// The handle must be sent to the importing process with the appropriate
// OS specific APIs.

// in importing process
 int fdHandle;
// The handle needs to be retrieved from the exporting process with the
// appropriate OS specific APIs.
// Create an imported pool from the shareable handle.
// Note that the handle is passed by value here. 
cudaMemPoolImportFromShareableHandle(&importedMemPool,
          (void*)fdHandle,
          CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR,
          0);

F.11.2. Set Access in the Importing Process

導(dǎo)入的內(nèi)存池最初只能從其常駐設(shè)備訪問。 導(dǎo)入的內(nèi)存池不繼承導(dǎo)出進(jìn)程設(shè)置的任何可訪問性。 導(dǎo)入過程需要啟用從它計劃訪問內(nèi)存的任何 GPU 的訪問（使用 cudaMemPoolSetAccess）。

如果導(dǎo)入的內(nèi)存池在導(dǎo)入過程中屬于不可見的設(shè)備，則用戶必須使用 cudaMemPoolSetAccess API 來啟用從將使用分配的 GPU 的訪問。

F.11.3. Creating and Sharing Allocations from an Exported Pool

共享池后，在導(dǎo)出進(jìn)程中使用 cudaMallocAsync（） 從池中進(jìn)行的分配可以與已導(dǎo)入池的其他進(jìn)程共享。由于池的安全策略是在池級別建立和驗證的，操作系統(tǒng)不需要額外的簿記來為特定的池分配提供安全性；換句話說，導(dǎo)入池分配所需的不透明 cudaMemPoolPtrExportData 可以使用任何機(jī)制發(fā)送到導(dǎo)入進(jìn)程。

雖然分配可以在不以任何方式與分配流同步的情況下導(dǎo)出甚至導(dǎo)入，但在訪問分配時，導(dǎo)入過程必須遵循與導(dǎo)出過程相同的規(guī)則。即，對分配的訪問必須發(fā)生在分配流中分配操作的流排序之后。以下兩個代碼片段顯示 cudaMemPoolExportPointer（） 和 cudaMemPoolImportPointer（） 與 IPC 事件共享分配，用于保證在分配準(zhǔn)備好之前在導(dǎo)入過程中不會訪問分配。

// preparing an allocation in the exporting process
cudaMemPoolPtrExportData exportData;
cudaEvent_t readyIpcEvent;
cudaIpcEventHandle_t readyIpcEventHandle;

// IPC event for coordinating between processes
// cudaEventInterprocess flag makes the event an IPC event
// cudaEventDisableTiming  is set for performance reasons

cudaEventCreate(
        &readyIpcEvent, cudaEventDisableTiming | cudaEventInterprocess)

// allocate from the exporting mem pool
cudaMallocAsync(&ptr, size,exportMemPool, stream);

// event for sharing when the allocation is ready.
cudaEventRecord(readyIpcEvent, stream);
cudaMemPoolExportPointer(&exportData, ptr);
cudaIpcGetEventHandle(&readyIpcEventHandle, readyIpcEvent);

// Share IPC event and pointer export data with the importing process using
//  any mechanism. Here we copy the data into shared memory
shmem->ptrData = exportData;
shmem->readyIpcEventHandle = readyIpcEventHandle;
// signal consumers data is ready

// Importing an allocation
cudaMemPoolPtrExportData *importData = &shmem->prtData;
cudaEvent_t readyIpcEvent;
cudaIpcEventHandle_t *readyIpcEventHandle = &shmem->readyIpcEventHandle;

// Need to retrieve the IPC event handle and the export data from the
// exporting process using any mechanism.  Here we are using shmem and just
// need synchronization to make sure the shared memory is filled in.

cudaIpcOpenEventHandle(&readyIpcEvent, readyIpcEventHandle);

// import the allocation. The operation does not block on the allocation being ready.
cudaMemPoolImportPointer(&ptr, importedMemPool, importData);

// Wait for the prior stream operations in the allocating stream to complete before
// using the allocation in the importing process.
cudaStreamWaitEvent(stream, readyIpcEvent);
kernel<<<..., stream>>>(ptr, ...);

釋放分配時，需要先在導(dǎo)入過程中釋放分配，然后在導(dǎo)出過程中釋放分配。 以下代碼片段演示了使用 CUDA IPC 事件在兩個進(jìn)程中的复制cudaFreeAsync操作之間提供所需的同步。 導(dǎo)入過程中對分配的訪問顯然受到導(dǎo)入過程側(cè)的自由操作的限制。 值得注意的是，复制cudaFree可用于釋放兩個進(jìn)程中的分配，并且可以使用其他流同步 API 代替 CUDA IPC 事件。

// The free must happen in importing process before the exporting process
kernel<<<..., stream>>>(ptr, ...); 

// Last access in importing process
cudaFreeAsync(ptr, stream); 

// Access not allowed in the importing process after the free
cudaIpcEventRecord(finishedIpcEvent, stream);

// Exporting process
// The exporting process needs to coordinate its free with the stream order 
// of the importing process’s free.
cudaStreamWaitEvent(stream, finishedIpcEvent);
kernel<<<..., stream>>>(ptrInExportingProcess, ...); 

// The free in the importing process doesn’t stop the exporting process 
// from using the allocation.
cudFreeAsync(ptrInExportingProcess,stream);

F.11.4. IPC Export Pool Limitations

IPC 池目前不支持將物理塊釋放回操作系統(tǒng)。 因此，cudaMemPoolTrimTo API 充當(dāng)空操作，并且 cudaMemPoolAttrReleaseThreshold 被有效地忽略。 此行為由驅(qū)動程序控制，而不是運行時控制，并且可能會在未來的驅(qū)動程序更新中發(fā)生變化。

F.11.5. IPC Import Pool Limitations

不允許從導(dǎo)入池中分配； 具體來說，導(dǎo)入池不能設(shè)置為當(dāng)前，也不能在 cudaMallocFromPoolAsync API 中使用。 因此，分配重用策略屬性對這些池沒有意義。

IPC 池目前不支持將物理塊釋放回操作系統(tǒng)。 因此，cudaMemPoolTrimTo API 充當(dāng)空操作，并且 cudaMemPoolAttrReleaseThreshold 被有效地忽略。

資源使用統(tǒng)計屬性查詢僅反映導(dǎo)入進(jìn)程的分配和相關(guān)的物理內(nèi)存。

F.12. Synchronization API Actions

作為 CUDA 驅(qū)動程序一部分的分配器帶來的優(yōu)化之一是與同步 API 的集成。 當(dāng)用戶請求 CUDA 驅(qū)動程序同步時，驅(qū)動程序等待異步工作完成。 在返回之前，驅(qū)動程序?qū)⒋_定什么釋放了保證完成的同步。 無論指定的流或禁用的分配策略如何，這些分配都可用于分配。 驅(qū)動程序還在這里檢查 cudaMemPoolAttrReleaseThreshold 并釋放它可以釋放的任何多余的物理內(nèi)存。

F.13. Addendums

F.13.1. cudaMemcpyAsync Current Context/Device Sensitivity

在當(dāng)前的 CUDA 驅(qū)動程序中，任何涉及來自 cudaMallocAsync 的內(nèi)存的異步 memcpy 都應(yīng)該使用指定流的上下文作為調(diào)用線程的當(dāng)前上下文來完成。 這對于 cudaMemcpyPeerAsync 不是必需的，因為引用了 API 中指定的設(shè)備主上下文而不是當(dāng)前上下文。

F.13.2. cuPointerGetAttribute Query

在對分配調(diào)用 cudaFreeAsync 后在分配上調(diào)用 cuPointerGetAttribute 會導(dǎo)致未定義的行為。 具體來說，分配是否仍然可以從給定的流中訪問并不重要：行為仍然是未定義的。

F.13.3. cuGraphAddMemsetNode

cuGraphAddMemsetNode 不適用于通過流排序分配器分配的內(nèi)存。 但是，分配的 memset 可以被流捕獲。

F.13.4. Pointer Attributes

cuPointerGetAttributes 查詢適用于流有序分配。 由于流排序分配與上下文無關(guān)，因此查詢 CU_POINTER_ATTRIBUTE_CONTEXT 將成功，但在 *data 中返回 NULL。 屬性 CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL 可用于確定分配的位置：這在選擇使用 cudaMemcpyPeerAsync 制作 p2h2p 拷貝的上下文時很有用。 CU_POINTER_ATTRIBUTE_MEMPOOL_HANDLE 屬性是在 CUDA 11.3 中添加的，可用于調(diào)試和在執(zhí)行 IPC 之前確認(rèn)分配來自哪個池。

關(guān)于作者

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

審核編輯：郭婷