CUDA - 如何让内核中的线程等待它的孩子答案

【问题标题】：CUDA - How to make thread in kernel wait for it's childrenCUDA - 如何让内核中的线程等待它的孩子
【发布时间】：2014-12-01 21:43:55
【问题描述】：

我正在尝试使用 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];
    }
}

我应该对我的代码进行哪些调整才能达到预期的效果？

感谢阅读。

【问题讨论】：

标签： sorting parallel-processing cuda dynamic-parallelism

【解决方案1】：

用于强制内核完成的典型障碍是cudaDeviceSynchronize()，它也适用于父内核，强制子内核完成。

如the documentation所示：

由于设备运行时不支持 cudaStreamSynchronize() 和 cudaStreamQuery()，当应用程序需要知道流启动的子内核已完成时，应使用 cudaDeviceSynchronize()。

【讨论】：