数据自动化系统的FMAD格式

Question

我找不到用数据自动化系统解释下列指令格式的文件

FMAD R6, -R6, c [0x1] [0x1], R5;

格式是什么(源，目的地，.)那-R6是什么

Answer 1

PTX参考指南对fma的描述如下

fma.rnd{.ftz}{.sat}.f32  d, a, b, c;
fma.rnd.f64              d, a, b, c;

执行

d = a*b + c;

无论是单精度还是双精度。

您正在查看拆卸的SASS，用于显示FMAD的指令集参考文献是(非IEEE754兼容的)来自GT200指令集的单精度表单。这有点问题，因为我目前还没有一个支持这个废弃指令集的工具链。但是，如果我使用费米指令集并编译这个内核：

__global__ void kernel(const float *x, const float *y, float *a)
{
    float xval = x[threadIdx.x];
    float yval = y[threadIdx.x];

    float aval = -xval * xval + yval;
    a[threadIdx.x] = aval;:
}

我知道这个混蛋：

code for sm_20
    Function : _Z6kernelPKfS0_Pf
.headerflags    @"EF_CUDA_SM20 EF_CUDA_PTX_SM(EF_CUDA_SM20)"
    /*0000*/         MOV R1, c[0x1][0x100];                       /* 0x2800440400005de4 */
    /*0008*/         S2R R3, SR_TID.X;                            /* 0x2c0000008400dc04 */
    /*0010*/         MOV32I R5, 0x4;                              /* 0x1800000010015de2 */
    /*0018*/         IMAD.U32.U32 R8.CC, R3, R5, c[0x0][0x20];    /* 0x200b800080321c03 */
    /*0020*/         IMAD.U32.U32.HI.X R9, R3, R5, c[0x0][0x24];  /* 0x208a800090325c43 */
    /*0028*/         IMAD.U32.U32 R6.CC, R3, R5, c[0x0][0x28];    /* 0x200b8000a0319c03 */
    /*0030*/         LD.E R0, [R8];                               /* 0x8400000000801c85 */
    /*0038*/         IMAD.U32.U32.HI.X R7, R3, R5, c[0x0][0x2c];  /* 0x208a8000b031dc43 */
    /*0040*/         IMAD.U32.U32 R4.CC, R3, R5, c[0x0][0x30];    /* 0x200b8000c0311c03 */
    /*0048*/         LD.E R2, [R6];                               /* 0x8400000000609c85 */
    /*0050*/         IMAD.U32.U32.HI.X R5, R3, R5, c[0x0][0x34];  /* 0x208a8000d0315c43 */
    /*0058*/         FFMA.FTZ R0, -R0, R0, R2;                    /* 0x3004000000001e40 */
    /*0060*/         ST.E [R4], R0;                               /* 0x9400000000401c85 */
    /*0068*/         EXIT;                                        /* 0x8000000000001de7 */
    ..................................

请注意，我还在FFMA.FTZ参数中使用了否定式寄存器。所以我猜你的：

FMAD R6, -R6, c [0x1] [0x1], R5;

等于

R6 = -R6 * const + R5

其中，c [0x1] [0x1]是编译时间常数，并且GPU具有某种指令修饰符，它可以将其设置为控制浮点值的否定，作为浮点操作的一部分，而无需在调用之前显式地旋转寄存器的符号位。

(我期待@njuffa将这个答案撕成碎片)。