论坛版：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 的说明文件，然后再把函式名称改掉)。
第一次发比较长篇的文章，如果版面有乱掉请见谅