4、CUDA内存管理：从cudaMalloc到统一内存

内存管理，说白了就是GPU编程的命门。

我见过太多新手，核函数写得飞起，结果一跑就崩。十有八九，问题出在内存上。今天咱们就把这块硬骨头啃下来。

4.1 传统内存管理三件套

cudaMalloc、cudaMemcpy、cudaFree，这三个API是CUDA编程的入门必修课。我习惯叫它们「铁三角」。

cudaMalloc：在GPU上分配显存。注意，它返回的是设备指针，不是CPU能直接访问的。

cudaMemcpy：在主机和设备之间搬运数据。你得指定方向——是H2D还是D2H。

cudaFree：释放显存。忘了它？等着OOM吧。

一个典型的数据流是这样的：

float *d_data;
cudaMalloc(&d_data, N * sizeof(float));   // GPU上分配
cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice); // 搬进去
// ... 跑核函数 ...
cudaMemcpy(h_result, d_data, N * sizeof(float), cudaMemcpyDeviceToHost); // 搬出来
cudaFree(d_data); // 释放

这里有个坑，我踩过好几次。cudaMalloc返回的是错误码，很多人不检查。一旦显存不够，程序直接崩。我的习惯是：每次调用都包一层check宏。

一个小技巧：cudaMallocPitch可以处理二维数组的对齐问题。如果你处理的是图像数据，用它性能更好。

4.2 统一内存：偷懒但高效

统一内存（Unified Memory）是CUDA 6.0引入的。说白了，就是让CPU和GPU共享一个地址空间。

你不需要手动搬数据了。用cudaMallocManaged分配一块内存，CPU和GPU都能直接访问。系统会自动帮你做页面迁移。

float *data;
cudaMallocManaged(&data, N * sizeof(float));
// CPU可以读写
for (int i = 0; i < N; i++) data[i] = i;
// GPU也能直接访问
kernel<<<grid, block>>>(data, N);
cudaDeviceSynchronize();
cudaFree(data);

听起来很完美对吧？但我要泼盆冷水。

统一内存有性能开销。页面迁移是隐式的，你控制不了它什么时候搬、搬多少。我遇到过这样的情况：一个简单的向量加法，用统一内存比手动管理慢了3倍。

注意：统一内存适合以下场景：

  • 代码原型快速验证
  • 数据访问模式不规则，难以手动优化
  • 内存占用小于GPU显存

不适合：高性能生产环境、超大显存需求、延迟敏感应用。

4.3 内存访问模式优化

这是真正拉开性能差距的地方。GPU的内存带宽很大，但如果你访问模式不对，带宽利用率可能不到10%。

核心原则：合并访问（Coalesced Access）

什么意思？相邻的线程应该访问相邻的内存地址。这样GPU可以把多次内存请求合并成一次，充分利用带宽。

我画个图帮你理解：

合并访问 vs 非合并访问 ✅ 合并访问（高效） 线程0 线程1 线程2 线程3 A[0] A[1] A[2] A[3] 一次内存事务搞定 ❌ 非合并访问（低效） 线程0 线程1 线程2 线程3 A[0] A[4] A[8] A[12] 需要4次内存事务 线程索引 = 内存索引 → 合并 线程索引 ≠ 内存索引 → 非合并

你看明白了吗？左边，线程0读A[0]，线程1读A[1]... 一次事务搞定。右边，线程0读A[0]，线程1读A[4]... 得四次事务，带宽浪费75%。

实战中的优化技巧

我总结了几条经验：

场景 优化方法 效果
二维数组按行访问 让threadIdx.x对应列索引 合并访问，带宽利用率高
二维数组按列访问 使用共享内存做转置 避免非合并，提升3-5倍
结构体数组 改为数组结构体（SoA） 连续访问，减少缓存miss
随机访问模式 用纹理内存或只读缓存 利用硬件缓存，降低延迟

我曾经优化过一个分子动力学模拟程序。原来用AoS（数组结构体），每个线程访问一个原子的位置、速度、力。改成SoA后，性能提升了4倍。说白了，就是让内存访问变得规整。

4.4 共享内存：程序员可控的缓存

共享内存是GPU上的片上存储，速度比全局内存快两个数量级。但容量有限——每个SM通常只有几十KB到上百KB。

它的典型用法是做数据复用。比如矩阵乘法，每个线程块把子矩阵加载到共享内存，然后反复使用，避免重复访问全局内存。

__global__ void matmul_shared(float *A, float *B, float *C, int N) {
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
    
    int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;
    int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;
    
    float sum = 0.0f;
    for (int t = 0; t < N / BLOCK_SIZE; t++) {
        // 协同加载到共享内存
        As[threadIdx.y][threadIdx.x] = A[row * N + t * BLOCK_SIZE + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(t * BLOCK_SIZE + threadIdx.y) * N + col];
        __syncthreads();
        
        // 计算
        for (int k = 0; k < BLOCK_SIZE; k++)
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        __syncthreads();
    }
    C[row * N + col] = sum;
}

注意那个__syncthreads()。它是个同步屏障，确保所有线程都加载完了再计算。少了它，数据可能还没到位，结果就错了。

共享内存的bank conflict是个隐藏的坑。如果多个线程访问同一个bank的不同地址，就会发生冲突，导致串行化。解决办法是加padding，或者调整访问模式。

4.5 内存池与异步操作

CUDA 11.2引入了cudaMemPool，可以预分配一大块显存，然后按需分配子块。好处是减少cudaMalloc的开销，避免内存碎片。

我建议在高频调用的场景下使用。比如推理服务，每次请求都要分配临时显存。用内存池能省掉malloc/free的延迟。

另外，cudaMemcpyAsync配合流（stream）可以实现数据传输和计算的重叠。说白了，就是一边搬数据一边算，把PCIe带宽用满。

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// 在stream1上搬数据
cudaMemcpyAsync(d_in1, h_in1, size, cudaMemcpyHostToDevice, stream1);
// 在stream2上搬另一批
cudaMemcpyAsync(d_in2, h_in2, size, cudaMemcpyHostToDevice, stream2);

// 在stream1上计算
kernel1<<<grid, block, 0, stream1>>>(d_in1, d_out1);
// 在stream2上计算
kernel2<<<grid, block, 0, stream2>>>(d_in2, d_out2);

这样两个流可以并行执行，吞吐量翻倍。我在做视频处理时经常用这招，效果立竿见影。

4.6 避坑指南

最后，分享几个我踩过的坑：

  • 忘记同步：cudaMemcpy是同步的，但cudaMemcpyAsync不是。我曾经在异步拷贝后直接读结果，读到的全是垃圾数据。
  • 过度使用统一内存：小数据量无所谓，大数据量性能惨不忍睹。我见过有人用统一内存处理10GB数据，结果页面迁移把系统拖垮了。
  • 忽略内存对齐：cudaMalloc默认128字节对齐，但如果你自己用指针偏移，可能破坏对齐。后果是访存性能下降。
  • 共享内存溢出：每个SM的共享内存有限。你声明了太大的共享内存数组，编译器不会报错，但运行时可能启动失败。

嗯，内存管理这块内容不少。但说白了，核心就三点：选对分配方式、优化访问模式、用好共享内存。把这三点吃透了，你的CUDA程序性能不会差。

想要更多资料，联系可可

← 上一章 📖 返回目录 下一章 →