0

ここに示されているように、ストラム圧縮操作を実行したい float の配列があります: Parallel Prefix Sum (Scan) with CUDA、および値とアドレスまたは元の要素に基づいて変換を適用します。

たとえば、値が {10,-1, -10, 2} の配列があり、絶対値が 5 より大きいすべての要素を返し、値とそのアドレスを取得する変換を適用したいとします。配列。ここでの結果は {transform(10,0),transform(-10,2)} になります。

これで推力を使用しようとしていますが、このコードは大きな配列で頻繁に実行されるため、理想的には、バッファーと配列の複数のトラバーサルを使用しません。

セカンダリ配列を割り当てたり、複数のトラバーサルを実行したりせずに、やりたいことを実行することは可能ですか? はいの場合、そのようなコードは実際に存在しますか? または、少なくとも、私の目標を達成するために作成できる推力またはその他のライブラリの機能へのポインタを誰かが持っていますか?

4

1 に答える 1

4

はい、単一の推力アルゴリズム呼び出しで推力で可能です（「複数のトラバーサルを実行せずに」という意味だと思います）および「セカンダリ配列を割り当てる」ことなく。

1つのアプローチは、データ配列とインデックス/「アドレス」配列を（thrust::counting_iterator割り当てを回避する を介して） thrust::transform_iterator「変換」操作（適切なファンクターと組み合わせて）を作成する に渡すことです。

次に、上記の変換反復子を適切なスラスト ストリーム圧縮アルゴリズムに渡して、必要な値を選択します。

可能なアプローチは次のとおりです。

$ cat t1044.cu
#include <thrust/device_vector.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/copy.h>
#include <math.h>

#include <iostream>

__host__ __device__ int my_transform(int data, int idx){
  return (data - idx);  //put whatever transform you want here
}

struct my_transform_func : public thrust::unary_function<thrust::tuple<int, int>, int>
{

  __host__ __device__
  int operator()(thrust::tuple<int, int> &t){
    return my_transform(thrust::get<0>(t), thrust::get<1>(t));
    }
};

struct my_test_func
{
  __host__ __device__
  bool operator()(int data){
    return (abs(data) > 5);
    }
};



int main(){

  int data[] = {10,-1,-10,2};
  int dsize = sizeof(data)/sizeof(int);

  thrust::device_vector<int> d_data(data, data+dsize);
  thrust::device_vector<int> d_result(dsize);
  int rsize = thrust::copy_if(thrust::make_transform_iterator(thrust::make_zip_iterator(thrust::make_tuple(d_data.begin(), thrust::counting_iterator<int>(0))), my_transform_func()), thrust::make_transform_iterator(thrust::make_zip_iterator(thrust::make_tuple(d_data.end(), thrust::counting_iterator<int>(dsize))), my_transform_func()), d_data.begin(), d_result.begin(), my_test_func()) - d_result.begin();
  thrust::copy_n(d_result.begin(), rsize, std::ostream_iterator<int>(std::cout, ","));
  std::cout << std::endl;
  return 0;
}
$ nvcc -o t1044 t1044.cu
$ ./t1044
10,-12,
$

このアプローチに対するいくつかの考えられる批判:

  1. 要素が 2 回読み込まれているように見えますd_data(変換操作で 1 回、ステンシルで 1 回)。ただし、CUDA 最適化コンパイラが、最終的に生成されるスレッド コードの冗長な負荷を認識し、最適化する可能性があります。

  2. 結果に保存するかどうかに関係なく、すべてのデータ要素に対して変換操作を実行しているように見えます。繰り返しますが、推力のcopy_if実装は、ステンシルの決定が行われるまで、データの読み込み操作を実際に延期する可能性があります。その場合、変換は必要に応じてのみ行われる可能性があります。常に行われる場合でも、多くのスラスト操作は計算バウンドではなく、ロード/ストアまたはメモリ帯域幅バウンドになる傾向があるため、これは取るに足らない問題である可能性があります。ただし、興味深い代替アプローチは、ここで@ms によって作成された適応を使用して、出力反復子ステップに適用される変換を作成することです。おそらく、結果に実際に保存されているデータ要素に対してのみ変換操作が実行されるように制限されますが、それも詳しく調べていません。

  3. 以下のコメントで述べたように、このアプローチは一時的なストレージを割り当てます (スラストはcopy_if操作の一部としてフードの下で行います)。もちろん、結果に O(n) ストレージを明示的に割り当てています。推力の割り当て（単一のcudaMalloc) もおそらく O(n) ストレージ用です。要求されたすべてのこと (並列プレフィックスの合計、ストリーム圧縮、データ変換) を、いかなる種類の追加ストレージもまったく使用せずに実行できる可能性があります (したがって、おそらく要求はインプレース操作に対するものです)。そのようなアルゴリズムは、実行可能である場合、パフォーマンスに重大な悪影響を与える可能性があります（ストリーム圧縮、つまりデータ平行移動）。スラストは使用するすべての一時ストレージを解放するため、このメソッドを頻繁に使用することに関連するストレージの問題はほとんどありません。残っている唯一の懸念 (私が推測する) は、パフォーマンスです。性能が気になるなら、カスタム アロケータ(こちらも参照) を推力すると、必要な最大ストレージ バッファが一度割り当てられ、上記のアルゴリズムが使用されるたびにそのバッファが再利用されます。

于 2016-01-14T15:19:14.803 に答える