关卡阅读材料 // 00
可编程并行
把一个串行循环改写成许多线程，并理解 Grid、Block 与 Thread 不是硬件参数表，而是程序描述并行工作的方式。
建议边读边写下答案；时长包含代码推演与练习。
进入互动关卡→- 01导读与目标05 MIN
- 02心智模型08 MIN
- 03概念深挖15 MIN
- 04代码推演12 MIN
- 05历史生态08 MIN
- 06练习复盘12 MIN
通过动手实践掌握 GPU 加速基础：识别可并行工作、编写 Kernel、配置执行层次，并验证结果与性能。
互动关卡先用视觉模型建立 Grid / Block / Thread 心智模型；可选实验再把同一映射迁移到真实 Vector Add，形成从概念到可运行程序的第一座桥。
- 基础 C/C++ 或 Python
- 数组、循环和函数
- 能使用命令行运行示例
- CUDA Toolkit / nvcc
- CUDA Samples
- Nsight Systems 或 CUDA Events
导读与目标
先记住这一句
先找出彼此独立的工作，再决定用多少线程表达它们。
完成后你应该能够
- 01
判断一段循环中的迭代是否彼此独立，是否适合映射到 GPU 线程。
- 02
根据数据规模与 Threads / Block 计算 Grid，并解释为什么需要边界检查。
- 03
区分 Grid、Block、Warp、Thread 的程序语义与它们在硬件上的执行关系。
- 04
说出至少三个让“线程很多”却没有获得等比例加速的原因。
建立心智模型
- 01工作
数组中的每个元素都要执行同一个加法，元素之间没有依赖。
- 02映射
一个线程负责一个元素；Block 把一组线程组织起来，Grid 再组织所有 Block。
- 03边界
启动线程数通常会向上取整，因此 kernel 必须检查全局索引是否仍在数据范围内。
概念深挖
全局线程索引
blockIdx.x × blockDim.x + threadIdx.x 把 Block 内局部编号变成整个 Grid 中的唯一编号。
层次不是装饰
同一 Block 的线程可以同步并共享片上 Shared Memory；不同 Block 应被设计为可独立调度。
并行度不等于加速比
线程数只描述可并行的工作。实际速度仍受访存、分支、资源占用和任务规模限制。
逻辑层次不是 SM 排布图
Grid 与 Block 是 kernel 的逻辑工作分解。一个 Grid 可以包含远多于物理 SM 数量的 Block；运行时会在资源允许时把 Block 分批调度到 SM。程序不应假设 Block 0 一定先于 Block 1 完成，也不应依赖不同 Block 之间的隐含同步。
Block 内线程会被组织成 Warp 执行。当前 CUDA 编程模型中 Warp 通常包含 32 个线程，但写正确程序时仍应从 Thread 和 Block 的语义出发；Warp 主要影响分支、访存与性能推理。Block 内可以使用 Shared Memory 和同步原语，这也是为什么 Block 不是任意切开的编号集合。
停下来想一想如果 Grid 有 10,000 个 Block，而 GPU 只有几十个 SM，程序还能正确运行吗？为什么？＋
参考答案可以。Block 会分批调度；只要 Block 之间不依赖未定义的执行顺序，物理并发数量只影响执行时间，不影响结果。
索引就是数据所有权
全局索引不仅是一个公式，它定义了哪一个线程拥有哪一份数据。一维数组常用 blockIdx.x × blockDim.x + threadIdx.x；二维图像通常分别计算 x 与 y，再用 y × width + x 得到线性地址。索引设计错误会带来重复写、遗漏或越界。
启动配置向上取整是常态，因为数据长度很少恰好是 Block 大小的整数倍。边界分支通常只影响最后一个 Block 的少量线程，却能让同一 kernel 安全处理任意 n。先保证所有权唯一、覆盖完整和边界安全，再谈性能。
停下来想一想二维图像 kernel 只检查 x < width，却不检查 y < height，会留下什么问题？＋
参考答案最后一行之后的线程仍可能计算出线性地址并越界访问。二维映射必须同时保护两个维度。
Block 大小是一组资源交换
Threads / Block 会决定每个 Block 包含多少 Warp，也会和每线程寄存器、每 Block Shared Memory 一起决定一个 SM 能同时驻留多少 Block。更大的 Block 可能减少可同时驻留的 Block 数；更小的 Block 又可能无法提供足够的 Warp 隐藏延迟。
Occupancy 是分析工具，不是唯一目标。一个内存带宽受限的 kernel、一个寄存器密集型 kernel 与一个计算密集型 kernel 可能偏好不同配置。课程中的 64 线程只是满足特定映射题目，不是所有 GPU 程序的最佳答案。
停下来想一想为什么把 Block 从 128 线程改成 1024 线程可能反而变慢？＋
参考答案单个 Block 可能消耗更多寄存器或 Shared Memory，减少同一 SM 上可驻留的 Block 与 Warp，并放大尾部或分支浪费。
把概念放回代码
01__global__ void add(float* a, float* b, float* c, int n) {02 int i = blockIdx.x * blockDim.x + threadIdx.x;03 if (i < n) c[i] = a[i] + b[i];04}05int blocks = (n + threads - 1) / threads;06add<<<blocks, threads>>>(a, b, c, n);代码逐段推演
- 01KERNEL__global__ 定义执行位置
函数由 Host 发起，在 Device 上执行。它描述每个线程都会运行的同一份程序。
- 02INDEX计算唯一的 i
blockIdx.x 选择 Block，blockDim.x 给出 Block 宽度，threadIdx.x 选择 Block 内线程；三者组成全局所有权。
- 03GUARD只让有效线程写回
向上取整会多启动线程。if (i < n) 保护数组，不能依赖数据长度刚好整除 Block 大小。
- 04LAUNCH配置是问题规模的函数
blocks = ceil(n / threads) 覆盖全部元素。Threads / Block 应在正确性成立后，通过资源分析与测量选择。
知识图谱
补充参考 · 不要求一次读完，也不计入核心 60 分钟
核心术语
- Kernel
- 由 Host 发起、在 Device 上由许多线程执行的函数；每个线程运行同一程序，但拥有不同索引。
- Grid
- 一次 Kernel 启动产生的全部 Block，是问题的逻辑工作空间，不是 GPU 物理单元图。
- Block
- 可共享 Shared Memory、执行 Block 级同步的一组线程，也是运行时调度到 SM 的基本工作单元。
- Warp
- SM 执行线程的基本分组，通常为 32 个线程；分支和地址模式会影响它的执行效率。
- Occupancy
- 一个 SM 上活跃 Warp 数与其最大可能活跃 Warp 数的比率；它帮助解释延迟隐藏，但不是性能分数。
带数字推演
- 01算 Block
ceil(1000 ÷ 256) = 4，因此 Grid 需要 4 个 Block。
- 02算线程
4 × 256 = 1024 个线程被启动，其中前 1000 个拥有真实数据。
- 03处理尾部
最后 24 个线程没有对应元素，必须由 if (i < n) 拦住。
诊断手册
结果偶发错误或报非法地址
- 先检查
- 索引唯一性、向上取整后的尾部线程、数组长度与错误检查
- 需要的证据
- Compute Sanitizer 报告、最小失败规模、CPU 参考输出差异
GPU 比 CPU 还慢
- 先检查
- 问题规模、冷启动、Host↔Device 搬运与 Kernel 启动占比
- 需要的证据
- 分段计时和多次预热后的中位数，而不是一次运行
增大 Block 后性能下降
- 先检查
- 每线程寄存器、每 Block Shared Memory、活跃 Warp 与尾部 Block
- 需要的证据
- 编译资源报告、Occupancy 数据和多个 Block 配置的可比测量
从硬件到生态
- 01固定功能图形流水线
程序围绕图形阶段组织计算。
通用问题需要伪装成图形工作，表达成本很高。 - 02G80 + CUDA
统一着色器硬件与 Kernel / Thread 编程模型出现。
开发者可以直接描述数据并行问题。 - 03现代 CUDA 生态
库、编译器、Profiler 与框架围绕同一执行模型协作。
学习线程映射成为进入科学计算和 AI 工作流的共同基础。
硬件与生态坐标
统一着色器架构让大量处理单元能够执行更通用的程序；CUDA 随后提供 kernel、线程层次、内存与同步等编程抽象。这里的转折不是“并行计算诞生”，而是 GPU 并行变得更直接、更可编程。
练习与复盘
处理 1000 个元素，每个 Block 256 个线程时，为什么 kernel 仍需要 if (i < n)？
- A. 让线程执行得更快
- B. 因为会启动 1024 个线程，最后 24 个索引越界
- C. 因为一个 Warp 只有 16 个线程
查看答案＋
Block 数向上取整为 4，因此一共启动 1024 个线程。边界检查让多出的线程安全退出。
n = 10,000，Threads / Block = 256。需要多少个 Block？会多启动多少线程？
提示＋
对 10,000 / 256 向上取整。
参考答案＋
需要 40 个 Block，共启动 10,240 个线程，多出的 240 个线程由边界检查退出。
kernel 直接执行 c[i] = a[i] + b[i]，但没有 if (i < n)。什么时候最容易暴露问题？
提示＋
考虑 n 不是 Block 大小整数倍。
参考答案＋
最后一个 Block 会包含无效索引，可能越界读写。加入边界检查，并确保所有输入输出数组都按相同范围保护。
处理 1920×1080 图像，Block 为 16×16。Grid 的 x、y 各是多少？
提示＋
两个维度分别向上取整。
参考答案＋
Grid 为 120×68。y 方向会覆盖 1088 行，因此 kernel 需要同时检查 x < 1920 与 y < 1080。
为什么不能让 Block 1 等待 Block 0 写入一个普通全局变量后再继续？
提示＋
考虑 Block 的调度顺序与驻留资源。
参考答案＋
Block 没有保证的先后顺序，等待可能造成死锁或读取竞态。需要拆成多个 kernel、使用受支持的协作启动，或重新设计为 Block 独立。
可选实践实验
不计入核心 60 分钟 · 需要相应 CUDA / GPU 环境
实验目标
把第一个 Vector Add 从 CPU 迁移到 GPU
得到与 CPU 参考结果一致的 CUDA Kernel，并能解释启动配置、边界检查和 Host↔Device 搬运分别承担什么职责。建议步骤
- 01
准备 CPU 参考实现和 10,000 个元素的确定性输入，保存输出作为 correctness oracle。
- 02
实现一个线程处理一个元素的 Kernel，先固定 256 Threads / Block，再用向上取整计算 Grid。
- 03
分别测试 n = 10,000、10,001 和小于一个 Block 的输入，确认边界保护有效。
- 04
把分配、H2D、Kernel、D2H 分段计时，并与只测 Kernel 的结果区分开。
提交 Kernel、三个规模的正确性比较，以及一张标明端到端时间与 Kernel 时间的表。不要把单次冷启动作为最终结论。
改写成二维图像映射，使用 16×16 Block，并解释为何 x、y 都需要边界检查。