CUDAを使用したMergesort

Question

ボトムアップ/反復マージソートアルゴリズムに基づいて独自のMergesortを実装しようとしました。このアルゴリズムは、データを2つの要素で分割し、並べ替えます。次に、すべてのデータがソートされるまで4要素とソートされます。だから、私の計画は2つの要素で各スレッドを割り当てています。

__global__ void mergeBU(int *d_a, int *d_aux, int sz, int N) 
{ 
    int idk = blockIdx.x*blockDim.x+threadIdx.x; 
    int lo = 2 * sz * idk; 
    int mid = lo + sz - 1; 
    float hi = fminf(lo + sz + sz - 1, N - 1); 
    merge(d_a, d_aux, lo, mid, hi); 
} 

__device__ void merge(int *d_a, int *d_aux, int lo, int mid, float hi) 
{ 
int i = lo; 
int j = mid + 1; 

    for (int k = lo; k <= hi; k++) 
    { 
     d_aux[k] = d_a[k]; 
    } 

    for (int k = lo; k <= hi; k++) 
    { 
     if (i > mid)     { d_a[k] = d_aux[j]; j++; } 
     else if (j > hi)    { d_a[k] = d_aux[i]; i++; } 
     else if (d_aux[j] < d_aux[i]) { d_a[k] = d_aux[j]; j++; } 
     else        { d_a[k] = d_aux[i]; i++; } 
    } 
} 

のは、私は（8つのスレッドです）私のカーネル< < < 2,4 >>>を起動しましょう、私ができる唯一の並べ替え16の要素最大：だから私はこれを行います。 32個の要素を入力すると、残りのデータインデックスは変更されません（16-31）。スレッドインデックスを作成する方法は、残りのデータインデックスを処理し続けますか？私はスレッドインデックス（0,1,2,3,4,5,6,7）が残りのデータインデックスを処理し続けていることを意味し、threadindex（dataindex、dataindex） - > 0（16、 17）。 1（18,19）; 2（20,21）;等々。コメントは大歓迎です。

Answer 1

実際のコードを見ずに：マージソートはマルチパスアルゴリズムです。 デバイス全体のアトミックを使用していない限り、カーネルを実行するときには通常、異なるブロックはまったく同期しないため、複数の連続したカーネルの起動を考慮する必要があります。たとえば、最初の起動では、スレッドの各ブロックはn_1個の要素をソートします。 2回目の起動では、ブロックの各ペアが2 * n_1個の要素をマージします。もちろん、それはそれほど簡単ではありません：どのブロックが正確に何をすべきかをどのように伝えることができますか？

その他のアイデアについてはapproach used in the ModernGPU libraryをご覧ください。

Answer 2

あなたのアプローチは、サイズnのアレイをn/2サブアレイに分割し、それらのサブアレイのペアをマージしてn/4サブアレイで終わるように見えることです。しかし、このアプローチはおそらくメモリ帯域幅に制限があります。

8つのスレッドを使用するとします。 8つのスレッドを使用してサブアレイをソートし、次に4つのスレッドを使用してソートされたサブの4つのペアをマージします（最後のサブアレイはサイズが異なる可能性があります） 2つのマージされたサブアレイをマージする2つのスレッド、最後の2つのマージをマージする1つのスレッド。

マルチスレッドソートでの経験に基づいて、CPUの8スレッドでメモリ帯域幅の制限に達しますが、gpuのメモリを使用してアレイの大きな部分を保持できる場合、8スレッド以上有益である。私は、どのような操作（比較、移動、...）は、GPUとそのメモリ内で可能なのか分からない。