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

總結

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

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。