java动态同步,CUDA动态并行和全局内存同步

__syncthreads() 执行的唯一操作是您在CUDA C编程指南中描述的引用 . 除了在多个内核启动中划分内核执行的天真方法之外，CUDA中没有办法跨块进行同步，并且在性能方面存在所有缺点 . 因此，你自己猜到的第一个问题的答案是否定的 .

在帖子的第二部分中，您指的是CUDA C编程指南的一个特定示例，即

__global__ void child_launch(int *data) {

data[threadIdx.x] = data[threadIdx.x]+1;

}

__global__ void parent_launch(int *data) {

data[threadIdx.x] = threadIdx.x;

__syncthreads();

if (threadIdx.x == 0) {

child_launch<<< 1, 256 >>>(data);

cudaDeviceSynchronize();

}

__syncthreads();

}

void host_launch(int *data) {

parent_launch<<< 1, 256 >>>(data);

}

这里， parent_launch 内核的所有 256 个线程在 data 中写入了一些东西 . 之后，线程 0 调用 child_launch . 需要第一个 __syncthreads() 来确保在子内核调用之前完成所有内存写入 . 在这一点上引用指南：

由于第一次__syncthreads()调用，子进程将看到数据[0] = 0，数据[1] = 1，...，数据[255] = 255(没有__syncthreads()调用，只有数据[0 ]将保证被孩子看到) .

关于第二个 __syncthreads() ，指南解释了这一点

当子网格返回时，线程0保证看到其子网格中的线程所做的修改 . 只有在第二次__syncthreads()调用之后，这些修改才可用于父网格的其他线程 .

在该特定示例中，第二个 __syncthreads() 是冗余的，因为由于内核终止而存在隐式同步，但是在子内核启动之后必须执行其他操作时需要第二个 __syncthreads() .

最后，关于你在帖子中引用的句子：

仅在第二次__syncthreads()调用之后，这些修改才可用于父网格的其他线程

请注意，在具体示例中， host_launch 函数只启动了一个线程块 . 这可能有点误导了你 .

在NVIDIA论坛上有一个有趣的讨论(可能甚至不止一个)，涉及跨越块的线程同步