2

2 つのベクトル (oldvectorとnewvector) があります。次の擬似コードで定義される残差の値を計算する必要があります。

residual = 0;
forall i : residual += (oldvector[i] - newvector[i])^2

現在、基本的に行っている 2 つの CUDA Thrust 操作でこれを計算しています。

forall i : oldvector[i] = oldvector[i] - newvector[i]

thrust::transform_reduce単項演算子として四角形の a が続きます。これは次のことを行っています。

residual = 0;
forall i : residual += oldvector[i]^2;

これに伴う問題は、明らかに の前のグローバル メモリへの中間ストアtransform_reduceです。これらの 2 つのステップを融合させる、この問題に対するより効率的なアプローチはありますか? 独自の CUDA カーネルを作成する以外に、他に選択肢はありますか?

私が考えた 1 つのアプローチは、thrust::reducewith zip イテレータを記述することでした。これの問題は、演算子の戻り値の型が入力と同じ型でなければならないことです。私によると、これが意味することは、リダクション演算子が余分な追加を意味するタプルを返すということです。

リダクション CUDA カーネルを作成する場合、リダクション カーネルの CUDA 1.1 の例よりも改善されていますか?

4

2 に答える 2

3

Thrust::inner_productは、単一の関数呼び出しでそれを行います。元のアイデアも機能させることができます (2 つのベクトルを一緒に圧縮して を使用thrust::transform_reduce)。このコードは両方の方法を示しています。

#include <iostream>

#include <thrust/tuple.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/transform.h>
#include <thrust/device_vector.h>
#include <thrust/inner_product.h>
#include <thrust/functional.h>

#define N 2

struct zdiffsq{
template <typename Tuple>
  __host__ __device__ float operator()(Tuple a)
  {
    float result = thrust::get<0>(a) - thrust::get<1>(a);
    return result*result;
  }
};

struct diffsq{
  __host__ __device__ float operator()(float a, float b)
  {
    return (b-a)*(b-a);
  }
};

int main(){

  thrust::device_vector<float> oldvector(N);
  thrust::device_vector<float> newvector(N);
  oldvector[0] = 1.0f;  oldvector[1] = 2.0f;
  newvector[0] = 2.0f;  newvector[1] = 5.0f;

  float result = thrust::inner_product(oldvector.begin(), oldvector.end(), newvector.begin(), 0.0f, thrust::plus<float>(), diffsq());
  std::cout << "Result: " << result << std::endl;

  float result2 = thrust::transform_reduce(thrust::make_zip_iterator(thrust::make_tuple(oldvector.begin(), newvector.begin())), thrust::make_zip_iterator(thrust::make_tuple(oldvector.end(), newvector.end())), zdiffsq(), 0.0f, thrust::plus<float>());
  std::cout << "Result2: " << result2 << std::endl;
}

スラストプレースホルダーを使用して、内積の例で使用されているファンクター定義を削除することも調査できます。

独自の CUDA コードを作成する場合でも、並列リダクションや並べ替えなどの頻繁に使用されるアルゴリズムに対する現在の標準的な推奨事項は、cubを使用することです。

はい、CUDA 並列リダクションのサンプルとそれに付随するプレゼンテーションは、高速な並列リダクションの優れた基本的な紹介です。

于 2014-05-11T23:17:48.043 に答える