CUDA中从全局内存加载数据到共享内存的指令是什么?
问题解答
你提到的这行shared_mem[i] = global_mem[i];赋值操作,会被CUDA编译器拆解成两条独立的PTX指令:
- 首先用
ld.global指令从全局内存读取数据到线程寄存器 - 接着用
st.shared指令把寄存器里的数据写入共享内存
这是因为该操作本质是「全局内存读取 + 共享内存写入」的两步动作,CUDA编译器会自动拆分这两个步骤,分别对应到对应的内存访问指令。
你可以通过nvcc编译时添加-ptx参数生成PTX代码来验证这一点,比如对你提供的代码执行:
nvcc -ptx your_code.cu
生成的.ptx文件里,就能看到循环内对应位置的ld.global和st.shared指令对,类似这样的片段:
ld.global.s32 %r1, [%rd+4]; st.shared.s32 [%rs+4], %r1;
内容的提问来源于stack exchange,提问作者Tae
相关产品推荐
相关产品推荐

