未知全局数组索引的约简或原子算子

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每一个对应的指数。

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

不过，我想减少这种情况，唯一能想到的方法是让每个线程都有自己的out版本(可以是大型的，10000+，X 10000+线程)。这意味着我设置了10000双10000调用我的内核，然后对这些数组进行求和，也许是在另一个内核中。肯定有更优雅的方法来做这件事吧？

值得注意的是，x、y、z和out驻留在全局内存中。由于我的内核(我有这样的其他内核)非常简单，所以我还没有决定跨位复制到共享(内核上的nvvp显示相同的计算和内存，所以我认为在增加将数据从全局迁移到共享并再次返回的开销时，没有多少性能可获得，有想法吗？)

Answer 1

方法1：

1. 构建一组“事务”。因为每个线程只有一个更新，所以您可以轻松地构建一个固定大小的“事务”记录，每个线程一个条目。假设我在out表中有8个线程(为了简化表示)和一些任意数目的条目。让我们假设我的8个线程想要执行这样的8个事务： 线程ID (i)：0 1 2 3 5 6 7 zi: 2 3 4 4 3 2 3 xi: 1.5 0.5 1.0 0.5 0.1 -0.2 -0.1“事务”：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]的顺序排列它们： 排序: 2,1.5 2，-0.2，0.5 3，-0.1，0.1，1.0，0.5
3. 现在对事务执行reduce_by_key操作： 键:2 3 4值: 1.3 0.5 1.5
4. 现在，根据键更新out[i]： out2 += 1.3 out3 += 0.5 out4 += 1.5

推力和/或幼崽可能是预构建的排序和减少操作的选项。

方法2：

正如您所说，全局内存中有数组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

zi组：

    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中重复执行排序操作(同样，如果您重复使用相同的“映射”向量)。