分配器支持在進程間輕松安全地共享分配

F.1. Introduction

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

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

F.2. Query for Support

用戶可以通過使用設備屬性 cudaDevAttrMemoryPoolsSupported 調用 cudaDeviceGetAttribute（） 來確定設備是否支持流序內存分配器。

從 CUDA 11.3 開始，可以使用 cudaDevAttrMemoryPoolSupportedHandleTypes 設備屬性查詢 IPC 內存池支持。 以前的驅動程序將返回 cudaErrorInvalidValue，因為這些驅動程序不知道屬性枚舉。

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í)行驅動程序版本檢查可避免在尚未定義屬性的驅動程序上遇到 cudaErrorInvalidValue 錯誤。 可以使用 cudaGetLastError 來清除錯誤而不是避免它。

F.3. API Fundamentals （cudaMallocAsync and cudaFreeAsync）

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

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()分配的內存。 在自由操作開始之前，用戶必須對訪問完成做出同樣的保證。

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

用戶可以使用cudaFree()釋放使用cudaMallocAsync分配的內存。 通過cudaFree()API 釋放此類分配時，驅動程序假定對分配的所有訪問都已完成，并且不執(zhí)行進一步的同步。 用戶可以使用cudaStreamQuery / cudaStreamSynchronize / cudaEventQuery / cudaEventSynchronize / cudaDeviceSynchronize來保證適當的異步工作完成并且GPU不會嘗試訪問分配。

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

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

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

注意：設備的內存池當前將是該設備的本地。因此，在不指定內存池的情況下進行分配將始終產生流設備本地的分配。

注意：cudaMemPoolSetAttribute 和 cudaMemPoolGetAttribute 控制內存池的屬性。

F.5. Default/Impicit Pools

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

F.6. Explicit Pools

API cudaMemPoolCreate 創(chuàng)建一個顯式池。 目前內存池只能分配設備分配。 分配將駐留的設備必須在屬性結構中指定。 顯式池的主要用例是 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

默認情況下，分配器嘗試最小化池擁有的物理內存。 為了盡量減少分配和釋放物理內存的操作系統(tǒng)調用，應用程序必須為每個池配置內存占用。 應用程序可以使用釋放閾值屬性 （cudaMemPoolAttrReleaseThreshold） 執(zhí)行此操作。

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

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

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

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 來查詢池的內存使用情況。

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

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

// 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

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

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

F.9.1. cudaMemPoolReuseFollowEventDependencies

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

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

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

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)分配和映射更多物理內存，驅動程序將尋找其可用性取決于另一個流的待處理進度的內存。 如果找到這樣的內存，驅動程序會將所需的依賴項插入分配流并重用內存。

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

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

F.10. Device Accessibility for Multi-GPU Support

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

值得注意的是，cudaMemPoolSetAccess 會影響內存池中的所有分配，而不僅僅是未來的分配。此外，cudaMemPoolGetAccess 報告的可訪問性適用于池中的所有分配，而不僅僅是未來的分配。建議不要頻繁更改給定 GPU 的池的可訪問性設置；一旦池可以從給定的 GPU 訪問，它應該在池的整個生命周期內都可以從該 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 的內存池允許在進程之間輕松、高效和安全地共享 GPU 內存。 CUDA 的 IPC 內存池提供與 CUDA 的虛擬內存管理 API 相同的安全優(yōu)勢。

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

F.11.1. Creating and Sharing IPC Memory Pools

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

// 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

導入的內存池最初只能從其常駐設備訪問。 導入的內存池不繼承導出進程設置的任何可訪問性。 導入過程需要啟用從它計劃訪問內存的任何 GPU 的訪問（使用 cudaMemPoolSetAccess）。

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

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

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

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

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

釋放分配時，需要先在導入過程中釋放分配，然后在導出過程中釋放分配。 以下代碼片段演示了使用 CUDA IPC 事件在兩個進程中的cudaFreeAsync操作之間提供所需的同步。 導入過程中對分配的訪問顯然受到導入過程側的自由操作的限制。 值得注意的是，cudaFree可用于釋放兩個進程中的分配，并且可以使用其他流同步 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 充當空操作，并且 cudaMemPoolAttrReleaseThreshold 被有效地忽略。 此行為由驅動程序控制，而不是運行時控制，并且可能會在未來的驅動程序更新中發(fā)生變化。

F.11.5. IPC Import Pool Limitations

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

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

資源使用統(tǒng)計屬性查詢僅反映導入進程的分配和相關的物理內存。

F.12. Synchronization API Actions

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

F.13. Addendums

F.13.1. cudaMemcpyAsync Current Context/Device Sensitivity

在當前的 CUDA 驅動程序中，任何涉及來自 cudaMallocAsync 的內存的異步 memcpy 都應該使用指定流的上下文作為調用線程的當前上下文來完成。 這對于 cudaMemcpyPeerAsync 不是必需的，因為引用了 API 中指定的設備主上下文而不是當前上下文。

F.13.2. cuPointerGetAttribute Query

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

F.13.3. cuGraphAddMemsetNode

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

F.13.4. Pointer Attributes

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

關于作者

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

審核編輯：郭婷