在共享内存中暂存尴尬的长度数组

Question

x 是全局内存中长度为 N 的数组，由 k 块每个 w 线程的 cuda/opencl 内核操作（因此 k = ceil(N/w)）。内核中的每个块都有一个长度为 w 的本地共享数组 xlocal。任务是让每个块将它们的 x 块加载到 xlocal 中。

如果 w 正好整除 N，那么我们可以这样做：

int lid = threadIdx.x;
int gid = threadIdx.x + (blockIdx.x * blockDim.x);
xlocal[lid] = x[gid];

如果不是，那么我们在最后一个块中有 (N%w) 个冗余线程。我们应该如何处理它们？我可以想到以下选项：

1. 为 x 分配更大的长度。即，分配 k*w 元素而不是 N。这很有用，因为上面的代码可以正常工作。不幸的是，我认为 cuda 或 opencl 中没有 realloc 等价物。

2. 在加载之前进行范围检查。这很好，因为我们不需要搞乱 x 的分配。但是仅仅因为边缘条件就给大多数线程添加工作很烦人。

   if (gid < N) xlocal[lid] = x[gid];
   

3. 从 x 模 N 加载，以便冗余线程环绕：

   xlocal[lid] = x[gid%N];
   

对于解决这个问题还有其他想法吗？

---

一些基准测试

以下是选项 (2) 范围检查（蓝色）与选项 (3) 加载模 N（红色）的比较结果。

我们将块大小固定为 32 个线程，并将 N 从 45.6k 变为 45.6k+32，以分别在最后一个块中提供 0 到 32 个冗余线程。该测试运行一个简单的内核，该内核从全局内存中预加载一个共享数组。左边（/右边）的图表为每个线程加载一个（/三个）元素。我使用 cuda 3.2.16 标志 -O2 编译并在 Tesla M2070 卡上运行。

Answer 1

1. 您可以从主机分配更大的 x。然后，您应该考虑可能引入的无用的额外复制时间，以及内存空间。这也会让你的代码失去意义和结构。

2. ：

1234563合并，但仍然）。

在我看来 2（或者 3）是你最好的选择 您只需向每个线程添加几条指令。考虑到您的代码将保持清晰且不言自明，因此无需担心太多。

你应该避免选项 1。