cuda - CUDA 中的 FMAD 格式

Question

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

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;

单精度或 double 。

您正在查看反汇编的SASS，指令集references为此，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 将这个答案撕成碎片)。