关于NVIDIA SASS中R1寄存器赋值的技术疑问咨询

背景信息 #

原CUDA内核代码：

__global__ void do_nothing(void) { }    

使用CUDA 13.2针对SM 89架构编译后，生成的SASS代码：

do_nothing():
 MOV R1, c[0x0][0x28] 
 EXIT 

疑问解答 #

1. 从常量内存加载至寄存器R1的内容是什么？

   • 常量内存地址c[0x0][0x28]对应的是线程块的内核启动参数基地址，加载的是存储内核启动配置信息（如网格/块维度、传入的内核参数等）的内存区域起始地址。

2. 完成加载后，寄存器R1会有何变化？该加载值的用途是什么？

   • 加载完成后，R1寄存器的值会被更新为上述的参数基地址。
   • R1是CUDA硬件规范约定的内核参数专用寄存器，所有CUDA内核都会隐式执行这个加载操作。它的作用是提供内核访问启动参数的入口：如果内核有输入参数，后续会通过R1加上指定偏移量读取对应参数；即使是空内核，编译器也会生成这段代码，保证内核运行环境符合CUDA硬件的标准要求。

内容的提问来源于stack exchange，提问作者einpoklum