fma中每个循环的指令号是多少？

Question

如果在cuda中使用fma(a，b，c)，就意味着公式a_b+c是在一次三元运算中计算出来的。但是，如果我想计算-a_b+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，这将给出类似的结果。