向量加法：第一个并行 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(inti=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(inti=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]

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

inti=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__voidvectorAdd(constfloat*a,constfloat*b,float*c,intn){// 一维网格下的全局坐标公式inti=blockIdx.x*blockDim.x+threadIdx.x;// 边界保护：线程总数可能多于数据量if(i<n){c[i]=a[i]+b[i];}}intmain(){constintN=1'000'003;// 故意取非 256 整数倍，让边界保护真实触发constsize_t bytes=N*sizeof(float);// ① Host Memoryfloat*h_a=(float*)malloc(bytes);float*h_b=(float*)malloc(bytes);float*h_c=(float*)malloc(bytes);for(inti=0;i<N;i++){h_a[i]=(float)i;h_b[i]=(float)(i*2);}// ② Device Memoryfloat*d_a,*d_b,*d_c;cudaMalloc(&d_a,bytes);cudaMalloc(&d_b,bytes);cudaMalloc(&d_c,bytes);// ③ Host → DevicecudaMemcpy(d_a,h_a,bytes,cudaMemcpyHostToDevice);cudaMemcpy(d_b,h_b,bytes,cudaMemcpyHostToDevice);// ④ Kernel Launch：向上取整保证覆盖全部元素intthreadsPerBlock=256;// 256 只是常见教学配置，不是 CUDA 固定值intblocksPerGrid=(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));return1;}// ⑤ 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));return1;}// ⑥ 验证结果：程序跑完 ≠ 算对// 本例输入为小整数，float 可精确表示，故直接比较；// 一般浮点计算结果应使用容差（如 |got - expected| < 1e-5）比较interrors=0;for(inti=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)return1;// ⑦ 释放cudaFree(d_a);cudaFree(d_b);cudaFree(d_c);free(h_a);free(h_b);free(h_c);return0;}

一个非常重要的区别：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，因此：

intblocksPerGrid=(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 程序时，是卡在了「数据搬不过去」，还是卡在了「结果搬不回来」？欢迎在评论区聊聊你踩的第一个坑。