CUDA

CUDA

01- CUDA C 基础

什么是 CUDA？（WHAT IS CUDA?）
CUDA 架构（CUDA Architecture）
为通用计算（general-purpose computing）开放 GPU 的并行能力
开放 / 释放（GPU 的）性能
CUDA C++
基于业界标准的 C++
是一套用于实现异构编程（heterogeneous programming） 的扩展
提供直观的 API 来管理设备、内存等
本讲介绍 CUDA C++
还支持其他语言 / 绑定（bindings）：Fortran、Python、Matlab 等

CUDA C++ 入门（INTRODUCTION TO CUDA C++）
本讲（Session）你将学到什么？
从向量相加（vector addition） 入手
编写并启动 CUDA C++ 内核（kernel）
管理 GPU 内存（GPU memory）
（通信与同步的管理（communication and synchronization）→ 留到下一次课程）
【前置要求】
假设你已具备一定的 C 或 C++ 编程基础

cudaMalloc 函数 会返回一个值，来表示函数执行的错误码，可以通过检查错误码，来检查程序是否有问题。

random_ints(a, N) 初始化 host 端申请内存的数据， 该函数接收一个指向长度为 N 的向量的指针。

实际计算的数据数量可能不是恰好等于 线程块的 数量 * 线程数 数量，

调用 GPU kernel 函数 时， GPU 线程块采用向上取整的方式，（N + M - 1）/ M 向上取整

同一个线程块内的线程 具有 通信 和同步的功能

· CUDA 入门介绍
https://devblogs.nvidia.com/easy-introduction-cuda-c-and-c/
· CUDA 另一篇入门介绍
https://devblogs.nvidia.com/even-easier-introduction-cuda/
· CUDA 编程指南（CUDA Programming Guide）
https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html
· CUDA 官方文档
https://docs.nvidia.com/cuda/index.html
https://docs.nvidia.com/cuda/cuda-runtime-api/index.html （Runtime API 运行时接口）

git clone git@github.com:olcf/cuda-training-series.git
按照 readme.md 中的说明完成作业
https://github.com/olcf/cuda-training-series/blob/master/exercises/hw1/readme.md

2026 年10月 ，他们的 作业仓库依然存在，

02- CUDA 共享内存

"stencil"（模板计算）是 CUDA / 高性能计算里很常见的一种并行模式——每个输出点由它邻域内多个输入点加权求和得到。典型例子就是图像模糊、卷积、有限差分（如热扩散模拟）这类网格迭代计算。
核心套路通常是：
把数据切成块（tile），分到多个 thread / 多 block
共享内存（shared memory） 里放中心块 + 一圈 halo（边界预留数据，需要从相邻块拿）
各线程算自己那个点的 stencil 加权和，同时用 __syncthreads() 做同步防止越界读
相比逐点暴力读写，正确的 stencil 能大幅减少全局内存访问。

由上述7个橙色方块计算得到 下方橙色方块的结果

直观理解：
radius = 1：每个点的输出 = 它自己 + 左边 1 个 + 右边 1 个，一共 3 个点参与（一维时窗口宽度 = 2×radius + 1 = 3）。这就像 1D 的模糊，只和紧邻的邻居有关。
radius = 2：窗口变成 5 个点（2×2+1 = 5），会用到更远一点的邻居。
通用公式：
一维窗口大小 = 2·radius + 1
二维（网格/图像）就是 (2·radius+1)² 个点，半径越大，每个点要考虑的邻域越大，通常结果越"平滑"、但计算量和内存访问也越大。
在 CUDA 里的关键影响：
radius 直接决定你要往共享内存里加载的 halo（边界）大小——每个块除了自己那部分数据，每个方向还得额外读 radius 圈邻居；线程同步和边界处理也是按这个半径算的

向左移动计算窗口，原本的蓝色方块（线程）参与了 计算

蓝色方块会以某种方式影响7个不同输出点的计算结果

在向量加法中， 一个输入点只影响一个输出点， 一个输出点也只依赖一个输入点

由此 一维 窗口计算模板，引出共享内存的概念

线程间共享数据

术语：在一个 block（块）内部，线程之间通过共享内存（shared memory） 来共享数据
它是一种极快的片上内存，由用户自行管理（手动分配、手动释放）
用 shared 关键字声明，每个 block 单独分配一份
该数据对其他 block 的线程不可见

全局内存 和 共享内存

全局内存（Global Memory）

• 位置：位于 GPU 芯片外部的显存（device memory / VRAM）
• 可访问性：所有 block、所有线程都能读写，也经常需要和 CPU（host）之间通过 cudaMemcpy 拷贝
• 容量：大（几个 GB 到几十 GB）
• 速度：慢，延迟高（相对而言）。每次访问要经过较远的片上/片外路径
• 生命周期：由你手动用 cudaMalloc 分配、cudaFree 释放，一直存活到你释放为止
• 典型用途：存放主数据和最终结果（比如 stencil 的输入输出数组）

共享内存（Shared Memory）

• 位置：位于 GPU 芯片内部的片上存储（on-chip），离计算单元非常近
• 可访问性：仅同一个 block 内的线程可读写，其他 block 看不到（这正是你之前看的那页 "SHARING DATA BETWEEN THREADS"）
• 容量：很小，受硬件限制（现代 GPU 每个 SM 通常只有几十到一两百 KB）
• 速度：极快，比全局内存快一个数量级
• 生命周期：随 block 创建/销毁，运行时在你 __shared__ 声明后于 kernel 内自动分配，block 结束即消失
• 典型用途：做数据复用和中间缓存，避免反复读慢的全局内存

	全局内存	共享内存
位置	芯片外（显存）	芯片内（on-chip）
速度	慢	快（1 个数量级差）
容量	大	很小
谁可见	所有线程	仅本 block 内线程
生命周期	手动分配/释放	随 block 生灭
用途	存主数据	复用+缓存，加速

为什么 stencil 要用共享内存

坦白讲：stencil 的每个输出点要读它邻域的好几个点。如果每个线程都直接去全局内存读自己那几个邻居，大量数据会被反复读很多遍，全局内存访问就变成瓶颈。

做法是：每个 block 只把「自己那部分加上一圈 halo」一次性搬进共享内存，之后 block 内线程计算时都从快的共享内存里取，大幅减少撞全局内存的次数——以为代价的是用 __syncthreads() 同步，防止线程没写完就被人读。

全局内存是"大而慢的仓库"，共享内存是"小而快的临时工作台"

共享内存和缓存的区别

核心区别：显式管理 vs 硬件自动管理

• 缓存（Cache） ：硬件自动管理（硬件管理） 。你写代码时不直接控制缓存——CPU/GPU 自动把刚读过的数据存进去，下次再读更快。对你来说是第三，要不要命中、淘汰哪块，都由硬件自己决定。
• 共享内存（Shared memory） ：软件（用户）显式管理。你在代码里手动 __shared__ 声明、自己负责往里面存什么、怎么对齐、何时同步。是程序员显式控制的。

缓存是自动的、透明的；共享内存是手动的、你操盘的。

GPU 上的具体情况

重点：在 NVIDIA GPU 上，L1 缓存和共享内存在很多代架构里共用同一块片上存储（on-chip SRAM） 。也就是说，那几 KB 的快速存储，一部分可以当硬件 L1 缓存用，一部分划分出来当显式共享内存用（可以配置比例，如共享内存 48KB + L1 缓存 16KB 之类）。

这也解释了为什么二者"同源"却用途不同：

• L1 cache 帮你在反复读取全局内存时自动提速；
• 共享内存 是你主动占用那块片上空间，做跨线程的数据复用和协作——比如 stencil 里把 halo 搬进共享内存，让 block 内所有线程共享这些值，省掉重复访存。

缓存 = 硬件偷偷帮你加速的"自动缓存"；共享内存 = 你亲手管理、能在线程间共享数据的"手动缓存"。

它们常内部共用同一片片上 SRAM，但一个全自动、一个全手动。

什么是 halo

stencil 计算时，每个输出点要用它邻域里的几个输入点。你为了并行，会把大数组切成一个个 block，每个 block 负责一块连续区域。

问题是：算一个 block 边缘的点时，需要用到隔壁 block 的数据（因为邻域"伸"到了相邻块里）。这些"伸到邻居地界去借来用的那一圈数据"，就叫 halo（也叫 ghost zone、边界留边） 。

一个具体例子

假设一维 stencil，radius = 1（算一个点要看它左边 1 个 + 右边 1 个，共 3 个点），数组切成每块 16 个点：

块 A（16个点） 块 B（16个点）

[················] [················]

​ ↑

​ 这个点是块A的边界点

​ 但要算它，需要看「右边1个」——

​ 这个"右边1个"属于块B

所以块 A 在搬运数据进共享内存时，除了自己的 16 个点，还得从块 B 那边额外多搬 1 个点进来（halo）。这样计算时就不必再去全局内存读邻居了。

halo 数量 = radius × 维度方向数

• 一维 stencil：每个方向多搬 radius 个 → 每块多 2×radius（左右各一圈）

• 二维（图像）：每块多 (2×radius)² 那一圈

shared int temp[BLOCK_SIZE + 2 * RADIUS];

// 共享内存 = 本块 BLOCK_SIZE 个点 + 左右各 RADIUS 圈的 halo

为什么 halo 一定要

1. 保证并行正确性：没有 halo，块边界点就缺邻居数据，算出来是错的。
2. 避免去慢的全局内存读：halo 是搬进来"借"的，让边界计算也只需访问快的共享内存。
3. 代价是同步：搬完 halo 后所有线程要 __syncthreads() 对齐，确保数据齐了再开始算。

通俗理解

把 stencil 想成"每个村的每家要看看自己和隔壁一里内的邻居"。村里自己那一片算核心数据，但村边界那户人家还想看村外邻居——于是村子得从隔壁村借一圈地（halo） 进来，才能让边界人家的计算在"本地快速完成"。

SYNCTHREADS()（块内同步）

线程的执行顺序没有明确规定

__SYNCTHREADS()（块内同步）

• 签名：void __syncthreads();

• 同步一个 block 内的所有线程

  • 用来防止 RAW / WAR / WAW 数据竞争风险

• 所有线程都必须到达这个屏障

  才能继续

  • 如果用在条件代码里，则那个条件必须在整个 block 内一致（否则会死锁）

补充解释几个术语：

• RAW（读后写） ：一个线程读别人刚写的值，顺序错了就读到旧数据
• WAR（写后读） / WAW（写后写） ：读写/写写的顺序乱掉产生的冲突
• 它们正是 stencil 用共享内存时的坑：线程 A 往共享内存写 halo，线程 B 还没等 A 写完就去读，就会出错——__syncthreads() 就是设一道「人齐了再干活」的屏障。

"条件必须 block 内一致"的意思： 如果只有一个线程 if(x) __syncthreads();，那其他线程不等就往下走，屏障永远凑不齐 → 死锁。所以要么所有人都进，要么所有人都不进。

简单说：__syncthreads() 是站在共享内存写和读之间的一道闸门，保证块内所有线程数据同步好了再继续。

03- CUDA 基础优化

GPU 架构

🎯 Kepler SMX 单元核心介绍

SMX 是 NVIDIA Kepler 架构中升级后的增强型流式多处理器，是GK110等Kepler核心的核心计算单元，相比前代Fermi架构的SM单元实现了性能与能效的大幅跃升。

📋 核心硬件参数

• ‌CUDA核心‌：单单元集成192个SP流处理器，是Fermi SM单元的6倍规模
• ‌双精度单元‌：内置64个DP计算单元，保障高性能计算场景的浮点算力
• ‌线程调度‌：配备4组Warp调度器，每组支持双发射机制，调度器与分派单元比例达到1:2，大幅缓解大计算规模下的线程分发压力
• ‌寄存器规格‌：总寄存器容量达65536×32bit，相比Fermi架构直接翻倍
• ‌缓存配置‌：支持L1缓存可配置，范围在16KB到48KB之间，搭配统一的共享内存/缓存体系

🔍 架构设计亮点

SMX单元彻底重构了传统SM的ALU团簇规模，在保留完整几何前端、线程仲裁机制、纹理阵列和统一缓存结构的基础上，通过放大计算资源和调度资源，实现了2倍于Fermi同频下的计算性能，同时将架构的功耗效率提升到前代的3倍，是Kepler架构“性耗比第一功臣”。
在GK110旗舰计算核心中，SMX单元以3个为一组绑定进GPC单元，全核心最多可集成15组SMX，对应Tesla K40等专业计算卡的规格配置。

🎯 Maxwell/Pascal 架构 SMM 核心介绍

这张图展示的是NVIDIA Maxwell（计算能力CC5.2）与Pascal（计算能力CC6.1）两代GPU架构的增强型流式多处理器SMM的硬件规格与结构设计。

📋 核心硬件参数

• ‌SMM定位‌：Maxwell架构推出的增强版SM单元，相比前代Kepler的SMX单元重新优化了控制逻辑与调度机制，能效比大幅提升
• ‌SP核心数量‌：单SMM集成128个流处理器（CUDA核心）
• ‌DP双精度单元‌：仅配置4个双精度计算单元，该代架构大幅削减了双精度算力，优先保障图形与通用计算场景的能效表现
• ‌专属特性‌：CC6.1的Pascal架构新增了INT8整数运算支持，大幅提升了AI推理场景的吞吐能力
• ‌线程调度‌：内置4组Warp调度器，每组调度器都支持双发射机制，可同时分发两条独立指令，提升指令级并行效率
• ‌配套单元‌：集成完整的LD/ST内存加载存储单元，保障数据访问的调度效率

🖥️ 对应专业计算卡规格

• ‌Tesla M40‌：集成24个SMM单元，显存容量提供12GB/24GB两个版本
• ‌Tesla P40‌：集成30个SM单元，标配24GB大容量显存
• ‌Tesla P4‌：集成20个SM单元，标配8GB显存，主打低功耗推理场景

🔍 架构设计亮点

Maxwell的SMM采用了基于象限的全新设计，通过精简冗余控制逻辑，让单个CUDA核心的实际执行效率相比Kepler架构提升了35%，整体能效达到Kepler的2倍。而后续Pascal架构在继承该SMM基础设计的前提下，新增INT8整数运算能力，成为早期深度学习推理场景的主力硬件，凭借高性价比在AI部署领域被广泛使用。

SM

这是NVIDIA Pascal（CC6.0）与Volta（CC7.0）两代GPU架构的流式多处理器SM单元硬件规格演示页。

📋 核心硬件参数

• ‌计算核心‌：单SM集成64个SP流处理器、32个DP双精度单元，FP16算力达到2倍SP速率
• ‌架构创新‌：CC7.0的Volta架构首次引入TensorCore张量核心，新增独立int32整数运算单元
• ‌线程调度‌：P100配置2组Warp调度器，V100升级为4组Warp调度器
• ‌对应产品‌：P100集成56个SM，搭配16GB显存；V100集成80个SM，显存规格为16/32GB。

线程束 warps

启动配置（LAUNCH CONFIGURATION）

理解的关键：

• 指令按顺序发出
• 当某个操作数还没就绪时，线程会停顿（stall）
  • 单纯的内存读取本身不会让执行停顿
• 延迟通过在线程之间切换来隐藏
  • 全局内存（GMEM）延迟：> 100 个时钟周期（因架构/设计而异）
  • 算术运算延迟：< 100 个时钟周期（因架构/设计而异）

该启动多少个线程 / 线程块（threadblocks）？

结论：

• 需要足够多的线程来隐藏延迟

GPU 靠"一批线程卡住了，就立刻换另一批上来算"来掩盖内存访问的慢。因为全局内存延迟很大（超过 100 个周期），如果同时只有很少线程在跑，线程一旦等内存，GPU 核心就空转了，性能就浪费。

所以关键问题就是：到底启动多少个线程才够？ ——答案：要多到能把内存延迟的"空窗期"填满，让 GPU 核心时刻有事可做。这也是你之前问"该启动多少个 block/线程"的答案方向：不是越多越好，而是要足以隐藏延迟。

启动配置（LAUNCH CONFIGURATION）

隐藏算术延迟：

• 每个 SM 大约需要几十个 warp（约 320 个线程）
• 或者，同一 warp 内相互独立的指令也能隐藏延迟
  • 只要指令不依赖前面指令的输出，那么只需要约 5 个 warp 就够了，依此类推

最大化全局内存吞吐：

• 取决于访问模式和字长（word size）

• 需要有

  足够多的内存操作在飞行中，以打满总线带宽

  • 同一线程内相互独立的加载（load）和存储（store）
  • 不同线程的加载和存储
  • 更大的字长也有帮助（例如 float2 的吞吐相当于两倍的 float）

1. 算数延迟怎么藏：GPU 核心在等一条算术结果时，可以切到别的 warp。每个 SM 放约 320 个线程（几十个 warp）就足够填满算术延迟。如果指令之间本来就互不依赖（一条不用等前一条的结果），连 warp 都可以少很多。
2. 全局内存怎么打满带宽：想喂饱内存总线，就要让很多 load/store 同时"在路上"。手段包括：一个线程里连续发多条互相独立的访存、多个线程同时访存，以及用更大的数据类型（如 float2 一次读两个 float，占用的交易数翻倍，效率更高）。

总的就是讲：用足够的线程 + 并发/独立的访存，把 GPU 的算数和内存延迟都"填饱"，让硬件尽量满负荷工作。

附注：什么是占用率（Occupancy）？

• 衡量 SM 里实际的线程负载，与理论上限/可达到上限相比
• CUDA 自带一个占用率计算表格（occupancy calculator） （Excel 表）
• 可达到的占用率受一些 "限制因素" 影响
• 主要限制因素：
  • 每个线程的寄存器数（Registers per thread） （可用 profiler 报告，也可在编译期获得）
  • 每个线程块的线程数（Threads per threadblock）
  • 共享内存的使用量（Shared memory usage）

通俗理解：

占用率 = 一个 SM（流式多处理器）上真正"住进去"的线程，占它能住的最大容量的百分比——它越高，说明线程越"挤满"，越有利于隐藏延迟、喂饱硬件。

但要想占满，得看三样东西有没有"卡脖子"（limiter）：

1. 寄存器个数：每个线程分到的寄存器越多，能同时驻留的线程就越少（寄存器是稀缺资源）；
2. 每块线程数：block 太大的话，会受每 SM 能容纳的 block 数限制；
3. 共享内存用量：共享内存占得越多，能同时放的 block 就越少。

这三样只要有一个用 "太狠"，占用率就被它拉低了。

占用率高不高，取决于寄存器、block 线程数、共享内存这三个"天花板"谁先顶到。

全局内存

共享内存

04- 原子操作 归约与线程束洗牌

课程议程（AGENDA）

• 变换（Transformations）vs. 规约（Reductions），线程策略（Thread Strategy）
• 原子操作（Atomics），原子规约（Atomic Reductions）
• 原子操作的小技巧与窍门（Atomic Tips and Tricks）
• 经典并行规约（Classical Parallel Reduction）
• 并行规约 + 原子操作（Parallel Reduction + Atomics）
• 线程束洗牌（Warp Shuffle），用 Warp Shuffle 做规约（Reduction with Warp Shuffle）
• Warp Shuffle 的其他用途（Other Warp Shuffle Uses）
• 延伸学习（Further Study）
• 课后作业（Homework）

原子操作

使用循环 计算 数组a 中所有元素的和

变换 vs. 规约（TRANSFORMATION VS. REDUCTION）

这可以用来指导线程策略——每个线程到底做什么？

变换（Transformation）：

• 例如：c[i] = a[i] + 10;
• 线程策略：每个输出点对应一个线程（one thread per output point）

规约（Reduction）：

• 例如：*c = Σ a[i]（求数组所有元素之和，输出是一个数）
• 线程策略： ？？？

这是在引出两类完全不同的问题：

1. 变换（Transformation） ——输入输出一一对应，有多少个元素，就开多少个线程，每个线程算一个输出点，逻辑很直观（你第一讲学的 vector addition 就属于这类）。
2. 规约（Reduction） ——是把整个数组"收"成一个数（比如求和、求最大、求内积）。这就有个难题：多个输入最后变成一个结果，该怎么开线程？不能简单地"每点一线程"了。这就是接下来的内容要解决的：怎么用并行策略把 N 个元素高效归约成 1 个结果。

变换是各干各的（每点一线程），规约是众人合力收敛成一个结果（线程策略要专门设计）。

规约：朴素的线程策略（REDUCTION: NAÏVE THREAD STRATEGY）

每个输入点用一个线程（One thread per input point）

*c += a[i]; // 每个线程都执行"把 a[i] 加到结果 c 上"

（行不通） ——看 GPU 实际执行的指令：

LD R2, a[i] // 读输入（各线程独立）

LD R1, c // 读结果 c（READ）

ADD R3, R1, R2 // 加（MODIFY）

ST c, R3 // 写回 c（WRITE）

但每个线程都在同时尝试做这件事

CUDA 编程模型不保证线程的执行顺序

核心问题（也是"规约为什么难"的根源）：

朴素的直觉是"每个线程把 a[i] 加到 c 上"，但这会出错，原因是经典的竞态条件（data race） ：

• 所有线程都在做 读 c → 加 → 写回 c 这四步；
• 但线程之间没有任何执行顺序保证——两个线程可能同时"读"到同一个 c。
• 举例：c 本来是 0，线程 A 和线程 B 同时读到 c=0，各自加完都想写成 5，最后 c 只变成 5，而正确结果应该是 10——有一个线程的加法被覆盖丢了。

这正是你在前面「__syncthreads」那页看到的 RAW（读后写）冲突的具体体现：多个线程对同一地址 c 又读又写，顺序错乱。

这一页的结论就是： 单靠"每点一线程直接把结果累加到一个共享变量"行不通，因为并行写同一个变量会互相覆盖。接下来你就会学到解决办法——要么用原子操作（atomic）串起来，要么用"规约树/两两合并"的方式，把冲突改造成无共享的并行求和。

原子操作来救场（ATOMICS TO THE RESCUE）

不可分割的「读-改-写」（indivisible READ-MODIFY-WRITE）

atomicAdd(&c, a[i]);

对比之前的指令——原来会分成好几条：

LD R2, a[i] // 读输入（线程独立）

LD R1, c // 读结果 c（READ）

ADD R3, R1, R2 // 加（MODIFY）

ST R3, c // 写回（WRITE）

现在变成一条不可分割的原子指令：

RED.E.ADD.F32.FTZ.RN [c], R2; // 在 [c] 地址上做一次原子加法

• 由 L2 缓存里的特殊硬件提供支持
• 可能带来性能影响

核心要点：

上一页说的"读-改-写被多个线程劈开、互相覆盖"的问题，解决办法就是原子操作（atomic） ：把「读 c → 加 → 写回 c」这三步合并成一条不可分割（indivisible）的原子指令，硬件保证同一时刻只有一个线程能完成这一整套操作，其他线程必须排队。所以不会再有覆盖丢失。

• 例子里的 atomicAdd(&c, a[i]) 就是"原子地把 a[i] 加到 c 上"，一条底层指令 RED.E.ADD... 搞定。
• 底层是 L2 cache 里的专用硬件单元来协调这个"读写串行化"。
• 代价：因为要串行化、要排队，大量线程同时抢同一地址时会有性能开销——但至少结果是对的。

原子操作把"又读又改"打包成不可分割的一步，靠 L2 硬件保证不丢数据，解决了规约的并发冲突，代价是串行排队带来的性能损耗。

其他原子操作（OTHER ATOMICS）

• atomicMax / atomicMin —— 取最大值 / 最小值
• atomicAdd / atomicSub —— 加 / 减
• atomicInc / atomicDec —— 递增 / 递减，并处理回绕（rollover/underflow）
• atomicExch / atomicCAS —— 交换值 / 条件交换（比较并交换）
• atomicAnd / atomicOr / atomicXor —— 按位与 / 或 / 异或
• 不同的原子操作支持不同的数据类型（如 int、unsigned、float 等）

这页做了个归纳：CUDA 提供了一套完整的原子运算全家桶，覆盖了各种规约/并发场景：

• 数值型：atomicAdd/Sub（加减）、atomicMax/Min（找极值——很常用，比如并行求数组最大值）
• 计数型：atomicInc/Dec（自增自减，还处理溢出回绕——常用于给任务分配编号、计数统计）
• 通用型：atomicExch（无条件交换）、atomicCAS（compare-and-swap，条件交换，这是实现"无锁同步/自定义原子逻辑"的基础）
• 位运算型：atomicAnd/Or/Xor（按位操作，常用于标志位、状态位聚合）

共同点：都是"读-改-写"不可分割的原子指令，只是运算算子不同；且各自支持的数据类型（int/unsigned/float/bool 等）有差异，用之前要确认对应的数据版本。

CUDA 的原子操作是个完整工具箱，规约、计数、标志位、无锁同步都能用上。你上一轮已经会了 atomicAdd，这页帮你拓宽到其他算子。

原子操作的小技巧与窍门（ATOMIC TIPS AND TRICKS）

确定自己在顺序中的位置（Determine my place in an order）

• 可用于分配下一个工作项、抢占队列槽位等场景

  int my_position = atomicAdd(order, 1);

• 大多数原子操作会返回一个"旧值"——即操作发生前、该位置原本存的那个值

经典技巧（一个很重要的 trick）：

原子操作会返回"旧值"（在它执行之前那个位置原来的值），这一点非常有用。用 atomicAdd(order, 1) 举的例子：

• 假设有一个计数器 order（比如任务流水线的序号）。

• 每个线程执行

  atomicAdd(order, 1)
  

  时，同时

  • 返回给它一个唯一的旧值 my_position（比如 0、1、2、3……各不相同，因为是串行加一的）；
  • 并把 order 加 1，供下一个线程用。

• 这样，每个线程都拿到一个不重复的编号，就能唯一地确定自己在队列/任务里的位置——下一个工作项、下一个槽位就归谁，清清楚楚，且不需要额外加锁。

关键点：正因为原子操作"读到旧值 + 原子更新"是不可分割的一步，每个线程拿到的 my_position 都互不相同，天然实现了无锁排队/无锁分配。

用原子操作返回的"旧值"可以无锁地给每个线程/任务发一个唯一编号，用于排队、分配工作槽位等。

原子操作的小技巧与窍门（ATOMIC TIPS AND TRICKS）

在缓冲区中预留空间（Reserve space in a buffer）

问题：我的 kernel 里每个线程可能产生数量不等的数据。怎样才能把它们并行地收集到一个缓冲区里？

示意图：

• buffer_ptr指向缓冲区首地址（一串格子┌───┐`）
• buffer_idx 是一个"当前写入位置"的偏移量指针（▲ 指向下一个可用位置）

int my_dsize = var;                     // 我这个线程要写入的数据量（可变）
float local_buffer[my_dsize] = {...};   // 我先算好自己要写的数据

// ★ 原子操作：原子地"领取"一段空间，拿到自己这段的起始偏移
int my_offset = atomicAdd(buffer_idx, my_dsize);
// buffer_ptr + my_offset 现在就指向给我预留的第一段位置，长度是 my_dsize

// 把自己那段数据拷进缓冲区
memcpy(buffer_ptr + my_offset, local_buffer, my_dsize * sizeof(float));

这一页讲的场景（很实用）：

问题是"每个线程产出的数据长度不一样，怎么高效拼到一个数组里"——典型如：每个线程过滤出数量不定的结果、或者产生变长的输出。

解法核心就是用 atomicAdd 当"取票机" ：

1. 维护一个共享偏移量 buffer_idx（记录缓冲区当前已写到哪里）；
2. 每个线程用 atomicAdd(buffer_idx, my_dsize) 原子地拿回一个"旧值"作为自己的起始位置 my_offset，同时把游标推进 my_dsize；
3. 因为原子性，每个线程拿到的起始位置都互不重叠、自然连续，谁也不会覆盖谁；
4. 然后 memcpy 把自己那段数据塞进 buffer_ptr + my_offset。

这样无需加锁，就能让"长度各异的多个线程的输出"并行地、无缝地填进同一个缓冲区。

把 atomicAdd 当"无锁取票"，每个线程领一段连续空间往大缓冲区里写，专门用来并行收集"每线程长度不一"的数据。 （这是 GPU 里做 Stream Compaction / 变长结果收集的常用手法。）

经典并行规约（Classical Parallel Reduction）

经典并行规约（THE CLASSICAL PARALLEL REDUCTION）

原子操作无法跑满内存带宽……

• 我们希望有一种不受原子操作吞吐量限制的规约方法
• 我们希望尽可能充分地利用所有线程
• 并行规约是一种常见且重要的数据并行原语（data parallel primitive）
• 朴素的实现往往会遇到瓶颈
• 基本方法是基于「树」的方式：

这一页讲的核心思想（承上启下）：

前几页讲了原子操作能解决并发冲突，但每个线程都抢同一个原子去加，会排队、性能拉胯，跑不满内存带宽。所以需要一种更好的规约方式——经典并行规约。

它的思路本质是一棵归约树（tree-based / 多趟两两合并） ：

• 第一轮：把相邻的两个数两两相加，8 个数 → 4 个数；
• 第二轮：4 个 → 2 个；
• 第三轮：2 个 → 1 个结果。

好处：

• 没有共享原子冲突——每一轮都是互不相关的两两加法，线程之间不抢同一个变量；
• 每个线程在每一轮都很忙（充分用线程）；
• 总计算是 O(log N) 趟，比串行 O(N) 快得多。

最终关键点：并行规约是 GPU 编程里最经典、最常见的数据并行原语，几乎所有需要"把数组归成一个数"的场景（求和、求极值、点积、范数等）都靠它，而且它是后面优化（warp shuffle、共享内存）的最佳练习场。

经典并行规约用"树形两两合并"代替"原子排队"，让每一轮加法都能并行满吞吐，是规约的核心算法。 你现在看的正是它最基础的原理示意图。

问题：全局同步（PROBLEM: GLOBAL SYNCHRONIZATION）

• 如果我们能在所有线程块之间同步，就能很轻松地规约超大数组，对吧？
  • 每个块算出自己的结果后做一次全局同步
  • 一旦所有块都到达同步点，就递归地继续归约
• 一个可行的解法：拆分成多个 kernel
  • kernel 启动本身可以作为一次全局同步点（因为上一个 kernel 必须全部跑完，下一个才会开始）
  • kernel 启动的软件开销较低（但不是零）
• 其他可行解法：
  • 在线程块级规约的末尾用原子操作（atomics）
  • 用线程块排空（threadblock-draining） 的方式（参考 threadFenceReduction 示例代码）
  • 用 Cooperative Groups —— 协作式内核启动（cooperative kernel launch）

这一页讲的核心问题：

做规约时会遇到一个痛点：__syncthreads() 只能同步一个 block 内的线程，没法跨 block 全局同步。而前面提到的树形规约，第一趟每个 block 各自归约出自己的中间结果后，还需要把这些块结果再归约成一——那一轮就需要跨 block 协作。

所以这一页在讨论"怎么实现全局同步"：

• 因为每个 kernel 启动都隐含一次全局同步（下一个 kernel 必须等前一个全部完成），所以最简单的做法就是拆成多个 kernel：第一个 kernel 各块归约 → 第二个 kernel 把块的中间结果再归约 → 依次下去。
• 也提供了不用多 kernel 的其他路子：块内归约完最后用原子操作汇总；或者用 threadFence 做线程块排空；以及更高级的 Cooperative Groups 协作内核（允许跨块显式同步）。

跨 block 无法直接用 __syncthreads() 同步，所以"块间全局同步"要么靠拆多个 kernel（天然是同步点），要么靠末尾原子/threadFence/Cooperative Groups。 这是经典并行规约落地时绕不开的一环。

顺序寻址（SEQUENTIAL ADDRESSING）

核心循环代码：

for (unsigned int s = blockDim.x/2; s > 0; s >>= 1) {
  if (tid < s) {
    sdata[tid] += sdata[tid + s];
  }
  __syncthreads();   // 注意：放在 if 外面，所有线程都要经过
}

逐步推演（共享内存中 16 个值）：

10 1 8 -1 0 -2 3 5 -2 -3 2 7 0 11 0 2

• 第 1 步 · 步长 8（Stride 8）

  线程 0-7 各自加相距 8 的元素 →

  8 -2 10 6 0 9 3 7（前 8 位被更新）-2 -3 2 7 0 11 0 2

  • 过程：10+(-2)=8, 1+(-3)=-2, 8+2=10, -1+7=6, 0+0=0, -2+11=9, 3+0=3, 5+2=7

• 第 2 步 · 步长 4（Stride 4）

  线程 0-3 加相距 4 的元素 →

  8 7 13 13（前 4 位更新）0 9 3 7 -2 -3 2 7 0 11 0 2

  • 8+0=8, -2+9=7, 10+3=13, 6+7=13

• 第 3 步 · 步长 2（Stride 2）

  线程 0-1 加相距 2 的元素 →

  21 20 13 13 … 0 9 3 7 …

  • 8+13=21, 7+13=20

• 第 4 步 · 步长 1（Stride 1）

  线程 0 加相距 1 的元素 →

  41 20 13 13 …

  • 21+20=41

最后：顺序寻址（sequential addressing）是无 bank-conflict（存储体冲突）的

这一页讲的核心（两个重点）：

1. 代码逻辑：用共享内存 sdata 做树形归约，s 是步长，从 blockDim.x/2 起每轮减半（8→4→2→1）。每个"仍活跃"的线程 tid 把自己和相距 s 的元素相加。__syncthreads() 放在 if 外，保证每一轮所有线程都同步、共享内存写完再进入下一轮，避免上轮没写完就被读。

2. 顺序寻址 = 无 bank-conflict：每个线程读的地址是 tid 和 tid+s——相邻线程访问的是相邻的共享内存地址（顺序连续性），这正好让同一 warp 的线程访问不同的存储体（bank） ，从而避免 bank conflict，访存效率最高。（这正是"sequential addressing"这个优化名字的由来——和交错式 s+2*tid 索引相比，它更利于并行访存无冲突。）

这一段代码 + 顺序寻址，就是经典并行规约的核心实现——共享内存两两合并 + 步长减半循环 + 每轮同步，顺序索引还顺带避免了 bank conflict。（sdata[0] 最终就是整个 block 的和。）

插曲：网格步长循环（DETOUR: GRID-STRIDE LOOP）

我们希望能设计出能高效加载、处理任意大小数据的 kernel：

• 希望保持合并访问（coalesced loads/stores）
• 希望高效使用共享内存
• 还可以用于大神级调优——根据 GPU 选择合适数量的 block（让 block 数量与 GPU 硬件规模匹配）

示意图（gdata[0..N-1] 被按"网格宽度"步长分段访问）：

| grid-width stride | grid-width stride | grid-width stride ...

int idx = threadIdx.x + blockDim.x * blockIdx.x;   // 当前线程的初始全局索引
while (idx < N) {
    sdata[tid] += gdata[idx];                       // 处理当前这个元素
    idx += gridDim.x * blockDim.x;                 // 跳到下一个"网格宽度"的偏移处
}

一页讲的核心（grid-stride loop 是什么）：

平时我们写 kernel 常假设"一个线程对应一个数据"，但当数据量 N 远大于线程总数（或反过来）时，这种一一对应就不好用了。网格步长循环就是解法：

• 先算出线程的全局索引 idx = threadIdx.x + blockDim.x*blockIdx.x；
• 用 while(idx < N) 循环处理自己负责的数据，每处理完一个就按"网格总宽度"（gridDim.x\*blockDim.x，也就是全部线程数）跳一步，去处理下一个"同相位"的元素；
• 这样一组线程就能覆盖任意大的数据量（只要多循环几趟）。

三个好处（这页特意列的）：

1. 合并访问保持良好：相邻线程仍访问相邻地址，gdata[idx] 的访问是合并的；
2. 共享内存用得更高效：可以先算一部分、存进 sdata 累加；
3. 灵活调优：你可以选一个和 GPU 硬件（SM/线程数）匹配的启动规模，而不是被数据规模死死绑住——这正是"N 大神级调优要靠它"的原因。

网格步长循环让"一个线程固定负责多段间隔的数据"，既能处理任意大小数组、保持合并访问，又能自由选 block 规模去匹配 GPU。 是写规约/任务 kernel 时很常用的基础模式。

整合起来（PUTTING IT ALL TOGETHER）

__global__ void reduce(float *gdata, float *out){
  __shared__ float sdata[BLOCK_SIZE];   // 块内共享内存，存本块部分和
  int tid = threadIdx.x;                 // 线程在块内的编号

  sdata[tid] = 0.0f;                     // 初始化
  size_t idx = threadIdx.x + blockDim.x * blockIdx.x;  // 全局索引

  // ——阶段一：网格步长循环加载数据，先做一部分规约——
  while (idx < N) {
    sdata[tid] += gdata[idx];            // 累加到共享内存
    idx += gridDim.x * blockDim.x;       // 跳到下一个网格宽度偏移处
  }

  // ——阶段二：经典并行扫描规约（顺序寻址）——
  for (unsigned int s = blockDim.x/2; s > 0; s >>= 1) {
    __syncthreads();                     // 同步，确保 sdata 写全
    if (tid < s)                         // 前半部分线程两两相加
      sdata[tid] += sdata[tid + s];
  }

  // ——阶段三：把本块结果写到输出——
  if (tid == 0) out[blockIdx.x] = sdata[0];  // 每块的和写到对应位置
}

这段代码把前面所有技巧都串起来了，分三阶段：

1. 加载 + 初步归约：网格步长循环（上一页学的）把 gdata 的元素按相位累加到 sdata[tid]，任意大数据都能处理，且保持合并访问。
2. 块内并行归约：顺序寻址的两两合并循环（步长 blockDim.x/2 起逐次减半），__syncthreads() 保证每轮共享内存同步。结束时 sdata[0] 就是本 block 所有元素的和。
3. 写出块结果：线程 0 把 sdata[0] 写到 out[blockIdx.x]——每个 block 各写一个位置，互不冲突。

这还没算完最后一个大数组：这个 kernel 只是完成了"每 block 归约出一个小结果"，输出 out 里有 gridDim.x 个部分和。要想得到全局总和，还得把这 ${gridDim.x} 个值再规约一次（可以：

• 用第二个 kernel 再归约一遍；
• 或像前面"全局同步"那页说的，末尾用一个原子 atomicAdd 到全局 out；
• 或递归到数据足够小为止。）

而 out[blockIdx.x] 这种"每块各占一个槽位、互不干扰"的写法，正是为了便于后续这些汇总方案。

这就是经典并行规约的完整 kernel——网格步长加载 + 共享内存两两合并 + 每块结果独立写出，把一个 block 内的元素归约成一个数；全局总和还要靠第二趟（多 kernel / 原子 / 递归）来完成。

线程束洗牌

线程间通信：到目前为止（INTER-THREAD COMMUNICATION: SO FAR）

目前用共享内存（Using shared memory）：

（线程 2 要跟线程 1 共享数据时，得先写进共享内存，再让对方读出来）

如果能这样多方便（Wouldn't this be convenient）：

（线程之间能直接把数据传给对方，不用经过共享内存的读写了）

前面做规约，线程之间传数据都是走共享内存（比如 sdata[tid] += sdata[tid+s]，一个线程写共享内存、另一个线程读出来）。这虽然可行，但要经过"写共享内存 + __syncthreads() 同步 + 读共享内存"三道工序，有点绕。

这一页在引出下文：如果能把数据直接在同一个 warp 的线程之间"递"过去，多方便——即不需要共享内存、不需要同步，一条指令直接从一个线程拿值给另一个线程。这正是接下来要讲的 Warp Shuffle（线程束洗牌）。

（图上第一行 "2→1" 表示现在线程 2 传给线程 1 要绕共享内存；下面 "1" 表示想要的理想状态——一个线程直接拿另一个线程的值。）

这是引入 Warp Shuffle 的铺垫——"能不能不用共享内存、直接在线程间传数据？

引入 Warp Shuffle（INTRODUCING WARP SHUFFLE）

• 实现 warp 内线程之间的通信（intra-warp communication）

一个 warp 就是 32 个线程（GPU 调度和执行的基本单位，一次性一起跑）。这 32 个线程在 warp 内部各有一个固定编号，这个编号就叫 lane ID（车道号）。

范围固定是 0 ~ 31，永不变。你可以把它理解成"这个线程在自己那个 32 人队伍里的座位号"。

• 支持

  多种移动模式：

  • __shfl_sync()：从指定的 lane ID 拷贝数据（任意模式）
  • __shfl_xor_sync()：从计算出的 lane ID 拷贝数据（计算模式）
  • __shfl_up_sync()：从更小的 lane（delta/offset 向低） 拷贝
  • __shfl_down_sync()：从更大的 lane（delta/offset 向高） 拷贝

• 源线程和目标线程（warp 内的双方）都必须"参与" （participate）

• 用同步掩码（sync mask） 来标识并行需要恢复同步的线程

两个注意点：

• warp 内的源和目标都要"在场/参与" ：被 shfl 用到的 lane 必须同样执行这条 shfl，否则取不到值；
• 用掩码（mask） 指明哪些线程参与、需要在哪同步（reconverge）。

Warp Shuffle 让一个 warp 内的线程不用共享内存、直接互相拿数据，有任意/上/下/异或四种取数模式，适合做 warp 级别的规约优化。 接下来你马上会看到怎么用 __shfl_down_sync 做规约（替代共享内存两两合并）。

Warp Shuffle 规约（WARP SHUFFLE REDUCTION）

__global__ void reduce_ws(float *gdata, float *out){
  __shared__ float sdata[32];      // 放每个 warp 的规约结果（一个大块有32个warp时）
  int tid = threadIdx.x;
  int idx = threadIdx.x + blockDim.x * blockIdx.x;
  float val = 0.0f;
  unsigned mask = 0xFFFFFFFFU;     // 参与同步的线程掩码（全部）
  int lane    = threadIdx.x % warpSize;   // lane：线程在 warp 内 0~31
  int warpID  = threadIdx.x / warpSize;   // warpID：这是第几个 warp

  // ——阶段1：网格步长循环加载数据——
  while (idx < N) {
    val += gdata[idx];
    idx += gridDim.x * blockDim.x;
  }

  // ——阶段2：第一次 warp-shuffle 规约——
  // 每个 warp 内部 32 个线程用 shfl_down 两两归约
  for (int offset = warpSize/2; offset > 0; offset >>= 1)
    val += __shfl_down_sync(mask, val, offset);
  if (lane == 0) sdata[warpID] = val;   // 每个 warp 的 lane 0 把自己的结果写进共享内存
  __syncthreads();                       // 同步，让所有 warp 结果都到位

  // ——阶段3：只用 warp 0 做最终归约——
  if (warpID == 0){
    // 从共享内存把各个 warp 的结果取回 val（若该 warp 存在）
    val = (tid < blockDim.x/warpSize) ? sdata[lane] : 0;
    // 最后一次 warp-shuffle 归约
    for (int offset = warpSize/2; offset > 0; offset >>= 1)
      val += __shfl_down_sync(mask, val, offset);

    if (tid == 0) atomicAdd(out, val);  // 用原子操作累加到全局输出
  }
}

这段代码怎么用 warp shuffle 规约、替代共享内存树形合并？分三阶段讲：

阶段1 · 加载数据：网格步长循环（前面学过），把全局 gdata 按相位累加到每个线程自己的 val 寄存器里。

阶段2 · warp 内规约（warp shuffle 的核心） ：在每个 warp 里，用 __shfl_down_sync(mask, val, offset) 把更高 lane 的数据"往下"递过来加：

• offset 从 16 → 8 → 4 → 2 → 1 减半；
• 每次 lane L 从 lane L+offset 拿值加到自己。
• 一轮轮下来，lane 0 就拿到了整个 warp（32 线程）的和。
• 这里完全不需要共享内存和 __syncthreads()——shfl 直接在线程寄存器间递数据！

然后 if (lane == 0) sdata[warpID] = val; 把每个 warp 的 lane 0 结果存进共享内存，__syncthreads() 让所有 warp 结果齐了。

阶段3 · warp 0 汇总所有 warp：一个 block 里可能有多个 warp，让 warp 0 把共享内存里各 warp 的结果取回来（val = sdata[lane]），再做最后一次 warp shuffle 归约，最终 tid==0 用 atomicAdd(out, val) 累加到全局输出（解决跨 block 汇总）。

相比经典共享内存规约，好处在哪：

• 省共享内存空间：warp shuffle 在寄存器间传数据，不需要那块大 sdata[BLOCK_SIZE]；
• 省 __syncthreads() 同步：warp 内天然同步（同一 warp 同时执行），shfl 内部即时拿值；
• 更快：避免了共享内存的读写 + 同步开销。

代价是逻辑分 warp、更精妙、要理解 lane / mask。

Warp shuffle 规约 = 用 __shfl_down_sync 在寄存器里累加（取代共享内存+同步），每个 warp 算一个小和、再交给 warp 0 二次归约、最后原子累加到全局。 是 GPU 规约的高阶优化版。

Warp Shuffle 的优势（WARP SHUFFLE BENEFITS）

• 减少甚至消除共享内存的使用（Reduce or eliminate shared memory usage）
• 单条指令，对比原来的 2 条或更多条指令（Single instruction vs. 2 or more instructions）
• 降低显式同步的层级（Reduce level of explicit synchronization）

这页就三条，都是性能优化点：

1. 省共享内存：warp shuffle 在寄存器之间直接递数据，不需要为 warp 内规约开辟共享内存、也省下本来要放的中间数组 → 占用率（occupancy）能更高。
2. 更少的指令：原来"写共享内存 + 同步 + 读共享内存"要 2 条以上访存指令；shfl 一条指令就完成线程间传值，指令更省、执行更快。
3. 更少的显式同步：warp 内线程天然一起执行（单 warp 同步），不用在 warp 内调 __syncthreads()；只有跨 warp 汇总时才需要一次同步。

Warp Shuffle 省共享内存、省指令、省显式同步，是规约/交换数据时更高效的手段。 这也是为什么第二讲专门引入它来优化经典的共享内存规约。

Warp Shuffle 的小技巧与窍门（WARP SHUFFLE TIPS AND TRICKS）

我们还能用它做什么？

• 单条指令把某个值广播给 warp 内的所有线程（Broadcast a value to all threads in the warp in a single instruction）
• 执行 warp 级的 prefix sum（前缀和 / 扫描） （Perform a warp-level prefix sum）
• 原子聚合（Atomic aggregation）

这页介绍 warp shuffle 除了规约之外的三种常见用途：

1. 广播（Broadcast） ：用 __shfl_sync(mask, val, srcLane) 指定从某个 lane 取值，warp 内所有线程一条指令都拿到同一个值——常用在需要所有线程读同一份数据、又不想走共享内存广播的场景。
2. warp 级前缀和（prefix sum / scan） ：前缀和是"每个位置 = 前面所有元素之和"（如 1 2 3 → 1 3 6）。用 warp shuffle 可以在一串 lane 上做并行扫描，是很多算法（如流压缩、桶排序）的底子。你课程里第二讲往后可能专门讲扫描，shfl 是实现它的利器。
3. 原子聚合（Atomic aggregation） ：可以把本来要多次原子操作的更新，先在 warp 内用 shuffle 聚合成少数几个值，再统一做一次原子操作 → 减少原子争用、提升性能。

Warp Shuffle 不止能规约，还能广播单值、做 warp 级前缀和、减少原子操作次数（原子聚合）。 是高性能 CUDA 里很灵活的一个原语。

05- 托管内存

课程议程（AGENDA）

• 托管内存（Managed Memory）：基本概念、目标、优势
• 按需分页（Demand-Paging）、超额订阅（Oversubscription）、并发（Concurrency）、原子操作（Atomics）
• 使用场景：深拷贝（Deep Copies）、链表（Linked Lists）、C++ 对象（C++ Objects）、图遍历（Graph Traversal）
• 性能：预取（Prefetching）、提示（Hints）
• 多 GPU 考量（Multi-GPU Considerations）
• 延伸学习（Further Study）
• 课后作业（Homework）

这一讲的核心主题——托管内存（Managed Memory / Unified Memory） ，简单说就是让 CPU 和 GPU 共享同一块虚拟地址空间，由系统自动在 CPU/GPU 间按需搬运数据，从而省掉你手动 cudaMemcpy 的麻烦。

CUDA 三步处理流程（THE CUDA 3-STEP PROCESSING SEQUENCE）

回顾第一讲……

简单处理流程（SIMPLE PROCESSING FLOW）

结构图：CPU ⟷ 桥（Bridge）/ PCIe 或 NVLink 总线 ⟷ GPU（GigaThread、L2、DRAM），两侧分别是 CPU Memory 和 GPU 的 DRAM。

1. 把输入数据从 CPU 内存拷贝到 GPU 内存
2. 加载 GPU 程序并执行，为了性能在片上缓存数据
3. 把结果从 GPU 内存拷贝回 CPU 内存

如果能不用（写代码实现）第 1 步和第 3 步，那该多好？

（Wouldn't it be nice if we didn't have to do steps 1 and 3?）

这一页在讲什么（托管内存的动机）：

回顾第一讲学的传统 CUDA，要分成个个显式的拷贝：程序里必须先 cudaMemcpy 把数据从 CPU 内存搬到 GPU 显存（第 1 步），算完再 cudaMemcpy 拷回 CPU（第 3 步）。这带来一堆麻烦：

• 你得自己管理两份内存（host 和 device） ，容易出错（拷贝错方向、漏释放、同步问题）；
• 数据太大、结构复杂（比如链表、对象）时，手动深拷贝特别痛苦。

这一页的标题就是整讲的核心目标——如果把第 1、3 步（手动拷贝）省掉，让系统自动搬数据，编程会轻松很多。这就是接下来要讲的 Managed Memory（托管内存 / Unified Memory） 想解决的问题：数据由系统统一管理，CPU/GPU 按需自动迁移，你不再手写那两步拷贝。

这页用"能省掉手动拷贝吗？"问出整讲主题，托管内存就是让系统替你自动管理数据在 CPU/GPU 间的搬运，省掉第 1、3 步。

统一内存（UNIFIED MEMORY）

减少开发者工作量（Reduce Developer Effort）

• 自 CUDA 6 / Kepler+ 架构起提供

核心：CPU 和 GPU 共享一个统一内存

• 分配上限可达 GPU 内存大小（Allocate Up To GPU Memory Size）

更简单的编程和内存模型（Simpler Programming & Memory Model）

• 一次分配、一个指针，随处可访问（Single allocation, single pointer, accessible anywhere）
• 无需显式拷贝（Eliminate need for explicit copy）
• 简化代码移植（Simplifies code porting）

通过数据局部性维持性能（Maintain Performance through Data Locality）

• 数据迁移到正在访问它的处理器（Migrate data to accessing processor）
• 保证全局一致性（Guarantee global coherence）
• 仍然允许显式手动调优（Still allows explicit hand tuning）

(结构中 CPU 与 GPU 之间以「统一内存」连接)

这一页把托管内存（Unified Memory）的核心价值讲全了，分三点：

1. 更省事（编程模型） ：
   • CPU 和 GPU 共用同一份统一内存地址空间；
   • 你只需一次分配、拿一个指针，CPU 端和 GPU 内核都能直接用同一个指针访问，不用再手写 cudaMemcpy 拷贝；
   • 代码移植（把串行程序改成用 GPU）也简单，因为它"看起来像普通内存"。
2. 仍然快（性能靠数据局部性） ：
   • 系统会自动把数据迁移到"正在访问它的那一边"（CPU 或 GPU） ，保证访问者是"近水楼台"，性能不至于因为"统一"就变慢；
   • 自动维护全局一致性（数据在所有端看到的都是一致的最新值）。
3. 仍可手动调优：你如果需要极致性能，还是可以显式控制数据放置/预取（对应你课程后面会讲的 Prefetching、Hints），统一内存不禁止手动优化。

一句话对比第一讲：以前要 cudaMalloc + cudaMemcpy 手动搬；现在一个 cudaMallocManaged 拿一个指针，系统自动帮你搬，既省事（少写拷贝代码）又能靠"数据放对地方"保住性能。

统一内存让 CPU/GPU 共享一块内存、一个指针随处可用，省掉手动拷贝，同时靠自动迁移保证性能、还支持手动调优。 这是第三讲的核心概念。

🎯 基础架构说明

这套架构打通了Pascal及以上架构GPU和CPU的内存空间，通过统一内存层实现了跨硬件的内存共享，支持分配超出GPU物理显存大小的内存空间，打破了传统GPU显存容量的硬限制。

✨ 三大核心能力

• ‌支持超大模型运行‌
  可以超额订阅GPU显存，内存分配上限最高可达到整个系统的物理内存总大小，让大尺寸的AI模型、科学计算数据集也能正常加载运行。
• ‌简化数据访问逻辑‌
  自动维护CPU和GPU之间的数据一致性，同时支持统一内存原子操作，开发者无需再手动编写大量数据拷贝代码，大幅降低异构编程的复杂度。
• ‌精细化性能调优‌
  提供了两类优化接口：通过cudaMemAdvise API传递内存使用提示，还可以借助显式预取API主动控制数据迁移，进一步规避自动页迁移带来的性能开销。

这张幻灯片深入到了底层机制，标题是 “统一内存在 Pascal+ 架构上是如何工作的”，副标题点明了核心机制：处理 CPU 和 GPU 的缺页中断（Servicing CPU and GPU Page Faults）。

它通过一段具体的代码示例，展示了 GPU 和 CPU 是如何通过“按需分页”协同工作的。

以下是对这张图的详细解读：

一、 代码逻辑梳理

幻灯片上方展示了 CPU 端和 GPU 端的代码，这是一段典型的异构计算程序：

1. CPU Code (主机端代码):
   • cudaMallocManaged(&array, size);：这是关键。使用统一内存分配函数，系统会分配一块既可以被 CPU 访问，也可以被 GPU 访问的托管内存。
   • memset(array, size);：CPU 初始化这块内存。
   • setValue<<<...>>>(array, size/2, 5);：启动一个 GPU 内核（Kernel），注意参数 size/2，说明 GPU 只准备修改数组前半部分的值。
2. GPU Code (设备端代码):
   • void setValue(char *ptr, int index, char val) { ptr[index] = val; }：这是一个简单的 GPU 内核函数，作用是往指针 ptr 的指定 index 位置写入值 val。

二、 内存映射状态（Memory Mapping）与缺页中断（Page Fault）

幻灯片下半部分用色块直观地展示了在 GPU 执行 ptr[index] = val 这一瞬间，内存的状态：

• CPU Memory Mapping（右侧）：
  • 内存被划分成了多个页面（方块）。
  • 虚线框代表 CPU 视角下分配给该数组的整体虚拟内存空间。
  • 所有页面都是浅绿色和淡蓝色，表示在 CPU 视角下，这些数据当前驻留在系统内存（System Memory）中。
• GPU Memory Mapping（左侧）：
  • 同样有代表整体虚拟内存空间的虚线框。
  • 大部分页面是深绿色，并且是实线，意味着这部分数据目前不在 GPU 的物理显存中（或者被标记为无效）。
  • 关键点：Page Fault（缺页中断）。箭头指向了一个亮绿色的页面。当 GPU 尝试写入数组的前半部分（size/2）时，发现所需的数据不在显存里，于是触发了 GPU 的缺页中断。
• Interconnect（互连总线）：
  • 底部粗大的灰色双向箭头代表 PCIe 总线或 NVLink。
  • 当缺页中断发生后，系统必须通过这个互连总线，按需将 CPU 内存中对应的页面（亮绿色部分）迁移到 GPU 显存中。

三、 这张图传达的核心技术原理

这张图完美解释了 Pascal 架构引入的硬件级页面迁移机制：

1. 按需迁移（On-Demand Migration）：GPU 不需要一次性把整个数组都拷贝过去。它只拷贝真正需要访问的页面（图例中的亮绿色页面）。
2. 统一寻址（Unified Addressing）：CPU 和 GPU 使用同一个指针 array，硬件通过页表将虚拟地址映射到不同的物理位置（系统内存或显存）。
3. 透明性：程序员不需要写任何显式的 cudaMemcpy 代码。ptr[index] = val 这一行代码触发的缺页中断，由底层硬件和驱动自动完成数据的搬运。

总结：
这张图生动地展示了“按需分页”的微观过程。它告诉我们，统一内存的便利性（代码简洁）背后，是硬件在默默处理昂贵的缺页中断和通过 PCIe 总线进行的数据迁移。这也解释了为什么上一张幻灯片中需要提供 cudaMemAdvise 等性能调优 API——因为如果频繁触发缺页中断，性能损耗是非常大的。

📌 Pre-Pascal 架构统一内存规则总结

这张图梳理了Pascal架构之前的GPU，以及特定系统环境下统一内存（UM）的运行机制与使用规范。

🎯 生效场景

这套受限的统一内存机制会在两类环境下自动启用：

• 设备架构早于Pascal（Jetson系列属于特殊例外情况）
• 运行在Windows系统且CUDA版本为9.x及以上

🔍 核心行为特性

• ‌批量数据迁移‌：托管内存会在Kernel启动时整体迁移到GPU显存，哪怕Kernel代码里没有显式访问的部分数据，也会被一并搬运。
• ‌同步触发回迁‌：Kernel执行完成后，必须调用cudaDeviceSynchronize()，CUDA运行时才会把数据同步回CPU侧，让主机代码可以正常访问更新后的内容。
• ‌功能存在限制‌：不支持CPU与GPU并发访问，没有硬件级的按需页迁移能力，也无法实现GPU显存的超额订阅，不能分配超出物理显存大小的统一内存空间。

💡 极简使用规范

这套环境下的统一内存上手门槛极低：

• 原本使用malloc()或者new分配内存的位置，直接替换为cudaMallocManaged()即可
• 原本使用free()或者delete释放内存的位置，替换为cudaFree()完成资源回收

🎯 Pascal+ 统一内存显存超分特性

这张图展示了Pascal及更新架构GPU的统一内存核心能力——显存超额订阅，突破物理显存容量的硬限制。

📋 核心代码示例

cppvoid foo() {
    // Assume GPU has 16 GB memory
    // Allocate 64 GB
    char *data;
    // be careful with size type:
    size_t size = 64ULL*1024*1024*1024;
    cudaMallocManaged(&data, size);
}

这段代码尝试在仅拥有16GB物理显存的GPU上，分配总计64GB的统一内存空间，这里特意使用64ULL确保数值按64位无符号长整型处理，避免32位整数溢出导致的内存计算错误。

🔍 按需分页迁移机制

Pascal架构支持仅将部分内存页常驻在GPU物理显存中，其余内存页托管在系统内存里。当GPU代码首次访问某块不在显存中的内存页时，硬件会自动触发页错误，驱动程序按需将对应页面从系统内存迁移到GPU显存中，全程无需开发者手动编写数据拷贝逻辑。

⚠️ 旧架构兼容性限制

这套显存超分能力在Kepler和Maxwell架构的GPU上无法正常运行，这两类旧架构没有硬件级的按需页迁移引擎，统一内存只能在物理显存容量范围内分配，无法实现超出显存大小的大内存分配。

🎯 Pascal+ 统一内存并发访问特性

这张图展示了Pascal及更新架构GPU在统一内存上的核心升级：支持CPU与GPU对托管内存的并发访问，彻底改变了旧架构的内存访问行为。

这段代码在启动GPU核函数后，没有调用任何同步接口，直接在CPU侧访问统一内存指针，是典型的无同步并发访问场景。

🔍 新旧架构行为差异

• ‌Pascal及更新架构‌：这种无同步并发访问只会触发硬件级的页错误，由GPU的页迁移引擎自动处理，程序不会直接崩溃，硬件会自动完成内存页的按需迁移，保证两边访问的正确性。
• ‌Pascal之前的旧架构‌：在核函数运行期间，CPU直接访问托管内存会直接触发致命的段错误，程序直接终止，完全不支持这种并发访问模式。

⚠️ 关键注意事项

统一内存的并发访问本身不会提供任何内存排序和数据可见性的自动保证，CPU和GPU的读写操作顺序没有默认约束，可能出现数据可见性异常。如果需要保证跨设备的操作顺序和数据一致性，可以借助Pascal架构支持的系统级原子操作，来实现跨CPU和GPU的原子同步，避免数据竞争问题。

🎯 Pascal+ 系统级原子操作特性

这张图展示了Pascal及更新架构GPU在统一内存体系下的全新能力——系统级原子操作，打破了传统原子操作仅能在单一设备内生效的限制。

这段代码通过cudaMallocManaged分配了4字节的统一内存，初始值设为0；GPU核函数使用atomicAdd_system对该地址执行原子加10操作，同时CPU侧也通过系统原子指令__sync_fetch_and_add对同一地址执行原子加10操作，最终得到正确的累加结果20。

🔍 核心运行机制

Pascal架构对系统级原子操作提供了分层支持：

• ‌NVLink直连场景‌：通过NVLink高速互联链路直接完成跨设备原子操作，硬件原生支持，性能几乎无损耗
• ‌PCIe连接场景‌：通过软件辅助的方式实现跨CPU和GPU的原子操作，保证操作的原子性与正确性

⚠️ 架构兼容性限制

系统级原子操作是Pascal及更新架构的专属特性，Kepler和Maxwell等旧架构GPU完全不支持该能力，无法实现CPU与GPU之间跨设备的原子同步，这类场景下很容易出现数据竞争导致计算结果错误。

🎯 CUDA 嵌套结构体深拷贝场景解析

这张图展示了CUDA异构编程中处理带二级指针结构体的经典深拷贝方案，解决了主机侧嵌套指针无法直接在GPU侧生效的核心痛点。

📋 场景核心背景

图中定义的dataElem是典型的嵌套指针数据结构：

struct dataElem {
  int key;
  int len;
  char *name;
}

结构体内部的name指针指向一段独立的字符缓冲区，这类带二级指针的数据结构如果直接整体拷贝到GPU，设备侧拿到的主机指针是完全无效的，会直接触发非法内存访问错误。

🔍 传统深拷贝执行流程

图中给出的launch函数完整演示了手动深拷贝的标准步骤：

• 先在GPU侧分配N个dataElem结构体的存储空间，将主机侧的结构体数组整体拷贝到设备
• 遍历每一个结构体元素，单独为每个元素的name指针指向的字符缓冲区分配设备显存
• 将主机侧的字符内容拷贝到新分配的设备缓冲区中
• 最后修改设备侧结构体里的name指针，让它指向刚分配好的设备缓冲区地址，完成指针的“修正”

⚠️ 关键注意事项

• 主机与GPU是两个完全独立的地址空间，主机侧的指针地址在GPU侧没有任何意义，直接使用必然报错
• 深拷贝过程中必须同时拷贝结构体对象本身和它指向的所有嵌套缓冲区，不能遗漏任何一层指针的地址修正
• 所有通过cudaMalloc分配的显存都必须配对调用cudaFree释放，否则会造成显存泄漏

💡 统一内存简化方案

在Pascal及更新架构的CUDA环境中，使用cudaMallocManaged分配统一内存可以彻底省略这套繁琐的手动深拷贝流程：

• 统一内存的指针同时对CPU和GPU可见，不需要手动维护主机指针到设备指针的映射关系
• 嵌套的二级指针在统一内存地址空间中天然有效，不需要逐一遍历修正指针地址
• 硬件自动按需完成内存页的迁移，大幅降低带复杂指针数据结构的开发出错概率

🎯 CUDA C++ 统一内存对象管理方案

这张图展示了CUDA环境下通过C++面向对象特性，重载new/delete实现统一内存自动化管理的优雅方案，彻底简化异构编程中的内存操作。

📋 基类Managed实现解析

这个基类通过重载全局operator new和operator delete，让所有继承它的子类自动获得统一内存分配能力：

• ‌分配逻辑‌：重载的new运算符内部直接调用cudaMallocManaged分配统一内存，返回的指针天然可同时被CPU和GPU访问
• ‌同步保障‌：分配和释放后都调用cudaDeviceSynchronize，确保内存操作完成后再返回，避免异步操作带来的访问异常
• ‌自动释放‌：重载的delete运算符自动调用cudaFree释放统一内存，无需开发者手动处理CUDA显存释放逻辑

🔍 派生类umString实现解析

通过继承Managed基类，快速构建支持统一内存的自定义字符串类：

• ‌继承复用‌：直接复用基类的统一内存分配/释放逻辑，无需重复编写内存管理代码
• ‌深拷贝构造‌：自定义拷贝构造函数，为字符串内部的data指针单独分配统一内存，实现完整的深拷贝，避免浅拷贝带来的野指针问题
• ‌跨设备兼容‌：整个对象的所有内存都在统一内存地址空间中，CPU和GPU都可以直接访问字符串内容，不需要额外的指针修正操作

💡 核心优势与注意事项

• ✅ ‌代码极简‌：彻底告别传统CUDA中繁琐的手动深拷贝流程，嵌套指针天然在统一内存中生效
• ✅ ‌面向对象友好‌：完全符合C++面向对象编程习惯，普通C++开发者几乎零学习成本就能上手异构内存管理
• ⚠️ ‌同步开销‌：基类中强制的cudaDeviceSynchronize会带来少量同步开销，适合对开发效率优先的场景，极致性能场景可根据需求调整同步策略
• ⚠️ ‌对齐保障‌：重载的new运算符天然满足CUDA统一内存的内存对齐要求，不会出现非对齐访问导致的未定义行为

🚀 Pascal+ 按需分页性能调优要点

这张图聚焦Pascal及更新架构GPU中统一内存按需分页机制的性能陷阱，讲解了页故障带来的额外开销问题。

⚠️ 按需分页的性能痛点

• ‌运行速度显著下降‌：完全依赖自动按需页迁移的核函数，运行速度会远慢于Pascal之前架构的统一内存实现，甚至比非统一内存的常规实现还要慢。
• ‌逐页故障开销‌：每一次GPU访问不在显存中的内存页，都会触发一次页故障，带来对应的服务处理开销。
• ‌批量数据迁移低效‌：如果完全依赖页故障逐页搬运大量数据，每一页都要承担故障处理开销，整体效率会非常低下。
• ‌批量拷贝更高效‌：对于大块数据的迁移，单次类似memcpy的批量操作，远比逐页触发故障迁移的综合效率高得多。

这段代码直接分配统一内存后启动核函数，没有做任何数据预取操作，GPU首次访问每一个内存页时都会触发缺页中断，产生大量不必要的性能开销。

✅ 针对性优化方案

• 优先使用cudaMemPrefetchAsync接口，在启动核函数之前就把即将访问的统一内存页批量预取到GPU显存中，避免运行时动态缺页。
• 配合cudaMemAdvise接口给驱动传递内存访问提示，告知驱动该内存主要由GPU访问，让驱动提前优化页面驻留策略。
• 避免随机访问统一内存，尽量采用连续顺序的访问模式，减少页故障的触发次数。

🚀 Pascal+ 统一内存显式预取优化

这张图讲解了Pascal及更新架构GPU中，通过显式预取接口彻底解决按需分页性能开销的核心调优手段。

📋 核心API基础信息

cudaMemPrefetchAsync是统一内存体系下的关键性能接口，完全可以替代传统的异步内存拷贝操作：

• ‌函数原型‌：cudaMemPrefetchAsync(ptr, length, destDevice, stream)
• ‌核心特性‌：支持将统一内存数据异步预取到任意指定GPU，也可以直接预取到CPU侧
• ‌性能收益‌：提前批量完成数据迁移，彻底消除运行时逐页缺页的额外开销，把统一内存的性能恢复到和传统显式拷贝方案相当的水平

这段代码的执行流程非常清晰：

1. 先分配统一内存空间
2. 启动核函数之前，就把全部数据异步预取到编号为0的GPU设备显存中
3. 启动GPU核函数，此时所有数据已经全部在显存内，不会触发任何缺页中断
4. 核函数执行完成后，再把修改后的数据异步预取回CPU侧，主机后续可以直接访问

✅ 最佳实践要点

• 优先在独立的CUDA流中执行预取操作，可以和其他计算任务并行执行，进一步隐藏数据迁移延迟
• 配合cudaMemAdvise接口给驱动传递内存访问提示，比如标记数据为GPU优先访问，进一步优化驱动的内存调度策略
• 避免CPU和GPU之间频繁交替访问同一块内存，否则会导致数据在两侧反复迁移，反而带来额外的带宽开销

🚀 Pascal+ 显式内存提示性能调优

这是Pascal及更新架构GPU中统一内存体系下的高级性能调优手段，通过向CUDA运行时传递内存访问行为提示，在不触发额外数据迁移的前提下，大幅优化内存调度策略。

📋 核心API基础信息

cudaMemAdvise是统一内存配套的非侵入式优化接口，函数原型为：

cudaMemAdvise(ptr, count, hint, device);

它的核心特性是‌本身不会触发任何数据移动‌，仅作为访问行为提示传递给CUDA驱动，让驱动在后续的内存调度中做出更贴合实际访问模式的决策，减少不必要的缺页中断和数据往返迁移。

🔍 三类关键提示详解

• ‌cudaMemAdviseSetReadMostly‌：指定内存区域以读为主，驱动会在各个访问该数据的处理器上创建只读副本，避免数据在不同设备间反复迁移，非常适合多GPU共享只读参数的场景。
• ‌cudaMemAdviseSetPreferredLocation‌：建议驱动将该内存区域的物理存储优先放置在指定设备上，让数据尽可能驻留在访问频率最高的位置，降低跨设备访问延迟。
• ‌cudaMemAdviseSetAccessedBy‌：提示指定设备将会访问该内存区域，驱动会提前为该设备建立直接的内存映射，减少首次访问时的映射建立开销。

⚠️ 关键使用注意事项

• 所有提示都是非强制性的，CUDA驱动会根据实际硬件状态、内存负载情况决定是否采纳，极端场景下提示可能被忽略。
• 提示需要和cudaMemPrefetchAsync预取操作配合使用，才能发挥最大的优化效果，单独使用内存提示通常不会带来立竿见影的性能提升。
• 该优化手段仅针对统一内存分配的区域生效，传统手动分配的显存/主机内存无法使用这套提示机制。

🚀 cudaMemAdviseSetReadMostly 特性解析

这是Pascal及更新架构CUDA统一内存体系下，专门针对读密集型数据的高级优化提示接口。

📋 核心特性说明

• ‌适用场景‌：专门针对几乎不会被修改的只读数据，是多GPU共享只读参数场景的最优优化手段。
• ‌自动副本机制‌：统一内存系统会为每一个访问该数据的处理器，在本地创建一份独立的只读副本，彻底避免数据在不同设备之间反复迁移。
• ‌写失效规则‌：如果有任意一个处理器对该数据执行写入操作，除了正在执行写入的处理器上的副本之外，其他所有副本都会被自动标记为无效，保证数据一致性。
• ‌特殊属性‌：该提示的device参数会被直接忽略，不需要指定特定目标设备，驱动会自动在所有访问该数据的设备上创建本地副本。

✅ 优化收益

• 完全消除多GPU共享只读数据时的跨设备数据迁移开销，大幅降低缺页中断触发概率。
• 在显存负载超过50%的场景下，实测可以带来最高25%左右的性能提升，非常适合大模型多卡分布式推理、训练场景下的共享权重参数优化。
• 该提示本身不会触发任何数据移动，仅作为访问模式提示传递给CUDA驱动，不会引入额外的传输开销。

⚠️ 注意事项

• 该提示仅适合读多写少的场景，如果数据会被频繁修改，反而会因为副本失效同步带来额外的性能损耗。
• 提示属于非强制性优化建议，CUDA驱动会根据当前硬件负载、内存占用情况决定是否采纳，极端高负载场景下可能被驱动忽略。

🚀 cudaMemAdviseSetPreferredLocation 特性解析

这是Pascal及更新架构CUDA统一内存体系下，用于指定数据最优物理存储位置的高级性能调优接口。

📋 核心基础特性

• ‌定位作用‌：向CUDA运行时建议该内存区域的最优驻留处理器，可以是CPU也可以是任意指定GPU。
• ‌非侵入性‌：该提示本身不会自动触发任何数据迁移操作，仅作为访问模式提示传递给驱动。
• ‌迁移触发规则‌：数据不会立刻被搬运到目标位置，只有当后续发生按需缺页访问、或者主动调用cudaMemPrefetchAsync预取时，才会完成向首选位置的迁移。

🔍 跨设备访问优化机制

当其他非首选处理器访问该数据时，驱动会优先尝试建立P2P远程映射，让访问方直接通过映射地址读取数据，不需要执行完整的数据拷贝迁移。
只有当硬件不支持远程P2P映射时，驱动才会 fallback 到传统的数据迁移逻辑，把数据完整搬运到访问方的内存中。

✅ Volta+ 架构增强特性

从Volta架构开始，GPU新增了硬件访问计数器，能够自动追踪不同处理器对内存页的访问频率，结合用户传入的首选位置提示，驱动可以做出更智能的内存调度决策，进一步减少不必要的数据往返迁移，降低缺页中断开销。

⚠️ 适用场景与注意事项

• 非常适合CPU-GPU协同处理、多GPU分布式计算场景，提前指定数据优先驻留的位置，避免数据在不同设备间反复迁移产生的抖动。
• 该提示属于非强制性优化建议，CUDA驱动会根据当前硬件负载、显存占用情况决定是否采纳，极端高负载场景下可能被驱动忽略。
• 仅对cudaMallocManaged分配的统一内存区域生效，传统手动分配的显存/主机内存无法使用该优化。

🚀 cudaMemAdviseSetAccessedBy 特性解析

这是Pascal及更新架构CUDA统一内存体系下，专门用于提前建立跨设备内存映射的高级性能调优接口。

📋 核心基础特性

• ‌无侵入性‌：该提示本身不会触发任何数据移动，也不会改变数据当前的物理存储位置，仅作为访问模式提示传递给CUDA驱动。
• ‌预建P2P映射‌：被指定的目标处理器会提前获得该内存区域的P2P远程访问映射，不需要等到首次访问时再动态建立映射。
• ‌自动映射更新‌：如果后续该内存区域发生数据迁移，驱动会自动同步更新所有已建立的P2P映射，保证访问始终有效。
• ‌核心优化目标‌：让目标处理器访问该内存区域时，完全避免触发缺页中断，消除首次访问的映射建立开销。

🔍 典型适用场景

• 多GPU协同计算场景下，多个GPU需要频繁交互访问同一块共享内存，提前为所有访问方建立映射，避免反复触发缺页中断。
• CPU和GPU交替访问同一块统一内存的场景，提前建立双向映射，减少跨设备访问的延迟抖动。
• 适合数据不会频繁迁移、但多个处理器都需要高频访问的场景，在不移动数据的前提下大幅降低访问开销。

⚠️ 注意事项

• 该提示属于非强制性优化建议，CUDA驱动会根据当前硬件负载、P2P支持情况决定是否采纳，极端高负载场景下可能被驱动忽略。
• 仅对cudaMallocManaged分配的统一内存区域生效，传统手动分配的显存/主机内存无法使用该优化。
• 提前建立P2P映射会占用少量驱动侧的映射表资源，不建议对大量小内存区域无差别使用该提示，避免资源浪费。

🚀 CUDA统一内存核心总结要点

这是CUDA统一内存（UM）技术的最终经验总结，清晰点明了统一内存的定位、优势与使用边界。

📋 统一内存核心定位

统一内存的核心设计目标从来不是让已经写得非常优秀的CUDA代码跑得更快，它的第一优先级是‌降低异构编程门槛、大幅提升开发者的生产效率‌，让开发者不用再耗费大量精力手动维护主机和设备两侧的内存镜像、处理繁琐的显式数据拷贝逻辑。

⚠️ 关键使用边界

• 在绝大多数场景下，统一内存的性能上限无法超过资深开发者手动编写的精细化数据搬运代码，盲目的“开箱即用”不做任何优化，反而可能因为大量缺页中断让程序性能出现大幅下跌。
• 统一内存的并发行为更难精准控制，开发者很难像手动管理内存那样，精确调度计算和数据传输的重叠逻辑。
• 显存超额订阅能力并不是“魔法”，它无法让任意数据集、任意算法自动获得GPU级别的性能，频繁的页换入换出反而会带来严重的性能抖动。

✅ 不可替代的适用场景

• 它可以支撑很多传统CUDA内存模型下很难实现的设计模式，比如图遍历这类存在大量随机指针跳转的算法，统一内存可以天然让嵌套指针在CPU和GPU侧同时生效，大幅降低实现难度。
• 针对大量第三方库混合调用的场景，每个库都会独立占用GPU显存且互不感知对方的内存占用，统一内存的超额订阅能力可以在不重写整个代码库的前提下，直接解决显存不足的问题，大幅降低改造工作量。

🚀 CUDA多GPU统一内存机制解析

📋 Pre-Pascal 架构运行机制

• ‌分配位置规则‌：统一内存分配的物理空间，和传统cudaMalloc行为完全一致，直接在当前被选中的GPU设备上完成分配。
• ‌P2P集群映射‌：所有处于同一个P2P连通集群内的其他GPU，都会自动获得这块内存的对等访问映射，无需开发者手动调用P2P启用接口。
• ‌非P2P场景降级‌：如果GPU之间不支持P2P直连，统一内存分配会自动降级到主机零拷贝内存上，这种模式下GPU访问数据需要直接走PCIe链路读取主机内存，会带来明显的性能开销。

🔍 Pascal+ 按需分页架构特性

• ‌全处理器按需可见‌：不再强制依赖P2P硬件能力，无论GPU是否处于同一个P2P连通集群，统一内存都可以被任意处理器按需访问。
• ‌灵活调度优化‌：开发者可以通过cudaMemPrefetchAsync显式预取、cudaMemAdvise内存提示接口，精准引导系统的内存调度行为，避免不必要的数据迁移开销。

⚠️ 多GPU场景使用注意事项

• Pre-Pascal架构下非P2P场景的零拷贝模式，会让GPU访问数据的延迟显著上升，带宽远低于直接访问显存，不适合高频访问的计算密集型场景。
• Pascal+架构下如果完全不使用预取和提示，多GPU跨设备访问统一内存时会触发大量缺页中断，反而会拖慢整体并行性能。
• 多GPU共享只读数据时，配合cudaMemAdviseSetReadMostly提示可以让每个GPU本地生成只读副本，彻底消除跨设备数据迁移开销。

🚀 POWER9 异构内存特性解析

📋 核心内存一致性能力

• ‌全处理器动态共享‌：CPU通过malloc、new动态分配的普通内存，可以在所有处理器之间实现硬件级的一致性访问。
• ‌无需固定内存‌：GPU核函数可以直接读取CPU普通动态分配的内存，不需要提前将内存标记为pinned固定内存，大幅降低异构编程的门槛。
• ‌硬件级地址翻译‌：CPU和GPU的MMU单元可以直接交互，虚拟地址到物理地址的转换完全在硬件层面完成，不需要驱动介入做额外的地址映射转换。

🔍 数据迁移特性

• 该架构下不会自动触发按需数据迁移，开发者必须手动调用cudaMemPrefetchAsync接口完成数据迁移操作。
• 手动迁移的性能会低于传统CUDA统一内存的自动迁移性能，整体带宽表现略逊于原生UM方案。

💡 HMM 技术背景

NVIDIA联合行业伙伴正在推进HMM（异构内存管理）技术在Linux内核中的通用落地，这套软件实现方案可以让更多硬件平台获得类似POWER9的跨设备内存透明访问能力，打破硬件平台的兼容性限制。

06- 并发

课程议程（AGENDA）

• 并发——动机（Concurrency - Motivation）
• 锁页内存（Pinned Memory）
• CUDA 流（CUDA Streams）
• 拷贝与计算的叠加（Overlap of Copy and Compute）
• 使用场景：向量计算 / 视频处理流水线（Vector Math / Video Processing Pipeline）
• 流的额外考量（Additional Stream Considerations）
• 托管内存下的拷贝-计算叠加（Copy-Compute Overlap with Managed Memory）
• 多 GPU 并发（Multi-GPU Concurrency）
• 其他并发场景：内核并发、主机/设备并发（Kernel Concurrency, Host/Device Concurrency）
• 延伸学习（Further Study）
• 课后作业（Homework）

这一讲的核心主题并发（Concurrency） ，简单说就是：让 GPU 上的不同操作（拷贝数据、跑计算、不同 kernel、多卡）同时进行，而不是傻等一件做完再做下一件，从而把硬件利用率提上去。

几个你马上会认识的关键概念就在提纲里：

• Pinned Memory（锁页内存） ：用在 CPU 端加速向 GPU 拷贝的固定内存，是"拷贝和计算能重叠"的前提之一；
• CUDA Streams（流） ：把不同任务放进不同流，让它们彼此并行跑，是实现并发的基础工具；
• Copy-Compute Overlap：数据拷贝和计算在不同流里错开，让 GPU 一边搬数据一边算，效率大增。

固定内存（pinned memory)

锁页（不可分页）内存（PINNED / NON-PAGEABLE MEMORY）

锁页内存带来：

• 更快的主机(CPU) ↔ 设备(GPU) 拷贝
• 可异步于 CPU 的 memcpy（拷贝不被 CPU 阻塞）
• 可异步于 GPU 的 memcpy

用法：

• cudaHostAlloc
  

  cudaFreeHost
  

  • 用它替代普通的 malloc/free 或 new/delete

• cudaHostRegister
  

  cudaHostUnregister
  

  • 对已分配好的普通内存（比如用 malloc 申请的）事后"钉住"（pin）

影响（Implication）：

• 锁页内存本质上从主机虚拟（可分页）内存中被移除了（即这部分内存不再参与操作系统的分页交换）

这一页在讲：为什么"锁页内存"是实现并发的关键铺垫。

普通的 CPU 内存是可分页的（pageable） ——操作系统可能把其中某些页换到磁盘（虚拟内存/页交换）。当你把这样的数据拷到 GPU 时，底层可能先要把数据换回物理内存，导致慢且不可预期；而且这种拷贝往往是同步的，堵住 CPU。

锁页内存（pinned）就不同：

• 用 cudaHostAlloc 分配（或 cudaHostRegister 把已有 malloc 内存"钉住"），把这块内存永久锁定在物理内存里、不参与分页交换；
• 好处就是这页列的三个：拷贝更快、可以异步于 CPU（CPU 不用傻等拷贝完成）、可以异步于 GPU（一边搬数据一边让 GPU 干别的）——这正是下一步 Stream 重叠（copy + compute） 的基础。

代价：锁页内存占用物理 RAM 且不可换出，所以不要滥用——只把需要反复/高速拷贝的数据 pin 起来。

锁页内存就是"钉在物理内存里、不参与分页交换"的主机内存，让 CPU↔GPU 拷贝更快且能异步，为 CUDA 流做计算/拷贝重叠打基础。

CUDA Stream

这里的kernel 指 cuda kernel 函数， D2H device to host H2D host to device

流与异步 API 概览（STREAMS AND ASYNC API OVERVIEW）

默认 API 的行为：

• kernel（内核）启动异步于 CPU（下发后 CPU 不等待）
• cudaMemcpy（D2H、H2D）会阻塞 CPU 线程（拷贝时 CPU 停住等它）
• CUDA 调用被驱动串行化到默认流（default stream） 中顺序执行

有了 Streams 和异步函数，则能：

• cudaMemcpyAsync（D2H、H2D）异步于 CPU（拷贝时 CPU 不阻塞）
• 能让一个 kernel 和一个 memcpy 并发执行
• 在大多数 GPU 上，两个方向的拷贝（D2H、H2D）可以同时进行

流（Stream）= 一串按照"提交顺序（issue-order）"在 GPU 上执行的操作

• 不同流里的操作可以交错（interleaved）执行
• 来自不同流的一个 kernel 和一个 memcpy 可以被重叠（overlapped）

这一页把"为什么要用 CUDA Stream"讲透了：

默认情况下（不用流）的痛点：

• kernel 启动虽然是异步的，但 cudaMemcpy（D2H/H2D 的同步版）会堵住 CPU——拷贝数据时 CPU 只能干等；
• 所有调用都被塞进同一个默认流，一条一条串行执行，GPU 只能在"先拷贝→再算→再拷贝"里来回，空闲时间很多。

用了 Stream（流）和异步 API 之后：

1. cudaMemcpyAsync 让拷贝不再阻塞 CPU（CPU 可以继续发别的命令、做别的活）；
2. 把不同的操作放进不同的流，它们就能在 GPU 上并发/重叠——比如"流 A 在算某批数据，流 B 同时在拷下一批数据"，GPU 的搬运引擎和计算单元都不闲着。

关键概念——流：一个流里的操作按提交顺序执行；但不同流之间可以交错、重叠。所以"并行"= 把相互独立的操作拆到不同流里，让 GPU 同时干多件事。

流 = 一条按序执行的命令队列；把独立的拷贝和计算放进不同流，就能让它们并发/重叠，避免 GPU 在不同环节来回空等。 （而异步 memcpy 是让 CPU 不被拷贝阻塞的关键。）

流语义（STREAM SEMANTICS）

1. 提交到同一个流的两个操作，按提交顺序（issue-order）执行。

• 操作 B 是在操作 A 之后提交的，B 必须等 A 完成之后才会开始执行。

2. 提交到两个不同流的操作，CUDA 不规定它们之间的任何顺序。

• 提交到流 1 的操作 A，可以在操作 B（提交到流 2）之前、之中或之后执行——全看硬件和调度。

（注：这里的"操作"，通常指 cudaMemcpyAsync 或一次 kernel 调用；更宽泛地说，是绝大多数带 stream 参数的 CUDA API，包括流回调。）

这一页用两个语义规则讲清了"流"最核心的约定：

规则 1（同流内串行） ：同一个流里的命令是按你提交的先后顺序一条条来的，后面的必须等前面的完成。所以同流内部天然有序、安全——你不用担心并发错乱。

规则 2（跨流无保证顺序） ：不同流之间，CUDA 不承诺任何顺序——A 在 B 前、后、或同时跑都"合法"。这正是并发的来源：正因为不同流之间没有先后约束，GPU 才可以把它们重叠/并行执行。

这对写代码的含义：

• 要保证顺序的（比如"先拷数据再算"）→ 放进同一个流，靠规则 1 保证。
• 想并行/重叠的（比如"拷下一批"和"算这一批"互相独立）→ 放进不同流，靠规则 2 让它们重叠。
• 如果跨流之间有数据依赖（比如流 2 要用流 1 算的结果），你得自己加同步（如 cudaStreamSynchronize、event），CUDA 不会替你管。

同一流内按提交顺序串行（保证有序安全）；不同流之间无顺序约束（所以能并发），但跨流依赖要自己同步。 这就是 CUDA 并发编程的结构化基础。

流创建 + 拷贝/计算重叠（STREAM CREATION AND COPY/COMPUTE OVERLAP）

实现重叠的两个前提（Requirements）：

• D2H 或 H2D 的拷贝必须从锁页内存（pinned memory） 发起
• kernel 和 memcpy 必须放在不同、且非 0 的流里

代码示例：

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);          // 创建流1
cudaStreamCreate(&stream2);          // 创建流2

cudaMemcpyAsync( dst, src, size, dir, stream1 );   ⎫
kernel<<<grid, block, 0, stream2>>>(...);          ⎬ 这两个操作
                                                    ⎭ 可以"重叠"

cudaStreamQuery(stream1);            // 查询：流1是否空闲（不阻塞）
cudaStreamSynchronize(stream2);      // 同步：强制让 CPU 线程等待流2完成
cudaStreamDestroy(stream2);          // 销毁流2

这一页给的是"拷贝和计算重叠"的标准模板：

两个必要条件（这是关键）：

1. 要拷贝的数据得是锁页内存（pinned） ——上一页讲过，只有锁页内存才能让 memcpy 异步、才能和 kernel 重叠；
2. kernel 和 memcpy 要放进两个不同的、非 0 的流（默认流 0 会和其它流有隐式同步，会破坏重叠），本例就是 stream1（拷数据）和 stream2（跑 kernel）。

关键 API 动作：

• cudaStreamCreate(&s) —— 创建流；
• 把 cudaMemcpyAsync 放进 stream1，把 kernel<<<..., stream2>>> 放进 stream2——因为它们在不同流，GPU 可以把"拷下一批数据"和"算这批数据"重叠起来；
• cudaStreamQuery(s) —— 非阻塞地"查"流是否空闲（可用来轮询）；
• cudaStreamSynchronize(s) —— 阻塞：让 CPU 线程等该流里的全部操作做完（确保用结果前数据已就绪）；
• cudaStreamDestroy(s) —— 用完释放流。

把"拷贝数据"和"跑内核"分别放进两个流，配合锁页内存 + 异步 API，就能在 GPU 上重叠 data copy 和 compute；再用 Query/Synchronize 控制等待、用完后销毁流。 这是 CUDA 并发性能优化的标准写法。

📊 CUDA流执行时序示例图

📋 图例基础说明

图中右侧标注了所有元素的含义：

• ‌K‌：Kernel，代表GPU上执行的核函数计算任务
• ‌M‌：Memcopy，代表主机与设备之间的内存拷贝任务
• ‌Integer‌：Stream ID，代表任务所属的CUDA流编号
  底部的绿色箭头代表时间推进方向，所有橙色色块的长度对应任务的执行耗时。

🔍 5种流调度时序解析

图中展示了5种不同任务分配到不同流后的并行执行效果：

1. ‌K1,M1,K2,M2‌
   流1执行K1，流2并行执行M1，之后流1执行K2，流2并行执行M2，实现了计算与拷贝的两次重叠。
2. ‌K1,K2,M1,M2‌
   流1连续执行K1、K2两个核函数，流2先后执行M1、M2两次内存拷贝，两段任务完全并行，最大化重叠时长。
3. ‌K1,M1,M2‌
   流1执行K1，流2先后执行M1、M2，K1执行结束后M1才启动，仅实现部分重叠。
4. ‌K1,M2,M1‌
   流1执行K1，流2先执行M2再执行M1，M2与K1完全并行，M1在K1结束后启动。
5. ‌K1,M2,M2‌
   流1执行一个长耗时的K1核函数，流2先后执行两次M2内存拷贝，两次拷贝完全被K1的执行时长覆盖，实现了全重叠。

💡 核心设计原理

这组示例直观展示了CUDA流的核心调度规则：

• 同流内的任务严格按提交顺序串行执行
• 不同流的独立任务可以利用GPU的计算引擎和DMA拷贝引擎并行运行
• 合理分配任务到不同流，能够让核函数计算和内存拷贝相互重叠，大幅减少GPU空闲时间，提升硬件整体利用率。

前提（assumes algorithm decomposability）：假设算法可分解——任务能切成互相独立的块，块和块之间没有依赖，才能流水线化。

❌ non-streamed（不用流）——串行三段：

cudaMemcpy(d_x, h_x, size_x, cudaMemcpyHostToDevice);  // 1. 全部拷上去
Kernel<<<b, t>>>(d_x, d_y, N);                          // 2. 全部算完
cudaMemcpy(h_y, d_y, size_y, cudaMemcpyDeviceToHost);   // 3. 全部拷回来

对应图上方三个大色块：H→D 拷贝 → kernel → D→H 拷贝，严格排队，任何时刻 GPU 只干一件事，拷贝时计算单元闲着，计算时拷贝引擎闲着。

✅ streamed（用流）——切块 + 交替入流：

for (int i = 0; i < c; i++){                    // 数据切成 c 块
  size_t offx = (size_x/c)*i;                   // 第 i 块的偏移
  size_t offy = (size_y/c)*i;
  cudaMemcpyAsync(d_x+offx, h_x+offx, size_x/c, cudaMemcpyHostToDevice, stream[i%ns]);   // 拷第 i 块上去
  Kernel<<<b/c, t, 0, stream[i%ns]>>>(d_x+offx, d_y+offy, N/c);                          // 算第 i 块
  cudaMemcpyAsync(h_y+offy, d_y+offy, size_y/c, cudaMemcpyDeviceToHost, stream[i%ns]);   // 拷第 i 块回来
}

图里那排 0 1 0 1 0 1 0 就是精髓：

• 每块数据的“拷→算→拷回”放进同一个流（保证块内顺序正确）；
• 相邻块交替放进 stream 0 / stream 1（i%ns）；
• 于是形成图下方的阶梯状重叠：stream 0 在算第 0 块时，stream 1 正在拷第 1 块；stream 0 拷回第 0 块时，stream 1 正在算第 1 块……拷贝和计算互相填满对方的空闲，总时间 ≈ 串行版的 1/2 左右。

“大任务切块、块内同流保序、块间异流重叠”——这就是 CUDA 流水线的标准套路。

底部那行 "Similar: video processing pipeline"（同理：视频处理流水线）——视频逐帧解码→处理→编码，天然就是这种可分块任务，一模一样的写法。

默认流（DEFAULT STREAM）

什么是默认流：

• 指定流（或显式传 0）的 kernel / cudaMemcpy，用的都是默认流（stream 0） 。

默认流的"legacy"（旧版）行为——会强制同步（synchronizing on the device）：

• 默认流里的那项操作开始前，之前发出的所有设备活动都必须先完成；
• 默认流里的那项操作结束后，之后发出的所有设备活动都要等它结束才能开始。
• 所有宿主线程（host threads）共享同一个默认流。

为什么要避开：

• 在复杂并发场景下，建议避免使用默认流——因为它这种"必须等前面做完、后面做完全部等它"的行为，等于一道全局屏障，会把你想重叠的流操作强行串行化。

可以改造它：

• 通过编译选项把默认流从 legacy 改成"普通流"：

  nvcc --default-stream per-thread ...
  

• 这样每个宿主线程会有自己独立的默认流，不再全局共享、不再强制同步。

直白理解：

默认流（stream 0）是 "老爷流" ——它一旦干活，GPU 上别的活都得停下等它，它做完了别人才能动。所以上一页例子特意强调 kernel 和 memcpy 要放"非 0 流"，就是为了绕开这道全局屏障，让拷贝和计算能真正重叠。

如果你有多个宿主线程（多线程程序），--default-stream per-thread 能让每个线程各拿一个普通流，互不干扰，也不会互相强制同步。

一句话：默认流=全局同步的拦路虎，复杂并发下要么避开它（进普通流），要么用 --default-stream per-thread 把它"降级"成普通流。

cudaLaunchHostFunc()（流回调）

• 允许定义一个宿主（CPU）函数，把它排进某个 CUDA 流。
• 遵循流的语义：这个函数不会提前执行——只有流执行到它那个位置时才会被调用。换句话说，它排在哪个操作后面，就等前面的活全干完才执行。
• 执行时是用 GPU 驱动额外拉起来的一个线程来做这件事的（不是主线程）。
• 有局限：这个函数里不能用任何 CUDA 运行时 API 调用（也不能发 kernel） ——因为它跑在驱动派的线程里，环境受限。
• 用途：把"等 GPU 结果就绪后才做的 CPU 收尾工作"延迟到合适时机，让 CPU 主线程先去干别的。
• cudaLaunchHostFunc() 是新版，取代了旧的 cudaStreamAddCallback()。

直白理解：

平常你写 cudaStreamSynchronize(stream) 是让 CPU 干等着流干完。而 cudaLaunchHostFunc 是反过来的——把它当成"排进流里的一个 CPU 小任务"：

• 你正常发 kernel、memcpy，最后把这个回调函数也 cudaLaunchHostFunc() 进同一个流；
• GPU 按流顺序跑完前面的活，轮到这个回调时，驱动派个线程在 CPU 上执行它（比如把结果存盘、算个统计、打个日志）；
• 主线程不用傻等，可以继续提交别的流的工作。

限制要记住：这个回调里不能再调 CUDA API 或发 kernel，只能做纯 CPU 的收尾工作——因为它的执行线程不是普通主线程环境。

流回调 = "把 CPU 收尾活排进流里，GPU 干完自动触发"，新版 API 是 cudaLaunchHostFunc()，取代老接口 cudaStreamAddCallback()。

📊 托管内存下拷贝计算重叠说明

🎯 核心主题

这页内容讲解了‌CUDA托管内存（Managed Memory）结合按需分页机制‌，实现数据拷贝与GPU计算重叠的相关特性与注意事项。

📋 基础实现逻辑

它沿用了传统异步流实现计算与拷贝重叠的经典模式，核心区别是用cudaMemPrefetchAsync替代了传统的cudaMemcpyAsync：

• 托管内存基于统一地址空间，CPU和GPU可以直接访问同一份虚拟地址的数据
• 借助按需分页机制，系统会在数据真正被访问时自动完成页迁移，无需开发者手动管理主机与设备内存的指针映射
• 流的语义会严格保证所有必要的内存迁移操作按照正确的顺序执行，避免出现数据访问依赖错误

⚠️ 关键性能特性

cudaMemPrefetchAsync相比普通的异步拷贝接口有特殊的性能表现：

• 它需要额外完成CPU和GPU端的页表更新工作，执行的内部工作量远大于普通的cudaMemcpyAsync
• 这会导致该API的主机侧返回耗时明显更长，甚至可能在执行时间线上引入意料之外的空闲间隙，破坏原本预期的计算与拷贝完全重叠效果
• 它的运行表现会根据流的负载状态发生变化：反直觉的是，处于忙碌状态的流反而可能带来更优的整体吞吐量

🔗 参考延伸资源

页面中提供了NVIDIA官方深度优化指南链接，可进一步了解托管内存的性能调优细节：

https://devblogs.nvidia.com/maximizing-unified-memory-performance-cuda/

📊 CUDA cudaEvent 核心特性解析

🎯 基础定义与状态规则

cudaEvent 是CUDA编程中用于在流中插入标记的核心对象，它的状态流转有明确的定义：

• ‌recorded（已记录）‌：当事件被调用提交到流中的瞬间，就进入已记录状态
• ‌completed（已完成）‌：当流的执行进度推进到该事件被标记的位置时，事件才会变为完成状态
  它最核心的基础用途是对GPU上的任务进行精确计时。

🔍 扩展能力与注意事项

除了基础计时功能外，它还能支撑更复杂的CUDA并发场景：

• 可以搭配cudaStreamWaitEvent()接口，实现不同流之间的执行依赖编排，灵活构建复杂的异步流水线
• 存在明显的使用边界：如果直接用它统计主机侧的CPU活动耗时，或是在多流高度并发的复杂场景下使用，很容易得到不符合预期的计时结果，无法准确反映真实的执行时长。

CUDA 多GPU

📊 CUDA多GPU设备管理核心说明

🎯 基础定位说明

多GPU设备管理接口并不是OpenMP、MPI这类并行框架的替代品，它更偏向底层的GPU资源调度能力，通常会和上层并行框架搭配使用，实现多卡场景下的算力协同。

📋 核心基础API能力

应用程序可以通过官方提供的原生接口，自主查询和选择可用的GPU设备：

cppcudaGetDeviceCount(int *count)
cudaSetDevice(int device)
cudaGetDevice(int *device)
cudaGetDeviceProperties(cudaDeviceProp *prop, int device)

• cudaGetDeviceCount：查询当前系统中可用的GPU总数量
• cudaSetDevice：指定当前主机线程后续操作绑定的目标GPU
• cudaGetDevice：查询当前线程当前绑定的GPU编号
• cudaGetDeviceProperties：获取指定GPU的硬件属性参数，比如算力、显存大小等信息

🔍 多卡灵活调度特性

CUDA多卡管理支持两种非常实用的资源共享模式：

• 多个主机线程可以同时共享同一个GPU设备，合理分配不同任务到不同线程，充分利用单卡的算力资源
• 单个主机线程也可以通过调度接口同时管理多个GPU设备，实现轻量场景下的多卡任务编排

⚡ 进阶跨卡操作能力

针对单线程管理多设备的场景，还可以使用专属接口实现高效跨卡交互：

cppcudaSetDevice(i) to select current device
cudaMemcpyPeerAsync(...) for peer-to-peer copies

通过cudaMemcpyPeerAsync可以直接在两个GPU的显存之间完成异步点对点数据拷贝，无需经过CPU内存中转，大幅降低跨卡数据传输的开销。

📊 多GPU流核心特性解析

🎯 基础关联规则

CUDA中的流和cudaEvent对象会自动和当前线程绑定的GPU建立隐式关联，每一个GPU设备都拥有自己独立的默认流，不同设备的流之间完全隔离。

• 如果你向一个和当前设备没有关联的流中提交核函数启动指令，该操作会直接执行失败，这是多流多卡场景下最常见的入门错误。

🔍 跨设备同步能力

CUDA提供了专门的接口实现跨设备的流协同：

• cudaStreamWaitEvent() 可以让当前设备的流等待另一个完全不同设备上的事件完成，实现跨设备的执行依赖编排
• cudaEventQuery() 可以随时查询位于其他设备上的事件是否已经执行完成，无需主机线程阻塞等待

这段代码分别在两张GPU上创建独立流，最终两个核函数可以在不同的GPU硬件上完全并行执行，充分利用多卡的总算力。

📊 多GPU设备间数据拷贝核心说明

🎯 核心传输能力

只要系统的硬件拓扑支持，GPU之间就可以通过PCIE或者NVLink这类高速互联 fabric 直接完成数据拷贝，全程不需要经过CPU内存中转。
这种传输模式就是GPUDirect P2P点对点传输，相比传统的“GPU→CPU→GPU”中转拷贝，能大幅降低跨卡数据传输的延迟和带宽损耗。

📋 前置配置规则

跨卡直连拷贝不能直接使用，必须先完成显式的对等关系配置：

• 首先需要把目标设备加入到同一个对等关系组（clique）中
• 如果需要双向传输，必须分别在两个设备上开启双向的对等访问权限，只配置单向会导致反向拷贝失败
• 完成配置后，两张GPU之间的内存拷贝就不会再经过系统内存缓冲区中转，实现真正的直连传输

⚠️ 关键使用限制

对等关系组（clique）中可支持的对等设备数量是存在硬件上限的，不同架构的GPU支持的最大对等设备数量不同，实际使用前需要提前查询硬件规格，避免超出上限导致配置失败。

📊 CUDA其他并发场景说明

🎯 主机与设备执行并发

这是CUDA最基础的异步并行模式，核函数在GPU上启动后，主机侧的CPU代码不需要阻塞等待核函数执行完成，两者可以同时并行运行。

只要核函数启动后没有强制同步操作，CPU就可以立刻执行后续的业务逻辑，充分利用主机和设备的算力资源，避免单一侧硬件处于空闲等待状态。

🔍 多核函数并发执行

通过将不同核函数提交到不同的独立流中，就具备了让多个核函数在同一张GPU上同时执行的可能性：

⚠️ 实际落地的关键约束

在真实的GPU运行环境中，同设备上的多核函数并发执行很难稳定复现，需要满足多个严格条件：

• ✅ 核函数自身的硬件资源占用率要相对较低，不能完全占满GPU的计算单元、共享内存等资源
• ✅ 核函数的执行时长要足够长，避免单个核函数快速执行完成，还没等第二个核函数调度上来就已经结束
• ✅ 单张GPU支持的同时并发核函数数量存在明确的硬件上限，不同架构的GPU上限数值不同
• ✅ 这种模式的整体效率，远低于用单个核函数充分占满GPU全部算力的执行效率，属于特殊场景下的优化手段，不建议作为常规优化方案使用。

📊 CUDA流优先级核心说明

🎯 基础优先级特性

CUDA流支持为流设置可选的优先级属性，该属性仅会对‌并发执行的核函数‌产生调度影响：

• GPU的块调度器会优先调度高优先级流中的核函数块，再调度低优先级流中的核函数块
• 当前主流CUDA实现中仅提供高低2档优先级，该规格属于非强制的实现细节，不同硬件版本可能存在差异
• 当前实现不会触发核函数块的抢占操作，同样该行为也属于非强制的实现细节

⚠️ 实际使用注意事项

• 不要依赖固定的优先级数值，必须先调用cudaDeviceGetStreamPriorityRange查询当前设备支持的优先级范围，不同架构GPU的优先级数值可能不同
• 优先级调度仅在多个核函数并发运行时生效，单一流独占GPU时优先级属性不会带来任何性能收益
• 低优先级流的任务不会被高优先级流直接抢占，仅会在调度顺序上获得优先，无法实现硬实时的任务抢占效果

CUDA 图

📊 CUDA Graphs 核心概述

🎯 基础定位与能力

CUDA Graphs 是CUDA 10版本推出的全新特性，它支持将流中的一系列工作任务定义为完整的有向图结构，可承载的任务类型覆盖核函数、内存拷贝操作、回调函数、主机函数甚至嵌套子图。

🔍 图的核心结构规则

• 图中的每一项独立工作任务都对应图中的一个节点
• 支持自定义节点之间的依赖关系，比如指定某3个节点全部执行完成后，下一个节点才可以开始运行
• 这些任务依赖关系本质上就是图结构中的边，完整构成了任务调度的有向无环图

💡 执行与使用特性

• 完成图的定义之后，只需要将图提交到流中就可以启动执行
• 图一旦定义完成就可以被反复复用，大幅降低重复提交大量细碎任务带来的主机侧开销
• 它同时提供了两种构建方式：手动逐个定义节点的手动创建方法，以及通过流捕获自动生成图的便捷捕获方法。

07- GPU 性能分析

📊 CUDA性能编码核心目标回顾

🎯 核心主题说明

这页内容是CUDA高性能编程的顶层核心目标总结，梳理了GPU性能优化最核心的两大方向，是CUDA开发中所有性能调优动作的总纲领。

📋 高效利用内存子系统

这是CUDA性能优化的第一核心维度，重点围绕GPU内存体系的特性最大化内存访问效率：

• ‌全局内存高效利用‌：核心要求是实现合并访问（coalesced access），让多个线程的内存访问请求合并为连续的大块传输，大幅降低全局内存的访问开销
• ‌智能运用内存层次结构‌：充分利用GPU多层级的内存资源，包括共享内存、常量内存、纹理内存以及各级硬件缓存，根据不同场景选择最合适的内存类型，减少不必要的高延迟访问。

🔍 暴露足够并行度隐藏延迟

通过充分释放并行算力，让GPU硬件始终处于高负载状态，用计算任务掩盖各类操作的延迟，最终打满硬件算力：

• 合理规划线程与块的分配策略，提供足够多的并行任务
• 提升GPU占用率（Occupancy），让更多的线程块同时驻留在SM中
• 优化单线程的计算负载，避免线程任务过轻或过重
• 保障整体执行效率，减少硬件流水线的空转等待。

📊 分析驱动优化闭环流程

🎯 核心定位

这是一套面向GPU/高性能计算场景的标准化性能优化闭环方法论，核心逻辑是完全基于实际性能数据而非主观猜测来推进优化，避免无效调优。

🔍 四步闭环执行逻辑

整个流程形成一个可循环迭代的完整闭环，每一步都有明确的执行目标：

1. ‌Profile（性能剖析）‌

   借助性能分析工具完整采集程序运行的全量性能数据，精准记录每一个环节的耗时、资源占用情况，拿到真实的运行画像。

2. ‌Determine Limiter（定位瓶颈）‌
   基于采集到的性能数据，从众多运行环节中精准识别出真正限制性能提升的核心瓶颈点，避免盲目优化非关键路径。

3. ‌Study, Reflect, Learn, Inspect（复盘分析）‌
   针对定位到的性能瓶颈深入研究，结合硬件架构特性、算法逻辑细节，梳理出可行的优化方向，排查当前代码中存在的低效问题。

4. ‌Optimize（落地优化）‌
   按照梳理好的优化方案针对性修改代码，完成性能调优动作，之后重新回到性能剖析环节，开启下一轮迭代优化。

💡 方法论核心优势

这套流程可以彻底避免“凭感觉优化”的误区，每一轮优化都瞄准真实的性能短板，用最小的开发成本获得最大的性能收益，是工业级高性能程序优化的标准实践思路。

📊 分析驱动优化执行流程

🎯 核心定位

这是一套面向高性能计算场景的标准化性能瓶颈定位与优化执行流程，承接上一轮的分析驱动优化闭环，给出了可直接落地的分支判断逻辑，帮助开发者精准定位程序的性能短板。

🔍 分步判断逻辑

整个流程从性能剖析完成后启动，通过两层判断精准分流到对应的瓶颈分析方向：

1. ‌第一层判断：是否为内存绑定‌
   如果判定程序是内存访问瓶颈，直接进入内存限制专项分析，针对性排查全局内存访问、内存层级利用的低效点。
2. ‌第二层判断：是否为计算绑定‌
   若程序不属于内存绑定场景，继续判断是否为计算瓶颈；如果是计算绑定，进入计算限制专项分析，排查算力利用不足的问题。
3. ‌兜底分支：延迟类瓶颈‌
   若程序既不属于内存绑定也不属于计算绑定，说明性能瓶颈来自指令调度、依赖等待等延迟类问题，进入延迟专项分析。

💡 优化落地环节

完成对应方向的瓶颈分析后，所有分支的结论都会汇总到统一的优化环节，针对性输出调优方案，优化完成后会重新回到最开始的性能剖析步骤，开启下一轮迭代优化，持续推进性能提升。

这套流程可以彻底避免盲目优化的问题，每一次调优动作都精准瞄准真实的性能短板，大幅提升优化效率。

📊 性能瓶颈顶层分类说明

🎯 三大核心性能限制器定义

这部分内容明确了高性能程序中三类最核心的性能瓶颈判定标准：

• ‌内存绑定（Memory Bound）‌：当程序实测的内存子系统性能达到或接近硬件理论最大值，也就是完全打满了内存总线带宽时，就属于内存绑定场景。
• ‌计算绑定（Compute Bound）‌：当程序的计算指令吞吐率达到或接近硬件理论最大值，算力资源被充分占满时，就属于计算绑定场景。
• ‌延迟绑定（Latency Bound）‌：如果程序既不属于内存绑定，也不属于计算绑定，就可以判定为延迟绑定场景，这类场景的性能瓶颈通常来自指令调度等待、依赖链过长等延迟类问题。

🔍 分析驱动优化的落地逻辑

分析驱动的优化方法论，会在优化的第一阶段就通过上述判定标准精准定位程序的瓶颈类型，以此来指导后续的代码重构工作，避免盲目优化非关键路径，大幅提升调优效率。

⚠️ 关键实践注意事项

• 程序的性能限制类型并不是固定不变的，在整个程序的执行周期中，不同阶段的瓶颈类型可能会发生动态切换。
• 推荐的优化实践是拆分出程序的小代码片段单独分析，比如一次只针对一个核函数做性能剖析，这样可以更精准地定位每个局部环节的真实瓶颈，避免整体分析掩盖局部的性能短板。

📊 计算与内存绑定判定指标说明

🎯 基础信息

这是CUDA性能分析中用于区分程序是计算绑定还是内存绑定的官方指标集合，对应的官方参考文档地址为：
https://docs.nvidia.com/nsight-compute/NsightComputeCli/index.html#nvprof-metric-comparison

⏱️ 延迟类指标

用于判定程序是否属于延迟绑定场景，核心指标为：

• ‌sm efficiency‌：对应指标名 smsp__cycles_active.avg.pct_of_peak_sustained_elapsed，统计SM活跃周期占峰值持续运行周期的占比，反映流多处理器的整体运行效率。

🧠 内存类指标

用于判定程序是否属于内存绑定场景，覆盖GPU全层级内存的利用率统计：

• ‌dram utilization‌：对应指标名 dram__throughput.avg.pct_of_peak_sustained_elapsed，统计显存带宽利用率占峰值持续带宽的比例
• ‌L2 utilization‌：对应指标名 lts__t_sectors.avg.pct_of_peak_sustained_elapsed，统计L2缓存的访问利用率占峰值的比例
• ‌shared utilization‌：对应指标名 l1tex__data_pipe_lsu_wavefronts_mem_shared.avg.pct_of_peak_sustained_elapsed，统计共享内存的访问利用率占峰值的比例

⚙️ 计算类指标

用于判定程序是否属于计算绑定场景，覆盖不同精度算力单元的利用率统计：

• ‌DP Utilization‌：对应指标名 smsp__inst_executed_pipe_fp64.avg.pct_of_peak_sustained_active，统计双精度浮点算力单元的利用率
• ‌SP Utilization‌：对应指标名 smsp__pipe_fma_cycles_active.avg.pct_of_peak_sustained_active，统计单精度浮点算力单元的利用率
• ‌HP Utilization‌：对应指标名 smsp__inst_executed_pipe_fp16.avg.pct_of_peak_sustained_active，统计半精度浮点算力单元的利用率
• ‌TC Utilization‌：对应指标名 sm__pipe_tensor_op_hmma_cycles_active.avg.pct_of_peak_sustained_active，统计Tensor Core张量算力单元的利用率
• ‌Integer Utilization‌：对应指标名 smsp__sass_thread_inst_executed_op_integer_pred_on.avg.pct_of_peak_sustained_active，统计整数运算单元的利用率

这些指标可以直接在Nsight Compute性能分析工具中采集，精准定位程序的性能瓶颈类型，为后续针对性优化提供数据支撑。

📊 内存绑定场景核心说明

🎯 内存绑定的基础定义

程序出现内存绑定，本质是性能受限于内存带宽或者内存访问延迟。这里会把内存延迟绑定的场景统一归类到通用延迟绑定场景中单独处理，当前页面重点聚焦内存带宽绑定的优化方向。

🔍 内存带宽绑定优化核心思路

针对内存带宽绑定的程序，核心优化目标是充分利用GPU的多层级内存子系统，最大化发挥内存层次结构的优势，减少对高延迟显存的直接访问：

• 首先优化全局内存的使用方式，通过合并访问等基础手段降低显存访问开销
• 在存在数据复用的场景下，高效利用L2缓存、共享内存等更高层级的内存资源，把高频访问的数据放到更低延迟的存储位置
• 利用GPU专属的特殊缓存实现“缓存多元化”，包括常量缓存、只读缓存、纹理缓存/纹理内存、表面内存等专用存储资源，针对性适配不同访问模式的数据

📈 性能量化分析方法

对于内存带宽绑定的程序，可以直接统计程序实际运行时的内存吞吐率，将其和硬件的理论峰值带宽做对比，就能直观量化当前程序的内存资源利用效率，精准判断优化空间还有多大。

⚙️ 计算绑定场景核心说明

🎯 计算绑定的基础定义

当程序中某一类特定计算指令/操作的性能，达到或接近负责执行该类指令的硬件功能单元的极限时，程序就进入了计算绑定状态，此时性能瓶颈完全来自算力资源的供给上限。

🔍 核心优化策略

针对计算绑定场景的优化核心分为两个方向：

• 针对性优化当前占满硬件资源的功能单元的使用方式，减少无效的算力浪费，提升该类硬件单元的执行效率
• 也可以尝试将部分计算负载转移到其他类型的硬件功能单元上，分散算力压力，打破单一功能单元的性能瓶颈

📊 性能量化分析方法

如果程序的性能完全由某一类特定计算操作主导，比如单精度浮点乘加运算占绝对主导，就可以直接统计程序实际运行时的该类操作吞吐率，将其和硬件的理论峰值算力做对比，直观量化当前算力资源的利用效率，清晰判断剩余优化空间。

⏱️ 延迟绑定场景核心说明

🎯 延迟绑定的基础定义

当GPU无法通过当前可用的、已暴露的并行工作任务保持满负载运行时，程序就进入了延迟绑定状态，此时性能瓶颈既不是内存带宽也不是计算算力，而是并行任务量不足导致硬件出现大量空转等待。

🔍 核心优化策略

针对延迟绑定场景，核心优化方向就是尽可能暴露更多的并行工作任务，充分填满GPU的流水线，用并行任务掩盖各类操作的延迟：

• ✅ 确保启动足够大规模的线程数量，提供充足的并行基础
• ✅ 提升单线程的处理工作量，比如通过循环遍历输入元素的方式，让单个线程承载更多计算任务
• ✅ 使用“向量加载”机制，让单个线程可以一次性处理多个输入元素，进一步提升单线程的任务密度
• ✅ 尽可能追求最高的GPU占用率，让更多的线程块驻留在SM中，减少硬件的空闲周期

💡 优化补充提示

延迟绑定场景很容易被误判为内存或计算绑定，这类场景下硬件的各类资源利用率都不会达到峰值，通过增加并行度的优化手段，往往可以用极低的代码改动成本获得非常可观的性能提升。

🧵 CUDA Occupancy 核心定义说明

🎯 基础概念定义

Occupancy（GPU占用率）是用来衡量SM（流多处理器）中实际线程负载，和该SM理论峰值/可达到的最大线程负载之间的比值指标，是CUDA性能优化中非常核心的参考维度。

📋 配套工具与优化特性

• CUDA官方提供了专门的occupancy calculator电子表格工具，可以帮助开发者提前预估不同配置下的GPU占用率数值
• 更高的GPU占用率往往是获得更高程序性能的有效路径，充足的活跃线程可以更好地掩盖各类操作的延迟，让硬件流水线保持满负载运行
• 程序最终能达到的实际占用率，会被多个限制因素约束，无法随意达到理论最大值

🔍 核心占用率限制因素

• ‌每个线程使用的寄存器数量‌：可以通过性能分析工具直接统计，也可以在编译阶段提前获取该数值
• ‌每个线程块包含的线程总数‌：线程块的大小配置会直接影响SM中可同时驻留的块数量
• ‌共享内存的使用量‌：每个线程块占用的共享内存大小，会限制SM中可同时容纳的线程块上限

💡 优化实践提示

在实际CUDA开发中，并非盲目追求100%占用率就一定能获得最优性能，需要结合程序的实际瓶颈类型做权衡，比如计算绑定场景下，适当降低占用率换取单线程更多的寄存器资源，反而可能获得更好的算力利用效率。

📊 CUDA示例算法流程说明

🎯 核心算法逻辑

这是CUDA性能优化教程中的典型示例算法流程，完整的运算链路如下：

1. 输入数据为M个长度为L的向量，先对这组向量做平均计算，得到一个长度为L的平均向量
2. 将这个平均向量和一个固定的L×L矩阵执行矩阵向量乘法，得到最终的输出结果向量
3. 整个流程会重复执行N次，全程复用同一个L×L矩阵，使用N组不同的输入向量，最终生成N个对应的结果向量。

🔍 场景设计目的

这个示例是CUDA分析驱动优化流程中的经典教学用例，专门用来演示如何通过性能剖析工具定位程序的瓶颈类型，针对性完成内存、计算或延迟类的优化，非常适合作为CUDA性能调优的入门实践案例。

💡 优化切入点参考

基于这个算法的特性，常见的优化方向包括：

• 利用共享内存缓存固定的L×L矩阵，减少重复的显存访问开销
• 优化向量平均和矩阵向量乘的合并内存访问模式，提升显存带宽利用率
• 调整线程块配置提升GPU占用率，暴露足够的并行度掩盖操作延迟。

Nvidia Nsight compute 性能分析工具， 可以对代码的性能进行分析

分析 性能瓶颈， 重构代码，

连接机器使用 Nvidia Nsight compute ui 工具 在线分析

08-协作组

Cooperative Groups：一个更灵活的线程组同步/通信模型

一眼扫过（At a glance）：

• 可扩展的线程组协作（Scalable cooperation among groups of threads）
• 灵活的并行分解（Flexible parallel decompositions）
• 能跨软件边界组合（Composition across software boundaries）
• 最直接的收益：grid 级同步（grid-wide sync）——也就是让整个网格（grid）里的所有线程能一起同步，这正是第二讲里"跨 block 全局同步"难题的官方解决手段之一。

举例应用（Examples include）：

• Persistent RNNs（常驻型循环神经网络）
• Reductions（归约）
• Search Algorithms（搜索算法）
• Sorting（排序）

普通 CUDA 里你能做的同步其实很有限——__syncthreads() 只能同步一个 block 内部，跨 block 同步需要拆 kernel、原子、fence 等绕来绕去的笨办法。Cooperative Groups（协作组） 就是把"同步/通信"这件事抽象成一个更通用、更灵活的模型：

• 它让你把线程组织成各种层次、各种形状的"组" （thread_block、tile、甚至整个 grid），并且定义了组内如何同步；
• 关键能力是 grid sync（整个 grid 同步） ：把整个 kernel 的所有线程当成一个组，一次同步全体到位——前提是 kernel 得用专门方式启动（后面页会讲 cudaLaunchCooperativeKernel）；
• 靠这套抽象，你就可以写跨层组合的并行代码，比如在 RNN、大归约、搜索、排序这些需要全局协调的算法里，直接做 grid 级协作，不再靠 hack。

Cooperative Groups 是把"线程同步/通信"升级成可跨 block、跨软件边界统一处理的官方框架，招牌能力是 grid 级全局同步，专治第二讲那个"跨 block 同步"痛点。

协作的层次：CUDA 9.0 之前

当时只有这一个同步工具：__syncthreads()——块级（block level）同步屏障

对照表（能同步 ✓ / 不能 ✗）：

协作范围	能否同步
Warp 内	✗（没有专门同步 API，靠天然锁步）
SM 内	✗
GPU（整个 block/grid）	✗
Multi-GPU（多卡）	✗

除了 __syncthreads() 这一个块内屏障，CUDA 9.0 之前没有任何跨 warp、跨 block、跨 GPU 的同步机制。

它在给你看"老底子"有多简陋——在 CUDA 9.0 之前，整个同步工具箱里只有 __syncthreads() 这一个，而且它只能管住自己那个 block：

• warp 内部：没有专门同步 API（warp 里的线程天然按锁步执行，所以没太需要）；
• 跨 warp（SM 内多个 block）：✗ 不行；
• 整个 GPU（grid 里所有 block）：✗ 更不行——这就是第二讲那个"跨 block 全局同步"问题的由来；
• 跨多张 GPU（multi-GPU）：✗ 完全不行。

协作的层次：CUDA 9.0

针对每一种粒度的线程集合，都有对应的组对象：

// 当前"合并(coalesced)"的那一组线程（warp 内活跃的统一线程）
auto g = coalesced_threads();

// warp 大小的线程组（把 block 切成 32 线程的 tile）
auto block = this_thread_block();
auto g = tiled_partition<32>(block);

// CUDA 线程块（整个 block）
auto g = this_thread_block();

// 覆盖整个设备的 grid（整个网格的所有 block）
auto g = this_grid();

// 跨多张 GPU 的多网格
auto g = this_multi_grid();

所有 Cooperative Groups 的功能都在 cooperative_groups:: 命名空间里。

CUDA 9.0 前只有 __syncthreads()（只管一个 block），而 CUDA 9.0 的 Cooperative Groups 把协作层次完整补齐：

协作范围	CUDA 9.0 前的 API	CUDA 9.0 的 Cooperative Groups
warp 内小tile	✗ 无	tiled_partition<32>(block) / coalesced_threads()
thread block	__syncthreads()	this_thread_block()
整个 grid	✗ 无	this_grid()
多 GPU grid	✗ 无	this_multi_grid()

几个要点：

• 每个层级返回一个 "组对象（group）" ，后面就可以对组对象调用 .sync() 等协作方法，也就是用组代替"脚本式的同步调用"来管理协作；
• coalesced_threads() 表示当前做事的一组活跃线程（比如条件分支里只有一部分线程进来时，它们构成一个 coalesced 组）；
• tiled_partition<32>(block) 是把 block 按大小 32（=一个 warp）切成小片，得到 warp 级组；
• this_grid() 就是grid 级同步的入口，this_multi_grid() 更进一步到多卡；
• 全部放在 cooperative_groups 命名空间下，避免污染。

CUDA 9.0 的 Cooperative Groups 用统一的"组对象"把 warp→block→grid→multi-grid 全部层次都覆盖了，你不再局限于 __syncthreads()，最吸引人的就是能 this_grid() 做 grid 全局同步。

线程组（THREAD GROUP）

• 基类型：线程组是一个基类型（base type） ，具体实现取决于它怎么被构造（this_thread_block() / tiled_partition<> / this_grid() 等返回的都是它的具体形态）。
• 统一：它把各种不同的组类型统一成一种通用的、集体的（collective）线程组。
• 为什么需要：我们要给 CUDA 编程模型扩展出"句柄（handle）" ，用它来表示——哪些线程能互相通信/同步。

它能表示的组类型（从左到右由小到大）：

• Thread Block Tile（块内小块，如 warp 切片）
• Thread Block（线程块）
• Coalesced Group（合并组，条件分支下活跃线程）
• Grid Group（整个网格）
• Multi-Grid Group（多网格）

（层级图：Thread Group 在最外层作为总基类型，下面展开这些具体形态）

这页讲什么：

一句话概括—— "线程组(Thread Group)"是 Cooperative Groups 的一等公民抽象：

• 它把前面那几种不同粒度的组（warp tile、block、coalesced、grid、multi-grid）统一收拢成一个类型，cuda::cooperative_groups::thread_group；
• 这些组对象就是"句柄" ，用来回答一个问题： "我，作为这个组的一员，能跟谁同步、能跟谁通信？" ——组的边界划定了协作的边界；
• 有了这个统一基类型，你就能用同一套写法（组对象方法）去操作不同粒度的协作，不用每种都记不同的 API。

线程块（THREAD BLOCK）

它是启动时那个线程块的隐式组（implicit group of all the threads in the launched thread block），也就是整个 block 的全体线程。

它实现了和 thread_group 相同的（通用）接口：

void sync();              // 同步组内所有线程（等价于 __syncthreads()）

unsigned size();          // 组内线程总数

unsigned thread_rank();   // 当前调用线程在组内的编号，范围 [0, size)

bool is_valid();          // 该组是否违反了任何 API 约束（返回是否有效）

以及 thread_block 专属的额外函数：

dim3 group_index();       // 本 block 在 grid 中的三维索引（= 原来的 blockIdx）

dim3 thread_index();      // 本线程在 block 内的三维索引（= 原来的 threadIdx）

它在给你演示 "线程组"统一模型实际长什么样——以最常见的 thread_block 为例，把 API 摆出来：

• 一个 thread_block 组 = 你启动 block 时的全部线程（隐含存在，this_thread_block() 直接拿到）；
• 它实现了通用 thread_group 的四个方法，其中最核心的是 sync()——作用等价于 __syncthreads()，但现在是"组对象方法"的写法；
• size() 给线程总数、thread_rank() 给出当前线程在组内第几个（在广播、归约这些集体操作里很有用）；
• is_valid() 是一个调试辅助，检查你有没有违反 API 约束（比如某些组方法得所有线程一起调，没满足就会被标记无效）；
• 额外还有 group_index() / thread_index()——其实就是把 blockIdx / threadIdx 收编成了组的方法，统一风格。

关键理解：__syncthreads() 只是"顺手的函数调用"，而 thread_block.sync() 是 "组对象的方法" ——同样的语义，但它是整个统一"组"框架里的一员，后面 grid、tile 等所有层级都共用同一种 sync()/size()/thread_rank() 写法。

程序定义的分解

// 整个 kernel 启动的所有线程
CUDA KERNEL
thread_block g = this_thread_block();       // 全部线程 → block 组

foobar(thread_block g)                      // 传进函数，仍指整个 block 的所有线程

thread_group tile32 = tiled_partition(g, 32);  // 把 block 切成 32 线程一组的 tile

thread_group tile4  = tiled_partition(tile32, 4); // 再把 tile32 切成 4 线程一组

限制：

• 切片大小必须是 2 的幂；
• 初始版本里最大为 32（= 一个 warp 的大小）。

上一页说 block 是"隐式组"，这页展示你自己定义更细的分解——通过 tiled_partition（分区）函数，把一个组按大小切成更小的组：

• this_thread_block() 拿到整个 block 组 g；
• 你可以把块组传给函数（foobar(thread_block g)），组对象随函数传，组内协作不受函数边界影响，跨函数照样能 sync；
• tiled_partition(g, 32) → 把 block 切成 32 线程一组（正好一个 warp）的 tile；
• 还能继续嵌套：tiled_partition(tile32, 4) → 再把每个 32 线程 tile 切成 4 线程一组——所以切片是可以递归/层级化的；
• 切出来的照样是 thread_group 类型，接口通用。

限制：切分大小必须2 的幂，且初始版≤ 32（因为最小的硬件同步单位就是 warp=32，tile 不该超过一个 warp）。

为什么有用：比如你想让每 4 个线程作为一个协作小组做局部归约、局部广播，就可以这样手动定义组的形状—— "分解粒度由程序自己定" ，这正是 Cooperative Groups 比"只有 block 一级同步"灵活得多的地方。

通用并行算法

按块调用（Per-Block）：

g = this_thread_block();          // 整个 block 作为一个组
reduce(g, ptr, myVal);            // block 级归约

按 warp 调用（Per-Warp）：

g = tiled_partition(this_thread_block(), 32);   // 切成 warp 组
reduce(g, ptr, myVal);                          // warp 级归约

核心通用函数：

__device__ int reduce(thread_group g, int *x, int val) {
    int lane = g.thread_rank();           // 我这个线程在组内的编号
    for (int i = g.size()/2; i > 0; i /= 2) {   // 树形规约: size/2 → ... → 1
        x[lane] = val;                    // 把我的值写进共享数组
        g.sync();                         // 同步 1
        if (lane < i) val += x[lane + i]; // 我这边负责的部分，累加配对的另一半
        g.sync();                         // 同步 2
    }
    return val;
}

这页讲什么——这才是 Cooperative Groups 最大的卖点：

同一个函数，两种粒度复用。 reduce() 完全不知道自己是 block 还是 warp——它只依赖组对象 g：

• g.size() —— 组里多少线程；
• g.thread_rank() —— 我在组内第几个；
• g.sync() —— 组内同步。

所以：

• 传 this_thread_block()，它就是一个 block 级归约（用 block 的共享内存数组 x，整块归约）；
• 传 tiled_partition(..., 32)，它自动就变成 warp 级归约——代码一行不用改！

归约逻辑（经典树形）： 每轮把组减半：先全体把自己的 val 写进位置 lane，同步，然后编号 < i 的前半线程把 x[lane + i] 累加进来，同步；i 从 size/2 一路除到 1；最后 lane=0 里就是和。

注意两点（这写法和 standard __syncthreads() 版本的区别）：

• 这里是通用的，靠 g.sync()——它内部根据组类型自动用 __syncthreads() 或 warp 级同步；
• 前面几页反复出现的"集体操作（collective） "含义在这体现：x[lane] = val; g.sync(); if (lane < i) ...——所有线程都要执行，不能有的跳过去（所以 sync 常在 if 外）。

Cooperative Groups 让你写"跟粒度无关"的并行算法——同一个 reduce()，传 block 组就是块归约，传 warp 组就是 warp 归约，靠 group 的 size/rank/sync 自适应。

线程块切片（THREAD BLOCK TILE）

定义：thread_block 的一个子集，按行优先（row-major）顺序切成大小相等的 tile。

thread_block_tile<32> tile32 = tiled_partition<32>(this_thread_block());

thread_block_tile<4>  tile4  = tiled_partition<4>(this_thread_block());

额外功能（比通用 thread_group 多，因为 size 编译期已知 = 快）：

.shfl()        // 任意 shuffle

.shfl_down()   // 向下广播（第二讲见过）

.shfl_up()     // 向上广播

.shfl_xor()    // 按 XOR 换 lane

.any()         // 组内是否任一线程满足条件

.all()         // 组内是否所有线程满足条件

.ballot()      // 返回条件掩码

.match_any()   // 哪些线程值相同

.match_all()   // 是否全部值相同

它把"warp 级协作"用 thread_block_tile<N> 这个模板组类型单独拎了出来：

• tiled_partition<32>(...) 得到大小 32 的 tile（就是一个 warp），tiled_partition<4>(...) 得到大小 4 的 tile；
• 注意这里用的是模板参数 <32>/<4>——所以组大小在编译期就是确定的，编译器能生成更快的指令（不用运行时判断），这是它比通用 thread_group 强的地方；
• 它额外暴露了一整套 warp 级原语：shfl/shfl_down/shfl_up/shfl_xor（就是我们第二讲学的 __shfl 系列，现在变成组方法），外加 any/all/ballot/match_any/match_all 这些投票/匹配类操作。

关键对照（呼应你第二讲的 Warp Shuffle）：

• 之前你用 __shfl_down_sync(mask, val, offset) 这种函数级 API；
• 现在可以用 tile.shfl_down(val, offset) 这种组对象方法，更清爽，而且组大小是模板参数，编译器替你保证都在一起执行（不需要手动传 mask）。

thread_block_tile<N> 是编译期定长、能直接调 shfl/ballot 等 warp 原语的组类型——把第二讲的 Warp Shuffle 升级成了更干净、更快的组对象写法。 （模板大小还能是 4、8、16 等小于 32 的 2 的幂，处理非整 warp 场景。）

静态切片归约

按 16 线程的 tile 调用：

g = tiled_partition<16>(this_thread_block());   // 切成每 16 线程一个组

tile_reduce(g, myVal);                           // tile 级归约

模板化的 tile 归约函数：

template <unsigned size>

__device__ int tile_reduce(thread_block_tile<size> g, int val) {

  for (int i = g.size()/2; i > 0; i /= 2) {   // 16 → 8 → 4 → 2 → 1

​    val += g.shfl_down(val, i);               // 用 shuffle 向下取值相加

  }

  return val;

}

把前两页的东西合起来用：用 thread_block_tile 组 + warp shuffle 做归约——这其实就是第二讲你学的 Warp Shuffle Reduction，只不过现在是组对象写法：

• 模板函数 tile_reduce<size> 接受任意大小的 tile 组（这里是 16）；
• 循环里 i = size/2 递减到 1，每轮调 g.shfl_down(val, i)——把比自己高 i 号位的线程的值拉下来累加（shfl_down 返回的是对方线程的 val，所以写法是 val += g.shfl_down(val, i)）；
• 为什么不用共享内存、不用同步？因为 shuffle 是直接在线程间交换寄存器值，不需要共享内存，也不需要 sync()（同一 warp 内天然锁步）——这就是第二讲"Warp Shuffle 好处"里的"省共享内存/省同步"。

调用它： tiled_partition<16> 把 block 切成每组 16 线程的 tile，每组各做各的归约，16 个线程最后 lane 0 里就是这一组的和。

对照第二讲：你当时写的版本是

for (offset=warpSize/2; offset>0; offset>>=1) val += __shfl_down_sync(0xffffffff, val, offset);

而这页就是它的 Cooperative Groups 版——g.shfl_down(val, i) 一行搞定，还不用手传 mask，因为模板大小保证组内 = 一个 warp 内锁步。

Static Tile Reduce = 用 thread_block_tile.shfl_down 做 warp 级归约，比第二讲的裸 __shfl_down_sync 更简洁——shuffle 直连寄存器、零共享内存、零显式同步。

网格组（GRID GROUP）

• Grid Group = 同一个 grid 内的所有线程的集合，并且保证这些线程此刻都驻留在设备上（resident） ——这是能不能做全局同步的前提。

要使用它，得用新的 CUDA 启动 API 显式选择（opt-in）：

cudaLaunchCooperativeKernel(...)   // 合作式 kernel 启动

kernel 内用法：

__global__ kernel() {

​    grid_group grid = this_grid();   // 拿到整个 grid 的组

​    // load data                     // 阶段1：加载数据

​    // loop - compute, share data    // 阶段2：计算、共享数据

​    grid.sync();                     // ← 设备级全局执行屏障！

}

两个前提条件：

• 设备必须支持 cooperativeLaunch 属性（不是所有 GPU 都支持）；
• 用 cudaOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocksPerSm, kernel, numThreads, 0) 查询每个 SM 最多能同时驻留多少 block——因为 grid sync 要求所有 block 必须能同时驻留，否则会死锁。

这就是你第二讲"跨 block 全局同步"那个难题的终极正式解法：

1. 以前：想跨 block 同步只能拆 kernel / 原子 / fence（第二讲列的笨办法），因为 grid 里 block 可能不是同时运行的，没法安全地全局屏障；
2. 现在：Cooperative Groups 让你 this_grid() 拿到 grid 组，然后 grid.sync() 一次性把整个 grid 的所有线程同步到位——跨 block 就像 __syncthreads() 一样简单。

但有两个关键代价/前提：

• 必须用新启动 API cudaLaunchCooperativeKernel 启动 kernel（这是"opting in"合作执行模式），普通 <<<>>> 不行；
• 所有 block 必须能同时驻留（resident）在 GPU 上才能安全同步——所以你要先用 cudaOccupancyMaxActiveBlocksPerMultiprocessor 算出每 SM 能放几个 block，再用它限制 grid 尺寸（保证全部 block 同时驻留）。如果 grid 太大、有些 block 不在设备上，grid.sync() 就会死锁。

Grid Group + cudaLaunchCooperativeKernel 让你用 this_grid() + grid.sync() 实现跨 block 全局同步；前提是设备支持 cooperativeLaunch，且用 occupancy API 保证所有 block 同时驻留，避免死锁。 这就是官方给出的"跨 block 同步"标准答案。

🎯 GRID GROUP 并行计算场景说明

这张图是并行计算领域的经典分类图示，核心目标是尽可能让计算状态常驻内存，以此提升计算效率，图中展示了三类典型的网格组并行计算场景。

🛤️ 最短路径/搜索场景

• ‌核心特征‌：图中展示了带权无向图的Dijkstra最短路径求解示例，起点为红色标记的节点a，终点为节点b，图中标注了各边的权重数值。
• ‌设计思路‌：权重数组天然适配持久化存储需求，通过对顶点进行迭代遍历完成计算，最终可以将多轮迭代的计算逻辑进行融合优化，减少不必要的重复计算开销。

🧬 主从驱动算法场景

• ‌核心特征‌：以遗传算法为代表的主从架构并行模式，中心节点维护全局最优解global best，多个从节点slave各自维护本地最优解local best和种群中的个体最优解an individual in population。
• ‌设计思路‌：核心逻辑是完成主节点和所有从节点之间的状态同步，让全局最优信息可以高效地在分布式节点间流转，保障算法的收敛效率。

⚛️ 粒子模拟场景

• ‌核心特征‌：网格空间内分布大量运动粒子，粒子带有运动方向箭头，在网格边界和内部会发生碰撞交互。
• ‌设计思路‌：核心要解决粒子状态更新逻辑和碰撞模拟逻辑之间的同步问题，避免出现粒子位置计算冲突、碰撞判定错误等问题，保障模拟过程的准确性。

🎯 多网格组（MULTI GRID GROUP）概念解析

这是NVIDIA CUDA并行计算体系中面向多GPU协同的高级线程组织模型，核心目标是让一组线程在同一系统内的多个GPU设备上常驻运行，实现跨设备的高效并行协作。

📋 核心基础定义

• ‌线程组特性‌：这组线程被保证驻留在同一个主机系统中，可分布在多个不同的GPU硬件设备上执行，突破了传统单GPU单Grid的线程调度限制。
• ‌跨设备同步能力‌：区别于传统CUDA中仅支持单GPU内线程块同步的机制，多网格组提供了原生的跨设备同步接口，让多GPU的计算流程可以精准协同。

1. 通过this_multi_grid()获取当前跨设备的多网格组实例
2. 完成数据加载后，多GPU并行执行计算并互相共享数据
3. 调用multi_grid.sync()触发跨设备全局同步，确保所有参与的GPU都完成当前阶段计算
4. 同步完成后，所有设备继续执行后续的计算任务

🖥️ 硬件协作架构

图中直观展示了双GPU的协作模式：

• GPU A内部包含Block 0和Block 1两组线程块，GPU B同样包含Block 0和Block 1两组线程块
• 两个GPU之间通过专门的同步通道完成全局状态对齐，所有跨设备的线程块都能在同步点达成执行状态一致，避免多GPU计算出现数据依赖错误。

🎯 多网格组多设备内核启动方案

这是CUDA中支持一次性在多个GPU设备上启动协作内核的高级编程方案，实现跨多GPU的线程协同运行。

📋 核心API与硬件要求

• ‌新增专属启动API‌：cudaLaunchCooperativeKernelMultiDevice(...)
  这是CUDA专门为多设备协作启动开放的新接口，区别于传统单设备内核启动方式，支持同时向多个GPU派发同一个内核任务。
• ‌硬件兼容性要求‌：所有参与协作的GPU设备，必须支持cooperativeMultiDeviceLaunch属性，否则无法使用该多设备启动能力。

每个设备都需要独立配置一套启动参数，分别指定内核函数、网格维度、线程块维度、共享内存大小、关联流以及内核参数。

⚠️ 关键使用约束

• 不能使用NULL空流来关联内核任务，必须为每个设备指定有效的非空CUDA流。
• 网格维度配置建议结合CUDA占用率计算器来优化，最大化GPU硬件资源利用率。
• 所有参与启动的设备必须提前校验多设备协作属性支持，避免运行时启动失败。

🎯 合并线程组（Coalesced Group）概念解析

这是CUDA协作组体系中面向分支收敛场景的高级线程分组机制，专门用于处理SIMD执行模式下的活跃线程管理。

📋 核心定义说明

Discover the set of coalesced threads, i.e. a group of converged threads executing in SIMD
译文：识别出合并后的线程集合，也就是一组已经完成分支收敛、以SIMD模式执行的线程

它是CUDA 9引入的协作组扩展能力，专门解决传统线程束执行中因分支分化导致的部分线程不活跃的问题。

🔍 关键特性

• ‌动态发现机制‌：通过coalesced_threads()接口可以直接获取当前所有处于活跃状态的线程，图中示例返回的合并组大小为8。
• ‌SIMD执行保障‌：组内所有线程都已经完成分支收敛，能够以单指令多数据的方式同步执行，不会出现分支分化带来的性能损耗。
• ‌架构兼容性‌：该机制可以跨不同NVIDIA GPU架构稳定工作，为线程束聚合原子操作等优化场景提供可靠支持。

💡 典型使用场景

• ‌减少原子操作竞争‌：让合并组内仅第一个活跃线程代表全组执行原子操作，一次性累加组内总线程数，避免每个线程都发起原子请求，大幅降低冲突开销。
• ‌不规则数据处理‌：在点云遍历、八叉树访问、稀疏矩阵计算这类存在大量分支提前退出的场景中，精准筛选出仍在运行的活跃线程，避免无效计算。
• ‌动态线程协作‌：处理分支执行后剩余的活跃线程集合，实现更灵活的细粒度线程同步与数据共享。

🎯 SIMD 核心定义

SIMD 是单指令多数据并行计算架构，是现代GPU高性能计算的核心执行基础。

📋 基础概念

• ‌全称‌：Single Instruction Multiple Data，单指令流多数据流
• ‌核心逻辑‌：仅发射一条指令，就能同时对多个不同的数据元素执行相同的操作，大幅提升数据并行处理效率。
• ‌典型应用场景‌：图形渲染、向量运算、AI推理、科学计算等存在大量重复数据操作的场景。

🔍 GPU 中的 SIMD 实现

NVIDIA GPU 中最典型的SIMD执行单元就是线程束（Warp）：

• 单条SIMD指令可以同时驱动32个线程同步执行相同的运算逻辑
• 前文提到的合并线程组（Coalesced Group），就是分支收敛后完全以SIMD模式运行的活跃线程集合，避免了分支分化带来的性能浪费。
• 相比传统单指令单数据（SISD）架构，SIMD在批量数据处理场景下可以获得数倍到数十倍的性能提升。

💡 常见SIMD扩展

除了GPU硬件层面的SIMD，通用CPU也有对应的软件级SIMD指令集：

• x86架构下的SSE、AVX、AVX-512指令集
• ARM架构下的NEON指令集
  这些指令集都可以让CPU单条指令同时处理多个数据，加速音视频编解码、数值计算等场景。