计算机视觉 基于CUDA编程的入门与实践 线程及同步三

一、线程同步

        对于需要线程之间互相交换数据才能完成任务的场景，必须存在某种能让线程彼此交流的机制。就需要共享内存，当很多线程并行工作并且访问相同的数据或者存储器位置的时候，线程间必须正确的同步。

        不过，线程间交换数据并不一定需要使用共享内存，只是共享内存较快而已。使用全局内存同样可以。例如配合正确的同步操作或者原子操作（原子操作也支持全局内存），依然可以正确地完成任务。只是使用共享内存，很多情况下较快（延迟较低，带宽较大）而已。

        为了确保并行线程协作时的正确结果，我们必须同步线程。CUDA 提供了一个简单的屏障同步原语__syncthreads()。线程的执行只有在其块中的所有线程都执行了。因此，我们可以通过在存储之后和从共享内存加载任何线程之前调用上述争用条件来避免上述争用条件。

二、共享内存

        共享内存位于芯片内部，因此它比全局内存快得多。（CUDA里面存储器的快慢有两方面，一个是延迟低，一个是带宽大。这里特指延迟低），相比没有经过缓存的全局内存访问，共享内存大约在延迟上低100倍。同一个块中的线程可以访问相同的一段共享内存（注意：不同块中的线程所见到的共享内存中的内容是不相同的），这在许多线程需要与其他线程共享它们的结果的应用程序中非常有用。但是如果不同步，也可能会造成混乱或错误的结果。如果某线程的计算结果在写入到共享内存完成之前被其他线程读取，那么将会导致错误。因此，应该正确地控制或管理内存访问。这是由__syncthreads()指令完成的，该指令确保在继续执行程序之前完成对内存的所有写入操作。这也被称为barrier。barrier的含义是块中的所有线程都将到达该代码行，然后在此等待其他线