没有显式内在函数的 CUDA 半浮点运算

如何解决没有显式内在函数的 CUDA 半浮点运算

我使用的是 CUDA 11.2，我使用 __half 类型对 16 位浮点值进行运算。

我很惊讶 nvcc 编译器在我这样做时不会正确调用融合乘法加法指令：

__half a,b,c;
...
__half x = a * b + c;

它发出单独的 mul 和 add 指令，而不是发出融合乘加指令。

mul.f16 %rs164,%rs1,%rs306;
add.f16 %rs167,%rs164,%rs65;

请注意，尽管使用了 --fmad=true 编译器选项。

而显式 __hfma( a,c ) 将发出：

fma.rn.f16 %rs164,%rs300,%rs65;

是使用 16 位浮点乘加来使用显式内在函数的唯一方法吗？

解决方法

GPU实际执行的指令是SASS，而不是PTX。 PTX 是一种中间格式，将 PTX 转换为 SASS 的工具是一个优化编译器。

当我按照您的建议执行操作并研究 SASS 时，我看到正在生成一条融合乘加指令：

$ cat t111.cu
#include <cuda_fp16.h>
__global__ void k(__half *x,__half a,__half b,__half c){
        *x = a*b+c;
}
$ nvcc -arch=sm_75 -c t111.cu
$ cuobjdump -ptx t111.o

Fatbin elf code:
================
arch = sm_75
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit

Fatbin ptx code:
================
arch = sm_75
code version = [7,1]
producer = <unknown>
host = linux
compile_size = 64bit
compressed

.version 7.1
.target sm_75
.address_size 64



.visible .entry _Z1kP6__halfS_S_S_(
.param .u64 _Z1kP6__halfS_S_S__param_0,.param .align 2 .b8 _Z1kP6__halfS_S_S__param_1[2],.param .align 2 .b8 _Z1kP6__halfS_S_S__param_2[2],.param .align 2 .b8 _Z1kP6__halfS_S_S__param_3[2]
)
{
.reg .b16 %rs<7>;
.reg .b64 %rd<3>;


ld.param.u64 %rd1,[_Z1kP6__halfS_S_S__param_0];
ld.param.u16 %rs2,[_Z1kP6__halfS_S_S__param_1];
ld.param.u16 %rs3,[_Z1kP6__halfS_S_S__param_2];
ld.param.u16 %rs6,[_Z1kP6__halfS_S_S__param_3];
cvta.to.global.u64 %rd2,%rd1;

        {mul.f16 %rs1,%rs2,%rs3;
}


        {add.f16 %rs4,%rs1,%rs6;
}

        st.global.u16 [%rd2],%rs4;
ret;
}


$ cuobjdump -sass t111.o

Fatbin elf code:
================
arch = sm_75
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit

        code for sm_75
                Function : _Z1kP6__halfS_S_S_
        .headerflags    @"EF_CUDA_SM75 EF_CUDA_PTX_SM(EF_CUDA_SM75)"
        /*0000*/                   MOV R1,c[0x0][0x28] ;                                /* 0x00000a0000017a02 */
                                                                                         /* 0x000fd00000000f00 */
        /*0010*/                   LDC.U16 R0,c[0x0][0x168] ;                           /* 0x00005a00ff007b82 */
                                                                                         /* 0x000e220000000400 */
        /*0020*/                   ULDC.64 UR4,c[0x0][0x160] ;                          /* 0x0000580000047ab9 */
                                                                                         /* 0x000fce0000000a00 */
        /*0030*/                   LDC.U16 R3,c[0x0][0x16a] ;                           /* 0x00005a80ff037b82 */
                                                                                         /* 0x000e240000000400 */
        /*0040*/                   HFMA2 R0,R0.H0_H0,R3.H0_H0,c[0x0] [0x16c].H0_H0 ;  /* 0x20005b0000007631 */
                                                                                         /* 0x001fd00000040803 */
        /*0050*/                   STG.E.U16.SYS [UR4],R0 ;                             /* 0x00000000ff007986 */
                                                                                         /* 0x000fe2000c10e504 */
        /*0060*/                   EXIT ;                                                /* 0x000000000000794d */
                                                                                         /* 0x000fea0003800000 */
        /*0070*/                   BRA 0x70;                                             /* 0xfffffff000007947 */
                                                                                         /* 0x000fc0000383ffff */
                ..........



Fatbin ptx code:
================
arch = sm_75
code version = [7,1]
producer = <unknown>
host = linux
compile_size = 64bit
compressed
$

(CUDA 11.1)

我不建议使用 PTX 分析来回答此类问题。