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必須是對齊的，也就是說首地址必須是32、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;
}

編譯運行：

$ 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

理解個大概，以后熟悉了回來補。。。