CUDA之CUDA編程模型概述（二）

LibraryPKU 2018-03-28

展開(kāi)全文

Abstract: 本文繼續(xù)上文介紹CUDA編程模型關(guān)于核函數(shù)以及錯(cuò)誤處理部分
Keywords: CUDA核函數(shù)，CUDA錯(cuò)誤處理

開(kāi)篇廢話

今天的廢話就是人的性格一旦形成，那么就會(huì)成為最大的指向標(biāo)，或者說(shuō)一個(gè)人的性格思維方式能夠決定這個(gè)人的全部生命軌跡，比如有人真的愛(ài)學(xué)習(xí)（比如我，嘻嘻嘻）有人真的不愛(ài)學(xué)習(xí)，沒(méi)有優(yōu)劣，只是兩種生活態(tài)度，因?yàn)閷W(xué)習(xí)這個(gè)事你學(xué)一輩子也學(xué)不完人類智慧的九牛一毛，而不學(xué)習(xí)可以有更多的時(shí)間進(jìn)行社會(huì)實(shí)踐，融入社會(huì)，榮華富貴，享受生命。這是兩種性格，沒(méi)有好壞，畢竟每個(gè)人評(píng)價(jià)生活的標(biāo)尺不一。努力追求自己想要的，不要嘲笑別人所追求的，能一起走的就一起走，不能一起走的就各自安好，這就是目前我理解的聲明的真諦。
繼續(xù)CUDA編程模型的后半部分，關(guān)于核函數(shù)以及錯(cuò)誤處理。

核函數(shù)

啟動(dòng)核函數(shù)
編寫(xiě)核函數(shù)
驗(yàn)證核函數(shù)

錯(cuò)誤處理

核函數(shù)概述

核函數(shù)就是在CUDA模型上諸多線程中運(yùn)行的那段串行代碼，這段代碼在設(shè)備上運(yùn)行，用NVCC編譯，產(chǎn)生的機(jī)器碼是GPU的機(jī)器碼，所以我們寫(xiě)CUDA程序就是寫(xiě)核函數(shù)，第一步我們要確保核函數(shù)能正確的運(yùn)行產(chǎn)生正切的結(jié)果，第二優(yōu)化CUDA程序的部分，無(wú)論是優(yōu)化算法，還是調(diào)整內(nèi)存結(jié)構(gòu)，線程結(jié)構(gòu)都是要調(diào)整核函數(shù)內(nèi)的代碼，來(lái)完成這些優(yōu)化的。
我們一直把我們的CPU當(dāng)做一個(gè)控制者，運(yùn)行核函數(shù)，要從CPU發(fā)起，那么我們開(kāi)始學(xué)習(xí)如何啟動(dòng)一個(gè)核函數(shù)

啟動(dòng)核函數(shù)

啟動(dòng)核函數(shù)，通過(guò)的以下的ANSI C 擴(kuò)展出的CUDA C指令：

1	kernel_name<<<grid,block>>>(argument list);

其標(biāo)準(zhǔn)C的原型就是C語(yǔ)言函數(shù)調(diào)用

1	function_name(argument list);

這個(gè)三個(gè)尖括號(hào)’>>’內(nèi)是對(duì)設(shè)備代碼執(zhí)行的線程結(jié)構(gòu)的配置（或者簡(jiǎn)稱為對(duì)內(nèi)核進(jìn)行配置），也就是我們上一篇中提到的線程結(jié)構(gòu)中的網(wǎng)格，塊。回憶一下上文，我們通過(guò)CUDA C內(nèi)置的數(shù)據(jù)類型dim3類型的變量來(lái)配置grid和block（上文提到過(guò)：在設(shè)備端訪問(wèn)grid和block屬性的數(shù)據(jù)類型是uint3不能修改的常類型結(jié)構(gòu)，這里反復(fù)強(qiáng)調(diào)一下）。
通過(guò)指定grid和block的維度，我們可以配置：

內(nèi)核中線程的數(shù)目
內(nèi)核中使用的線程布局

我們可以使用dim3類型的grid維度和block維度配置內(nèi)核，也可以使用int類型的變量，或者常量直接初始化：

1	kernel_name<<<4,8>>>(argument list);

上面這條指令的線程布局是：

我們的核函數(shù)是同時(shí)復(fù)制到多個(gè)線程執(zhí)行的，上文我們說(shuō)過(guò)一個(gè)對(duì)應(yīng)問(wèn)題，多個(gè)計(jì)算執(zhí)行在一個(gè)數(shù)據(jù)，肯定是浪費(fèi)時(shí)間，所以為了讓多線程按照我們的意愿對(duì)應(yīng)到不同的數(shù)據(jù)，就要給線程一個(gè)唯一的標(biāo)識(shí)，由于設(shè)備內(nèi)存是線性的（基本市面上的內(nèi)存硬件都是線性形式存儲(chǔ)數(shù)據(jù)的）我們觀察上圖，可以用threadIdx.x 和blockIdx.x 來(lái)組合獲得對(duì)應(yīng)的線程的唯一標(biāo)識(shí)（后面我們會(huì)看到，threadIdx和blockIdx能組合出很多不一樣的效果）

接下來(lái)我們就是修改代碼的時(shí)間了，改變核函數(shù)的配置，產(chǎn)生運(yùn)行出結(jié)果一樣，但效率不同的代碼：

一個(gè)塊：

1	kernel_name<<<1,32>>>(argument list);

32個(gè)塊

1	kernel_name<<<32,1>>>(argument list);

上述代碼如果沒(méi)有特殊結(jié)構(gòu)在核函數(shù)中，執(zhí)行結(jié)果應(yīng)該一致，但是有些效率會(huì)一直比較低。

上面這些是啟動(dòng)部分，當(dāng)主機(jī)啟動(dòng)了核函數(shù)，控制權(quán)馬上回到主機(jī)，而不是主機(jī)等待設(shè)備完成核函數(shù)的運(yùn)行，這一點(diǎn)我們上一篇文章也有提到過(guò)（就是等待hello world輸出的那段代碼后面要加一句）

想要主機(jī)等待設(shè)備端執(zhí)行可以用下面這個(gè)指令：

1	cudaError_t cudaDeviceSynchronize(void);

這是一個(gè)顯示的方法，對(duì)應(yīng)的也有隱式方法，隱式方法就是不明確說(shuō)明主機(jī)要等待設(shè)備端，而是設(shè)備端不執(zhí)行完，主機(jī)沒(méi)辦法進(jìn)行，比如內(nèi)存拷貝函數(shù)：

1 2	cudaError_t cudaMemcpy(void* dst,const void * src, size_t count,cudaMemcpyKind kind);

這個(gè)函數(shù)上文已經(jīng)介紹過(guò)了，當(dāng)核函數(shù)啟動(dòng)后的下一條指令就是從設(shè)備復(fù)制數(shù)據(jù)回主機(jī)端，那么主機(jī)端必須要等待設(shè)備端計(jì)算完成。

所有CUDA核函數(shù)的啟動(dòng)都是異步的，這點(diǎn)與C語(yǔ)言是完全不同的

編寫(xiě)核函數(shù)

我們會(huì)啟動(dòng)核函數(shù)了，但是核函數(shù)哪里來(lái)的？當(dāng)然我們寫(xiě)的，核函數(shù)也是一個(gè)函數(shù)，但是聲明核函數(shù)有一個(gè)比較模板化的方法：

1	__global__ void kernel_name(argument list);

注意：聲明和定義是不同的，這點(diǎn)CUDA與C語(yǔ)言是一致的

在C語(yǔ)言函數(shù)前沒(méi)有的限定符global ,CUDA C中還有一些其他我們?cè)贑中沒(méi)有的限定符，如下：

限定符	執(zhí)行	調(diào)用	備注
__global__	設(shè)備端執(zhí)行	可以從主機(jī)調(diào)用也可以從計(jì)算能力3以上的設(shè)備調(diào)用	必須有一個(gè)void的返回類型
__device__	設(shè)備端執(zhí)行	設(shè)備端調(diào)用
__host__	主機(jī)端執(zhí)行	主機(jī)調(diào)用	可以省略

而且這里有個(gè)特殊的情況就是有些函數(shù)可以同時(shí)定義為 device 和 host ，這種函數(shù)可以同時(shí)被設(shè)備和主機(jī)端的代碼調(diào)用，主機(jī)端代碼調(diào)用函數(shù)很正常，設(shè)備端調(diào)用函數(shù)與C語(yǔ)言一致，但是要聲明成設(shè)備端代碼，告訴nvcc編譯成設(shè)備機(jī)器碼，同時(shí)聲明主機(jī)端設(shè)備端函數(shù)，那么就要告訴編譯器，生成兩份不同設(shè)備的機(jī)器碼。

Kernel核函數(shù)編寫(xiě)有以下限制

只能訪問(wèn)設(shè)備內(nèi)存
必須有void返回類型
不支持可變數(shù)量的參數(shù)
不支持靜態(tài)變量
顯示異步行為

并行程序中經(jīng)常的一種現(xiàn)象：把串行代碼并行化時(shí)對(duì)串行代碼塊for的操作，也就是把for并行化。
例如：
串行：

1

2

3

4

void sumArraysOnHost(float *A, float *B, float *C, const int N) {

for (int i = 0; i < n;="">

C[i] = A[i] + B[i];

}

并行：

1

2

3

4

__global__ void sumArraysOnGPU(float *A, float *B, float *C) {

int i = threadIdx.x;

C[i] = A[i] + B[i];

}

這兩個(gè)簡(jiǎn)單的段不能執(zhí)行，但是我們可以大致的看一下for展開(kāi)并行化的樣子，沒(méi)吃過(guò)雞肉的時(shí)候可以先看看雞跑（我的博客是清真的）。

驗(yàn)證核函數(shù)

驗(yàn)證核函數(shù)就是驗(yàn)證其正確性，下面這段代碼上文出現(xiàn)過(guò)，但是同樣包含驗(yàn)證核函數(shù)的方法：
代碼庫(kù)：https://github.com/Tony-Tan/CUDA_Freshman

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

49

50

51

52

53

54

55

56

57

58

59

60

61

62

63

64

65

66

67

68

69

70

/*

* https://github.com/Tony-Tan/CUDA_Freshman

* 3_sum_arrays

*/

#include

#include 'freshman.h'

void sumArrays(float * a,float * b,float * res,const int size)

{

for(int i=0;i<>4)

{

res[i]=a[i]+b[i];

res[i+1]=a[i+1]+b[i+1];

res[i+2]=a[i+2]+b[i+2];

res[i+3]=a[i+3]+b[i+3];

}

__global__ void sumArraysGPU(float*a,float*b,float*res)

{

int i=threadIdx.x;

res[i]=a[i]+b[i];

}

int main(int argc,char **argv)

{

int dev = 0;

cudaSetDevice(dev);

int nElem=32;

printf('Vector size:%d\n',nElem);

int nByte=sizeof(float)*nElem;

float *a_h=(float*)malloc(nByte);

float *b_h=(float*)malloc(nByte);

float *res_h=(float*)malloc(nByte);

float *res_from_gpu_h=(float*)malloc(nByte);

memset(res_h,0,nByte);

memset(res_from_gpu_h,0,nByte);

float *a_d,*b_d,*res_d;

CHECK(cudaMalloc((float**)&a_d,nByte));

CHECK(cudaMalloc((float**)&b_d,nByte));

CHECK(cudaMalloc((float**)&res_d,nByte));

initialData(a_h,nElem);

initialData(b_h,nElem);

CHECK(cudaMemcpy(a_d,a_h,nByte,cudaMemcpyHostToDevice));

CHECK(cudaMemcpy(b_d,b_h,nByte,cudaMemcpyHostToDevice));

dim3 block(nElem);

dim3 grid(nElem/block.x);

sumArraysGPU>>(a_d,b_d,res_d);

printf('Execution configuration<<%d,%d>>>\n',block.x,grid.x);

check(cudamemcpy(res_from_gpu_h,res_d,nbyte,cudamemcpydevicetohost));

sumarrays(a_h,b_h,res_h,nelem);

checkresult(res_h,res_from_gpu_h,nelem);

cudafree(a_d);

cudafree(b_d);

cudafree(res_d);

free(a_h);

free(b_h);

free(res_h);

free(res_from_gpu_h);

return 0;

}

在開(kāi)發(fā)階段，每一步都進(jìn)行驗(yàn)證是絕對(duì)高效的，比把所有功能都寫(xiě)好，然后進(jìn)行測(cè)試這種過(guò)程效率高很多，同樣寫(xiě)cuda也是這樣的每個(gè)代碼小塊都進(jìn)行測(cè)試，看起來(lái)慢，實(shí)際會(huì)提高很多效率。
cuda小技巧，當(dāng)我們進(jìn)行調(diào)試的時(shí)候可以把核函數(shù)配置成單線程的：

1	kernel_name<<<1,1>>>(argument list)

錯(cuò)誤處理

所有編程都需要對(duì)錯(cuò)誤進(jìn)行處理，早起的編碼錯(cuò)誤，編譯器會(huì)幫我們搞定，內(nèi)存錯(cuò)誤也能觀察出來(lái)，但是有些邏輯錯(cuò)誤很難發(fā)現(xiàn)，甚至到了上線運(yùn)行時(shí)才會(huì)被發(fā)現(xiàn)，而且有些厲害的bug復(fù)現(xiàn)會(huì)很難，不總出現(xiàn)，但是很致命，而且cuda基本都是異步執(zhí)行的，當(dāng)錯(cuò)誤出現(xiàn)的時(shí)候，不一定是哪一條指令觸發(fā)的，這一點(diǎn)非常頭疼；這時(shí)候我們就需要對(duì)錯(cuò)誤進(jìn)行防御性處理了，例如我們代碼庫(kù)頭文件里面的這個(gè)宏：

1

2

3

4

5

6

7

8

9

10

#define check(call)\

{\

const cudaerror_t error=call;\

if(error!=cudasuccess)\

{\

printf('error: %s:%d,',__file__,__line__);\

printf('code:%d,reason:%s\n',error,cudageterrorstring(error));\

exit(1);\

}\

}

就是獲得每個(gè)函數(shù)執(zhí)行后的返回結(jié)果，然后對(duì)不成功的信息加以處理，cuda c 的api每個(gè)調(diào)用都會(huì)返回一個(gè)錯(cuò)誤代碼，這個(gè)代碼我們就可以好好利用了，當(dāng)然在release版本中可以去除這部分，但是開(kāi)發(fā)的時(shí)候一定要有的。

編譯執(zhí)行

接下來(lái)我們對(duì)上面那段向量加法的代碼編譯執(zhí)行，觀察結(jié)果，這里需要一張圖，我明天早上連接服務(wù)器后運(yùn)行給出，今晚網(wǎng)太差，所以請(qǐng)見(jiàn)諒，
編譯指令：

1	nvcc xxxx.cu -o xxxx

總結(jié)

這篇可以說(shuō)寫(xiě)的比前幾篇流暢很多，因?yàn)檫@篇知識(shí)很連貫，不想漆面的概覽，瑣碎的知識(shí)，我們明天繼續(xù)。。。

machineln 交流群請(qǐng)掃碼加machinelp為好友：

>>>\n',block.x,grid.x);

check(cudamemcpy(res_from_gpu_h,res_d,nbyte,cudamemcpydevicetohost));

sumarrays(a_h,b_h,res_h,nelem);

checkresult(res_h,res_from_gpu_h,nelem);

cudafree(a_d);

cudafree(b_d);

cudafree(res_d);

free(a_h);

free(b_h);

free(res_h);

free(res_from_gpu_h);

return 0;

}

在開(kāi)發(fā)階段，每一步都進(jìn)行驗(yàn)證是絕對(duì)高效的，比把所有功能都寫(xiě)好，然后進(jìn)行測(cè)試這種過(guò)程效率高很多，同樣寫(xiě)cuda也是這樣的每個(gè)代碼小塊都進(jìn)行測(cè)試，看起來(lái)慢，實(shí)際會(huì)提高很多效率。
cuda小技巧，當(dāng)我們進(jìn)行調(diào)試的時(shí)候可以把核函數(shù)配置成單線程的：

1	kernel_name<<<1,1>>>(argument list)

錯(cuò)誤處理

所有編程都需要對(duì)錯(cuò)誤進(jìn)行處理，早起的編碼錯(cuò)誤，編譯器會(huì)幫我們搞定，內(nèi)存錯(cuò)誤也能觀察出來(lái)，但是有些邏輯錯(cuò)誤很難發(fā)現(xiàn)，甚至到了上線運(yùn)行時(shí)才會(huì)被發(fā)現(xiàn)，而且有些厲害的bug復(fù)現(xiàn)會(huì)很難，不總出現(xiàn)，但是很致命，而且cuda基本都是異步執(zhí)行的，當(dāng)錯(cuò)誤出現(xiàn)的時(shí)候，不一定是哪一條指令觸發(fā)的，這一點(diǎn)非常頭疼；這時(shí)候我們就需要對(duì)錯(cuò)誤進(jìn)行防御性處理了，例如我們代碼庫(kù)頭文件里面的這個(gè)宏：

1

2

3

4

5

6

7

8

9

10

#define check(call)\

{\

const cudaerror_t error=call;\

if(error!=cudasuccess)\

{\

printf('error: %s:%d,',__file__,__line__);\

printf('code:%d,reason:%s\n',error,cudageterrorstring(error));\

exit(1);\

}\

}

就是獲得每個(gè)函數(shù)執(zhí)行后的返回結(jié)果，然后對(duì)不成功的信息加以處理，cuda c 的api每個(gè)調(diào)用都會(huì)返回一個(gè)錯(cuò)誤代碼，這個(gè)代碼我們就可以好好利用了，當(dāng)然在release版本中可以去除這部分，但是開(kāi)發(fā)的時(shí)候一定要有的。

編譯執(zhí)行

接下來(lái)我們對(duì)上面那段向量加法的代碼編譯執(zhí)行，觀察結(jié)果，這里需要一張圖，我明天早上連接服務(wù)器后運(yùn)行給出，今晚網(wǎng)太差，所以請(qǐng)見(jiàn)諒，
編譯指令：

1	nvcc xxxx.cu -o xxxx

總結(jié)

這篇可以說(shuō)寫(xiě)的比前幾篇流暢很多，因?yàn)檫@篇知識(shí)很連貫，不想漆面的概覽，瑣碎的知識(shí)，我們明天繼續(xù)。。。

本站是提供個(gè)人知識(shí)管理的網(wǎng)絡(luò)存儲(chǔ)空間，所有內(nèi)容均由用戶發(fā)布，不代表本站觀點(diǎn)。請(qǐng)注意甄別內(nèi)容中的聯(lián)系方式、誘導(dǎo)購(gòu)買等信息，謹(jǐn)防詐騙。如發(fā)現(xiàn)有害或侵權(quán)內(nèi)容，請(qǐng)點(diǎn)擊一鍵舉報(bào)。

小男孩‘自慰网亚洲一区二区,亚洲一级在线播放毛片,亚洲中文字幕av每天更新,黄aⅴ永久免费无码,91成人午夜在线精品,色网站免费在线观看,亚洲欧洲wwwww在线观看

CUDA之CUDA編程模型概述（二）

開(kāi)篇廢話

核函數(shù)概述

啟動(dòng)核函數(shù)

編寫(xiě)核函數(shù)

驗(yàn)證核函數(shù)

錯(cuò)誤處理

編譯執(zhí)行

總結(jié)

錯(cuò)誤處理

編譯執(zhí)行

總結(jié)