3

デバイスにソートされた整数配列があります。例： 

[0,0,0,1,1,2,2]

そして、別の配列の各要素へのオフセットが必要です。

[0,3,5]

（最初の0は位置0にあり、最初の1は位置3にあるなど）、事前にいくつの異なる要素があるかを知っています。これをCUDAでどのように効率的に実装しますか？私はコードを求めていませんが、この変換を計算するために実装するアルゴリズムの高レベルの説明を求めています。推力名前空間のさまざまな関数を見るのはすでに嫌いですが、これを実現するための推力関数の組み合わせは考えられませんでした。また、この変換には広く受け入れられている名前がありますか？

4

4 に答える 4

4

私はスラストライブラリを使用したことはありませんが、この可能なアプローチ（単純ですがおそらく効果的）はどうですか？

int input[N];  // your sorted array
int offset[N]; // the offset of the first values of each elements. Initialized with -1

// each thread will check an index position
if (input[id] > input[id-1]) // bingo! here begins a new value
{
    int oid = input[id];  // use the integer value as index
    offset[oid] = id;     // mark the offset with the beginning of the new value
}

あなたの例では、出力は次のようになります。

[0,3,5]

ただし、入力配列が次の場合：

[0,0,0,2,2,4,4]

次に、出力は次のようになります。

[0,-1, 3, -1, 5]

ここで、thrustがそれを実行できる場合は、remove_if（offset [i] == -1）を実行して、配列を圧縮します。

このアプローチでは、オフセット配列に多くのメモリが浪費されますが、検出するオフセットの数がわからないため、最悪の場合、入力配列と同じ量のメモリが使用されます。

一方、グローバルメモリ負荷と比較してスレッドあたりの命令数が少ないと、メモリ帯域幅によってこの実装が制限されます。スレッドごとにいくつかの値を処理するため、この場合にはいくつかの最適化があります。

私の2セント！

于 2011-11-15T13:05:25.870 に答える
4

thrust::unique_by_key_copyを使用してスラストでこれを解決できますthrust::counting_iterator。アイデアは、整数配列をkeysへの引数として扱いunique_by_key_copy、昇順の整数のシーケンス（つまり、counting_iterator）をとして使用することvaluesです。 unique_by_key_copy値の配列をそれぞれの一意のインデックスに圧縮しますkey：

#include <thrust/device_vector.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/discard_iterator.h>
#include <thrust/unique.h>
#include <thrust/copy.h>
#include <iterator>
#include <iostream>

int main()
{
  thrust::device_vector<int> keys(7);
  keys[0] = 0; keys[1] = 0; keys[2] = 0;
  keys[3] = 1; keys[4] = 1; keys[5] = 2; keys[6] = 2;

  std::cout << "keys before unique_by_key_copy: [ ";
  thrust::copy(keys.begin(), keys.end(), std::ostream_iterator<int>(std::cout," "));
  std::cout << "]" << std::endl;

  thrust::device_vector<int> offsets(3);

  thrust::unique_by_key_copy(keys.begin(), keys.end(),          // keys
                             thrust::make_counting_iterator(0), // [0, 1, 2, 3, ...] are the values
                             thrust::make_discard_iterator(),   // discard the compacted keys
                             offsets.begin());                  // the offsets are the values

  std::cout << "offsets after unique_by_key_copy: [ ";
  thrust::copy(offsets.begin(), offsets.end(), std::ostream_iterator<int>(std::cout," "));
  std::cout << "]" << std::endl;

  return 0;
}

出力は次のとおりです。

$ nvcc test.cu -run
keys before unique_by_key_copy: [ 0 0 0 1 1 2 2 ]
offsets after unique_by_key_copy: [ 0 3 5 ]
于 2011-11-15T20:06:44.740 に答える
1

スキャンはあなたが探しているアルゴリズムです。実装がない場合は、Thrustライブラリが優れたリソースになります。（推力を探す::スキャン）

スキャン（または「並列プレフィックス合計」）は、入力配列を受け取り、各要素がそのポイントへの入力の合計である出力を生成します。[1 5 3 7] => [1 6 9 16]

特定の要素が前の要素と同じであるかどうかを述語がチェックする述語（評価された条件に応じて0または1）をスキャンする場合、問題の要素の出力インデックスを計算します。サンプル配列

[0 0 0 1 1 2 2] [0 0 0 1 0 10]<=述語[000 1 1 22]<=スキャンされた述語

これで、スキャンした述語をインデックスとして使用して、出力を書き込むことができます。

于 2011-11-15T13:06:29.080 に答える
0

良い質問と答えは、後でそれをどうする必要があるかによって異なります。説明させてください。

この問題がCPUのO（n）（nは入力長）で解決できるとすぐに、メモリの割り当てとコピー（ホスト->デバイス（入力）およびデバイス->ホスト（結果））の欠点が発生します。 。これにより、単純なCPUソリューションに対するパフォーマンスが低下します。

配列がすでにデバイスメモリにある場合でも、各計算ブロックはそれをローカルまたはレジスタに読み取る必要があり（少なくともデバイスメモリにアクセスする）、CPUよりも大幅に高速に実行することはできません。

一般に、CUDAは次の場合にパフォーマンスを向上させます。

  1. 計算の漸近的な複雑さは、入力データの長さに比べて高くなります。たとえば、入力データの長さはnで、複雑さはO（n ^ 2）またはO（n ^ 3）です。

  2. タスクを独立したサブタスクまたは弱い依存のサブタスクに分割する方法があります。

ですから、もし私があなたなら、可能であれば、CUDAでそのような種類の計算をしようとはしません。また、スタンドアロン関数または他の関数の出力形式変換である必要がある場合は、CPUで実行します。

それがより複雑なアルゴリズムの一部である場合、答えはより複雑です。もし私があなたのところにいたら[0,3,5]、CUDAの計算能力の利用に制限を加えるので、どういうわけかフォーマットを変更しようとします。独立したブロックでタスクを効果的に分割することはできません。たとえば、ある計算スレッドで10個の整数を処理し、別のスレッドで次の10個の整数を処理する場合です。2番目の人は、最初の人が終了しないまで、出力をどこに配置するかを知りません。配列をサブ配列に分割し、各サブ配列の回答を個別に保存する可能性があります。これは、実行している計算に大きく依存します。

于 2011-11-15T11:54:42.887 に答える