<<<CUDA C++ 学习路线grid · block · warp · lane
EN中文
阶段 1 · 起步03 / 15约 20 分钟cuda/02_vector_add.cu

显存管理与错误检查：写出不骗人的 CUDA 程序

cudaMalloc / cudaMemcpy 的正确姿势、统一内存、以及那个你必须抄一遍的宏

学完这一课你会
  • 掌握显式显存管理的完整流程和常见陷阱
  • 写出一个可复用的 CUDA 错误检查宏并理解它为什么必要
  • 知道统一内存（Unified Memory）适合和不适合什么场景
  • 理解 pinned memory 为什么能让拷贝更快

在经典的 CUDA 模型里，host 和 device 各有各的地址空间。float* d_a 里存的是显存地址，你在 CPU 代码里对它解引用会直接段错误。反过来，把 host 指针传进 kernel 也一样会崩。管好这两类指针，是 CUDA 编程的基本功。

完整的四步流程

02_vector_add.cu 的主干
1int n = 1 << 20;2size_t bytes = n * sizeof(float);3 4std::vector<float> h_a(n, 1.0f), h_b(n, 2.0f), h_c(n, 0.0f);5 6// 1. 在显存上分配7float *d_a, *d_b, *d_c;8CUDA_CHECK(cudaMalloc(&d_a, bytes));      // 注意是 &d_a：要改的是指针本身9CUDA_CHECK(cudaMalloc(&d_b, bytes));10CUDA_CHECK(cudaMalloc(&d_c, bytes));11 12// 2. 输入搬到显存13CUDA_CHECK(cudaMemcpy(d_a, h_a.data(), bytes, cudaMemcpyHostToDevice));14CUDA_CHECK(cudaMemcpy(d_b, h_b.data(), bytes, cudaMemcpyHostToDevice));15 16// 3. 算17int threads = 256, blocks = (n + threads - 1) / threads;18vecAdd<<<blocks, threads>>>(d_a, d_b, d_c, n);19CUDA_CHECK(cudaGetLastError());           // 抓 launch 失败（配置非法等）20 21// 4. 结果搬回来。同步版 cudaMemcpy 会隐式等待 kernel 完成22CUDA_CHECK(cudaMemcpy(h_c.data(), d_c, bytes, cudaMemcpyDeviceToHost));23 24CUDA_CHECK(cudaFree(d_a));25CUDA_CHECK(cudaFree(d_b));26CUDA_CHECK(cudaFree(d_c));

那个你必须抄一遍的宏

几乎所有 CUDA Runtime API 都返回 cudaError_t。不检查返回值是 CUDA 调试痛苦的头号原因：错误会一直沉默地向后传播，直到在一个和根因毫无关系的地方爆出来。下面这个宏是所有正经 CUDA 项目的标配，建议直接放进你的 common.cuh。

cuda/common.cuh
1#pragma once2#include <cstdio>3#include <cstdlib>4 5#define CUDA_CHECK(call)                                                   \6    do {                                                                   \7        cudaError_t err_ = (call);                                         \8        if (err_ != cudaSuccess) {                                         \9            fprintf(stderr, "CUDA error %s:%d: '%s' -> %s\n",              \10                    __FILE__, __LINE__, #call, cudaGetErrorString(err_));  \11            std::exit(EXIT_FAILURE);                                       \12        }                                                                  \13    } while (0)14 15// kernel launch 不返回错误码，要单独抓两类错误：16//   cudaGetLastError()      —— 配置非法之类的同步错误17//   cudaDeviceSynchronize() —— kernel 执行期间的异步错误（越界、非法指令）18#define CUDA_CHECK_KERNEL()                     \19    do {                                        \20        CUDA_CHECK(cudaGetLastError());         \21        CUDA_CHECK(cudaDeviceSynchronize());    \22    } while (0)

统一内存：省事，但不是免费的

cudaMallocManaged 分配的内存 host 和 device 都能直接访问，缺页时由驱动在后台迁移页面。代码能短一大截，很适合原型和教学。但它的代价是页面迁移的开销不可见——一个访问模式不好的 kernel 可能在后台疯狂搬页，而你从代码里完全看不出来。

统一内存版本
1float *a, *b, *c;2CUDA_CHECK(cudaMallocManaged(&a, bytes));3CUDA_CHECK(cudaMallocManaged(&b, bytes));4CUDA_CHECK(cudaMallocManaged(&c, bytes));5 6for (int i = 0; i < n; ++i) { a[i] = 1.0f; b[i] = 2.0f; }   // CPU 直接写7 8// 有明确访问意图时给驱动一点提示，能显著减少缺页抖动9int device = 0;10CUDA_CHECK(cudaMemPrefetchAsync(a, bytes, device));11CUDA_CHECK(cudaMemPrefetchAsync(b, bytes, device));12 13vecAdd<<<blocks, threads>>>(a, b, c, n);14CUDA_CHECK(cudaDeviceSynchronize());        // 必须同步后 CPU 才能安全读 c15 16printf("c[0] = %f\n", c[0]);17cudaFree(a); cudaFree(b); cudaFree(c);
方式适合不适合
cudaMalloc + cudaMemcpy生产代码、需要精确控制传输时机、要做流水线重叠快速原型（样板代码多）
cudaMallocManaged原型、教学、数据结构含大量指针、超出显存的超大数据集对延迟敏感的热路径（迁移开销难以预测）
cudaHostAlloc（pinned）需要异步拷贝、想榨干 PCIe 带宽分配巨量内存（锁页会挤压系统可用内存）

pinned memory 为什么更快

普通 malloc 出来的 host 内存是可分页的，操作系统随时可能把它换出去。GPU 的 DMA 引擎没法直接读这种内存，所以驱动的实际做法是：先把数据拷到一块内部的锁页缓冲区，再由 DMA 搬到显存——白白多了一次 CPU 拷贝。用 cudaHostAlloc 直接分配锁页内存，DMA 就能一步到位，带宽通常能提升到接近 PCIe 理论值，而且这是使用 cudaMemcpyAsync 的前提条件。

pinned memory
1float* h_pinned;2CUDA_CHECK(cudaHostAlloc(&h_pinned, bytes, cudaHostAllocDefault));3 4// 只有 pinned 内存才能真正异步拷贝；传可分页内存给 Async 版本5// 不会报错，但驱动会退化成同步行为，重叠优化直接失效。6CUDA_CHECK(cudaMemcpyAsync(d_a, h_pinned, bytes, cudaMemcpyHostToDevice, stream));7 8CUDA_CHECK(cudaFreeHost(h_pinned));

自测

先自己在心里回答一遍，再展开对照。答不上来的说明这一段值得重读。