4、CUDA编程基础（上）：内存模型（Global/Shared/Local）、核函数编写与调用、<<<>>>语法

好，咱们今天正式进入CUDA编程的核心环节。说实话，很多初学者一上来就被各种内存搞晕了——Global、Shared、Local，还有那个长得像外星符号的 <<<>>> 语法。别急，我当年也是这么过来的。咱们一个一个拆开讲，保证你听完就能上手写代码。

4.1 内存模型：Global / Shared / Local

先说说CUDA的内存模型。你想想看，GPU里那么多线程，它们怎么访问数据？总不能像CPU那样搞个大缓存吧？实际上，CUDA把内存分成了好几个层次，每个层次的速度、容量、作用都不一样。

4.1.1 Global Memory（全局内存）

这是最常用的内存。说白了，就是GPU板子上的显存。所有线程都能读写它，容量最大（几个GB到几十个GB），但速度最慢。我在项目中遇到过一个问题：一开始把数据全放在Global Memory里，结果跑起来发现带宽根本不够用，整个程序被内存访问拖慢了10倍。

关键点：Global Memory的访问延迟大约是几百个时钟周期。如果你不优化，程序大部分时间都在等数据。

Global Memory的生命周期由主机端（CPU）控制。你用 cudaMalloc 分配，用 cudaMemcpy 传输，最后用 cudaFree 释放。嗯，这里要注意：千万别忘了释放内存，否则显存泄漏了，你只能重启机器。

4.1.2 Shared Memory（共享内存）

Shared Memory是块上的内存，同一个Block里的所有线程都能访问。它的速度比Global Memory快得多——大概快100倍左右。但容量很小，每个Block通常只有几十KB。

我个人的习惯是：把频繁访问的、需要线程间共享的数据放到Shared Memory里。比如做矩阵乘法时，把子矩阵先加载到Shared Memory，然后让线程从Shared Memory里读数据，而不是反复去Global Memory里取。

小技巧：Shared Memory的访问速度很快，但要注意Bank Conflict。简单说，就是多个线程同时访问同一个Bank时，会被串行化。我曾经因为这个坑，性能直接腰斩。解决办法？调整数据布局，让线程访问不同的Bank。

4.1.3 Local Memory（局部内存）

Local Memory这个名字容易让人误解。它其实不是“本地”的，而是存在Global Memory里的，只不过每个线程私有的。当寄存器不够用时，编译器会自动把变量溢出到Local Memory里。

为什么会这样？因为每个线程的寄存器数量有限（比如Kepler架构每个线程最多63个寄存器）。如果你声明了太多局部变量，寄存器放不下，编译器就会把它们放到Local Memory里。但Local Memory的访问速度跟Global Memory一样慢，所以尽量避免使用过多的局部变量。

内存类型 作用域 生命周期 速度 容量
Global Memory 所有线程 由主机控制 慢（~400 cycles） 大（GB级）
Shared Memory 同一个Block Block生命周期 快（~5 cycles） 小（KB级）
Local Memory 单个线程 线程生命周期 慢（~400 cycles） 有限

4.2 核函数编写与调用

核函数（Kernel）就是跑在GPU上的函数。它的写法跟C函数很像，但有几个关键区别。

4.2.1 核函数的定义

核函数必须用 __global__ 修饰符。返回值必须是 void。举个例子：

__global__ void vecAdd(float* A, float* B, float* C, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

你看，这个函数跟普通C函数差不多，但多了 __global__ 关键字。函数体里用 threadIdx、blockIdx、blockDim 这些内置变量来获取线程的索引。

注意：核函数不能调用C标准库函数（比如 printf 在计算能力2.0以上才支持），也不能使用静态变量。我曾经试图在核函数里用 malloc，结果编译通过了，但运行时直接崩溃。嗯，血的教训。

4.2.2 核函数的调用：<<<>>> 语法

调用核函数时，要用 <<>> 语法来指定线程配置。这个语法看起来有点怪，但用多了就习惯了。

int N = 1024;
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

这里 blocksPerGrid 是网格中的Block数量，threadsPerBlock 是每个Block中的线程数。我建议你把Block大小设为32的倍数，因为GPU的warp大小就是32。如果你设成33，最后一个warp里只有1个线程在工作，其他31个都闲着，多浪费啊。

4.2.3 线程索引计算

在核函数里，你需要自己计算全局线程ID。常用的公式是：

int tid = threadIdx.x + blockIdx.x * blockDim.x;

如果是二维网格，那就更复杂一点：

int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
int idx = y * width + x;  // 假设是行优先存储

我个人习惯把线程索引计算放在核函数的第一行，这样后面用起来方便。而且一定要做边界检查，因为数据量不一定正好是Block大小的整数倍。

4.3 实战：向量加法

光说不练假把式。咱们写一个完整的向量加法程序，把刚才讲的知识点串起来。

#include <stdio.h>
#include <cuda_runtime.h>

__global__ void vecAdd(float* A, float* B, float* C, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

int main() {
    int N = 1 << 20;  // 1M个元素
    size_t size = N * sizeof(float);

    // 分配主机内存
    float *h_A = (float*)malloc(size);
    float *h_B = (float*)malloc(size);
    float *h_C = (float*)malloc(size);

    // 初始化数据
    for (int i = 0; i < N; i++) {
        h_A[i] = i * 1.0f;
        h_B[i] = i * 2.0f;
    }

    // 分配设备内存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);

    // 拷贝数据到设备
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    // 配置线程
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

    // 调用核函数
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

    // 拷贝结果回主机
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    // 验证结果
    for (int i = 0; i < 10; i++) {
        printf("C[%d] = %f\n", i, h_C[i]);
    }

    // 释放内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    free(h_A);
    free(h_B);
    free(h_C);

    return 0;
}

这个程序很简单，但包含了CUDA编程的完整流程：分配内存、拷贝数据、配置线程、调用核函数、取回结果、释放内存。你想想看，如果把这个流程记熟了，后面学更复杂的操作就轻松多了。

调试建议：每次调用CUDA API后，最好检查一下返回值。可以用 cudaGetLastError() 来捕获错误。我曾经花了一整天找bug，最后发现是 cudaMalloc 返回了 cudaErrorMemoryAllocation，因为显存不够了。嗯，从那以后我每次分配内存都加检查。

4.4 避坑指南

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

  • 忘记同步：核函数是异步的，调用后CPU会继续执行。如果你在核函数返回后立即读取结果，可能会读到旧数据。记得用 cudaDeviceSynchronize() 等待GPU完成。
  • Block大小不合理：每个Block的线程数不能超过1024（计算能力2.0以上），而且最好是32的倍数。我见过有人设成1000，结果最后一个warp浪费了24个线程。
  • 内存越界：核函数里访问数组时，一定要做边界检查。否则你可能会读到垃圾数据，或者更糟——导致程序崩溃。
  • 忘记释放内存：这个前面说过了，但还是要强调。显存泄漏比内存泄漏更难排查，因为GPU不会像CPU那样给你报错。

好了，这一章的内容就到这里。下一章我们会深入讲Shared Memory的优化技巧，以及如何用CUDA事件来测量性能。到时候我会分享一个我实际项目中用到的优化案例，保证让你大开眼界。

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