4.10. 异步屏障

高级同步原语中介绍的异步屏障将 CUDA 的同步能力扩展到了__syncthreads()和__syncwarp()之外，能够实现细粒度的非阻塞协调，并使通信与计算更好地重叠。

本节详细介绍如何主要通过cuda::barrier API 使用异步屏障（在适用之处提供指向cuda::ptx及原语的参考）。

4.10.1. 初始化

初始化必须在任何线程开始参与屏障之前完成。

CUDA C++ cuda::barrier

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

__global__ void init_barrier()
{
  __shared__ cuda::barrier<cuda::thread_scope_block> bar;
  auto block = cooperative_groups::this_thread_block();

  if (block.thread_rank() == 0)
  {
    // A single thread initializes the total expected arrival count.
    init(&bar, block.size());
  }
  block.sync();
}

CUDA C++ cuda::ptx

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

__global__ void init_barrier()
{
  __shared__ uint64_t bar;
  auto block = cooperative_groups::this_thread_block();

  if (block.thread_rank() == 0)
  {
    // A single thread initializes the total expected arrival count.
    cuda::ptx::mbarrier_init(&bar, block.size());
  }
  block.sync();
}

CUDA C 原语

#include <cuda_awbarrier_primitives.h>
#include <cooperative_groups.h>

__global__ void init_barrier()
{
  __shared__ uint64_t bar;
  auto block = cooperative_groups::this_thread_block();

  if (block.thread_rank() == 0)
  {
    // A single thread initializes the total expected arrival count.
    __mbarrier_init(&bar, block.size());
  }
  block.sync();
}

在任何线程可以参与屏障之前，必须使用友元函数cuda::barrier::init()初始化该屏障。这必须在线程到达屏障之前完成。这带来了一个自举难题：线程必须在参与屏障之前同步，但线程创建屏障正是为了同步。在此示例中，将参与的线程属于一个协作组，并使用block.sync()来引导初始化。由于整个线程块都参与该屏障，因此也可以使用__syncthreads()。

init()的第二个参数是预期到达计数，即在一个参与线程从其对bar.wait(std::move(token))的调用中解除阻塞之前，参与线程将调用bar.arrive()的次数。在本示例和前面的示例中，屏障使用线程块中的线程数（即cooperative_groups::this_thread_block().size()）进行初始化，以便线程块内的所有线程都能参与屏障。

异步屏障可灵活指定线程如何参与（拆分到达与等待）以及哪些线程参与。相比之下，this_thread_block.sync()或__syncthreads()适用于整个线程块，而__syncwarp(mask)适用于线程束的指定子集。不过，如果用户的意图是同步整个线程块或整个线程束，我们建议分别使用__syncthreads()和__syncwarp()以获得更好的性能。

4.10.2. 屏障的阶段：到达、递减计数、完成与重置

随着参与线程调用bar.arrive()，异步屏障从预期到达计数递减至零。当计数减至零时，屏障的当前阶段就已完成。当最后一次bar.arrive()调用使计数减至零时，该递减计数会以原子方式自动重置。重置会将递减计数设为预期到达计数，并使屏障进入下一阶段。

由token=bar.arrive()返回的cuda::barrier::arrival_token类的token对象与屏障的当前阶段相关联。当屏障处于当前阶段时，即当令牌关联的阶段与屏障的阶段相匹配时，调用bar.wait(std::move(token))会阻塞调用线程。如果阶段在调用bar.wait(std::move(token))之前推进了（由于计数减至零），那么线程不会阻塞；如果阶段在线程阻塞于bar.wait(std::move(token))期间推进了，那么该线程会解除阻塞。

必须清楚何时可能发生重置、何时不可能发生重置，尤其是在复杂的到达/等待同步模式中。

线程对token=bar.arrive()和bar.wait(std::move(token))的调用必须按顺序安排，使token=bar.arrive()发生在屏障的当前阶段，而bar.wait(std::move(token))发生在同一阶段或下一阶段。
线程对bar.arrive()的调用必须在屏障计数器非零时发生。屏障初始化后，如果线程对bar.arrive()的调用使计数减至零，那么在屏障能够被后续的bar.arrive()调用复用之前，必须先调用bar.wait(std::move(token))。
只能使用当前阶段或紧邻的前一阶段的token对象调用bar.wait()。对于token对象的任何其他值，其行为均未定义。

对于简单的到达/等待同步模式，遵守这些使用规则很容易。

4.10.2.1. 线程束纠缠

线程束分歧会影响到达操作更新屏障的次数。如果发起调用的线程束完全收敛，则屏障会更新一次。如果发起调用的线程束完全分歧，则会对屏障执行 32 次单独的更新。

注意

建议由收敛的线程调用arrive-on(bar)，以尽量减少对屏障对象的更新。当这些操作之前的代码导致线程分歧时，应在调用到达操作之前通过__syncwarp使线程束重新收敛。

4.10.3. 显式阶段跟踪

异步屏障可以有多个阶段，具体取决于它用于同步线程和内存操作的次数。我们可以使用通过cuda::ptx和原语 API 提供的mbarrier_try_wait_parity()系列函数来直接跟踪阶段，而不必使用令牌来跟踪屏障的阶段翻转。

在最简单的形式中，cuda::ptx::mbarrier_try_wait_parity(uint64_t* bar, const uint32_t& phaseParity)函数等待具有特定奇偶性的阶段。phaseParity操作数是屏障对象当前阶段或紧邻的前一阶段的整数奇偶值。偶数阶段的整数奇偶值为 0，奇数阶段的整数奇偶值为 1。初始化屏障时，其阶段的奇偶值为 0。因此，phaseParity的有效值为 0 和 1。跟踪异步内存操作时，显式阶段跟踪可能很有用，因为它允许仅由一个线程到达屏障并设置事务计数，而其他线程只需等待基于奇偶性的阶段翻转。这可能比让所有线程都到达屏障并使用令牌更高效。此功能仅适用于线程块和簇作用域的共享内存屏障。

CUDA C++ cuda::barrier

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

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

__global__ void split_arrive_wait(int iteration_count, float *data)
{
  using barrier_t = cuda::barrier<cuda::thread_scope_block>;
  __shared__ barrier_t bar;
  int parity = 0; // Initial phase parity is 0.
  auto block = cooperative_groups::this_thread_block();

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

  for (int i = 0; i < iteration_count; ++i)
  {
    /* code before arrive */

    // This thread arrives. Arrival does not block a thread.
    // Get a handle to the native barrier to use with cuda::ptx API.
    (void)cuda::ptx::mbarrier_arrive(cuda::device::barrier_native_handle(bar));

    compute(data, i);

    // Wait for all threads participating in the barrier to complete mbarrier_arrive().
    // Get a handle to the native barrier to use with cuda::ptx API.
    while (!cuda::ptx::mbarrier_try_wait_parity(cuda::device::barrier_native_handle(bar), parity)) {}
    // Flip parity.
    parity ^= 1;

    /* code after wait */
  }
}

CUDA C++ cuda::ptx

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

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

__global__ void split_arrive_wait(int iteration_count, float *data)
{
  __shared__ uint64_t bar;
  int parity = 0; // Initial phase parity is 0.
  auto block = cooperative_groups::this_thread_block();

  if (block.thread_rank() == 0)
  {
    // Initialize barrier with expected arrival count.
    cuda::ptx::mbarrier_init(&bar, block.size());
  }
  block.sync();

  for (int i = 0; i < iteration_count; ++i)
  {
    /* code before arrive */

    // This thread arrives. Arrival does not block a thread.
    (void)cuda::ptx::mbarrier_arrive(&bar);

    compute(data, i);

    // Wait for all threads participating in the barrier to complete mbarrier_arrive().
    while (!cuda::ptx::mbarrier_try_wait_parity(&bar, parity)) {}
    // Flip parity.
    parity ^= 1;

    /* code after wait */
  }
}

CUDA C 原语

#include <cuda_awbarrier_primitives.h>
#include <cooperative_groups.h>

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

__global__ void split_arrive_wait(int iteration_count, float *data)
{
  __shared__ __mbarrier_t bar;
  bool parity = false; // Initial phase parity is false.
  auto block = cooperative_groups::this_thread_block();

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

  for (int i = 0; i < iteration_count; ++i)
  {
    /* code before arrive */

    // This thread arrives. Arrival does not block a thread.
    (void)__mbarrier_arrive(&bar);

    compute(data, i);

    // Wait for all threads participating in the barrier to complete __mbarrier_arrive().
    while(!__mbarrier_try_wait_parity(&bar, parity, 1000)) {}
    parity ^= 1;

    /* code after wait */
  }
}

4.10.4. 提前退出

当参与一系列同步的线程必须提前退出该同步序列时，该线程必须在退出之前显式停止参与。其余参与线程可以正常继续后续的到达和等待操作。

CUDA C++ cuda::barrier

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

__device__ bool condition_check();

__global__ void early_exit_kernel(int N)
{
  __shared__ cuda::barrier<cuda::thread_scope_block> bar;
  auto block = cooperative_groups::this_thread_block();

  if (block.thread_rank() == 0)
  {
    init(&bar, block.size());
  }
  block.sync();

  for (int i = 0; i < N; ++i)
  {
    if (condition_check())
    {
      bar.arrive_and_drop();
      return;
    }
    // Other threads can proceed normally.
    auto token = bar.arrive();

    /* code between arrive and wait */

    // Wait for all threads to arrive.
    bar.wait(std::move(token));

    /* code after wait */
  }
}

CUDA C 原语

#include <cuda_awbarrier_primitives.h>
#include <cooperative_groups.h>

__device__ bool condition_check();

__global__ void early_exit_kernel(int N)
{
  __shared__ __mbarrier_t bar;
  auto block = cooperative_groups::this_thread_block();

  if (block.thread_rank() == 0)
  {
    __mbarrier_init(&bar, block.size());
  }
  block.sync();

  for (int i = 0; i < N; ++i)
  {
    if (condition_check())
    {
      __mbarrier_token_t token = __mbarrier_arrive_and_drop(&bar);
      return;
    }
    // Other threads can proceed normally.
    __mbarrier_token_t token = __mbarrier_arrive(&bar);

    /* code between arrive and wait */

    // Wait for all threads to arrive.
    while (!__mbarrier_try_wait(&bar, token, 1000)) {}

    /* code after wait */
  }
}

bar.arrive_and_drop()操作到达屏障，履行参与线程在当前阶段到达的义务，然后递减下一阶段的预期到达计数，从而不再预期此线程到达屏障。

4.10.5. 完成函数

cuda::barrier API 支持一个可选的完成函数。cuda::barrier<Scope, CompletionFunction>的CompletionFunction在每个阶段执行一次，时间是在最后一个线程到达之后、任何线程从wait解除阻塞之前。在该阶段到达barrier的线程所执行的内存操作，对执行CompletionFunction的线程可见；在CompletionFunction内执行的所有内存操作，则对所有在barrier处等待的线程可见，前提是这些线程已从wait解除阻塞。

CUDA C++ cuda::barrier

#include <cuda/barrier>
#include <cooperative_groups.h>
#include <functional>
namespace cg = cooperative_groups;

__device__ int divergent_compute(int *, int);
__device__ int independent_computation(int *, int);

__global__ void psum(int *data, int n, int *acc)
{
  auto block = cg::this_thread_block();

  constexpr int BlockSize = 128;
  __shared__ int smem[BlockSize];
  assert(BlockSize == block.size());
  assert(n % BlockSize == 0);

  auto completion_fn = [&]
  {
    int sum = 0;
    for (int i = 0; i < BlockSize; ++i)
    {
      sum += smem[i];
    }
    *acc += sum;
  };

  /* Barrier storage.
     Note: the barrier is not default-constructible because
           completion_fn is not default-constructible due
           to the capture. */
  using completion_fn_t = decltype(completion_fn);
  using barrier_t = cuda::barrier<cuda::thread_scope_block,
                                  completion_fn_t>;
  __shared__ std::aligned_storage<sizeof(barrier_t),
                                  alignof(barrier_t)>
      bar_storage;

  // Initialize barrier.
  barrier_t *bar = (barrier_t *)&bar_storage;
  if (block.thread_rank() == 0)
  {
    assert(*acc == 0);
    assert(blockDim.x == blockDim.y == blockDim.y == 1);
    new (bar) barrier_t{block.size(), completion_fn};
    /* equivalent to: init(bar, block.size(), completion_fn); */
  }
  block.sync();

  // Main loop.
  for (int i = 0; i < n; i += block.size())
  {
    smem[block.thread_rank()] = data[i] + *acc;
    auto token = bar->arrive();
    // We can do independent computation here.
    bar->wait(std::move(token));
    // Shared-memory is safe to re-use in the next iteration
    // since all threads are done with it, including the one
    // that did the reduction.
  }
}

4.10.6. 跟踪异步内存操作

异步屏障可用于跟踪异步内存复制。当异步复制操作绑定到屏障时，该复制操作会在启动时自动递增当前屏障阶段的预期计数，并在完成时递减该计数。此机制确保屏障的wait()操作会阻塞，直到所有关联的异步内存复制都已完成，从而提供一种同步多个并发内存操作的便捷方式。

从计算能力 9.0 开始，位于共享内存中、具有线程块或簇作用域的异步屏障可以显式跟踪异步内存操作。我们将这些屏障称为异步事务屏障。除了预期到达计数外，屏障对象还可以接受一个事务计数，该计数可用于跟踪异步事务的完成情况。事务计数跟踪尚未完成的异步事务数量，单位由异步内存操作指定（通常为字节）。当前阶段要跟踪的事务计数可以在到达时通过cuda::device::barrier_arrive_tx()设置，也可以直接通过cuda::device::barrier_expect_tx()设置。当屏障使用事务计数时，它会在等待操作处阻塞线程，直到所有生产者线程都执行了到达操作，并且所有事务计数之和达到预期值。

CUDA C++ cuda::barrier

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

__global__ void track_kernel()
{
  __shared__ cuda::barrier<cuda::thread_scope_block> bar;
  auto block = cooperative_groups::this_thread_block();

  if (block.thread_rank() == 0)
  {
    init(&bar, block.size());
  }
  block.sync();

  auto token = cuda::device::barrier_arrive_tx(bar, 1, 0);

  bar.wait(cuda::std::move(token));
}

CUDA C++ cuda::ptx

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

__global__ void track_kernel()
{
  __shared__ uint64_t bar;
  auto block = cooperative_groups::this_thread_block();

  if (block.thread_rank() == 0)
  {
    cuda::ptx::mbarrier_init(&bar, block.size());
  }
  block.sync();

  uint64_t token = cuda::ptx::mbarrier_arrive_expect_tx(cuda::ptx::sem_release, cuda::ptx::scope_cluster, cuda::ptx::space_shared, &bar, 1, 0);

  while (!cuda::ptx::mbarrier_try_wait(&bar, token)) {}
}

在此示例中，cuda::device::barrier_arrive_tx()操作构造一个与当前阶段的阶段同步点相关联的到达令牌对象。随后，将到达计数减 1，并将预期事务计数加 0。由于事务计数的更新量为 0，该屏障不跟踪任何事务。后续的使用张量内存加速器（TMA）一节包含跟踪异步内存操作的示例。

4.10.7. 使用屏障的生产者-消费者模式

可以对线程块进行空间划分，以允许不同的线程执行独立的操作。最常见的做法是将线程块内不同线程束中的线程分配给特定任务。这项技术称为线程束专门化。

本节展示了生产者-消费者模式中的空间划分示例，其中一个线程子集生成数据，另一个与其不相交的线程子集并发消费这些数据。生产者-消费者空间划分模式需要两次单边同步，以管理生产者和消费者之间的数据缓冲区。

生产者	消费者
等待缓冲区准备好被填充	发出缓冲区已准备好被填充的信号
生成数据并填充缓冲区
发出缓冲区已填满的信号	等待缓冲区被填满
	消费已填满缓冲区中的数据

生产者线程等待消费者线程发出缓冲区已准备好被填充的信号；但是，消费者线程不等待此信号。消费者线程等待生产者线程发出缓冲区已填满的信号；但是，生产者线程不等待此信号。为了实现生产者与消费者的完全并发，此模式采用（至少）双缓冲，其中每个缓冲区需要两个屏障。

CUDA C++ cuda::barrier

#include <cuda/barrier>

using barrier_t = cuda::barrier<cuda::thread_scope_block>;

__device__ void produce(barrier_t ready[], barrier_t filled[], float *buffer, int buffer_len, float *in, int N)
{
  for (int i = 0; i < N / buffer_len; ++i)
  {
    ready[i % 2].arrive_and_wait(); /* wait for buffer_(i%2) to be ready to be filled */
    /* produce, i.e., fill in, buffer_(i%2)  */
    barrier_t::arrival_token token = filled[i % 2].arrive(); /* buffer_(i%2) is filled */
  }
}

__device__ void consume(barrier_t ready[], barrier_t filled[], float *buffer, int buffer_len, float *out, int N)
{
  barrier_t::arrival_token token1 = ready[0].arrive(); /* buffer_0 is ready for initial fill */
  barrier_t::arrival_token token2 = ready[1].arrive(); /* buffer_1 is ready for initial fill */
  for (int i = 0; i < N / buffer_len; ++i)
  {
    filled[i % 2].arrive_and_wait(); /* wait for buffer_(i%2) to be filled */
    /* consume buffer_(i%2) */
    barrier_t::arrival_token token3 = ready[i % 2].arrive(); /* buffer_(i%2) is ready to be re-filled */
  }
}

__global__ void producer_consumer_pattern(int N, float *in, float *out, int buffer_len)
{
  constexpr int warpSize = 32;

  /* Shared memory buffer declared below is of size 2 * buffer_len
     so that we can alternatively work between two buffers.
     buffer_0 = buffer and buffer_1 = buffer + buffer_len */
  __shared__ extern float buffer[];

  /* bar[0] and bar[1] track if buffers buffer_0 and buffer_1 are ready to be filled,
     while bar[2] and bar[3] track if buffers buffer_0 and buffer_1 are filled-in respectively */
  #pragma nv_diag_suppress static_var_with_dynamic_init
  __shared__ barrier_t bar[4];

  if (threadIdx.x < 4)
  {
    init(bar + threadIdx.x, blockDim.x);
  }
  __syncthreads();

  if (threadIdx.x < warpSize)
  { produce(bar, bar + 2, buffer, buffer_len, in, N); }
  else
  { consume(bar, bar + 2, buffer, buffer_len, out, N); }
}

CUDA C++ cuda::ptx

#include <cuda/ptx>

__device__ void produce(barrier ready[], barrier filled[], float *buffer, int buffer_len, float *in, int N)
{
  for (int i = 0; i < N / buffer_len; ++i)
  {
    uint64_t token1 = cuda::ptx::mbarrier_arrive(ready[i % 2]);
    while(!cuda::ptx::mbarrier_try_wait(&ready[i % 2], token1)) {} /* wait for buffer_(i%2) to be ready to be filled */
    /* produce, i.e., fill in, buffer_(i%2)  */
    uint64_t token2 = cuda::ptx::mbarrier_arrive(&filled[i % 2]); /* buffer_(i%2) is filled */
  }
}

__device__ void consume(barrier ready[], barrier filled[], float *buffer, buffer_len, float *out, int N)
{
  uint64_t token1 = cuda::ptx::mbarrier_arrive(&ready[0]); /* buffer_0 is ready for initial fill */
  uint64_t token2 = cuda::ptx::mbarrier_arrive(&ready[1]); /* buffer_1 is ready for initial fill */
  for (int i = 0; i < N / buffer_len; ++i)
  {
    uint64_t token3 = cuda::ptx::mbarrier_arrive(&filled[i % 2]);
    while(!cuda::ptx::mbarrier_try_wait(&filled[i % 2], token3x)) {} /* wait for buffer_(i%2) to be filled */
    /* consume buffer_(i%2) */
    uint64_t token4 = cuda::ptx::mbarrier_arrive(&ready[i % 2]); /* buffer_(i%2) is ready to be re-filled */
  }
}

__global__ void producer_consumer_pattern(int N, float *in, float *out, int buffer_len)
{
  constexpr int warpSize = 32;

  /* Shared memory buffer declared below is of size 2 * buffer_len
     so that we can alternatively work between two buffers.
     buffer_0 = buffer and buffer_1 = buffer + buffer_len */
  __shared__ extern float buffer[];

  /* bar[0] and bar[1] track if buffers buffer_0 and buffer_1 are ready to be filled,
     while bar[2] and bar[3] track if buffers buffer_0 and buffer_1 are filled-in respectively */
  #pragma nv_diag_suppress static_var_with_dynamic_init
  __shared__ uint64_t bar[4];

  if (threadIdx.x < 4)
  {
    cuda::ptx::mbarrier_init(bar + block.thread_rank(), block.size());
  }
  __syncthreads();

  if (threadIdx.x < warpSize)
  {  produce(bar, bar + 2, buffer, buffer_len, in, N); }
  else
  {  consume(bar, bar + 2, buffer, buffer_len, out, N); }
}

CUDA C 原语

#include <cuda_awbarrier_primitives.h>

__device__ void produce(__mbarrier_t ready[], __mbarrier_t filled[], float *buffer, int buffer_len, float *in, int N)
{
  for (int i = 0; i < N / buffer_len; ++i)
  {
    __mbarrier_token_t token1 = __mbarrier_arrive(&ready[i % 2]); /* wait for buffer_(i%2) to be ready to be filled */
    while(!__mbarrier_try_wait(&ready[i % 2], token1, 1000)) {}
    /* produce, i.e., fill in, buffer_(i%2)  */
    __mbarrier_token_t token2 = __mbarrier_arrive(filled[i % 2]);  /* buffer_(i%2) is filled */
  }
}

__device__ void consume(__mbarrier_t ready[], __mbarrier_t filled[], float *buffer, int buffer_len, float *out, int N)
{
  __mbarrier_token_t token1 = __mbarrier_arrive(&ready[0]); /* buffer_0 is ready for initial fill */
  __mbarrier_token_t token2 = __mbarrier_arrive(&ready[1]); /* buffer_1 is ready for initial fill */
  for (int i = 0; i < N / buffer_len; ++i)
  {
    __mbarrier_token_t token3 = __mbarrier_arrive(&filled[i % 2]);
    while(!__mbarrier_try_wait(&filled[i % 2], token3, 1000)) {}
    /* consume buffer_(i%2) */
    __mbarrier_token_t token4 = __mbarrier_arrive(&ready[i % 2]); /* buffer_(i%2) is ready to be re-filled */
  }
}

__global__ void producer_consumer_pattern(int N, float *in, float *out, int buffer_len)
{
  constexpr int warpSize = 32;

  /* Shared memory buffer declared below is of size 2 * buffer_len
     so that we can alternatively work between two buffers.
     buffer_0 = buffer and buffer_1 = buffer + buffer_len */
  __shared__ extern float buffer[];

  /* bar[0] and bar[1] track if buffers buffer_0 and buffer_1 are ready to be filled,
     while bar[2] and bar[3] track if buffers buffer_0 and buffer_1 are filled-in respectively */
  #pragma nv_diag_suppress static_var_with_dynamic_init
  __shared__ __mbarrier_t bar[4];

  if (threadIdx.x < 4)
  {
    __mbarrier_init(bar + threadIdx.x, blockDim.x);
  }
  __syncthreads();

  if (threadIdx.x < warpSize)
  { produce(bar, bar + 2, buffer, buffer_len, in, N); }
  else
  { consume(bar, bar + 2, buffer, buffer_len, out, N); }
}

在此示例中，第一个线程束专门充当生产者，其余线程束专门充当消费者。所有生产者线程和消费者线程都参与这四个屏障中的每一个（调用bar.arrive()或bar.arrive_and_wait()），因此预期到达计数均等于block.size()。

生产者线程等待消费者线程发出共享内存缓冲区可以填充的信号。为了等待某个屏障，生产者线程必须先到达该ready[i%2].arrive()以获取令牌，然后使用该令牌执行ready[i%2].wait(token)。为简单起见，ready[i%2].arrive_and_wait()将这些操作合并在一起。

bar.arrive_and_wait();
/* is equivalent to */
bar.wait(bar.arrive());

生产者线程进行计算并填充已就绪的缓冲区，然后通过到达表示已填充的屏障filled[i%2].arrive()，发出缓冲区已填充的信号。生产者线程此时并不等待，而是等待下一次迭代的缓冲区（双缓冲）准备好接受填充。

消费者线程首先发出两个缓冲区均已准备好接受填充的信号。消费者线程此时并不等待，而是等待本次迭代的缓冲区被填充，即filled[i%2].arrive_and_wait()。消费者线程消费缓冲区后，会发出缓冲区已准备好再次接受填充的信号，即ready[i%2].arrive()，然后等待下一次迭代的缓冲区被填充。