CUDA ---- Memory Model

Memory

kernel性能高低是不能單純的從warp的執行上來解釋的。好比以前博文涉及到的，將block的維度設置爲warp大小的一半會致使load efficiency下降，這個問題沒法用warp的調度或者並行性來解釋。根本緣由是獲取global memory的方式不好勁。

衆所周知，memory的操做在講求效率的語言中佔有極重的地位。low-latency和high-bandwidth是高性能的理想狀況。可是購買擁有大容量，高性能的memory是不現實的，或者不經濟的。所以，咱們就要儘可能依靠軟件層面來獲取最優latency和bandwidth。CUDA將memory model unit分爲device和host兩個系統，充分暴露了其內存結構以供咱們操做，給予用戶充足的使用靈活性。

Benefits of a Memory Hierarchy

通常來講，程序獲取資源是有規律的，也就是計算機體系結構常常提到的局部原則。其又分爲時間局部性和空間局部性。 相信你們對計算機內存方面的知識都很熟悉了，這裏就很少說了，只簡單提下。

 

GPU和CPU的主存都是用DRAM實現，cache則是用lower-latency的SRAM來實現。GPU和CPU的存儲結構基本同樣。並且CUDA將memory結構更好的呈現給用戶，從而能更靈活的控制程序行爲。

CUDA Memory Model

對於程序員來講，memory能夠分爲下面兩類：

• Programmable：咱們能夠靈活操做的部分。
• Non-programmable：不能操做，由一套自動機制來達到很好的性能。

在CPU的存儲結構中，L1和L2 cache都是non-programmable的。對於CUDA來講，programmable的類型很豐富：

• Registers
• Shared memory
• Local memory
• Constant memory
• Texture memory
• Global memory

下圖展現了memory的結構，他們各自都有不用的空間、生命期和cache。

 

其中constant和texture是隻讀的。最下面這三個global、constant和texture擁有相同的生命週期。

Registers

寄存器是GPU最快的memory，kernel中沒有什麼特殊聲明的自動變量都是放在寄存器中的。當數組的索引是constant類型且在編譯期能被肯定的話，就是內置類型，數組也是放在寄存器中。

寄存器變量是每一個線程私有的，一旦thread執行結束，寄存器變量就會失效。寄存器是稀有資源。在Fermi上，每一個thread限制最多擁有63個register，Kepler則是255個。讓本身的kernel使用較少的register就可以容許更多的block駐留在SM中，也就增長了Occupancy，提高了性能。

使用nvcc的-Xptxas -v,-abi=no（這裏Xptxas表示這個是要傳給ptx的參數，不是nvcc的，v是verbose，abi忘了，好像是application by interface）選項能夠查看每一個thread使用的寄存器數量，shared memory和constant memory的大小。若是kernel使用的register超過硬件限制，這部分會使用local memory來代替register，即所謂的register spilling，咱們應該儘可能避免這種狀況。編譯器有相應策略來最小化register的使用而且避免register spilling。咱們也能夠在代碼中顯式的加上額外的信息來幫助編譯器作優化：

__global__ void
__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor)
kernel(...) {
    // your kernel body
}

maxThreadsPerBlock指明每一個block能夠包含的最大thread數目。minBlocksPerMultiprocessor是可選的參數，指明必要的最少的block數目。

咱們也可使用-maxrregcount=32來指定kernel使用的register最大數目。若是使用了__launch_bounds__，則這裏指定的32將失效。

Local Memory

有時候，若是register不夠用了，那麼就會使用local memory來代替這部分寄存器空間。除此外，下面幾種狀況，編譯器可能會把變量放置在local memory：

• 編譯期沒法決定確切值的本地數組。
• 較大的結構體或者數組，也就是那些可能會消耗大量register的變量。
• 任何超過寄存器限制的變量。

local memory這個名字是有歧義的：在local memory中的變量本質上跟global memory在同一塊存儲區。因此，local memory有很高的latency和較低的bandwidth。在CC2.0以上，GPU針對local memory會有L1（per-SM）和L2（per-device）兩級cache。

Shared Memory

用__shared__修飾符修飾的變量存放在shared memory。由於shared memory是on-chip的，他相比localMemory和global memory來講，擁有高的多bandwidth和低不少的latency。他的使用和CPU的L1cache很是相似，可是他是programmable的。

按慣例，像這類性能這麼好的memory都是有限制的，shared memory是以block爲單位分配的。咱們必須很是當心的使用shared memory，不然會無心識的限制了active warp的數目。

不一樣於register，shared memory儘管在kernel裏聲明的，可是他的生命週期是伴隨整個block，而不是單個thread。當該block執行完畢，他所擁有的資源就會被釋放，從新分配給別的block。

shared memory是thread交流的基本方式。同一個block中的thread經過shared memory中的數據來相互合做。獲取shared memory的數據前必須先用__syncthreads()同步。L1 cache和shared memory使用相同的64KB on-chip memory，咱們也可使用下面的API來動態配置兩者：

cudaError_t cudaFuncSetCacheConfig(const void* func, enum cudaFuncCachecacheConfig);

func是分配策略，可使用下面幾種：

cudaFuncCachePreferNone: no preference (default)

cudaFuncCachePreferShared: prefer 48KB shared memory and 16KB L1 cache

cudaFuncCachePreferL1: prefer 48KB L1 cache and 16KB shared memory

cudaFuncCachePreferEqual: Prefer equal size of L1 cache and shared memory, both 32KB

Fermi僅支持前三種配置，Kepler支持所有，注意，在Maxwell以後，L1被捨棄了，因此這64KB就徹底屬於shared Memory了，也就沒有了上面這個分配一說。

Constant Memory

Constant Memory駐留在device Memory，而且使用專用的constant cache（per-SM）。該Memory的聲明應該以__connstant__修飾。constant的範圍是全局的，針對全部kernel，對於全部CC其大小都是64KB。在同一個編譯單元，constant對全部kernel可見。

kernel只能從constant Memory讀取數據，所以其初始化必須在host端使用下面的function調用：

cudaError_t cudaMemcpyToSymbol(const void* symbol, const void* src,size_t count);

這個function拷貝src指向的count個byte到symbol的地址，symbol指向的是在device中的global或者constant Memory。

當一個warp中全部thread都從同一個Memory地址讀取數據時，constant Memory表現最好。例如，計算公式中的係數。若是全部的thread從不一樣的地址讀取數據，而且只讀一次，那麼constant Memory就不是很好的選擇，由於一次讀constant Memory操做會廣播給全部thread知道。

Texture Memory

texture Memory駐留在device Memory中，而且使用一個只讀cache（per-SM）。texture Memory實際上也是global Memory在一塊，可是他有本身專有的只讀cache。這個cache在浮點運算頗有用（具體還沒弄懂）。texture Memory是針對2D空間局部性的優化策略，因此thread要獲取2D數據就可使用texture Memory來達到很高的性能，D3D編程中有兩種重要的基本存儲空間，其中一個就是texture。

Global Memory

global Memory是空間最大，latency最高，GPU最基礎的memory。「global」指明瞭其生命週期。任意SM均可以在整個程序的生命期中獲取其狀態。global中的變量既能夠是靜態也能夠是動態聲明。可使用__device__修飾符來限定其屬性。global memory的分配就是以前頻繁使用的cudaMalloc，釋放使用cudaFree。global memory駐留在devicememory，能夠經過32-byte、64-byte或者128-byte三種格式傳輸。這些memory transaction必須是對齊的，也就是說首地址必須是3二、64或者128的倍數。優化memory transaction對於性能提高相當重要。當warp執行memory load/store時，須要的transaction數量依賴於下面兩個因素：

1. Distribution of memory address across the thread of that warp 就是前文的連續
2. Alignment of memory address per transaction 對齊

通常來講，所需求的transaction越多，潛在的沒必要要數據傳輸就越多，從而致使throughput efficiency下降。

對於一個既定的warp memory請求，transaction的數量和throughput efficiency是由CC版本決定的。對於CC1.0和1.1來講，對於global memory的獲取是很是嚴格的。而1.1以上，因爲cache的存在，獲取要輕鬆的多。

GPU Cache

跟CPU的cache同樣，GPU cache也是non-programmable的。在GPU上包含如下幾種cache，在前文都已經提到：

• L1
• L2
• Read-only constant
• Read-only texture

每一個SM都有一個L1 cache，全部SM共享一個L2 cache。兩者都是用來緩存local和global memory的，固然也包括register spilling的那部分。在Fermi GPus 和 Kepler K40或者以後的GPU，CUDA容許咱們配置讀操做的數據是否使用L1和L2或者只使用L2。

在CPU方面，memory的load/store均可以被cache。可是在GPU上，只有load操做會被cache，store則不會。

每一個SM都有一個只讀constant cache和texture cache來提高性能。

CUDA Variable Declaration Summary

下表是以前介紹的幾種memory的聲明總結：

 

 

Static Global Memory

下面的代碼介紹了怎樣靜態的聲明global variable（以前的博文其實都是global variable）。大體過程就是，先聲明瞭一個float全局變量，在checkGlobal-Variable中，該值被打印出來，隨後，其值便被改變。在main中，這個值使用cudaMemcpyToSymbol來初始化。最終當全局變量被改變後，將值拷貝回host。

#include <cuda_runtime.h>
#include <stdio.h>
__device__ float devData;
__global__ void checkGlobalVariable() {
    // display the original value
    printf("Device: the value of the global variable is %f\n",devData);
    // alter the value
    devData +=2.0f;
}

int main(void) {
    // initialize the global variable
    float value = 3.14f;
    cudaMemcpyToSymbol(devData, &value, sizeof(float));
    printf("Host: copied %f to the global variable\n", value);
    // invoke the kernel
    checkGlobalVariable <<<1, 1>>>();
    // copy the global variable back to the host
    cudaMemcpyFromSymbol(&value, devData, sizeof(float));
    printf("Host: the value changed by the kernel to %f\n", value);
    cudaDeviceReset();
    return EXIT_SUCCESS;
}                                    

編譯運行：

$ nvcc -arch=sm_20 globalVariable.cu -o globalVariable
$ ./globalVariable

輸出：

Host: copied 3.140000 to the global variable
Device: the value of the global variable is 3.140000
Host: the value changed by the kernel to 5.140000

熟悉了CUDA的基本思想後，不難明白，儘管host和device的代碼是寫在同一個源文件，可是他們的執行卻在徹底不一樣的兩個世界，host不能直接訪問device變量，反之亦然。

咱們可能會反駁說，用下面的代碼就能得到device的全局變量：

cudaMemcpyToSymbol(devD6ata, &value, sizeof(float));

可是，咱們應該還注意到下面的幾點：

• 該函數是CUDA的runtime API，使用的GPU實現。
• devData在這兒只是個符號，不是device的變量地址。
• 在kernel中，devData被用做變量。

並且，cudaMemcpy不能用&devData這種方式來傳遞變量,正如上面所說，devData只是個符號，取址這種操做自己就是錯誤的：

cudaMemcpy(&devData, &value, sizeof(float),cudaMemcpyHostToDevice);         // It’s wrong!!!

無論怎樣，CUDA仍是爲咱們提供了，利用devData這種符號來獲取變量地址的方式：

cudaError_t cudaGetSymbolAddress(void** devPtr, const void* symbol);

獲取地址以後，就可使用cudaMemcpy了：

float *dptr = NULL;
cudaGetSymbolAddress((void**)&dptr, devData);
cudaMemcpy(dptr, &value, sizeof(float), cudaMemcpyHostToDevice);

咱們只有一種方式可以直接獲取GPU memory，即便用pinned memory，下文將詳細介紹。

Memory Management

CUDA很是接近C的編程風格，以便可以快速上手掌握，在內存管理這點上，CUDA區別於C最明顯的操做就是在device和host之間不停的傳遞數據。很麻煩的一個過程，不過Unified Memory出現後，程序編寫就沒那麼複雜了，可是目前，Unified Memory的使用並未普及，咱們仍是要關注Memory的顯式的操做過程：

• Allocate and deallocate device Memory
• Transfer data between the host and device

爲了達到最好的性能，CUDA提供了五花八門的接口供程序員顯式的在device和host之間傳遞數據。

Memory Allocation and Deallocation

前面的博文已經提到一部份內存分配函數了，在分配global Memory時，最經常使用的就是下面這個了：

cudaError_t cudaMalloc(void **devPtr, size_t count);

若是分配出錯則返回cudaErrorMemoryAllocation。分配成功後，就得對該地址初始化值，要麼從host調用cudaMemcpy賦值，要麼調用下面的API初始化：

cudaError_t cudaMemset(void *devPtr, int value, size_t count);

釋放資源就是：

cudaError_t cudaFree(void *devPtr);

device資源分配是個很是昂貴的操做，因此，device Memory應該儘量的重用，而不是從新分配。

Memory Transfer

一旦global Memory分配好後，若是不用cudaMemset就得用下面這個：

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

這個你們應該也很熟悉了，kind就是下面這幾種：

cudaMemcpyHostToHost

cudaMemcpyHostToDevice

cudaMemcpyDeviceToHost

cudaMemcpyDeviceToDevice

下圖是CPU和GPU之間傳輸關係圖，能夠看出來，CPU和GPU之間傳輸速度相對不好（NVLink技術能提升5~10倍），GPU和on-board Memory傳輸速度要快得多，因此對於編程來講，要時刻考慮減小CPU和GPU之間的數據傳輸。

 

Pinned Memory

Host Memory的分配默認狀況下是pageable的，也就是說，咱們要承受因pagefault致使的操做，，這個操做要將host virtual Memory的數據轉移到由OS決定的不物理位置。GPU沒法安全的獲取host的pageable Memory，由於GPU沒有辦法控制host OS物理上轉移數據的時機。所以，當將pageable host Memory數據送到device時，CUDA驅動會首先分配一個臨時的page-locked或者pinned host Memory，並將host的數據放到這個臨時空間裏。而後GPU從這個所謂的pinned Memory中獲取數據，以下左圖所示：

 

左圖是默認的過程，咱們也能夠顯式的直接使用pinned Memory，以下：

cudaError_t cudaMallocHost(void **devPtr, size_t count);

因爲pinned Memory可以被device直接訪問（不是指不經過PCIE了，而是相對左圖咱們少了pageable Memory到pinned Memory這一步），因此他比pageable Memory具備至關高的讀寫帶寬，固然像這種東西依然不能過分使用，由於這會下降pageable Memory的數量，影響整個虛擬存儲性能，咱們不能因小失大。

cudaError_t status = cudaMallocHost((void**)&h_aPinned, bytes);
if (status != cudaSuccess) {
    fprintf(stderr, "Error returned from pinned host memory allocation\n");
    exit(1);
}    

Pinned Memory的釋放也比較特殊：

cudaError_t cudaFreeHost(void *ptr);

Pinned Memory比pageable Memory的分配操做更加昂貴，可是他對大數據的傳輸有很好的表現。還有就是，pinned Memory效果的高低也是跟CC有關的。

將許多小的傳輸合併到一次大的數據傳輸，並使用pinned Memory將下降很大的傳輸消耗。這裏說起下，數據傳輸的消耗有時候是能夠被kernel的執行覆蓋的。

Zero-Copy Memory

通常來講，host和device是不能直接訪問對方的數據的，前文也有提到，可是Zero-Copy Memory是個特例。

該Memory是位於host的，可是GPU thread能夠直接訪問，其優勢有：

• 當device Memory不夠用時，可以利用host Memory。
• 避免device和host之間顯式的數據傳輸。
• 提升PCIe傳輸效率。

當使用zero-copy來共享host和device數據時，咱們必須同步Memory的獲取，不然，device和host同時訪問該Memory會致使未定義行爲。

Zero-copy自己實質就是pinned memory而且被映射到了device的地址空間。下面是他的分配API：

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

其資源釋放固然也是cudaFreeHost，至於flag則是下面幾個選項：

• cudaHostAllocDefault
• cudaHostAllocPortable
• cudaHostAllocWriteCombined
• cudaHostAllocMapped

當使用cudaHostAllocDefault時，cudaHostAlloc和cudaMallocHost等價。cudaHostAllocPortable則說明，分配的pinned memory對全部CUDA context都有效，而不是單單執行分配此操做的那個context或者說線程。cudaHostAllocWriteCombined是在特殊系統配置狀況下使用的，這塊pinned memory在PCIE上的傳輸更快，可是對於host本身來講，卻沒什麼效率。因此該選項通常用來讓host去寫，而後device讀。最經常使用的是cudaHostAllocMapped，就是返回一個標準的zero-copy。能夠用下面的API來獲取device端的地址：

cudaError_t cudaHostGetDevicePointer(void **pDevice, void *pHost, unsigned int flags);

flags是保留參數，留待未來使用，目前必須設置爲零。

使用zero-copy memory來做爲device memory的讀寫很頻繁的那部分的補充是很不明智的，pinned這一類適合大數據傳輸，不適合頻繁的操做，究其根本緣由仍是GPU和CPU之間低的可憐的傳輸速度，甚至，頻繁讀寫狀況下，zero-copy表現比global memory也要差很多。

下面一段代買是比較頻繁讀寫狀況下，zero-copy的表現： 

int main(int argc, char **argv) {
// part 0: set up device and array
// set up device
int dev = 0;
cudaSetDevice(dev);
// get device properties
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, dev);
// check if support mapped memory
if (!deviceProp.canMapHostMemory) {
printf("Device %d does not support mapping CPU host memory!\n", dev);
cudaDeviceReset();
exit(EXIT_SUCCESS);
}
printf("Using Device %d: %s ", dev, deviceProp.name);
// set up date size of vectors
int ipower = 10;
if (argc>1) ipower = atoi(argv[1]);
int nElem = 1<<ipower;
size_t nBytes = nElem * sizeof(float);
if (ipower < 18) {
printf("Vector size %d power %d nbytes %3.0f KB\n", nElem,\
ipower,(float)nBytes/(1024.0f));
} else {
printf("Vector size %d power %d nbytes %3.0f MB\n", nElem,\
ipower,(float)nBytes/(1024.0f*1024.0f));
}
// part 1: using device memory
// malloc host memory
float *h_A, *h_B, *hostRef, *gpuRef;
h_A = (float *)malloc(nBytes);
h_B = (float *)malloc(nBytes);
hostRef = (float *)malloc(nBytes);
gpuRef = (float *)malloc(nBytes);
// initialize data at host side
initialData(h_A, nElem);
initialData(h_B, nElem);
memset(hostRef, 0, nBytes);
memset(gpuRef, 0, nBytes);
// add vector at host side for result checks
sumArraysOnHost(h_A, h_B, hostRef, nElem);
// malloc device global memory
float *d_A, *d_B, *d_C;
cudaMalloc((float**)&d_A, nBytes);
cudaMalloc((float**)&d_B, nBytes);
cudaMalloc((float**)&d_C, nBytes);
// transfer data from host to device
cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);
// set up execution configuration
int iLen = 512;
dim3 block (iLen);
dim3 grid ((nElem+block.x-1)/block.x);
// invoke kernel at host side
sumArrays <<<grid, block>>>(d_A, d_B, d_C, nElem);
// copy kernel result back to host side
cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);
// check device results
checkResult(hostRef, gpuRef, nElem);
// free device global memory
cudaFree(d_A);
cudaFree(d_B);
free(h_A);
free(h_B);
// part 2: using zerocopy memory for array A and B
// allocate zerocpy memory
unsigned int flags = cudaHostAllocMapped;
cudaHostAlloc((void **)&h_A, nBytes, flags);
cudaHostAlloc((void **)&h_B, nBytes, flags);
// initialize data at host side
initialData(h_A, nElem);
initialData(h_B, nElem);
memset(hostRef, 0, nBytes);
memset(gpuRef, 0, nBytes);
// pass the pointer to device
cudaHostGetDevicePointer((void **)&d_A, (void *)h_A, 0);
cudaHostGetDevicePointer((void **)&d_B, (void *)h_B, 0);
// add at host side for result checks
sumArraysOnHost(h_A, h_B, hostRef, nElem);
// execute kernel with zero copy memory
sumArraysZeroCopy <<<grid, block>>>(d_A, d_B, d_C, nElem);
// copy kernel result back to host side
cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);
// check device results
checkResult(hostRef, gpuRef, nElem);
// free memory
cudaFree(d_C);
cudaFreeHost(h_A);
cudaFreeHost(h_B);
free(hostRef);
free(gpuRef);
// reset device
cudaDeviceReset();
return EXIT_SUCCESS;
}

View Code

編譯運行：

$ nvcc -O3 -arch=sm_20 sumArrayZerocpy.cu -o sumZerocpy
$ nvprof ./sumZerocpy
Using Device 0: Tesla M2090 Vector size 1024 power 10 nbytes 4 KB
Time(%) Time Calls Avg Min Max Name
27.18% 3.7760us 1 3.7760us 3.7760us 3.7760us sumArraysZeroCopy
11.80% 1.6390us 1 1.6390us 1.6390us 1.6390us sumArrays
25.56% 3.5520us 3 1.1840us 1.0240us 1.5040us [CUDA memcpy HtoD]
35.47% 4.9280us 2 2.4640us 2.4640us 2.4640us [CUDA memcpy DtoH]

下表是嘗試不一樣數組長度後的結果：

./sumZerocopy <size-log-2>

 

所以，對於共享host和device之間的一小塊內存空間，zero-copy是很好的選擇，由於他簡化的編程並且提供了合理的性能。

Unified Virtual Addressing

在CC2.0以上的設備支持一種新特性：Unified Virtual Addressing（UVA）。這個特性在CUDA4.0中首次介紹，並被64位Linux系統支持。以下圖所示，在使用UVA的狀況下，CPU和GPU使用同一塊連續的地址空間：

 

在UVA以前，咱們須要分別管理指向host memory和device memory的指針。使用UVA以後，實際指向內存空間的指針對咱們來講是透明的，咱們看到的是同一塊連續地址空間。

這樣，使用cudaHostAlloc分配的pinned memory得到的地址對於device和host來講是通用的。咱們能夠直接在kernel裏使用這個地址。回看前文，咱們對於zero-copy的處理過程是：

• 分配已經映射到device的pinned memory。
• 根據得到的host地址，獲取device的映射地址。
• 在kernel中使用該映射地址。

使用UVA以後，就不必來獲取device的映射地址了，直接使用一個地址就能夠，以下代碼所示：

// allocate zero-copy memory at the host side
cudaHostAlloc((void **)&h_A, nBytes, cudaHostAllocMapped);
cudaHostAlloc((void **)&h_B, nBytes, cudaHostAllocMapped);
// initialize data at the host side
initialData(h_A, nElem);
initialData(h_B, nElem);
// invoke the kernel with zero-copy memory
sumArraysZeroCopy<<<grid, block>>>(h_A, h_B, d_C, nElem);

能夠看到，cudaHostAlloc返回的指針直接就使用在了kernel裏面，編譯指令;

$ nvcc -O3 -arch=sm_20 sumArrayZerocpyUVA.cu -o sumArrayZerocpyUVA

修改後的代碼執行效率和以前的效率是相差無幾的，你們能夠本身動手試試。

Unified Memory

理解個大概，之後熟悉了回來補。。。