6.1 CUDA: pinned memory固定存儲

CPU和GPU內存交互

在CUDA編程中,內存拷貝是非常費時的一個動作.

從上圖我們可以看出:
1. CPU和GPU之間的總線bus是PCIe,是雙向傳輸的.

2. CPU和GPU之間的數據拷貝使用DMA機制來實現,非常容易理解,為了更快的傳輸速度. 

 

虛擬內存(virtual memory)

我們都知道，雖然在運行速度上硬盤不如內存，但在容量上內存是無法與硬盤相提並論的。當運行一個程序需要大量數據、占用大量內存時，內存就會被“塞滿”，並將那些暫時不用的數據放到硬盤中，而這些數據所占的空間就是虛擬內存。

分頁（英語：Paging），是一種操作系統里存儲器管理的一種技術，可以使電腦的主存可以使用存儲在輔助存儲器中的數據。操作系統會將輔助存儲器（通常是磁盤）中的數據分區成固定大小的區塊，稱為“頁”（pages）。當不需要時，將分頁由主存（通常是內存）移到輔助存儲器；當需要時，再將數據取回，加載主存中。相對於分段，分頁允許存儲器存儲於不連續的區塊以維持文件系統的整齊。^[1]分頁是磁盤和內存間傳輸數據塊的最小單位.

 

固定內存(pinned memory)

我們用cudaMalloc()為GPU分配內存,用malloc()為CPU分配內存.除此之外,CUDA還提供了自己獨有的機制來分配host內存:cudaHostAlloc(). 這個函數和malloc的區別是什么呢?
malloc()分配的標准的,可分頁的主機內存(上面有解釋到),而cudaHostAlloc()分配的是頁鎖定的主機內存,也稱作固定內存pinned memory,或者不可分頁內存,它的一個重要特點是操作系統將不會對這塊內存分頁並交換到磁盤上,從而保證了內存始終駐留在物理內存中.也正因為如此,操作系統能夠安全地使某個應用程序訪問該內存的物理地址,因為這塊內存將不會被破壞或者重新定位.

由於GPU知道內存的物理地址,因此就可以使用DMA技術來在GPU和CPU之間復制數據.當使用可分頁的內存進行復制時(使用malloc),CUDA驅動程序仍會通過dram把數據傳給GPU,這時復制操作會執行兩遍,第一遍從可分頁內存復制一塊到臨時的頁鎖定內存,第二遍是再從這個頁鎖定內存復制到GPU上.當從可分頁內存中執行復制時,復制速度將受限制於PCIE總線的傳輸速度和系統前段速度相對較低的一方.在某些系統中,這些總線在帶寬上有着巨大的差異,因此當在GPU和主機之間復制數據時,這種差異會使頁鎖定主機內存比標准可分頁的性能要高大約2倍.即使PCIE的速度於前端總線的速度相等,由於可分頁內訓需要更多一次的CPU參與復制操作,也會帶來額外的開銷.

當我們在調用cudaMemcpy(dest, src, ...)時，程序會自動檢測dest或者src是否為Pinned Memory，若不是，則會自動將其內容拷入一不可見的Pinned Memory中，然后再進行傳輸。可以手動指定Pinned Memory，對應的API為：cudaHostAlloc(address, size, option)分配地址，cudaFreeHost(pointer)釋放地址。注意，所謂的Pinned Memory都是在Host端的，而不是Device端。

有的人看到這里,在寫代碼的過程中把所有的malloc都替換成cudaHostAlloc()這樣也是不對的.

固定內存是一把雙刃劍.當時使用固定內存時,虛擬內存的功能就會失去,尤其是,在應用程序中使用每個頁鎖定內存時都需要分配物理內存,而且這些內存不能交換到磁盤上.這將會導致系統內存會很快的被耗盡,因此應用程序在物理內存較少的機器上會運行失敗,不僅如此,還會影響系統上其他應用程序的性能.
綜上所述,建議針對cudaMemcpy()調用中的源內存或者目標內存,才使用頁鎖定內存,並且在不在使用他們的時候立即釋放,而不是在應用程序關閉的時候才釋放.我們使用下面的測試實例:

float cuda_malloc_test( int size, bool up ) {
    cudaEvent_t start, stop;
    int *a, *dev_a;
    float elapsedTime;
    HANDLE_ERROR( cudaEventCreate( &start ) );
    HANDLE_ERROR( cudaEventCreate( &stop ) );
    a = (int*)malloc( size * sizeof( *a ) );
    HANDLE_NULL( a );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_a,
    size * sizeof( *dev_a ) ) );
    HANDLE_ERROR( cudaEventRecord( start, 0 ) );
    for (int i=0; i<100; i++) {
    if (up)
        HANDLE_ERROR( cudaMemcpy( dev_a, a,size * sizeof( *dev_a ),cudaMemcpyHostToDevice ) );
    else
        HANDLE_ERROR( cudaMemcpy( a, dev_a,size * sizeof( *dev_a ),cudaMemcpyDeviceToHost ) );
    }
    HANDLE_ERROR( cudaEventRecord( stop, 0 ) );
    HANDLE_ERROR( cudaEventSynchronize( stop ) );
    HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime,start, stop ) );
    free( a );
    HANDLE_ERROR( cudaFree( dev_a ) );
    HANDLE_ERROR( cudaEventDestroy( start ) );
    HANDLE_ERROR( cudaEventDestroy( stop ) );
    return elapsedTime;
}

float cuda_host_alloc_test( int size, bool up ) {
    cudaEvent_t start, stop;int *a, *dev_a;
    float elapsedTime;
    HANDLE_ERROR( cudaEventCreate( &start ) );
    HANDLE_ERROR( cudaEventCreate( &stop ) );
    HANDLE_ERROR( cudaHostAlloc( (void**)&a,size * sizeof( *a ),cudaHostAllocDefault ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_a,size * sizeof( *dev_a ) ) );
    HANDLE_ERROR( cudaEventRecord( start, 0 ) );
    for (int i=0; i<100; i++) {
        if (up)
        HANDLE_ERROR( cudaMemcpy( dev_a, a,size * sizeof( *a ),cudaMemcpyHostToDevice ) );
        else
        HANDLE_ERROR( cudaMemcpy( a, dev_a,size * sizeof( *a ),cudaMemcpyDeviceToHost ) );
    }
    HANDLE_ERROR( cudaEventRecord( stop, 0 ) );
    HANDLE_ERROR( cudaEventSynchronize( stop ) );
    HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime,start, stop ) );
    HANDLE_ERROR( cudaFreeHost( a ) );
    HANDLE_ERROR( cudaFree( dev_a ) );
    HANDLE_ERROR( cudaEventDestroy( start ) );
    HANDLE_ERROR( cudaEventDestroy( stop ) );
    return elapsedTime;
}

 Main 函數代碼:

#include "../common/book.h"
#define SIZE    (10*1024*1024)
int main( void ) {
    float elapsedTime;
    float MB = (float)100*SIZE*sizeof(int)/1024/1024;
    
    elapsedTime = cuda_malloc_test( SIZE, true );
    printf( "Time using cudaMalloc:%3.1f ms\n",elapsedTime );
    printf( "\tMB/s during copy up:%3.1f\n",MB/(elapsedTime/1000) );
    
    elapsedTime = cuda_malloc_test( SIZE, false );
    printf( "Time using cudaMalloc:%3.1f ms\n",elapsedTime );
    printf( "\tMB/s during copy down:%3.1f\n",MB/(elapsedTime/1000) );
    
    elapsedTime = cuda_host_alloc_test( SIZE, true );
    printf( "Time using cudaHostAlloc:%3.1f ms\n",elapsedTime );
    printf( "\tMB/s during copy up:%3.1f\n",MB/(elapsedTime/1000) );
    
    elapsedTime = cuda_host_alloc_test( SIZE, false );
    printf( "Time using cudaHostAlloc:%3.1f ms\n",elapsedTime );
    printf( "\tMB/s during copy down:%3.1f\n",MB/(elapsedTime/1000) );
}

cuda_malloc_test()的參數up為true,因此前一次調用將測試從主機到設備的復制性能.false則測試相反方向設備到主機的性能.

同時也執行了相同的步驟來測試cudaHostAlloc()的性能,在GeForce GTX 285上,當使用固定內存而不是可分頁內存時,從主機拷貝到設備的性能從2.77GB/s 提升到5.11GB/s.當從設備復制到主機時,性能從2.43GB/s提升到5.46GB/s.因此對於大多數PCIE寬帶有限的應用程序,當使用固定內存而不是標准分頁內存時,可以看到顯著的性能提升.