0 前言

翻了下以前關于CUDA的UVA、零拷貝、統一內存的筆記，感覺順序有些亂，而且里面有個描述還是錯的，這次重新整理一下。

1 swap內存跟鎖頁內存

Swap 是操作系統提供的一種“虛擬內存擴展機制”。當物理內存（比如一根 4GB 內存條）不夠用時，操作系統會將某些暫時不活躍的內存頁（比如后臺程序的數據）換出（swap out）到磁盤上的交換空間（Swap 分區或 Swap 文件），從而釋放物理內存，給當前活躍的程序使用。

通俗地說：

“房子床位不夠，就把出差的人行李先收拾塞進倉庫（磁盤），騰出床位給新來的程序。老住戶回來時，再從倉庫拿回來換入（swap in）。”

與此相反，鎖頁內存就是告訴操作系統，這塊內存是我“強占”的，不能隨便給我換到磁盤上去。

2 UVA(Unified Virtual Addressing)統一虛擬地址

“以前 CPU 和 GPU 各自管理自己的虛擬地址空間，彼此之間的指針不能通用。而有了統一虛擬地址（UVA）后，CPU 內存和 GPU 顯存共享同一個虛擬地址空間，指針在 CPU 和 GPU 間可以直接傳遞、訪問更自然、管理更統一。”

3 先看最普通的cuda內存分配、釋放、傳輸

先看內存分配和釋放

float * devMem=NULL;
cudaError_t cudaMalloc((float**) devMem, count)
cudaError_t cudaMemset(void * devPtr,int value,size_t count)
cudaError_t cudaFree(void * devPtr)

然后是內存傳輸

cudaError_t cudaMemcpy(void *dst,const void * src,size_t count,enum cudaMemcpyKind kind)

然后最后一個參數有下面四個枚舉值

cudaMemcpyHostToHost
cudaMemcpyHostToDevice
cudaMemcpyDeviceToHost
cudaMemcpyDeviceToDevice

4 申請鎖頁內存

cudaError_t cudaHostAlloc(void ** pHost,size_t count,unsigned int flags)

這里的第三個參數flags可以有下面四個選項，這個 flags 可以是以下這些值之一或多個（可 OR 組合）

cudaHostAllocDefalt
cudaHostAllocPortable
cudaHostAllocWriteCombined
cudaHostAllocMapped

4.1 cudaHostAllocDefault

• 說明：這是默認行為，分配常規的鎖頁內存。
• 特性：
  • 分配的內存是鎖頁內存（pinned），不允許被操作系統換出（swap）。
  • 適合用于 CPU?GPU 傳輸的緩沖區，支持高效的 DMA 傳輸。
  • 不保證能被所有 CUDA 上下文共享，也不保證是寫結合（write-combined）內存。
• 使用場景：
  • 一般常規的主機緩沖區分配，兼顧傳輸性能和通用性

4.2 cudaHostAllocPortable

• 說明：分配的內存是可被所有 CUDA 上下文共享的鎖頁內存。
• 特性：
  • 與 cudaHostAllocDefault 類似，但確保這塊內存在所有 CUDA 上下文中都可用。
  • 適合多 GPU 或多上下文應用程序。
• 使用場景：
  • 多 GPU 環境下，多個上下文都需要訪問同一塊主機內存。

4.3 cudaHostAllocWriteCombined

• 說明：分配寫結合（Write-Combined, WC）內存。
• 特性：
  • 寫結合內存不保證 CPU 讀取效率高，但對 GPU 寫入性能有利。
  • 適合CPU 主要寫入、GPU 主要讀取的場景。
  • CPU 讀這塊內存時速度可能較慢（因為寫結合內存是針對寫優化）。
  • 仍是鎖頁內存，支持高速傳輸。
• 使用場景：
  • CPU 向緩沖區寫數據，GPU 讀取數據的流式處理場景，如視頻解碼后處理。

4.3 cudaHostAllocMapped

• 說明：分配映射的鎖頁內存，允許 GPU 直接訪問這塊主機內存。
• 特性：
  • 這塊內存同時映射到 CPU 和 GPU 地址空間。
  • GPU 可以通過特定設備指針直接訪問主機內存，實現零拷貝（Zero-Copy）。
  • 減少了顯存占用和顯存間的顯式數據拷貝，但訪問速度受限于 PCIe 帶寬。
• 使用場景：
  • 適合小數據量、對延遲敏感、不想顯式拷貝的場景。
  • 需要調用 cudaHostGetDevicePointer() 獲取對應的 GPU 設備指針。
  • 零拷貝場景。

4.4 幾種鎖頁內存總結

標志	特點與說明	使用建議
cudaHostAllocDefault	普通鎖頁內存，等價于 cudaMallocHost()	最常用，適合常規 H?D 拷貝
cudaHostAllocPortable	多 context 多 GPU 共享主機內存	多 GPU / 多線程環境
cudaHostAllocWriteCombined	主機只寫，優化 CPU→GPU 傳輸性能，讀很慢	圖像、音頻、傳感器流式寫緩沖區
cudaHostAllocMapped	支持 ZeroCopy，GPU 可訪問主機內存，需要配合 cudaHostGetDevicePointer() 使用	小數據共享、無需頻繁 memcpy 場景

4.5 cudaHostAllocDefault補充說明

這個cudaHostAllocDefault也是比較常用的一個flag.

特性	說明
? 分配主機鎖頁（Pinned）內存	比普通 malloc 分配的 pageable memory 更適合 cudaMemcpy
? 提高 H2D / D2H 的數據傳輸速率	DMA 傳輸，繞過頁交換機制，避免內核拷貝中斷
? 適用于大多數單 GPU、單 context 應用	也是最不容易踩坑的分配方式
? 行為與 cudaMallocHost() 完全一致	所以也可以用它來替代后者

當我們在host端申請內存，而我們可能需要再host和device相互之間memcpy這塊內存的時候，用這個申請內存要比用malloc申請的內存更快。

因為 malloc() 分配的是 pageable memory（可分頁內存）：

• 操作系統可以把它 swap 到磁盤；
• 在進行 cudaMemcpy() 時，驅動需要：
  1. 臨時創建一塊鎖頁緩沖區；
  2. 先從 malloc 的內存拷貝到臨時鎖頁內存；
  3. 再拷貝到 GPU 顯存；
• 整個過程是 雙拷貝 + page fault 風險，速度較慢。

而 cudaHostAllocDefault() 分配的是 pinned memory（鎖頁內存）：

• 操作系統保證這塊內存 不會被分頁；
• 可以被 CUDA 驅動直接用于 DMA（直接內存訪問）拷貝；
• 是真正的 單次、高速拷貝。

4.6 cudaMallocHost

cudaError_t cudaMallocHost(void ** devPtr,size_t count)
cudaError_t cudaFreeHost(void *ptr)

一句話總結，cudaMallocHost相當于是cudaHostAlloc的第三個參數選cudaHostAllocDefault。

4.7 零拷貝內存

    /*** Allocate ZeroCopy mapped memory, shared between CUDA and CPU.** @note although two pointers are returned, one for CPU and GPU, they both resolve to the same physical memory.** @param[out] cpuPtr Returned CPU pointer to the shared memory.* @param[out] gpuPtr Returned GPU pointer to the shared memory.* @param[in] size Size (in bytes) of the shared memory to allocate.** @returns `0` if the allocation succeeded, otherwise faield.* @ingroup cudaMemory*/int cudaAllocMapped(void** cpuPtr, void** gpuPtr, size_t size) {if (!cpuPtr || !gpuPtr || size == 0)return -1;CUDA_SAFECALL(cudaHostAlloc(cpuPtr, size, cudaHostAllocMapped), "cudaHostAlloc failed", -1);CUDA_SAFECALL(cudaHostGetDevicePointer(gpuPtr, *cpuPtr, 0), "cudaHostGetDevicePointer failed", -1);memset(*cpuPtr, 0, size);VLOG(3) << "[InferServer] cudaAllocMapped " << size << " bytes, CPU " << *cpuPtr << " GPU " << *gpuPtr;return 0;}/*** Allocate ZeroCopy mapped memory, shared between CUDA and CPU.** @note this overload of cudaAllocMapped returns one pointer, assumes that the*       CPU and GPU addresses will match (as is the case with any recent CUDA version).** @param[out] ptr Returned pointer to the shared CPU/GPU memory.* @param[in] size Size (in bytes) of the shared memory to allocate.** @returns `0` if the allocation succeeded, otherwise failed.* @ingroup cudaMemory*/int cudaAllocMapped(void** ptr, size_t size) {void* cpuPtr{};void* gpuPtr{};if (!ptr || size == 0)return cudaErrorInvalidValue;auto error = cudaAllocMapped(&cpuPtr, &gpuPtr, size);if (error != cudaSuccess)return error;CUDA_SAFECALL(cpuPtr != gpuPtr, "cudaAllocMapped() - addresses of CPU and GPU pointers don't match", cudaErrorMemoryAllocation);*ptr = gpuPtr;return cudaSuccess;}

這個零拷貝內存其實就是在申請鎖頁內存的基礎上，用cudaHostGetDevicePointer獲取了跟鎖頁內存對應的GPU設備內存指針。cpuPtr 和 gpuPtr 實際上指向的是同一塊物理內存。這是通過CUDA的統一虛擬尋址（Unified Virtual Addressing, UVA）實現的。

4.7.1 補充說明：ZeroCopy 的注意事項

• cudaHostAllocMapped 產生的內存雖然 CPU 和 GPU 都能訪問，但仍然是主機內存，GPU 訪問時通過 PCIe 遠程訪問(不是cuda驅動自動背后memcpy)，性能遠不如顯存；
• 即使有了 UVA（統一虛擬尋址），也仍需調用 cudaHostGetDevicePointer() 獲取 GPU 可訪問的地址；
• 不適合大數據頻繁訪問，用在小數據共用、高效開發場景更好。

4.8 malloc、cudaHostAllocDefault()、cudaHostAllocMapped() 對比

方法	是否鎖頁內存	GPU是否可直接訪問	性能
malloc()	? 否	? 否	普通 CPU 內存，傳輸慢
cudaHostAllocDefault()	? 是	? 否	高效 H2D/D2H 拷貝
cudaHostAllocMapped()	? 是	? 是（需映射）	可 ZeroCopy，但訪問慢

主機內存類型	拷貝方式	帶寬性能（相對）
malloc()	cudaMemcpy()	1.0x
cudaHostAlloc()	cudaMemcpy()	🔺 1.5x ~ 2.5x
cudaHostAllocMapped() + 直接訪問	ZeroCopy	?? 慢，適合小數據

5 統一內存（Unified Memory） cudaMallocManaged

統一內存是從 CUDA 6.0 引入的一項機制，其核心目標是：

? 簡化內存管理 —— 讓 CPU 和 GPU 使用同一個指針訪問數據，CUDA 運行時自動在主機和設備之間遷移數據，無需手動調用 cudaMemcpy。

特點：

• 使用 cudaMallocManaged() 分配的托管內存，可以被 CPU 和 GPU 共同訪問；
• 背后會在 CPU/GPU 之間 自動分頁遷移（通過頁錯誤機制），無需手動拷貝；
• 依賴于 UVA（統一虛擬地址）實現統一指針；
• 內存不再需要分別分配 host 和 device 內存再同步內容；
• 更適合新手開發、代碼更簡潔，但有時性能不如手動拷貝。

6 匯總比較

類型	分配方式	是否鎖頁內存	是否需 memcpy	GPU 是否直接訪問	性能表現	適合場景
普通主機內存	malloc()	? 否	? 需要	? 否	?? 最慢，H2D需拷貝	最普通的內存，不推薦傳輸用
鎖頁主機內存	cudaHostAllocDefault() 或 cudaMallocHost()	? 是	? 需要	? 否	? 快速拷貝（H2D/D2H）	高效拷貝用，推薦用于傳輸
零拷貝內存（ZeroCopy）	cudaHostAllocMapped()	? 是	? 不需要	? 是（映射）	?? 延遲高，帶寬低	小數據共享、開發階段調試
統一內存（Unified Memory）	cudaMallocManaged()	? 是（托管）	? 不需要	? 是（自動遷移）	? 自動遷移但性能波動	開發方便，復雜數據結構共享等

補充說明

• ? 鎖頁內存（Pinned Memory）：不能被操作系統 swap，提高了傳輸效率。
• ?? ZeroCopy：雖然不需要顯式拷貝，但實際通過 PCIe 總線遠程訪問，延遲和帶寬都劣于顯存。
• ? Unified Memory：托管內存在訪問時由 CUDA 運行時系統自動分頁遷移，適合開發快速驗證，性能不易控制。
• cudaMemcpy()：適用于大數據高吞吐傳輸，配合顯存使用效率最佳。