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);

  1. 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/

推薦閱讀:

相關文章