Stream-Ordered Memory Allocator

重點總覽

配置與釋放語意（Stream-Ordered Semantics）

cudaMallocAsync 觸發綁定到特定 stream 的非同步配置，cudaFreeAsync 則以 stream-ordered 方式釋放。兩者都不阻塞 host 或其他 stream，避免 cudaMalloc/cudaFree 的昂貴同步。

void *ptr;
size_t size = 512;
cudaMallocAsync(&ptr, size, cudaStreamPerThread);
// do work using the allocation
kernel<<<..., cudaStreamPerThread>>>(ptr, ...);
// 不需同步 CPU 與 GPU 即可指定非同步釋放
cudaFreeAsync(ptr, cudaStreamPerThread);

• 同一 stream 內「配置 → 使用 → 釋放」是最基本的正確用法。
• cudaMallocAsync 忽略 current device/context，改依「指定的 memory pool 或 supplied stream」決定 allocation 落在哪個 device。

時序（單一 stream）：
 stream: [ mallocAsync ]──[ kernel use ]──[ freeAsync ]
            位址有效起點 ────────────────► 位址失效起點

跨 stream 釋放需用 event/stream 同步

當配置、使用、釋放分散在不同 stream，必須用 event 或 stream 同步保證順序。

cudaMallocAsync(&ptr, size, stream1);
cudaEventRecord(event1, stream1);
cudaStreamWaitEvent(stream2, event1);   // stream2 等配置就緒才存取
kernel<<<..., stream2>>>(ptr, ...);
cudaEventRecord(event2, stream2);
cudaStreamWaitEvent(stream3, event2);   // stream3 等存取完成才釋放
cudaFreeAsync(ptr, stream3);

重點：free 前所有對該記憶體的存取都必須完成；event 把「就緒」與「用完」這兩個時間點接力傳遞。

與 cudaMalloc/cudaFree 的互通

cudaMallocAsync(&ptr, size, stream);
kernel<<<..., stream>>>(ptr, ...);
cudaStreamSynchronize(stream);  // 必須同步，否則會過早釋放
cudaFree(ptr);

Memory Pools

memory pool 封裝 virtual address 與 physical memory 資源，依 pool 屬性與性質配置管理。pool 最主要的面向是它管理的記憶體種類與位置。

• 所有 cudaMallocAsync 都從某個 pool 取用資源；未指定時用 supplied stream 之 device 的 current pool。
• current pool 以 cudaDeviceSetMempool 設定、cudaDeviceGetMempool 查詢；未呼叫 set 前 default pool 為作用中。
• cudaMallocFromPoolAsync（及 cudaMallocAsync 的 C++ overload）可在不設為 current pool 的情況下指定 pool。
• cudaDeviceGetDefaultMempool / cudaMemPoolCreate 回傳 pool handle；cudaMemPoolSetAttribute / cudaMemPoolGetAttribute 控制屬性。

Default / Implicit Pools

• 以 cudaDeviceGetDefaultMempool 取得。
• 配置物為 non-migratable 的 device allocation，位於該 device，永遠可從該 device 存取。
• 存取性可用 cudaMemPoolSetAccess 修改、cudaMemPoolGetAccess 查詢。
• 因不需顯式建立，常稱為 implicit pool；default pool 不支援 IPC。

Explicit Pools

cudaMemPoolCreate 建立 explicit pool，可要求 default/implicit pool 無法提供的屬性（IPC 能力、最大池大小、駐於特定 CPU NUMA node 等）。

// 類似 device 0 implicit pool 的 pool
int device = 0;
cudaMemPoolProps poolProps = { };
poolProps.allocType     = cudaMemAllocationTypePinned;
poolProps.location.id   = device;
poolProps.location.type = cudaMemLocationTypeDevice;
cudaMemPoolCreate(&memPool, &poolProps);

建立駐於 CPU NUMA node 且可經 file descriptor 做 IPC 的 pool，只需改 location.type 為 cudaMemLocationTypeHostNuma 並設 handleType 為 cudaMemHandleTypePosixFileDescriptor。

Device Accessibility for Multi-GPU Support

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;
  return cudaMemPoolSetAccess(memPool, &accessDesc, 1);  // 開放位址可存取
}

Enabling Memory Pools for IPC

CUDA 的 IPC memory pool 提供與 virtual memory management API 相同的安全保障。跨 process 共享需兩步：

步驟一：共享 pool（建立並強制安全邊界）
  export proc:  cudaMemPoolExportToShareableHandle ──OS IPC──► import proc:
                                                              cudaMemPoolImportFromShareableHandle

步驟二：共享個別 allocation（協調虛擬位址與映射時機）
  export proc:  cudaMemPoolExportPointer ──任意機制──► import proc:
                                                       cudaMemPoolImportPointer

• pool 必須以「pool properties 中指定的 handle type」建立，cudaMemPoolExportToShareableHandle 才會成功。
• handleTypes 設為非零即讓 pool 可匯出（IPC capable），例如 CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR（FD 為整數型 handle）。

// exporting process：建立可匯出之 IPC pool
poolProps.handleTypes = CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR;
cudaMemPoolCreate(&memPool, &poolProps);
int fdHandle = 0;
cudaMemPoolExportToShareableHandle(&fdHandle, memPool,
        CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR, 0);  // 傳指標進去
// importing process：由 shareable handle 建立 imported pool（handle 以 value 傳入）
cudaMemPoolImportFromShareableHandle(&importedMemPool, (void*)fdHandle,
        CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR, 0);

共享與釋放 allocation（IPC event 同步）

pool 共享後，exporting process 用 cudaMallocAsync 從 pool 配置的記憶體即可分享給已匯入的 process。安全策略在 pool 層級建立驗證，故 opaque 的 cudaMemPoolPtrExportData 可用任意機制傳遞。

// exporting：配置並記錄 ready event（cudaEventInterprocess + cudaEventDisableTiming）
cudaMallocAsync(&ptr, size, exportMemPool, stream);
cudaEventRecord(readyIpcEvent, stream);
cudaMemPoolExportPointer(&exportData, ptr);
cudaIpcGetEventHandle(&readyIpcEventHandle, readyIpcEvent);
// importing：匯入後等 ready event 才使用（import 不阻塞於 allocation 就緒）
cudaMemPoolImportPointer(&ptr, importedMemPool, importData);
cudaStreamWaitEvent(stream, readyIpcEvent);
kernel<<<..., stream>>>(ptr, ...);

importing: kernel ─ freeAsync ─ ipcEventRecord(finished) ─┐
                                                          ▼ (跨 process)
exporting:                      streamWaitEvent(finished) ─ kernel ─ freeAsync

IPC Export / Import Pool 限制

Best Practices and Tuning

Query for Support

int driverVersion = 0, deviceSupportsMemoryPools = 0, poolSupportedHandleTypes = 0;
cudaDriverGetVersion(&driverVersion);
if (driverVersion >= 11020)
  cudaDeviceGetAttribute(&deviceSupportsMemoryPools,
                         cudaDevAttrMemoryPoolsSupported, device);
if (driverVersion >= 11030)
  cudaDeviceGetAttribute(&poolSupportedHandleTypes,
                         cudaDevAttrMemoryPoolSupportedHandleTypes, device);

• cudaDevAttrMemoryPoolsSupported：是否支援 stream-ordered allocator。
• cudaDevAttrMemoryPoolSupportedHandleTypes：IPC 支援的 handle type（CUDA 11.3 新增；舊 driver 查詢會回 cudaErrorInvalidValue）。
• 先做 driver version 檢查可避免該錯誤；或用 cudaGetLastError 清除錯誤。

Physical Page Caching Behavior

預設下 allocator 盡量縮小 pool 持有的 physical memory。要減少 OS 配置/釋放呼叫，應為每個 pool 設定記憶體足跡——即 release threshold。

• cudaMemPoolAttrReleaseThreshold：pool 在試圖把記憶體歸還 OS 前可持有的 bytes。
• 超過門檻時，allocator 會在下一次 stream / event / device synchronize 嘗試歸還。
• 設為 UINT64_MAX 可防止 driver 在每次同步後縮小 pool。

cuuint64_t setVal = UINT64_MAX;
cudaMemPoolSetAttribute(memPool, cudaMemPoolAttrReleaseThreshold, &setVal);
// 需要時顯式縮小：minBytesToKeep 保留下一階段預期需要量
cudaStreamSynchronize(stream);     // 讓 trim 知道 allocation 已不再使用
cudaMemPoolTrimTo(memPool, 0);

重點：把 threshold 設高等於停用自動縮小；此時若想釋放足跡須顯式呼叫 cudaMemPoolTrimTo，minBytesToKeep 讓你保留一定量供後續階段使用。

Resource Usage Statistics

• *MemHigh watermark 可用 cudaMemPoolSetAttribute 把它設回 current 值（傳入 0 即重置為 current）。
• 可寫 helper 用 cudaMemPoolGetAttribute 一次抓 reserved/reservedHigh/used/usedHigh。

Memory Reuse Policies

服務配置請求時，driver 會先嘗試重用先前 cudaFreeAsync 釋放的記憶體，再向 OS 配置新記憶體。同一 stream 釋放的記憶體可立即被該 stream 後續配置重用；當 stream 與 CPU 同步後，先前在該 stream 釋放的記憶體即可供任意 stream 配置重用。policy 對 default 與 explicit pool 皆適用。

配置請求 ─► 可重用已釋放記憶體？
              ├─ Follow event deps：循 event 依賴跨 stream 重用
              ├─ Opportunistic：free 之 stream-order 已達成即重用
              └─ 否則向 OS 配置；失敗時 Internal deps：插入依賴後重用

• 三者皆透過 cudaMemPoolSetAttribute 啟用/停用；升級 driver 可能變更/增補/重排 reuse policy 列舉。
• FollowEventDependencies 不受 Opportunistic 停用影響，兩者獨立。

Synchronization API Actions

allocator 身為 CUDA driver 的一部分，與 synchronize API 整合：

• 使用者請求 driver 同步時，driver 等非同步工作完成。
• 回傳前，driver 判定哪些 free 已「保證完成」，這些 allocation 即可供配置——不論指定 stream 或停用的配置策略。
• driver 同時檢查 cudaMemPoolAttrReleaseThreshold，釋放可釋放的多餘 physical memory。

附錄要點（Addendums）

考試/測驗重點

Related Notes

• 04-CUDA-Features/05-CUDA-Graphs-Memory-Nodes-and-Device-Launch
• 04-CUDA-Features/19-Interprocess-Communication
• 04-CUDA-Features/20-Virtual-Memory-Management
• 04-CUDA-Features/21-Extended-GPU-Memory
• 04-CUDA-Features/01-Unified-Memory-Full-Support
• 04-CUDA-Features/02-Unified-Memory-Platforms-and-Hints
• 03-Advanced-CUDA/10-Tour-of-CUDA-Features
• 02-Programming-GPUs/16-Unified-and-System-Memory
• 04-CUDA-Features/Practice-CUDA-Features
• 00-Dashboard/Exam-Traps