3.2。編寫 CUDA 內核 · Numba 0.44 中文文檔

# 3.2。編寫 CUDA 內核 > 原文： [http://numba.pydata.org/numba-doc/latest/cuda/kernels.html](http://numba.pydata.org/numba-doc/latest/cuda/kernels.html) ## 3.2.1。簡介與用于編程 CPU 的傳統順序模型不同，CUDA 具有執行模型。在 CUDA 中，您編寫的代碼將由多個線程同時執行（通常為數百或數千）。您的解決方案將通過定義 _grid_ ， _blocks_ 和 _threads_ 的線程層次結構來建模。 Numba 的 CUDA 支持公開了用于聲明和管理這種線程層次結構的工具。這些設施與 NVidia 的 CUDA C 語言大致相似。 Numba 還暴露了三種 GPU 內存：全局[設備內存](memory.html#cuda-device-memory)（連接到 GPU 本身的大型，相對較慢的片外內存），片上[共享內存](memory.html#cuda-shared-memory)和 []本地記憶](memory.html#cuda-local-memory)。除了最簡單的算法外，您必須仔細考慮如何使用和訪問內存，以最大限度地減少帶寬需求和爭用。 ## 3.2.2。內核聲明 _ 內核函數 _ 是一個 GPU 函數，用于從 CPU 代碼（*）調用。它賦予它兩個基本特征： * 內核無法顯式返回值;所有結果數據必須寫入傳遞給函數的數組（如果計算一個標量，你可能會傳遞一個單元素數組）; * 內核在被調用時顯式聲明它們的線程層次結構：即線程塊的數量和每個塊的線程數（注意，當內核被編譯一次時，可以使用不同的塊大小或網格大小多次調用它）。乍一看，使用 Numba 編寫 CUDA 內核看起來非常像為 CPU 編寫 [JIT 函數](../glossary.html#term-jit-function)： ```py @cuda.jit def increment_by_one(an_array): """ Increment all array elements by one. """ # code elided here; read further for different implementations ``` （*）注意：較新的 CUDA 設備支持設備端內核啟動;此功能稱為 _ 動態并行 _，但 Numba 目前不支持它 ## 3.2.3。內核調用內核通常以以下方式啟動： ```py threadsperblock = 32 blockspergrid = (an_array.size + (threadsperblock - 1)) // threadsperblock increment_by_one[blockspergrid, threadsperblock](an_array) ``` 我們在這里注意兩個步驟： * 通過指定多個塊（或“每個網格的塊”）以及每個塊的多個線程來實例化內核。兩者的乘積將給出啟動的線程總數。內核實例化是通過編譯內核函數（此處為`increment_by_one`）并使用整數元組對其進行索引來完成的。 * 運行內核，通過傳遞輸入數組（以及任何必要的單獨輸出數組）。默認情況下，運行內核是同步的：當內核完成執行并且數據被同步回來時，函數返回。 ### 3.2.3.1。選擇塊大小在聲明內核所需的線程數時，擁有兩級層次結構似乎很奇怪。塊大小（即每個塊的線程數）通常至關重要： * 在軟件方面，塊大小決定共享[共享內存](memory.html#cuda-shared-memory)的給定區域的線程數。 * 在硬件方面，塊大小必須足夠大才能完全占用執行單元;建議可在 [CUDA C 編程指南](http://docs.nvidia.com/cuda/cuda-c-programming-guide)中找到。 ### 3.2.3.2。多維塊和網格為了幫助處理多維數組，CUDA 允許您指定多維塊和網格。在上面的示例中，您可以使`blockspergrid`和`threadsperblock`元組為一個，兩個或三個整數。與等效大小的 1D 聲明相比，這不會改變生成代碼的效率或行為，但可以幫助您以更自然的方式編寫算法。 ## 3.2.4。螺紋定位運行內核時，每個線程執行一次內核函數的代碼。因此，它必須知道它所在的線程，以便知道它負責哪個數組元素（復雜的算法可能定義更復雜的責任，但基本原理是相同的）。一種方法是讓線程定位它在網格和塊中的位置，并手動計算在數組中對應的位置： ```py @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`](../cuda-reference/kernel.html#numba.cuda.threadIdx "numba.cuda.threadIdx") ， [`blockIdx`](../cuda-reference/kernel.html#numba.cuda.blockIdx "numba.cuda.blockIdx") ， [`blockDim`](../cuda-reference/kernel.html#numba.cuda.blockDim "numba.cuda.blockDim") 和 [`gridDim`](../cuda-reference/kernel.html#numba.cuda.gridDim "numba.cuda.gridDim") 是 CUDA 后端為鞋底提供的特殊對象了解線程層次結構的幾何以及當前線程在該幾何中的位置的目的。這些對象可以是 1D，2D 或 3D，具體取決于內核[調用的方式](#cuda-kernel-invocation)。要訪問每個維度的值，請分別使用這些對象的`x`，`y`和`z`屬性。 ```py numba.cuda.threadIdx ``` 當前線程塊中的線程索引。對于 1D 塊，索引（由`x`屬性給出）是一個整數，范圍從 0 到包括 [`numba.cuda.blockDim`](../cuda-reference/kernel.html#numba.cuda.blockDim "numba.cuda.blockDim") 不包括。當使用多個維度時，每個維度都存在類似的規則。 ```py numba.cuda.blockDim ``` 線程塊的形狀，在實例化內核時聲明。對于給定內核中的所有線程，該值是相同的，即使它們屬于不同的塊（即每個塊都是“滿”）。 ```py numba.cuda.blockIdx ``` 線程網格中的塊索引啟動了一個內核。對于 1D 網格，索引（由`x`屬性給出）是一個整數，范圍從 0 到包括 [`numba.cuda.gridDim`](../cuda-reference/kernel.html#numba.cuda.gridDim "numba.cuda.gridDim") 不包括。當使用多個維度時，每個維度都存在類似的規則。 ```py numba.cuda.gridDim ``` 塊網格的形狀，即由內核調用啟動的塊的總數，在實例化內核時聲明。 ### 3.2.4.1。絕對位置簡單的算法傾向于始終以與上面示例中所示相同的方式使用線程索引。 Numba 提供額外的設施來自動進行這樣的計算： ```py numba.cuda.grid(ndim) ``` 返回整個塊網格中當前線程的絕對位置。 _ndim_ 應該對應于實例化內核時聲明的維數。如果 _ndim_ 為 1，則返回單個整數。如果 _ndim_ 為 2 或 3，則返回給定數量的整數的元組。 ```py numba.cuda.gridsize(ndim) ``` 返回整個塊網格的線程中的絕對大小（或形狀）。 _ndim_ 具有與上述 [`grid()`](../cuda-reference/kernel.html#numba.cuda.grid "numba.cuda.grid") 相同的含義。使用這些函數，增量示例可以變為： ```py @cuda.jit def increment_by_one(an_array): pos = cuda.grid(1) if pos < an_array.size: an_array[pos] += 1 ``` 2D 陣列和線程網格的相同示例是： ```py @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 ``` 請注意，實例化內核時的網格計算仍必須手動完成，例如： ```py from __future__ import division # for Python 2 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_array[blockspergrid, threadsperblock](an_array) ``` ### 3.2.4.2。進一步閱讀有關 CUDA 編程的詳細討論，請參見 [CUDA C 編程指南](http://docs.nvidia.com/cuda/cuda-c-programming-guide)。