1

問題

2 つの配列があるとします。

   const int N = 1000000;
   float A[N];
   myStruct *B[N];

A の数値は正または負にすることができます (例: A[N]={3,2,-1,0,5,-2})。GPU で配列 A を部分的に並べ替えるにはどうすればよいですか (最初にすべての正の値、並べ替える必要がなく、次に負の値) (たとえばA[N]={3,2,5,0,-1,-2}またはA[N]={5,2,3,0,-2,-1})。配列 B は A に従って変更する必要があります (A はキー、B は値)。

A,B のスケールは非常に大きくなる可能性があるため、ソート アルゴリズムは GPU に実装する必要があると思います (特に CUDA では、このプラットフォームを使用しているため)。確かにthrust::sort_by_keyこの作業ができることはわかっていますが、配列 A&B を完全にソートする必要がないため、余分な作業が必要になります。

この種の問題に遭遇した人はいますか？

推力例

thrust::sort_by_key(thrust::device_ptr<float> (A), 
            thrust::device_ptr<float> ( A + N ),  
            thrust::device_ptr<myStruct> ( B ),  
            thrust::greater<float>() );
4

3 に答える 3

1

Github 上の Thrust のドキュメントは最新ではありません。@JaredHoberock が言ったように、 stencils をサポートするthrust::partitionようになったので、行くべき道です。Github リポジトリからコピーを取得する必要がある場合があります。

git クローン git://github.com/thrust/thrust.git

次にscons doc、Thrust フォルダーで実行して更新されたドキュメントを取得し、コードをコンパイルするときにこれらの更新された Thrust ソースを使用します ( nvcc -I/path/to/thrust ...)。新しいステンシル パーティションを使用すると、次のことができます。

#include <thrust/partition.h>
#include <thrust/execution_policy.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/tuple.h>

struct is_positive
{
__host__ __device__
bool operator()(const int &x)
{
  return x >= 0;
}
};


thrust::partition(thrust::host, // if you want to test on the host
                  thrust::make_zip_iterator(thrust::make_tuple(keyVec.begin(), valVec.begin())),
                  thrust::make_zip_iterator(thrust::make_tuple(keyVec.end(), valVec.end())),
                  keyVec.begin(),
                  is_positive());

これは以下を返します:

Before:
  keyVec =   0  -1   2  -3   4  -5   6  -7   8  -9
  valVec =   0   1   2   3   4   5   6   7   8   9
After:
  keyVec =   0   2   4   6   8  -5  -3  -7  -1  -9
  valVec =   0   2   4   6   8   5   3   7   1   9

2 つのパーティションは必ずしもソートされていないことに注意してください。また、元のベクトルとパーティションの間で順序が異なる場合があります。これが重要な場合は、次を使用できますthrust::stable_partition。

stable_partition は、stable_partition が相対順序を保持することが保証されているという点で partition とは異なります。つまり、x と y が pred(x) == pred(y) のような [first, last) の要素であり、x が y に先行する場合、stable_partition の後でも x が y に先行することは true です。

完全な例が必要な場合は、次のとおりです。

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/partition.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/tuple.h>

struct is_positive
{
__host__ __device__
bool operator()(const int &x)
{
  return x >= 0;
}
};

void print_vec(const thrust::host_vector<int>& v)
{
  for(size_t i = 0; i < v.size(); i++)
    std::cout << "  " << v[i];
  std::cout << "\n";
}

int main ()
{
  const int N = 10;

  thrust::host_vector<int> keyVec(N);
  thrust::host_vector<int> valVec(N);

  int sign = 1;
  for(int i = 0; i < N; ++i)
  {
    keyVec[i] = sign * i;
    valVec[i] = i;
    sign *= -1;
  }

  // Copy host to device
  thrust::device_vector<int> d_keyVec = keyVec;
  thrust::device_vector<int> d_valVec = valVec;

  std::cout << "Before:\n  keyVec = ";
  print_vec(keyVec);
  std::cout << "  valVec = ";
  print_vec(valVec);

  // Partition key-val on device
  thrust::partition(thrust::make_zip_iterator(thrust::make_tuple(d_keyVec.begin(), d_valVec.begin())),
                    thrust::make_zip_iterator(thrust::make_tuple(d_keyVec.end(), d_valVec.end())),
                    d_keyVec.begin(),
                    is_positive());

  // Copy result back to host
  keyVec = d_keyVec;
  valVec = d_valVec;

  std::cout << "After:\n  keyVec = ";
  print_vec(keyVec);
  std::cout << "  valVec = ";
  print_vec(valVec);
}

アップデート

thrust::sort_by_keyバージョンと簡単に比較したところ、thrust::partition実装の方が速いようです (これは当然のことです)。NVIDIA Visual Profiler で取得したものを次に示します。左N = 1024 * 1024がソート バージョン、右がパーティション バージョンです。同じ種類のテストを自分で実行したい場合があります。

ソートとパーティション

于 2013-05-17T05:45:04.933 に答える