編寫HSA內(nèi)核
介紹
HSA提供類似于OpenCL的執(zhí)行模型。指令由一組硬件線程并行執(zhí)行。在某種程度上，這類似于 單指令多數(shù)據(jù)（SIMD）模型，但具有這樣的便利：細(xì)粒度調(diào)度對(duì)于程序員而言是隱藏的，而不是使用SIMD向量作為數(shù)據(jù)結(jié)構(gòu)進(jìn)行編程。在HSA中，編寫的代碼將同時(shí)由多個(gè)線程（通常成百上千個(gè)）執(zhí)行。解決方案將通過(guò)定義網(wǎng)格，工作組 和工作項(xiàng)的線程層次結(jié)構(gòu)進(jìn)行建模。
Numba的HSA支持提供了用于聲明和管理此線程層次結(jié)構(gòu)的工具。
CUDA程序簡(jiǎn)介
HSA執(zhí)行模型類似于CUDA。HSA在ROC GPU上采用的內(nèi)存模型也與CUDA相似。ROC的GPU具有專用于GPU存儲(chǔ)器，因此，to_device()與copy_to_host()等需要按照CUDA。
這是CUDA術(shù)語(yǔ)到HSA的快速映射：
? Aworkitem等效于CUDA線程。
? Aworkgroup等效于CUDA線程塊。
? Agrid等效于CUDA網(wǎng)格。
? Awavefront等效于CUDA warp。
內(nèi)核聲明
一個(gè)核心功能是指從CPU代碼稱為GPU功能。它具有兩個(gè)基本特征：
? 內(nèi)核無(wú)法顯式返回值；所有結(jié)果數(shù)據(jù)都必須寫入傳遞給函數(shù)的數(shù)組中（如果計(jì)算標(biāo)量，則可能傳遞一個(gè)單元素?cái)?shù)組）；
? 內(nèi)核在被調(diào)用時(shí)顯式聲明其線程層次結(jié)構(gòu)：即工作組的數(shù)量和每個(gè)工作組的工作項(xiàng)的數(shù)量（注意，雖然內(nèi)核僅編譯一次，但可以使用不同的工作組大小或網(wǎng)格大小多次調(diào)用）。
用Numba編寫HSA內(nèi)核非常類似于為CPU編寫JIT函數(shù)：
@roc.jit
def increment_by_one(an_array):
“”"
Increment all array elements by one.
“”"
# code elided here; read further for different implementations
內(nèi)核調(diào)用
通常以以下方式啟動(dòng)內(nèi)核：
itempergroup = 32
groupperrange = (an_array.size + (itempergroup - 1)) // itempergroup
increment_by_onegroupperrange, itempergroup
注意到兩個(gè)步驟：
? 通過(guò)指定多個(gè)工作組（或“每個(gè)網(wǎng)格的工作組”）和每個(gè)工作組的多個(gè)工作項(xiàng)來(lái)實(shí)例化內(nèi)核。兩者的乘積將給出啟動(dòng)的工作項(xiàng)總數(shù)。內(nèi)核實(shí)例化是通過(guò)采用已編譯的內(nèi)核函數(shù)（在此處increment_by_one）并用整數(shù)元組對(duì)其進(jìn)行索引來(lái)完成的。
? 通過(guò)將輸入數(shù)組（如果需要，以及任何單獨(dú)的輸出數(shù)組）傳遞給內(nèi)核來(lái)運(yùn)行內(nèi)核。默認(rèn)情況下，運(yùn)行內(nèi)核是同步的：當(dāng)內(nèi)核完成執(zhí)行并且數(shù)據(jù)被同步時(shí)，該函數(shù)返回。
選擇工作組大小
在聲明內(nèi)核所需的工作項(xiàng)數(shù)量時(shí)，具有兩級(jí)層次結(jié)構(gòu)似乎很奇怪。工作組的大小（即每個(gè)工作組的工作項(xiàng)數(shù)）通常很關(guān)鍵：
? 在軟件方面，工作組的大小確定了多少線程共享內(nèi)存的給定區(qū)域。
? 在硬件方面，工作組的大小必須足夠大
獨(dú)占執(zhí)行單位。
多維工作組和網(wǎng)格
為了幫助處理多維數(shù)組，HSA指定多維工作組和網(wǎng)格。在上面的示例中，可以使itempergroupandgroupperrange元組為一個(gè)，兩個(gè)或三個(gè)整數(shù)。與等效大小的一維聲明相比，這不會(huì)改變所生成代碼的效率或行為，但可以幫助以更自然的方式編寫算法。
工作項(xiàng)定位
運(yùn)行內(nèi)核時(shí)，內(nèi)核函數(shù)的代碼由每個(gè)線程執(zhí)行一次。因此，它必須知道它在哪個(gè)線程中，以便知道它負(fù)責(zé)哪個(gè)數(shù)組元素（復(fù)雜算法可以定義更復(fù)雜的任務(wù)，但是基本原理是相同的）。
一種方法是讓線程確定其在網(wǎng)格和工作組中的位置，然后手動(dòng)計(jì)算相應(yīng)的數(shù)組位置：
@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
注意
除非確定工作組大小和網(wǎng)格大小是陣列大小的除數(shù)，否則必須如上所述檢查邊界。
get_local_id()，get_local_size()，get_group_id()和 get_global_id()是由HSA后端為知道thread層次結(jié)構(gòu)的幾何形狀和該幾何形狀內(nèi)的當(dāng)前工作項(xiàng)的位置的唯一目的提供特殊功能。
numba.roc.get_local_id（dim）
取得要查詢的維度的索引
返回給定維度的當(dāng)前工作組中的本地工作項(xiàng)ID。對(duì)于一維工作組，索引是一個(gè)整數(shù)，范圍從0（含）到numba.roc.get_local_size()異（exclusive）。
numba.roc.get_local_size（dim）
取得要查詢的維度的索引
返回給定維度上工作組的大小。實(shí)例化內(nèi)核時(shí)聲明該值。對(duì)于給定內(nèi)核中的所有工作項(xiàng)，該值相同，即使屬于不同的工作組（即，每個(gè)工作組“已滿”）也是如此。
numba.roc.get_group_id（dim）
取得要查詢的維度的索引
在啟動(dòng)了內(nèi)核的工作組網(wǎng)格中返回工作組ID。
numba.roc.get_global_id（dim）
取得要查詢的維度的索引
返回給定維度的全局工作項(xiàng)ID。與numba.roc .get_local_id（）不同，此數(shù)字對(duì)于網(wǎng)格中的所有工作項(xiàng)都是唯一的。

總結(jié)

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

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