3. 内存管理：cudaMalloc、cudaMemcpy、cudaFree、统一内存入门
说到CUDA编程，内存管理是绕不开的第一道坎。我刚开始学的时候，总觉得这玩意儿不就是分配内存、拷贝数据、释放内存嘛，跟C语言里的malloc/memcpy/free差不多。但实际一上手，坑还真不少。
今天咱们就把这块彻底捋清楚。你想想看，GPU和CPU是两套独立的物理内存，数据不能直接共享。你要在GPU上算东西，得先把数据搬过去，算完再搬回来。这个搬运过程，就是内存管理的核心。
3.1 三大基本操作：cudaMalloc、cudaMemcpy、cudaFree
这三个函数，说白了就是CUDA版的malloc、memcpy和free。但细节上有很多不同，咱们一个一个说。
3.1.1 cudaMalloc —— 在GPU上分配内存
函数签名长这样：
cudaError_t cudaMalloc(void **devPtr, size_t size);注意第一个参数是二级指针。为什么？因为你要修改指针本身的值，让它指向GPU上的地址。我见过不少新手写成这样：
float *d_data;
cudaMalloc(d_data, N * sizeof(float)); // 错误！少了个&正确的写法是：
float *d_data;
cudaMalloc((void**)&d_data, N * sizeof(float));我个人习惯在变量命名上做区分：CPU上的指针加前缀h_（host），GPU上的加d_（device）。这样一眼就能看出来数据在哪儿。
3.1.2 cudaMemcpy —— 在CPU和GPU之间搬运数据
函数签名：
cudaError_t cudaMemcpy(void *dst, const void *src, size_t count, cudaMemcpyKind kind);最后一个参数指定方向，有四种取值：
| 枚举值 | 含义 |
|---|---|
| cudaMemcpyHostToDevice | 从CPU拷贝到GPU |
| cudaMemcpyDeviceToHost | 从GPU拷贝到CPU |
| cudaMemcpyDeviceToDevice | GPU内部拷贝 |
| cudaMemcpyHostToHost | CPU内部拷贝（基本不用） |
这里有个容易犯的错：方向搞反了。我见过有人把计算结果从CPU拷到GPU，结果算出来全是错的。嗯，方向错了，数据根本没回来。
一个完整的例子：
float *h_data = (float*)malloc(N * sizeof(float));
float *d_data;
cudaMalloc((void**)&d_data, N * sizeof(float));
// 初始化CPU数据
for (int i = 0; i < N; i++) h_data[i] = i * 1.0f;
// 拷到GPU
cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice);
// 在GPU上计算...（省略kernel调用）
// 拷回CPU
cudaMemcpy(h_data, d_data, N * sizeof(float), cudaMemcpyDeviceToHost);
// 清理
cudaFree(d_data);
free(h_data);3.1.3 cudaFree —— 释放GPU内存
这个最简单：
cudaError_t cudaFree(void *devPtr);但有个坑：释放之后，记得把指针置为NULL。不然你后面不小心又用它，就会访问非法地址。我习惯在free之后加一句：
cudaFree(d_data);
d_data = NULL;#define CUDA_CHECK(call) \
do { \
cudaError_t err = call; \
if (err != cudaSuccess) { \
fprintf(stderr, "CUDA error at %s:%d: %s\n", \
__FILE__, __LINE__, cudaGetErrorString(err)); \
exit(EXIT_FAILURE); \
} \
} while(0)3.2 统一内存（Unified Memory）入门
手动管理两套内存，说实话挺烦的。你想想看，每次都要malloc、cudaMalloc、memcpy、cudaMemcpy、free、cudaFree，代码里一半都是内存操作。统一内存就是来解决这个问题的。
统一内存的核心思想：让CPU和GPU共享同一个地址空间。你只需要分配一次内存，系统自动帮你处理数据迁移。
3.2.1 基本用法
用cudaMallocManaged代替cudaMalloc：
float *data;
cudaMallocManaged(&data, N * sizeof(float));
// 在CPU上初始化
for (int i = 0; i < N; i++) data[i] = i * 1.0f;
// 直接在GPU上使用，不需要显式拷贝
my_kernel<<<blocks, threads>>>(data, N);
// 在CPU上读取结果
for (int i = 0; i < N; i++) printf("%f\n", data[i]);
// 释放
cudaFree(data);你看，代码简洁多了。没有cudaMemcpy，没有方向参数，清爽。
3.2.2 工作原理
统一内存背后是CUDA驱动在管理页表。当CPU访问某个页面时，如果页面在GPU上，驱动会自动触发缺页中断，把页面迁回CPU。反过来也一样。
我用一张图来说明：
说白了，统一内存就是让CUDA驱动帮你做数据搬运的脏活累活。你只管用指针，剩下的交给系统。
3.2.3 什么时候用统一内存？
我个人经验是：
- 原型开发阶段：统一内存能让你快速验证算法，不用纠结内存管理细节
- 数据访问模式复杂：如果CPU和GPU频繁交替访问同一块数据，手动管理会很痛苦
- 代码可读性优先：统一内存让代码更简洁，更容易维护
但要注意，统一内存不是银弹。它有一些性能开销：
- 手动管理：你精确控制数据迁移时机，可以做到零额外开销
- 统一内存：驱动按需迁移，可能触发缺页中断，有延迟
- 建议：性能敏感的场景，还是用手动管理
3.3 避坑指南
我这些年踩过的坑，总结一下：
- 忘记同步：cudaMemcpy是同步操作，但kernel启动是异步的。如果你在kernel启动后立刻在CPU上读数据，读到的可能是旧值。记得加cudaDeviceSynchronize()。
- 内存泄漏：cudaMalloc了不cudaFree，显存会慢慢耗尽。我见过一个同事跑训练任务，显存占用越来越高，最后OOM了。查了半天，发现是循环里每次分配了没释放。
- 统一内存的页错误：在统一内存上频繁交替访问，会导致大量页迁移，性能反而比手动管理差。我曾经在一个项目里用统一内存，结果性能比手动管理慢了3倍。后来改成手动管理，速度就上来了。
- 指针作用域：cudaMallocManaged分配的指针，在CPU和GPU上都能用。但如果你在kernel里访问了还没迁移过来的页面，会触发隐式同步。这个同步可能发生在你意想不到的地方。
好了，内存管理这块就讲这么多。记住三个核心函数和一个统一内存的概念，后面写kernel的时候，你就能把精力放在算法本身，而不是被内存搬运折腾得焦头烂额。