IADD.X GPU指令

Question

查看为 NVIDIA Fermi 架构生成的 SASS 输出时，观察到指令 IADD.X。来自NVIDIAdocumentation，IADD是integer add的意思，但是不明白IADD.X是什么意思。有人可以帮忙吗...这是否意味着具有扩展位数的整数加法？

指令片段是：

IADD.X R5, R3, c[0x0][0x24];   /* 0x4800400090315c43 */

Answer 1

是的，.X 代表扩展精度。您会看到 IADD.X 与 IADD.CC 一起使用，其中后者将较低有效位相加，并产生一个进位标志（因此 .CC），然后将此进位标志合并到加法中IADD.X.

执行的更重要的位

由于 NVIDIA GPU 基本上是具有 64 位寻址能力的 32 位处理器，因此在地址（指针）算术中经常使用此成语。使用 64 位整数类型，例如 long long int 或 uint64_t 将同样导致使用这些指令。

这是内核执行 64 位整数加法的工作示例。此 CUDA 代码是针对 CUDA 7.5 的计算能力 3.5 编译的，机器代码使用 cuobjdump --dump-sass.

转储

__global__ void addint64 (long long int a, long long int b, long long int *res)
{
    *res = a + b;
}

MOV     R1, c[0x0][0x44];         
MOV     R2, c[0x0][0x148];        // b[31:0]
MOV     R0, c[0x0][0x14c];        // b[63:32]
IADD    R4.CC, R2, c[0x0][0x140]; // tmp[31:0] = b[31:0] + a[31:0]; carry-out
MOV     R2, c[0x0][0x150];        // res[31:0]
MOV     R3, c[0x0][0x154];        // res[63:32]
IADD.X  R5, R0, c[0x0][0x144];    // tmp[63:32] = b[63:32] + a[63:32] + carry-in
ST.E.64 [R2], R4;                 // [res] = tmp[63:0]
EXIT

IADD.X GPU指令

IADD.X GPU instruction

cuda

gpu

nvidia