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

如何解决%f, %rd 在 ptx 程序集中是什么意思

嗨，我刚接触 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,%r8;

[已编辑]

Million 感谢 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

解决方法

PTX 寄存器命名汇总here。 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 指的是 predicate register。可以将谓词寄存器视为保存单个布尔真/假数量。

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

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