c - 如何有效地将大数组加载到 GPU 共享内存？

Question

我想将一个大数组加载到 GPU 共享内存。当我像下面这样使用它时：

1：int索引=threadidx.x；

2：共享无符号字符 x[1000]；

3: x[i] = 数组[i];

那么如果我们调用一个有 1000 个线程和一个块的内核代码，每个线程都会发生一次内存访问？是否可以通过一次内存访问加载此数组并将其加载到共享内存？

任何建议将不胜感激。

Answer 1

不，它不能通过单一访问来完成。

正如您所展示的，并行使用线程来加载共享内存是最快的方法。共享内存只能由 CUDA 内核中的线程执行的内存操作加载。没有用于加载共享内存的 API 函数。

如果您的数组大于线程块中的线程数，则可以使用类似于此处概述的循环方法。

Answer 2

通常，范例是在块上使用 for 循环。请记住，共享内存仅在块之间共享，因此如果您的共享内存大于块的数量，您将需要一个花哨的循环。

在内核的开头，有这样的东西。

//loadbuff is a T*, T is whatever type you want
//sharedmemsize is some (compile time) constant
__shared__ T sharedmem[sharedmemsize]; 
int index = threadidx.x;
int blocksize = blockDim.x;

for (int i = index; i < sharedmemsize; i+=blocksize)
{
   sharedmem[i]=loadbuff[i];
}
__syncthreads();