c - 未知全局数组索引的归约或原子运算符

Question

我有以下算法:

__global__ void Update(int N, double* x, double* y, int* z, double* out)
{
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < N)
    {
      x[i] += y[i];
      if (y[i] >= 0.)
        out[z[i]] += x[i];
      else
        out[z[i]] -= x[i];
    }
}

重要的是要注意 out 小于 x。假设 x、y 和 z 总是相同的大小，比如 1000，out 总是更小，比如 100。z 是每个 x 和 y 对应的 in out 的索引。

除了out的更新外，这都是find。线程之间可能会发生冲突，因为 z 不仅包含唯一值而且有重复值。因此，我目前使用 atomicAdd 的原子版本实现了这一点。并使用比较和交换减去。这显然很昂贵，意味着我的内核运行时间要长 5-10 倍。

我想减少这个，但是我能想到的唯一方法是让每个线程都有自己的 out 版本(可以很大，10000+，X 10000+ 线程)。这意味着我设置 10000 double[10000](可能在共享中？)调用我的内核，然后对这些数组求​​和，也许在另一个内核中。肯定有更优雅的方法来做到这一点吧？

可能值得注意的是，x、y、z 和 out 驻留在全局内存中。由于我的内核(我有其他类似的内核)非常简单，我还没有决定跨位复制到共享(内核上的 nvvp 显示相等的计算和内存，所以我认为添加从移动数据的开销时不会获得太多性能全局到共享，然后再回来，有什么想法吗？)。

Answer 1

方法一:

1. 构建一组“事务”。由于每个线程只有一个更新，因此您可以轻松地构建一个固定大小的“事务”记录，每个线程一个条目。假设我的 out 表中有 8 个线程(为简单起见)和一些任意数量的条目。假设我的 8 个线程想要执行 8 个这样的事务:

   thread ID (i):  0      1      2      3      5      6      7
   z[i]:           2      3      4      4      3      2      3
   x[i]:           1.5    0.5    1.0    0.5    0.1    -0.2   -0.1
   "transaction":  2,1.5  3,0.5  4,1.0  4,0.5  3,0.1  2,-0.2 3,-0.1
   

2. 现在对交易做一个 sort_by_key，按照 z[i] 的顺序排列它们:

   sorted:         2,1.5  2,-0.2 3,0.5  3,-0.1 3,0.1  4,1.0  4,0.5
   

3. 现在对事务进行 reduce_by_key 操作:

   keys:           2      3      4    
   values:         1.3    0.5    1.5
   

4. 现在根据 key 更新out[i]:

             out[2] += 1.3
             out[3] += 0.5
             out[4] += 1.5
   

thrust和/或 cub可能是排序和归约操作的预置选项。

方法二:

如您所说，全局内存中有数组x、y、z 和out。如果您打算重复使用作为“映射”的 z，您可能需要按照 z 的顺序重新排列(分组)或排序数组:

    index (i): 0      1      2       3      4       5       6      7
         z[i]: 2      8      4       8      3       1       4      4
         x[i]: 0.2    0.4    0.3     0.1   -0.1    -0.4     0.0    1.0

按 z[i] 分组:

    index (i): 0      1      2       3      4       5       6      7
         z[i]: 1      2      3       4      4       4       8      8
         x[i]:-0.4    0.2   -0.1     0.3    0.0     1.0     0.4    0.1

这个，或者它的一些变体，可以让你不必在方法 1 中重复执行排序操作(同样，如果你重复使用相同的“映射” vector )。