我使用的是 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 分析来回答此类问题。