在CUDA中，用什么指令将数据从全局内存加载到共享内存？
在CUDA中，用什么指令将数据从全局内存加载到共享内存？
提问于 2022-11-15 05:05:08
我目前正在学习数据自动化系统，并了解到有全局记忆和共享记忆。
我检查了CUDA文档，发现GPU可以分别使用ld.shared/st.shared和ld.global/st.global指令访问共享内存和全局内存。
我好奇的是，用什么指令将数据从全局内存加载到共享内存？
如果有人能告诉我那就太好了。
谢谢!
__global__ void my_function(int* global_mem)
{
__shared__ int shared_mem[10];
for(int i = 0; i < 10; i++) {
shared_mem[i] = global_mem[i]; // What instrcuton is used for this load operation?
}
}回答 3
Stack Overflow用户
回答已采纳
发布于 2022-11-15 05:09:18
在以下情况下
__shared__ float smem[2];
smem[0] = global_memory[0];那么操作是
LDG Rx, [Ry]
STS [Rz], Rx摘要
LDS是来自共享空间的负载，最不发达国家是从常量空间加载的，LDG是从全局空间加载的，LD是从提供的地址导出的通用负载->。
Stack Overflow用户
发布于 2022-11-15 08:04:29
在NVIDIA的安培微体系结构中，引入了流水线功能，以提高从全局内存复制到共享内存的性能。因此，我们不再需要每个元素加载两个指令，这使得线程比所需的更繁忙。相反，您可以这样写：
#define NO_ZFILL 0
// ...
for(int i = 0; i < 10; i++) {
__pipeline_memcpy_async(&shared_mem[i], &global_mem[i], sizeof(int), NO_ZFILL);
}
__pipeline_commit();
__pipeline_wait_prior(0); // wait for the first commited batch of pipeline ops生成的PTX代码看起来如下所示：
{
ld.param.u64 %rd1, [my_function(int*)_param_0];
mov.u32 %r1, my_function(int*)::shared_mem;
cp.async.ca.shared.global [%r1], [%rd1], 4, 4;
add.s64 %rd2, %rd1, 4;
add.s32 %r2, %r1, 4;
cp.async.ca.shared.global [%r2], [%rd2], 4, 4;
add.s64 %rd3, %rd1, 8;
add.s32 %r3, %r1, 8;
cp.async.ca.shared.global [%r3], [%rd3], 4, 4;
add.s64 %rd4, %rd1, 12;
add.s32 %r4, %r1, 12;
cp.async.ca.shared.global [%r4], [%rd4], 4, 4;
add.s64 %rd5, %rd1, 16;
add.s32 %r5, %r1, 16;
cp.async.ca.shared.global [%r5], [%rd5], 4, 4;
add.s64 %rd6, %rd1, 20;
add.s32 %r6, %r1, 20;
cp.async.ca.shared.global [%r6], [%rd6], 4, 4;
add.s64 %rd7, %rd1, 24;
add.s32 %r7, %r1, 24;
cp.async.ca.shared.global [%r7], [%rd7], 4, 4;
add.s64 %rd8, %rd1, 28;
add.s32 %r8, %r1, 28;
cp.async.ca.shared.global [%r8], [%rd8], 4, 4;
add.s64 %rd9, %rd1, 32;
add.s32 %r9, %r1, 32;
cp.async.ca.shared.global [%r9], [%rd9], 4, 4;
add.s64 %rd10, %rd1, 36;
add.s32 %r10, %r1, 36;
cp.async.ca.shared.global [%r10], [%rd10], 4, 4;
cp.async.commit_group;
cp.async.wait_group 0;
ret;
}关于PTX的说明：
- 关键指令是以
cp.async开头的指令，而add的指令是地址计算。 - 使用目标虚拟体系结构sm_80编译。
- 编译器已经展开了循环(尽管它不必这样做)。
- 这仍然需要进一步编译成实际的装配指令。
有关更多细节，请参见CUDA编程指南中的节CUDA编程指南。
有一种使用捆绑在一起的“合作组”C++接口的更高级、但更不透明的方法。
Stack Overflow用户
发布于 2022-11-15 07:50:00
正如Kryrene已经说过的，将数据从全局内存加载到共享内存通常需要两个指令。
但是，由于安培体系结构(CC >= 8.0)，GPU也支持使用单一指令将数据直接从全局内存加载到共享内存中，从而发出异步副本。运营
在ptx中，这可能是cp.async.ca.shared.global，在SASS中可能是LDGSTS。
页面原文内容由Stack Overflow提供。腾讯云小微IT领域专用引擎提供翻译支持
原文链接：
https://stackoverflow.com/questions/74440827
复制相关文章
相似问题
腾讯云开发者
