CUDA 内存管理：从 cudaMalloc/cudaMemcpy 到 CUDA_CHECK

让错误在第一现场暴露——建立 CUDA API → 内核（Kernel）→ 净化器（Sanitizer）的错误观察链，补齐 CUDA 数据链路的最后一块地基。

核心判断：前文走通了 Host↔Device 数据闭环，但那个闭环少了一道保险——出错时怎么办。CUDA 的显式内存管理把分配、搬运、释放交给程序员，而内核启动（Kernel launch）又具有异步执行语义，因此「错误发生的位置」和「Host 观察到错误的位置」可能不同。工程上真正需要建立的不是「多写几个 if」，而是一套明确的错误观察链：API 返回值检查 → Kernel Launch 检查 → 执行完成检查 → Sanitizer 深度诊断。本篇就用一个宏（CUDA_CHECK）加一个模型（错误状态），把这条链立起来。

① CUDA 最难查的 Bug，往往出在内存上

上一篇结束时，读者已经能写出一个「能跑」的 vectorAdd。但「能跑」和「可靠」之间，隔着整整一类问题：内存错误。

先看一个典型认知误区：

「程序没崩，就是没出错。」

对 CPU 程序，这句话也只在「大体」层面成立。CPU 越界访问同样不保证立即崩溃——越界写可能悄悄破坏相邻对象，几十行代码之后才以莫名其妙的方式崩掉。差别在于：CUDA 的主机（Host）与设备（Device）执行异步、地址空间不同，错误的发生与报告之间更容易产生时间和空间上的分离。Kernel 在 GPU 上异步执行，越界写显存、读未初始化内存这些错误发生的那一刻，CPU 端毫无知觉。程序继续「正常」运行，直到某个检查点才把错误吐出来；更糟的情况是永远不吐，只是结果悄悄错了。

为什么 CUDA 会这样？两个机制叠加：

统一根源是：错误的发生时刻 ≠ 错误的报告时刻。错误先发生在 GPU 内部，报告却在之后的检查点。等错误终于浮出时，现场早被后续计算覆盖了——这就是内存 bug 难查的本质。

💡 第一性原理：CPU 与 GPU 具有不同的内存访问域和物理内存层次（物理约束）→ 在经典显式 Device Memory 模型下，数据位置需要显式管理，分配、搬运、生命周期都成为程序责任（必然需求）→ 每一步都可能失败，而异步执行又使错误观察存在延迟（工程代价）→ 因此需要统一错误检查链。

② CPU 思维 vs GPU 思维：Host 内存与 Device 内存

上一节讲的是「错误为什么难查」，这一节回答更基础的问题：数据为什么必须分开管理？ 答案藏在地址空间里。

CPU 程序里，malloc 分到的内存和普通变量住在同一个地址空间，读写无需特别手续。GPU 侧多了一个「搬运」步骤：在本文采用的经典显式 Device Memory 模型下，Host 代码不能把设备全局内存（Device Global Memory）当作普通 Host 指针直接解引用，Kernel 也不能把普通可分页主机内存（pageable Host Memory）当作 Device Global Memory 使用——因此需要通过 CUDA 提供的内存访问/拷贝机制完成数据交换。CUDA 还提供统一内存（Unified Memory）、映射主机内存（Mapped Host Memory）等其他内存模型，本篇暂不展开。

三个 API 对应关系，以及每个 API 真正要防的问题：

对初学者来说，cudaFree 最隐蔽的坑其实是 use-after-free：释放 d_a 之后 Kernel 还在用它，例如：

cudaFree(d_a);
doubleArray<<<blocks, threads>>>(d_a, N); // 用已释放的指针

这种错误比 double-free 更常见——程序不崩、错误不报，结果却是错的。所以释放前要确认：不再有排队的 Kernel 或异步操作引用这块内存。

最关键的一个差异，也是经典显式 Device Memory 模型的核心：Host 与 Device 使用各自可访问的内存空间，h_a 指向 Host Memory，d_a 指向 Device Memory，二者不是同一块可直接互换的存储。在这个模型下，输入数据通常需要同时存在于 Host 与 Device 两侧，因此 cudaMemcpy 承担数据搬运职责。本篇在此基础上补一句：搬运和分配本身都可能失败，而失败只通过返回值告诉你，不会中断程序。

以 cudaMalloc 为例：

float* d_a = nullptr;
cudaError_t err = cudaMalloc(&d_a, bytes);
if (err != cudaSuccess) {
fprintf(stderr, "cudaMalloc 失败: %s\\n", cudaGetErrorString(err));
return 1;
}

这段代码很容易写错两处：一是忘记检查返回值，二是 &d_a 写成了 d_a。后者编译器不一定报错，运行时才炸。

注意，这里的错误检查只覆盖 CUDA Runtime API。Host 侧自己的内存分配（malloc）同样可能失败，返回 nullptr——需要用 C/C++ 自己的机制检查。两类 API 各有各的错误约定：

不要用 CUDA_CHECK 去包 malloc，也不要用 if (ptr == nullptr) 去检查 cudaMalloc——各归各的。

但问题来了：CUDA Runtime API 的检查每个都要手写一遍 if (err != cudaSuccess)，既啰嗦又容易漏。而且手写检查一旦漏掉一个调用，错误的防线就出现缺口。接下来用一个宏统一封装，把「检查」变成每个调用的默认动作。

现代 CUDA 内存管理：你现在学的是经典模型

本文使用 cudaMalloc/cudaFree + cudaMemcpy，这是理解 CUDA Device Memory 生命周期最重要的经典模型。但现代 CUDA 已经提供流有序内存分配器（Stream-Ordered Memory Allocator）：cudaMallocAsync / cudaFreeAsync，并通过 cudaMemPool_t 实现内存池复用——它允许内存分配/释放与流（Stream）建立顺序关系，减少传统 cudaMalloc/cudaFree 带来的同步和分配开销。此外，现代平台存在统一虚拟地址空间（UVA），Unified Memory 可以自动在 CPU/GPU 间迁移，cudaMemcpyAsync + 锁页主机内存（pinned host memory） 才能实现真正的异步拷贝。本文以经典模型为主，这些现代机制在 Stream / 异步执行专题展开。

③ 最小实现：完整内存生命周期 + CUDA_CHECK

先看 vectorAdd 的内存部分，用宏统一包裹后长什么样：

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

// CUDA_CHECK：统一错误检查宏，出错时打印 文件:行号 + 错误描述
#define CUDA_CHECK(call) \\
do { \\
cudaError_t err = (call); \\
if (err != cudaSuccess) { \\
std::fprintf(stderr, "CUDA 错误 %s:%d: %s\\n", __FILE__, __LINE__, \\
cudaGetErrorString(err)); \\
std::exit(1); \\
} \\
} while (0)

__global__ void doubleArray(int* d_array, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) d_array[i] *= 2;
}

int main() {
const int N = 1'000'003;
const size_t bytes = N * sizeof(int);

// ① Host 内存：malloc + 初始化（Host API 用 if 检查）
int* h_array = (int*)malloc(bytes);
if (!h_array) {
std::fprintf(stderr, "Host allocation failed\\n");
return 1;
}
for (int i = 0; i < N; i++) h_array[i] = i + 1;

// ② Device 内存：cudaMalloc（注意 &d_array）
int* d_array = nullptr;
CUDA_CHECK(cudaMalloc(&d_array, bytes));

// ③ Host → Device：把输入搬上 GPU
CUDA_CHECK(cudaMemcpy(d_array, h_array, bytes, cudaMemcpyHostToDevice));

// ④ Kernel launch：launch 后立刻查 launch 阶段错误
int threads = 256;
int blocks = (N + threads – 1) / threads;
doubleArray<<<blocks, threads>>>(d_array, N);
CUDA_CHECK(cudaGetLastError());

// ⑤ 同步并检查：执行期错误在这里浮出
CUDA_CHECK(cudaDeviceSynchronize());

// ⑥ Device → Host：结果搬回
CUDA_CHECK(cudaMemcpy(h_array, d_array, bytes, cudaMemcpyDeviceToHost));

// ⑦ 验证
int errors = 0;
for (int i = 0; i < N && errors < 5; i++) {
if (h_array[i] != (i + 1) * 2) {
std::fprintf(stderr, "Mismatch at %d\\n", i);
errors++;
}
}
if (errors > 0) return 1;

// ⑧ 释放
CUDA_CHECK(cudaFree(d_array));
free(h_array);
return 0;
}

这篇代码的骨架，其实是一条贯穿全文的「错误观察链」：

Runtime Error Checking 负责发现错误，Sanitizer 负责诊断错误。 链上的每一环观察一段不同来源的错误，环环相扣才构成完整防线。下面拆开最核心的两个检查点：

CUDA Kernel 的两个检查点

启动检查（Launch Check）（④，launch 后立刻检查）

CUDA_CHECK(cudaGetLastError());

cudaGetLastError() 本质上是查询并清除当前 Host 线程的 CUDA Runtime last-error 状态。把它紧跟 Kernel launch 使用，是为了尽早观察 launch 阶段已经产生的错误（例如无效的执行配置、参数错误等）。注意两点：第一，它检查的是当前 Host 线程的 error state，理论上也可能观察到 launch 之前遗留的异步错误——「查完即清」正是为了减少这种干扰；第二，它不等待 Kernel 执行完成，因此不能证明「GPU 已经成功开始执行 Kernel」。

执行检查（Execution Check）（⑤，取结果前检查）

CUDA_CHECK(cudaDeviceSynchronize());

cudaDeviceSynchronize() 会等待当前 Device 上此前由该 Host 线程提交的相关工作完成，并返回执行过程中观察到的错误——它是一个明确的执行完成边界 + 异步错误观察点。把执行期错误（如非法内存访问）留到这个检查点，它们在这里浮出。

这个 API 同时承担两个职责，值得拆开看：

理解了「同步」与「错误观察」是两个职责，就不会再把它误当成「单纯的等待函数」——它既是并发控制原语，又是错误检查点。

Kernel launch 用 <<<>>> 语法，不能像普通函数那样返回错误码，所以 launch 之后要主动查一次错误状态（④）。需要澄清的是，这两个检查点不是「各管一类错误的唯一通道」——cudaDeviceSynchronize() 本身也可能返回此前异步 launch 或异步操作产生的错误；cudaGetLastError() 也可能观察到此前遗留的 error state。更准确的说法是：两个检查点侧重不同——Launch Check 尽早观察当前错误状态、重点定位立即可报告的启动问题；Execution Check 建立明确的执行完成边界、观察执行阶段产生的异步错误。 两者成对使用，才构成完整的错误防线。

用一张时序图看清「错误什么时候真正浮出」：

图里最该记住的是：错误发生在 GPU 执行 Kernel 时，报告却要等 CPU 下一次触碰错误状态。cudaGetLastError() 读取并清除当前 Host 线程的 Runtime error state——放在 launch 后，主要用于尽早发现立即可报告的 launch 问题；但返回 cudaSuccess 不代表 Kernel 已经执行成功，更不代表它已经开始执行。执行期错误则可能在后续 Runtime API 中被观察到；需要明确建立执行完成边界时，应使用同步机制。

关于「同步点」需要说得更准确一点：不是任意 Runtime API 都能当作可靠的执行完成检查点。某些后续 Runtime API 也可能观察到此前异步操作产生的错误，但不同 API 的同步语义各不相同（有的阻塞、有的不阻塞、有的与相关 Device 工作毫无同步关系）。正确的编程模型是：需要确认执行完成时，用 cudaDeviceSynchronize() 这类「明确等待相关 Device 工作完成」的 API，而不是依赖「碰巧某个 API 返回了错误」。

把这条原则固化下来：

注意，这是调试阶段的「明确语义」——让错误尽早、稳定地暴露；它不是性能代码的最终同步策略。到了 Stream、Event、异步拷贝、流水线这些阶段，同步时机本身就是性能设计的一部分，那是后面的内容。

调试的第二种手段：CUDA_LAUNCH_BLOCKING

除了在代码里显式加 cudaDeviceSynchronize()，CUDA 还提供了一个环境变量，让 Kernel launch 对 Host 变成同步语义：

CUDA_LAUNCH_BLOCKING=1 ./double_array

默认情况下 launch 是异步的——Host 提交后继续跑，错误要等后续同步点才暴露；设置该变量后，每个 Kernel launch 都会等待完成，错误在 launch 附近就能被观察到，非常适合排查「错误到底发生在哪个 Kernel」。它主要用于调试，不是生产可靠性的解决方案——生产代码不能靠环境变量来保证正确性。

一个现在就要记住的性能边界

cudaDeviceSynchronize() 是非常粗粒度的同步方式：它会等待当前 Device 上此前由当前 Host 线程发起的相关工作完成。它适合作为调试阶段的明确错误检查点，但不应该被理解成「每个 Kernel 后都应该调用」的标准性能写法。后续进入 Stream、Event 和异步流水线后，我们会把「为了检查错误而同步」与「为了协调工作而同步」两种完全不同的需求拆开。

④ 为什么是宏：把错误处理从「可选」变成「默认」

在宏出现之前，每个 API 都要先接返回值、再 if 判断、出错就打印退出——三行样板在代码里重复多次，既啰嗦又容易漏。对比一下：

写宏之后，一行搞定：

CUDA_CHECK(cudaMalloc(&d_a, bytes));

宏的价值不只是少打字，而是：

关于宏本身有两个工程细节值得注意：

• do { … } while (0) 不是装饰。它让宏在 if / else 语句里也能安全展开，避免悬空 else 的经典宏陷阱。
• CUDA_CHECK 在出错时直接 exit(1)。对小示例够用；生产代码往往换成「记录日志 + 清理资源 + 优雅退出」，那是后面工程化文章的内容。

还有一个容易踩的坑要提前说清：CUDA_CHECK 不是「所有 CUDA 语法都能包」的万能宏。 它只能包「返回 cudaError_t 的表达式」：

CUDA_CHECK(cudaMalloc(&d_a, bytes)); // ✅ Runtime API，返回值是 cudaError_t
CUDA_CHECK(cudaMemcpy(...)); // ✅
CUDA_CHECK(cudaFree(d_a)); // ✅

而 Kernel launch 用的是 <<<>>> 语法，它不是表达式、不返回错误码，直接包会编译报错。正确写法是 launch 之后单独查一次：

doubleArray<<<blocks, threads>>>(d_array, N); // launch 本身不返回错误码
CUDA_CHECK(cudaGetLastError()); // 单独检查 launch 阶段错误

这也是为什么 ③ 节里 launch 和错误检查是两行，而不是像 cudaMalloc 那样一行搞定。

💡 隐性知识：CUDA_CHECK(cudaGetLastError()) 与 CUDA_CHECK(cudaDeviceSynchronize()) 不是二选一，而是各管一段。漏掉 launch 后那次检查，启动期错误（如 Block 数超过上限）会延迟到同步才暴露，误导排查方向；漏掉同步后那次检查，执行期错误（如越界）可能永远不浮出。两者成对使用，才构成完整错误覆盖。

⑤ 为什么本篇不做基准测试（Benchmark）

按常规结构，这里通常会放一节性能对比，但本篇明确不做，原因有二：

所以本篇跳过 Benchmark，把篇幅留给错误处理本身。

⑥ 初识 compute-sanitizer：把隐藏的越界揪出来

CUDA_CHECK 能发现「Runtime 已经报告出来的错误」——比如越界访问在同步点返回 cudaErrorIllegalAddress。但它提供的定位信息有限：只能告诉你是哪个 Host 检查点观察到了错误（宏打印的文件:行号），无法仅凭这个错误码定位到 Kernel 内部具体哪条内存访问、哪个线程或哪个地址。真正把「GPU 行为错在哪里」挖出来的，是 compute-sanitizer：它通过动态检查进一步诊断 GPU 内存访问、竞争、初始化和同步问题。

一个必须亲手做一遍的实验：故意越界

把 ③ 的 Kernel 边界条件从 i < n 改成 i <= n，故意制造越界：

__global__ void doubleArray(int* d_array, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
// 故意越界：最后一个 block 的尾部线程会访问 d_array[n]
if (i <= n) {
d_array[i] *= 2;
}
}

d_array 只分配了 n 个元素，d_array[n] 是越界访问。现在分别观察两层防线能看到什么：

第一层：CUDA_CHECK。 运行程序（N 卡环境）：

./a.out

launch 之后的 CUDA_CHECK(cudaGetLastError()) 返回正常——launch 阶段没有错误；等到 CUDA_CHECK(cudaDeviceSynchronize())，执行期的非法访问浮出，程序报错并退出。CUDA_CHECK 能告诉你错误在哪个 Host 检查点被观察到（比如宏打印的 cudaDeviceSynchronize() 所在行），但通常不能仅凭这个 Runtime 错误码定位到 Kernel 内部具体哪条内存访问、哪个线程或哪个地址。

第二层：compute-sanitizer。 在它下面跑同一个程序。为了让报告关联到源码位置，编译时带上 -lineinfo：

nvcc -lineinfo double_array.cu -o double_array
compute-sanitizer –tool memcheck ./double_array
compute-sanitizer –tool memcheck –leak-check full ./double_array # 额外检测显存泄漏

-lineinfo 与完整调试模式 -G 是两个不同目的：

-lineinfo
保留源码行信息
→ 便于 sanitizer 报告关联到源码
→ 对优化/性能影响小

-G
设备调试模式
→ 会显著降低优化与性能
→ 用于需要单步调试等更深入场景

日常用 sanitizer 排查，-lineinfo 是更合适的基础实践。内存检查（Memcheck）会对 Kernel 中的 GPU 内存访问进行动态检查，发现非法访问时给出精确报告：

========= Invalid __global__ write of size 4
========= at 0x… in doubleArray(int*, int)
========= by thread (3,0,0) in block (4,0,0)
========= Address 0x… is out of bounds

报告尝试把错误归因到 Kernel、源码位置（配合 -lineinfo）、线程/Block 和访问类型——具体能给出多少细节，取决于错误类型与编译信息。对上例，它能定位到哪个 Kernel、哪个线程、哪个 Block、访问了哪个非法地址。

更准确地说，compute-sanitizer 不只是「告诉你错误在哪里」，它本身是一套功能正确性检查工具：Memcheck（非法/越界/未对齐内存访问，含 leak）、Racecheck（主要检测共享内存（shared memory）中的线程访问 hazard）、Initcheck（device global memory 未初始化访问）、Synccheck（同步原语使用问题）。本篇只初识 Memcheck。

把全文的防线画成一张主图，就是「CUDA Error Handling」的完整路线：

这张图的关键点是：CUDA_CHECK 不是一个「错误层级」，而是 Host/API 层的统一错误检查封装——它把「每个调用检查返回值」变成默认动作；真正的分层在它两侧：API 层错误由 CUDA_CHECK 拦下，Kernel 层错误由 cudaGetLastError() + cudaDeviceSynchronize() 观察，再由 compute-sanitizer 深度诊断。

一句话记忆（传播版）：CUDA_CHECK 负责把错误拦在现场，compute-sanitizer 负责把现场还原出来。

更严谨一点（准确版）：CUDA_CHECK 是 CUDA Runtime 错误观察/处理机制——在错误观察点捕获并处理 cudaError_t；compute-sanitizer 则进一步对 GPU 运行时行为进行动态诊断，解释「GPU 行为为什么错」。

这里值得把两个概念正式区分开：

错误检查 ≠ 错误诊断。

CUDA_CHECK 是错误检查（Error Observation）——它告诉你「某个错误被观察到了」，比如 cudaErrorIllegalAddress 意味着发生了非法地址访问；
compute-sanitizer 是错误诊断（Error Diagnosis）——它告诉你「哪个 Kernel、哪个线程、哪个 Block、哪条访问、哪个地址」。

类比：CUDA_CHECK 像报警器，响了你只知道「出事了」；sanitizer 像监控回放，能告诉你「谁、在哪个房间、干了什么」。

到这里，一个重要的认知要立起来：CUDA 的错误不是只有一种，CUDA_CHECK 也不是发现所有错误的万能机制。

\\*：非法内存访问通常在后续错误观察点（如同步点）才被 CUDA_CHECK 报告。

这张表回答了一个关键问题：为什么最终结果验证（⑦ 的 for 循环）永远不能省？ 因为数学逻辑错误和精度错误既不会让 CUDA_CHECK 报错，通常也不会让 sanitizer 报错——只有把 GPU 结果和 CPU 黄金结果逐一对上，它们才会现形。三层防线各管一段，缺一层就漏一类错误。

本篇只做「初识」：记住两层防线各自的职责、以及「故意越界」实验里它们各自能看到什么。compute-sanitizer 除了 memcheck（内存访问错误）还有 racecheck（竞争）、initcheck（初始化）等模式，覆盖不同类别的问题——这些留给调试专题，本篇不展开。

待 N 卡验证：本文代码未在当前 Apple Silicon 环境编译运行；compute-sanitizer 报告格式以 NVIDIA 工具实际输出为准。

⑦ 工业应用：生产代码的错误处理纪律

生产环境的 CUDA 代码，错误处理不是「应该做」，而是「必须做」。几条经验：

1. 正常的生产 CUDA 代码路径，应对每个可能失败的 CUDA Runtime API 建立明确的错误处理策略。 手写检查的版本在教程里可以出现，在可交付的代码库里不允许。但要澄清：「明确策略」不意味着必须用这个具体的 CUDA_CHECK + exit(1) 宏。教程里的宏适合快速定位问题；生产代码的策略可以是「统一检查 + 返回错误码」，也可以是「记录日志 + 清理资源 + 优雅退出」，甚至由框架层统一封装、上层负责重试/降级——那是 L2 → L3 的内容。本篇先立起「凡是返回 cudaError_t 的调用都要有明确策略」的纪律。

2. 调试阶段默认采用 Launch Check + Execution Check。 这是 L2 调试 / 正确性验证（Debug / Correctness Validation）阶段的默认姿势——只做 Launch Check、不做 Execution Check，等于只买了半份保险。注意：这是调试姿势，不是生产性能代码的固定模板。 进入异步流水线后，Execution Check 会根据 Stream/Event 的同步语义换成更细粒度的 cudaStreamSynchronize / cudaEventSynchronize，甚至依赖后续异步操作的流排序（stream ordering）（见第 5 条）。

3. 在 GPU CI 的专项验证阶段运行 compute-sanitizer。 很多内存错误在真实数据上偶发、在测试数据上不触发。但 compute-sanitizer 有明显性能开销，GPU CI 本身也是成本——所以不是「每个普通 CI 都跑」，而是把 memcheck（必要时 racecheck / initcheck / synccheck）放进 GPU 专项验证阶段，在合入之前把越界和竞争扼杀掉。

4. 错误信息要带上下文。 出错时打印「哪个操作、哪个缓冲区、多大尺寸」，比打印一个裸错误码有用得多。宏帮你带上了文件和行号，剩下的是操作语义。

5. 别把 cudaDeviceSynchronize() 写进生产热路径。 它是错误检查和调试的重要工具，但每次调用都会让 CPU 停下来等 GPU，破坏异步并行机会。调试阶段可以放心用；性能代码里是否同步，需要结合 Stream、Pipeline 设计决定——那是后面专题的内容。本篇教你在调试时建立检查点，不是教你在生产代码里无条件同步。

反面教材也很常见：把 CUDA_CHECK 包在函数里，然后在错误路径 return; 而不是退出——错误被吞掉，程序带着坏状态继续跑，问题被推迟到更晚才爆炸。错误处理最忌讳的是「吞掉错误」。

⑧ 三个面试问题

为什么异步错误往往要等同步点才被发现？

因为 Kernel 是异步执行的：CPU 提交 launch 后立刻返回，不等待 GPU 完成。如果 Kernel 在执行期发生错误（如越界），那一刻 CPU 正在跑别的代码，没有任何 API 调用在执行，错误只能先记录在运行时的「错误状态」里。直到 CPU 下一次触碰这个状态——cudaDeviceSynchronize()、cudaMemcpy 或其他返回 cudaError_t 的 API——错误才可能被取回。所以不是「错误等同步才发生」，而是「错误发生了，但要等检查点才被报告」。注意「可能被取回」不是「任意 API 都可靠」：不同 API 的同步语义各不相同，需要确认执行完成时，用 cudaDeviceSynchronize() 这类明确等待的 API，而不是碰运气。

cudaGetLastError() 和 cudaPeekAtLastError() 有什么区别？

两者都查询当前 host 线程关联的 CUDA Runtime error state，区别在于「查询后是否清除」：

在 Kernel launch 后，常见的检查方式是 cudaGetLastError()——它查询并清除当前 Host 线程的最后错误（last-error）状态。将其紧跟 Kernel launch 使用，可以减少之前遗留错误状态对当前检查的干扰，因此常被用作 launch 后的错误检查点。如果某个流程需要保留错误状态供后续逻辑继续观察，则用 cudaPeekAtLastError()。两者不是「二选一的正确 API」，而是「清除或不清除」的两种取法。

需要特别强调两点：第一，「清除」指的是 Runtime 的 last-error 状态，不等于撤销已经发生的 GPU 错误，也不意味着 GPU 工作被取消。 把 cudaGetLastError() 理解成「把 GPU 错误恢复掉了」是常见误区——它只是把 Host 侧的错误状态归位，GPU 上已经发生的错误不会因此消失。第二，这两个 API 都不会主动等待异步 Kernel 完成，因此都不能替代 cudaDeviceSynchronize() 作为执行完成检查。

把错误状态的完整生命周期画出来，比任何文字都直观：

左半部分是 launch 阶段错误，右半部分是执行期异步错误——两条路径最终都在「查询错误状态」的检查点被 Host 观察到。

为什么每个返回 cudaError_t 的 CUDA Runtime API 都要检查返回值？

因为错误只记录、不主动弹出。任何一次调用失败，如果不检查，程序会带着坏状态继续跑：cudaMalloc 失败后，后续 cudaMemcpy 可能收到空指针；cudaMemcpy 失败后，结果缓冲区里是垃圾数据。错误在发生时刻不报告，就会在更晚、更隐蔽的位置以更糟的方式暴露——检查返回值是把「迟到的问题」拉回「发生的位置」。

⑨ 延伸阅读

收藏速查表

口诀：分配检查，搬运检查，Launch 检查，同步检查，释放检查——没有一行调用可以裸奔。

延伸阅读

• 前文：向量加法数据闭环（Host↔Device 数据生命周期）；
• 前文：Kernel 生命周期与异步发射；
• 下一篇：一维线程索引——全局坐标公式的工程化；
• 后续：CUDA Stream 与异步拷贝——cudaMemcpyAsync + pinned host memory + Event，建立真正的 Host↔Device overlap；
• 后续：CUDA 异步内存管理——cudaMallocAsync + 内存池（Memory Pool） + Stream-Ordered Allocator；
• 后续：compute-sanitizer 与调试工作流；
• 后续：CUDA Event 计时方法论。

💡 本篇真正需要记住的不是某个 API，而是这条原则：

每个返回 cudaError_t 的 CUDA Runtime API 都可能失败，且失败不会自动报告——主动检查是唯一防线。

核心口诀：API 返回值要查，Kernel Launch 要查，执行完成要查，GPU 内存行为还要用 Sanitizer 查。

技术演进路线（为后续专题定位）

本篇
├── cudaMalloc / cudaMemcpy / cudaFree
├── CUDA_CHECK / cudaGetLastError / cudaDeviceSynchronize
└── Compute Sanitizer
↓
后续
├── CUDA_LAUNCH_BLOCKING / CUDA Event / Stream / 异步拷贝
├── cudaMallocAsync / cudaFreeAsync / Memory Pool
├── Unified Memory / 多 GPU / GPUDirect
└── Nsight Systems / Nsight Compute

下一篇：一维线程索引——把 blockIdx.x * blockDim.x + threadIdx.x 从公式变成直觉。

你踩过最隐蔽的 CUDA 内存坑是哪一个？是忘了检查返回值，还是错误被「吞掉」后找了两天？欢迎在评论区聊聊你的经历。