第 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）

flowchart LR
    subgraph Grid["CUDA 网格层次与一维映射"]
        direction TB
        G["Grid (dim3: blocksPerGrid)"] --> B["Block (dim3: threadsPerBlock = 256/512)"]
        B --> W["Warp 0..W-1 (每个 Warp 32 线程)"]
        W --> T["Thread (全局索引: idx = blockIdx.x * blockDim.x + threadIdx.x)"]
    end

图 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）

flowchart TD
    Start["线程启动: idx = blockIdx.x * blockDim.x + threadIdx.x"] --> Cond{"idx < n ?"}
    Cond -->|是| Compute["计算 c[idx] = a[idx] + b[idx]"]
    Compute --> Stride["idx += blockDim.x * gridDim.x (步进一个网格)"]
    Stride --> Cond
    Cond -->|否| End["线程退出"]

图 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];
}

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

解耦硬件规模与问题规模：
- 步长 $\text{stride} = \text{blockDim.x} \times \text{gridDim.x}$ 严格等于整个 Grid 在单次迭代中能够同时处理的线程总数。
- 无论向量长度 $N$ 是一千还是十亿，Grid 尺寸可以任意配置并加上限保护（例如限制 gridDim.x <= 4096），线程在循环中自动以固定步长前移，天然处理剩余尾块。
保证全局内存合并访存（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 突发事务，榨干显存总线带宽。
极佳的可移植性与调试性：同一个 Kernel 在只有 1 个 SM 的嵌入式 GPU 或拥有 132 个 SM 的 H100 上均可零修改无缝编译运行。

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

边界防护 i < n 严禁遗漏：
- 无论 $N$ 是否能被 Block 尺寸整除，尾块未对齐部分必须依靠 i < n 进行严格的越界拦截。
- 遗漏边界检查会导致非对齐的最后几个线程访问未分配内存，触发非法内存访问（cudaErrorIllegalAddress）导致 GPU 驱动崩溃。
指针别名与 __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/______ 所在位置；不要删除断言或放宽误差。代码注释说明了每个边界条件。

Code

%%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 分；
一票否决：结果碰巧正确但核心因果链错误、删除边界检查、把未运行结果说成实测。

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

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

参考答案（仅 answer 分支）

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

Code

%%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 的发射排队开销，使硬件调度开销降至最低。

参考资料

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

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

--- title: "第 1 课：线程映射与 Grid-Stride Loop" jupyter: python3 --- ::: {.callout-note appearance="simple"} **参考答案版**：包含完整代码实现与问答题深度参考解析。建议先独立完成练习题后再展开核对。 ::: ## 本课目标与完成标准学完后你应能：从 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） ```{mermaid} flowchart LR subgraph Grid["CUDA 网格层次与一维映射"] direction TB G["Grid (dim3: blocksPerGrid)"] --> B["Block (dim3: threadsPerBlock = 256/512)"] B --> W["Warp 0..W-1 (每个 Warp 32 线程)"] W --> T["Thread (全局索引: idx = blockIdx.x * blockDim.x + threadIdx.x)"] end ``` 图 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） ```{mermaid} flowchart TD Start["线程启动: idx = blockIdx.x * blockDim.x + threadIdx.x"] --> Cond{"idx < n ?"} Cond -->|是| Compute["计算 c[idx] = a[idx] + b[idx]"] Compute --> Stride["idx += blockDim.x * gridDim.x (步进一个网格)"] Stride --> Cond Cond -->|否| End["线程退出"] ``` 图 1-2：Grid-Stride Loop 网格跨步循环执行控制流 在朴素的初学者代码中，通常假设线程总数严格大于等于向量长度 $N$，直接使用 `if (idx < N) c[idx] = a[idx] + b[idx];`。这种写法在工业级生产环境中存在严重缺陷（无法应对超大规模 $N$，且启动过多 Block 会浪费驱动调度开销）。工业级标准范式是采用 **Grid-Stride Loop（网格跨步循环）**： ```cuda 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`/`______` 所在位置；不要删除断言或放宽误差。代码注释说明了每个边界条件。 ```{python} #| course_role: exercise %%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 分； - 一票否决：结果碰巧正确但核心因果链错误、删除边界检查、把未运行结果说成实测。需要提示时按四级机制请求：概念区域 → 具体方向 → 关键局部 → 完整答案。 ::: {.callout-tip collapse="true" title="💡 点击展开参考答案与详细代码实现" icon="false"} ## 参考答案（仅 answer 分支）先完成题目再核对。即使代码一致，也要能解释关键步骤，并尝试更换一个输入规模。 ```{python} #| course_role: solution %%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](https://docs.nvidia.com/cuda/cuda-programming-guide/) - [CUDA Best Practices Guide](https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/) 资料用于建立事实基线；面试回答仍需用自己的语言组织。 ::: ::: {.course-pagination} [**← 上一课** 第 16 课：非均质流水线调度](../train/lesson16_microbatch_throughput.qmd){.pagination-link .pagination-prev} [**下一课 →** 第 2 课：Warp/Block 归约与两阶...](lesson02_reduce_sum.qmd){.pagination-link .pagination-next} :::