CUDA 的块间同步方法

简介: 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() 也会引起一定的额外开销,所以需要进行实际测试和权衡。



相关实践学习
部署Stable Diffusion玩转AI绘画(GPU云服务器)
本实验通过在ECS上从零开始部署Stable Diffusion来进行AI绘画创作,开启AIGC盲盒。
目录
打赏
0
0
0
0
5
分享
相关文章
推荐场景GPU优化的探索与实践:CUDA Graph与多流并行的比较与分析
RTP 系统(即 Rank Service),是一个面向搜索和推荐的 ranking 需求,支持多种模型的在线 inference 服务,是阿里智能引擎团队沉淀多年的技术产品。今年,团队在推荐场景的GPU性能优化上又做了新尝试——在RTP上集成了Multi Stream,改变了TensorFlow的单流机制,让多流的执行并行,作为增加GPU并行度的另一种选择。本文详细介绍与比较了CUDA Graph与多流并行这两个方案,以及团队的实践成果与心得。
【AI系统】Tensor Core 基本原理
本文深入介绍了英伟达GPU中的Tensor Core,一种专为加速深度学习设计的硬件单元。文章从发展历程、卷积计算、混合精度训练及基本原理等方面,详细解析了Tensor Core的工作机制及其在深度学习中的应用,旨在帮助读者全面理解Tensor Core技术。通过具体代码示例,展示了如何在CUDA编程中利用Tensor Core实现高效的矩阵运算,从而加速模型训练和推理过程。
780 0
【AI系统】Kernel 层架构
推理引擎的Kernel层负责执行底层数学运算,如矩阵乘法、卷积等,直接影响推理速度与效率。它与Runtime层紧密配合,通过算法优化、内存布局调整、汇编优化及调度优化等手段,实现高性能计算。Kernel层针对不同硬件(如CPU、GPU)进行特定优化,支持NEON、AVX、CUDA等技术,确保在多种平台上高效运行。
222 32
【AI系统】NVLink 原理剖析
随着AI技术的发展,大模型参数量激增,对底层硬件和网络架构提出新挑战。高效训练这些模型需要大规模GPU集群及高速网络连接,以实现快速数据交换。然而,网络瓶颈限制了GPU性能的充分发挥,表明单纯增加GPU数量不能线性提升算力。因此,算存互连和算力互连技术成为关键,如PCIe、NVLink和NVSwitch等,它们通过提高数据传输速度和效率,支持大规模并行计算,解决了大规模GPU集群中的通信延迟问题,推动了万亿级模型训练的实现。
372 2
技术分享 | 软件跨架构迁移(X86->ARM)的原理及实践
针对阿里云倚天实例的软件迁移,阿里云为开发者提供了迁移工具EasyYitian和性能调优工具KeenTune,能够帮助用户解决软件迁移评估分析过程中人工分析投入大、准确率低、代码兼容性人工排查困难、迁移经验欠缺、反复依赖编译调错定位等痛点,实现业务在ARM ECS的快速适配。EasyYitian支持主流开发语言,通过系统自动化扫描可以一键生成分析报告。KeenTune通过AI算法与专家知识库的有效结合,为软件应用提供动态和静态协同调优的能力。
技术分享 | 软件跨架构迁移(X86->ARM)的原理及实践
【CUDA学习笔记】第四篇:线程以及线程同步(附案例代码下载方式)(一)
【CUDA学习笔记】第四篇:线程以及线程同步(附案例代码下载方式)(一)
775 0
技术改变AI发展:CUDA Graph优化的底层原理分析(GPU底层技术系列一)
随着人工智能(AI)的迅速发展,越来越多的应用需要巨大的GPU计算资源。CUDA是一种并行计算平台和编程模型,由Nvidia推出,可利用GPU的强大处理能力进行加速计算。
105754 1
AI助理

你好,我是AI助理

可以解答问题、推荐解决方案等