c++ - maxCUDAを使用した絶対差

Question

次のシリアルCコードが動作しています

2つのベクトル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

を使用して、この推力の低減を実装することも可能thrust::transform_reduceです。talonmiesが示唆するように、このソリューションは操作全体を融合します。

#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;
}

a[i] = r * a[i] + (1 - r) * b[i]このソリューションでは割り当てを実行しないことに注意してください。ただし、を使用して縮小した後に実行するのは簡単thrust::transformです。transform_reduceどちらのファンクターでもの引数を変更することは安全ではありません。

Answer 2

ループのこの2行目：

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

削減として知られています。幸い、CUDA SDKには削減サンプルコードがあります。これを見て、アルゴリズムのテンプレートとして使用してください。

これを2つの別々の操作（おそらく2つの別々のカーネルとして）に分割することをお勧めします。1つは並列部分（bp[]値の計算）用で、もう1つは縮小（calculate errors）用です。