選択Cudaで並べ替え

Question

Cudaで選択ソートを実装しようとしていますが、これまでのところ成功していません。

__device__ void selection_sort(int *data, int left, int right){ 

    for(int i = left ; i <= right ; ++i){ 
     int min_val = data[i]; 
     int min_idx = i; 

    // Find the smallest value in the range [left, right]. 
     for(int j = i+1 ; j <= right ; ++j){ 
      int val_j = data[j]; 
      if(val_j < min_val){ 
       min_idx = j; 
       min_val = val_j; 
      } 
     } 

     // Swap the values. 
     if(i != min_idx){ 
      data[min_idx] = data[i]; 
      data[i] = min_val; 
     } 
    } 
} 

ここに私の主な試みは、最小を見つけ、解決策を並列化することです。今、私はコードがC++のように見えることを理解していますが、私はCudaの熟練者ではありません。

ソリューションを並列化する方法はありますか？これ以上の追加はありませんか？ 還元と呼ばれる問題の広く知られており、十分に文書クラスに

for i from N-1 down to 0 
    find the maximum element among data[0] ~ data[i] 
    swap that maximum element with data[i] within the data array 

最初の部分（最大の要素を見つける）下がる：N番号の

Answer 1

選択ソートアルゴリズムは概ねとして記述することができます。ただし、2番目の部分（スワッピング）を実行するには、値を比較しながら最大要素のインデックスを追跡する必要があります。削減を実行する際にそれを行うことはあまり自然ではありません。これは、選択ソートがパラレル・アーキテクチャーにうまく機能しない理由の1つです。

また、ループごとに問題のサイズが1ずつ小さくなることがわかります。これは、選択ソートアルゴリズムのもう1つの側面であり、並列アーキテクチャには十分に対応していません。 CUDAの場合、32スレッドはワープを同時に実行します。任意の数のスレッドをワープ内で実行するように指示することはできますが、コンピューティング能力が失われているため、通常はそうすることはお勧めしません。

私はCUDAバージョンの選択ソートを自分で作成しようとしましたが、CUDAに適したアルゴリズムがより優れていると思われるため、選択のソートをやめました。しかし、これまで選択したソートがCUDAには適していない理由を説明するために、これまでに行ったことを紹介します。

まず、小さくて単純な問題から始めます：32個の要素をソートします。 32スレッドはワープを形成するため、最大値を見つけるにはshuffle instructionsを使用できます。 （Full code）

// Finds the maximum element within a warp and gives the maximum element to 
// thread with lane id 0. Note that other elements do not get lost but their 
// positions are shuffled. 
__inline__ __device__ int warpMax(int data, unsigned int threadId) 
{ 
    for (int mask = 16; mask > 0; mask /= 2) { 
     int dual_data = __shfl_xor(data, mask, 32); 
     if (threadId & mask) 
      data = min(data, dual_data); 
     else 
      data = max(data, dual_data); 
    } 
    return data; 
} 

__global__ void selection32(int* d_data, int* d_data_sorted) 
{ 
    unsigned int threadId = blockIdx.x * blockDim.x + threadIdx.x; 
    unsigned int laneId = threadIdx.x % 32; 

    int n = N; 
    while(n-- > 0) { 
     // get the maximum element among d_data and put it in d_data_sorted[n] 
     int data = d_data[threadId]; 
     data = warpMax(data, threadId); 
     d_data[threadId] = data; 

     // now maximum element is in d_data[0] 
     if (laneId == 0) { 
      d_data_sorted[n] = d_data[0]; 
      d_data[0] = INT_MIN; // this element is ignored from now on 
     } 
    } 
} 

int main() 
{ 
    // ... build data and trasfer to d_data ... 
    selection32<<<1, 32>>>(d_data, d_data_sorted); 
    // ... get the sorted array stored at d_data_sorted ... 
} 

はソートされていない領域の配列要素がシャッフルを保つ（一部は、これは1が正確選択ソートではないと主張してもよい）、および2）それはインプレースの並べ替えではありません。私はちょうど選択ソートがCUDAに適合していないことを示していることに注意してください。また、warpMaxには分岐が多いため、CUDAの最適化には向いていません。

要素の反りが1つしかない場合は平行に見えますが、問題のサイズが複数の反りになると悪化します。 1024要素の場合を見てみましょう。 （私はブロック内のスレッドの最大数制限であるため1024を選択しました）今度は32個のワープがあり、各ワープに対してwarpMaxを呼び出した後、最大の要素を得るために各ワープの最大要素を比較しなければなりません1024要素の中でデータ配列の最後の要素と最大値を入れ替えるために最大値が来たワープを追跡する必要があるため、32ワープ・最大値を比較するこの問題はwarpMaxでは実行できません。私がこれを行うために考えることの1つの方法は、ワープ - 最大値を比較するために1つのスレッドを使用することです。これは、ブロック内の他の1023スレッドがアイドル状態になるため、CUDAの実装には適していません。

さらに、問題のサイズがブロックでカバーできるよりも大きくなる場合は、各ブロックの最大値を比較する必要があります。ブロック間で同期をとる必要があるため、別々のカーネルを起動する必要があります。そして、最大値がどのブロックに由来するかを追跡する必要があると言うのは冗長です。これらすべては、CUDAの選択ソートの実装は良い考えではないことを示しています。