CUDA -- cudaMalloc / cudaMallocHost

學習代碼時，遇到了cudaMalloc 和 cudaMallocHosts 同時出現的情景，所以學習一下二者的區別。

參考資料1：cudaMallocHost函數詳解

參考資料2：How to Optimize Data Transfers in CUDA C/C++   中文翻譯：中文翻譯

 

host內存：分為pageable memory 和 pinned memory

pageable memory： 通過操作系統API（malloc（），new（））分配的存儲器空間；

pinned memory ：始終存在於物理內存中，不會被分配到低速的虛擬內存中，能夠通過DMA加速與設備端進行通信；cudaHostAlloc(), cudaFreeHost()來分配和釋放pinned memory；

使用Malloc分配的內存都是Pageable(交換頁)的，而另一個模式就是Pinned(Page-locked)，實質是強制讓系統在物理內存中完成內存申請和釋放的工作，不參與頁交換，從而提高系統效率，需要使用cudaHostAlloc和cudaFreeHost(cudaMallocHost的內存也這樣釋放)來分配和釋放。

 

Pageable(交換頁)與Pinned(Page-locked)都是“Write-back”，現在X86/X64CPU，會直接在內部使用一個特別的緩沖區，將寫入合並，等沒滿64B（一個cache line)，集中直接寫入一次，越過所有的緩存，而讀取的時候會直接從內存讀取，同樣無視各級緩存。
這種最大的用途可以用來在CUDA上准備輸入數據，因為它在跨PCI-E傳輸的時候，可能會更快一些（因為不需要詢問CPU的cache數據是否在里面）。

 

使用pinned memory優點：主機端-設備端的數據傳輸帶寬高；某些設備上可以通過zero-copy功能映射到設備地址空間，從GPU直接訪問，省掉主存與顯存間進行數據拷貝的工作；

使用pinned memory缺點：pinned memory 不可以分配過多：導致操作系統用於分頁的物理內存變少， 導致系統整體性能下降；通常由哪個cpu線程分配，就只有這個線程才有訪問權限；

*************************************************************************************************************************************************

主機(CPU)數據分配的內存默認是可分頁的。GPU不能直接訪問可分頁的主機內存，所以當從可分頁內存到設備內存的進行數據傳輸時，CUDA驅動必須首先分配一個臨時的不可分頁的或者固定的主機數組，然后將主機數據拷貝到固定數組里，最后再將數據從固定數組轉移到設備內存，如下圖所示：

 

 正如你在圖中所看到的那樣，固定內存被用作數據傳輸的暫存區。我們可以通過直接分配固定內存的主機數組來避免這一開銷。在CUDA C/C++中，我們可以使用cudaMallocHost()或者cudaHostAlloc()來分配固定內存，使用 cudaFreeHost()來釋放內存。

固定內存的分配有可能會失敗，所以你應該總是檢查錯誤。下面的代碼片段演示了如何分配固定內存並進行錯誤檢查。

cudaError_t status = cudaMallocHost((void**)&h_aPinned, bytes);
if (status != cudaSuccess)
  printf("Error allocating pinned host memoryn");

固定內存的數據傳輸和可分頁內存一樣，使用相同的cudaMemcpy()語法。我們可以使用下面的“bandwidthtest”(帶寬測試)程序(同樣可以在Github上找到)來對比可分頁內存和固定內存的傳輸速度。

#include <stdio.h>
#include <assert.h>
 
// Convenience function for checking CUDA runtime API results
// can be wrapped around any runtime API call. No-op in release builds.
inline
cudaError_t checkCuda(cudaError_t result)
{
#if defined(DEBUG) || defined(_DEBUG)
  if (result != cudaSuccess) {
    fprintf(stderr, "CUDA Runtime Error: %sn",
            cudaGetErrorString(result));
    assert(result == cudaSuccess);
  }
#endif
  return result;
}
 
void profileCopies(float        *h_a,
                   float        *h_b,
                   float        *d,
                   unsigned int  n,
                   char         *desc)
{
  printf("n%s transfersn", desc);

 

(24條消息) CUDA:cudaMalloc vs cudaMallocHost_微風❤水墨的博客-CSDN博客_cuda_malloc