論壇版：https://forum.community.tw/t/topic/68
新嘗試的論壇，目前以電腦相關的主題為主，能自動將程式碼上色，也可使用 Markdown
並且可以作為其他部落格的留言系統，歡迎大家來發表/參與討論
如果有關於這篇的相關感想或問題，發在 PTT 或者該論壇上都可
HIP
AMD 前幾年開始推廣 HIP，其中一個賣點就是 HIP 的程式碼是可以跑在 A
MD GPU (但非全部 GPU，例如 rx5700xt 因為架構不一樣所以還不行) 跟 NVIDIA GPU。另
外 AMD 也有提供 HIPIFY 幫助你把 cuda 程式碼轉成 HIP 程式碼。等等會介紹他大致上
有哪些差異。
HIP 和 CUDA 的差別
- 函式庫 (library) : 例如 cusparse 會變成 hipsparse、cuda_runtime 變
成 hip_runtime 等，那相對應的呼叫函數也會相對應的改變，但大多都是 cuda -> hip
這樣的轉換
- warp/wavefront size 以及 launch bound 等一些可能會影響程式的實作。另外 HIP 並
沒有直接支援 cooperative group ，必須直接使用 shuffle ，當然也可以另外實作相對
應的 cooperative group 來方便使用。
- 呼叫 gpu kernel 的差異，CUDA 是使用 <<<>>> ， cuda_kerenl<<<grid, block,
dynamic shared memory, stream>>>(args...) ，之中 stream 跟 dynamic shared
memory 是可以省略的；HIP 則是使用 hipLaunchKernelGGL(hip_kernel, grid, block,
dynamic shared memory, stream, args...) 那這邊的 stream 跟 dynamic shared
memory 就不能省略了。另外確保他 kernel 名稱沒有解析錯誤可以用上
HIP_KERNEL_NAME
如果在程式編寫上有用的 warp 的觀念的話，最大的差別是在 NVIDIA GPU 是用 32 但
是 AMD GPU 是用 64 ( AMD 稱為 wavefront = CUDA warp)，但如果原先在 CUDA 是使
用 blocksize = (32, 32) 這種，在 AMD 上也是只能用 blocksize = (32, 32) 並非
(64, 64) ，因為 blocksize 一樣最多只能 1024。另外 AMD 的 shuffle 是利用
shared memory 去做，並不像 NVIDIA 是直接從 register 操作。但他們也有提供 AMD
GCN Assembly: Cross-Lane Operations，就不會從 shared memory 處理。
另外一個差異點是在 __launch_bounds__ ， CUDA 的 __launch_bounds__(max threads
per block, min blocks per multipreocessor) ， 但 HIP 是 __launch_bounds__(max
threads per block, min warps per eu) 。如果只使用第一個參數的話那就是一致的，
但第二個參數就要修改，HIP 也有提供公式轉換 min warps per eu = (min blocks per
multiprocessor * max threads per block)/32
範例
// CUDA
template <int value>
__global__ void set_value(int num, int *__restrict__ array) {
// set_value;
}
int main() {
// ...
set_value<4><<<grid, block>>>(num, array);
// ...
}
// HIP
template <int value>
__global__ void set_value(int num, int *__restrict__ array) {
// set_value;
}
int main() {
// ...
hipLaunchKernelGGL(HIP_KERNEL_NAME(set_value<4>), grid, block, 0, 0,
num, array);
// ...
}
基本上 kernel 本身不太會需要變動，通常更動呼叫 kernel 的方式
hipify 常見錯誤
- namespace: 假如 kernel 有放在 namespace 的話，hipify 會給出錯誤的
結果，例如變成 namespace::hipLaunchKernelGGL(HIP_KERNEL_NAME(kernel), ...) 但
應該是要改成 hipLaunchKernelGGL(HIP_KERNEL_NAME(namespace::kernel), ...) 才對
- 有逗號 , 在 kernel 名稱或是 syntax (grid, block 那些 <<<>>> 中的參數):
hipify 類似將前面東西放完看有幾個逗號，然後再補差項上去。所以當有逗號在裡面時
他可能就會補錯地方。 例如 kernel<<<calc(num, block), block>>>(...) 可能就會變
成 hipLaunchKernelGGL(kernel, calc(num, block), block, 0, ...) ，中間的
calc(num, block) 被當成兩項了，應該要是 hipLaunchKernelGGL(kernel, calc(num,
block), block, 0, 0, ...) 才對。
個人心得 (或趣事?)
我們在將 CUDA 轉成 HIP 時，整體上沒有遇到太多困難，大多東西都有對應，hipify 也
可以只用在單檔上，但在整體使用上會相較 CUDA 本身比較麻煩的在於沒有一致性、且有
一些變動是沒有考慮到以前版本。例如在使用 CMake 時， find_pacakage 要找 HIP,
hipblas, hipsparse 但在 hiprand 的時候還要額外找 rocrand ，而在
target_link_libraries 時是用 roc::hipblas, roc::hipsparse, hip::hiprand,
roc::rocrand 等總總不能從上一個的經驗推出下一個的時常發生
又或者 hip 4.1 版之前 HIP_PLATFORM(可拿來區別 amd/nvidia 的程式碼) 是用 hcc 或
者 nvcc 來做編譯期間的區別，但在 4.1 後，改成 amd 或 nvidia 來作區別，所以程式
本身要自己另外定義使之支援不同版。
但好消息是，AMD 有在跟 CMake 一起處理這些，希望之後再 CMake 上使用體驗會再更順
暢些。
另外一些是跟 CUDA 有關，有些函式在 CUDA 的說明文件是沒有標明 const 的，但在使
用上你是可以丟 const 的參數進去，且那個函數本身的確是不會更動那個參數；但 HIP
為了符合該說明文件，使用上是要去掉 const 才能使用函數，即使兩邊都不會動到該參
數的值。
HIP 也算是有推一陣子，在 CUDA 10.2 之前是真的可以兩邊都跑；但 CUDA 更新到 11
時，這個就不成立了，因為 CUDA 11 把很多函式如 csr_spmv 等砍掉，改用 generic
api 來操作，所以 10.2 的函數就不能直接在 cuda 11 用且 HIP 尚未有 generic api
的對應函式，因此目前是只有在 cuda 10.2 版以前才能直接用 HIP 跑在 NVIDIA GPU 上
。
對於 cross line operation 試起來的確比較快，但是寫法上並不是直接從 shuffle 直
接對應，且他是用 compile time 的參數，可能一個加總原先是五個 shuffle 但改用
cross line operation 可能就要七到八個指令組合，且不能直接使用 for (因為要
compile time 的參數)
總結
HIP 整體上使用還算可以，且 CUDA (10.2 前) 的函式大多都已有相對應的，所以不用擔
心移到 HIP 時要自己寫一些函式庫的東西，一些參數也相對好轉換很多。且 HIP 本身是
開源，所以有興趣的話也可以去看他們怎麼實作一些函式的，但缺點可能就是在編譯上或
者一些說明文件並沒有像 CUDA 那麼一致，需要一些時間去克服(我很多時候都是先去看
CUDA 的說明文件，然後再把函式名稱改掉)。
第一次發比較長篇的文章，如果版面有亂掉請見諒