CUDA 的块间同步方法

CUDA 块间同步方法有以下三种

① Lock-Based Synchronization

② Lock-Free Synchronization

③ __threadfence(）

基于锁的块间同步

CUDA 基于锁的同步的基本思想是使用一个全局互斥量变量来计算到达同步点的线程块的数量。如下代码所示，在 barrier 函数 __gpu_sync() 中，在一个块完成它的计算之后，它的一个线程 (这里人为设置为 0 号线程，我们称之为主导线程) 将自动地向 g_mutex 添加 1 （原子操作）。然后，主导线程将重复将 g_mutex 和一个目标值 goalVal 进行比较。如果 g_mutex 等于 goalVal，那么就意味着同步完成，即每个线程块都可以进行下一阶段的计算。在该设计中，当第一次调用 barrier 函数时，将 goalVal 设置为内核中的块数 N 。然后，当连续调用 barrier 函数时，goalVal的值每次递增 N 。这种设计比保持 goalVal 常量并在每个 barrier 之后重新设置 g_mutex 更有效，因为前者节省了指令的数量并避免了条件分支 。

// lock-based
__device__ volatile int g_mutex;
// GPU lock-based synchronization function
__device__ void __gpu_sync(int goalVal)
{
  // thread ID in a block
  int tid_in_block = threadIdx.x * blockDim.y + threadIdx.y;
  // only thread 0 is used for synchronization
  if (tid_in_block == 0)
  {
    atomicAdd((int*) &g_mutex, 1);
    // only when all blocks add 1 go g_mutex
    // will g_mutex equal to goalVal
    while (g_mutex != goalVal)
    {
      // Do nothing here
    }
  }
  __syncthreads();
}

无锁的块间同步

在 CUDA 基于锁的同步中，互斥量变量 g_mutex 是通过原子函数 atomicAdd() 添加的。这意味着 g_mutex 的添加只能按顺序执行，即使这些操作是在不同的块中执行的。因此，提出一种完全避免使用原子操作的无锁同步方法。这种方法的基本思想是为每个线程块分配一个同步变量，这样每个块就可以独立地记录其同步状态，而不必争用单个全局互斥锁变量。如下代码所示，我们的无锁同步方法使用两个数组 Arrayin 和 Arrayout 来协调来自不同块的同步请求。在这两个数组中，每个元素都映射到内核中的一个线程块，即，将元素 i 映射到线程块 i ，算法大致分为三个步骤:

一开始，当 block i 准备好通信时，它的主导线程 (线程 0) 将 Arrayin 中的元素 i 设置为目标值 goalVal。block i 中的主导线程然后忙等 Arrayout 的元素 i 被设置为 goalVal 。

然后人为的让 block 1中的前 N 个（N 等于块数）线程重复检查 Arrayin 中的所有元素是否等于 goalVal ，线程 i 负责检查 Arrayin 中的第 i 个元素（一对一检查）。将 Arrayin 中的所有元素设置为 goalVal 后，每个检查线程将 Arrayout 中的相应元素设置为 goalVal 。注意，在更新 Arrayout的元素之前，block 1 的每个检查线程都会调用块内 barrier 函数  __syncthreads()。

最后，当 block i 的主导线程看到 Arrayout 中的对应元素被设置为 goalVal 时，该 block 将继续执行。

// lock-free
__device__ void __gpu_sync(int goalVal, volatile int *Arrayin, volatile int *Arrayout)
{
  // thread ID in a block
  int tid_in_blk = threadIdx.x * blockDim.y + threadIdx.y;
  int nBlockNum = gridDim.x * gridDim.y;
  int bid = blockIdx.x * gridDim.y + blockIdx.y;
  // only thread 0 is used for synchonization
  if (tid_in_blk == 0)
  {
    Arrayin[bid] = goalVal;
  }
  if (bid == 1)
  {
    if (tid_in_blk < nBlockNum)
    {
      while (Arrayin[tid_in_blk] != goalVal)
      {
        // Do nothing here
      }
    }
    __syncthreads();
    if (tid_in_blk < nBlockNum)
    {
      Arrayout[tid_in_blk] = goalVal;
    }
  }
  if (tid_in_blk = 0)
  {
    while (Arrayout[bid] != goalVal)
    {
      // Do nothing here
    }
  }
  __syncthreads();
}

从以上代码可以看出，CUDA 无锁同步中没有原子操作。所有的操作都可以并行执行。不同线程块的同步由单个块 (block 1) 中的 N 个线程来控制，可以通过调用块内 barrier 函数 __syncthreads() 来有效地同步。

__threadfence()

最后，值得注意的是，另外一种保证 CUDA 块间同步通信的正确性的办法是使用 __threadfence() （CUDA 2.2中引入了一个新的函数 ）。这个函数将阻塞调用线程，直到之前对 全局内存 或 共享内存 的写入对其他线程可见为止。但是使用 __threadfence() 也会引起一定的额外开销，所以需要进行实际测试和权衡。