并行编程实战——CUDA编程的同步

原创于 2025-09-21 09:31:35 发布 · 771 阅读

9 ·

CC 4.0 BY-SA版权

文章标签：

#CUDA #c++

C++ 同时被 2 个专栏收录

463 篇文章

订阅专栏

并行编程

48 篇文章

订阅专栏

部署运行你感兴趣的模型镜像

一、线程同步

说句不负责任的话，有线程，线程同步就必然会出现。毕竟那种完全异步的通信情况是非常少数的应用场景，直接堆叠硬件即可解决问题，何必搞什么这个框架那个框架的。当然，有人会提出硬件的发展速度远跟不上软件的发展脚步，充分利用硬件的核心也是一种重要的手段。跑偏了啊，扯回来。
在CPU上的同步开发者都比较熟悉，那么在CUDA框架中的同步如何进行呢？本文将对其进行一个整体的分析和说明。

二、CUDA中的同步

CUDA针对并行编程的同步手段，与CPU不同的是，它面对的情况比较复杂，大致可以分成以下几种情况：

主机和设备之间同步
主机和设备间在原来分析中已经明白，它们基本上是异步执行的，原因就是为了提高并行效率。但为了实现同步的目标，其实也是可以使用同步机制来控制的，在它们之间，CUDA提供了显式和隐式两种同步机制：
显式同步：即使用cudaDeviceSynchronize()这个函数，它直接阻塞主机线程，直到前面所有发送到设备上的任务执行完成
隐式同步：这种情况一般在CUDA提供的API中，在某些情况下会导致同步的发生，比如cudaMalloc、cudaMemset以及在数据从设备端拷贝主机端时，也会发生同步。另外，在一些设备内存访问时也可能出现同步现象，如对异常的查询时
设备内整体的同步
CUDA并没有提供设备内的同步的机制，但这种需求是天然存在的。CUDA提供了一种变通的方式来实现了这种同步机制，即将不同的同步需求的任务分解成多个核函数来进行操作即可。此时，CUDA会依次启动不同的核函数，从而达到了同步的目的（串行）。
块内的同步
其实从线程的角度来看，块的线程才算得上真正的同步。同样也只有块内的线程才可以同步。它是CUDA中最常用的安全同步机制，即使用__syncthreads()接口来实现。它类似于上层CPU线程中的barrier（没接触的可以看Java或C++新标准中的相关内容），即所有块的线程都执行到此处时，才能继续进行下一步的操作。
注意它的限制，首先，不允许线程分支的出现，否则可能引起死锁；其次，必须保障内存的一致性。
Warp线程束的同步
在前面的分析中可以知道，CUDA中Warp是基本的执行单元（有点线程在CPU中的味道）。在Warp中线程是设计同步的，但有的时候仍然需要进行显式的同步，这时可以使用__syncwarp()来实现，它比上面提到的__syncthreads()同步开销更小且粒度更小。
流的同步
流内操作是一个串行操作的，自然就是同步的。
流间操作是并行执行的，如果需要实现同步可以使用事件来实现，通过事件的API创建、操作并最终销毁。可以实现同步机制。
内存的同步
在CUDA的内存中存在着原子操作这种同步原语，其与上层应用开发的原子操作类似
协作组同步和网格同步
在CUDA9之后，提供了协作组（Cooperative Groups ）的同步原语。其实就是为了给开发者提供更细粒度的原语控制，毕竟开发者做为一个人，一般都是想要又要的存在。
协作组可以是一个块Block或簇也可以是一个网格Grid，它提供了barrier_arrive and barrier_wait以及Sync函数。而如果想跨网格同步则可使用grid.sync()进行处理。

需要注意的是，并行机制中，同步现象是并行机制的一个重大的堵点，它的出现往往意味着效率会飞速的下降。一般来说，在应用CUDA的同步时，要尽量避免同步的使用，特别是全局的同步，如果无法避免则尽量减小同步的粒度，通过利用事件和流来改进同步的效果。同时，要对CUDA可能产生隐式同步的接口API认真使用，防止出现意想不到的结果。

三、CUDA中的同步函数

在CUDA编程中提供了几个同步接口，主要有以下几个：

__syncthreads函数
它用来用于协调同一块的线程之间的通信。一直等待到线程块中的所有线程都到达此点，且这些线程在__syncthreads()之前进行的所有全局和共享内存访问对块中的全部线程都是可见的。比如块内的线程共享某个内存相同地址时（类似访问一个相同的指针）。在条件代码中可以应用syncthreads()，但必须保证整个线程块中的计算结果相同，否则可能产生死锁等副作用。
__syncthreads_count(int predicate)函数
它的功能与__syncthreads()相同，不过提供了一个额外的特性，即统计结果为predicate的线程执行结果的数量并返回。这有一点条件过滤的意思。
__syncthreads_and(int predicate)函数
这个功能与上面的类似，不过它返回的是所有线程返回的结果与predicate保持一致时，才回返回非零值
__syncthreads_or(int predicate)函数
与and的正好相反，只要所有线程中有任何一个返回的结果与predicate一样，即返回非零值
__syncwarp(unsigned mask=0xffffffff)函数
它同样会导致执行线程等待，直到掩码中指定的所有Warp通道都执行了__syncwarp()（具有相同的掩码）后才会继续继续执行。每个调用线程必须在掩码中设置自己的位，掩码中命名的所有未退出线程必须使用相同的掩码执行相应的__syncwarp()。否则可能导致不确定的结果。
__syncwarp()保证了参与屏障的线程之间的内存顺序。因此，希望通过内存进行通信的Warp中的线程可以在存储数据到内存后调用__syncwarp()与其它线程安全的进行数据交互。

后面三个函数应用的环境是：具有2.x及更高计算能力的设备。

四、例程

看一下CUDA框架说明文档中相关的例程：

#include <cuda/barrier>
#include <cooperative_groups.h>

__device__ void compute(float* data, int curr_iteration);

__global__ void split_arrive_wait(int iteration_count, float *data) {
    using barrier = cuda::barrier<cuda::thread_scope_block>;
    __shared__  barrier bar;
    auto block = cooperative_groups::this_thread_block();

    if (block.thread_rank() == 0) {
        init(&bar, block.size()); // Initialize the barrier with expected arrival count
    }
    block.sync();

    for (int curr_iter = 0; curr_iter < iteration_count; ++curr_iter) {
        /* code before arrive */
       barrier::arrival_token token = bar.arrive(); /* this thread arrives. Arrival does not block a thread */
       compute(data, curr_iter);
       bar.wait(std::move(token)); /* wait for all threads participating in the barrier to complete bar.arrive()*/
        /* code after wait */
    }
}

上面的代码是一个通过五个阶段来实现同步的过程，即：

到达前的代码：执行将在等待之后被读取的内存更新。
带有隐式内存栅栏的到达点(即相当于 atomic_thread_fence(memory_order_seq_cst, thread_scope_block))。
到达与等待之间的代码。
等待点。
等待后的代码：能够看到在到达点之前执行的更新。
其中，同步点（block.sync()）被拆分为两步即到达点（bar.arrive()）和等待点（bar.wait(std::move(token)))。在线程调用arrive()时进入barrier，当线程调用wait时，其被阻塞直到参与的线程到达初始化中指定的计数。注意，它可以保证在wait前更新的内存为后续其它相关线程的可见性。arrive()不会阻塞线程。