向量加法：第一个并行 Kernel

从主机内存（Host Memory）到设备内存（Device Memory），第一次走通 CUDA 数据闭环。

核心判断：向量加法真正教给你的不是 c[i] = a[i] + b[i]，而是 CUDA 的第一条完整数据链路：Host 数据如何进入 Device、线程如何通过全局索引找到自己的数据、内核（Kernel）如何产生 Device 结果，以及结果如何回到 Host。

后续绝大多数 CUDA 优化，都会不断回到两个问题：数据在哪里？数据搬了几次？

本篇写作目标

本篇不追求性能优化，而是让读者第一次完整走通 CUDA 程序的数据生命周期：

读完本篇，读者应该能够回答三个问题：

1. GPU Kernel 的输入数据从哪里来？
2. 一个线程如何找到自己负责的数据？
3. GPU 算完之后，结果如何回到 CPU？

① 为什么“能算出来”才算入门

前面几篇已经解决了两个问题：

• Thread 如何找到自己的位置；
• 内核启动（Kernel Launch）后 GPU 如何执行。

但还有一个更基础的问题：

数据在哪里？

CPU 程序里，数组就像放在手边的书桌上，伸手就能读：

for (int i = 0; i < N; i++) {
    c[i] = a[i] + b[i];
}

而 CUDA 里，数据在另一栋楼的柜子里——CPU 书桌和 GPU 那栋楼是两套不同的地址空间，得先有人把书搬过去、算完再搬回来。

严格地说，本文采用经典 CUDA 显式内存管理模型：Kernel 使用 设备全局内存（Device Global Memory） 中的数据，而输入数据最初位于 Host Memory，因此需要通过 cudaMemcpy 完成 Host↔Device 数据传输。CUDA 还有统一内存（Unified Memory）、映射主机内存（Mapped Host Memory）等其他内存模型，本篇暂不展开。

本文为了建立最基础、最清晰的 CUDA 心智模型，采用最常见的路径。

因此，向量加法真正需要解决的不是“加法怎么算”，而是：

这就是本文第一次完整建立的 Host↔Device 数据闭环。

💡 第一性原理：CUDA 程序不是简单地把 CPU 的 for 循环搬到 GPU，而是首先要解决“数据在哪里、谁来处理、结果在哪里”的完整生命周期。

本篇边界：学 Host Memory / Device Global Memory / cudaMalloc / cudaMemcpy / Global Index / 边界检查（Boundary Check）；暂不学锁页内存（Pinned Memory）、Unified Memory、流（Streams）、共享内存（Shared Memory）、合并访存、占用率（Occupancy）、Nsight 等——一次只解决一个层次的问题。

② 串行思维 vs 并行思维

先看 CPU：

for (int i = 0; i < N; i++) {
    c[i] = a[i] + b[i];
}

一个 CPU 线程负责从 0 到 N-1 的全部元素。

GPU 则把任务拆开：

Thread 0   → c[0]
Thread 1   → c[1]
Thread 2   → c[2]
...
Thread N-1 → c[N-1]

前文已经知道全局坐标公式：

int i = blockIdx.x * blockDim.x + threadIdx.x;

因此可以直接建立：

线程坐标
   ↓
数组下标
   ↓
数据元素

每个线程只处理自己的元素：

c[i] = a[i] + b[i];

对于这个向量加法，每个线程读写独立的 c[i]，线程之间不存在写同一个结果位置的竞争，因此不需要锁或原子操作。

下面这张图把「网格（Grid）→ 线程块（Block）→ 线程（Thread）」的层级，和「线程坐标如何对应到数据位置」放在了一起，帮你把前文的线程层次和 Kernel 执行模型接到同一个画面里：

图里最该记住的是：每个线程拿到自己的坐标后，用坐标当数组下标去读写数据；线程之间靠「各算各的元素」来避免冲突，而不是靠锁。

这就是最简单的 GPU 并行模式：一个线程负责一个独立元素。

坐标会算了，但坐标指向的内存此刻还躺在 CPU 那栋楼里。下一步要把数据搬过去，再让每个线程按坐标取数——这正是 ③ 要走的完整闭环。

③ 最小可运行实现

完整闭环可以拆成七步（比「六步」多出来的，是结果校验——程序跑完不等于算对）：

Host 与 GPU 两侧的数据流向可以这样看：

#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>

__global__ void vectorAdd(
    const float* a,
    const float* b,
    float* c,
    int n
) {
    // 一维网格下的全局坐标公式
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    // 边界保护：线程总数可能多于数据量
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

int main() {
    const int N = 1'000'003;               // 故意取非 256 整数倍，让边界保护真实触发
    const size_t bytes = N * sizeof(float);

    // ① Host Memory
    float* h_a = (float*)malloc(bytes);
    float* h_b = (float*)malloc(bytes);
    float* h_c = (float*)malloc(bytes);
    for (int i = 0; i < N; i++) {
        h_a[i] = (float)i;
        h_b[i] = (float)(i * 2);
    }

    // ② Device Memory
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);

    // ③ Host → Device
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

    // ④ Kernel Launch：向上取整保证覆盖全部元素
    int threadsPerBlock = 256;             // 256 只是常见教学配置，不是 CUDA 固定值
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, N);

    // 在关键节点检查错误：先看 launch 是否返回错误（统一封装留到下一篇）
    cudaError_t e = cudaGetLastError();
    if (e != cudaSuccess) {
        std::fprintf(stderr, "Kernel launch failed: %s\n",
                     cudaGetErrorString(e));
        return 1;
    }

    // ⑤ Device → Host：阻塞式 D2H 会等待相关 Device 工作完成；
    // 检查返回值，及时发现拷贝失败（统一封装留到下一篇）
    cudaError_t e2 = cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);
    if (e2 != cudaSuccess) {
        std::fprintf(stderr, "D2H memcpy failed: %s\n",
                     cudaGetErrorString(e2));
        return 1;
    }

    // ⑥ 验证结果：程序跑完 ≠ 算对
    // 本例输入为小整数，float 可精确表示，故直接比较；
    // 一般浮点计算结果应使用容差（如 |got - expected| < 1e-5）比较
    int errors = 0;
    for (int i = 0; i < N && errors < 5; i++) {
        if (h_c[i] != h_a[i] + h_b[i]) {
            std::fprintf(stderr, "Mismatch at %d: got=%f expected=%f\n",
                         i, h_c[i], h_a[i] + h_b[i]);
            errors++;
        }
    }
    if (errors > 0) return 1;

    // ⑦ 释放
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);
    return 0;
}

一个非常重要的区别：h_a 和 d_a 不是同一块内存

h_a  → Host Memory    CPU 可以直接访问
d_a  → Device Memory  Kernel 可以访问

cudaMemcpy() 不是「改变指针指向」，而是把数据从一块内存复制到另一块内存：

所以 Kernel 用的是 d_a / d_b / d_c，而不是 h_a / h_b / h_c：

vectorAdd(d_a, d_b, d_c, N);

这是 CUDA 初学者要建立的第一个内存心智模型：CPU 侧一套指针、Device 侧另一套指针，二者靠 cudaMemcpy 传数据。

这里先埋一个最常见的坑：新手写完 vectorAdd<<<...>>> 常直接 cudaMemcpy 往回搬，却忘了前面的 cudaMemcpy(H2D)——Kernel 读取的是尚未正确初始化的 Device Memory，结果是未定义的；这种错误尤其容易让初学者误以为「Kernel 算错了」，其实算的是空数据。这正是「数据闭环」断在第一步的典型症状。

Kernel 没报错 ≠ 结果正确：即便 Kernel 正常结束，索引算错、Block 数错、数据初始化错，都会让程序退出但结果错误。所以 ③ 多了「验证结果」这一步。

注意这里没有额外调用 cudaDeviceSynchronize()。本文使用的是从 Device 到 pageable Host Memory（可分页内存——由操作系统按需换页的普通 malloc 内存，区别于锁在物理内存、支持真正异步拷贝的 pinned memory）的同步式 cudaMemcpy()。在这种执行路径下，D2H 拷贝返回前会确保此前相关的 Device 工作已经完成，因此 Host 在拿到 h_c 时可以直接读取结果。

但要注意：cudaMemcpy() 的核心职责是数据复制，不是通用的 GPU 同步原语。 不同内存类型、异步 API、Stream 使用方式下，它的同步语义会有所不同——后续学习 Streams 与异步内存传输时再展开。

而 cudaGetLastError() 常用于在 Kernel launch 后检查启动阶段的错误（例如启动配置（launch configuration）无效、device function 无效等）；它不会等待 GPU 执行完成，因此 Kernel 执行阶段的错误（如 Kernel 内部越界访问显存）无法靠单独一次 cudaGetLastError() 可靠捕获。这正是下一篇统一错误检查要覆盖的两类问题：

CUDA_CHECK(cudaGetLastError());       // 覆盖 launch 阶段错误
CUDA_CHECK(cudaDeviceSynchronize());  // 覆盖执行完成后的错误

但需要强调：教学代码有意省略了 cudaMalloc() 和 H2D 拷贝的逐项错误检查，生产代码不应该这样写。 如果 cudaMalloc() 失败，d_a 等指针状态异常，错误往往要等到更晚的 H2D 或 Kernel 阶段才暴露。本文只在最关键的两个节点（Kernel launch、D2H 拷贝）检查了返回值，是为了保持主线简洁；下一篇会统一用 CUDA_CHECK 包装所有 Runtime API 调用，并用 compute-sanitizer 检查非法内存访问。

待 N 卡验证：本文代码未在当前 Apple Silicon 环境编译运行。

④ 两个必须记住的 CUDA 惯用法（idiom）

4.1 Grid 向上取整

线程总数通常不需要刚好等于数据量。例如 N = 10、T = 4，需要 ceil(10 / 4) = 3 个 Block，因此：

int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

这是 CUDA 中极其常见的向上取整写法 (N + T - 1) / T。不要写成 N / T——否则 10 / 4 = 2，只能覆盖 0～7，最后两个元素永远不会被处理。

4.2 边界保护

因为「线程数量 ≥ 数据数量」是非常常见的情况，所以 Kernel 中必须保护数组边界：

if (i < n) {
    c[i] = a[i] + b[i];
}

例如 N = 10、每 Block 4 个线程时：

Block 0 → 0 1 2 3
Block 1 → 4 5 6 7
Block 2 → 8 9 10 11

其中 10、11 没有对应的数据，if (i < 10) 会让它们直接跳过，避免越界写显存。

💡 隐性知识：(N + T - 1) / T 与 if (i < N) 通常应该成对出现。前者保证覆盖全部数据，后者保证多出来的线程不会越界。单独写一个都会埋雷。

两个 idiom 记牢，代码在「正确性」上就稳了。但正确性不等于「该不该用 GPU」——接下来看性能。

⑤ GPU 为什么不一定更快

为了建立第一版性能心智模型，可以把本文这种简单、串行的执行路径近似写成：

T_gpu ≈ T_H2D + T_launch + T_kernel + T_D2H

这是本文用来建立心智模型的简化端到端时间模型：实际 CUDA 程序还可能受到 Host 调度、内存分配、初始化等因素影响，且采用异步执行时 H2D、Kernel、D2H 并不总是串行排列，中间可能重叠。

也就是说：

GPU 的并行计算优势不能脱离数据传输成本讨论。

当前向量加法实际上发生了：

也就是「2 次 H2D + 1 次 D2H + 1 次 Launch + Kernel 执行」。GPU 是否更快，不由数据规模单独决定，而取决于并行计算收益能否覆盖数据传输与 Kernel Launch 成本（D2H 等待结果的时间已包含在传输时间里，不应再单列一笔「同步成本」）：

当 并行计算收益 > 数据传输 + Launch + 计算本身成本
GPU 才可能取得端到端优势

数据规模只是影响因素之一。向量加法的计算本身非常简单——每个元素只做一次加法，却要读两个输入、写一个输出，它通常不是一个「计算特别重」的 Kernel。但本篇暂不展开算术强度与内存带宽分析，后续学习 Global Memory 与屋顶线模型（Roofline Model）时再回答「为什么 vector add 很容易受内存带宽限制」。

下面这张图展示了典型离散 GPU 系统中 Host 与 Device 之间的互连：在常见的离散 GPU 系统里，Host Memory 与 Device Memory 位于不同的处理器/内存系统，数据通常经 PCIe、NVLink 等互连路径传输；具体路径取决于 GPU、CPU 与平台架构（例如集成 GPU、Grace Hopper 等并不严格符合「CPU → PCIe → GPU」）。

图里要观察的是「CPU 内存 ↔ 互连路径 ↔ GPU 显存」这条链路：数据搬运通常要经过这条互连路径，但具体走哪条路取决于平台架构——并非所有 CUDA 程序都严格符合「CPU → PCIe → GPU」。

可以把这条互连路径想成一座收费桥：不管你运 1 箱还是 1000 箱货，每次过桥都要先付一笔固定的「起步费」（传输延迟 + launch 开销）。数据规模只是影响因素之一，真正决定 GPU 是否值得使用的，是并行计算收益能否覆盖数据传输、Kernel Launch 以及计算本身的成本。

工程判断：判断 GPU 是否值得使用，首先要问“数据需要搬几次”，其次才是“Kernel 算得多快”。

⑥ 为什么本文暂时不做 Nsight

本文的目标是建立：

Host → Device → Kernel → Device → Host

的数据闭环。因此暂时不展开 Nsight Systems、Nsight Compute、内存吞吐（Memory Throughput）、占用率（Occupancy）、线程束效率（Warp Efficiency）这些指标。

原因很简单：

在读者还不知道数据在哪里之前，讨论 GPU 带宽利用率没有意义。

后续会逐步进入「能跑 → 跑对 → 跑得快 → 知道为什么快」的工程路线，Nsight 会在那时自然登场。

⑦ 工业应用：向量加（Vector Add）为什么是逐元素（element-wise）的起点

向量加法 c[i] = a[i] + b[i] 是最典型的 element-wise 模式之一。这类算子的通用形态是：

output[i] = f(input0[i], input1[i], ...)

只看对应位置的输入，逐元素计算输出。ReLU、SiLU、逐元素乘法、逐元素乘加、向量加法都是同一模式。它们都有一个共同特点：output[i] 只依赖对应位置的输入，不同 i 之间没有前后依赖（不存在「c[i] 依赖 c[i-1]」的情形），因此大量线程可以同时执行，天然适合 GPU 并行——这也是本篇最简单直接的实现方式。

注意：「一线程一元素」只是 L1 的教学映射，不是 CUDA 的唯一写法。工业实现里一个线程往往要处理多个元素（如 float4 向量化读写、一次搬 4 个元素），或一个线程只算输出的一个分量（如张量核心（Tensor Core）的矩阵计算），映射关系取决于算子的访存与计算特征，后续各篇会逐步展开。

注意：LayerNorm 并不是纯粹的 element-wise 算子。它需要先通过归约（reduction）算出 mean/variance，再进行逐元素归一化；后续学习 Reduction 时再展开。

⑧ 三个面试问题

为什么小数据下 GPU 不一定比 CPU 快？ 因为 GPU 端到端时间不仅包含 Kernel 执行，还包含 Host↔Device 数据传输和 Kernel Launch 开销；只有当并行收益盖过这些固定成本时，GPU 才可能更快（详见 ⑤）。

(N + T - 1) / T 在干什么？ 整数除法里的向上取整 ceil(N / T)，保证线程数量至少覆盖全部数据；再配合 if (i < N) 挡住多余线程，避免越界。两者通常成对出现（见 ④）。

CUDA Kernel 为什么通常通过输出缓冲区返回结果？ Kernel launch 本质是 Host 向 Device 发起一段异步工作，不是普通的同步函数调用，因此没有「返回值」可用。常见路径是：Host launch → GPU Kernel 把结果写入 Device Memory → cudaMemcpy(D2H) 搬回 Host。所以「Kernel → Device Output → Host」是最常见的结果返回路径。

⑨ 延伸阅读

收藏速查表

口诀：数据先上 GPU，线程再找位置；Kernel 算完以后，结果再回来。

延伸阅读

• 前文：线程层次与 Block/Thread 坐标；
• 前文：Kernel 生命周期与异步发射；
• 下一篇：CUDA Memory Management——cudaMalloc、cudaMemcpy、错误检查与 compute-sanitizer；
• 后续：Global Memory 与合并访存；
• 后续：Shared Memory 与数据复用；
• 后续：Element-wise Kernel 与算子融合。

💡 本篇真正需要记住的不是 vectorAdd，而是这条完整链路：

CUDA 的绝大多数工程优化，最终都可以追溯到一个问题：数据在哪里，以及数据搬了几次。

下一篇：CUDA Memory Management——从 cudaMalloc 到 compute-sanitizer，把“数据在哪里”彻底搞清楚。

你第一次写 CUDA 程序时，是卡在了「数据搬不过去」，还是卡在了「结果搬不回来」？欢迎在评论区聊聊你踩的第一个坑。