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

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 的索引。

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

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

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

Answer 1

方法一：

建立一组"transactions"。由于每个线程只有一个更新，因此您可以轻松地构建一个固定大小的 "transaction" 记录，每个线程一个条目。假设我有 8 个线程（为简单起见）和我的 out table 中的一些任意数量的条目。假设我的 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
```

现在对交易执行 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

现在对交易执行 reduce_by_key 操作：

keys:           2      3      4    
values:         1.3    0.5    1.5

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

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

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

方法二：

如您所说，全局内存中有数组 x、y、z 和 out。如果您要重复使用 z 这是一个 "mapping"，您可能需要重新排列（分组）或按 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 中重复执行排序操作（同样，如果你重复使用相同的 "mapping" 向量）。

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

Reduction or atomic operator on unknown global array indices

c

cuda

atomic

reduction