一：cuda程式設計模型
1：主機與裝置
主機---CPU 裝置/處理器---GPU
CUDA程式設計模型如下：
ＧＰＵ多層存儲空間結構如圖：
2：Kernel函數的定義與調用
A:運作在GPU上
相關限定符：__global__能在主機端和device端調用中調用；__device__隻能在device端調用
__host__隻能在host端調用 _host__device_:都可調用，單元測試是可隻使用一份代碼---應該編譯了兩份binary.
B:在調用時必須聲明核心函數的執行參數----<<<>>>。
C:先為核心函數中用到的變量配置設定好足夠空間再調用kernel函數
D:每個線程都有自己對應的id----由裝置端的寄存器提供的内建變量儲存，且是隻讀的。
E:CUDA C通過允許程式員定義稱為核心的 C函數來擴充C，這些函數在被調用時由N個不同的CUDA線程并行執行N次，而不是像正常C函數那樣僅執行一次。
3：線程結構
1）線程辨別
dim3類型（基于uint3定義的矢量類型----由三個unsigned int組成的結構體）的内建變量threadIdx和blockIdx。
2）一維block
線程threadID----threadIdx.x.
3）二維block---(Dx,Dy)
線程threadID----threadIdx.x+threadIdx.y*Dx;
4）三維block---(Dx,Dy,Dz)
線程threadID----threadIdx.x+threadIdx.y*Dx+threadIdx.z*Dx*Dy;
4：硬體映射
1）計算單元
SM---流多處理器 SP---流處理器
A：一個SM包含8個SP，共用一塊共享存儲器
2）warp
線程束在采用Tesla架構的gpu中：一個線程束由32個線程組成，且其線程隻和threadID有關
A：warp才是真正的執行機關，當在一個warp内不需要__syncwarp.
3）執行模型
SIMT---單指令多線程 SIMD---單指令多資料
4）deviceQuery執行個體
1 #include <stalib.h>
2 #include <stdio.h>
3 #include<string.h>
4 #include <cutil.h>
5
6 int main()
7 {
8 int deviceCount;
9 CUDA_SAFE_CALL(cudaGetDeviceCount(&deviceCount));
10 if(0 == deviceCount)
11 {
12 printf("no deice\n");
13 }
14 int dev;
15 for(dev = 0;dev <deviceCount;dev++)
16 {
17 cudaDeviceProp deviceProp;
18 CUDA_SAFE_CALL(cudaGetDeviceProperties(&deviceProp,dev));
19 print();
20 }
21 } View Code
5)cuda程式編寫流程
Ａ：主機端
1 啟動ＣＵＤＡ，使用多卡時需加上裝置号，或使用ｃｕｄａＳｅｔＤｅｖｉｃｅ（）設定
2 為輸入資料配置設定空間
3 初始化輸入資料
4 為ＧＰＵ配置設定顯存，用于存放輸入資料
5 将記憶體中的輸入資料拷貝到顯存
6 為ＧＰＵ配置設定顯存，用于存放輸出資料
7 調用ｄｅｖｉｃｅ端的ｋｅｒｎｅｌ進行計算，将結果寫到顯存中對應區域
8 為ＣＰＵ配置設定記憶體，用于存放ＧＰＵ傳回來的輸出資料
9 使用ＣＰＵ對資料進行其他處理
10 釋放記憶體和顯存空間
11 退出ＣＵＤＡ Ｂ：裝置端
1 從顯存讀資料到ＧＰＵ片内 2 對資料進行處理 3 将處理後的資料寫回顯存
（１）在顯存全局記憶體配置設定線性空間－－cudaMalloc()/cudaFree()
（２）拷貝存儲器中的資料 －－cudaMemcpy()
拷貝操作類型：cudaMemcpyDeiceToHost cudaMemcpyHostToDevice cudaMemcpyDeviceToDevice
（３）網格定義
<<<Dg,Db,Ns,S>>>
Dg----grid緯度與尺寸 Db---block次元與尺寸 Ns－－可配置設定動态共享記憶體大小 s－－stream_t類型的可選參數
（４）裝置端内建變量
gridDim blockIdx blockDim threadIdx warpSize
６）核心執行個體
Ａ：與shared memory有關-----shared memory 與opencl local memory 類似，隻對device可見，hardwork隻知道shared 大小且都認為從0開始編号
1 __global__ void
2 testKernel(float* g_idata,float* g_odata)
3 {
//配置設定共享記憶體 将全局記憶體的資料寫入共享記憶體 進行計算，将結果寫入共享記憶體 将結果寫回全局記憶體
4 extern __shared__ float sdata[];//動态配置設定共享記憶體空間－－__device__ __global__函數中
//動态配置設定大小是執行參數中的第三個參數。當靜态配置設定時必須指明大小
5
6 const unsigned int bid = blockIdx.x;
7 const unsigned int tid_in_block = threadIdx.x;
8 const unsigned int tid_in_grid = blockIdx.x*blockDim.x+threadIdx.x;
9 sdata[tid_in_block] = g_idata[tid_in_grid];
10 __syncthreads();
11
12 sdata[tid_in_block] *= (float)bid;
13
14 __syncthreads();
g_odata[tid_in_grid] = sdata[tid_in_block];
15 } B：靜态共享記憶體與二維網格
1 __gloabal__ void
2 testKernel(float* g_idata,int width,int height )
3 {
4 __shared__ float sdata[4];
5
6 //__mul24()-----cuda中兩數相乘函數
7 //block在網格中的索引-----如同元素在二維數組中的下标
8 //blockIdx.x----線程所在塊索引号
9 unsigned int bid_in_grid = __mul24(blockIdx.y,griDim.x)+blockIdx.x;
10
11 //線程線上程快中的索引---與上一個類似，但嵌套在網格中
12 //threadIdx.x---線程在塊中的索引号
13 unsigned int tid_in_block = __mul24(threadIdx.y,blockDim.x)+threadIdx.x;
14 unsigned int tid_in_grid_x = __mul24(blockDim.x,blockIdx.x)+threadIdx.x;
15 unsigned int tid_in_grid_y = __mul24(blockDim.y,blockIdx.y)+threadIdx.y;
16 unsigned int tid_in_grid = __mul24(tid_in_grid_y,width)+tid_in_grid_x;
17
18 SDATA(tid_in_block) = (float)bid_in_grid*SDATA(tid_in_block);
19 __syncthreads();
20
21 g_odata[tid__in_grid] = SDATA(tid_in_block);
22 __syncthreads();
23 } 二：CUDA軟體體系
1：CUDA C語言---我不熟悉的
A：blockIdx threadIdx-----索引線程塊和線程 gridDim blockDim---描述線程網格和線程塊的次元
warpSize----warp中的線程數量
B：初始化
在初始化過程中，運作時建立系統中的每個裝置。此上下文是此裝置的主要上下文，并且在應用程式的所有主機線程之間共享。作為此上下文建立的一部分，裝置代碼會在必要時進行實時編譯，并加載到裝置記憶體中。這一切都是在背景進行的，并且運作時不會向應用程式公開主要上下文。當主機線程調用 cudaDeviceReset（），這将破壞主機線程目前在其上操作的裝置的主要上下文。具有該裝置作為目前裝置的任何主機線程進行的下一個運作時函數調用都将為此裝置建立一個新的主上下文。
2：CUDA驅動API
A：調用任何一個驅動API函數之前，必須先調用cuInit()完成初始化，建立一個CUDA上下文。
B：上下文
封裝驅動程式API中執行的 所有資源和操作------管理相關資源，當上下文被銷毀時，系統自動清理這些資源。
一個主機端線程在一個時刻隻能擁有一個目前裝置上下文。
cuCtxCreate()建立上下文；cuCtxPopCurrent()---解除或恢複主機端線程與上下文關系
cuCtxAttach()---上下文使用計數遞增 cuCtxDetach()--使用計數遞減
cuCtxDetach()/cuCtxDestroy()---使用數為0是上下文被銷毀
計數目的-----實作在同一上下文中與第三方代碼進行互操作。當有庫使用時上下文計數加1，使用完後計數減1。特殊情況---庫單獨使用一個上下文：庫初始化調用-cuCtxCreate()-初始化上下文--cuCtxPopCurrent()将該庫上下文壓入。庫調用---cuCtxPushCurrent()---使用上下文--cuCtxPopCurrent()釋放上下文。
C：Kernel執行
cuFuncSetBlockShape()----為給定函數設定每個塊的線程數以及塊中threadID的設定方式
cuFuncSetSharedSize()----為函數設定每個block***享存儲器的大小
實作矢量加法
1 int main()
2 {
//初始化裝置
3 if(cuInit() != CUDA_SUCCESS)
4 exit(0);
5
//獲得支援裝置的數目
6 int deiceCount = 0;
7 cuDeviceGetCount(&deviceCount);
8 if(0 == deviceCount)
9 exit(0);
10
//獲得裝置0的句柄
11 CUdevice cuDevice = 0;
12 cuDeviceGet(&cuDevice,0);
13
//建立上下文
14 CUcontext cuContext;
15 cuCtxCreate(&cuContex,0,cuDevice);
16
//從二進制檔案生成模闆
17 CUmodule cuModule;
18 cuModuleLoad(&cuModule,"VecAdd.cubin");
//從模闆取得函數句柄
CUfunction vecAdd;
cuModuleGetFunction(&vecAdd,cuModule,"VecAdd");
//啟動Kernel
19 //cuParam*()函數用于指定調用cuLaunchGrid()/cuLaunch()調用啟動核心時為核心提供的參數
20 int threadsPerBlock = 256;
21 int threadsPerGrid = (N+threadsPerBlock-1)/threadsPerBlock;
22 int offset = 0;
23 cuParamSeti(vecAdd.offset.A);offset += sizeof(A);
24 cuParamSeti(vecAdd.offset.B);offset += sizeof(B);
25 cuParamSeti(vecAdd.offset.C);offset += sizeof(C);
26 cuParamSetSize(vecAdd.offset);
27 cuFuncSetBlockShape(vecAdd,threadsPerBlock,1,1);
28 cuLaunchGrid(vecAdd,threadsPerGrid,1);
29 } 三：CUDA存儲器模型
相關圖檔在随筆開頭已經給出；
| 存儲器 | 位置 | 擁有緩存 | 通路權限 | 變量生存周期 |
| register | GPU片内 | N/A | device可讀/寫 | 與thread相同 |
| local memory | 闆載顯存 | 無 | device可讀/寫 | 與thread相同 |
| shared memory | GPU片内 | N/A | device可讀/寫 | 與block相同 |
| constant memory | 闆載顯存 | 有 | device可讀，host可讀/寫 | 可在程式中保持 |
| global memory | 闆載顯存 | 無 | device可讀/寫，host可讀/寫 | 可在程式中保持 |
| texture memory | 闆載顯存 | 有 | device可讀，可讀/ | 可在程式中保持 |
| host memory | host記憶體 | 無 | host可讀/寫 | 可在程式中保持 |
| pinned memory | host記憶體 | 無 | host可讀/寫 | 可在程式中保持 |
1：裝置端記憶體
A：寄存器記憶體
每個寄存器檔案大小為32bits,當私有變量不大時，将其配置設定為寄存器變量，否則為局部變量--通路速度慢
B：常數存儲器
空間較小（64kb）
用于存放需要頻繁通路的隻讀參數
在所有函數外聲明定義
C；全局記憶體
runtime API 使用global memory
1 __constant__ float constData[256];
2 float data[256];
3 cudaMemcpyToSymbol(constData,data,sizeof(data));
4 cudaMemcpyFromSymbol(data,constData,sizeof(data));
5
6 __device__ float devData;
7 float value = 3.14;
8 cudaMemcpyToSymbol(devData,&value,sizeof(float));
9
10 __device__ float* devPointer;
11 float* ptr;
12 cudaMalloc(&ptr,256*sizeof(float));
13 cudaMemcpyToSymbol(devPointer,&ptr,256*sizeof(ptr));
14
15 //cudaGetSymbolAddress()用于檢索指向為全局記憶體空間中聲明的變量配置設定的記憶體的位址。配置設定的記憶體大小通過以下方式獲得 cudaGetSymbolSize（）。 D：執行個體
１）運作時ＡＰＩ完成資料計算
線性記憶體通常使用 cudaMalloc（） 并使用釋放 cudaFree（） 主機記憶體和裝置記憶體之間的資料傳輸通常使用 cudaMemcpy（）
1 __global__ void VecAdd(float* A,float* B,float* C)
2 {
3 int i = threadIdx.x;
4 if(i<N)
5 C[i] = A[i]+B[i];
6 }
7
8 int main()
9 {
10 //顯存中配置設定向量空間
11 size_t size = N*sizeof(float);
12 float* d_A;
13 cudaMalloc(（void**)&d_A,size);
14 float* d_B;
15 cudaMalloc((void**)&d_B,size);
16 flaot* d_C;
17 cudaMalloc((void**)&d_C,size);
18
19 //從記憶體向顯存拷貝向量
20 cudaMemcoy(d_A,h_A,size,cudaMemcpyHostToDevice);
21 cudaMemcpy(d_B,h_B,size,cudaMemcpyHostToDevice);
22
23 //啟動ｋｅｒｎｅｌ
24 int threadsPerBlock = 256;
25 int threadsPerGrid = (N+threadsPerBlock-1)/threadsPerBlock;
26 VecAdd<<<threadsPerGrid,threadsPerBlock>>>(d_A,d_B,d_C);
27
28 //從顯存向記憶體考回結果
29 cudaMemcpy(h_C,d_C,size,cudaMemcpyDeviceToHost);
30
31 //釋放緩存空間
32 cudaFree(d_A);
33 cudaFree(d_B);
34 cudaFree(d_C);
35 } 對于二三維數組使用cudaMallocPitch() cudaMalloc3D()配置設定線性存儲空間，cudaMemcpy2D() cudaMemcpy3D()進行拷貝。
２）驅動ＡＰＩ完成資料運算
1 int main()
2 {
3 //初始化裝置
4 if(cuInit(0) != CUDA_SUCCESS)
5 exit(0);
6
7 //獲得支援ｃｕｄａ的裝置數目
8 int deviceCount = 0;
9 cuDeviceGetCount(&deviceCount);
10 if(deviceCount == 0)
11 exit(0);
12
13 //獲得裝置０句柄
14 CUdevice cuDevice = 0;
15 cuDeviceGet(&cuDevice,0);
16
17 //建立上下文
18 CUcontext cuContext;
19 cuCtxCreate(&cuContext,0,cuDevice);
20
21 //從二進制檔案建立子產品
22
23 CUmoudle cuMoudle;
24 cuMoudleLoad(&cuMoudle,"VecAdd.cubin");
25
26 //從子產品獲得函數句柄
27 CUfunction VecAdd;
28 cuModuleGetFunction(&VecAdd,cuModule,"VecAdd");
29
30 //顯存中配置設定向量空間
31 size_t ize = N*sizeof(float);
32 CUdeviceptr d_A;
33 cuMalloc(&d_A,size);
34 CUdeviceptr d_B;
35 cuMalloc(&d_B,size);
36 CUdeviceptr d_C;
37 cuMalloc(&d_C,size);
38
39 //從記憶體向顯存拷貝向量
40 cuMemcpyHtoD(d_A,h_A,size);
41 cuMemcpyHtoD(d_B,h_B,size);
42
43 // 啟動ｋｅｒｎｅｌ
44 int threadsPerBlock = 256;
45 int threadsPerGrid = (N+threadsPerBlock-1)/threadsPerBlock;
46 int offset = 0;
47 cuParamSeti(VecAdd,offset,d_A);offset += sizeof(A);
48 cuParamSeti(VecAdd,offset,d_B);offset += sizeof(B);
49 cuParamSeti(VecAdd,offset,d_C);offset += sizeof(C);
50 cuParamSetSize(VecAdd,offset);
51 cuLaunchGrid(VecAdd,threadsPerGrid,1);
52
53 //從顯存向記憶體拷回結果
54 cuMemcpyDtoH(h_C,d_C,size);
55
56 //釋放顯存空間
57 cuMemFree(d_A);
58 cuMemFree(d_B);
59 CUMemFree(d_C);
60
61 } ２：主機端記憶體
可分頁記憶體－－通過作業系統ＡＰＩ配置設定的存儲器空間
頁鎖定記憶體－－保證存在于實體記憶體中，不被配置設定到虛拟記憶體中
１）頁鎖定記憶體
Ａ：運作時ＡＰＩ實作
cudaHostAlloc() cudaFreeHost()配置設定釋放pinned memory
<1>portable memory
cudaHostAllocPortable:可以讓控制不同ＧＰＵ的幾個ＣＰＵ線程共享同一塊pinned memory,減少ＣＰＵ線程間的資料傳輸和通信,頁鎖定記憶體預設為此标志
<2>write-combined Memory
cudaostAllocWriteCombined:減少緩存機制－對記憶體的監視，在總線傳輸期間不會被來自ＣＰＵ的監聽打斷。由于沒有緩存機制，ＣＰＵ在讀資料時速度有所降低，最好隻将從ＣＰＵ端隻寫的資料存放在此類記憶體。cache---當多個使用一個位址空間資料時，會多次重新整理，導緻浪費記憶體（當寫入一定量的資料一次性重新整理到記憶體，而寫記憶體不用cache可直接寫入memory）。
<3>mapped memory
cudaHostAllocMapped:兩個位址：主機端位址 裝置端位址
可以在核心程式中直接通路此類記憶體中的資料：zero-copy 經常用于量少資料
cudaHostAlloc()配置設定 cudaHostGetDevicePointer()擷取裝置端指針
cudaGetDeviceProperties()傳回的cuMapHostMemory屬性檢視是否支援此類記憶體
必須通過同步保證ＣＰＵ和ＧＰＵ對同一塊存儲器操作的順序一緻性－－流與事件等
當多個主機端線程操作一塊ｐｉｎｎｅｄ 記憶體時，每個線程必須擷取裝置端指針
必須調用cudaSetDeviceFlags()--cudaDeviceMapHost标志，再擷取裝置端指針
四：compiling ----暫時用不着
1:offline compilation
2:just in time compliation