5.2。編寫 HSA 內核 · Numba 0.44 中文文檔

# 5.2。編寫 HSA 內核 > 原文： [http://numba.pydata.org/numba-doc/latest/roc/kernels.html](http://numba.pydata.org/numba-doc/latest/roc/kernels.html) ## 5.2.1。簡介 HSA 提供類似于 OpenCL 的執行模型。指令由一組硬件線程并行執行。在某種程度上，這類似于 _ 單指令多數據 _（SIMD）模型，但方便的是，細粒度調度對程序員是隱藏的，而不是用 SIMD 向量作為數據結構進行編程。在 HSA 中，您編寫的代碼將由多個線程同時執行（通常為數百或數千）。您的解決方案將通過定義 _grid_ ， _workgroup_ 和 _workitem_ 的線程層次結構來建模。 Numba 的 HSA 支持公開了用于聲明和管理這種線程層次結構的工具。 ## 5.2.2。 CUDA 程序員簡介 HSA 執行模型類似于 CUDA。 HSA 在 ROC GPU 上采用的存儲器模型也類似于 CUDA。 ROC GPU 專用于 GPU 內存，因此根據 CUDA 需要`to_device()`和`copy_to_host()`等。以下是 CUDA 術語與 HSA 的快速映射： * `workitem`相當于 CUDA 線程。 * `workgroup`相當于 CUDA 線程塊。 * `grid`相當于 CUDA 網格。 * `wavefront`相當于 CUDA `warp`。 ## 5.2.3。內核聲明 _ 內核函數 _ 是一個 GPU 函數，可以從 CPU 代碼中調用。它賦予它兩個基本特征： * 內核無法顯式返回值;所有結果數據必須寫入傳遞給函數的數組（如果計算一個標量，你可能會傳遞一個單元素數組）; * 內核在調用時顯式聲明其線程層次結構：即工作組的數量和每個工作組的工作項數量（請注意，雖然內核編譯一次，但可以使用不同的工作組大小或網格大小多次調用它）。乍一看，用 Numba 編寫 HSA 內核看起來非常像為 CPU 編寫 [JIT 函數](../glossary.html#term-jit-function)： ```py @roc.jit def increment_by_one(an_array): """ Increment all array elements by one. """ # code elided here; read further for different implementations ``` ## 5.2.4。內核調用內核通常以以下方式啟動： ```py itempergroup = 32 groupperrange = (an_array.size + (itempergroup - 1)) // itempergroup increment_by_one[groupperrange, itempergroup](an_array) ``` 我們在這里注意兩個步驟： * 通過指定多個工作組（或“每個網格的工作組”）以及每個工作組的多個工作項來實例化內核。兩者的產品將給出推出的工作項目總數。內核實例化是通過編譯內核函數（此處為`increment_by_one`）并使用整數元組對其進行索引來完成的。 * 運行內核，通過傳遞輸入數組（以及任何必要的單獨輸出數組）。默認情況下，運行內核是同步的：當內核完成執行并且數據被同步回來時，函數返回。 ### 5.2.4.1。選擇工作組大小在聲明內核所需的工作項數時，擁有兩級層次結構似乎很奇怪。工作組大小（即每個工作組的工作項數）通常至關重要： * 在軟件方面，工作組大小決定共享[共享內存](memory.html#roc-shared-memory)的給定區域的線程數。 * ```py On the hardware side, the workgroup size must be large enough for full ``` 占領執行單位。 ### 5.2.4.2。多維工作組和網格為了幫助處理多維數組，HSA 允許您指定多維工作組和網格。在上面的示例中，您可以使`itempergroup`和`groupperrange`元組為一個，兩個或三個整數。與等效大小的 1D 聲明相比，這不會改變生成代碼的效率或行為，但可以幫助您以更自然的方式編寫算法。 ## 5.2.5。 WorkItem 定位運行內核時，每個線程執行一次內核函數的代碼。因此，它必須知道它所在的線程，以便知道它負責哪個數組元素（復雜的算法可能定義更復雜的責任，但基本原理是相同的）。一種方法是讓線程確定它在網格和工作組中的位置，并手動計算相應的數組位置： ```py @roc.jit def increment_by_one(an_array): # workitem id in a 1D workgroup tx = roc.get_local_id(0) # workgroup id in a 1D grid ty = roc.get_group_id(0) # workgroup size, i.e. number of workitem per workgroup bw = roc.get_local_size(0) # Compute flattened index inside the array pos = tx + ty * bw # The above is equivalent to pos = roc.get_global_id(0) if pos < an_array.size: # Check array boundaries an_array[pos] += 1 ``` 注意除非您確定工作組大小和網格大小是數組大小的除數，否則**必須**檢查邊界，如上所示。 [`get_local_id()`](#numba.roc.get_local_id "numba.roc.get_local_id") ， [`get_local_size()`](#numba.roc.get_local_size "numba.roc.get_local_size") ， [`get_group_id()`](#numba.roc.get_group_id "numba.roc.get_group_id") 和 [`get_global_id()`](#numba.roc.get_global_id "numba.roc.get_global_id") 是 HSA 后端為鞋底提供的特殊功能了解線程層次結構的幾何以及當前工作項在該幾何中的位置的目的。 ```py numba.roc.get_local_id(dim) ``` 獲取要查詢的維度的索引返回給定維度的當前工作組中的本地工作項 ID。對于 1D 工作組，索引是一個整數，范圍從 0 到包括 [`numba.roc.get_local_size()`](#numba.roc.get_local_size "numba.roc.get_local_size") 不包括。 ```py numba.roc.get_local_size(dim) ``` 獲取要查詢的維度的索引返回給定維度的工作組的大小。在實例化內核時聲明該值。對于給定內核中的所有工作項，此值都是相同的，即使它們屬于不同的工作組（即每個工作組都是“完整”）。 ```py numba.roc.get_group_id(dim) ``` 獲取要查詢的維度的索引返回啟動內核的工作組網格中的工作組 ID。 ```py numba.roc.get_global_id(dim) ``` 獲取要查詢的維度的索引返回給定維度的全局工作項 ID。與 <cite>numba.roc .get_local_id（）</cite>不同，此數字對于網格中的所有工作項都是唯一的。