第一个CUDA程序：向量加法

说实话，我教了这么多年CUDA，向量加法这个例子，就像编程界的"Hello World"。但它比Hello World有意思多了——它让你真正感受到GPU并行计算的威力。

我记得第一次写这个程序时，心里还在嘀咕：这玩意儿真能比CPU快？结果一跑，好家伙，100万个元素加起来，CPU要跑好几毫秒，GPU眨眼就完事了。今天我们就来亲手写一个，把CUDA的骨架搭起来。

先看代码，再讲道理

别急着看理论，咱们先把代码跑起来。下面是一个完整的向量加法核函数：

// 核函数声明
__global__ void vecAdd(float *A, float *B, float *C, int n) {
    // 计算当前线程的全局索引
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 边界检查，防止越界
    if (i < n) {
        C[i] = A[i] + B[i];
    }
}

// 主函数调用
int main() {
    int n = 1000000;
    size_t size = n * sizeof(float);
    
    // 分配主机内存
    float *h_A = (float*)malloc(size);
    float *h_B = (float*)malloc(size);
    float *h_C = (float*)malloc(size);
    
    // 初始化数据...
    
    // 分配设备内存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);
    
    // 拷贝数据到GPU
    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);
    
    // 清理...
}

代码不长，但信息量很大。咱们拆开揉碎了讲。

核函数声明：__global__ 到底啥意思？

你看那个 __global__ 关键字，它告诉编译器：这个函数要在GPU上跑，而且能被CPU调用。CUDA有三种函数修饰符：

修饰符 调用方 执行位置 典型用途
__global__ CPU调用 GPU执行 核函数入口
__device__ GPU调用 GPU执行 辅助函数
__host__ CPU调用 CPU执行 普通函数（可省略）

我个人习惯，__global__ 函数返回类型必须是 void。为什么？因为GPU上成千上万个线程同时跑，你让谁去return？所以核函数只管干活，结果通过指针参数传回来。

语法糖：<<<grid, block>>> 的真相

这个尖括号语法，我第一次看到时觉得挺别扭。但用多了就发现，它其实很直观。

<<<blocksPerGrid, threadsPerBlock>>> 定义了线程的组织方式：

  • grid：一个网格，包含多个block
  • block：一个线程块，包含多个thread
  • thread：最小的执行单元

说白了，就是三层结构。GPU把任务拆成网格，网格里放线程块，线程块里放线程。每个线程执行相同的代码，但处理不同的数据。

关键理解：blockIdx.x 是当前线程所在块的编号，threadIdx.x 是线程在块内的编号，blockDim.x 是每个块的线程数。三者组合起来，就能算出全局唯一的线程ID。

画个图，一目了然

光说可能不够直观，我画了个图，帮你理解线程是怎么组织的：

Grid（网格） Block (0,0) Block (1,0) Block (2,0) T0 T1 T2 T3 T4 T5 ... blockDim.x = 256 gridDim.x = 3

你看，每个Block里有多个Thread，每个Grid里有多个Block。实际运行时，GPU会把Block分配到不同的SM（流多处理器）上执行。这就是CUDA的并行精髓。

线程索引计算：那个公式是怎么来的？

代码里有一行：int i = blockIdx.x * blockDim.x + threadIdx.x;

这个公式其实很简单。你想想看：

  • 每个Block有 blockDim.x 个线程
  • 当前是第 blockIdx.x 个Block
  • 所以前面已经有 blockIdx.x * blockDim.x 个线程处理过了
  • 再加上当前线程在Block内的偏移 threadIdx.x
  • 就得到了全局唯一的线程ID

举个例子：假设 blockDim.x = 256，blockIdx.x = 2，threadIdx.x = 10。那么 i = 2 * 256 + 10 = 522。这个线程就负责处理数组的第522个元素。

小技巧：我习惯把线程ID打印出来调试。用 printf 在核函数里输出 blockIdx.x、threadIdx.x 和计算出的 i，能帮你快速理解线程的映射关系。

边界检查：为什么一定要加？

你看代码里有个 if (i < n) 的判断。这个很重要。

为什么？因为 blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock 这个公式，算出来的Block数可能会多出一些线程。比如 n = 1000，threadsPerBlock = 256，那么 blocksPerGrid = 4。总线程数 = 4 * 256 = 1024，比1000多了24个。

这24个线程如果不加判断，就会访问到数组边界之外，轻则读到垃圾数据，重则段错误。

我曾经在项目里吃过这个亏。有一次写矩阵乘法，忘了加边界检查，结果跑出来的结果时对时错。查了两天才发现，原来是边界外的线程在捣乱。从那以后，我写核函数的第一件事就是加边界判断。

内存管理：主机和设备之间的搬运工

CUDA编程里，内存管理是绕不开的坎。你得记住：CPU和GPU有各自的内存空间，数据不能直接共享。

基本流程就三步：

  1. cudaMalloc：在GPU上分配显存
  2. cudaMemcpy：在主机和设备之间拷贝数据
  3. cudaFree：释放显存

这里有个坑：cudaMemcpy 的最后一个参数，指定拷贝方向。常见的有：

  • cudaMemcpyHostToDevice：从CPU到GPU
  • cudaMemcpyDeviceToHost：从GPU到CPU
  • cudaMemcpyDeviceToDevice：GPU内部拷贝

方向搞反了，数据就传不过去。我刚开始学的时候，经常忘了把结果拷回来，然后盯着CPU上的数组发呆——怎么全是0？

性能小贴士：Block大小怎么选？

代码里我用了 threadsPerBlock = 256。这个数字不是随便选的。

一般来说，Block大小建议是32的倍数（因为warp大小是32）。常用的值有128、256、512。太小了（比如32）会导致SM利用率低，太大了（比如1024）又会受寄存器数量限制。

我个人经验：对于大多数计算密集型任务，256是个不错的起点。然后你可以用profiler去调优，找到最适合你GPU型号的值。

完整代码示例

最后，给你一个可以直接编译运行的完整代码：

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

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

int main() {
    int n = 1000000;
    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);
    
    // 拷贝到GPU
    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;
}

编译命令：nvcc -o vecAdd vecAdd.cu

跑一下试试。如果看到输出 C[0] = 0.0, C[1] = 3.0, C[2] = 6.0...，恭喜你，第一个CUDA程序跑通了！

嗯，向量加法虽然简单，但它包含了CUDA编程的所有核心要素：核函数声明、线程组织、内存管理。把这个吃透了，后面的矩阵乘法、归约、卷积，都是在这个基础上加花样。

可可精品资料库

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