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 IN着色器组装指令，但是这给出了相同的结果）。 您也可以尝试使用双精度操作数和其他计算功能，而不是sm_60，这会给您类似的结果。