同步在CUDA

Question

我读的CUDA有关同步CUDA参考手册，但我不知道很清楚。例如为什么我们使用cudaDeviceSynchronize()或__syncthreads()？如果不使用它们会发生什么情况，程序无法正常工作？什么cudaMemcpy和行动cudaMemcpyAsync区别？
你能展示一个显示这种差异的例子吗？

Answer 1

cudaDeviceSychronize()在主机代码使用（即在CPU上运行）当期望的是CPU活动等待任何未决GPU活动的完成。在许多情况下，这是没有必要明确地做到这一点，作为颁发给单流GPU操作是自动序列，像cudaMemcpy()某些其他操作都内置了固有的闭锁装置同步。但出于其他目的，例如调试代码，强制设备完成任何未完成的活动可能会很方便。

__syncthreads()用于设备代码（即在GPU上运行），并且在具有独立并行操作的代码中（例如逐个添加两个向量）可能根本不需要。但是，常用的一个例子是运行在共享内存之外的算法。在这种情况下，它是经常需要从全局内存值加载到共享内存中，我们希望在threadblock每个线程有机会加载它是适当的共享存储位置（S），发生任何实际的处理之前。在这种情况下，我们要使用__syncthreads()发生的处理之前，确保共享内存完全填充。这只是一个例子。 __syncthreads()可能用于任何需要线程块内的任何时间同步。它不允许块之间的同步。

cudaMemcpy和cudaMemcpyAsync之间的区别在于该调用的非异步版本只能发送到流0，并会阻止调用CPU线程，直到复制完成。异步版本可以选择使用流参数，并在复制完成之前立即将控制权返回给调用线程。异步版本通常在我们想要拥有asynchronous concurrent execution的情况下找到用法。

如果您有关于CUDA编程的基本问题，建议您采用一些webinars available。

Answer 2

此外，__syncthreads()当你在你的代码中的一些有条件的路径，然后你想运行取决于几个数组元素的操作变得很必要的。 考虑下面的例子：

int n = threadIdx.x; 

if(myarray[n] > 0) 
{ 
    myarray[n] = - myarray[n]; 
} 
double y = myarray[n] + myarray[n+1]; // Not all threads reaches here at the same time 

在上述例子中，不是所有的线程将具有相同的执行顺序。根据if条件，某些线程需要更长的时间。在考虑示例的最后一行时，您需要确保所有线程都完全正确完成了if条件并更新了myarray。如果不是这种情况，y可能会使用一些更新和未更新的值。 在这种情况下，它成为一个必须要评估y来克服这个问题前添加__syncthreads()：

if(myarray[n] > 0) 
{ 
    myarray[n] = - myarray[n]; 
} 
__syncthreads(); // All threads will wait till they come to this point 
// We are now quite confident that all array values are updated. 
double y = myarray[n] + myarray[n+1];