CUDA - 如何使線程在內核中等待它的孩子

Question

我想實現一個非常簡單的合併排序使用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]; 
    } 
} 

我應該對我的代碼進行哪些改編，以便達到預期效果？

感謝您的閱讀。

Answer 1

用於強制內核完成的典型屏障是cudaDeviceSynchronize()，它也在父內核中工作，強制子內核完成。

如the documentation表示：

由於cudaStreamSynchronize（）和的cudaStreamQuery（）是不支持的設備運行時間，cudaDeviceSynchronize（）應改爲當應用程序需要知道流啓動子的內核已經使用完成。