在CUDA中，使用什么指令将数据从全局内存加载到共享内存？

Question

我目前正在研究CUDA，了解到有全局内存和共享内存。

我查看了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?
    }
}

Answer 1

如果是

__shared__ float smem[2];
smem[0] = global_memory[0];

那么操作是（在SASS中）

LDG  Rx, [Ry]
STS  [Rz], Rx

要进一步扩展，请阅读https://forums.developer.nvidia.com/t/whats- Different- Between-ld-and-ldg-load-from-generic-memory-vs-load-from-global-memory /40856/2

概括：

操作说明	意义
LDS	从共享空间加载
LDC	从恒定空间加载
LDG	从全局空间加载
LD	通用负载 - 从提供的地址推导出的空间