5
运行时
DTK中half2类型成员变量x,y类型问题
2023-09-20 13:17:13
109

在DTK中,half2类型的两个成员变量x,y被定义为usigned short类型,而在CUDA中则是half类型。这会导致如下的代码编译出来的汇编代码产生FP16到uint16的转换指令(比如汇编的第4~7行将FP16转成了uint16),从而导致整个结果发生错误。正确的定义应该是hip_fp16.h中half2的两个x,y成员变量的类型为half,而不是usigned short。

global void add(float *d_A, float *d_B, float *d_C) {
half2 a, b, c;
a.x = __float2half(d_A[0]);
a.y = __float2half(d_A[1]);
b.x = __float2half(d_B[0]);
b.y = __float2half(d_B[1]);
c = a * b;
d_C[0] = __half2float(c.x);
d_C[1] = __half2float(c.y);
}
#编译出来的汇编代码为
0: v_cvt_f16_f32_e32 v1, s0
1: v_cvt_f16_f32_e32 v3, s8
2: v_cvt_f16_f32_e32 v2, s1
3: v_cvt_f16_f32_e32 v4, s9
4: v_cvt_u16_f16_e32 v1, v1
5: v_cvt_u16_f16_e32 v3, v3
6: v_cvt_u16_f16_e32 v2, v2
7: v_cvt_u16_f16_e32 v4, v4
8: v_and_b32_e32 v1, v0, v1
9: v_and_b32_e32 v0, v0, v3
10: v_lshl_or_b32 v1, v2, 16, v1
11: v_lshl_or_b32 v0, v4, 16, v0
12: v_pk_mul_f16 v1, v1, v0
13: v_cvt_f32_f16_e32 v0, v1
14: v_cvt_f32_f16_sdwa v1, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
15: v_mov_b32_e32 v2, 0
16: global_store_dwordx2 v2, v[0:1], s[2:3]

全部评论(3)
只看楼主
网
网友bpkSEX
3楼
您好,请问,您通过反编译得到汇编代码的命令是什么?是extractkernel吗?
2024/1/23
J
jiaojiankun
2楼

您好,经过测试发现,由于底层实现上的差异,cuda部分用法目前暂时无法直接在hip上使用,该段示例建议做如下转换

global void add(float *d_A, float *d_B, float *d_C)

{ half2 a, b, c;

a.x = __half_as_ushort(__float2half(d_A[0]));

a.y = __half_as_ushort(__float2half(d_A[1]));

b.x = __half_as_ushort(__float2half(d_B[0]));

b.y = __half_as_ushort(__float2half(d_B[1]));

c = a * b;

d_C[0] = __half2float(__ushort_as_half(c.x));

d_C[1] = __half2float(__ushort_as_half(c.y));

}

2023/10/13
L
lizhigong
1楼
问题已经记录,后续版本更正,当前可以通过自己修改头文件amd_hip_fp16.h中struct __half2中的x,y类型为_Float16来修复此问题。感谢您提供的问题信息。
2023/9/28
没有更多内容了
--
楼主
Level.1
0 帖子
楼主热帖
暂无数据