【hip程序的几个坑】

hip作为对标cuda的编程语言在开源社区较为活跃。不过因为架构问题，他也有一些坑的地方。在这里做一些总结。

半精度的使用

hip 中我们通过half和half2来使用半精度，half2并不能看作half[2]。如果想要获取half2的两个half值，需要使用__low2half()和__high2half()来获取。（不能使用.x和.y获取）

同时hip中没有__hmax和__hmax2函数，需要自己定义。示例如下：

__device__ half __hmax(half a, half b) {
    return __hgt(a, b) ? a : b;
}

__device__ half2 __hmax2(half2 a, half2 b) {
    half tmp1;
    half tmp2;
    tmp2 = __hgt(__high2half(a), __high2half(b)) ? __high2half(a) : __high2half(b);
    tmp1 = __hgt(__low2half(a), __low2half(b)) ? __low2half(a) : __low2half(b);
    return __halves2half2(tmp1, tmp2);
}

大概解释一下上述代码。__hgt(half a, half b) 函数等于 a > b, 最好使用这种比较的内置函数。__halves2half2(half a, half b) 表示将两个 half 合并，并输出。

warp指令的使用

warp指令用于在线程束中交换数据。AMD的线程束指令相比cuda有很大的限制。

1. warp指令不能用于条件语句中。详见这里

2. warp指令只能用于单精度。这是底层指令问题，没法解决。

warp的大小

在AMD上warp的大小一直是64（cuda为32），可是在最新的RDNA3架构中，warp的大小变为32。望注意！