PAYATHON 2026

NVIDIA SASS 中为什么要对 R1 进行看似多余的赋值？

技术老齐

结论

MOV R1, c[0x0][0x28] 用于初始化栈指针：从系统常量区读取栈指针初值，写入编译器用作栈指针的寄存器 R1。

题目中的空 kernel 执行完 MOV 后就直接执行 EXIT，没有栈访问或函数调用，也没有使用这个值。从这段代码的数据流看，这次赋值是多余的。

c[0x0][0x28] 中存放的是什么？

这条指令中各部分的含义是：

  • c[0x0]：常量 bank 0，包含驱动、编译器使用的系统数据以及 kernel 参数。
  • [0x28]：该 bank 内的字节偏移。
  • R1：接收读取结果的 32 位寄存器。

按这里采用的编译约定，偏移 0x28 处存放的是栈指针初值。系统常量区与这一初始化序列的关系，可参见 GPU 内存布局研究；常量 bank 0 的用途可参见 NVIDIA 开发者论坛中的说明。

这个值用于线程的 local memory 栈寻址，不能直接当作普通的 64 位全局内存指针。

0x28 是读取位置的偏移，不是加载到 R1 的值。要知道实际加载的数值，需要在运行时观察，仅凭静态反汇编无法确定。

加载之后，R1 通常怎样使用？

当 kernel 需要栈空间时，后续代码可以调整 R1，再以它为基址访问栈中的数据。下面的 SASS 示意省略了无关指令：

MOV   R1, c[0x0][0x28]   // 初始化栈指针
IADD3 R1, R1, -0x20, RZ  // 调整栈指针，留出栈空间
STL   [R1], R2          // 写入 local memory
LDL   R3, [R1]          // 从 local memory 读取

这里的 0x20 只是示例，实际栈帧大小取决于编译结果。实际程序中的类似栈指针调整和 STL、LDL 访问序列，可参见 GPU Memory Exploitation for Fun and Profit。

题目中的代码只有：

MOV R1, c[0x0][0x28]
EXIT

后续没有指令读取 R1 或通过它访问内存，因此这个值没有参与后续计算。

为什么编译器仍然保留这条指令？

从生成结果看，可以把它理解为编译器保留的入口序言：先初始化栈指针，而这个空 kernel 恰好用不到栈。

这条指令的用途与编译器保留它的原因是两个问题：

  1. 指令用于初始化栈指针。
  2. 编译器为何没有消除这次无用赋值，取决于具体的代码生成实现，单凭这两条指令无法确定。

这两条指令也不足以证明该初始化是硬件强制要求，或 EXIT 必须依赖 R1。“为了调试器而保留”的说法同样需要额外证据。

如何自行检查？

把空 kernel 保存为 empty.cu：

__global__ void do_nothing(void) {}

然后生成目标架构的 cubin，检查指令和资源使用情况：

nvcc -arch=sm_89 -cubin empty.cu -o empty.cubin
cuobjdump --dump-sass empty.cubin
cuobjdump --dump-resource-usage empty.cubin

--dump-sass 用于查看机器指令，--dump-resource-usage 用于查看包括栈在内的资源使用情况。选项含义见 CUDA Binary Utilities 文档。

检查时要分别看指令和资源报告。仅有栈指针初始化，不能证明 kernel 实际使用了非零大小的栈帧。上述命令提供了复查方法，这里尚未复现 CUDA 13.2 的编译输出。

R1 的用途和系统常量区偏移，需要结合目标架构、工具链和具体反汇编来理解。对于这个空 kernel，这条指令没有改变计算结果，无需为此修改 CUDA 源码。

备注：内容仅供参考。