CUDA中的流序內(nèi)存分配

1. Introduction

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

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

2. Query for Support

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

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

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ū)動(dòng)程序版本檢查可避免在尚未定義屬性的驅(qū)動(dòng)程序上遇到 cudaErrorInvalidValue 錯(cuò)誤。 可以使用 cudaGetLastError 來(lái)清除錯(cuò)誤而不是避免它。

3. API Fundamentals (cudaMallocAsync and cudaFreeAsync)

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

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)存。 在自由操作開(kāi)始之前，用戶必須對(duì)訪問(wèn)完成做出同樣的保證。

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

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

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

4. Memory Pools and the cudaMemPool_t

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

所有對(duì) cudaMallocAsync 的調(diào)用都使用內(nèi)存池的資源。在沒(méi)有指定內(nèi)存池的情況下，cudaMallocAsync API 使用提供的流設(shè)備的當(dāng)前內(nèi)存池。設(shè)備的當(dāng)前內(nèi)存池可以使用 cudaDeviceSetMempool 設(shè)置并使用 cudaDeviceGetMempool 查詢。默認(rèn)情況下（在沒(méi)有 cudaDeviceSetMempool 調(diào)用的情況下），當(dāng)前內(nèi)存池是設(shè)備的默認(rèn)內(nèi)存池。 cudaMallocFromPoolAsync 的 API cudaMallocFromPoolAsync 和 c++ 重載允許用戶指定要用于分配的池，而無(wú)需將其設(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)存池的屬性。

5. Default/Impicit Pools

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

6. Explicit Pools

API cudaMemPoolCreate 創(chuàng)建一個(gè)顯式池。 目前內(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));

7. Physical Page Caching Behavior

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

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

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

將 cudaMemPoolAttrReleaseThreshold 設(shè)置得足夠高以有效禁用內(nèi)存池收縮的應(yīng)用程序可能希望顯式收縮內(nèi)存池的內(nèi)存占用。 cudaMemPoolTrimTo 允許此類應(yīng)用程序這樣做。 在修剪內(nèi)存池的占用空間時(shí)，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.

8. Resource Usage Statistics

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

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

cudaMemPoolAttr*MemHigh 屬性是記錄自上次重置以來(lái)各個(gè) 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);
}

9. Memory Reuse Policies

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

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

9.1. cudaMemPoolReuseFollowEventDependencies

在分配更多物理 GPU 內(nèi)存之前，分配器會(huì)檢查由 CUDA 事件建立的依賴信息，并嘗試從另一個(gè)流中釋放的內(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);

9.2. cudaMemPoolReuseAllowOpportunistic

根據(jù) cudaMemPoolReuseAllowOpportunistic 策略，分配器檢查釋放的分配以查看是否滿足釋放的流序語(yǔ)義（即流已通過(guò)釋放指示的執(zhí)行點(diǎn)）。 禁用此功能后，分配器仍將重用在流與 cpu 同步時(shí)可用的內(nèi)存。 禁用此策略不會(huì)阻止 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);

9.3. cudaMemPoolReuseAllowInternalDependencies

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

9.4. Disabling Reuse Policies

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

10. Device Accessibility for Multi-GPU Support

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

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

// 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 accessiblereturn cudaMemPoolSetAccess(memPool, &accessDesc, 1);
}

11. IPC Memory Pools

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

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

11.1. Creating and Sharing IPC Memory Pools

共享對(duì)池的訪問(wèn)涉及檢索池的 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)中指定的請(qǐng)求句柄類型創(chuàng)建內(nèi)存池。 請(qǐng)參考示例以了解在進(jìn)程之間傳輸操作系統(tǒng)本機(jī)句柄的適當(dāng) IPC 機(jī)制。 該過(guò)程的其余部分可以在以下代碼片段中找到。

// 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 processint 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);

11.2. Set Access in the Importing Process

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

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

11.3. Creating and Sharing Allocations from an Exported Pool

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

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

// 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 reasonscudaEventCreate(&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, ...);

釋放分配時(shí)，需要先在導(dǎo)入過(guò)程中釋放分配，然后在導(dǎo)出過(guò)程中釋放分配。 以下代碼片段演示了使用 CUDA IPC 事件在兩個(gè)進(jìn)程中的 cudaFreeAsync 操作之間提供所需的同步。 導(dǎo)入過(guò)程中對(duì)分配的訪問(wèn)顯然受到導(dǎo)入過(guò)程側(cè)的自由操作的限制。 值得注意的是，cudaFree 可用于釋放兩個(gè)進(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);

11.4. IPC Export Pool Limitations

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

11.5. IPC Import Pool Limitations

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

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

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

12. Synchronization API Actions

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

13. Addendums

13.1. cudaMemcpyAsync Current Context/Device Sensitivity

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

13.2. cuPointerGetAttribute Query

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

13.3. cuGraphAddMemsetNode

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

13.4. Pointer Attributes

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

更多精彩內(nèi)容:
https://www.nvidia.cn/gtc-global/?ncid=ref-dev-876561