點選上方AI人工智能初學者，訂閱我！此刻開始我們一起學習進步！
目錄
1、CUDA線程
2、存儲器架構
2.1、全局記憶體
2.2、本地記憶體和寄存器堆
2.3、告訴緩沖存儲器
3、線程同步
3.1、共享記憶體
3.2、原子操作
筆記來源書籍推薦
1、CUDA線程
CUDA關于并行執行具有分層結構。每次核心啟動時可以被切分成多個并行執行的塊，而每個塊又可以進一步地被切分成多個線程。
在上一推文我們已經知道，maxThreadPerBlock屬性限制了每個塊能啟動的線程數量。這個值對于最新的GPU卡來說是1024。類似地，第二種方式能最大啟動的塊數量被限制成2^31-1個。
更加理想的則是，我們并不單獨啟動1個塊，裡面多個線程；也不啟動多個塊，每個裡面1個線程。我們一次并行啟動多個塊，每個塊裡面多個線程(最多可以是maxThread-PerBlock那麼多哦)。是以，假設上一章的那個向量加法例子你需要啟動N=50000這麼多的線程，我們可以這樣調用核心：
最大的塊能有1024個線程。不過我們這裡舉例，對于N個線程來說，每個塊有512個線程，則需要有N/512個塊。但是如果N不是512的整數倍，那麼N除以512會計算得到錯誤的塊數量，比實際的塊數量少1個。是以為了計算得到下一個最小的能滿足要求的整數結果，N需要加上511，然後再除以512。這基本上是一個除法的向上取整操作。
還是直接撸代碼吧：
#include "stdio.h"#include#include #include //Defining number of elements in Array#define N50000//Defining Kernel function for vector addition__global__ void gpuAdd(int *d_a, int *d_b, int *d_c) {//Getting block index of current kernelint tid = threadIdx.x + blockIdx.x * blockDim.x;while (tid < N){d_c[tid] = d_a[tid] + d_b[tid];tid += blockDim.x * gridDim.x;}}int main(void) {//Defining host arraysint h_a[N], h_b[N], h_c[N];//Defining device pointersint *d_a, *d_b, *d_c;// allocate the memorycudaMalloc((void**)&d_a, N * sizeof(int));cudaMalloc((void**)&d_b, N * sizeof(int));cudaMalloc((void**)&d_c, N * sizeof(int));//Initializing Arraysfor (int i = 0; i < N; i++) {h_a[i] = 2 * i*i;h_b[i] = i;}// Copy input arrays from host to device memorycudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice);cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice);//Calling kernels with N blocks and one thread per block, passing device pointers as parametersgpuAdd << <512, 512 >> >(d_a, d_b, d_c);//Copy result back to host memory from device memorycudaMemcpy(h_c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost);cudaDeviceSynchronize();int Correct = 1;printf("Vector addition on GPU \n");//Printing result on consolefor (int i = 0; i < N; i++) {if ((h_a[i] + h_b[i] != h_c[i])){Correct = 0;}}if (Correct == 1){printf("GPU has computed Sum Correctly\n");}else{printf("There is an Error in GPU Computation\n");}//Free up memorycudaFree(d_a);cudaFree(d_b);cudaFree(d_c);return 0;}
本核心的代碼和上一文寫過的那個很相似。但是有兩處不同：
(1)計算初始的tid的時候；
(2)是添加了while循環部分。
計算初始的tid的變化，是因為我們現在是啟動多個塊，每個裡面有多個線程，直接看成ID的結構，多個塊橫排排列，每個塊裡面有N個線程，那麼自然計算tid的時候是用：
目前塊的ID*目前塊裡面的線程數量+目前線程在塊中的ID
即tid=blockIdx.x(目前塊的ID)*blockDim.x(目前塊裡面的線程數量)+threadIdx.x(目前線程在塊中的ID)。
而while部分每次增加現有的線程數量(因為你沒有啟動到N)，直到達到N。這就如同你有一個卡，一次最多隻能啟動100個塊，每個塊裡有7個線程，也就是一次最多能啟動700個線程。但N的規模是8000，遠遠超過700怎麼辦？答案是直接啟動K個(K≥700)，這樣就能安全啟動。然後裡面添加一個while循環，這700個線程第一次處理[0,699)，第二次處理[700,1400)，第三次處理[1400,2100)……直到這8000個元素都被處理完。這就是我們本例中看到的代碼。初始化時候的tid=threadIdx.x+blockDim.x*blockIdx.x，每次while循環的時候tid+=blockDim.x*gridDim.x(注意一個是=，一個是+=，後者是增加的由來)。下面的2D表格用來輔助了解。
對于任意一個線程，使用blockIdx.x指令可以得到目前的塊的ID，而使用threadIdx.x指令可以得到本線程在該塊中的ID。例如，對于表格中綠色标記的線程，它的塊ID是2，線程ID是1，如果想将這兩個數字進行ID化，得到每個線程唯一的總ID，可以用塊的ID乘以塊中的線程總數，然後加上線程在這個塊中的ID。數學表達式如下：
這次的main函數，和我們上次寫過的那個非常類似。唯一的不同點在于核心的啟動方式。現在我們用512個塊，每個塊裡面有512個線程啟動該核心。這樣N非常大的問題就得到了解決。此外，我們不再将很長的結果數組中的每個值都列印出來，隻列印結果是否正确。
2、存儲器架構
在GPU上的代碼執行被劃分為流多處理器、塊和線程。GPU有幾個不同的存儲器空間，每個存儲器空間都有特定的特征和用途以及不同的速度和範圍。這個存儲空間按層次結構劃分為不同的組塊，比如全局記憶體、共享記憶體、本地記憶體、常量記憶體和紋理記憶體，每個組塊都可以從程式中的不同點通路。此存儲器架構如圖所示：
如圖所示，每個線程都有自己的本地存儲器和寄存器堆。與處理器不同的是，GPU核心有很多寄存器來存儲本地資料。當線程使用的資料不适合存儲在寄存器堆中或者寄存器堆中裝不下的時候，将會使用本地記憶體。寄存器堆和本地記憶體對每個線程都是唯一的。寄存器堆是最快的一種存儲器。同一個塊中的線程具有可由該塊中的所有線程通路的共享記憶體。全局記憶體可被所有的塊和其中的所有線程通路。它具有相當大的通路延遲，但存在緩存這種東西來給它提速。如下表，GPU有一級和二級緩存(即L1緩存和L2緩存)。常量記憶體則是用于存儲常量和核心參數之類的隻讀資料。最後，存在紋理記憶體，這種記憶體可以利用各種2D和3D的通路模式。
所有存儲器特征總結如下。
上表表述了各種存儲器的各種特性。作用範圍欄定義了程式的哪個部分能使用該存儲器。而生存期定義了該存儲器中的資料對程式可見的時間。除此之外，L1和L2緩存也可以用于GPU程式以便更快地通路存儲器。
總之，所有線程都有一個寄存器堆，它是最快的。共享記憶體隻能被塊中的線程通路，但比全局記憶體塊。全局記憶體是最慢的，但可以被所有的塊通路。常量和紋理記憶體用于特殊用途。存儲器通路是程式快速執行的最大瓶頸。
2.1、全局記憶體
所有的塊都可以對全局記憶體進行讀寫。該存儲器較慢，但是可以從你的代碼的任何地方進行讀寫。緩存可加速對全局記憶體的通路。所有通過cudaMalloc配置設定的存儲器都是全局記憶體。下面的簡單代碼示範了如何從程式中使用全局記憶體：
#include #define N 5__global__ void gpu_global_memory(int *d_a){// "array" is a pointer into global memory on the deviced_a[threadIdx.x] = threadIdx.x;}int main(int argc, char **argv){// Define Host Arrayint h_a[N];//Define device pointerint *d_a;cudaMalloc((void **)&d_a, sizeof(int) *N);// now copy data from host memory to device memory cudaMemcpy((void *)d_a, (void *)h_a, sizeof(int) *N, cudaMemcpyHostToDevice);// launch the kernel gpu_global_memory << <1, N >> > (d_a);// copy the modified array back to the host memorycudaMemcpy((void *)h_a, (void *)d_a, sizeof(int) *N, cudaMemcpyDeviceToHost);printf("Array in Global Memory is: \n");//Printing result on consolefor (int i = 0; i < N; i++) {printf("At Index: %d --> %d \n", i, h_a[i]);}return 0;}
這段代碼示範了如何從裝置代碼中進行全局記憶體的寫入，以及如何從主機代碼中用cudaMalloc進行配置設定，如何将指向該段全局記憶體的指針作為參數傳遞給核心函數。核心函數用不同的線程ID的值來填充這段全局記憶體。然後(用cudaMemcpy)複制到記憶體以便顯示内容。最終結果如圖所示：
2.2、本地記憶體和寄存器堆
本地記憶體和寄存器堆對每個線程都是唯一的。寄存器是每個線程可用的最快存儲器。當核心中使用的變量在寄存器堆中裝不下的時候，将會使用本地記憶體存儲它們，這叫寄存器溢出。
請注意使用本地記憶體有兩種情況：
(1)、寄存器不夠了
(2)、某些情況根本就不能放在寄存器中
例如對一個局部數組的下标進行不定索引的時候。基本上可以将本地記憶體看成是每個線程的唯一的全局記憶體部分。相比寄存器堆，本地記憶體要慢很多。雖然本地記憶體通過L1緩存和L2緩存進行了緩沖，但寄存器溢出可能會影響你的程式的性能。
下面示範一個簡單的程式：
#include #define N 5__global__ void gpu_local_memory(int d_in){int t_local;t_local = d_in * threadIdx.x;printf("Value of Local variable in current thread is: %d \n", t_local);}int main(int argc, char **argv){printf("Use of Local Memory on GPU:\n");gpu_local_memory << <1, N >> > (5);cudaDeviceSynchronize();return 0;}
代碼中的t_local變量是每個線程局部唯一的，将被存儲在寄存器堆中。用這種變量計算的時候，計算速度将是最快速的。以上代碼的輸出如圖所示：
2.3、高速緩沖存儲器
在較新的GPU上，每個流多處理器都含有自己獨立的L1緩存，以及GPU有L2緩存。L2緩存是被所有的GPU中的流多處理器都共有的。所有的全局記憶體通路和本地記憶體通路都使用這些緩存，因為L1緩存在流多處理器内部獨有，接近線程執行所需要的硬體機關，是以它的速度非常快。一般來說，L1緩存和共享記憶體共用同樣的存儲硬體，一共是64KB(注意：這是和計算能力有關，不一定共用相同的存儲硬體，也不一定可以配置互相占用的比例，例如計算能力5.X和6.X的GPU卡就不能。同時L1緩存和共享記憶體在這兩個計算能力上也不是共用的，但舊的計算能力和7.X GPU卡是如此)，你可以配置L1緩存和共享記憶體分别在這64KB中的比例。所有的全局記憶體通路通過L2緩存進行。紋理記憶體和常量記憶體也分别有它們獨立的緩存。
3、線程同步
3.1、共享記憶體
共享記憶體位于晶片内部，是以它比全局記憶體快得多。(CUDA裡面存儲器的快慢有兩方面，一個是延遲低，一個是帶寬大。這裡特指延遲低)，相比沒有經過緩存的全局記憶體通路，共享記憶體大約在延遲上低100倍。同一個塊中的線程可以通路相同的一段共享記憶體(注意：不同塊中的線程所見到的共享記憶體中的内容是不相同的)，這在許多線程需要與其他線程共享它們的結果的應用程式中非常有用。但是如果不同步，也可能會造成混亂或錯誤的結果。如果某線程的計算結果在寫入到共享記憶體完成之前被其他線程讀取，那麼将會導緻錯誤。是以，應該正确地控制或管理記憶體通路。這是由__syncthreads()指令完成的，該指令確定在繼續執行程式之前完成對記憶體的所有寫入操作。這也被稱為barrier。barrier的含義是塊中的所有線程都将到達該代碼行，然後在此等待其他線程完成。當所有線程都到達了這裡之後，它們可以一起繼續往下執行。
#include __global__ void gpu_shared_memory(float *d_a){// Defining local variables which are private to each threadint i, index = threadIdx.x;float average, sum = 0.0f;//Define shared memory__shared__ float sh_arr[10];sh_arr[index] = d_a[index];__syncthreads(); // This ensures all the writes to shared memory have completedfor (i = 0; i<= index; i++) { sum += sh_arr[i]; }average = sum / (index + 1.0f);d_a[index] = average; sh_arr[index] = average;}int main(int argc, char **argv){//Define Host Arrayfloat h_a[10]; //Define Device Pointerfloat *d_a; for (int i = 0; i < 10; i++) {h_a[i] = i;}// allocate global memory on the devicecudaMalloc((void **)&d_a, sizeof(float) * 10);// now copy data from host memory to device memory cudaMemcpy((void *)d_a, (void *)h_a, sizeof(float) * 10, cudaMemcpyHostToDevice);gpu_shared_memory << <1, 10 >> >(d_a);// copy the modified array back to the host memorycudaMemcpy((void *)h_a, (void *)d_a, sizeof(float) * 10, cudaMemcpyDeviceToHost);printf("Use of Shared Memory on GPU: \n");//Printing result on consolefor (int i = 0; i < 10; i++) {printf("The running average after %d element is %f \n", i, h_a[i]);}return 0;}
在main函數中，當配置設定好主機和裝置上的數組後，用0.0到9.0填充主機上的數組，然後将這個數組複制到顯存。核心将對顯存中的資料進行讀取，計算并儲存結果。最後結果從顯存中傳輸到記憶體，然後在控制台上輸出。控制台上的輸出結果如圖所示：
這個程式還含有額外的一個CUDA函數調用：cudaDeviceSynchronize()。為何要加這句？這是因為啟動核心是一個異步操作，隻要釋出了核心啟動指令，不等核心執行完成，控制權就會立刻傳回給調用核心的CPU線程。在上述的代碼中，CPU線程傳回，繼續執行的下一句是printf()。而再之後，在核心完成之前，程序就會結束，終止控制台視窗。是以，如果不加上這句同步函數，你就看不到任何的核心執行結果輸出。在程式退出後核心生成的輸出結果，将沒有地方可去，你沒法看到它們，是以，如果我們不包含這個指令，你将不會看到任何核心執行的printf語句的輸出結果。要能看到核心生成的輸出結果，我們必須包含這句同步函數。這樣，核心的結果将通過可用的标準輸出顯示，而應用程式則會在核心執行完成之後才退出。
3.2、原子操作
考慮當大量的線程需要試圖修改一段較小的記憶體區域的情形，這是(在日常的算法實作中)常發生的現象。當我們試圖進行“讀取-修改-寫入”操作序列的時候，這種情形經常會帶來很多麻煩。
一個例子是代碼d_out[i]++，這代碼首先将d_out[i]的原值從存儲器中讀取出來，然後執行了+1操作，再将結果回寫到存儲器。然而，如果多個線程試圖在同一個記憶體區域中進行這個操作，則可能會得到錯誤的結果。
假設某記憶體區域中有初始值6，兩個線程p和q分别試圖将這段區域中的内容+1，則最終的結果應當是8。但是在實際執行的時候，可能p和q兩個線程同時讀取了這個初始值，兩者都得到了6，執行+1操作都得到了7，然後它們将7寫回這個記憶體區域。這樣，和正确的結果8不同，我們得到的最終結果是7，這是錯誤的。這種錯誤是如何的危險，我們通過ATM取現操作來示範。假設你的賬戶餘額為5000盧比，你的賬戶下面開了兩張銀行卡，你和你的朋友同時去2個不同的ATM上取現4000盧比，你倆在同一瞬間刷卡取現。是以，當兩個ATM檢查餘額的時候，都将顯示5000盧比的餘額。當你倆同時取現4000盧比的時候，兩個ATM機都隻根據初始值5000盧比判斷，要取的現金4000盧比小于目前餘額。是以兩個機器将會給你們每人4000盧比。即使你之前隻有5000盧比的餘額，你們也能得到8000盧比，這很危險。為了示範一下這種情形，做了一個很多線程試圖同時通路一個小數組的例子：
#include #define NUM_THREADS 10000#define SIZE 10#define BLOCK_WIDTH 100__global__ void gpu_increment_without_atomic(int *d_a){// Calculate thread id for current threadint tid = blockIdx.x * blockDim.x + threadIdx.x;// each thread increments elements wrapping at SIZE variabletid = tid % SIZE;d_a[tid] += 1;}int main(int argc, char **argv){printf("%d total threads in %d blocks writing into %d array elements\n",NUM_THREADS, NUM_THREADS / BLOCK_WIDTH, SIZE);// declare and allocate host memoryint h_a[SIZE];const int ARRAY_BYTES = SIZE * sizeof(int);// declare and allocate GPU memoryint * d_a;cudaMalloc((void **)&d_a, ARRAY_BYTES);//Initialize GPU memory to zerocudaMemset((void *)d_a, 0, ARRAY_BYTES);gpu_increment_without_atomic <> >(d_a);// copy back the array to host memorycudaMemcpy(h_a, d_a, ARRAY_BYTES, cudaMemcpyDeviceToHost);printf("Number of times a particular Array index has been incremented without atomic add is: \n");for (int i = 0; i < SIZE; i++){printf("index: %d --> %d times\n ", i, h_a[i]);}cudaFree(d_a);return 0;}
可能如同已經猜到的那樣，每次運作你的程式，每個記憶體區域中的元素值都可能會不同。這是裝置上不定順序的多線程執行導緻的。
為了解決這個問題，CUDA提供了atomicAdd這種原子操作函數。該函數會從邏輯上保證，每個調用它的線程對相同的記憶體區域上的“讀取舊值-累加-回寫新值”操作是不可被其他線程擾亂的原子性的整體完成的。使用atomicAdd進行原子累加的核心函數代碼如下：
#include #define NUM_THREADS 10000#define SIZE 10#define BLOCK_WIDTH 100__global__ void gpu_increment_atomic(int *d_a){// Calculate thread id for current threadint tid = blockIdx.x * blockDim.x + threadIdx.x;// each thread increments elements wrapping at SIZE variabletid = tid % SIZE;atomicAdd(&d_a[tid], 1);}int main(int argc, char **argv){printf("%d total threads in %d blocks writing into %d array elements\n",NUM_THREADS, NUM_THREADS / BLOCK_WIDTH, SIZE);// declare and allocate host memoryint h_a[SIZE];const int ARRAY_BYTES = SIZE * sizeof(int);// declare and allocate GPU memoryint * d_a;cudaMalloc((void **)&d_a, ARRAY_BYTES);//Initialize GPU memory to zerocudaMemset((void *)d_a, 0, ARRAY_BYTES);gpu_increment_atomic << > >(d_a);// copy back the array to host memorycudaMemcpy(h_a, d_a, ARRAY_BYTES, cudaMemcpyDeviceToHost);printf("Number of times a particular Array index has been incremented is: \n");for (int i = 0; i < SIZE; i++) { printf("index: %d --> %d times\n ", i, h_a[i]); }cudaFree(d_a);return 0;}
在main函數中，具有10個元素的數組被初始化成0值，然後傳遞給了核心，但現在，核心中的代碼将執行原子累加操作。是以，這個程式輸出的結果将是對的，數組中的每個元素将被累加1000。運作結果顯示如圖：
如果你測量一下這個程式的運作時間，相比之前的那個簡單地在全局記憶體上直接進行加法操作的程式它用的時間更長。這是因為使用原子操作後程式具有更大的執行代價。可以通過使用共享記憶體來加速這些原子累加操作。如果線程規模不變，但原子操作的元素數量擴大，則這些同樣次數的原子操作會更快地完成。這是因為更廣泛的分布範圍上的原子操作有利于利用多個能執行原子操作的單元，以及每個原子操作單元上面的競争性的原子事務也相應減少了。
筆記來源書籍：
《基于GPU加速的計算機視覺程式設計》——使用OpenCV和CUDA實時處理複雜圖像資料
比較好的一本書，推薦給大家，希望可以幫助到你，後續小編也會持續摘錄本書内容以及調試好代碼并上傳，并在最後以視覺項目的形式進行收尾。小編也是第一次接觸CUDA程式設計，做筆記知識希望可以和大家一起交流學習。
關注【AI人工智能初學者】公衆号，回複【CUDA4】建議長按複制，即可獲得完整的項目代碼檔案。
希望您可以關注公衆号，也非常期待您的打賞。
聲明：轉載請說明出處
下方為小生公衆号，還望包容接納和關注，非常期待與您的美好相遇，讓我們以夢為馬，砥砺前行。
希望技術與靈魂可以一路同行
長按識别二維碼關注一下
更多精彩内容可回複關鍵詞
每篇文章的主題即可
點“在看”給我一朵小黃花