CUDA inline PTX ld.shared触发cudaErrorIllegalAddress错误相关问题

1. 你编写的inline PTX存在的问题 #

你的inline PTX写法有两处直接触发cudaErrorIllegalAddress错误的问题：

• 地址传递错误：你将共享内存指针&As[TY][k]强制转换为int类型传递，在64位CUDA架构下设备指针为64位宽度，直接截断为32位int会生成非法地址，是触发700错误的直接原因之一。
• 缺少依赖声明：asm块没有添加volatile修饰符，也没有声明memory clobber。编译器不会解析asm字符串内部的指令逻辑，默认只会跟踪显式声明的输入输出操作数的依赖，不知道这条PTX指令会读取As数组的内存，因此会将这条ld指令和之前的As数组写入、__syncthreads()同步操作做非法重排，导致LDS指令执行时As对应位置还没有完成写入，甚至地址计算都未完成。

2. inline PTX对编译器依赖跟踪的影响 #

inline PTX本身不会干扰编译器的依赖跟踪能力，你遇到的依赖丢失是约束声明不全导致的：

• 编译器不会主动解析asm指令字符串的语义，所有内存访问、依赖关系都需要开发者通过输入输出约束、clobber列表显式告知编译器。
• 只要正确声明所有输入、输出、以及内存访问侧效，编译器会完全按照依赖关系进行指令调度，不会出现非法重排的问题。

修复方案 #

修改asm语句为如下写法即可解决问题：

// 加volatile禁止编译器优化/重排该指令，声明memory clobber告知编译器该指令访问内存
asm volatile("ld.shared.f32 %0, [%1];" : "=f"(t) : "r"(&As[TY][k]) : "memory");

如果不需要全局内存屏障，也可以用"m"约束直接传递共享内存变量，编译器会自动处理地址生成和依赖跟踪：

asm volatile("ld.shared.f32 %0, %1;" : "=f"(t) : "m"(As[TY][k]));

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