c++ - 使用 CUDA 的最大绝对差

Question

我们有以下串行 C 代码在运行

两个向量 a[] 和 b[]：

double a[20000],b[20000],r=0.9;

for(int i=1;i<=10000;++i)
{
    a[i]=r*a[i]+(1-r)*b[i]];
    errors=max(errors,fabs(a[i]-b[i]);
    b[i]=a[i];
}

请告诉我们如何将此代码移植到 CUDA 和 cublas？

Answer 1

It's also possible to implement this reduction in Thrust using thrust::transform_reduce. This solution fuses the entire operation, as talonmies suggests:

#include <thrust/device_vector.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/transform_reduce.h>
#include <thrust/functional.h>

// this functor unpacks a tuple and then computes
// a weighted absolute difference of its members
struct weighted_absolute_difference
{
  double r;

  weighted_absolute_difference(const double r)
    : r(r)
  {}

  __host__ __device__
  double operator()(thrust::tuple<double,double> t)
  {
    double a = thrust::get<0>(t);
    double b = thrust::get<1>(t);

    a = r * a + (1.0 - r) * b;

    return fabs(a - b);
  }
};

int main()
{
  using namespace thrust;

  const std::size_t n = 20000;

  const double r = 0.9;

  device_vector<double> a(n), b(n);

  // initialize a & b
  ...

  // do the reduction
  double result =
    transform_reduce(make_zip_iterator(make_tuple(a.begin(), b.begin())),
                     make_zip_iterator(make_tuple(a.end(),   b.end())),
                     weighted_absolute_difference(r),
                     -1.f,
                     maximum<double>());

  // note that this solution does not set
  // a[i] = r * a[i] + (1 - r) * b[i]

  return 0;
}

Note that we do not perform the assignment a[i] = r * a[i] + (1 - r) * b[i] in this solution, though it would be simple to do so after the reduction using thrust::transform. It is not safe to modify transform_reduce's arguments in either functor.

Answer 2

循环中的第二行：

errors=max(errors,fabs(a[i]-b[i]);

被称为减少。幸运的是，CUDA SDK 中有简化示例代码 - 看看这个并将其用作您的算法的模板。

您可能希望将其拆分为两个单独的操作（可能作为两个单独的内核） - 一个用于并行部分（计算bp[]值），另一个用于减少（计算errors）。