嵌套循环中数组的二维累积和——CUDA实现？

Question

我一直在考虑如何使用归约在 CUDA 上执行此操作，但对于如何完成它我有点茫然。 C代码如下。要记住的重要部分——变量 precalculatedValue 取决于 both 循环迭代器。此外，变量 ngo 并不是每个 m 的值都是唯一的......例如m = 0,1,2 可能有 ngo = 1，而 m = 4,5,6,7,8 可能有 ngo = 2 等。我已经包含了循环迭代器的大小，以防它有助于提供更好的实现建议。

// macro that translates 2D [i][j] array indices to 1D flattened array indices
#define idx(i,j,lda) ( (j) + ((i)*(lda)) )

int Nobs = 60480;
int NgS  = 1859;
int NgO  = 900;
// ngo goes from [1,900]

// rInd is an initialized (and filled earlier) as:
// rInd = new long int [Nobs];

for (m=0; m<Nobs; m++) {        
    ngo=rInd[m]-1;

    for (n=0; n<NgS; n++) {
            Aggregation[idx(n,ngo,NgO)] += precalculatedValue;
    }
}

在前面的例子中，当 precalculatedValue 只是内部循环变量的函数时，我将值保存在唯一的数组索引中，并在事后通过并行缩减（推力）将它们添加。然而，这种情况让我很困惑：m 的值并没有唯一地映射到 ngo 的值。因此，我看不出有一种方法可以使此代码高效（甚至可行）以使用减少。欢迎任何想法。

Answer 1

只是一个刺...

我怀疑转置你的循环可能会有所帮助。

for (n=0; n<NgS; n++) {
    for (m=0; m<Nobs; m++) {            
        ngo=rInd[m]-1;
        Aggregation[idx(n,ngo,NgO)] += precalculatedValue(m,n);
    }
}

我这样做的原因是因为idx 与ngo（m 的函数）的变化比n 的变化更快，因此使 m 成为内循环提高了连贯性。请注意，我还将 precalculatedValue 设为 (m, n) 的函数，因为您说它是 - 这使伪代码更清晰。

然后，您可以先将外循环留在主机上，然后为内循环制作内核（64,480 路并行性足以满足大多数当前 GPU）。

在内部循环中，首先使用 atomicAdd() 来处理冲突。如果它们不常见，那么在 Fermi GPU 上的性能应该不会太差。在任何情况下，您都将受到带宽限制，因为该计算的算术强度很低。因此，一旦这工作正常，请测量您正在实现的带宽，并与您的 GPU 的峰值进行比较。如果您还差得远，然后考虑进一步优化（也许通过并行化外循环 - 每个线程块一次迭代，并使用一些共享内存和线程协作优化来执行内循环）。 p>

关键：从简单开始，衡量性能，然后决定如何优化。

请注意，此计算看起来与直方图计算非常相似，但也有类似的挑战，因此您可能需要在 Google 上搜索 GPU 直方图以了解它们是如何实现的。

一个想法是使用 thrust::sort_by_key() 对 (rInd[m], m) 对进行排序，然后（因为 rInd 重复项将被组合在一起），您可以迭代它们并进行减少而不会发生冲突。 （这是制作直方图的一种方法。）您甚至可以使用 thrust::reduce_by_key() 来做到这一点。