CUDA中的FMAD格式

我找不到解释CUDA中以下指令格式的文档

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

格式是什么(来源,目的地......),那是-R6什么？

PTX参考指南描述了如下fma

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

施行

d = a*b + c;

单精度或双精度.

您正在查看反汇编的SASS,该节目FMAD 的指令集引用是GT200指令集中的(非IEEE 754兼容)单精度形式.这有点问题,因为我目前没有支持不推荐使用的指令集的工具链.但是,如果我使用Fermi指令集而不是编译这个内核:

__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;:
}

我得到这个SASS:

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将这个答案撕成碎片).