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

【问题标题】：how to efficiently load a big array to GPU shared memory?如何有效地将大数组加载到 GPU 共享内存？
【发布时间】：2014-11-16 13:28:17
【问题描述】：

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

1: int index = threadidx.x;

2: 共享无符号字符 x[1000];

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

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

任何建议将不胜感激。

【问题讨论】：

标签： c cuda shared-memory memory-access

【解决方案1】：

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

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

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

【讨论】：

通过这种方式，将数组元素加载到共享内存中是一次并并行发生的吗？
唯一的“并行”“一次”活动将在经线级别。每个warp（32个线程）会将它们对全局内存（array[i]）的访问组合成一个访问/事务，并将它们对共享内存（x[i]）的访问组合成一个访问。多个warp也可以并行执行，但它们的活动不会合并到一个内存事务中。
每个线程至少加载 4 个字节通常会获得更好的加载/存储效率。即使您的底层数据结构是面向字节的，这也可以使用适当的指针算法通过基本的指针转换（例如char * 到int *）来完成。
如果我通过每个线程加载 4 字节数据与从全局内存加载 1 字节数据有区别吗？我可以通过线程从全局内存中加载的最大数据量是多少？如果我可以并行加载 32 个数据值，我可以按每个线程加载最大的数据值吗？它会影响内存访问吗？
也许你应该提出一个新问题。 programming guide 中涵盖了单个线程在单个访问周期中可以检索的最大数据量，它是 16 字节（例如 uint4 向量类型）。由于 32 个线程在一个 warp 中，每个线程加载 4 个字节将产生 128 个字节的负载，这将充分利用（100% 效率）全局加载请求。与此相比，通过将每线程负载增加到 4 个字节以上，没有额外的好处。

【解决方案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();

【讨论】：

blocksize[i] 在您的代码中没有意义。概述了一种处理线程块和数据数组任意大小的方法here
@RobertCrovella 很抱歉，这是一个错字。
@RobertCrovella 链接中的方法也等同于我的方法，只是使用了 C 语法。如果你喜欢这种事情，那很好。