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)

现在，前缀和可以并行完成（这不是真正完全独立的操作，但可以并行化。 我不想在这里争论太多关于术语或纯度的问题。 我为您陈述的问题提供了一种可能的并行化方法，而不是唯一的方法。 要在 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节。 这个问题/答案给出了一个示例，说明如何在更一般的一阶递归情况下实现该论文中的方程 1.5。这个问题考虑的是二阶重复情况。