CUDA强制指令执行顺序

Question

我正在尝试将一些数据操作从 CPU 传输到 GPU (CUDA)，但有一小部分需要指令运行以特定顺序进行。原则上我可以做前几个并行部分，然后将串行部分的结果传输到主机，然后再将其余并行部分传回主机，但我试图避免内存传输开销。

计算的序列部分的形式为：

for (int i = 2; i<size; i++)
{
    result[i] = oldArray[i] + result[i-1];
}

除了在单个线程上为该计算启动内核之外，是否有办法强制线程或计算按特定顺序运行？

编辑：

这个问题比我最初展示的稍微复杂一些，据我所知，它不能作为前缀和问题解决。

循环实际上采用以下形式：

for (int i = 2; i<size; i++)
{
    result[i] = oldArray[i] + k * result[i-1];
}

我一直在查看 Thrust 库的文档，但似乎没有解决方案。但是，我可能只是不明白我在看什么。有类似的解决方案吗？

Answer 1

我们可以给这样的问题一个可能的描述是将它们放在递归关系.

的类别中

原题：

for (int i = 2; i<size; i++)
{
    result[i] = oldArray[i] + result[i-1];
}

如果需要遵循 .

中给出的描述，

可以通过 oldArray 上的前缀和轻松解决

经过编辑的修改：

for (int i = 2; i<size; i++)
{
    result[i] = oldArray[i] + k * result[i-1];
}

我们必须做额外的工作。参考先前链接的答案，在该答案的底部，Blelloch 给出了 this paper 的参考。如果我们研究该论文的第 1.4 节，我们可以观察到这个新问题公式符合第 1.4.1 节中描述的“一阶递归”模式，特别是公式 1.5。如果我们仔细指定 input/output 数据以及扫描运算符，则提供了如何使用 scan 操作实现该公式的解决方案的证明。

Thrust 能够支持对所提供的基本扫描进行这些泛化。 s和c所指的那篇论文中的对集合可以实现为thrust::tuple，并且可以将特定的运算符传递给推力扫描操作，以概括操作行为。

我不会试图涵盖那篇论文的所有内容；我们主要只需要关心第 48 和 49 页上的 material。

接下来是一个使用 thrust 的例子，证明我们可以完全按照论文中的描述使用 thrust 扫描操作来解决这个问题公式。下面的代码注释了引用 Blelloch 论文中特定公式的注释：

$ cat t1929.cu
#include <iostream>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/scan.h>
#include <thrust/iterator/zip_iterator.h>
#include <cstdlib>

template <typename T>
void cpufunction(T *result, T *oldArray, size_t size, T k){
  for (int i = 1; i<size; i++)
  {
    result[i] = oldArray[i] + k * result[i-1];
  }
}

struct scan_op // as per blelloch (1.7)
{
  template <typename T1, typename T2>
  __host__ __device__
  T1 operator()(const T1 &t1, const T2 &t2){
    T1 ret;
    thrust::get<0>(ret) = thrust::get<0>(t1)*thrust::get<0>(t2);
    thrust::get<1>(ret) = thrust::get<1>(t1)*thrust::get<0>(t2)+thrust::get<1>(t2);
    return ret;
    }
};

typedef float mt;
const size_t ds = 1048576;
const mt k = 1.01;
int main(){

  mt *b  = new mt[ds]; // b as in blelloch (1.5)
  mt *a  = new mt[ds]; // a as in blelloch (1.5)
  mt *cr = new mt[ds]; // cpu result
  for (int i = 0; i < ds; i++) { a[i] = k; b[i] = rand()/(float)RAND_MAX;}
  cr[0] = b[0];
  cpufunction(cr, b, ds, k);
  for (int i = 0; i < 10; i++) std::cout << cr[i] << ",";
  std::cout << std::endl;
  thrust::device_vector<mt> db(b, b+ds);
  thrust::device_vector<mt> da(a, a+ds);
  thrust::device_vector<mt> dy(ds);
  thrust::device_vector<mt> dx(ds);
  thrust::inclusive_scan(thrust::make_zip_iterator(thrust::make_tuple(da.begin(), db.begin())), thrust::make_zip_iterator(thrust::make_tuple(da.end(), db.end())), thrust::make_zip_iterator(thrust::make_tuple(dy.begin(), dx.begin())), scan_op());
  thrust::host_vector<mt> hx = dx;
  thrust::copy_n(hx.begin(), 10, std::ostream_iterator<mt>(std::cout, ","));
  std::cout << std::endl;
}
$ nvcc -std=c++14 t1929.cu -o t1929
$ ./t1929
0.840188,1.24297,2.0385,2.85733,3.79755,4.03307,4.40863,5.22094,5.55093,6.16041,
0.840188,1.24297,2.0385,2.85733,3.79755,4.03307,4.40863,5.22094,5.55093,6.16041,

Blelloch 描述的一阶递归允许或多或少任意 a 数组的可能性。在这个问题中，a 数组简单地由 k、k、k、...我们可以通过消除 a 来进一步简化它数组并将其替换为 thrust::constant_iterator。该练习相当机械，留给 reader.

CUDA强制指令执行顺序

CUDA force instruction execution order

cuda

thrust