《GPU高效能程式設計 CUDA實戰》(CUDA By Example)讀書筆記
寫在最前 這本書是2011年出版的，按照計算機的發展速度來說已經算是上古書籍了，不過由於其簡單易懂，仍舊被推薦為入門神書。先上封面： 由於書比較老，而且由於學習的目的不同，這裡只介紹了基礎程式碼相關的內容，跳過了那些影象處理的內容。 另外這本書的程式碼這裡：csdn資源
前兩章 科普 就各種講CUDA的變遷，然後第二章講如何安裝CUDA。不會安裝的請移步這裡:安裝CUDA.
第三章 CUDA C簡介 輸出hello world
#include<stdio.h>
__global__ void kernel() { printf("hello world"); }
int main() { kernel<<<1, 1>>>(); return 0; } 1 2 3 4 5 6 7 8 9 10 11 12 這個程式和普通的C程式的區別值得注意
函式的定義帶有了__global__這個標籤，表示這個函式是在GPU上執行 函式的呼叫除了常規的引數之外，還增加了<<<>>>修飾。而其中的數字將傳遞個CUDA的執行時系統，至於能幹啥，下一章會講。 進階版
#include<stdio.h>
__global__ void add(int a,int b,int *c){ *c = a + b; } int main(){ int c; int *dev_c; cudaMalloc((void**)&dev_c,sizeof(int)); add<<<1,1>>>(2,7,dev_c); cudaMemcpy(&c,dev_c,sizeof(int),cudaMemcpyDeviceToHost); printf("2 + 7 = %d",c); return 0; } 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 這裡就涉及了GPU和主機之間的記憶體交換了，cudaMalloc是在GPU的記憶體裡開闢一片空間，然後通過操作之後，這個記憶體裡有了計算出來內容，再通過cudaMemcpy這個函式把內容從GPU複製出來。就是這麼簡單。
第四章 CUDA C並行程式設計 這一章開始體現CUDA並行程式設計的魅力。 以下是一個數組求和的程式碼
#include<stdio.h>
#define N 10
__global__ void add( int *a, int *b, int *c ) { int tid = blockIdx.x; // this thread handles the data at its thread id if (tid < N) c[tid] = a[tid] + b[tid]; }
int main( void ) { int a[N], b[N], c[N]; int *dev_a, *dev_b, *dev_c;
// allocate the memory on the GPU cudaMalloc( (void**)&dev_a, N * sizeof(int) ); cudaMalloc( (void**)&dev_b, N * sizeof(int) ); cudaMalloc( (void**)&dev_c, N * sizeof(int) );
// fill the arrays 'a' and 'b' on the CPU for (int i=0; i<N; i++) { a[i] = -i; b[i] = i * i; }
// copy the arrays 'a' and 'b' to the GPU cudaMemcpy( dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice ); cudaMemcpy( dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice );
add<<<N,1>>>( dev_a, dev_b, dev_c );
// copy the array 'c' back from the GPU to the CPU cudaMemcpy( c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost );
// display the results for (int i=0; i<N; i++) { printf( "%d + %d = %d\n", a[i], b[i], c[i] ); }
// free the memory allocated on the GPU cudaFree( dev_a ); cudaFree( dev_b ); cudaFree( dev_c ); return 0; } 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 重點也是對於初學者最難理解的就是kernel函數了：
__global__ void add( int *a, int *b, int *c ) { int tid = blockIdx.x; if (tid < N) c[tid] = a[tid] + b[tid]; } 1 2 3 4 5 GPU程式設計和CPU程式設計的最大區別也就在這裡體現出來了，就是陣列求和竟然不要迴圈！為什麼不要迴圈，就是因為這裡的tid可以把整個迴圈的工作做了。這裡的tid也就是thread的id，每個thread負責陣列一個數的操作，所以將10個迴圈操作拆分成了十個執行緒同時搞定。這裡的kernel函式也就是可以同時併發執行，而裡面的tid的數值是不一樣的。
第五章 執行緒協作 GPU邏輯結構 這章就開始介紹執行緒塊和網格的相關知識了，也就是<<<>>>這裡面數字的含義。首先講一下什麼叫執行緒塊，顧名思義就是執行緒組成的塊咯。GPU的邏輯結構如下圖所示： 這個圖來自NVIDIA官方文件，其中CTA就是執行緒塊，Grid就是執行緒塊組成的網格，每個執行緒塊裡有若干執行緒束warp，然後執行緒束內有最小的單位執行緒(文件裡會稱其為lanes，翻譯成束內執行緒)。 基礎知識稍微介紹一下，就開始介紹本章的內容了，本章的內容主要基於以下這個事實:
我們注意到硬體將執行緒塊的數量限制為不超過65535.同樣，對於啟動核函式每個執行緒塊中的執行緒數量，硬體也進行了限制。
由於這種限制的存在，我們就需要一些更復雜的組合來操作更大長度的陣列，而不僅僅是使用threadIdx這種naive的東西了。 我們提供了以下的kernel來操作比較長的陣列：
__global__ void add(int *a, int *b, int *c) { int tid = threadIdx.x + blockIdx.x * blockDim.x; while (tid < N) { c[tid] = a[tid] + b[tid]; tid += blockDim.x * gridDim.x; } } 1 2 3 4 5 6 7 嗯，理解透了int tid = threadIdx.x + blockIdx.x * blockDim.x;這句話，這章就算勝利完工了。首先，為啥是x，那有沒有y,z呢，答案是肯定的，但是這裡（對，就這本書裡），用不上。其實執行緒塊和網格都並不是只有一維，執行緒塊其實有三個維度，而網格也有兩個維度。因此存在.x的現象。當然我們不用管這些事，就當做它們只有一維好了。那就看下面這個圖：
這就是隻有一維的執行緒網格。其中，threadIdx.x就是每個執行緒在各自執行緒塊中的編號，也就是圖中的thread 0,thread 1。但是問題在於，每個block中都有thread 0，但是想讓這不同的thread 0操作不同的位置應該怎麼辦。引入了blockIdx.x，這個就表示了執行緒塊的標號，有了執行緒塊的標號，再乘上每個執行緒塊中含有執行緒的數量blockDim.x，就可以給每個執行緒賦予依次遞增的標號了，程式猿們就可以操作比較長的陣列下標了。
但是問題又來了，要是陣列實在太大，我用上所有的執行緒都沒辦法一一對應咋辦，這裡就用tid += blockDim.x * gridDim.x;這句話來讓一個執行緒操作很好幾個下標。具體是怎麼實現的呢，就是在處理過當前的tid位置後，讓tid增加所以執行緒的數量，blockDim.x是一塊中執行緒總數，而gridDim.x則是一個網格中所有塊的數量，這樣乘起來就是所有執行緒的數量了。
至此，執行緒協作也講完了。再上一個更直觀的圖：
共享記憶體 共享記憶體是個好東西，它只能在block內部使用，訪問速度巨快無比，好像是從離運算器最近的L1 cache中分割了一部分出來給的共享記憶體，因此巨快。所以我們要把這玩意用起來。 這裡的例子是點積的例子，就是： 最後得到一個和。主要思想如下：
前一半加後一半：
要同步，別浪 把最後的並行度小的工作交給CPU 具體程式碼是醬嬸兒的： __global__ void dot(float *a, float *b, float *c) { //建立一個thread數量大小的共享記憶體陣列 __shared__ float cache[threadsPerBlock]; int tid = threadIdx.x + blockIdx.x * blockDim.x; int cacheIndex = threadIdx.x; float temp = 0; while (tid < N) { temp += a[tid] * b[tid]; tid += blockDim.x * gridDim.x; } //把算出的數存到cache裡 cache[cacheIndex] = temp; //這裡的同步，就是說所有的thread都要達到這裡之後程式才會繼續執行 __syncthreads(); //下面的程式碼必須保證執行緒數量的2的指數，否則總除2會炸的 int i = blockDim.x / 2; while (i != 0) { if (cacheIndex < i) cache[cacheIndex] += cache[cacheIndex + i]; //這裡這個同步保證了0號執行緒不要一次浪到底就退出執行了，一定要等到都算好才行 __syncthreads(); i /= 2; } if (cacheIndex == 0) c[blockIdx.x] = cache[0]; } 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 其中這個陣列c其實只是所以結果中的一部分，最後會返回block數量個c，然後由cpu執行最後的加法就好了。
第九章 原子性操作 原子性操作，就是，像作業系統的PV操作一樣，同時只能有一個執行緒進行。好處自然是不會產生同時讀寫造成的錯誤，壞處顯而易見是增加了程式執行的時間。
計算直方圖 原理：假設我們要統計資料範圍是[0,255]，因此我們定義一個unsigned int histo[256]陣列，然後我們的資料是data[N]，我們遍歷data陣列，然後histo[data[i]]++，就可以在最後計算出直方圖了。這裡我們引入了原子操作
__global__ void histo_kernel(unsigned char *buffer, long size, unsigned int *histo) { int i = threadIdx.x + blockIdx.x * blockDim.x; int stride = blockDim.x * gridDim.x; while (i < size) { atomicAdd(&(histo[buffer[i]]), 1); i += stride; } } 1 2 3 4 5 6 7 8 9 這裡的atomicAdd就是同時只能有一個執行緒操作，防止了其他執行緒的騷操作。但是，巨慢，書裡說自從服用了這個，竟然比CPU慢四倍。因此我們需要別的。
升級版計算直方圖 使用原子操作很慢的原因就在於，當資料量很大的時候，會同時有很多對於一個數據位的操作，這樣操作就在排隊，而這次，我們先規定執行緒塊內部有256個執行緒(這個數字不一定)，然後線上程內部定義一個臨時的共享記憶體儲存臨時的直方圖，然後最後再將這些臨時的直方圖加總。這樣衝突的範圍從全域性的所有的執行緒，變成了執行緒塊內的256個執行緒，而且由於也就256個數據位，這樣造成的資料衝突會大大減小。具體見以下程式碼：
__global__ void histo_kernel(unsigned char *buffer, long size, unsigned int *histo) { __shared__ unsigned int temp[256]; temp[threadIdx.x] = 0; //這裡等待所有執行緒都初始化完成 __syncthreads(); int i = threadIdx.x + blockIdx.x * blockDim.x; int offset = blockDim.x * gridDim.x; while (i < size) { atomicAdd(&temp[buffer[i]], 1); i += offset; } __syncthreads(); //等待所有執行緒完成計算，講臨時的內容加總到總的直方圖中 atomicAdd(&(histo[threadIdx.x]), temp[threadIdx.x]); } 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 第十章 流 頁鎖定記憶體 這種記憶體就是在你申請之後，鎖定到了主機記憶體裡，它的實體地址就固定不變了。這樣訪問起來會讓效率增加。 CUDA流 流的概念就如同java裡多執行緒的概念一樣，你可以把不同的工作放入不同的流當中，這樣可以併發執行一些操作，比如在記憶體複製的時候執行kernel: 文後講了一些優化的方法，但是親測無效啊，可能是cuda對於流的支援方式變了，關於流的知識會在以後的博文裡再提及。 十一章 多GPU 這章主要看了是第一節零拷貝記憶體，也十分好理解就是，在CPU上開闢一片記憶體，而GPU可以直接訪問而不用複製到GPU的視訊記憶體裡。至於和頁鎖定記憶體效能上的差距和區別，需要實驗來驗證
===================2017.7.30更新======================== 在閱讀程式碼時發現有三種函式字首： （1）__host__ int foo(int a){}與C或者C++中的foo(int a){}相同，是由CPU呼叫，由CPU執行的函式 （2）__global__ int foo(int a){}表示一個核心函式，是一組由GPU執行的平行計算任務，以foo<<>>(a)的形式或者driver API的形式呼叫。目前global函式必須由CPU呼叫，並將平行計算任務發射到GPU的任務呼叫單元。隨著GPU可程式設計能力的進一步提高，未來可能可以由GPU呼叫。 （3）__device__ int foo(int a){}則表示一個由GPU中一個執行緒呼叫的函式。由於Tesla架構的GPU允許執行緒呼叫函式，因此實際上是將__device__ 函式以__inline形式展開後直接編譯到二進位制程式碼中實現的，並不是真正的函式。
具體來說，device字首定義的函式只能在GPU上執行，所以device修飾的函式裡面不能呼叫一般常見的函式；global字首，CUDA允許能夠在CPU，GPU兩個裝置上執行，但是也不能執行CPU裡常見的函式；host字首修飾的事普通函式，預設預設，可以呼叫普通函式。 --------------------- 作者：FishSeeker 來源：CSDN 原文：https://blog.csdn.net/fishseeker/article/details/75093166 版權宣告：本文為博主原創文章，轉載請附上博文連結！