Lecture 2: CUDA 编程模型基础 —— Kernel、线程层次与内存模型 (对应 Lab 0 / Lab 1)
Lecture 2: CUDA 编程模型基础 —— Kernel、线程层次与内存模型 (对应 Lab 0 / Lab 1)
概述
本讲要回答一个核心问题：如何把”每个输出元素彼此独立”的数据并行程序，翻译成 GPU 能高效执行的代码。CUDA 给出的答案是”主机（host）写串行控制流 + 设备（device）写 SPMD kernel”的异构执行模型：程序员用 __global__ 声明一个 kernel，用 <<<gridDim, blockDim>>> 启动它，GPU 用成千上万个拥有唯一编号的线程同时执行同一个函数体，各自用 blockIdx/blockDim/threadIdx 算出自己该处理的那一份数据。为了让这套抽象落到硅片上，硬件把 block 内部的线程按线性编号每 32 个切成一个 warp（线程束），warp 才是 SM（Streaming Multiprocessor，流式多处理器）真正的调度单位，这直接导致”分支发散”和”合并访问”两个贯穿全课程的性能主题。最后，本讲给出 CUDA 的六级内存层次（register / local / shared / global / constant / texture），它们的作用域、生命周期和访问延迟相差两到三个数量级，是后续所有优化（尤其是 Lab 2 的 shared memory 复用）的物理基础。
核心概念与 GPU 架构图解
Kernel 函数与 SPMD 数据并行模型（Kernel / SPMD）
定义与目的：
__global__修饰的函数叫 kernel（内核函数）。它在 GPU 上执行、由主机端代码启动，返回值必须是void。一个 kernel 的一次启动会产生一个 grid（线程网格），grid 中的每个线程执行同一份 kernel 代码，但用自己唯一的索引决定”读哪块内存、算哪份数据、走哪条分支”。这就是 SPMD（Single Program, Multiple Data，单程序多数据） 模型。它解决的问题是：数据并行程序（如向量加法、图像逐像素变换）里 CPU 上的for循环完全可以被”把循环体变成线程体、把循环变量变成线程编号”的机械变换替代，从而把 N 次串行迭代摊到 N 个物理并行单元上。直观解释（”它是什么？”）：想象一家公司给 5000 名员工发同一本操作手册（kernel 代码），每个员工手册内容完全一样，但工牌号（thread index）不同，工牌号决定了”你去仓库取第几号箱子、把结果放进第几号货架”。没有主管逐个下命令说明每一步——手册就是程序，工牌号就是数据索引。这就是 SPMD：一份程序，N 份数据。与之对比，CPU 的写法是一位主管（单线程）亲自把 5000 个箱子搬完。
讲义中的经典起点是图像灰度化：
for each pixel { pixel = gsConvert(pixel) } // Every pixel is independent of every other pixel因为”每个像素与其它像素无关”，所以这个循环的每一次迭代都可以变成一个独立的线程：这正是数据并行（data parallelism） 的判定标准——输出的每个元素只依赖输入，不依赖其它输出（无循环携带依赖）。
架构/机制图解：
Host 端 C 代码 (串行 / 适度并行) Device 端 kernel 代码 (高度并行, SPMD)
┌──────────────────────────────┐ ┌────────────────────────────────────┐
│ void vecAdd(float*A,float*B, │ │ __global__ void vecAddKernel( │
│ float*C,int N){ │ │ float* A_d, float* B_d, │
│ cudaMalloc(...) │ │ float* C_d, int N) { │
│ cudaMemcpy(H2D) │ │ int i = blockIdx.x*blockDim.x │
│ │ │ + threadIdx.x; │
│ vecAddKernel<<<nBlk,nTid>>>│ ───────► │ if (i < N) │
│ (A_d,B_d,C_d,N); │ C_d[i] = A_d[i] + B_d[i]; │
│ │ │ } │
│ cudaMemcpy(D2H) │ ◄─────── │ │
│ cudaFree(...) │ │ │
│ } │ │ │
└──────────────────────────────┘ └────────────────────────────────────┘
└─ 编译器: host 编译器 (gcc/clang) └─ 编译器: nvcc 前端 -> PTX -> SASS
一次 kernel 启动的逻辑结构（KernelA<<<nBlk, nTid>>>(args) 是”异步”的，主机发完命令立刻继续往下跑）：
kernel<<<nBlk, nTid>>>(args)
│
├─ gridDim = nBlk = (grid 每一维的 block 个数)
├─ blockDim = nTid = (每个 block 每一维的线程个数)
│
▼
┌───────────────────────── Grid（一次 kernel 启动 = 一个 grid）──────────────┐
│ Block 0 Block 1 Block 2 ... Block nBlk-1 │
│ ┌────────┐ ┌────────┐ ┌────────┐ ┌────────┐ │
│ │ 256 th │ │ 256 th │ │ 256 th │ │ 256 th │ │
│ └────────┘ └────────┘ └────────┘ └────────┘ │
│ blockIdx.x=0 blockIdx.x=1 blockIdx.x=2 ... blockIdx.x=nBlk-1 │
└────────────────────────────────────────────────────────────────────────────┘
性能视角：SMPD 的”零成本展开”是有代价的——网格里可能有大量线程因为越界检查而什么都不做（例如 N=1000、block=256 时需要 4 个 block 共 1024 个线程，24 个线程被 if (i<N) 判掉，浪费率 2.3%），同时 kernel 启动本身有约 3–10 µs 的固定开销，所以小数据量（< 10 万个元素）时往往跑不过 CPU。
异构计算与 host/device 代码分离（Heterogeneous Computing）
定义与目的：一个 CUDA 程序同时包含两种代码：host 代码在 CPU 上执行（做 I/O、内存分配、启动 kernel、串行部分），device 代码在 GPU 上执行（做计算密集的并行部分）。三个函数修饰符划清了边界：
__host__（只能在 CPU 调用、在 CPU 执行，默认）、__device__（只能在 GPU 调用、在 GPU 执行）、__global__（从 CPU 调用、在 GPU 执行，因此被叫做 kernel）。目的是让程序员用一门语言、一个源文件同时驾驭两个不同架构的处理器，而不必写两套工程。直观解释（”它是什么？”）：把 CPU 比作总经理，GPU 比作上千人的车间。总经理不做流水线活，只负责”下订单、把原材料从仓库（主机内存）运到车间（显存）、发号施令（kernel 启动）、把成品运回来”。
__global__就是”发给车间的一张工单”，__device__是”车间内部师傅之间互相调用的工序说明书”，__host__是”办公室内部的流程”。架构/机制图解：
源文件 prog.cu
│
▼
┌──────────┐ ① 分离 host / device 代码
│ NVCC │
└────┬─────┘
├───────────────────────────────► Host Code (纯 C/C++) ──► gcc/clang/msvc ──► .o
│ │
└──► Device Code ──► PTX (虚拟 ISA, 可移植) ──► ② ptxas 汇编 ──► SASS (sm_80 机器码)
│ │
└──► ③ JIT: 运行时由驱动把 PTX 编译成目标 GPU 的 SASS │
▼
链接成可执行文件
│
▼
┌───────────── 异构平台：CPU + GPU（通过 PCIe/NVLink 连接）──────┐
│ CPU: 控制器 + 大容量 DRAM GPU: 数千 ALU + 高带宽 HBM │
└────────────────────────────────────────────────────────────────┘
三种修饰符的调用/执行矩阵（对应讲义 “More on CUDA Function Declarations”）：
修饰符 执行位置 只能被谁调用 备注
---------- -------- ------------------ ----------------------------------
__host__ host 仅 host 默认，可省略
__global__ device 仅 host (kernel) 返回类型必须是 void，异步启动
__device__ device 仅 device 高效的内联小函数；不能单独出现在
函数体内（需与 __shared__/
__constant__ 等连用才合法）
__host__ __device__ 两者 两者 一份代码两种编译，常用于公共工具函数
性能视角：host 与 device 之间的每次 cudaMemcpy 都要走 PCIe 4.0 x16（双向理论 32 GB/s，实测约 25 GB/s）或 NVLink（A100 上约 600 GB/s 双向）。相对 A100 自身 1555 GB/s 的显存带宽，PCIe 要慢 48 倍。因此一条铁律是：数据搬移次数越少越好，单次搬移量越大越好；Lab 1 中把 A、B 各传一次、C 取回一次是标准做法。
线程层次：Grid → Block → Thread（Thread Hierarchy）
定义与目的：CUDA 线程被组织成三层：一个 grid 由若干 thread block（线程块） 组成，一个 block 由若干 thread 组成；每一层都可以是 1D/2D/3D 的（用
dim3描述）。每个 block 有唯一的三元组blockIdx（取值 0 … gridDim-1），每个线程在 block 内也有唯一的三元组threadIdx（取值 0 … blockDim-1）。这套层次不是为了好看，而是为了：① 让”多维数据”（图像、矩阵、体数据）的寻址在代码里直观可读；② 给硬件一个明确的可调度性与可同步性边界——同一个 block 内的线程可以用 shared memory 通信、用__syncthreads()同步，不同 block 之间不能这样合作，且block 的执行顺序是任意的（arbitrary order）。直观解释（”它是什么？”）：把 grid 想成整栋写字楼，block 是一间间独立的办公室，thread 是办公室里的员工。同一间办公室的员工可以共用一块白板（shared memory）并随时喊”大家停一下”（
__syncthreads()）；不同办公室之间没有共享白板，只能通过楼下的公告栏（global memory）传话，而且谁先上班谁后上班没有规定——你不能假设 12 楼的办公室一定比 7 楼先开工。这就是”block 执行顺序任意”的直观含义，也是为什么 kernel 里绝不能写”等别的 block 先算完”这种逻辑。架构/机制图解：
┌──────────────────── Grid: gridDim = (4, 3, 1) → 共 4x3x1 = 12 个 block ───────────┐
│ │
│ blockIdx=(0,2) blockIdx=(1,2) blockIdx=(2,2) blockIdx=(3,2) │
│ ┌──────────┐ ┌──────────┐ ┌──────────┐ ┌──────────┐ │
│ │ Block │ │ Block │ │ Block │ │ Block │ y = gridDim.y-1 │
│ └──────────┘ └──────────┘ └──────────┘ └──────────┘ │
│ blockIdx=(0,1) blockIdx=(1,1) blockIdx=(2,1) blockIdx=(3,1) │
│ blockIdx=(0,0) blockIdx=(1,0) blockIdx=(2,0) blockIdx=(3,0) y = 0 │
│ x = 0 x = 1 x = 2 x = 3 │
└────────────────────────────────────────────────────────────────────────────────────┘
│
取出其中一个 block，看它内部的线程布局
▼
┌──── Block blockIdx = (1,0), blockDim = (4, 2, 1) → 共 4x2x1 = 8 个线程 ────┐
│ │
│ threadIdx=(0,1) (1,1) (2,1) (3,1) ← threadIdx.y = 1 │
│ threadIdx=(0,0) (1,0) (2,0) (3,0) ← threadIdx.y = 0 │
│ ▲ │
│ threadIdx.x = 0..blockDim.x-1 │
│ │
│ 注意: threadIdx 只在 block 内唯一（WITHIN A BLOCK），跨 block 会重复！ │
└────────────────────────────────────────────────────────────────────────────────┘
│
硬件把 block 内线程按"x 变化最快"线性化，再每 32 个切一个 warp
▼
┌──────────────── Warp：SM 的真正调度单位（32 个线程一组）──────────────────────┐
│ linear tid = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z*blockDim.x*blockDim.y
│ │
│ warp 0 : linear tid 0 .. 31 ─┐ │
│ warp 1 : linear tid 32 .. 63 ├─ 每个 warp 占 SM 的一个调度槽位 │
│ warp 2 : linear tid 64 .. 95 ─┘ │
│ ... │
└──────────────────────────────────────────────────────────────────────────────┘
内置变量速查（全部是只读的 uint3/dim3 类型）：
变量 类型 含义 取值范围
-------- ------- ------------------------------------ ---------------------------
gridDim dim3 grid 每一维的 block 个数 x<=2^31-1, y,z<=65535
blockIdx uint3 当前 block 在 grid 中的坐标 0 .. gridDim-1 (每一维)
blockDim dim3 每个 block 每一维的线程个数 x<=1024, 且 x*y*z<=1024
threadIdx uint3 当前线程在 block 内的坐标 0 .. blockDim-1 (每一维)
warpSize int 一个 warp 的线程数（当前所有 NVIDIA GPU） 32
关键性能数字：一个 block 最多 1024 个线程、blockDim.y/blockDim.z 最多 1024、gridDim.y/gridDim.z 最多 65535（gridDim.x 可达 2^31-1）。这些上限在概念上”不是 CUDA 编程模型的一部分”——编程模型只要求”线程能分组、组内能同步”，1024、32 都是硬件实现决定（implementation decision），所以代码里应尽量用 warpSize 或 #define 而不是硬编码 32。
全局线程编号的推导与多维索引展平（Linearization / Index Flattening）
定义与目的：kernel 里最常用的一行就是
int i = blockIdx.x * blockDim.x + threadIdx.x;。它把”二维的 block 坐标 + 一维的线程坐标”压成一个沿 x 方向从 0 连续递增的全局编号，用来索引一维数组。推广到二维、三维时，就是把 (row, col) 映射成行主序（row-major）的一维偏移Row * width + Col。目的是让相邻线程访问相邻内存，这是达成合并访问（coalescing）的前提。直观解释（”它是什么？”）：像电影院找座位。”第 3 排第 5 号”是二维坐标，但影院经理清点人数时用的是”第 3 排之前的 2 排共 2×每排座位数，再加第 5 号”这样一个连续的一维编号。CUDA 的这行公式就是”排号 × 每排座位数 + 座号”；
blockIdx.x是”排号”（排很长，排内又有 block），blockDim.x是”每排有几个 block 的位置”。架构/机制图解：
N = 1000, block size = 256 → gridDim.x = ceil(1000/256) = 4
一维情况：
i = blockIdx.x * blockDim.x + threadIdx.x
│ │ │
│ │ └── 块内线程号 E （0..255）
│ └─────────────── 每块线程数 D （=256）
└─────────────────────────── 块号 C （0..3）
┌── Block 0 ──────────────┐┌── Block 1 ──────────────┐┌── Block 2 ───────┐┌─ Block 3 ──┐
│ tid: 0 1 2 ... 254 255││ tid: 0 1 2 ... 254 255││ 0 ... 255 ││ 0 ... 255 │
│ i : 0 1 2 ... 254 255││ i :256 257 ... 510 511 ││512 ... 767 ││768 ...1023│
└───────────────────────────┘└─────────────────────────┘└────────────────┘└───────────┘
▲
i = 1024 > N-1 = 999 → 被 if(i<N) 屏蔽
→ 启动 1024 个线程，只有 1000 个干活，浪费 2.34%
二维情况：把两个一维公式分别用在两个方向上
Col = blockIdx.x * blockDim.x + threadIdx.x; // 横向编号（列）
Row = blockIdx.y * blockDim.y + threadIdx.y; // 纵向编号（行）
offset = Row * width + Col; // 行主序展平（讲义 Page 12 的 2D 数组布局）
行主序示意图（width = 4）：M2,1 → 2*4 + 1 = 9
列: 0 1 2 3
行 0 ┌────┬────┬────┬────┐
│ 0 │ 1 │ 2 │ 3 │
行 1 ├────┼────┼────┼────┤
│ 4 │ 5 │ 6 │ 7 │
行 2 ├────┼────┼────┼────┤
│ 8 │ 9 │ 10 │ 11 │ ← 元素 (2,1) 的线性下标 = 2*4+1 = 9
行 3 ├────┼────┼────┼────┤
│ 12 │ 13 │ 14 │ 15 │
└────┴────┴────┴────┘
性能视角：block 数量必须”向上取整”，否则末尾元素没人处理：
gridDim.x = ceil(N / blockDim.x) = (N + blockDim.x - 1) / blockDim.x ← 整数写法
反例：N = 1000, blockDim.x = 256
C 语言里 N/256 = 1000/256 = 3 （整数除法向下取整！）
→ 只启动 3 个 block = 768 个线程 → 元素 768..999 共 232 个永远不被计算（静默错误）
正解： (1000 + 255)/256 = 1255/256 = 4 → 1024 个线程覆盖 1000 个元素
讲义写法： dim3 DimGrid(ceil(N/256.0), 1, 1) 或 dim3 DimGrid(N/256,1,1); if(N%256) DimGrid.x++;
Warp 与 SIMT 执行模型（Warp / SIMT）
定义与目的：一个 block 内的线程按照线性化线程编号每 32 个一组被切成 warp（线程束）：线程 0–31 是 warp 0，32–63 是 warp 1，依次类推。warp 是 SM 的调度单位，一个 warp 内的 32 个线程在同一时刻被同一套硬件用一个指令流驱动，这就是 SIMT（Single Instruction, Multiple Threads）。它解决的问题是：如果 32 个线程各自取指、译码，控制开销会大到不可接受；让它们共享取指/译码、只让数据通路并行，就能用极低的控制成本换来极高的吞吐。
直观解释（”它是什么？”）：像军训方队。教官（warp scheduler）只喊一句”向右转”，整个 32 人方队同时转身——取指令、译码只做一次，动作由 32 个人各自的身体完成。但如果教官喊”戴眼镜的向前一步”，队伍里只有一部分人动，另一部分人站着等，方队还是得按两条命令花两倍时间走完——这就是 warp 发散（divergence）。
架构/机制图解：
SIMT 执行：一个 warp 共享 PC，32 条数据通路并行
┌─────────────────────────── Warp Scheduler ───────────────────────────┐
│ 取指 ──► 译码 ──► 发射一条指令到 32 条 lane │
└───────────────────────────────┬───────────────────────────────────────┘
│ ADD R1, R2, R3
┌──────┬──────┬──────┬──────┬──────┬──────┬─ ··· ─┬──────┐
│lane0 │lane1 │lane2 │lane3 │lane4 │lane5 │ │lane31│ ← 32 条 lane
└──┬───┴──┬───┴──┬───┴──┬───┴──┬───┴──┬───┴─ ··· ─┴──┬───┘
│ │ │ │ │ │ │
Reg0-1 Reg1-1 Reg2-1 Reg3-1 Reg4-1 Reg5-1 ... Reg31-1 ← 每 lane 独立的寄存器组
(线程0) (线程1)(线程2)(线程3)(线程4)(线程5) ... (线程31)
---------------------------------------------------------------------------------
Warp 发散（branch divergence）：if (threadIdx.x > 2) { THEN } else { ELSE }
---------------------------------------------------------------------------------
lane: 0 1 2 │ 3 4 5 ... 31
predicate:0 0 0 │ 1 1 1 ... 1
│ │
▼ ▼
┌─────────────────────────────────────────────────────────────┐
│ 时间 → │
│ ① 执行 THEN 路径：只有 lane3..31 有效（结果写回）， │
│ lane0..2 被谓词屏蔽（predicated off） │
│ ② 执行 ELSE 路径：只有 lane0..2 有效，lane3..31 被屏蔽 │
│ ③ 重汇聚（reconverge）到 if 之后的代码 │
│ │
│ 代价：两条路径串行，本应 1 个 warp 时间单位完成的工作花了 2 个 │
│ 若 THEN/ELSE 各 100 条指令，则这个 warp 执行 200 条指令 │
└─────────────────────────────────────────────────────────────┘
发散的本质：执行路径依赖于"线程独有的量"（threadIdx、由 threadIdx 算出的下标）。
若依赖的量在同一 warp/block 内所有线程相同（如统一的内核参数 q），则不会发散。
避免发散的经典手法（讲义 Page 27）：让分支粒度是 warp size 的整数倍：
if (threadIdx.x > 2) { // 坏：同一 warp 内 lane0..2 与 lane3..31 走不同路径
/* THEN path (lots of lines) */
} else {
/* ELSE path (lots more lines) */
}
if (threadIdx.x / WARP_SIZE > 2) { // 好：warp 0 全体走 THEN，warp 1.. 全体走 ELSE
/* THEN path (lots of lines) */
} else {
/* ELSE path (lots more lines) */
}
现代 GPU（Volta sm_70 及以后）引入了独立线程调度（Independent Thread Scheduling）：每个线程拥有自己的 PC 和调用栈，发散的分支可以基于 per-thread 的 SIMT 栈交错推进，不再强制”先走完 THEN 再走完 ELSE”。这让 warp 内的线程级同步原语（如 __syncwarp()）和细粒度生产者–消费者模式变得可行（例如 __shfl_sync 交换需先 __syncwarp()）。但请注意：独立线程调度并不消除发散的性能代价——32 条 lane 的物理数据通路依然是共享的，两个分支仍然要串行占用发射槽，只是调度更灵活、不再有”死锁式”的假设。
性能数字（典型值）：
项目 A100 (sm_80) RTX 4090 (sm_89)
----------------------------- ------------ ----------------
每 SM 最多驻留 warp 64 48
每 SM 最多驻留线程 2048 1536
warp 大小 / 调度单位 32 32
每 SM warp scheduler 数量 4 4
每周期可发射的 warp 指令数 4 4
每条 warp 指令的 lane 数 32 32
一个 256 线程 block = 8 个 warp 8 个 warp
一个 1024 线程 block = 32 个 warp 32 个 warp
全 SM 满载时未完成的 warp 状态总量 64 48
→ 讲义 "Thread Scheduling (1/2)" 问题：3 个 256 线程 block 分到 1 个 SM，
共 3 * (256/32) = 24 个 warp
SM 硬件结构与 warp 调度（Streaming Multiprocessor）
定义与目的：SM（Streaming Multiprocessor）是 GPU 的”核心中的核心”，一个 GPU 有几十到一百多个 SM（A100 有 108 个，RTX 4090 有 128 个，H100 SXM 有 132 个）。SM 里装着：寄存器文件、L1 缓存/共享内存（统一的一块 SRAM，可分区配置）、若干 warp scheduler、以及执行单元（CUDA core、SFU、LD/ST unit、Tensor Core）。block 是以整块为单位被分配到 SM 上的，一个 SM 可以同时驻留多个 block，SM 负责维护它们的 block id / thread id 并调度其中 warp 的执行。
直观解释（”它是什么？”）：SM 像一间有 4 个工位的车间。车间里有 64 个工人（warp），每个人手里有活（指令），但只有 4 个工位（warp scheduler）。工位长每拍从”手头材料齐全”的工人里挑 4 个上工位干活；材料没到的（等待全局内存返回数据的 warp）就靠边站。这就是延迟隐藏（latency hiding）：只要待命的 warp 足够多，某几个 warp 卡在 500 周期的访存延迟上时，其它 warp 能继续占满工位，SM 的吞吐就不会掉。
架构/机制图解：
┌────────────────────────────── SM (Ampere A100, sm_80) ─────────────────────────────┐
│ │
│ ┌── Processing Block 0 ──┐ ┌── Processing Block 1 ──┐ ┌── PB 2 ──┐ ┌── PB 3 ──┐ │
│ │ warp scheduler 0 │ │ warp scheduler 1 │ │ sched 2 │ │ sched 3 │ │
│ │ │ (每周期选 1 个 warp)│ │ │ │ │ │ │ │ │
│ │ ▼ │ │ ▼ │ │ │ │ │ │
│ │ 16 x FP32 CUDA core │ │ 16 x FP32 CUDA core │ │ 16 FP32 │ │ 16 FP32 │ │
│ │ 16 x INT32 core │ │ 16 x INT32 core │ │ 16 INT32│ │ 16 INT32│ │
│ │ 4 x SFU (sin/cos/rsqrt)│ │ 4 x SFU │ │ 4 SFU │ │ 4 SFU │ │
│ │ 1 x LD/ST unit │ │ 1 x LD/ST unit │ │ 1 LD/ST │ │ 1 LD/ST │ │
│ │ 1 x Tensor Core │ │ 1 x Tensor Core │ │ 1 TC │ │ 1 TC │ │
│ └────────────────────────┘ └────────────────────────┘ └──────────┘ └──────────┘ │
│ │
│ ┌──────────────── Register File ────────────────┐ ┌─ L1 / Shared Memory ──────┐ │
│ │ 65536 x 32-bit = 256 KB, 按 warp 分配 │ │ 192 KB 统一 SRAM, │ │
│ │ 每线程最多 255 个寄存器 │ │ 最多 164 KB 划给 shared │ │
│ └───────────────────────────────────────────────┘ └────────────────────────────┘ │
│ │
│ 资源上限: 64 warp / 2048 thread / 32 block 同时驻留 │
└────────────────────────────────────────────────────────────────────────────────────┘
▲ ▲
│ 以 block 为粒度分配 │ miss 时向 L2 取
┌─────────────┴─────────────────────────────────────────────────────┴────────────────┐
│ L2 Cache (A100: 40 MB，全芯片所有 SM 共享) │
└────────────────────────────────────────────────────────────────────────────────────┘
▲
┌─────────────┴─────────────────────────────────────────────────────────────────────┐
│ Global Memory (HBM2, A100: 40/80 GB, 1555 GB/s, 延迟 400-800 cycles) │
└────────────────────────────────────────────────────────────────────────────────────┘
零开销 warp 调度（zero-overhead warp scheduling）的两条规则（讲义 Page 24）：① 只有下一条指令的操作数已就绪的 warp 才是 eligible（可执行）；② 在 eligible 的 warp 中按优先级策略挑选。因为”线程切换”所需的 PC、寄存器组都常驻在 SM 上，切换 warp 不保存/恢复上下文，所以是零开销的——这正是 GPU 用海量线程隐藏长延迟的根本机制。
性能数字与”占用率（occupancy）”的定量计算（讲义 Page 28 的经典问题，用 A100 的参数重做）：
A100: 每 SM 2048 线程、64 warp、寄存器 65536、共享内存 164 KB、每 SM 最多 32 个 block
问题：colorToGreyscaleConversion 该用 8x8、16x16 还是 32x32 的 block？
设每个线程用 20 个寄存器、block 不用共享内存。
(a) block = 8x8 = 64 线程 = 2 warp
线程限制: 2048/64 = 32 个 block
寄存器限制: 每 block 寄存器 = 64*20 = 1280 (向上对齐到 256 粒度 → 1280)
65536/1280 = 51.2 → 51 个 block
→ 受线程限制，32 个 block，恰好 32*64 = 2048 线程 = 64 warp = 100% 占用率
但若硬件上限是 8 个 block/SM（讲义中的假设），则只有 8*64 = 512 线程 = 16 warp = 25%
(b) block = 16x16 = 256 线程 = 8 warp
线程限制: 2048/256 = 8 个 block
寄存器限制: 256*20 = 5120 → 65536/5120 = 12.8 → 12 个 block
→ 8 个 block = 2048 线程 = 64 warp = 100% 占用率 ★ 最优
(c) block = 32x32 = 1024 线程 = 32 warp
线程限制: 2048/1024 = 2 个 block
→ 2 个 block = 2048 线程 = 64 warp = 100% 占用率
但若硬件上限是 1 个 block/SM（讲义中的假设），则只有 1024 线程 = 32 warp = 50%
讲义给出的硬件假设（每 SM 最多 1536 线程、最多 8 个 block）：
8x8 : 64 线程/块 → SM 可放 24 块，但受 8 块上限 → 只有 512 线程(16 warp) → 用不满
16x16: 256 线程/块 → 1536/256 = 6 块（≤8） → 1536 线程(48 warp) → 用满 ★
32x32: 1024 线程/块 → 1536/1024 = 1 块 → 只有 1024 线程(32 warp) → 只用 2/3
CUDA 内存模型（CUDA Memory Model）
定义与目的：CUDA 为线程提供六个逻辑存储空间，它们用变量声明的位置和修饰符区分，各自有明确的作用域（scope） 与生命周期（lifetime）。设计目的是用一块容量小但极快、且可被程序员显式管理的片上存储（register + shared memory）来对抗”处理器比内存快得多”的冯·诺依曼瓶颈——CPU 靠硬件缓存自动完成这件事，GPU 则把其中最关键的一层（shared memory）交给程序员显式控制，从而在数据复用上做到比硬件缓存更精准。
- 直观解释（”它是什么？”）：把 GPU 想象成一家餐厅：
- 寄存器（register） = 厨师手上的调料勺：最快（1 个周期），但每人只有几把（每线程最多 255 个）。
- 共享内存（shared memory） = 后厨的案板：一块 block 里所有厨师共用，从冰箱取出食材放案板上反复切（复读），比每次跑冰箱快 20 倍（20–30 周期 vs 400–800 周期）。
- 全局内存（global memory） = 楼下的大冷库：能装几十 GB，但跑一趟要几百个周期，带宽有限（A100 1555 GB/s）。
- 常量内存（constant memory） = 墙上贴的菜单/配方：所有人都看同一张，读起来极快（命中常量缓存约 5 周期），但只能读、不能改（从 device 侧看）。
- 局部内存（local memory） = 厨师自己的抽屉，但抽屉在楼下冷库里（物理上就是 global memory）——所以它的”局部”只是作用域局部，速度并不快，这是初学者最容易误解的地方。
- 纹理内存（texture memory） = 自助餐台：专门为”相邻线程读相邻数据 + 硬件插值”设计的只读通道，有独立缓存，适合有空间局部性的只读访问。
- 架构/机制图解：
逻辑视图（Programmer's view，讲义 Page 9 的图）：作用域与生命周期
┌──────────────────────────── Grid（一次 kernel 启动）─────────────────────────────┐
│ │
│ Global Memory ← 作用域: application(整个程序) 生命周期: application │
│ Constant Memory← 作用域: application 生命周期: application │
│ 可被 host 的 cudaMalloc/cudaMemcpy 读写与初始化 │
│ │
│ ┌──── Block (0,0) ────────────────┐ ┌──── Block (1,0) ────────────────┐ │
│ │ Shared Memory │ │ Shared Memory │ │
│ │ 作用域: block 生命周期: block │ │ 作用域: block 生命周期: block │ │
│ │ ┌─Thread(0,0)─┐ ┌─Thread(1,0)─┐ │ │ ┌─Thread(0,0)─┐ ┌─Thread(1,0)─┐ │ │
│ │ │ Registers │ │ Registers │ │ │ │ Registers │ │ Registers │ │ │
│ │ │ 作用域:thread│ │ 作用域:thread│ │ │ │ │ │ │ │ │
│ │ │ 生命周期:thread││生命周期:thread││ │ │ │ │ │ │ │
│ │ └─────────────┘ └─────────────┘ │ │ └─────────────┘ └─────────────┘ │ │
│ └─────────────────────────────────┘ └─────────────────────────────────┘ │
└─────────────────────────────────────────────────────────────────────────────────┘
▲ ▲
│ host 只能直接访问这一层 │
┌─────┴──────────────────────────────┴─────┐
│ Host (CPU + 主机 DRAM) │
└──────────────────────────────────────────┘
物理视图（芯片里的真实位置）：谁离执行单元多近、延迟多大
┌────────────────────────────── GPU 芯片 (die) ──────────────────────────────────────┐
│ │
│ ┌──────────────── SM 0 ────────────────────┐ ┌────────── SM 1 ──────────┐ │
│ │ ┌─────────────────────────────────────┐ │ │ │ │
│ │ │ Register File 256 KB (~1 cycle) │ │ │ （同左，全芯片 108 个） │ │
│ │ │ ↑ register / local(溢出时才走内存) │ │ │ │ │
│ │ └─────────────────────────────────────┘ │ │ │ │
│ │ ┌────────────────────┬────────────────┐ │ │ │ │
│ │ │ L1 / Tex Cache │ Shared Memory │ │ │ │ │
│ │ │ 命中 ~30 cycles │ ~20-30 cycles │ │ │ │ │
│ │ │ (192 KB 统一 SRAM, │ 作用域=block │ │ │ │ │
│ │ │ 最多 164KB 给 smem)│ 生命周期=block │ │ │ │ │
│ │ └────────────────────┴────────────────┘ │ │ │ │
│ │ ┌─────────────────────────────────────┐ │ │ │ │
│ │ │ Constant Cache 8 KB/SM (~5 cycle) │ │ │ │ │
│ │ └─────────────────────────────────────┘ │ │ │ │
│ └──────────────────────────────────────────┘ └──────────────────────────┘ │
│ │
│ ┌──────────────── L2 Cache (A100: 40 MB, ~200 cycles, 全 SM 共享) ──────────────┐ │
│ └───────────────────────────────────────────────────────────────────────────────┘ │
│ ┌──────────────── Global Memory / Local Memory (HBM2, 400-800 cycles) ─────────┐ │
│ │ A100: 40/80 GB, 1555 GB/s RTX 4090: 24 GB GDDR6X, 1008 GB/s │ │
│ │ 注：local memory 物理上就在 global memory 里，只是每个线程有独立私有段 │ │
│ └───────────────────────────────────────────────────────────────────────────────┘ │
│ ┌──────────────── Constant Memory 空间 (64 KB, 物理在 DRAM, 由常量缓存服务) ────┐ │
│ └───────────────────────────────────────────────────────────────────────────────┘ │
└────────────────────────────────────────────────────────────────────────────────────┘
▲
│ PCIe 4.0 x16 ≈ 25 GB/s / NVLink ≈ 600 GB/s
▼
┌──────────────────────────── Host (CPU + 主机 DRAM) ───────────────────────────────┐
│ 翻页（pageable）内存 ~6-8 GB/s；锁页（pinned）内存 ~12-25 GB/s（PCIe 4.0 x16） │
└───────────────────────────────────────────────────────────────────────────────────┘
讲义 Page 10 的变量类型修饰符总表（作用域 / 生命周期是考试重点）：
变量声明 内存空间 作用域 生命周期 物理位置
--------------------------------- -------- ---------- ---------- --------------
int LocalVar; register thread thread 寄存器文件
__device__ __shared__ int SharedVar; shared block block SM 片上 SRAM
__device__ int GlobalVar; global application application HBM
__device__ __constant__ int ConstVar; constant application application DRAM + 常量缓存
float LocalArray[16]; (带下标的自动 local thread thread HBM（溢出后！）
数组 / 寄存器不足时的溢出变量)
补充规则（讲义 Page 10）：
① 无修饰符的自动变量：基本类型与结构体 → 放进寄存器；
带下标的按线程私有数组 → 放进 global memory（即 local memory）
② __device__ 可以与 __shared__ / __constant__ 连用，
但不能单独出现在函数体内（函数体内的 __shared__ 变量默认就是 device 存储）
延迟与带宽的数量级（全课程统一口径）：
存储空间 典型延迟 相对寄存器倍率 备注
---------------- ----------------- -------------- --------------------------------
寄存器 1 cycle 1x 每线程最多 255 个
共享内存 20-30 cycles 20-30x 无 bank conflict 时；否则按冲突倍数
常量内存(命中缓存) ~5 cycles 5x 同一 warp 读同一地址时最快（广播）
L1 命中 ~30 cycles 30x 全局访存命中 L1 时
L2 命中 ~200 cycles 200x 全芯片共享 40 MB
全局内存 400-800 cycles 400-800x 一个 warp 合并访问 = 128 B = 1 cache line
= 4 个 32 B sector = 1 次 transaction
主机内存(PCIe 4.0) ~10,000+ cycles ~25 GB/s，比显存带宽慢 62 倍
复用全局内存的收益估算（讲义 Page 29 “150 GB/s Bandwidth Implies 37.5 GFLOPs”）：若每个浮点操作需要 4 字节内存（矩阵乘朴素版每个线程每对 M/N 元素做 1 次乘加 = 2 FLOP，却要读 4+4 = 8 字节，即 4 B/FLOP），则
可达性能 = 内存带宽 / 每 FLOP 需要的字节数
= 150 GB/s / (4 B/FLOP)
= 37.5 GFLOP/s ← 而同期 GPU 的计算能力是 1000 GFLOP/s
也就是说：只发挥了 3.75% 的算力，程序被内存带宽死死卡住。
实测还会更低（约 25 GFLOP/s），因为内存不可能 100% 时间都在传数据。
结论：必须大幅削减对 global memory 的访问 —— 后续 Lab 的核心思想。
在 A100 上把同样的公式重算一遍：1555 GB/s ÷ 4 B/FLOP = 388.75 GFLOP/s，而 A100 的 FP32 峰值是 19.5 TFLOPS，即只能发挥 2.0%——这就是为什么”复用内存访问”是 GPU 优化的第一性原理。
线程到数据的映射粒度：细粒度与粗粒度（Thread Mapping Granularity / Thread Coarsening）
定义与目的：细粒度（fine-grained） 指”1 个线程只处理 1 个数据元素”；粗粒度（coarse-grained） 指”1 个线程处理多个元素”。在讲义中体现为 vecAdd 的两种线程分配：
// 细粒度：一个线程一个元素，需要 ceil(N/256) 个 block vecAdd<<<ceil(N/256.0), 256>>>(...) i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) C[i] = A[i] + B[i]; // 粗粒度：一个线程两个元素，只需要 ceil(N/(2*256)) 个 block vecAdd<<<ceil(N/(2*256.0)), 256>>>(...) i = blockIdx.x * (2*blockDim.x) + threadIdx.x; if (i < n) C[i] = A[i] + B[i]; i = i + blockDim.x; if (i < n) C[i] = A[i] + B[i];目的：减少总线程数/block 数，从而摊薄每个线程的固定开销（索引运算、判断、kernel 启动与 block 调度开销），并给编译器更多指令级并行（ILP）的机会。
直观解释（”它是什么？”）：细粒度像邮局派 1000 个邮递员每人送 1 封信——派工的交通费（固定开销）比送信本身还贵。粗粒度像派 500 个邮递员每人送 2 封信，走一趟顺手送两封，派工成本减半。但如果每人送太多（比如 100 封），邮递员数量就不够，机器（SM）的并行度反而下降，还可能因为一次要拿着 100 封信而”手不够用”（寄存器溢出）。
架构/机制图解：
细粒度：1 thread = 1 element 粗粒度：1 thread = 2 elements
N = 1000, block = 256 N = 1000, block = 256
gridDim = ceil(1000/256) = 4 gridDim = ceil(1000/512) = 2
总线程 = 1024 总线程 = 512
┌─ Block 0 (i=0..255) ──┐ ┌─ Block 0 ─────────────────────────┐
│ t0 t1 t2 ... t254 t255│ │ t0 -> i=0 与 i=256 │
│ i0 i1 i2 ... i254 i255│ │ t1 -> i=1 与 i=257 │
└───────────────────────┘ │ ... │
┌─ Block 1 (i=256..511)─┐ │ t255 -> i=255 与 i=511 │
│ t0 -> i=256 │ └───────────────────────────────────┘
│ ... │ ┌─ Block 1 ─────────────────────────┐
└───────────────────────┘ │ t0 -> i=512 与 i=768 │
┌─ Block 2 (i=512..767)─┐ │ ... │
┌─ Block 3 (i=768..1023)┐ ← 24 个线程空转 │ t231 -> i=743 与 i=999 │
│ t232..t255 -> 两个都越界，空转 │
└───────────────────────────────────┘
空转 24 个线程（但只占 1 个 block 的一小部分）
合并访问仍然成立（关键！）：
细粒度 warp 0: 线程 0..31 → 地址 0..127 B → 1 条 128 B cache line
粗粒度 warp 0 第一次访问: 线程 0..31 → 地址 0..127 B → 1 条 cache line
粗粒度 warp 0 第二次访问: 线程 0..31 → 地址 256..383 B → 1 条 cache line
两次访问各自合并，只是"跨步"了 blockDim.x 个元素
性能数字（N = 1000 万，A100，block = 256）：
配置 线程数 每线程元素 索引计算次数 总指令数(粗估) 实测带宽
-------------- --------- ---------- ------------ ------------- ---------
细粒度 10,000,000 1 1 次/线程 ~5 条/元素 高
粗粒度(2 元素) 5,000,000 2 2 次/线程 ~4 条/元素 略高
粗粒度(4 元素) 2,500,000 4 4 次/线程 ~3.5 条/元素 更高
粗粒度(8 元素) 1,250,000 8 8 次/线程 ~3.2 条/元素 开始下降
粗粒度(16 元素) 625,000 16 16 次/线程 ~3.1 条/元素 下降(并行度不足)
→ 甜点区通常在每线程 2-8 个元素：既摊薄了开销，又保住了足够的 warp 数来隐藏延迟
→ 判断依据：线程总数至少要 >= 2 * SM 数 * 每 SM 线程上限，A100 即 2*108*2048 = 442K 线程
同步原语：__syncthreads() 与 cudaDeviceSynchronize()
定义与目的：GPU 上没有”全局屏障”这种廉价操作，所以 CUDA 把同步分成两级：① block 级——
__syncthreads()是一个 barrier（屏障），调用它的 block 内所有线程必须都到达这一点，才允许任何一个继续执行；它保证该点之前对 shared memory 和 global memory 的写对 block 内所有线程可见。② host-device 级——cudaDeviceSynchronize()让主机阻塞等待，直到此前提交到该设备的所有工作（所有 kernel、所有 memcpy）全部完成。直观解释（”它是什么？”）：
__syncthreads()像小组讨论里的“大家停一下，等所有人都把白板写完了再一起看”——不等到齐不许动。cudaDeviceSynchronize()像项目经理等整个车间下班才开车走。为什么需要前者？因为共享内存是 block 内所有线程共同的案板，如果 A 还没把菜放上去，B 就去拿，拿到的就是垃圾数据。架构/机制图解：
__syncthreads() 的行为（block 内屏障）
┌──────────────────────── Block (256 线程, 8 warp) ────────────────────────┐
│ │
│ warp0 ──■──┐ │
│ warp1 ──■──┤ │
│ warp2 ──■──┼──► 全部到达 ■ 之后才统一放行 ──► 后续代码 │
│ ... │ │
│ warp7 ──■──┘ │
│ ▲ │
│ barrier: 到达者计数 = blockDim，归零则全体放行 │
│ │
│ 没有 __syncthreads() 时的竞态（race）： │
│ 时刻 t1: warp0 写 shared[0..31] = in[...] │
│ 时刻 t1: warp7 读 shared[0..31] ← 可能读到还没写的旧值/垃圾 │
│ 时刻 t2: warp0 才写完 │
└─────────────────────────────────────────────────────────────────────────┘
两级同步的适用范围
┌────────────────────────────┬──────────────────────────┬─────────────────────┐
│ 同步方式 │ 作用范围 │ 谁在等待 │
├────────────────────────────┼──────────────────────────┼─────────────────────┤
│ __syncthreads() │ 同一个 block 内所有线程 │ device 线程 │
│ __syncwarp(mask) │ 同一个 warp 内指定 lane │ device 线程 │
│ atomic 操作 │ 同一地址的读改写原子性 │ device 线程（不阻塞） │
│ cudaDeviceSynchronize() │ 整个设备上所有已提交的工作 │ host 线程（阻塞） │
│ cudaStreamSynchronize(s) │ 某个 stream 上所有工作 │ host 线程（阻塞） │
│ cudaMemcpy(D2H) 默认同步 │ 隐式同步该次拷贝 │ host 线程（阻塞） │
└────────────────────────────┴──────────────────────────┴─────────────────────┘
★ 关键：没有任何原语能同步"不同 block"！不同 block 之间只能靠
kernel 结束（或 cooperative groups 的 grid.sync()，超出本讲范围）。
性能数字：__syncthreads() 本身开销很小（几十个周期），但它会让整个 block 等待最慢的 warp。如果各 warp 工作量不均（负载不均 load imbalance），屏障会把最慢的 warp 的耗时强加给所有人。例如 8 个 warp 中 7 个用 100 周期到达屏障、1 个因为一次 L2 miss 用了 400 周期，则整个 block 白等 300 周期，约等于损失 8 × 300 = 2400 warp-周期的发射机会。因此屏障应尽量少用，且最好放在循环外或负载均衡的位置。
另外，__syncthreads() 必须被 block 内所有线程以相同次数执行——把它放进 if (threadIdx.x < 16) { ... __syncthreads(); } 这类分支里，会导致某些 warp 永远等不到同伴（在 Volta 之前直接死锁/未定义行为；Volta 之后官方仍然明确规定这是未定义行为）。
代码示例与性能分析
程序 1（Lab 0）：device_query.cu —— 读懂你手里的 GPU
// 文件: device_query.cu
// 编译: nvcc -O3 -arch=sm_80 device_query.cu -o device_query
// 运行: ./device_query
//
// 对应 Lab 0 (Device Query)：把后面所有占用率计算所需的硬件参数一次性打印出来。
#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>
#define CUDA_CHECK(call) \
do { \
cudaError_t err__ = (call); \
if (err__ != cudaSuccess) { \
fprintf(stderr, "CUDA error %s:%d: %s\n", __FILE__, __LINE__, \
cudaGetErrorString(err__)); \
exit(EXIT_FAILURE); \
} \
} while (0)
static void printDeviceInfo(int dev)
{
cudaDeviceProp p;
CUDA_CHECK(cudaGetDeviceProperties(&p, dev));
int smemPerSM = 0, smemPerBlock = 0, maxBlocksPerSM = 0;
CUDA_CHECK(cudaDeviceGetAttribute(&smemPerSM,
cudaDevAttrMaxSharedMemoryPerMultiprocessor, dev));
CUDA_CHECK(cudaDeviceGetAttribute(&smemPerBlock,
cudaDevAttrMaxSharedMemoryPerBlock, dev));
CUDA_CHECK(cudaDeviceGetAttribute(&maxBlocksPerSM,
cudaDevAttrMaxBlocksPerMultiprocessor, dev));
printf("\n================ Device %d: %s ================\n", dev, p.name);
printf("compute capability : %d.%d (编译时用 -arch=sm_%d%d)\n",
p.major, p.minor, p.major, p.minor);
printf("multiProcessorCount (SM 数) : %d\n", p.multiProcessorCount);
printf("maxThreadsPerMultiProcessor : %d\n", p.maxThreadsPerMultiProcessor);
printf("maxThreadsPerBlock : %d\n", p.maxThreadsPerBlock);
printf("maxBlocksPerMultiProcessor : %d\n", maxBlocksPerSM);
printf("warpSize : %d\n", p.warpSize);
printf("regsPerMultiprocessor : %d (= %d KB)\n",
p.regsPerMultiprocessor, p.regsPerMultiprocessor * 4 / 1024);
printf("regsPerBlock : %d\n", p.regsPerBlock);
printf("sharedMemPerMultiprocessor : %d B (= %d KB)\n",
smemPerSM, smemPerSM / 1024);
printf("sharedMemPerBlock : %d B (= %d KB)\n",
smemPerBlock, smemPerBlock / 1024);
printf("maxGridSize (x, y, z) : (%d, %d, %d)\n",
p.maxGridSize[0], p.maxGridSize[1], p.maxGridSize[2]);
printf("maxThreadsDim (x, y, z) : (%d, %d, %d)\n",
p.maxThreadsDim[0], p.maxThreadsDim[1], p.maxThreadsDim[2]);
printf("l2CacheSize : %d B (= %d KB)\n",
p.l2CacheSize, p.l2CacheSize / 1024);
printf("memoryBusWidth / memoryClock : %d bit / %d kHz\n",
p.memoryBusWidth, p.memoryClockRate);
printf("totalGlobalMem : %.2f GB\n",
(double)p.totalGlobalMem / (1024.0 * 1024.0 * 1024.0));
printf("concurrentKernels : %d\n", p.concurrentKernels);
printf("unifiedAddressing : %d\n", p.unifiedAddressing);
double bw = 2.0 * (double)p.memoryClockRate * ((double)p.memoryBusWidth / 8.0)
/ 1.0e6; // DDR: 每个时钟周期传 2 次 -> GB/s
printf("theoretical DRAM bandwidth : %.1f GB/s\n", bw);
// 用读到的参数直接算出这台机器上一个 256 线程 block 的驻留上限
int blocksByThreads = p.maxThreadsPerMultiProcessor / 256;
int blocksByBlocks = maxBlocksPerSM;
int resident = blocksByThreads < blocksByBlocks ? blocksByThreads
: blocksByBlocks;
printf("---- 256 线程/block 时的驻留上限（不含寄存器与 smem 限制）----\n");
printf(" 受线程数限制 : %d blocks\n", blocksByThreads);
printf(" 受 block 数限制: %d blocks\n", blocksByBlocks);
printf(" => 实际驻留 : %d blocks = %d threads = %d warps\n",
resident, resident * 256, resident * 256 / p.warpSize);
printf(" => 占用率 : %.1f%%\n",
100.0 * resident * 256 / p.maxThreadsPerMultiProcessor);
}
int main(void)
{
int devCount = 0;
CUDA_CHECK(cudaGetDeviceCount(&devCount));
printf("CUDA devices found: %d\n", devCount);
if (devCount == 0) {
printf("No CUDA-capable device available.\n");
return EXIT_SUCCESS;
}
for (int d = 0; d < devCount; ++d) {
printDeviceInfo(d);
}
int cur = 0;
CUDA_CHECK(cudaGetDevice(&cur));
printf("\nCurrent device: %d\n", cur);
CUDA_CHECK(cudaDeviceReset());
printf("Device query done.\n");
return EXIT_SUCCESS;
}
- 【代码做什么？】
cudaGetDeviceCount拿到机器上可用的 GPU 数量；cudaGetDeviceProperties一次性填满cudaDeviceProp，其中就包含后面做占用率计算必需的multiProcessorCount、maxThreadsPerMultiProcessor、regsPerMultiprocessor、sharedMemPerMultiprocessor、warpSize。- 有三个属性（每 SM 最大共享内存、每 block 最大共享内存、每 SM 最大 block 数）在老的
cudaDeviceProp里没有直接成员或容易与”capability 相关数值”混淆，所以用cudaDeviceGetAttribute按枚举名精确获取。 - 用
2 × memoryClockRate(kHz) × memoryBusWidth(bytes) / 1e6估算理论峰值带宽：memoryClockRate单位是 kHz，乘busWidth/8得到字节/秒，再乘 2 是因为 GDDR/HBM 在一个时钟周期内传两次数据。A100 上是2 × 1215 MHz × 512 B = 1244 GB/s（官方标称 boost 后约 1555 GB/s，差异来自 boost 时钟）。 - 最后用读到的参数当场算一遍占用率：256 线程的 block，先看线程数允许几个 block，再看硬件 block 数上限允许几个，取小者；再换算成 warp 数与百分比。
cudaDeviceReset()显式归还所有设备资源，等价于进程退出时的隐式清理，但在同一进程里反复跑多轮测试时可以避免状态残留。
- 【并行机制与硬件映射解说】
- 这个程序本身不启动任何 kernel，所以没有 warp、没有 SM 调度。它属于”元信息程序”：它的价值在于把后续所有代码的调优参数从”查文档”变成”运行时读出来”。同一个二进制在 A100（108 SM）、RTX 4090（128 SM）、RTX 2080 Ti（68 SM）上跑出的数字完全不同，硬编码常量必然出错。
- 由打印结果可以立刻推出并行槽位规模：A100 = 108 SM × 2048 线程 = 221,184 个同时驻留线程 = 108 × 64 = 6912 个同时驻留 warp；RTX 4090 = 128 × 1536 = 196,608 线程 = 128 × 48 = 6144 个 warp。这就是”要跑满 GPU 至少需要多少线程”的答案。
- 如果某个 kernel 的 grid 只有 1 个 block（例如
kernel<<<1, 256>>>），那么无论 GPU 有多少 SM，都只有 1 个 SM 在干活，其余 107 个 SM 全部空闲——利用率 1/108 ≈ 0.93%。这是最典型的”GPU 加速比不满意”的原因。 - 观察
maxThreadsDim = (1024, 1024, 64)：x×y×z 的总乘积不得超过 1024，所以dim3(16,16)合法（256），dim3(64,64)非法（4096）；而maxGridSize.y = 65535意味着二维 grid 的 y 方向最多 65535 个 block，处理超大图像时若H/16 > 65535（即 H > 1,048,560 像素）就必须改用一维 grid 手工切分。
- 【性能优化分析】
- 占用率（occupancy）：占用率 = 实际驻留 warp 数 ÷ 硬件上限 =
(驻留 block 数 × block 线程数 ÷ 32) ÷ (maxThreadsPerMultiProcessor ÷ 32)。对 A100、256 线程 block：2048/256 = 8个 block，恰好8×256 = 2048线程 = 64 warp = 100%。但这是”不含寄存器/共享内存约束”的理论上限，真实值还要满足：寄存器约束: 驻留 block 数 <= floor(regsPerSM / (numRegs * threadsPerBlock 对齐到 256)) 共享内存约束: 驻留 block 数 <= floor(smemPerSM / smemPerBlock) 最终上限 = min(线程约束, 寄存器约束, smem 约束, maxBlocksPerSM) - 算术强度与 Roofline：本程序不涉及计算，只给出后续分析的”机器常数”：
A100 脊点 (ridge point) = 峰值算力 / 峰值带宽 = 19.5e12 / 1555e9 = 12.5 FLOP/Byte RTX 4090 脊点 = 82.6e12 / 1008e9 = 81.9 FLOP/Byte 含义：算术强度 < 12.5 FLOP/B 的 kernel 在 A100 上必然带宽受限； 而同一 kernel 在 RTX 4090 上要 > 81.9 FLOP/B 才算计算受限。 → 4090 的"内存墙"相对更高，因为它的算力涨得比带宽快得多。 - 瓶颈判定与优化方向：Lab 0 的”性能目标”不是跑得快，而是建立正确的机器模型。用这些数字做两件事：① 由
最大驻留线程数决定 grid 至少要有多少线程（建议 ≥ 2× 驻留量以覆盖尾部和调度抖动，A100 约 44 万线程）；② 由理论带宽作为 Roofline 的屋顶，Lab 1 的向量加法实测带宽一般能达到理论值的 80–92%，达不到就说明访存没合并好或并行度不足。
- 占用率（occupancy）：占用率 = 实际驻留 warp 数 ÷ 硬件上限 =
程序 2（Lab 1 风格）：rgba_channel_split.cu —— 二维线程索引与合并访问
// 文件: rgba_channel_split.cu
// 编译: nvcc -O3 -arch=sm_80 rgba_channel_split.cu -o rgba_split
// 运行: ./rgba_split (默认 3840x2160 的 4K 图)
// ./rgba_split 76 62 (讲义里的 76x62 小图，看网格如何向上取整)
//
// 对应内容: Lecture 3 的二维线程映射 (colorToGreyscaleConversion) + Lab 1 的向量化访存思想。
// 功能: 把 RGBA 交错存储的彩色图拆成 R/G/B/A 四个独立通道平面 (planar 格式)。
#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>
#define CUDA_CHECK(call) \
do { \
cudaError_t err__ = (call); \
if (err__ != cudaSuccess) { \
fprintf(stderr, "CUDA error %s:%d: %s\n", __FILE__, __LINE__, \
cudaGetErrorString(err__)); \
exit(EXIT_FAILURE); \
} \
} while (0)
#define BLOCK_X 16
#define BLOCK_Y 16
// ---------------- 版本 A：朴素版，4 次单字节 load（每次跨步 4 字节） ----------------
__global__ void rgbaSplitNaive(const unsigned char* __restrict__ rgba,
unsigned char* __restrict__ r,
unsigned char* __restrict__ g,
unsigned char* __restrict__ b,
unsigned char* __restrict__ a,
int width, int height)
{
int col = blockIdx.x * blockDim.x + threadIdx.x; // 列号
int row = blockIdx.y * blockDim.y + threadIdx.y; // 行号
if (row < height && col < width) { // 边界检查：挡住越界线程
int pix = row * width + col; // 像素的线性编号（行主序）
int o = 4 * pix; // RGBA 平面中的字节偏移
r[pix] = rgba[o + 0];
g[pix] = rgba[o + 1];
b[pix] = rgba[o + 2];
a[pix] = rgba[o + 3];
}
}
// ---------------- 版本 B：向量化版，1 次 uchar4 load 取回全部 4 个通道 ----------------
__global__ void rgbaSplitVec(const uchar4* __restrict__ rgba,
unsigned char* __restrict__ r,
unsigned char* __restrict__ g,
unsigned char* __restrict__ b,
unsigned char* __restrict__ a,
int width, int height)
{
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
if (row < height && col < width) {
int pix = row * width + col;
uchar4 p = rgba[pix]; // 一条 4 字节向量 load，warp 内正好拼成 128 字节
r[pix] = p.x;
g[pix] = p.y;
b[pix] = p.z;
a[pix] = p.w;
}
}
// ---------------- CPU 参考实现 ----------------
static void rgbaSplitCPU(const unsigned char* rgba,
unsigned char* r, unsigned char* g,
unsigned char* b, unsigned char* a,
int width, int height)
{
long n = (long)width * (long)height;
for (long i = 0; i < n; ++i) {
r[i] = rgba[4 * i + 0];
g[i] = rgba[4 * i + 1];
b[i] = rgba[4 * i + 2];
a[i] = rgba[4 * i + 3];
}
}
static int verify(const unsigned char* ref, const unsigned char* got, long n,
const char* name)
{
for (long i = 0; i < n; ++i) {
if (ref[i] != got[i]) {
printf(" [FAIL] %s: index %ld ref=%u got=%u\n",
name, i, (unsigned)ref[i], (unsigned)got[i]);
return 0;
}
}
printf(" [PASS] %s: %ld bytes identical to CPU reference\n", name, n);
return 1;
}
int main(int argc, char** argv)
{
int width = (argc > 1) ? atoi(argv[1]) : 3840;
int height = (argc > 2) ? atoi(argv[2]) : 2160;
long npix = (long)width * (long)height;
long nbyte = npix; // 每个通道平面的字节数
printf("Image: %d x %d = %ld pixels\n", width, height, npix);
// ---- 网格配置：核心是"向上取整"，否则图像右侧/底部会漏掉 ----
dim3 block(BLOCK_X, BLOCK_Y);
dim3 grid((width + BLOCK_X - 1) / BLOCK_X, // = ceil(width /16)
(height + BLOCK_Y - 1) / BLOCK_Y, // = ceil(height/16)
1);
long launchedThreads = (long)grid.x * grid.y * block.x * block.y;
printf("Grid : (%u, %u) blocks of (%u, %u) threads\n",
grid.x, grid.y, block.x, block.y);
printf("Total blocks = %u x %u = %u ; total threads = %ld ; useful = %ld (%.2f%%)\n",
grid.x, grid.y, grid.x * grid.y, launchedThreads, npix,
100.0 * (double)npix / (double)launchedThreads);
// ---- 主机端分配 ----
unsigned char* h_rgba = (unsigned char*)malloc(4 * nbyte);
unsigned char* h_r = (unsigned char*)malloc(nbyte);
unsigned char* h_g = (unsigned char*)malloc(nbyte);
unsigned char* h_b = (unsigned char*)malloc(nbyte);
unsigned char* h_a = (unsigned char*)malloc(nbyte);
unsigned char* c_r = (unsigned char*)malloc(nbyte);
unsigned char* c_g = (unsigned char*)malloc(nbyte);
unsigned char* c_b = (unsigned char*)malloc(nbyte);
unsigned char* c_a = (unsigned char*)malloc(nbyte);
if (!h_rgba || !h_r || !h_g || !h_b || !h_a || !c_r || !c_g || !c_b || !c_a) {
fprintf(stderr, "host malloc failed\n");
return EXIT_FAILURE;
}
for (long i = 0; i < 4 * nbyte; ++i) {
h_rgba[i] = (unsigned char)((i * 37 + 11) & 0xFF); // 确定性伪随机
}
// ---- 设备端分配 ----
unsigned char *d_rgba = NULL, *d_r = NULL, *d_g = NULL, *d_b = NULL, *d_a = NULL;
CUDA_CHECK(cudaMalloc((void**)&d_rgba, 4 * nbyte));
CUDA_CHECK(cudaMalloc((void**)&d_r, nbyte));
CUDA_CHECK(cudaMalloc((void**)&d_g, nbyte));
CUDA_CHECK(cudaMalloc((void**)&d_b, nbyte));
CUDA_CHECK(cudaMalloc((void**)&d_a, nbyte));
CUDA_CHECK(cudaMemcpy(d_rgba, h_rgba, 4 * nbyte, cudaMemcpyHostToDevice));
// ---- 计时 ----
cudaEvent_t t0, t1;
CUDA_CHECK(cudaEventCreate(&t0));
CUDA_CHECK(cudaEventCreate(&t1));
double traffic = (double)(4 * nbyte + 4 * nbyte); // 读 4B/像素 + 写 4B/像素
// ==== 版本 A ====
CUDA_CHECK(cudaEventRecord(t0));
rgbaSplitNaive<<<grid, block>>>(d_rgba, d_r, d_g, d_b, d_a, width, height);
CUDA_CHECK(cudaEventRecord(t1));
CUDA_CHECK(cudaEventSynchronize(t1));
float msA = 0.0f;
CUDA_CHECK(cudaEventElapsedTime(&msA, t0, t1));
CUDA_CHECK(cudaMemcpy(c_r, d_r, nbyte, cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaMemcpy(c_g, d_g, nbyte, cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaMemcpy(c_b, d_b, nbyte, cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaMemcpy(c_a, d_a, nbyte, cudaMemcpyDeviceToHost));
printf("\n--- Version A: naive (4 byte loads, stride 4) ---\n");
printf(" time = %.3f ms effective BW = %.1f GB/s\n",
msA, traffic / (msA * 1.0e-3) / 1.0e9);
int okA = verify(h_rgba + 0, c_r, nbyte, "R plane") &
verify(h_rgba + 1, c_g, nbyte, "G plane") &
verify(h_rgba + 2, c_b, nbyte, "B plane") &
verify(h_rgba + 3, c_a, nbyte, "A plane");
(void)okA;
// ==== 版本 B ====
CUDA_CHECK(cudaMemset(d_r, 0, nbyte));
CUDA_CHECK(cudaMemset(d_g, 0, nbyte));
CUDA_CHECK(cudaMemset(d_b, 0, nbyte));
CUDA_CHECK(cudaMemset(d_a, 0, nbyte));
CUDA_CHECK(cudaEventRecord(t0));
rgbaSplitVec<<<grid, block>>>((const uchar4*)d_rgba, d_r, d_g, d_b, d_a,
width, height);
CUDA_CHECK(cudaEventRecord(t1));
CUDA_CHECK(cudaEventSynchronize(t1));
float msB = 0.0f;
CUDA_CHECK(cudaEventElapsedTime(&msB, t0, t1));
CUDA_CHECK(cudaMemcpy(c_r, d_r, nbyte, cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaMemcpy(c_g, d_g, nbyte, cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaMemcpy(c_b, d_b, nbyte, cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaMemcpy(c_a, d_a, nbyte, cudaMemcpyDeviceToHost));
printf("\n--- Version B: vectorized (uchar4 load) ---\n");
printf(" time = %.3f ms effective BW = %.1f GB/s\n",
msB, traffic / (msB * 1.0e-3) / 1.0e9);
int okB = verify(h_rgba + 0, c_r, nbyte, "R plane") &
verify(h_rgba + 1, c_g, nbyte, "G plane") &
verify(h_rgba + 2, c_b, nbyte, "B plane") &
verify(h_rgba + 3, c_a, nbyte, "A plane");
(void)okB;
printf("\nSpeedup (A/B) = %.2fx\n", msA / msB);
// ---- CPU 参考结果（抽样验证，避免小图时 CPU 太慢） ----
rgbaSplitCPU(h_rgba, h_r, h_g, h_b, h_a, width, height);
printf("\nCPU reference check: %s\n",
(h_r[123 % nbyte] == h_rgba[4 * (123 % nbyte) + 0]) ? "consistent" : "MISMATCH");
// ---- 释放 ----
CUDA_CHECK(cudaEventDestroy(t0));
CUDA_CHECK(cudaEventDestroy(t1));
CUDA_CHECK(cudaFree(d_rgba));
CUDA_CHECK(cudaFree(d_r));
CUDA_CHECK(cudaFree(d_g));
CUDA_CHECK(cudaFree(d_b));
CUDA_CHECK(cudaFree(d_a));
free(h_rgba); free(h_r); free(h_g); free(h_b); free(h_a);
free(c_r); free(c_g); free(c_b); free(c_a);
CUDA_CHECK(cudaDeviceReset());
printf("\nDone.\n");
return EXIT_SUCCESS;
}
- 【代码做什么？】
- 决定网格形状。
dim3 block(16,16)一共 256 个线程；dim3 grid((W+15)/16, (H+15)/16)用整数向上取整覆盖整个图像。以讲义里的 76×62 为例：ceil(76/16) = 5、ceil(62/16) = 4，共 5×4 = 20 个 block，启动 20×256 = 5120 个线程，而图像只有 76×62 = 4712 个像素，408 个线程会被边界检查挡住（浪费 7.97%）。 - 每个线程算出自己负责的像素。
col由blockIdx.x与threadIdx.x合成，row由blockIdx.y与threadIdx.y合成——注意 x 管列（横向）、y 管行（纵向），x 变化最快，正好对应 C 的行主序内存布局。 - 边界检查
if (row < height && col < width)。因为网格是向上取整的，最右一列 block 的 col 可能 ≥ width（76×62 例子中 col 可取到 64…79，而 width=76，所以 col=76…79 的 4 列越界），最下一行 block 的 row 可能 ≥ height（row 可取到 48…63，而 height=62，所以 row=62、63 越界）。没有这个 if 就会写坏别人的内存（静默错误）。 - 算出字节偏移。
pix = row*width + col是像素的线性编号（行主序）；交错存储的 RGBA 中第 pix 个像素的 4 个字节位于4*pix + {0,1,2,3}；输出平面中第 pix 个字节位于pix。 - 两个版本做同一件事：A 版读 4 次单字节，B 版读 1 次
uchar4（4 字节向量）。二者的算法完全相同，差别只在内存指令的粒度。 - 主机端调度。先用
cudaMalloc申请 5 块显存，cudaMemcpy(H2D)把交错图送进去，启动 kernel，cudaEventRecord/cudaEventSynchronize计时，再把 4 个平面cudaMemcpy(D2H)取回，与 CPU 参考实现逐字节比对，最后cudaFree/cudaDeviceReset收尾。 - 计时必须包住同步点。kernel 启动是异步的，如果只测
cudaEventRecord(t0); kernel<<<>>>; cudaEventRecord(t1);而不做cudaEventSynchronize(t1)，读到的可能是”启动开销”而不是”执行时间”。
- 决定网格形状。
- 【并行机制与硬件映射解说】
- block → warp 的划分：
block = (16,16)共 256 线程 = 8 个 warp。线性化规则是linear = threadIdx.x + threadIdx.y * blockDim.x（x 最快），所以：warp 0 : linear 0..31 = threadIdx.y=0 (x=0..15) 与 threadIdx.y=1 (x=0..15) → 覆盖第 0、1 两行的各 16 个像素 warp 1 : linear 32..63 = threadIdx.y=2 与 y=3 两行 ... warp 7 : linear 224..255 = threadIdx.y=14 与 y=15 两行 → 每个 warp 横跨 2 行 × 16 列 = 32 个像素，行内的 16 个线程地址连续 - warp 到 SM 的调度：A100 有 4 个 warp scheduler/SM，每周期最多发射 4 条 warp 指令。这个 kernel 每个线程约用 18–24 个寄存器（
col/row/pix/o/p加上地址计算），以 20 个估算：每 block 寄存器 =256 × 20 = 5120，65536 / 5120 = 12.8 → 12个 block；线程上限给2048/256 = 8个 block；共享内存用量为 0；A100 每 SM 最多 32 个 block。取最小值 8 个 block/SM = 2048 线程 = 64 warp = 100% 占用率（A100 上限 64 warp）。占用率计算（A100, block=16x16, 20 regs/thread, 0 smem） blocks_by_threads = 2048 / 256 = 8 blocks_by_regs = 65536 / (256*20) = 65536/5120 = 12 (向下取整) blocks_by_smem = 164KB / 0 = 无限制 blocks_by_hw = 32 resident = min(8, 12, inf, 32) = 8 occupancy = (8*256)/2048 = 100% RTX 4090 (48 warp/SM, 1536 线程/SM, 100 KB smem) blocks_by_threads = 1536 / 256 = 6 blocks_by_regs = 65536 / 5120 = 12 resident = min(6,12,inf,32) = 6 → 6*256 = 1536 线程 = 48 warp = 100% - 共享内存与 bank conflict：本 kernel 完全不用 shared memory（数据没有复用，每个像素只被读一次），因此不存在 bank conflict。作为对照，如果误写成
__shared__ unsigned char sTile[256];再让线程t访问sTile[t]，由于 bank 宽度是 4 字节而元素是 1 字节，线程 t=0..3 会撞到同一个 bank 0，形成 4-way bank conflict，把本来 1 个周期的访问拖成 4 个周期（这正是字节型共享数组必须”padding 或按 int 组织”的原因）。 - 全局内存访问是否合并（coalesced）——这是本程序的核心：
【版本 A：朴素版，字节 load，跨步 4 字节】 warp 内 32 个线程读 rgba[4*(base+j) + c] (j = 0..31, c 固定 0/1/2/3) 字节地址: 4*base+c, 4*base+4+c, ..., 4*base+124+c 地址跨度: 128 字节 = 4 个 sector (每 sector 32 B) = 1 条 cache line 这 128 字节里本次指令真正用到的只有 32 字节 → sector 利用率 = 32/128 = 25% 4 条 load 指令各自都要走一遍这 128 字节（后 3 条通常命中 L1） → 对 L1/LSU 的压力是向量化版的 4 倍；每条 load 需要 4 个 sector wavefront → DRAM 侧因 L1 命中，实际流量接近 8 B/像素，但 L1 吞吐/MSHR 很快成为瓶颈， 实测带宽明显低于峰值（典型只有峰值的 30%~50%） 【版本 B：向量化版，uchar4 load】 warp 内 32 个线程读 rgba[base+j] (uchar4 = 4 字节) 字节地址: 4*base, 4*base+4, ..., 4*base+124 → 32 * 4 = 128 字节，且完全连续 = 1 条完整 cache line = 4 个 sector → 128/128 = 100% sector 利用率，一次 transaction 取回全部有用数据 ★ 【写侧：r[pix] = ...】 warp 内连续 pix → 写 32 个连续字节 = 恰好 1 个 32 字节 sector → 写合并良好（写最小粒度就是 sector，32 B 全用满） - warp 发散：kernel 里唯一的条件是
if (row < height && col < width)。对绝大多数 block 而言 32 个 lane 全部为真，几乎零发散；只有在图像最右一列（col ∈ [width, grid.x16)）和最下一行（row ∈ [height, grid.y16)）的 block 上，同一个 warp 内会出现部分 lane 为假的情况。以 76×62 为例：右边界 block（blockIdx.x = 4）覆盖 col 64…79，其中 col ≥ 76 的 4 列无效，而一个 warp 横跨 2 行 × 16 列 → 每行 16 个线程中有 4 个无效，发散比例为 4/16 = 25%，仅影响那 5 列中的 1 列 block，总体影响 < 2%。判断准则（讲义 Problem Solving）：发散的前提是”执行路径依赖线程独有的量”；下面这段代码因为q对所有线程相同（是 kernel 参数），不会发散：__global__ void do_work(int q, int *A) { int result = 0; if (q < 5) result = threadIdx.x; // q 与 threadIdx 无关 → 全 block 一致 → 无发散 A[threadIdx.x] = result; } - 寄存器使用量：朴素版因为要同时维护 4 个字节的地址计算，寄存器可能略高（约 20–24）；向量化版把 4 个字节打包进一个
uchar4寄存器对，寄存器数通常下降 2–4 个。若寄存器数超过阈值导致占用率掉档（例如从 20 涨到 33，65536/(256*33) = 7个 block → 87.5%），可以用__launch_bounds__(256, 8)让编译器主动限制寄存器用量。
- block → warp 的划分：
- 【性能优化分析】
- 算术强度（arithmetic intensity）：本 kernel 每个像素的”有用计算量”近似为 0（就是搬字节，至多算 4 次地址加法，按 4 FLOP/像素计），而内存流量是实打实的 8 字节/像素（读 4 + 写 4）：
AI = FLOPs / Bytes = 4 / 8 = 0.5 FLOP/Byte （若完全不计运算则为 ~0） - Roofline 模型：
性能上限 = min( 峰值算力, AI × 峰值带宽 ) = min( 19.5e12 FLOP/s , 0.5 × 1555e9 B/s ) = min( 19 500 GFLOP/s , 777 GFLOP/s ) = 777 GFLOP/s ← 比峰值算力低 25 倍 结论：纯带宽受限（memory-bound），与"算得快不快"完全无关。 - 代入具体数字的理想时间（4K 图，3840×2160 = 8,294,400 像素）：
总流量 = 8,294,400 × 8 B = 66.36 MB A100 理想时间 = 66.36e6 B / 1555e9 B/s = 42.7 µs → 100% 峰值 达到 80% 峰值 → 53.4 µs；达到 50% → 85.3 µs RTX 4090 理想时间 = 66.36e6 / 1008e9 = 65.8 µs 小图 76x62：总流量 = 4712 × 8 = 37.7 KB → 理想 24 ns， 远小于 kernel 启动开销(3-10 µs)，此时"性能"只反映启动开销。 - 瓶颈判定：内存带宽受限（AI = 0.5 « A100 脊点 12.5 FLOP/B）。因此优化方向只有三条：① 减少字节数；② 提高每次访存的字节效率（向量化）；③ 提高并行度以逼近峰值带宽。具体到本程序：
① 减少字节数：A 通道若下游不需要，读 uchar3 仍然要拉 4 字节的 cache line （RGB 交错时更糟：uchar3 是 3 字节对齐，warp 读 3B 跨步会覆盖 96 字节 但每条指令只用 32 字节 → 33% 效率）；可以考虑只做灰度化， 输出 1 字节/像素 → 流量从 8 B/像素降到 5 B/像素，理论提速 1.6x。 ② 提高访存效率：用 uchar4/float4 让 warp 访问 = 128 B 完整 cache line（本程序 B 版） 实测提速通常在 1.5x ~ 3x（取决于 L1 压力）。 ③ 提高并行度：4K 图有 32,400 个 block，远超 A100 的 108 SM × 8 block = 864 个驻留槽位，并行度充足；小图（76×62 只有 20 个 block）则 GPU 基本闲置， 此时应改用"一个线程处理多个像素"的粗粒度方案，或者干脆用 CPU。
- 算术强度（arithmetic intensity）：本 kernel 每个像素的”有用计算量”近似为 0（就是搬字节，至多算 4 次地址加法，按 4 FLOP/像素计），而内存流量是实打实的 8 字节/像素（读 4 + 写 4）：
程序 3：mem_space_demo.cu —— 六级内存空间的作用域与生命周期
// 文件: mem_space_demo.cu
// 编译: nvcc -O3 -arch=sm_80 mem_space_demo.cu -o mem_space_demo
// 运行: ./mem_space_demo
//
// 在一个 kernel 里同时使用 register / shared / constant / global(__device__)，
// 并在主机端打印 cudaFuncGetAttributes 得到的寄存器数与共享内存使用量。
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <cuda_runtime.h>
#define CUDA_CHECK(call) \
do { \
cudaError_t err__ = (call); \
if (err__ != cudaSuccess) { \
fprintf(stderr, "CUDA error %s:%d: %s\n", __FILE__, __LINE__, \
cudaGetErrorString(err__)); \
exit(EXIT_FAILURE); \
} \
} while (0)
#define N 4096
#define BLOCK 256
// ============ 四种存储空间的声明 ============
// (1) __device__ 变量 -> global memory，作用域 application，生命周期 application
__device__ float gScratch[N];
__device__ int gBlockCount = 0;
// (2) __constant__ 变量 -> constant memory，作用域 application，生命周期 application
// 可以在声明处初始化（编译期写入常量段），运行时也能被 host 改写
__constant__ float cWeight[4] = {0.5f, 1.0f, 1.5f, 2.0f};
__global__ void memSpaceDemo(const float* __restrict__ in,
float* __restrict__ out,
int n)
{
// (3) __shared__ 静态数组 -> shared memory，作用域 block，生命周期 block
__shared__ float sData[BLOCK];
// (4) extern __shared__ 动态数组 -> shared memory，大小由启动配置给出
extern __shared__ float sDyn[];
// (5) 无修饰符的自动变量 -> register，作用域 thread，生命周期 thread
int tid = blockIdx.x * blockDim.x + threadIdx.x;
float self = (tid < n) ? in[tid] : 0.0f; // 读 global memory（也可能是 L1 命中）
float acc = 0.0f;
float sum = 0.0f;
// ---- 把数据搬进 shared memory ----
sData[threadIdx.x] = self;
__syncthreads(); // 保证整个 block 的 sData 都写完了
// ---- 读 constant memory（只读，走常量缓存）----
float w = cWeight[threadIdx.x & 3];
// ---- 读邻居：block 内跨线程通信只能靠 shared memory ----
float nb = sData[(threadIdx.x + 1) % blockDim.x];
acc = w * self + 0.5f * nb; // 全部在寄存器里算
sDyn[threadIdx.x] = acc; // 写动态共享内存
__syncthreads(); // 保证 sDyn 全部就绪
// ---- 反复复用 shared memory 里的数据（这就是 shared memory 存在的意义）----
for (int k = 0; k < 4; ++k) {
sum += sDyn[(threadIdx.x + k * 7) % blockDim.x] * cWeight[k];
}
if (tid < n) {
out[tid] = sum; // 写 global memory
}
atomicAdd(&gScratch[blockIdx.x], sum); // 写 __device__ 全局变量（需原子）
if (threadIdx.x == 0) {
atomicAdd(&gBlockCount, 1); // 统计实际执行过的 block 数
}
}
// ---------------- CPU 参考实现（严格按 GPU 的执行顺序模拟 block 内行为）----------------
static void memSpaceDemoCPU(const float* in, float* out, const float* w4, int n)
{
float sData[BLOCK];
float sDyn[BLOCK];
int nBlocks = (n + BLOCK - 1) / BLOCK;
for (int b = 0; b < nBlocks; ++b) {
for (int t = 0; t < BLOCK; ++t) {
int tid = b * BLOCK + t;
sData[t] = (tid < n) ? in[tid] : 0.0f;
}
for (int t = 0; t < BLOCK; ++t) {
int tid = b * BLOCK + t;
float self = (tid < n) ? in[tid] : 0.0f;
float w = w4[t & 3];
float nb = sData[(t + 1) % BLOCK];
sDyn[t] = w * self + 0.5f * nb;
}
for (int t = 0; t < BLOCK; ++t) {
int tid = b * BLOCK + t;
float sum = 0.0f;
for (int k = 0; k < 4; ++k) {
sum += sDyn[(t + k * 7) % BLOCK] * w4[k];
}
if (tid < n) out[tid] = sum;
}
}
}
int main(void)
{
const float weights[2][4] = {
{0.5f, 1.0f, 1.5f, 2.0f}, // 与 __constant__ 声明处的初值一致
{1.0f, 2.0f, 3.0f, 4.0f} // 用来演示 cudaMemcpyToSymbol 改写常量内存
};
// ---- 1) 打印 kernel 的静态属性：寄存器数与共享内存用量 ----
cudaFuncAttributes attr;
CUDA_CHECK(cudaFuncGetAttributes(&attr, memSpaceDemo));
printf("========== cudaFuncGetAttributes(memSpaceDemo) ==========\n");
printf(" numRegs = %d (每个线程占用的寄存器数，register 空间)\n",
attr.numRegs);
printf(" sharedSizeBytes = %zu B (静态 __shared__ sData[%d] = %zu B)\n",
attr.sharedSizeBytes, BLOCK, BLOCK * sizeof(float));
printf(" localSizeBytes = %zu B (local memory 溢出量，0 表示没有溢出)\n",
attr.localSizeBytes);
printf(" constSizeBytes = %zu B\n", attr.constSizeBytes);
printf(" maxThreadsPerBlock = %d\n", attr.maxThreadsPerBlock);
printf(" numRegs per block = %d * %d = %d\n",
attr.numRegs, BLOCK, attr.numRegs * BLOCK);
int blocksByRegs = 65536 / (attr.numRegs * BLOCK);
int blocksBySmem = (164 * 1024) / (int)(attr.sharedSizeBytes
+ BLOCK * sizeof(float));
int blocksByThreads = 2048 / BLOCK;
int resident = blocksByRegs;
if (blocksBySmem < resident) resident = blocksBySmem;
if (blocksByThreads < resident) resident = blocksByThreads;
printf(" ---- A100 上的占用率估算 ----\n");
printf(" blocks_by_regs = 65536 / (%d*%d) = %d\n",
attr.numRegs, BLOCK, blocksByRegs);
printf(" blocks_by_smem = 164KB / (static %zu B + dynamic %zu B) = %d\n",
attr.sharedSizeBytes, (size_t)(BLOCK * sizeof(float)), blocksBySmem);
printf(" blocks_by_threads = 2048 / %d = %d\n", BLOCK, blocksByThreads);
printf(" resident blocks = %d -> %d threads = %d warps = %.1f%% occupancy\n",
resident, resident * BLOCK, resident * BLOCK / 32,
100.0 * resident * BLOCK / 2048.0);
// ---- 2) 主机端准备数据 ----
float* h_in = (float*)malloc(N * sizeof(float));
float* h_out = (float*)malloc(N * sizeof(float));
float* h_ref = (float*)malloc(N * sizeof(float));
float* h_scr = (float*)malloc(N * sizeof(float));
if (!h_in || !h_out || !h_ref || !h_scr) { fprintf(stderr, "malloc failed\n"); return EXIT_FAILURE; }
for (int i = 0; i < N; ++i) h_in[i] = (float)((i % 17) * 0.5) - 4.0f;
float* d_in = NULL;
float* d_out = NULL;
CUDA_CHECK(cudaMalloc((void**)&d_in, N * sizeof(float)));
CUDA_CHECK(cudaMalloc((void**)&d_out, N * sizeof(float)));
CUDA_CHECK(cudaMemcpy(d_in, h_in, N * sizeof(float), cudaMemcpyHostToDevice));
int nBlocks = (N + BLOCK - 1) / BLOCK;
dim3 grid(nBlocks);
dim3 block(BLOCK);
size_t dynSmem = BLOCK * sizeof(float); // 动态共享内存字节数
cudaEvent_t t0, t1;
CUDA_CHECK(cudaEventCreate(&t0));
CUDA_CHECK(cudaEventCreate(&t1));
// ---- 3) 两轮运行：第一轮用编译期常量，第二轮用 cudaMemcpyToSymbol 改写 ----
for (int round = 0; round < 2; ++round) {
if (round == 1) {
CUDA_CHECK(cudaMemcpyToSymbol(cWeight, weights[1], sizeof(weights[1])));
printf("\n[round 1] cudaMemcpyToSymbol(cWeight, {1,2,3,4}) done.\n");
} else {
printf("\n[round 0] using compile-time __constant__ cWeight = {0.5,1,1.5,2}\n");
}
// 初始化 __device__ 全局变量（host 侧只能通过 symbol 或 kernel 访问）
float zeros[N];
memset(zeros, 0, sizeof(zeros));
CUDA_CHECK(cudaMemcpyToSymbol(gScratch, zeros, nBlocks * sizeof(float)));
int zeroCount = 0;
CUDA_CHECK(cudaMemcpyToSymbol(gBlockCount, &zeroCount, sizeof(int)));
CUDA_CHECK(cudaMemset(d_out, 0, N * sizeof(float)));
CUDA_CHECK(cudaEventRecord(t0));
memSpaceDemo<<<grid, block, dynSmem>>>(d_in, d_out, N);
CUDA_CHECK(cudaEventRecord(t1));
CUDA_CHECK(cudaEventSynchronize(t1)); // 必须同步，否则计时不准、D2H 也可能读到半成品
float ms = 0.0f;
CUDA_CHECK(cudaEventElapsedTime(&ms, t0, t1));
double bytes = 2.0 * N * sizeof(float); // 读 N*4B + 写 N*4B
printf(" kernel time = %.4f ms, effective BW = %.2f GB/s\n",
ms, bytes / (ms * 1.0e-3) / 1.0e9);
CUDA_CHECK(cudaMemcpy(h_out, d_out, N * sizeof(float), cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaMemcpyFromSymbol(h_scr, gScratch, nBlocks * sizeof(float)));
int count = 0;
CUDA_CHECK(cudaMemcpyFromSymbol(&count, gBlockCount, sizeof(int)));
printf(" gBlockCount (__device__ global var) = %d (expected %d)\n",
count, nBlocks);
double gpuSum = 0.0, scrSum = 0.0;
for (int i = 0; i < nBlocks; ++i) scrSum += h_scr[i];
for (int i = 0; i < N; ++i) gpuSum += h_out[i];
printf(" sum(out) = %.4f ; sum(gScratch) = %.4f (两者应相等)\n",
gpuSum, scrSum);
memSpaceDemoCPU(h_in, h_ref, weights[round], N);
int ok = 1;
for (int i = 0; i < N; ++i) {
float diff = h_out[i] - h_ref[i];
if (diff < 0) diff = -diff;
if (diff > 1e-3f) { ok = 0; printf(" MISMATCH at %d: gpu=%f cpu=%f\n",
i, h_out[i], h_ref[i]); break; }
}
printf(" verification vs CPU : %s\n", ok ? "[PASS]" : "[FAIL]");
if (count != nBlocks) printf(" [WARN] gBlockCount != nBlocks\n");
}
// ---- 4) 释放 ----
CUDA_CHECK(cudaEventDestroy(t0));
CUDA_CHECK(cudaEventDestroy(t1));
CUDA_CHECK(cudaFree(d_in));
CUDA_CHECK(cudaFree(d_out));
free(h_in); free(h_out); free(h_ref); free(h_scr);
CUDA_CHECK(cudaDeviceReset());
printf("\nDone.\n");
return EXIT_SUCCESS;
}
- 【代码做什么？】
- 声明四种存储空间：
float acc等自动变量 → register；__shared__ float sData[256]→ 静态共享内存（1 KB/block）；extern __shared__ float sDyn[]→ 动态共享内存（大小在启动配置<<<grid, block, dynSmem>>>的第三个参数给出，本例 1 KB/block）；__constant__ float cWeight[4]→ 常量内存；__device__ float gScratch[N]与__device__ int gBlockCount→ global memory。 cudaFuncGetAttributes(&attr, memSpaceDemo)在主机端读出 kernel 的静态属性：numRegs（每线程寄存器数）、sharedSizeBytes（静态共享内存字节数）、localSizeBytes（local memory 溢出量，0 表示没有寄存器溢出）、maxThreadsPerBlock。这是 Lab 0/Lab 2 里判断”占用率被什么限制”的第一手数据。- 两轮运行演示常量内存的两种初始化方式：第 0 轮直接用声明处的编译期初值
{0.5, 1.0, 1.5, 2.0}；第 1 轮用cudaMemcpyToSymbol(cWeight, weights[1], sizeof(weights[1]))在运行时改写。注意__constant__变量的地址在主机端不是一个普通指针，必须用cudaMemcpyToSymbol/cudaMemcpyFromSymbol（cudaMemcpy传cWeight会尝试解引用一个设备地址而报错）。 __device__全局变量也必须用 symbol API 初始化和读回：cudaMemcpyToSymbol(gScratch, zeros, nBlocks*sizeof(float))清零、cudaMemcpyFromSymbol(h_scr, gScratch, ...)读回，最后用cudaMemcpyFromSymbol(&count, gBlockCount, sizeof(int))验证”确实有 nBlocks = 16 个 block 执行了”。- kernel 内部的 shared memory 生命线：
sData[threadIdx.x] = self;写入 →__syncthreads()→ 读邻居sData[(threadIdx.x+1) % 256]。如果没有这个__syncthreads()，warp 7 可能读到 warp 0 尚未写入的 sData，结果是随机错误（且往往在小数据量下”看起来是对的”，非常难查）。 - 本 kernel 刻意做成”共享内存有收益”的形状：每个 block 的 256 个元素被读进 sData 后被读 2 次（自己一次 + 邻居一次），再经 sDyn 被读 4 次（循环累加）。若不用共享内存，这 6 次访问全部要打到 L1/global，流量放大 6 倍。
- 同步与计时：
cudaEventSynchronize(t1)保证计时覆盖完整执行；cudaDeviceSynchronize在这里没显式出现，但cudaMemcpy(D2H)本身是同步操作，等效地起到了屏障作用——这正是”忘记同步”最隐蔽的形式：有时因为后面的 memcpy 掩盖了错误。 - 主机端逐元素比对：CPU 参考实现严格按”先整块写 sData → 再算 sDyn → 再累加”的顺序模拟，因此可以做到 1e-3 容差内的完全一致。
- 声明四种存储空间：
- 【并行机制与硬件映射解说】
- block/warp 划分：
block = 256线程 = 8 个 warp；grid = 16个 block。A100 有 108 个 SM，16 个 block 只够铺满 16 个 SM（约 15%），这是一次”并行度严重不足”的配置——它的目的是教学演示，不是性能。要跑满 A100（108 SM × 8 block = 864 个驻留槽位，建议 2 倍即 1728 个 block），N 至少要1728 × 256 = 442,368个元素。 - 占用率（以 A100 为例）：假设
numRegs = 32：blocks_by_regs = 65536 / (256 * 32) = 65536 / 8192 = 8 blocks_by_smem = 164 KB / (1 KB static + 1 KB dynamic) = 164/2 = 82 blocks_by_threads = 2048 / 256 = 8 blocks_by_hw = 32 resident = min(8, 82, 8, 32) = 8 → 8*256 = 2048 线程 = 64 warp = 100% occupancy 若编译器把寄存器用到 40 个： blocks_by_regs = 65536 / (256*40) = 65536 / 10240 = 6 (向下取整) resident = min(6, 82, 8, 32) = 6 → 6*256 = 1536 线程 = 48 warp = 75% occupancy 寄存器分配粒度是 256 个寄存器/warp：warp 需要的寄存器数会向上取整到 256 的倍数， 所以 8 warp × 32 线程 × 40 regs = 10240 已经对齐，直接按 10240 计。这就是”寄存器压力 → 占用率掉档 → 延迟隐藏能力下降“的完整因果链。可用
__launch_bounds__(256, 8)告诉编译器”我要 8 个 block/SM”，它会主动把寄存器压到 32 个以内（代价是可能产生 local memory 溢出，此时localSizeBytes > 0就是警号）。 - 共享内存 bank 与 bank conflict（逐条分析 bank 编号）：A100 的共享内存有 32 个 bank，每个 bank 宽 4 字节，地址 a 落在 bank
(a/4) % 32。① sData[threadIdx.x] = self; // 写 warp 内线程 t 访问 float 下标 t → 字节地址 4t → bank = (4t/4)%32 = t%32 → 32 个线程命中 32 个不同 bank → 无冲突 (conflict-free)，1 个周期完成 ② sData[(threadIdx.x+1) % 256] // 读邻居 地址 = 4*(t+1) → bank = (t+1)%32 → 仍然是 32 个不同 bank（只是整体旋转了 1） → 无冲突 ③ sDyn[(threadIdx.x + k*7) % 256] // k 是编译期常数，循环内被展开 对固定的 k，地址 = 4*(t+7k) → bank = (t+7k)%32 → 32 个不同 bank → 无冲突 ④ 反例：如果写成 sData[threadIdx.x * 2] 地址 = 8t → bank = (8t/4)%32 = (2t)%32 t=0..15 → bank 0,2,4,...,30；t=16..31 → bank 0,2,4,...,30（重复） → 每个 bank 被 2 个线程访问 → 2-way bank conflict，耗时从 1 周期变 2 周期 再如 sData[threadIdx.x * 32] → bank = (32t/4)%32 = (8t)%32 → 4 个 bank 被 8 线程 轮流访问 → 8-way conflict，耗时 8 倍 ★ 这类"跨步访问"是共享内存的头号陷阱现代 GPU 对同一 32 位字内的多线程访问会做广播（broadcast），所以”多个线程读同一个地址”不算冲突，只有”落到同一 bank 的不同地址”才算。
- 全局内存合并访问分析：
in[tid]与out[tid]中 tid 连续 → warp 读写 32×4 = 128 字节连续，正好 1 条 cache line（4 个 32 B sector，100% 利用），完美合并。gScratch[blockIdx.x]被同一 block 的 256 个线程原子累加，同一个地址被 32 个 lane 同时原子操作，硬件会串行化这 32 次原子（约 32 个周期/次），本例每 block 256 次 → 这是演示代码里唯一的性能瑕疵，换成”先做 block 内 reduction 再由线程 0 原子加一次”可以降低 32 倍原子压力。 - warp 发散：kernel 中唯一的分支是
if (tid < n)与if (threadIdx.x == 0)。前者在 N=4096 且 BLOCK=256 时恰好整除，全 false，零发散；后者只有一个 lane 为真，但它只包住一条指令，属于可接受的谓词化开销（predicated execution），代价约 1 个发射槽。 __syncthreads()的代价：本例有 2 个屏障，每个屏障让 8 个 warp 相互等待。因为所有 warp 的工作量完全对称（同样的访存模式、同样的指令数），负载均衡良好，屏障开销被压到最小——这也从反面说明：屏障的代价取决于负载均衡，而不是屏障本身。
- block/warp 划分：
- 【性能优化分析】
- 算术强度：每个元素读 4 B（
in[tid]）+ 写 4 B（out[tid]）= 8 B；运算量约 12 FLOP（1 次乘加算 2、共享内存复用的 4 次乘加算 8、邻居加权算 2）：AI = 12 FLOP / 8 B = 1.5 FLOP/Byte注意这里没有把 shared memory 流量算进 DRAM 流量——这正是 shared memory 的意义：把 6 次逻辑访问压缩成 1 次 DRAM 访问 + 5 次片上访问。
- Roofline：
上限 = min(19.5e12 , 1.5 × 1555e9) = min(19500 GFLOP/s , 2332 GFLOP/s) = 2332 GFLOP/s 仍然远低于 19.5 TFLOP/s → 依旧带宽受限，但带宽利用率的要求比程序 2 宽松（AI 高 3 倍） 理想时间 = (4096 × 8 B) / 1555e9 B/s = 32.8 KB / 1555 GB/s = 21 ns 实测时间会被 kernel 启动开销(3-10 µs)完全支配 → 加速比看上去只有 0.001x， 这不是 kernel 写得差，而是"数据量太小，不适合 GPU"。 - 占用率：见上，8 个 block/SM = 100%（32 寄存器）或 6 个 block/SM = 75%（40 寄存器）。要判断真实瓶颈，应在主机端打印
attr.numRegs、attr.localSizeBytes，再用cudaOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks, memSpaceDemo, BLOCK, dynSmem)让 runtime 直接告诉你答案。 - 瓶颈判定与优化方向：延迟受限 + 并行度不足（16 个 block 只填 16/108 个 SM；单次执行 21 ns 而启动开销 5 µs）。三条可执行方向：① 增大 grid——把 N 提到 100 万以上，让 block 数达到 4000+；② 粗粒度化——让每个线程处理 4~8 个元素，减少总 block 数反而能提高每线程 ILP 并减少索引开销（见”线程映射粒度”概念）；③ 批处理——把多次小 kernel 合并成一个 stream 上的连续启动，摊薄启动开销。
- 算术强度：每个元素读 4 B（
性能优化技巧总结
- 网格维度必须向上取整：
grid.x = (N + block.x - 1) / block.x，因为 C 的整数除法是向下取整，N/256会漏掉尾部元素，而漏算的元素不会报错，只会算错——这是最危险的 bug 类型。 - 每个 kernel 都要有边界检查：
if (row < height && col < width)挡住”向上取整带来的多余线程”；没有它，越界线程会写坏相邻缓冲区，且症状随机。 - block 大小选 256（16×16）作为默认起点：256 线程 = 8 个 warp，在 A100/RTX 4090 上都能整除达到 100% 占用率，又不会被单个 block 的 1024 线程上限或”每 SM 最多 8 个 block”这类硬件限制卡住。
- 让 warp 的访存连续且对齐：
threadIdx.x必须映射到”内存中变化最快的维度”（行主序中的列），这样 32 个线程 × 4 字节 = 128 字节 = 1 条 cache line，一次 transaction 拿满。 - 能向量化就向量化：
uchar4/float4/int4一条指令搬 4 个元素，把 warp 的有效载荷从 128 字节里的 32 字节提升到 128 字节，通常带来 1.5×~3× 的带宽提升。 - 分支粒度对齐 warp size：把
if (threadIdx.x > 2)改成if (threadIdx.x / warpSize > 2)，让每个 warp 内部走同一条路径，避免两条路径串行执行。 - 共享内存访问避免跨步：
sData[threadIdx.x]是无冲突的，sData[threadIdx.x * 2]是 2-way 冲突，sData[threadIdx.x * 32]是 8-way 冲突；必要时用 padding（sData[BLOCK + 1]）打破跨步。 - 用
__syncthreads()建立”生产者–消费者”边界，并且只在需要时用：写入共享内存之后、读取之前必须有一个屏障；但屏障也会让快 warp 等慢 warp，所以要放在负载均衡的位置，绝不放进条件分支。 - 先算占用率再谈优化：
resident = min(线程约束, 寄存器约束, smem 约束, block 数上限)，任何一项都可能成为瓶颈；attr.numRegs和cudaOccupancyMaxActiveBlocksPerMultiprocessor是最快的诊断工具。 - 小数据不要上 GPU：kernel 启动固定开销 3–10 µs，PCIe 传输约 25 GB/s（比显存慢 60 倍），当总计算量 < 10 万次操作时，GPU 版本通常比 CPU 还慢。
关键要点
- SPMD 是 CUDA 的全部哲学：一个 kernel 一份代码，靠
blockIdx/threadIdx决定”我是谁、我处理哪份数据”；把 CPU 的for循环变量替换成线程索引，就是数据并行的标准改写手法。 - 线程层次有三层且语义不同：grid 只是”启动配置”，block 是”可同步、可共享内存的合作单元”，warp 是”硬件调度单元”。不同 block 之间不能同步、执行顺序任意，跨 block 合作只能靠 kernel 边界或原子操作。
- warp（32 线程）是性能分析的最小单位：合并访问按 warp 的 128 字节 transaction 判定，分支发散按 warp 内是否走不同路径判定，占用率按 warp 数计算。看代码时必须以 warp 为单位心算，而不是逐线程。
- 内存层次的延迟跨越三个数量级：寄存器 1 周期、共享内存 20–30 周期、L1 约 30 周期、L2 约 200 周期、全局内存 400–800 周期、主机内存 10000+ 周期。优化的本质就是把访问尽量往上层挪。
__syncthreads()是 block 级屏障，cudaDeviceSynchronize()是 host-device 同步；前者保证共享内存的可见性，后者保证主机能安全读取设备结果。kernel 启动是异步的，没有同步的计时和读回都是错的。- 共享内存”有收益”的前提是数据复用：程序 2 的每个像素只读一次，用共享内存只会增加开销；程序 3 的每个元素被读 6 次，共享内存才有价值。盲目使用共享内存是新手最常见的”优化反而变慢”。
常见陷阱与注意事项
- 整数除法取整导致网格覆盖不足：写
dim3 grid(N/256)时N=1000得到 3 个 block，元素 768–999 无人处理，结果数组尾部全是未初始化的垃圾 → 正确做法是(N + 255)/256或ceil(N/256.0)，并且永远保留边界检查。 - 越界访问（out-of-bounds）：2D kernel 里忘记
if (row < height && col < width)，最右一列/最下一行 block 的多余线程会写到别的缓冲区，症状是”结果大部分正确、少数像素花屏”，且在compute-sanitizer关闭时毫无提示 → 正确做法是所有访问全局内存的语句都包在边界检查里。 - 忘记
cudaDeviceSynchronize()/ 同步点放错位置：kernel 是异步启动的，紧跟其后的cudaMemcpy(D2H)虽然本身同步，但如果用cudaMemcpyAsync或直接读 pinned 内存，就会读到半成品；计时时若不在cudaEventRecord(t1)后加cudaEventSynchronize(t1)，测到的是启动开销 → 正确做法是计时必同步、读回前必同步。 - 忘记
__syncthreads()：多个 warp 通过共享内存交换数据时，写完之后必须有一个屏障，否则快 warp 会读到慢 warp 还没写的旧值；典型表现是”数据量小时结果正确、数据量大时偶发错误” → 正确做法是在”写共享内存”与”读共享内存”之间放一个__syncthreads()。 - 把
__syncthreads()放进分支里：if (threadIdx.x < 16) { ...; __syncthreads(); }会让部分 warp 永远等不到同伴，行为未定义（可能挂死） → 正确做法是让屏障处于所有线程都必然执行的代码路径上。 - 共享内存 bank conflict：让线程按
sData[threadIdx.x * 32]这样的跨步访问，32 个线程会撞到 4 个 bank，形成 8-way conflict，shared 访问耗时直接乘以 8 → 正确做法是保证 warp 内连续线程访问连续的字（sData[threadIdx.x]），或加 padding（__shared__ float s[BLOCK + 1]）。 - warp 发散（branch divergence）：把
threadIdx.x当分支条件（如if (threadIdx.x > 2)）会让同一 warp 的两条路径串行执行，若两条路径各 100 条指令，warp 就要执行 200 条 → 正确做法是让分支粒度对齐warpSize（if (threadIdx.x / 32 > 2)），或者用查表/数学变换消除分支。注意：若分支条件依赖的是所有线程都相同的量（如 kernel 参数q），则不会发散。 - 不检查 CUDA API 返回值：
cudaMalloc失败返回cudaErrorMemoryAllocation、kernel 越界启动返回cudaErrorInvalidConfiguration，如果全部忽略，错误会在几百行之后以”结果不对”的形式出现，极难定位 → 正确做法是全程使用CUDA_CHECK宏，并且在 kernel 启动后补一次CUDA_CHECK(cudaGetLastError())（启动错误是异步上报的）。 cudaMemcpy方向写错：把cudaMemcpyHostToDevice和cudaMemcpyDeviceToHost弄反，行为是”立即报 invalid argument”或”拷回一堆垃圾”；从设备指针拷到设备指针必须用cudaMemcpyDeviceToDevice→ 正确做法是记住参数顺序(dst, src, bytes, kind)，并且从谁的角度看 dst/src：主机变量在前、设备指针在后是 H2D。- host/device 指针混用：把
cudaMalloc得到的设备指针直接传给 CPU 代码解引用，会触发段错误或读到无关数据；反之把主机指针传给 kernel 会得到非法地址访问 → 正确做法是用命名规范（h_前缀 /d_前缀）区分，并在cudaMalloc之后立刻cudaMemset清零以便早期发现。 - 共享内存容量超限：A100 每 block 静态共享内存上限默认 48 KB，超过会启动失败（
cudaErrorInvalidValue）；要用到 48 KB 以上必须调用cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, bytes)显式申请（A100 上限 164 KB/SM） → 正确做法是先查sharedMemPerBlockOptin，再 opt-in。
思考题（带答案）
Q1. 用 dim3 block(16,16) 覆盖一张 76×62 的灰度图，需要多少个 block？总共启动多少线程？其中有多少线程什么也不做？如果把 block 换成 8×8 呢？ 需要 ceil(76/16) = 5 列 × ceil(62/16) = 4 行 = 20 个 block，共启动 20 × 256 = 5120 个线程。图像实际只有 76 × 62 = 4712 个像素，因此 5120 − 4712 = 408 个线程被 if (row < height && col < width) 挡住（浪费 7.97%）。若换成 8×8：ceil(76/8) = 10、ceil(62/8) = 8，共 80 个 block、80 × 64 = 5120 个线程，仍然是 408 个空转线程，浪费比例完全相同（因为两种 block 的整除余数一样）；但 8×8 的 block 只有 2 个 warp，受”每 SM 最多 block 数”限制时更难达到高占用率，所以 16×16 是更好的选择。
Q2. 一个 kernel 每个线程用 32 个寄存器，block 为 256 线程，不使用共享内存。在 A100（每 SM 65536 寄存器、2048 线程、64 warp、最多 32 个 block）上，占用率是多少？如果把寄存器用到 40 个呢？ 32 寄存器时：每 block 寄存器 256 × 32 = 8192，65536 / 8192 = 8 个 block；线程上限给 2048/256 = 8 个 block；取 min 得 8 个 block = 8 × 256 = 2048 线程 = 64 个 warp = 100% 占用率。40 寄存器时：256 × 40 = 10240，65536 / 10240 = 6.4 → 6 个 block = 1536 线程 = 48 个 warp = 75% 占用率。可见寄存器每多 8 个，占用率就掉 25%；若占用率下降导致延迟隐藏不足（例如这个 kernel 频繁访问全局内存），可以用 __launch_bounds__(256, 8) 强制编译器把寄存器压回 32 个以内。
Q3. 讲义中的向量加法用”一个线程一个元素”，现在改成”一个线程两个元素”：vecAdd<<<ceil(N/(2*256.0)), 256>>>，kernel 里写 i = blockIdx.x*(2*blockDim.x) + threadIdx.x; C[i]=A[i]+B[i]; i = i + blockDim.x; C[i]=A[i]+B[i];。当 N = 1000 时，网格是多少？这样做的好处和风险是什么？为什么第二次访问仍然是合并的？ ceil(1000/512) = 2 个 block，共 2 × 256 = 512 个线程（细粒度版需要 4 个 block、1024 个线程）。好处：总线程数减半，每个线程承担 2 次乘加，索引计算与边界判断的固定开销被摊薄，同时每个线程有两条独立的 load 可以重叠，提高了内存级并行（MLP）。风险：线程总数太少会降低并行度，无法填满 SM 的 warp 槽位（A100 需要约 44 万线程才能跑满）；并且必须两次都做边界检查，否则 N 为奇数时第二次访问会越界。第二次访问仍然合并：warp 内 32 个线程的 i 是连续递增的（因为是 i + blockDim.x，对同一 warp 内所有 lane 加的常数相同），地址 32×4 = 128 字节连续，正好 1 条 cache line，只是整体偏移了 blockDim.x × 4 = 1024 字节。