第 1 课：线程映射与 Grid-Stride Loop

本课目标与完成标准

学完后你应能：从 thread/block/grid 映射一维数据，并写出处理任意 N 的合并访存 kernel。

通过必须同时满足：

• 独立补齐本课唯一的代码填空题，并通过给定检查；
• 三个问答题均说明因果链，而不是只报术语；
• 能指出至少一个正确性边界和一个性能取舍；
• 总分不低于 8/10，且没有一票否决级概念错误。

前置关系

• 课程前置：C/C++、线性代数、基本并行编程
• 本课在路线中的作用：Vector Add 是最小 CUDA 工作流：host launch 决定 grid，device thread 用全局索引定位元素。

核心心智模型

1. 向量加法在 GPU 体系中的定位与第一性原理

向量加法（Vector Add: $c[i] = a[i] + b[i]$）是所有高性能 GPU 编程学习者的“Hello World”，更是深入理解 CUDA 线程层次结构与硬件访存机制的最小工作原型：

• 数据并行性（Embarrassingly Parallel）：每个元素 $i$ 的加法运算完全独立，元素之间不存在任何数据依赖或跨线程通信需求。
• 访存受限特征（Memory-Bound Kernel）： 向量加法的算术强度极低：每个线程从显存中加载两个 float（$4 + 4 = 8$ 字节），执行 1 次加法（$1\text{ FLOP}$），然后写回一个 float（$4$ 字节）。 其算术强度仅为： $$I = \frac{1\text{ FLOP}}{12\text{ Bytes}} \approx 0.083\text{ FLOPs/Byte}$$ 这决定了其性能完全受限于 GPU 全局显存（HBM / GDDR）的物理带宽。编写高吞吐向量加法内核的核心，在于如何完美对齐硬件的内存访问合并（Memory Coalescing）。

2. 线程层次映射与 Grid-Stride Loop 执行拓扑（如图 1-1 与图 1-2 所示）

A. 层次化坐标映射体系（图 1-1）

CUDA 网格层次与一维映射

Grid (网格)
dim3: blocksPerGrid

Block (线程块)
dim3: 256 / 512

Warp (线程束)
每个 Warp 32 线程

Thread (单线程)
全局绝对索引 idx:
blockIdx.x * blockDim.x + threadIdx.x

图 1-1：GPU 线程层次（Grid/Block/Warp/Thread）与全局内存一维映射拓扑

CUDA 运行时通过网格（Grid）和线程块（Block）管理数十万级并发线程：

• threadIdx.x：当前线程在所属 Thread Block 内的一维局部编号（$0 \le \text{threadIdx.x} < \text{blockDim.x}$）。
• blockIdx.x：当前 Thread Block 在全局 Grid 内的一维块编号（$0 \le \text{blockIdx.x} < \text{gridDim.x}$）。
• 全局一维绝对线性索引（Global Linear Index）： $$\text{idx} = \text{blockIdx.x} \times \text{blockDim.x} + \text{threadIdx.x}$$

B. 网格跨步循环（Grid-Stride Loop，图 1-2）

线程启动
idx = blockIdx.x * blockDim.x + threadIdx.x

idx < n ?

计算 c[idx] = a[idx] + b[idx]

线程退出

网格步进:
idx += blockDim.x * gridDim.x

是否

图 1-2：Grid-Stride Loop 网格跨步循环执行控制流

在朴素的初学者代码中，通常假设线程总数严格大于等于向量长度 $N$，直接使用 if (idx < N) c[idx] = a[idx] + b[idx];。这种写法在工业级生产环境中存在严重缺陷（无法应对超大规模 $N$，且启动过多 Block 会浪费驱动调度开销）。

工业级标准范式是采用 Grid-Stride Loop（网格跨步循环）：

int idx = blockIdx.x * blockDim.x + threadIdx.x;
int stride = blockDim.x * gridDim.x;
for (int i = idx; i < n; i += stride) {
    c[i] = a[i] + b[i];
}

其执行机制与工程红利包含三重保证：

1. 解耦硬件规模与问题规模：
   • 步长 $\text{stride} = \text{blockDim.x} \times \text{gridDim.x}$ 严格等于整个 Grid 在单次迭代中能够同时处理的线程总数。
   • 无论向量长度 $N$ 是一千还是十亿，Grid 尺寸可以任意配置并加上限保护（例如限制 gridDim.x <= 4096），线程在循环中自动以固定步长前移，天然处理剩余尾块。
2. 保证全局内存合并访存（Coalesced Access）：
   • 同一个 Warp 内的 32 个连续线程（$\text{threadIdx.x} = 0 \dots 31$）在第 1 轮循环中访问的地址是连续的：$a[\text{idx}], a[\text{idx}+1], \dots, a[\text{idx}+31]$。
   • 在进入下一次迭代时，所有线程同步跨进 $\text{stride}$ 步，访问的地址依然是 $a[\text{idx}+\text{stride}], \dots, a[\text{idx}+\text{stride}+31]$，依然完全连续。
   • 硬件内存控制器可以将这 32 个连续的 4 字节读取合并为单个或两个 128 字节的 Cache Line 突发事务，榨干显存总线带宽。
3. 极佳的可移植性与调试性：同一个 Kernel 在只有 1 个 SM 的嵌入式 GPU 或拥有 132 个 SM 的 H100 上均可零修改无缝编译运行。

3. 正确性边界与常见陷阱

1. 边界防护 i < n 严禁遗漏：
   • 无论 $N$ 是否能被 Block 尺寸整除，尾块未对齐部分必须依靠 i < n 进行严格的越界拦截。
   • 遗漏边界检查会导致非对齐的最后几个线程访问未分配内存，触发非法内存访问（cudaErrorIllegalAddress）导致 GPU 驱动崩溃。
2. 指针别名与 __restrict__ 关键字：
   • 在函数参数中使用 const float* __restrict__ a 告知编译器指针 a, b, c 之间不存在内存重叠（Aliasing）。
   • 这允许 NVCC 编译器放心将全局内存加载优化为使用只读数据缓存（Read-Only Cache / LDG 算子），规避重复加载。

4. 性能与工程取舍（Block Size 与 Occupancy）

• Block 大小的选择：
  • Block Size 必须是 32（Warp 大小）的整数倍，常见取值为 128、256 或 512。
  • 太小（如 32）：无法充分掩盖指令流水线停顿，SM 内部调度器无法保持高饱和度；
  • 太大（如 1024）：单 Block 消耗过多的寄存器堆和共享内存配额，反倒导致每个 SM 上能同时驻留的 Active Blocks 数量锐减，降低硬件占用率（Occupancy）。
• 工业经验法则：对于纯访存密集型向量算子，通常推荐 BLOCK_SIZE = 256 或 512，在寄存器开销极小的情况下轻松达成 100% 理论 Occupancy。

具体演示与数值推演

设向量长度 $N = 1000$，配置 BLOCK_SIZE = 256，启动网格限制为 grid = 2：

• 单网格总并发线程数： $$\text{stride} = \text{gridDim.x} \times \text{blockDim.x} = 2 \times 256 = 512 \text{ 线程}$$

• 跨步循环执行演化：

  • 第 1 轮迭代（$k = 0$）： 全局索引 $\text{idx} \in [0, 511]$ 的 512 个线程同时启动。所有线程均满足 $i < 1000$，并行计算 $c[0 \dots 511] = a[0 \dots 511] + b[0 \dots 511]$。
  • 第 2 轮迭代（$k = 1$）： 所有线程跨步递增：$i = \text{idx} + 512 \in [512, 1023]$。
  • 线程索引 $0 \dots 487$ 对应的计算位置为 $i \in [512, 999]$，依然满足 $i < 1000$，顺利计算剩余的 488 个元素；
  • 线程索引 $488 \dots 511$ 对应的计算位置为 $i \in [1000, 1023]$，此时条件 $i < 1000$ 判定为假，这 24 个边界线程自然安全退出。

• 全量覆盖验证：共执行两轮迭代，精确完成 $512 + 488 = 1000$ 个元素的计算，无任何越界写入，且全流程维持 128 字节合并访存。

实践任务：唯一代码填空题

补齐全局下标与 stride，不改变边界循环。

规则：只能修改 TODO/______ 所在位置；不要删除断言或放宽误差。代码注释说明了每个边界条件。

%%writefile /tmp/01_vector_add.cu
#include <cuda_runtime.h>

// Vector Add: c[i] = a[i] + b[i]
//
// 这是 CUDA 面试里最基础的算子，重点不是算法复杂，
// 而是确认你理解 thread/block/grid 映射和 coalesced memory access。
__global__ void vector_add_kernel(
    const float* __restrict__ a,
    const float* __restrict__ b,
    float* __restrict__ c,
    int n
) {
    int idx = ______;  // TODO 1: 当前线程全局下标
    int stride = ______;  // TODO 2: 整个 grid 的线程跨度

    // grid-stride loop 的好处：
    // 1. kernel 可以处理任意 n，不要求 grid 覆盖全部元素后就停止。
    // 2. 如果 n 很大，每个线程可以处理多个元素，代码仍然简单。
    // 3. 相邻线程访问相邻 idx，读写都是合并访存。
    for (int i = idx; i < n; i += stride) {
        c[i] = a[i] + b[i];
    }
}

void launch_vector_add(
    const float* a,
    const float* b,
    float* c,
    int n,
    cudaStream_t stream
) {
    constexpr int BLOCK_SIZE = 256;
    int grid = (n + BLOCK_SIZE - 1) / BLOCK_SIZE;

    // 限制 grid 上限是常见写法，避免 n 极大时创建过多 block。
    // grid-stride loop 会保证剩余元素继续被处理。
    grid = grid > 4096 ? 4096 : grid;

    vector_add_kernel<<<grid, BLOCK_SIZE, 0, stream>>>(a, b, c, n);
}

检查方法

有 CUDA 环境时执行 nvcc -std=c++17 -c /tmp/01_vector_add.cu -o /tmp/01_vector_add.cu.o；无 CUDA 环境时只做静态审查并登记待验证。

提交时请给出：补齐后的代码、实际运行输出（环境不可用时注明“仅静态审查”）以及对失败用例的解释。

Q1

不要背定义：请从输入、状态变化和输出三个阶段解释“线程映射与 grid-stride loop”的工作机制。

你的答案：

Q2

删掉 i < n 后为什么小尺寸测试也可能偶尔看似正常？

你的答案：

Q3

若 N 远大于最大 grid，为什么 grid-stride loop 仍能覆盖全部元素？

你的答案：

评分与通过规则

• 代码 4 分：正常输入 2 分，边界输入 1 分，解释实现 1 分；
• Q1～Q3 各 2 分；
• 一票否决：结果碰巧正确但核心因果链错误、删除边界检查、把未运行结果说成实测。

需要提示时按四级机制请求：概念区域 → 具体方向 → 关键局部 → 完整答案。

Tip💡 点击展开参考答案与详细代码实现

参考答案（仅 answer 分支）

先完成题目再核对。即使代码一致，也要能解释关键步骤，并尝试更换一个输入规模。

%%writefile /tmp/01_vector_add.cu
#include <cuda_runtime.h>

// Vector Add: c[i] = a[i] + b[i]
//
// 这是 CUDA 面试里最基础的算子，重点不是算法复杂，
// 而是确认你理解 thread/block/grid 映射和 coalesced memory access。
__global__ void vector_add_kernel(
    const float* __restrict__ a,
    const float* __restrict__ b,
    float* __restrict__ c,
    int n
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;

    // grid-stride loop 的好处：
    // 1. kernel 可以处理任意 n，不要求 grid 覆盖全部元素后就停止。
    // 2. 如果 n 很大，每个线程可以处理多个元素，代码仍然简单。
    // 3. 相邻线程访问相邻 idx，读写都是合并访存。
    for (int i = idx; i < n; i += stride) {
        c[i] = a[i] + b[i];
    }
}

void launch_vector_add(
    const float* a,
    const float* b,
    float* c,
    int n,
    cudaStream_t stream
) {
    constexpr int BLOCK_SIZE = 256;
    int grid = (n + BLOCK_SIZE - 1) / BLOCK_SIZE;

    // 限制 grid 上限是常见写法，避免 n 极大时创建过多 block。
    // grid-stride loop 会保证剩余元素继续被处理。
    grid = grid > 4096 ? 4096 : grid;

    vector_add_kernel<<<grid, BLOCK_SIZE, 0, stream>>>(a, b, c, n);
}

Q1 参考答案

• 输入状态：GPU 接收到输入指针 a, b, c、数据规模 $n$ 以及发射配置 <<<grid, block>>>。每个线程硬件内建寄存器被赋予其唯一的坐标 (blockIdx.x, threadIdx.x)。
• 状态变化（执行演化）：
  1. 线程通过基准公式计算首轮全局线性索引：$\text{idx} = \text{blockIdx.x} \times \text{blockDim.x} + \text{threadIdx.x}$；
  2. 若 $\text{idx} < n$，线程执行合并读取 $a[\text{idx}], b[\text{idx}]$，计算加法并写回 $c[\text{idx}]$；
  3. 步长递增：$\text{idx} \leftarrow \text{idx} + \text{blockDim.x} \times \text{gridDim.x}$，进入下一轮跨步循环，直到所有超出 $n$ 的线程自然退出。
• 输出状态：所有 $n$ 个元素被精确且仅被计算一次，结果完整写回全局内存，无越界污染，且连续线程在全生命周期内维持连续访存事务。

Q2 参考答案

• 偶然看似正常的根本原因（内存填充与页对齐）：
  1. CUDA 驱动内存分配对齐：通过 cudaMalloc 申请显存时，底层驱动通常按照 256 字节或 4 KB（甚至更大页面）进行物理内存对齐分配。若实际数组长度 $n$ 较小（如 1000 个 float 占 4000 字节），操作系统分配的真实物理页可能容纳 4096 字节。
  2. 脏读脏写发生在有效页内：尾部超出 $n$ 的 24 个越界线程写入了这部分未被应用使用的对齐内衬区域（Padding Memory），未触发 MMU 缺页异常或段错误（cudaErrorIllegalAddress）。
  3. 输出校验盲区：宿主机在比对时通常只核验 c[0 ... n-1] 的数值，完全无法觉察越界非法写入。
• 潜在危害：一旦数据规模 $n$ 恰好触及物理页边界，或者在连续申请的多张量布局中越界踩踏了相邻张量的数据，将引发极其难以复现的数据损坏与静默崩溃。

Q3 参考答案

• 数学与硬件机制： 在网格跨步循环中，整个 Grid 的瞬时并发能力上限为 $W_{\text{grid}} = \text{gridDim.x} \times \text{blockDim.x}$。 无论 $N$ 达到数千万还是数亿（远大于硬件最大活跃线程数），循环将原本的“一维空间并行展平”转化为了“空间向量化 + 时间序列化”的双重正交执行： $$\text{Iteration } k \implies \text{处理元素区间 } [k \cdot W_{\text{grid}} + \text{idx}]$$ 每个物理线程类似于 CPU 上的一个工作进程（Worker），每轮循环吞吐一个步长的元素，不断轮询直至遍历整个线性空间。

• 性能优势：避免了向 GPU 驱动发射数十万个微型 Block 的发射排队开销，使硬件调度开销降至最低。

参考资料

• CUDA Programming Guide
• CUDA Best Practices Guide

资料用于建立事实基线；面试回答仍需用自己的语言组织。

← 上一课
第 16 课：非均质流水线调度 下一课 →
第 2 课：Warp/Block 归约与两阶…