CUDA 程式設計的基本模式

reproduced from: 

本文將介紹 cuda 程式設計的基本模式，所有 cuda 程式都基於此模式編寫，即使是呼叫庫，庫的底層也是這個模式實現的。

1. 定義需要在 device 端執行的核函式。( 函式宣告前加 _golbal_ 關鍵字 )

2. 在視訊記憶體中為待運算的資料以及需要存放結果的變數開闢視訊記憶體空間。( cudamalloc 函式實現 )

3. 將待運算的資料傳輸進視訊記憶體。( cudamemcpy，cublassetvector 等函式實現 )

4. 呼叫 device 端函式，同時要將需要為 device 端函式建立的塊數執行緒數等引數傳遞進 <<<>>>。( 注: <<<>>>下方編譯器可能顯示語法錯誤，不用管 )

5. 從視訊記憶體中獲取結果變數。( cudamemcpy，cublasgetvector 等函式實現 )

6. 釋放申請的視訊記憶體空間。( cudafree 實現 )

在 cuda 標準程式設計模式中，增加了一些程式設計規範，在這裡簡要說明：

函式宣告關鍵字：

1. __device__

表明此函式只能在 gpu 中被呼叫，在 gpu 中執行。這類函式只能被 __global__ 型別函式或 __device__ 型別函式呼叫。

2. __global__

表明此函式在 cpu 上呼叫，在 gpu 中執行。這也是以後會常提到的 "核心函式"，有時為了便於理解也稱 "device" 端函式。

3. __host__

表明此函式在 cpu 上呼叫和執行，這也是預設情況。

核心函式配置運算子 <<<>>> - 這個運算子在呼叫核心函式的時候使用，一般情況下傳遞進三個引數：

1. 塊數

2. 執行緒數

3. 共享記憶體大小 (此引數預設為0 )

核心函式中的幾個系統變數 - 這幾個變數可以在核心函式中使用，從而控制塊與執行緒的工作：

1. griddim：塊數

2. blockdim：塊中線程數

3. blockidx：塊編號 (0 - griddim-1)

4. threadidx：執行緒編號 (0 - blockdim-1)

該程式採用 cuda 並行化思想來對陣列進行求和：

1

//相關 cuda 庫
2 #include "
cuda_runtime.h
"3 #include "
cuda.h
"4 #include "
device_launch_parameters.h"5
6 #include 7 #include 8
9using
namespace
std;
1011
const
int n = 100;12
13//
塊數14
const
int block_data = 3
; 15
//各塊中的執行緒數
16const
int thread_data = 10
; 17
18//
cuda初始化函式
19bool
initcuda()
2031
32int
i;33
for (i=0; i)
3442}43
}4445if (i==devicecount)
4650
51 cudasetdevice(i); //
選定使用的顯示裝置
5253
return
exit_success;54}
5556
//此函式在主機端呼叫，裝置端執行。
57__global__ 
58static
void sum (int *data,int *result)
5972
73//
result 陣列存放各個執行緒的計算結果
74 result[bid*thread_data+tid] =sum; 75}
7677
intmain ()
7883 cout << "
成功建立 cuda 計算環境
"<< endl <8485
//建立，初始化，列印測試陣列
86int *data = new
int[n];
87 cout << "
測試矩陣: 
"<88for (int i=0; i)
8994 cout <9596
int *gpudata, *result; 
9798
//在視訊記憶體中為計算物件開闢空間
99 cudamalloc ((void**)&gpudata, sizeof(int)*n); 
100//
在視訊記憶體中為結果物件開闢空間
101 cudamalloc ((void**)&result, sizeof(int)*block_data*thread_data);
102103
//將陣列資料傳輸進視訊記憶體
104 cudamemcpy (gpudata, data, sizeof(int)*n, cudamemcpyhosttodevice); 
105//
呼叫 kernel 函式 - 此函式可以根據視訊記憶體位址以及自身的塊號，執行緒號處理資料。
106 sum<<0>>>(gpudata,result);
107108
//在記憶體中為計算物件開闢空間
109int *sumarray = new
int[thread_data*block_data];
110//
從視訊記憶體獲取處理的結果
111 cudamemcpy (sumarray, result, sizeof(int)*thread_data*block_data, cudamemcpydevicetohost);
112113
//釋放視訊記憶體
114cudafree (gpudata); 
115cudafree (result);
116117
//計算 gpu 每個執行緒計算出來和的總和
118int final_sum=0
;119
for (int i=0; i)
120123
124 cout << "
gpu 求和結果為: 
"<< final_sum <125126
//使用 cpu 對矩陣進行求和並將結果對照
127 final_sum = 0
;128
for (int i=0; i)
129132 cout << "
cpu 求和結果為: 
"<< final_sum <133134
getchar();
135136
return0;
137 }

1. 掌握本節知識的關鍵除了要掌握各個api，還要深刻理解核心函式中的塊及執行緒變數的控制，或者說施展 :) 

2. 一定要明確傳遞進 api 的是引數本身，還是引數的位址，這很關鍵。

UE基本模式

1 ue有五種狀態 a idle b rrc連線狀態 cell pch狀態 ura pch狀態 cell fach狀態 cell dch狀態 2 小區重選和切換 小區重選和切換是一對概念。在td系統中 a ue在idle狀態 cell pch狀態 ura pch狀態 cell fach狀態進行小區重...

面試的基本模式

面試是評價應聘者素質特徵的一種考試方式，根據招聘物件的水平，面試常採用不同的模式。面試的模式按應聘者的行為反應可分為言談面試和模擬操作面試。言談面試是通過主試與被試的口頭交流溝通，由主試提出問題，由被試口頭回答考察應聘者知識層次 業務能力 頭腦機敏性的一種測試方法。模擬操作面試是讓被試模擬在實際工作...

設計模式 3種基本模式

1 工廠模式 使用工廠方法或者類生成物件，而不是在 中直接new。這樣做可以更好地管理物件。2 單例模式 單例模式確保某個類只有乙個例項，而且自行例項化並向整個系統提供這個例項。namespace imooc class database 私有化 clone，防止被轉殖 private functi...