在 CUDA 中有效地初始化共享内存数组

【问题标题】：Efficiently Initializing Shared Memory Array in CUDA在 CUDA 中有效地初始化共享内存数组
【发布时间】：2014-06-25 22:45:50
【问题描述】：

请注意，这个共享内存数组永远不会被写入，只会被读取。

正如我所拥有的那样，我的共享内存被初始化为：

__shared__ float TMshared[2592]; 
for (int i = 0; i< 2592; i++)
{
TMshared[i] = TM[i];

}
__syncthreads();

（TM 从内核启动传递到所有线程）

您可能已经注意到这是非常低效的，因为没有进行并行化，并且同一块中的线程正在写入同一位置。

由于所讨论的共享数组相对较小，如果这个问题真的需要优化，有人可以推荐一种更有效的方法/评论吗？

谢谢！

【问题讨论】：

标签： c++ memory cuda shared

【解决方案1】：

使用所有线程写独立的位置，可能会更快。

示例假设一维线程块/网格：

#define SSIZE 2592

__shared__ float TMshared[SSIZE]; 

  int lidx = threadIdx.x;
  while (lidx < SSIZE){
    TMShared[lidx] = TM[lidx];
    lidx += blockDim.x;}

__syncthreads();

【讨论】：

不错。 “#define SSIZE 2592”到底去哪儿了？在 cu 文件的顶部，在 global 内核之外？
另外，使用#define 有什么意义？与仅在适当的位置显式编码数字 2592 相比，它是否具有优势？
是的，定义通常放在文件的顶部，尽管我很确定你可以把它放在任何地方（在代码中使用它之前的任何地方）。与 2592 相比，define 没有明显的代码或性能优势。但是，如果我更改共享内存数组的大小，我只需在一处更改即可。