ptx 汇编中的 %f, %rd 是啥意思

Posted 2023-02-16

【中文标题】ptx 汇编中的 %f, %rd 是啥意思

【英文标题】：What does %f, %rd mean in ptx assemblyptx 汇编中的 %f, %rd 是什么意思

【发布时间】：2021-05-15 20:25:12

【问题描述】：

您好，我刚接触 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。

并且参考文档的具体部分是here

【参考方案1】：

PTX 寄存器命名总结为here。 PTX 具有虚拟寄存器约定，这意味着寄存器实际上是变量名称，它们不一定对应于物理设备中的硬件寄存器。因此，正如那里所指出的，这些的实际解释需要比您在此处拥有的 sn-p 更多的 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 指的是predicate register。谓词寄存器可以被认为是保存一个布尔真/假数量。

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

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