3

我一直在考虑如何使用归约在 CUDA 上执行此操作，但我对如何完成它有点茫然。C代码如下。要记住的重要部分——变量precalculatedValue取决于两个循环迭代器。此外，变量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的值。因此，我看不出有一种方法可以使此代码高效（甚至可行）以使用减少。任何想法都是最受欢迎的。

4

1 回答 1

2

就一个刺...

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

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

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

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

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

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

请注意，此计算看起来与直方图计算非常相似，后者具有类似的挑战，因此您可能想在谷歌上搜索 GPU 直方图以了解它们是如何实现的。

一种想法是使用and 然后对 ( rInd[m], m) 对进行排序（因为重复项将被组合在一起），您可以遍历它们并在没有冲突的情况下进行归约。（这是制作直方图的一种方法。）您甚至可以使用.thrust::sort_by_key()rIndthrust::reduce_by_key()

于 2012-03-02T03:58:33.747 回答