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，這會給您類似的結果。