鑒於自己的畢設需要使用GPU CUDA這項技術，想找一本入門的教材，選擇了Jason Sanders等所著的書《CUDA By Example an Introduction to General Purpose GPU Programming》。這本書作為入門教材，寫的很不錯。自己覺得從理解與記憶的角度的出發，書中很多內容都可以被省略掉，於是就有了這篇博文。此博文記錄與總結此書的筆記和理解。注意本文並沒有按照書中章節的順序來寫。書中第8章圖像互操作性和第11章多GPU系統上的CUDA C，這兩章沒有看。等有時間了再看吧，趕緊碼字。
CUDA是什么
CUDA，Compute Unified Device Architecture的簡稱，是由NVIDIA公司創立的基於他們公司生產的圖形處理器GPUs（Graphics Processing Units,可以通俗的理解為顯卡）的一個並行計算平台和編程模型。
通過CUDA，GPUs可以很方便地被用來進行通用計算（有點像在CPU中進行的數值計算等等）。在沒有CUDA之前，GPUs一般只用來進行圖形渲染（如通過OpenGL，DirectX）。
開發人員可以通過調用CUDA的API，來進行並行編程，達到高性能計算目的。NVIDIA公司為了吸引更多的開發人員，對CUDA進行了編程語言擴展，如CUDA C/C++,CUDA Fortran語言。注意CUDA C/C++可以看作一個新的編程語言，因為NVIDIA配置了相應的編譯器nvcc,CUDA Fortran一樣。更多信息可以參考文獻。
64位Ubuntu12.04安裝CUDA5.5
具體步驟請點擊此處http://bookc.github.io/2014/05/08/my-summery-the-book-cuda-by-example-an-introduction-to-general-purpose-gpu-programming/。
[b]對CUDA C的個人懵懂感覺[/b]
如果粗暴的認為C語言工作的對象是CPU和內存條（接下來,稱為主機內存），那么CUDA C工作的的對象就是GPU及GPU上的內存（接下來,稱為設備內存），且充分利用了GPU多核的優勢及降低了並行編程的難度。一般通過C語言把數據從外界讀入，再分配數據，給CUDA C，以便在GPU上計算，然后再把計算結果返回給C語言，以便進一步工作，如進一步處理及顯示，或重復此過程。
主要概念與名稱
主機
將CPU及系統的內存（內存條）稱為主機。
設備
將GPU及GPU本身的顯示內存稱為設備。
線程(Thread)
一般通過GPU的一個核進行處理。（可以表示成一維，二維，三維，具體下面再細說）。
線程塊(Block)
1. 由多個線程組成（可以表示成一維，二維，三維，具體下面再細說）。
2. 各block是並行執行的，block間無法通信，也沒有執行順序。
3. 注意線程塊的數量限制為不超過65535（硬件限制）。
線程格(Grid)
由多個線程塊組成（可以表示成一維，二維，三維，具體下面再細說）。
線程束
在CUDA架構中，線程束是指一個包含32個線程的集合，這個線程集合被“編織在一起”並且“步調一致”的形式執行。在程序中的每一行，線程束中的每個線程都將在不同數據上執行相同的命令。
核函數（Kernel）
1. 在GPU上執行的函數通常稱為核函數。
2. 一般通過標識符__global__修飾，調用通過<<<參數1,參數2>>>，用於說明內核函數中的線程數量，以及線程是如何組織的。
3. 以線程格（Grid）的形式組織，每個線程格由若干個線程塊（block）組成，而每個線程塊又由若干個線程（thread）組成。
4. 是以block為單位執行的。
5. 叧能在主機端代碼中調用。
6. 調用時必須聲明內核函數的執行參數。
7. 在編程時，必須先為kernel函數中用到的數組或變量分配好足夠的空間，再調用kernel函數，否則在GPU計算時會發生錯誤，例如越界或報錯，甚至導致藍屏和死機。
|
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
|
/*
* @file_name HelloWorld.cu 后綴名稱.cu
*/
#include <stdio.h>
#include <cuda_runtime.h> //頭文件
//核函數聲明，前面的關鍵字__global__
__global__
void
kernel(
void
) {
}
int
main(
void
) {
//核函數的調用，注意<<<1,1>>>，第一個1，代表線程格里只有一個線程塊；第二個1，代表一個線程塊里只有一個線程。
kernel<<<1,1>>>();
printf
(
"Hello, World!\n"
);
return
0;
}
|
dim3結構類型
1. dim3是基亍uint3定義的矢量類型，相當亍由3個unsigned int型組成的結構體。uint3類型有三個數據成員unsigned int x; unsigned int y; unsigned int z;
2. 可使用亍一維、二維或三維的索引來標識線程，構成一維、二維或三維線程塊。
3. dim3結構類型變量用在核函數調用的<<<,>>>中。
4. 相關的幾個內置變量
4.1. threadIdx，顧名思義獲取線程thread的ID索引；如果線程是一維的那么就取threadIdx.x，二維的還可以多取到一個值threadIdx.y，以此類推到三維threadIdx.z。
4.2. blockIdx，線程塊的ID索引；同樣有blockIdx.x，blockIdx.y，blockIdx.z。
4.3. blockDim，線程塊的維度，同樣有blockDim.x，blockDim.y，blockDim.z。
4.4. gridDim，線程格的維度，同樣有gridDim.x，gridDim.y，gridDim.z。
5. 對於一維的block，線程的threadID=threadIdx.x。
6. 對於大小為（blockDim.x, blockDim.y）的 二維 block，線程的threadID=threadIdx.x+threadIdx.y*blockDim.x。
7. 對於大小為（blockDim.x, blockDim.y, blockDim.z）的 三維 block，線程的threadID=threadIdx.x+threadIdx.y*blockDim.x+threadIdx.z*blockDim.x*blockDim.y。
8. 對於計算線程索引偏移增量為已啟動線程的總數。如stride = blockDim.x * gridDim.x; threadId += stride。
函數修飾符
1. __global__，表明被修飾的函數在設備上執行，但在主機上調用。
2. __device__，表明被修飾的函數在設備上執行，但只能在其他__device__函數或者__global__函數中調用。
常用的GPU內存函數
cudaMalloc()
1. 函數原型： cudaError_t cudaMalloc (void **devPtr, size_t size)。
2. 函數用處：與C語言中的malloc函數一樣，只是此函數在GPU的內存你分配內存。
3. 注意事項：
3.1. 可以將cudaMalloc()分配的指針傳遞給在設備上執行的函數；
3.2. 可以在設備代碼中使用cudaMalloc()分配的指針進行設備內存讀寫操作；
3.3. 可以將cudaMalloc()分配的指針傳遞給在主機上執行的函數；
3.4. 不可以在主機代碼中使用cudaMalloc()分配的指針進行主機內存讀寫操作（即不能進行解引用）。
cudaMemcpy()
1. 函數原型：cudaError_t cudaMemcpy (void *dst, const void *src, size_t count, cudaMemcpyKind kind)。
2. 函數作用：與c語言中的memcpy函數一樣，只是此函數可以在主機內存和GPU內存之間互相拷貝數據。
3. 函數參數：cudaMemcpyKind kind表示數據拷貝方向，如果kind賦值為cudaMemcpyDeviceToHost表示數據從設備內存拷貝到主機內存。
4. 與C中的memcpy()一樣，以同步方式執行，即當函數返回時，復制操作就已經完成了，並且在輸出緩沖區中包含了復制進去的內容。
5. 相應的有個異步方式執行的函數cudaMemcpyAsync()，這個函數詳解請看下面的流一節有關內容。
cudaFree()
1. 函數原型：cudaError_t cudaFree ( void* devPtr )。
2. 函數作用：與c語言中的free()函數一樣，只是此函數釋放的是cudaMalloc()分配的內存。
下面實例用於解釋上面三個函數
|
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
|
#include <stdio.h>
#include <cuda_runtime.h>
__global__
void
add(
int
a,
int
b,
int
*c ) {
*c = a + b;
}
int
main(
void
) {
int
c;
int
*dev_c;
//cudaMalloc()
cudaMalloc( (
void
**)&dev_c,
sizeof
(
int
) );
//核函數執行
add<<<1,1>>>( 2, 7, dev_c );
//cudaMemcpy()
cudaMemcpy( &c, dev_c,
sizeof
(
int
),cudaMemcpyDeviceToHost ) ;
printf
(
"2 + 7 = %d\n"
, c );
//cudaFree()
cudaFree( dev_c );
return
0;
}
|
GPU內存分類
全局內存
通俗意義上的設備內存。
共享內存
1. 位置：設備內存。
2. 形式：關鍵字__shared__添加到變量聲明中。如__shared__ float cache[10]。
3. 目的：對於GPU上啟動的每個線程塊，CUDA C編譯器都將創建該共享變量的一個副本。線程塊中的每個線程都共享這塊內存，但線程卻無法看到也不能修改其他線程塊的變量副本。這樣使得一個線程塊中的多個線程能夠在計算上通信和協作。
常量內存
1. 位置：設備內存
2. 形式：關鍵字__constant__添加到變量聲明中。如__constant__ float s[10];。
3. 目的：為了提升性能。常量內存采取了不同於標准全局內存的處理方式。在某些情況下，用常量內存替換全局內存能有效地減少內存帶寬。
4. 特點：常量內存用於保存在核函數執行期間不會發生變化的數據。變量的訪問限制為只讀。NVIDIA硬件提供了64KB的常量內存。不再需要cudaMalloc()或者cudaFree(),而是在編譯時，靜態地分配空間。
5. 要求：當我們需要拷貝數據到常量內存中應該使用cudaMemcpyToSymbol()，而cudaMemcpy()會復制到全局內存。
6. 性能提升的原因：
6.1. 對常量內存的單次讀操作可以廣播到其他的“鄰近”線程。這將節約15次讀取操作。（為什么是15，因為“鄰近”指半個線程束，一個線程束包含32個線程的集合。）
6.2. 常量內存的數據將緩存起來，因此對相同地址的連續讀操作將不會產生額外的內存通信量。
紋理內存
1. 位置：設備內存
2. 目的：能夠減少對內存的請求並提供高效的內存帶寬。是專門為那些在內存訪問模式中存在大量空間局部性的圖形應用程序設計，意味着一個線程讀取的位置可能與鄰近線程讀取的位置“非常接近”。如下圖：
3. 紋理變量（引用）必須聲明為文件作用域內的全局變量。
4. 形式：分為一維紋理內存 和 二維紋理內存。
4.1. 一維紋理內存
4.1.1. 用texture<類型>類型聲明，如texture<float> texIn。
4.1.2. 通過cudaBindTexture()綁定到紋理內存中。
4.1.3. 通過tex1Dfetch()來讀取紋理內存中的數據。
4.1.4. 通過cudaUnbindTexture()取消綁定紋理內存。
4.2. 二維紋理內存
4.2.1. 用texture<類型,數字>類型聲明，如texture<float，2> texIn。
4.2.2. 通過cudaBindTexture2D()綁定到紋理內存中。
4.2.3. 通過tex2D()來讀取紋理內存中的數據。
4.2.4. 通過cudaUnbindTexture()取消綁定紋理內存。
固定內存
1. 位置：主機內存。
2. 概念：也稱為頁鎖定內存或者不可分頁內存，操作系統將不會對這塊內存分頁並交換到磁盤上，從而確保了該內存始終駐留在物理內存中。因此操作系統能夠安全地使某個應用程序訪問該內存的物理地址，因為這塊內存將不會破壞或者重新定位。
3. 目的：提高訪問速度。由於GPU知道主機內存的物理地址，因此可以通過“直接內存訪問DMA（Direct Memory Access)技術來在GPU和主機之間復制數據。由於DMA在執行復制時無需CPU介入。因此DMA復制過程中使用固定內存是非常重要的。
4. 缺點：使用固定內存，將失去虛擬內存的所有功能；系統將更快的耗盡內存。
5. 建議：對cudaMemcpy()函數調用中的源內存或者目標內存，才使用固定內存，並且在不再需要使用它們時立即釋放。
6. 形式：通過cudaHostAlloc()函數來分配；通過cudaFreeHost()釋放。
7. 只能以異步方式對固定內存進行復制操作。
原子性
1. 概念：如果操作的執行過程不能分解為更小的部分，我們將滿足這種條件限制的操作稱為原子操作。
2. 形式：函數調用，如atomicAdd（addr,y)將生成一個原子的操作序列，這個操作序列包括讀取地址addr處的值，將y增加到這個值，以及將結果保存回地址addr。
常用線程操作函數
1. 同步方法__syncthreads()，這個函數的調用，將確保線程塊中的每個線程都執行完__syscthreads()前面的語句后，才會執行下一條語句。
使用事件來測量性能
1. 用途：為了測量GPU在某個任務上花費的時間。CUDA中的事件本質上是一個GPU時間戳。由於事件是直接在GPU上實現的。因此不適用於對同時包含設備代碼和主機代碼的混合代碼設計。
2. 形式：首先創建一個事件，然后記錄事件，再計算兩個事件之差，最后銷毀事件。如：
|
1
2
3
4
5
6
7
8
9
10
|
cudaEvent_t start, stop;
cudaEventCreate( &start );
cudaEventCreate( &stop );
cudaEventRecord( start, 0 );
//do something
cudaEventRecord( stop, 0 );
float
elapsedTime;
cudaEventElapsedTime( &elapsedTime,start, stop );
cudaEventDestroy( start );
cudaEventDestroy( stop )；
|
流
1. 扯一扯：並發重點在於一個極短時間段內運行多個不同的任務；並行重點在於同時運行一個任務。
2. 任務並行性：是指並行執行兩個或多個不同的任務，而不是在大量數據上執行同一個任務。
3. 概念：CUDA流表示一個GPU操作隊列，並且該隊列中的操作將以指定的順序執行。我們可以在流中添加一些操作，如核函數啟動，內存復制以及事件的啟動和結束等。這些操作的添加到流的順序也是它們的執行順序。可以將每個流視為GPU上的一個任務，並且這些任務可以並行執行。
4. 硬件前提：必須是支持設備重疊功能的GPU。支持設備重疊功能，即在執行一個核函數的同時，還能在設備與主機之間執行復制操作。
5. 聲明與創建：聲明cudaStream_t stream;，創建cudaSteamCreate(&stream);。
6. cudaMemcpyAsync()：前面在cudaMemcpy()中提到過，這是一個以異步方式執行的函數。在調用cudaMemcpyAsync()時，只是放置一個請求，表示在流中執行一次內存復制操作，這個流是通過參數stream來指定的。當函數返回時，我們無法確保復制操作是否已經啟動，更無法保證它是否已經結束。我們能夠得到的保證是，復制操作肯定會當下一個被放入流中的操作之前執行。傳遞給此函數的主機內存指針必須是通過cudaHostAlloc()分配好的內存。（流中要求固定內存）
7. 流同步：通過cudaStreamSynchronize()來協調。
8. 流銷毀：在退出應用程序之前，需要銷毀對GPU操作進行排隊的流，調用cudaStreamDestroy()。
9. 針對多個流：
9.1. 記得對流進行同步操作。
9.2. 將操作放入流的隊列時，應采用寬度優先方式，而非深度優先的方式，換句話說，不是首先添加第0個流的所有操作，再依次添加后面的第1，2,…個流。而是交替進行添加，比如將a的復制操作添加到第0個流中，接着把a的復制操作添加到第1個流中，再繼續其他的類似交替添加的行為。
9.3. 要牢牢記住操作放入流中的隊列中的順序影響到CUDA驅動程序調度這些操作和流以及執行的方式。
技巧
1. 當線程塊的數量為GPU中處理數量的2倍時，將達到最優性能。
2. 核函數執行的第一個計算就是計算輸入數據的偏移。每個線程的起始偏移都是0到線程數量減1之間的某個值。然后，對偏移的增量為已啟動線程的總數。
實例程序
感興趣的讀者可以下載本書附帶的示例代碼點擊此處下載https://developer.nvidia.com/sites/default/files/akamai/cuda/files/cuda_by_example.zip。