向量加法：第一个并行 Kernel

从主机内存Host Memory到设备内存Device Memory第一次走通 CUDA 数据闭环。核心判断向量加法真正教给你的不是c[i] a[i] b[i]而是 CUDA 的第一条完整数据链路Host 数据如何进入 Device、线程如何通过全局索引找到自己的数据、内核Kernel如何产生 Device 结果以及结果如何回到 Host。后续绝大多数 CUDA 优化都会不断回到两个问题数据在哪里数据搬了几次本篇写作目标本篇不追求性能优化而是让读者第一次完整走通 CUDA 程序的数据生命周期读完本篇读者应该能够回答三个问题GPU Kernel 的输入数据从哪里来一个线程如何找到自己负责的数据GPU 算完之后结果如何回到 CPU① 为什么“能算出来”才算入门前面几篇已经解决了两个问题Thread 如何找到自己的位置内核启动Kernel Launch后 GPU 如何执行。但还有一个更基础的问题数据在哪里CPU 程序里数组就像放在手边的书桌上伸手就能读for(inti0;iN;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 并行思维先看 CPUfor(inti0;iN;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]前文已经知道全局坐标公式intiblockIdx.x*blockDim.xthreadIdx.x;因此可以直接建立线程坐标 ↓ 数组下标 ↓ 数据元素每个线程只处理自己的元素c[i]a[i]b[i];对于这个向量加法每个线程读写独立的c[i]线程之间不存在写同一个结果位置的竞争因此不需要锁或原子操作。下面这张图把「网格Grid→ 线程块Block→ 线程Thread」的层级和「线程坐标如何对应到数据位置」放在了一起帮你把前文的线程层次和 Kernel 执行模型接到同一个画面里图里最该记住的是每个线程拿到自己的坐标后用坐标当数组下标去读写数据线程之间靠「各算各的元素」来避免冲突而不是靠锁。这就是最简单的 GPU 并行模式一个线程负责一个独立元素。坐标会算了但坐标指向的内存此刻还躺在 CPU 那栋楼里。下一步要把数据搬过去再让每个线程按坐标取数——这正是 ③ 要走的完整闭环。③ 最小可运行实现完整闭环可以拆成七步比「六步」多出来的是结果校验——程序跑完不等于算对Host 与 GPU 两侧的数据流向可以这样看#includecstdio#includecstdlib#includecuda_runtime.h__global__voidvectorAdd(constfloat*a,constfloat*b,float*c,intn){// 一维网格下的全局坐标公式intiblockIdx.x*blockDim.xthreadIdx.x;// 边界保护线程总数可能多于数据量if(in){c[i]a[i]b[i];}}intmain(){constintN1000003;// 故意取非 256 整数倍让边界保护真实触发constsize_t bytesN*sizeof(float);// ① Host Memoryfloat*h_a(float*)malloc(bytes);float*h_b(float*)malloc(bytes);float*h_c(float*)malloc(bytes);for(inti0;iN;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向上取整保证覆盖全部元素intthreadsPerBlock256;// 256 只是常见教学配置不是 CUDA 固定值intblocksPerGrid(NthreadsPerBlock-1)/threadsPerBlock;vectorAddblocksPerGrid,threadsPerBlock(d_a,d_b,d_c,N);// 在关键节点检查错误先看 launch 是否返回错误统一封装留到下一篇cudaError_t ecudaGetLastError();if(e!cudaSuccess){std::fprintf(stderr,Kernel launch failed: %s\n,cudaGetErrorString(e));return1;}// ⑤ Device → Host阻塞式 D2H 会等待相关 Device 工作完成// 检查返回值及时发现拷贝失败统一封装留到下一篇cudaError_t e2cudaMemcpy(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比较interrors0;for(inti0;iNerrors5;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(errors0)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_cvectorAdd(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 惯用法idiom4.1 Grid 向上取整线程总数通常不需要刚好等于数据量。例如N 10、T 4需要ceil(10 / 4) 3个 Block因此intblocksPerGrid(NthreadsPerBlock-1)/threadsPerBlock;这是 CUDA 中极其常见的向上取整写法(N T - 1) / T。不要写成N / T——否则10 / 4 2只能覆盖07最后两个元素永远不会被处理。4.2 边界保护因为「线程数量 ≥ 数据数量」是非常常见的情况所以 Kernel 中必须保护数组边界if(in){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」是最常见的结果返回路径。⑨ 延伸阅读收藏速查表概念一句话理解跳过后果Host MemoryCPU 侧数据所在的内存不知道输入数据在哪里Device Global MemoryGPU Kernel 使用的显式存储空间本文用 Global Memory不知道 Kernel 从哪里读数据cudaMemcpy在 Host/Device 之间复制数据不知道数据如何跨设备移动Global IndexblockIdx.x * blockDim.x threadIdx.x线程不知道自己处理哪个元素向上取整(N T - 1) / T尾部元素可能漏算边界保护if (i N)多余线程可能越界Kernel Launch把计算提交给 GPU容易把提交误认为完成H2D / D2HHost→Device / Device→HostGPU 拿不到输入 / CPU 拿不到结果端到端时间≈ H2D Launch Kernel D2H简化模型容易误判 GPU 性能口诀数据先上 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 程序时是卡在了「数据搬不过去」还是卡在了「结果搬不回来」欢迎在评论区聊聊你踩的第一个坑。