【CUDA 基礎】2.0 CUDA編程模型概述（二）

04-19

Abstract: 本文繼續上文介紹CUDA編程模型關於核函數以及錯誤處理部分

Keywords: CUDA核函數，CUDA錯誤處理

開篇廢話

今天的廢話就是人的性格一旦形成，那麼就會成為最大的指向標，或者說一個人的性格思維方式能夠決定這個人的全部生命軌跡，比如有人真的愛學習（比如我，嘻嘻嘻）有人真的不愛學習，沒有優劣，只是兩種生活態度，因為學習這個事你學一輩子也學不完人類智慧的九牛一毛，而不學習可以有更多的時間進行社會實踐，融入社會，榮華富貴，享受生命。這是兩種性格，沒有好壞，畢竟每個人評價生活的標尺不一。努力追求自己想要的，不要嘲笑別人所追求的，能一起走的就一起走，不能一起走的就各自安好，這就是目前我理解的聲明的真諦。

繼續CUDA編程模型的後半部分，關於核函數以及錯誤處理。

- 核函數

- 啟動核函數

- 編寫核函數

- 驗證核函數

- 錯誤處理

核函數概述

核函數就是在CUDA模型上諸多線程中運行的那段串列代碼，這段代碼在設備上運行，用NVCC編譯，產生的機器碼是GPU的機器碼，所以我們寫CUDA程序就是寫核函數，第一步我們要確保核函數能正確的運行產生正切的結果，第二優化CUDA程序的部分，無論是優化演算法，還是調整內存結構，線程結構都是要調整核函數內的代碼，來完成這些優化的。

我們一直把我們的CPU當做一個控制者，運行核函數，要從CPU發起，那麼我們開始學習如何啟動一個核函數

啟動核函數

啟動核函數，通過的以下的ANSI C 擴展出的CUDA C指令：

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

其標準C的原型就是C語言函數調用

function_name(argument list);

這個三個尖括弧<<>>內是對設備代碼執行的線程結構的配置（或者簡稱為對內核進行配置），也就是我們上一篇中提到的線程結構中的網格，塊。回憶一下上文，我們通過CUDA C內置的數據類型dim3類型的變數來配置grid和block（上文提到過：在設備端訪問grid和block屬性的數據類型是uint3不能修改的常類型結構，這裡反覆強調一下）。

通過指定grid和block的維度，我們可以配置：

- 內核中線程的數目

- 內核中使用的線程布局

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

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

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

我們的核函數是同時複製到多個線程執行的，上文我們說過一個對應問題，多個計算執行在一個數據，肯定是浪費時間，所以為了讓多線程按照我們的意願對應到不同的數據，就要給線程一個唯一的標識，由於設備內存是線性的（基本市面上的內存硬體都是線性形式存儲數據的）我們觀察上圖，可以用threadIdx.x 和blockIdx.x 來組合獲得對應的線程的唯一標識（後面我們會看到，threadIdx和blockIdx能組合出很多不一樣的效果）

接下來我們就是修改代碼的時間了，改變核函數的配置，產生運行出結果一樣，但效率不同的代碼：

1. 一個塊：

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

32個塊

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

上述代碼如果沒有特殊結構在核函數中，執行結果應該一致，但是有些效率會一直比較低。

上面這些是啟動部分，當主機啟動了核函數，控制權馬上回到主機，而不是主機等待設備完成核函數的運行，這一點我們上一篇文章也有提到過（就是等待hello world輸出的那段代碼後面要加一句）

想要主機等待設備端執行可以用下面這個指令：

cudaError_t cudaDeviceSynchronize(void);

這是一個顯示的方法，對應的也有隱式方法，隱式方法就是不明確說明主機要等待設備端，而是設備端不執行完，主機沒辦法進行，比如內存拷貝函數：

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

這個函數上文已經介紹過了，當核函數啟動後的下一條指令就是從設備複製數據回主機端，那麼主機端必須要等待設備端計算完成。

所有CUDA核函數的啟動都是非同步的，這點與C語言是完全不同的

編寫核函數

完整內容參考https://face2ai.com/CUDA-F-2-1-CUDA編程模型概述2/