GPU实现中的顺序操作

Question

我必须在GPU中实现以下算法

for(int I = 0; I < 1000; I++){
    VAR1[I+1] = VAR1[I] + VAR2[2*K+(I-1)];//K is a constant
}

每次迭代都依赖于先前的,因此并行化很困难.我不确定原子操作在这里是否有效.我能做什么？

编辑:

的VAR1和VAR2都是一维数组.

VAR1[0] = 1

Answer 1

这是一类称为递归关系的问题.根据递归关系的结构,可能存在闭合形式的解决方案,其描述如何单独地(即并行地,没有递归地)计算每个元素.其中一篇早期的开创性论文(关于并行计算)是Kogge和Stone,并且存在用于并行化特定形式的配方和策略.

有时递归关系是如此简单,以至于我们可以通过一点点"检查"来识别封闭形式的公式或算法.这个简短的教程为这个想法提供了更多的处理.

在你的情况下,让我们看看我们是否可以通过绘制出前几个VAR1应该是什么样的术语,用以前的术语替换为更新的术语来发现任何东西:

i      VAR1[i]
___________________
0        1
1        1 + VAR2[2K-1]
2        1 + VAR2[2K-1] + VAR2[2K]
3        1 + VAR2[2K-1] + VAR2[2K] + VAR2[2K+1]
4        1 + VAR2[2K-1] + VAR2[2K] + VAR2[2K+1] + VAR2[2K+2]
...

希望突然发现的是,上述VAR2[]术语遵循前缀总和的模式.

这意味着可以通过以下方式给出一种可能的解决方

VAR1[i] = 1+prefix_sum(VAR2[2K + (i-2)])   (for i > 0) notes:(1) (2)
VAR1[i] = 1                                (for i = 0)

现在,前缀sum可以并行完成(这不是真正完全独立的操作,但它可以并行化.我不想在这里过多讨论术语或纯度.我提供了一种可能的并行化方法对于你陈述的问题,不是唯一的方法.)要在GPU上并行执行前缀和,我会使用像CUB或Thrust这样的库.或者你可以写自己的,虽然我不推荐它.

笔记:

1. 使用-1或-2作为i前缀和的偏移量可以通过使用inclusive或exclusive扫描或前缀求和操作来决定.

2. VAR2必须在适当的域上定义才能使其合理化.但是,您的问题陈述中隐含了该要求.

这是一个微不足道的实例.在这种情况下,由于VAR2索引项2K+(I-1)只是表示I(2K-1)的固定偏移量,我们只是使用偏移量0来进行演示,因此VAR2只是一个简单的数组在同一个域上VAR1.为了演示目的,我正在定义VAR2为所有的数组1.gpu并行计算发生在VAR1向量中,CPU等效计算只是在cpu变量中即时计算以进行验证:

$ cat t1056.cu
#include <thrust/scan.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/transform.h>
#include <iostream>

const int dsize = 1000;
using namespace thrust::placeholders;
int main(){

  thrust::device_vector<int> VAR2(dsize, 1);  // initialize VAR2 array to all 1's
  thrust::device_vector<int> VAR1(dsize);
  thrust::exclusive_scan(VAR2.begin(), VAR2.end(), VAR1.begin(), 0); // put prefix sum of VAR2 into VAR1
  thrust::transform(VAR1.begin(), VAR1.end(), VAR1.begin(),  _1 += 1);   // add 1 to every term
  int cpu = 1;
  for (int i = 1; i < dsize; i++){
    int gpu = VAR1[i];
    cpu += VAR2[i];
    if (cpu != gpu) {std::cout << "mismatch at: " << i << " was: " << gpu << " should be: " << cpu << std::endl; return 1;}
    }
  std::cout << "Success!" << std::endl;
  return 0;
}

$ nvcc -o t1056 t1056.cu
$ ./t1056
Success!
$

对于要扫描操作的使用的附加参考特定求解线性复发问题,请参阅Blelloch的纸在这里第1.4节.