带负号的 fma 中每个周期的指令数是多少？

Question

如果我在cuda中使用fma(a, b, c)，这意味着公式ab+c是在单个三元运算中计算的。但是，如果我想计算 -ab+c，调用 fma(-a, b, c) 是否需要再进行一次乘法运算？

Answer 1

不幸的是，着色器汇编语言在那个级别没有记录。

不过我们可以试试看：

#!/bin/bash
cat <<EOF > fmatest.cu
__global__ void fma_plus(float *res, float a, float b, float c)
{
    *res = fma(a, b, c);
}

__global__ void fma_minus(float *res, float a, float b, float c)
{
    *res = fma(-a, b, c);
}
EOF
nvcc -arch sm_60 -c fmatest.cu
cuobjdump -sass fmatest.o

给予

code for sm_60
    Function : _Z9fma_minusPffff
.headerflags    @"EF_CUDA_SM60 EF_CUDA_PTX_SM(EF_CUDA_SM60)"
                                                                 /* 0x001fc400fe2007f6 */
    /*0008*/                   MOV R1, c[0x0][0x20];             /* 0x4c98078000870001 */
    /*0010*/                   MOV R0, c[0x0][0x148];            /* 0x4c98078005270000 */
    /*0018*/                   MOV R5, c[0x0][0x14c];            /* 0x4c98078005370005 */
                                                                 /* 0x001fc800fe8007f1 */
    /*0028*/                   MOV R2, c[0x0][0x140];            /* 0x4c98078005070002 */
    /*0030*/                   MOV R3, c[0x0][0x144];            /* 0x4c98078005170003 */
    /*0038*/                   FFMA R0, R0, -R5, c[0x0][0x150];  /* 0x5181028005470000 */
                                                                 /* 0x001ffc00ffe000f1 */
    /*0048*/                   STG.E [R2], R0;                   /* 0xeedc200000070200 */
    /*0050*/                   EXIT;                             /* 0xe30000000007000f */
    /*0058*/                   BRA 0x58;                         /* 0xe2400fffff87000f */
                                                                 /* 0x001f8000fc0007e0 */
    /*0068*/                   NOP;                              /* 0x50b0000000070f00 */
    /*0070*/                   NOP;                              /* 0x50b0000000070f00 */
    /*0078*/                   NOP;                              /* 0x50b0000000070f00 */
    ..................................


    Function : _Z8fma_plusPffff
.headerflags    @"EF_CUDA_SM60 EF_CUDA_PTX_SM(EF_CUDA_SM60)"
                                                                /* 0x001fc400fe2007f6 */
    /*0008*/                   MOV R1, c[0x0][0x20];            /* 0x4c98078000870001 */
    /*0010*/                   MOV R0, c[0x0][0x148];           /* 0x4c98078005270000 */
    /*0018*/                   MOV R5, c[0x0][0x14c];           /* 0x4c98078005370005 */
                                                                /* 0x001fc800fe8007f1 */
    /*0028*/                   MOV R2, c[0x0][0x140];           /* 0x4c98078005070002 */
    /*0030*/                   MOV R3, c[0x0][0x144];           /* 0x4c98078005170003 */
    /*0038*/                   FFMA R0, R0, R5, c[0x0][0x150];  /* 0x5180028005470000 */
                                                                /* 0x001ffc00ffe000f1 */
    /*0048*/                   STG.E [R2], R0;                  /* 0xeedc200000070200 */
    /*0050*/                   EXIT;                            /* 0xe30000000007000f */
    /*0058*/                   BRA 0x58;                        /* 0xe2400fffff87000f */
                                                                /* 0x001f8000fc0007e0 */
    /*0068*/                   NOP;                             /* 0x50b0000000070f00 */
    /*0070*/                   NOP;                             /* 0x50b0000000070f00 */
    /*0078*/                   NOP;                             /* 0x50b0000000070f00 */
    .................................

所以FFMA指令确实可以带一个额外的符号来应用于乘积（注意它在着色器汇编指令中应用于b，但是这给出了相同的结果）。您也可以尝试使用双精度操作数和其他计算功能来代替 sm_60，这会给您类似的结果。

带负号的 fma 中每个周期的指令数是多少？

What is the instruction number per cycle in fma with minus?

cuda

fma