# 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>不同,此數字對于網格中的所有工作項都是唯一的。
- 1. 用戶手冊
- 1.1。 Numba 的約 5 分鐘指南
- 1.2。概述
- 1.3。安裝
- 1.4。使用@jit 編譯 Python 代碼
- 1.5。使用@generated_jit 進行靈活的專業化
- 1.6。創建 Numpy 通用函數
- 1.7。用@jitclass 編譯 python 類
- 1.8。使用@cfunc 創建 C 回調
- 1.9。提前編譯代碼
- 1.10。使用@jit 自動并行化
- 1.11。使用@stencil裝飾器
- 1.12。從 JIT 代碼 中回調到 Python 解釋器
- 1.13。性能提示
- 1.14。線程層
- 1.15。故障排除和提示
- 1.16。常見問題
- 1.17。示例
- 1.18。會談和教程
- 2. 參考手冊
- 2.1。類型和簽名
- 2.2。即時編譯
- 2.3。提前編譯
- 2.4。公用事業
- 2.5。環境變量
- 2.6。支持的 Python 功能
- 2.7。支持的 NumPy 功能
- 2.8。與 Python 語義的偏差
- 2.9。浮點陷阱
- 2.10。 Python 2.7 壽命終止計劃
- 3. 用于 CUDA GPU 的 Numba
- 3.1。概述
- 3.2。編寫 CUDA 內核
- 3.3。內存管理
- 3.4。編寫設備功能
- 3.5。 CUDA Python 中支持的 Python 功能
- 3.6。支持的原子操作
- 3.7。隨機數生成
- 3.8。設備管理
- 3.10。示例
- 3.11。使用 CUDA 模擬器 調試 CUDA Python
- 3.12。 GPU 減少
- 3.13。 CUDA Ufuncs 和廣義 Ufuncs
- 3.14。共享 CUDA 內存
- 3.15。 CUDA 陣列接口
- 3.16。 CUDA 常見問題
- 4. CUDA Python 參考
- 4.1。 CUDA 主機 API
- 4.2。 CUDA 內核 API
- 4.3。內存管理
- 5. 用于 AMD ROC GPU 的 Numba
- 5.1。概述
- 5.2。編寫 HSA 內核
- 5.3。內存管理
- 5.4。編寫設備功能
- 5.5。支持的原子操作
- 5.6。代理商
- 5.7。 ROC Ufuncs 和廣義 Ufuncs
- 5.8。示例
- 6. 擴展 Numba
- 6.1。高級擴展 API
- 6.2。低級擴展 API
- 6.3。示例:間隔類型
- 7. 開發者手冊
- 7.1。貢獻給 Numba
- 7.2。 Numba 建筑
- 7.3。多態調度
- 7.4。關于發電機的注意事項
- 7.5。關于 Numba Runtime 的注意事項
- 7.6。使用 Numba Rewrite Pass 獲得樂趣和優化
- 7.7。實時變量分析
- 7.8。上市
- 7.9。模板注釋
- 7.10。關于自定義管道的注意事項
- 7.11。環境對象
- 7.12。哈希 的注意事項
- 7.13。 Numba 項目路線圖
- 8. Numba 增強建議
- 9. 術語表