%f, %rd 在 ptx 程序集中是什么意思

Question

嗨，我是 CUDA 编程的新手。我从使用 OpenCL 构建程序中获得了这段汇编代码。

我开始想知道那些数字和字符是什么意思。比如 %f7, %f11, %rd3, %r3, %f, %p.

我猜 rd 可能指的是寄存器？数字是寄存器号？，也许百分比只是将操作数写入 ptx 命令的一种方式（即 ld.shared.f32）？如果我的猜测是正确的，那么 %r3 是什么意思，它像一个不同的 class 寄存器吗？还有 %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

Answer 1

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

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

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

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

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

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

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

p指的是predicate register。谓词寄存器可以被认为是保存单个布尔值 true/false 数量。

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

None 其中直接与 CUDA 相关，它是 PTX 的一部分，记录在案 here。

%f, %rd 在 ptx 程序集中是什么意思

What does %f, %rd mean in ptx assembly

assembly

opencl

cpu-registers

ptx