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

Question

我正在使用 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,b,c ) 将发出：

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

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

Answer 1

GPU真正执行的指令是SASS，不是PTX。 PTX 是一种中间格式，将 PTX 转换为 SASS 的工具是一个 optimizing 编译器。

当我按照您的建议执行操作并研究 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 分析来回答这样的问题。

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

CUDA half float operations without explicit intrinsics

cuda

intrinsics

nvcc

fma

half-precision-float