編寫CUDA內(nèi)核
介紹
與用于CPU編程的傳統(tǒng)順序模型不同，CUDA具有執(zhí)行模型。在CUDA中，編寫的代碼將同時(shí)由多個(gè)線程（通常成百上千個(gè)）執(zhí)行。解決方案將通過定義網(wǎng)格，塊和線程層次結(jié)構(gòu)進(jìn)行建模。
Numba的CUDA支持提供了用于聲明和管理此線程層次結(jié)構(gòu)的工具。這些功能與NVidia的CUDA C語言開放的功能非常相似。
Numba還開放了三種GPU內(nèi)存：全局設(shè)備內(nèi)存（連接到GPU本身的大型，相對(duì)較慢的片外內(nèi)存），片上 共享內(nèi)存和本地內(nèi)存。對(duì)于除最簡單算法以外的所有算法，務(wù)必仔細(xì)考慮如何使用和訪問內(nèi)存，以最大程度地減少帶寬需求和爭用，這一點(diǎn)很重要。
內(nèi)核聲明
一個(gè)核心功能是指從CPU代碼（）稱為GPU功能。它具有兩個(gè)基本特征：
? 內(nèi)核無法顯式返回值；所有結(jié)果數(shù)據(jù)都必須寫入傳遞給函數(shù)的數(shù)組中（如果計(jì)算標(biāo)量，則可能傳遞一個(gè)單元素?cái)?shù)組）；
? 內(nèi)核在調(diào)用時(shí)顯式聲明其線程層次結(jié)構(gòu)：即線程塊數(shù)和每個(gè)塊的線程數(shù)（請(qǐng)注意，雖然內(nèi)核僅編譯一次，但可以使用不同的塊大小或網(wǎng)格大小多次調(diào)用）。
用Numba編寫CUDA內(nèi)核看起來非常像為CPU編寫JIT函數(shù)：
@cuda.jit
def increment_by_one(an_array):
“”"
Increment all array elements by one.
“”"
# code elided here; read further for different implementations
（）注意：較新的CUDA支持設(shè)備端內(nèi)核啟動(dòng)；此功能稱為動(dòng)態(tài)并行性，但Numba當(dāng)前不支持它）
內(nèi)核調(diào)用
通常以以下方式啟動(dòng)內(nèi)核：
threadsperblock = 32
blockspergrid = (an_array.size + (threadsperblock - 1)) // threadsperblock
increment_by_oneblockspergrid, threadsperblock
注意到兩個(gè)步驟：
? 通過指定多個(gè)塊（或“每個(gè)網(wǎng)格的塊”）和每個(gè)塊的線程數(shù)來實(shí)例化內(nèi)核。兩者的乘積將給出啟動(dòng)的線程總數(shù)。內(nèi)核實(shí)例化是通過采用已編譯的內(nèi)核函數(shù)（在此處increment_by_one）并用整數(shù)元組對(duì)其進(jìn)行索引來完成的。
? 通過將輸入數(shù)組（如果需要，以及任何單獨(dú)的輸出數(shù)組）傳遞給內(nèi)核來運(yùn)行內(nèi)核。內(nèi)核異步運(yùn)行：啟動(dòng)將其在設(shè)備上的執(zhí)行排隊(duì)，然后立即返回。可以 cuda.synchronize()用來等待所有先前的內(nèi)核啟動(dòng)完成執(zhí)行。
注意
傳遞駐留在主機(jī)內(nèi)存中的數(shù)組，將隱式地導(dǎo)致將副本復(fù)制回主機(jī)，這將是同步的。在這種情況下，直到將數(shù)據(jù)復(fù)制回內(nèi)核啟動(dòng)才會(huì)返回，因此似乎是同步執(zhí)行的。
選擇塊大小
在聲明內(nèi)核所需的線程數(shù)時(shí)，具有兩級(jí)層次結(jié)構(gòu)似乎很奇怪。塊大小（即每個(gè)塊的線程數(shù)）通常很關(guān)鍵：
? 在軟件方面，塊大小確定多少線程共享內(nèi)存的給定區(qū)域。
? 在硬件方面，塊的大小必須足夠大以完全占用執(zhí)行單元。建議可在 CUDA C編程指南中找到。
多維塊和網(wǎng)格
為了幫助處理多維數(shù)組，CUDA允許指定多維塊和網(wǎng)格。在上面的示例中，可以使blockspergridandthreadsperblock元組為一個(gè)，兩個(gè)或三個(gè)整數(shù)。與等效大小的一維聲明相比，這不會(huì)改變所生成代碼的效率或行為，但可以幫助以更自然的方式編寫算法。
Thread線程定位
運(yùn)行內(nèi)核時(shí)，內(nèi)核函數(shù)的代碼由每個(gè)線程執(zhí)行一次。因此，它必須知道它在哪個(gè)線程中，以便知道它負(fù)責(zé)哪個(gè)數(shù)組元素（復(fù)雜算法可以定義更復(fù)雜的職責(zé)，但是基本原理是相同的）。
一種方法是讓線程確定其在網(wǎng)格和塊中的位置，然后手動(dòng)計(jì)算相應(yīng)的數(shù)組位置：
@cuda.jit
def increment_by_one(an_array):
# Thread id in a 1D block
tx = cuda.threadIdx.x
# Block id in a 1D grid
ty = cuda.blockIdx.x
# Block width, i.e. number of threads per block
bw = cuda.blockDim.x
# Compute flattened index inside the array
pos = tx + ty * bw
if pos < an_array.size: # Check array boundaries
an_array[pos] += 1
注意
除非確定塊大小和網(wǎng)格大小是陣列大小的除數(shù)，否則必須如上所述檢查邊界。
threadIdx，blockIdx，blockDim和gridDim 是由CUDA后端為知道Thread線程層次結(jié)構(gòu)的幾何形狀和當(dāng)前線程的該幾何形狀內(nèi)的位置，唯一目的提供特殊對(duì)象。
這些對(duì)象可以是1D，2D或3D，具體取決于調(diào)用內(nèi)核的方式 。在每個(gè)維度訪問該值，可使用x，y并z分別這些對(duì)象的屬性。
numba.cuda.threadIdx
當(dāng)前線程塊中的線程索引。對(duì)于1D塊，索引（由x屬性賦予）是一個(gè)整數(shù)，范圍從0（包括）到numba.cuda.blockDim排除（exclusive）。當(dāng)使用多個(gè)維度時(shí)，每個(gè)維度都存在類似的規(guī)則。
numba.cuda.blockDim
實(shí)例化內(nèi)核時(shí)聲明的線程塊的形狀。對(duì)于給定內(nèi)核中的所有線程，即使屬于不同的塊（即，每個(gè)塊“已滿”），該值也相同。
numba.cuda.blockIdx
線程網(wǎng)格中的塊索引啟動(dòng)了內(nèi)核。對(duì)于一維網(wǎng)格，索引（由x屬性賦予）是一個(gè)整數(shù)，范圍從0（含）到numba.cuda.gridDim不包含（exclusive）。當(dāng)使用多個(gè)維度時(shí)，每個(gè)維度都存在類似的規(guī)則。
numba.cuda.gridDim
實(shí)例化內(nèi)核時(shí)，聲明的塊網(wǎng)格形狀，即此內(nèi)核調(diào)用啟動(dòng)的塊總數(shù)。
絕對(duì)位置
簡單的算法將傾向于總是以與上例相同的方式使用線程索引。Numba提供了其它工具來自動(dòng)執(zhí)行此類計(jì)算：
numba.cuda.grid（ndim ）
返回當(dāng)前線程在整個(gè)塊網(wǎng)格中的絕對(duì)位置。 ndim應(yīng)該與實(shí)例化內(nèi)核時(shí)聲明的維數(shù)相對(duì)應(yīng)。如果ndim為1，則返回一個(gè)整數(shù)。如果ndim為2或3，則返回給定整數(shù)的元組。
numba.cuda.gridsize（ndim ）
返回整個(gè)塊網(wǎng)格中Thread線程的絕對(duì)尺寸（或形狀）。 ndim與grid()上述含義相同。
使用這些功能，遞增示例可以變成：
@cuda.jit
def increment_by_one(an_array):
pos = cuda.grid(1)
if pos < an_array.size:
an_array[pos] += 1
二維數(shù)組和線程網(wǎng)格的相同示例為：
@cuda.jit
def increment_a_2D_array(an_array):
x, y = cuda.grid(2)
if x < an_array.shape[0] and y < an_array.shape[1]:
an_array[x, y] += 1
注意，實(shí)例化內(nèi)核時(shí)，網(wǎng)格計(jì)算仍必須手動(dòng)完成，例如：
threadsperblock = (16, 16)
blockspergrid_x = math.ceil(an_array.shape[0] / threadsperblock[0])
blockspergrid_y = math.ceil(an_array.shape[1] / threadsperblock[1])
blockspergrid = (blockspergrid_x, blockspergrid_y)
increment_a_2D_arrayblockspergrid, threadsperblock
進(jìn)一步閱讀
請(qǐng)參閱《CUDA C編程指南》，以詳細(xì)了解CUDA編程。

總結(jié)

以上是生活随笔為你收集整理的编写CUDA内核的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯(cuò)，歡迎將生活随笔推薦給好友。