cuda - 避免内联 PTX 中不必要的 mov 操作

Question

在单独的文件中写入 PTX 时，可以将内核参数加载到寄存器中：

.reg .u32 test;
ld.param.u32 test, [test_param];

但是，当使用内联 PTX 时，在 CUDA（版本 01）中使用内联 PTX 程序集应用说明描述了一种语法，其中加载参数与另一个操作密切相关。它提供了这个例子：

asm("add.s32 %0, %1, %2;" : "=r"(i) : "r"(j), "r"(k));

生成：

ld.s32 r1, [j];
ld.s32 r2, [k];
add.s32 r3, r1, r2;
st.s32 [i], r3;

在许多情况下，有必要将这两个操作分开。例如，可能希望将参数存储在循环外的寄存器中，然后在循环内重用和修改寄存器。我发现这样做的唯一方法是使用额外的 mov 指令，将参数从隐式加载的寄存器移动到我以后可以使用的另一个寄存器。

从单独文件中的 PTX 移动到内联 PTX 时，有没有办法避免这个额外的 mov 指令？

score 4 · Accepted Answer

如果我是你，我不会太担心那些 mov 操作。

请记住，PTX 不是最终的汇编代码。PTX 在内核启动之前被进一步编译成 CUBIN。其中，最后一步执行寄存器分配并将删除所有不必要mov的操作。

特别是，如果您从%r1to移动到%r2然后根本不使用%r1，则该算法可能会将%r1和分配%r2给相同的硬件寄存器并删除移动。

1 回答 1