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數量依賴於下面兩個因素：

Distribution of memory address across the thread of that warp 就是前文的連續
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

熟悉了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也要差不少。

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

因此，對於共享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

修改後的程式碼執行效率和之前的效率是相差無幾的，大家可以自己動手試試。

轉自：http://www.cnblogs.com/1024incn/p/4564726.html