CUDA C++ 高效入门第四章 -- CUDA 同步技术 101-北京尧图网络

CUDA C 高效入门第四章 -- CUDA 同步技术 1011 资料2 正文2.1 C 同步技术粗盘点2.2 CUDA 同步技术盘点3 总结1 资料前面的第二第三章节我们系统介绍了 CUDA 的线程组织和内存组织本章我们用五篇文章系统介绍 CUDA 的同步技术。在 CUDA C 高效入门第一章 – CUDA 导学之为什么要学习 GPU CUDA 中我们提到 CPU 是顺序执行加有限并行其中有限并行主要使用的是多线程编程技术。多线程编程必然涉及到数据竞争Data race于是就有了线程同步技术同步的本质是排队用的最多的就是各种锁。相比 CPU 的有限并行GPU 专为高度并行计算而设计CUDA 编程技术的核心就是大规模的线程组织。因此CUDA 编程也会遇到数据竞争Data race也需要一套对应的同步技术。本文我们将对 CUDA 的同步技术进行盘点为后面几章的学习提供全景图。本文参考资料如下1樊哲勇CUDA编程基础与实践第九十十一章2cuda-programming-guide 2.2 编写 CUDA SIMT 内核3cuda-programming-guide 2.3 异步执行本系列博客汇总链接CUDA C 高效入门系列。2 正文2.1 C 同步技术粗盘点1C 搞线程同步用的最多的是互斥锁std::mutex并通常与 std::lock_guard RAII配合使用无需手动调用 lock() / unlock() std::mutex mtx;intdata0;// 1. lock_guard最简单作用域内持有{std::lock_guardstd::mutexlock(mtx);data;}2C 多线程编程中一个常用的模型是生产者-消费者模型此时需要用到条件变量 std::condition_variable并且条件变量必须与互斥锁配合才能使用下面是一个简单的生产者-消费者样例#includeiostream#includethread#includequeue#includemutex#includecondition_variablestd::mutex mtx;std::condition_variable cv;std::queueintdata_queue;voidconsumer(){// 必须使用 std::unique_lock因为它支持灵活的加锁和解锁std::unique_lockstd::mutexlock(mtx);// wait 函数会原子地执行以下操作// 1. 检查 lambda 表达式条件// 2. 若条件为 false原子地 - a) 释放锁 lock; b) 将线程置于等待状态// 3. 被唤醒后原子地 - a) 重新获取锁 lock; b) 再次检查条件cv.wait(lock,[]{return!data_queue.empty();});// 当 wait 返回时可以保证// a) 当前线程持有锁 lock。// b) 条件 !data_queue.empty() 为 true。intdatadata_queue.front();data_queue.pop();std::cout消费数据: datastd::endl;// lock 在作用域结束时自动释放}voidproducer(){// 使用 std::lock_guard 即可因为它只需在作用域内持有锁std::lock_guardstd::mutexlock(mtx);data_queue.push(42);cv.notify_one();}3使用锁的代价是增加性能开销有些情况下希望实现无锁此时可以使用 std::atomic 原子操作。原子操作指不可中断不可分割的操作在多线程环境中对 std::atomic 对象的操作总是无数据竞争的不需要额外的互斥锁保护。下面是用法举例请特别关注里面的 CAS他是很多无锁编程的基石。#includeatomic#includeiostreamintmain(){std::atomicintvalue{0};// 基本读写 value.store(42);// 原子写入intxvalue.load();// 原子读取// 常用操作 value100;// 等价于 store(100)intyvalue;// 等价于 load()// 读-改-写操作 intoldvalue.fetch_add(10);// 原子地old value; value 10;intprevvalue;// 原子地prev value; value 1;intnextvalue;// 原子地value 1; next value;value.fetch_sub(5);// 原子地 value - 5value0xFF;// 原子地 value value 0xFFvalue|0x80;// 原子地 value value | 0x80// 交换操作 intoldValuevalue.exchange(999);// 原子地oldValue value; value 999;// CAS (Compare-And-Swap) // CAS 是无锁编程的基石。它原子地执行“如果当前值等于我期望的值就更新为新值否则不更新。”intexpected999;boolsuccessvalue.compare_exchange_strong(expected,1000);// 如果 value expected则 value 1000返回 true// 否则 expected value返回 falsestd::coutFinal value: value.load()std::endl;}2.2 CUDA 同步技术盘点1CUDA 的并行计算分两个层次第一层同次调用的同一核函数在线程束和线程块内部的并行如果没有数据竞争就不需要同步同步的本质就是排队比如两个数组相加这样的例子。如果有数据竞争就需要同步比如数组归约。第二层任务级别的并行分为以下五种类型1. 核函数与主机计算之间的并行 2. 不同核函数之间的并行 3. 核函数与数据传输之间的并行 4. 主机计算与数据传输之间的并行 5. 不同的数据传输之间的并行2CUDA 两种层次的并行计算涉及不同的同步技术这里进行列举后面四篇文章会逐个详细介绍第一层并行计算的同步技术1. __syncthreads 是专用于线程块内的同步; 2. __syncwarp 和以 _sync 结尾的线程束基本函数是专用于线程束内的同步; 3. 原子函数 atomicAdd 可同时用于线程块和线程束内的同步; 4. 协作组cooperative groups可以实现线程网格内的同步但目前主要用于线程块内的同步;第二层并行计算的同步技术任务间的并行使用 cuda stream 实现如果任务间没有依赖关系就不需要进行流同步比如数组相加的例子。如果有依赖关系就需要进行流同步。主要有以下两种方式1. 主机等待 stream 完成用 cudaError_t cudaStreamSynchronize(cudaStream_t stream) 阻塞主机直到该 stream 所有操作完成。 2. 跨 stream 依赖用 EventcudaEventRecord cudaStreamWaitEvent可以让一个 Stream 等待另一个 Stream 中的某个操作完成而不必阻塞整个主机线程是任务级并行的核心。3 总结本文是本章的全景图后面会逐个详细介绍看完后面的文章后重新回来再看这篇理解会更深

CUDA C++ 高效入门第四章 -- CUDA 同步技术 101

相关推荐

【架构设计】业务逻辑总被技术细节污染？六边形架构给出了一种解法

告别复杂环境配置｜OpenClaw 3.1.0 一键整合包部署实操指南

零依赖纯手写：C++ 实现完整神经网络，张量反向传播全打通

更多资讯

基于PCA降维与BP神经网络的交通流量预测MATLAB项目实例

金融科技Java初级岗笔试复盘：集合并发、SQL优化与Spring核心考点

金融科技初级Java笔试全复盘：题型考点与编程题实战拆解

专科生AI论文工具实测：千笔与学术猹如何解决写作痛点

电力用户用电信息采集通信协议集合与测试工具实战指南

8个AI论文网站实测：是智商税还是论文写作加速器？

今日更新

MATLAB缩进快捷键与批量格式化实战指南

用Poco实现Qt风格信号与槽的跨平台C++方案

ESXi裸机虚拟化安装指南：从硬件兼容到虚拟机部署

本周精选

金融AI落地指南：从调研报告到技术选型与工程实践

大模型上下文成本与注意力优化技术深度解析——为何O(n²)平方复杂度无法被常规优化消除｜TaoToken 统一 Key 通道实测

基于GMM的语音识别系统原理与MATLAB实现详解

本月热门

我发现了一个新思路：用 Remotion + Claude Code 像写代码一样自动化生成短视频

Windows下 Codex 中 Chrome 和 Computer Use 插件不可用问题排查及解决参考方式：TaoToken 统一 Key 配置与验证

2026 大模型集体涨价：用 Python 做企业 Token 成本测算与选型避坑（附配置）

读完文章，想聊聊您的网站？