2

CUDA 再帰 (cm > 35 の場合) テクノロジを使用して非常に単純なマージソートを実装しようとしていますが、親スレッドに子を同時に起動し、子の計算を待つように指示する方法が見つかりません。 cudaEventSynchronize( ) と cudaStreamSynchronize() はホストのみです。親の次の行は、その子がすべての計算を完了した後にのみ実行する必要があるため、__syncthread() は目的の効果をアーカイブしません。

__global__ void simple_mergesort(int* data,int *dataAux,int begin,int end, int depth){
     int middle = (end+begin)/2;
     int i0 = begin;
     int i1 = middle;
     int index;
     int n = end-begin;

     cudaStream_t s,s1;

     //If we're too deep or there are few elements left, we use an insertion sort...
     if( depth >= MAX_DEPTH || end-begin <= INSERTION_SORT ){
         selection_sort( data, begin, end );
         return;
     }

     if(n < 2){
         return;
     }

    // Launches a new block to sort the left part.
    cudaStreamCreateWithFlags(&s,cudaDeviceScheduleBlockingSync);
    simple_mergesort<<< 1, 1, 0, s >>>(data,dataAux, begin, middle, depth+1);
    cudaStreamDestroy(s);

    // Launches a new block to sort the right part.
    cudaStreamCreateWithFlags(&s1,cudaDeviceScheduleBlockingSync);
    simple_mergesort<<< 1, 1, 0, s1 >>>(data,dataAux, middle, end, depth+1);
    cudaStreamDestroy(s1);

    // Waits until children have returned, does not compile.
    cudaStreamSynchronize(s);
    cudaStreamSynchronize(s1);


    for (index = begin; index < end; index++) {
        if (i0 < middle && (i1 >= end || data[i0] <= data[i1])){
            dataAux[index] = data[i0];
            i0++;
        }else{
            dataAux[index] = data[i1];
            i1++;
        }
    }

    for(index = begin; index < end; index ++){
        data[index] = dataAux[index];
    }
}

目的の効果を得るには、コードをどのように変更すればよいですか?

読んでくれてありがとう。

4

1 に答える 1

4

カーネルを強制的に完了させるために使用される典型的なバリアはcudaDeviceSynchronize()、親カーネルでも機能し、子カーネルを強制的に完了させます。

ドキュメントに示されているように：

cudaStreamSynchronize() と cudaStreamQuery() はデバイス ランタイムでサポートされていないため、ストリーム起動された子カーネルが完了したことをアプリケーションが認識する必要がある場合は、代わりに cudaDeviceSynchronize() を使用する必要があります。

于 2014-12-01T23:18:59.943 に答える