TW Community

AMD HIP

HIP: Heterogeneous-Computing Interface for Portability (HIP) is a C++ dialect designed to ease conversion of CUDA applications to portable C++ code. It provides a C-style API and a C++ kernel language. The C++ interface can use templates and classes across the host/kernel boundary.

AMD 前幾年開始推廣 HIP,其中一個賣點就是 HIP 的程式碼是可以跑在 AMD GPU (但非全部 GPU,例如 rx5700xt 因為架構不一樣所以還不行) 跟 NVIDIA GPU。另外 AMD 也有提供 HIPIFY 幫助你把 cuda 程式碼轉成 HIP 程式碼。等等會介紹他大致上有哪些差異。

如果在程式編寫上有用的 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 ,更詳細的內容可以參照 HIP GUIDE

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

簡單範例

template <int value>
__global__ void set_value(int num, int *__restrict__ array) {
	// set_value;
}
int main() {
	// ...
    set_value<4><<<grid, block>>>(num, array);
    // ...
}
cuda simple code
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);
    // ...
}
hip simple code

可以看到其實在 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 的說明文件,然後再把函式名稱改掉)。

把 PTT 的問題回覆也放在這裡

  1. 轉 HIP 的原因
    主要是讓程式也能在 AMD GPU 上跑,但由於 CUDA 有一些功能還是沒有在 HIP 裡頭,所以我們還是有留著 CUDA 專門給 NVIDIA GPU 使用,但在 HIP 的部分,就可以跑在 NVIDIA 或 AMD GPU 上了
  2. 之後維護該怎麼確保兩邊一致
    這個可以分兩個面向來講
    – 測試:我們有用 google test 做單元測試,我們先寫出 reference (只用一執行緒) 然後再寫 CUDA 跟 HIP,去檢查 CUDA/HIP 的結果跟 reference 的結果在誤差範圍內。
    – 減少重複性:像文章所提到的一樣,大多數的情況下 kernel 在兩邊是一致的,所以我們把共用的部分寫在一個檔案,在 CUDA/HIP 個別放進去 (#include "file"),有點像是大型 MACRO,但我們直接弄成一個檔案,方便查寫跟排版;另外我們也有考慮到 AMD NVIDIA 始終是不同的 GPU,像是 warp/wavefront size 就不一樣 (NVIDIA - 32, AMD - 64),所以如果我們要盡量使用 GPU 的效能,就要有一點彈性來設定參數,#include "file" 這個方式也可以給我們這樣的彈性。下面用範例可能會比較清楚
  • kernel.inc
__global__ void func(...) {
    // some functions needs some parameters to get good performance
    reduction<warp_size>(...); 
}
  • nvidia.cu
constexpr int warp_size = 32;

#include "kernel.inc"

void call_func() {
    func<<<...>>>(...);
}
  • amd.hip.cpp
#ifdef __HIP_PLATFORM_NVIDIA__
constexpr int warp_size = 32;
#else // __HIP_PLATFORM_AMD__
constexpr int warp_size = 64;
#endif

#include "kernel.inc"

void call_func() {
    // launch func
}

例如這樣子, nvidia.cu 和 amd.hip.cpp 都直接使用 kernel.inc,所以可以改參數來配合不同的 GPU,而 kernel 本身也只需要改一個地方,但由於參數不一樣,所以我們還有 gtest 來做另一層把關;這是目前我們根據維護性跟高效能所需的彈性,來使用的方式,可能未來還會再變動也說不定。