ptx 汇编中 %f, %rd 是什么意思

Question

你好，我是 CUDA 编程新手。我通过使用 OpenCL 构建程序获得了这段汇编代码。

我开始想知道这些数字和字符意味着什么。如%f7、%f11、%rd3、%r3、%f、%p。

我猜这rd可能指的是寄存器？数字是寄存器号？也许百分比只是将操作数写入 ptx 命令的一种方式（即 ld.shared.f32）？如果我的猜测是正确的，那么 %r3 是什么意思，它就像不同类别的寄存器吗？还有 %p 和 %f7。

先感谢您。

    ld.global.f32   %f7, [%rd16];
    st.shared.f32   [%rd2], %f7;
    bar.sync    0;
    ld.shared.f32   %f8, [%rd4];
    ld.shared.f32   %f9, [%rd3];
    fma.rn.f32  %f10, %f9, %f8, %f32;
    ld.shared.f32   %f11, [%rd4+32];
    ld.shared.f32   %f12, [%rd3+4];
    fma.rn.f32  %f13, %f12, %f11, %f10;
    ld.shared.f32   %f14, [%rd4+64];
    ld.shared.f32   %f15, [%rd3+8];
    fma.rn.f32  %f16, %f15, %f14, %f13;
    ld.shared.f32   %f17, [%rd4+96];
    ld.shared.f32   %f18, [%rd3+12];
    fma.rn.f32  %f19, %f18, %f17, %f16;
    ld.shared.f32   %f20, [%rd4+128];
    ld.shared.f32   %f21, [%rd3+16];
    fma.rn.f32  %f22, %f21, %f20, %f19;
    ld.shared.f32   %f23, [%rd4+160];
    ld.shared.f32   %f24, [%rd3+20];
    fma.rn.f32  %f25, %f24, %f23, %f22;
    ld.shared.f32   %f26, [%rd4+192];
    ld.shared.f32   %f27, [%rd3+24];
    fma.rn.f32  %f28, %f27, %f26, %f25;
    ld.shared.f32   %f29, [%rd4+224];
    ld.shared.f32   %f30, [%rd3+28];
    fma.rn.f32  %f32, %f30, %f29, %f28;
    bar.sync    0;
    add.s32     %r37, %r37, 8;
    add.s32     %r36, %r36, %r11;
    add.s32     %r38, %r38, 1;
    setp.lt.s32 %p5, %r38, %r8;

---

[编辑]

非常感谢 Robert Crovella 的详尽回答！以防万一有人可能想知道，这是我的 ptx 文件顶部的寄存器声明部分（？）

    .reg .pred  %p<6>;
    .reg .f32   %f<33>;
    .reg .b32   %r<39>;
    .reg .b64   %rd<19>;
    .shared .align 4 .b8 sgemm$blockA[256];
    // demoted variable
    .shared .align 4 .b8 sgemm$blockB[256];

共享寄存器大小为 256，我将其设置为 16 * 16。

参考文档的具体部分在这里

Answer 1

PTX 寄存器命名总结如下。PTX 具有虚拟寄存器约定，这意味着寄存器实际上是变量名称，它们不一定对应于物理设备中的硬件寄存器。因此，如此处所示，对这些内容的实际解释需要比此处的代码片段更多的 PTX 代码。（虚拟寄存器在使用之前已正式声明。）具体来说，您通常会找到一组如下所示的声明：

    .reg .pred      %p<11>;
    .reg .f32       %f<3075>;
    .reg .b32       %r<54>;
    .reg .b64       %rd<10>;

在任何完整 PTX 代码的“顶部”，它将定义实际的虚拟寄存器命名/定义。

但是我们可以依靠编译器前端通常用来生成这些虚拟寄存器名称的一些“约定”来回答您的问题，用于指导目的，而不是陈述实际的“规范”。

%rXY当用作指令的操作数时，指的是这些寄存器之一，其中XY是寄存器编号，例如 30。根据下面的变化，r通常指的是用于表示用于保存整数的 32 位寄存器的寄存器、二进制或地址信息。

rd指双寄存器，即寄存器对，即64位寄存器。您会注意到代码中的使用rd主要与寻址有关，因此它是 64 位是有意义的。

f指浮点寄存器。（f通常用于指代 32 位浮点寄存器，而fd通常用于指代 64 位浮点寄存器。）

p指谓词寄存器。谓词寄存器可以被认为保存单个布尔真/假量。

是的，该数字指的是（该类型的）特定寄存器。

这些都与 CUDA 没有直接关系，它是 PTX 的一部分，记录在此处。