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 总结本文是本章的全景图后面会逐个详细介绍看完后面的文章后重新回来再看这篇理解会更深 
长河君 · 2026/10/24 2:15:41 · 阅读 1.2k 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 总结本文是本章的全景图后面会逐个详细介绍看完后面的文章后重新回来再看这篇理解会更深