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

Posted 2023-02-16

【中文标题】fma 中每个周期的指令数是多少，带有负数？

【英文标题】：What is the instruction number per cycle in fma with minus?

【发布时间】：2017-09-02 07:11:59

【问题描述】：

如果我在cuda中使用fma(a, b, c)，则表示公式ab+c是在一次三元运算中计算出来的。但是如果我想计算-ab+c，调用fma(-a, b, c) 是否需要再进行一次乘法运算？

【参考方案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，这将得到类似的结果。

【讨论】：

由于汇编语言表示的限制（FNMA 没有单独的助记符），在反汇编代码中，对乘积 a*b 的否定始终显示为对 b 操作数的否定。跨度>