一、CUDA编程模型入门：从CPU到GPU的思维转变

说实话，我第一次接触CUDA的时候，脑子里全是CPU那套串行思维。写个for循环，一个线程跑完所有数据——这在GPU世界里，简直就是暴殄天物。今天我就带你从零开始，把CUDA这套编程模型掰开揉碎了讲清楚。

1.1 CUDA架构简介：GPU不是单纯的“多核CPU”

很多人以为GPU就是CPU的“加强版”——核心多嘛。但你要这么想，后面写代码会吃大亏。

我个人的理解是这样的：CPU是“特种兵”，单兵作战能力极强；GPU是“步兵方阵”，单个士兵很弱，但人海战术无敌。CUDA架构的精髓，就是让成千上万个“弱鸡”线程协同作战，完成海量数据的并行处理。

CUDA架构里，GPU被组织成一个个流多处理器（SM）。每个SM里有多个CUDA核心，共享一套控制单元和缓存。嗯，这里要注意：线程不是直接跑在CUDA核心上的，而是由SM调度执行。这个调度机制，我后面会细说。

1.2 主机与设备：CPU和GPU的分工协作

在CUDA编程模型里，CPU叫主机（Host），GPU叫设备（Device）。它们各干各的活，通过PCIe总线交换数据。

我刚开始写CUDA程序时，犯过一个低级错误：直接在GPU上访问CPU的内存。结果程序崩溃，查了半天才发现——主机和设备的内存是物理隔离的。

角色	职责	内存
主机（CPU）	控制逻辑、串行任务、I/O	系统内存（DDR）
设备（GPU）	大规模并行计算	显存（VRAM）

典型的CUDA程序流程是这样的：

1. 主机分配内存并初始化数据
2. 把数据从主机拷贝到设备
3. 在设备上启动核函数（并行计算）
4. 把结果从设备拷贝回主机
5. 释放内存

说白了，CPU是“项目经理”，GPU是“施工队”。项目经理把图纸（数据）交给施工队，施工队干完活再把成果交回来。

1.3 核函数（Kernel）：GPU上的“并行任务”

核函数，就是跑在GPU上的函数。用__global__修饰符声明，调用时用三括号语法<<<grid, block>>>指定线程配置。

来看个最简单的例子：

// 核函数定义
__global__ void vec_add(float *a, float *b, float *c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

// 主机端调用
int main() {
    // ... 数据准备和拷贝 ...
    vec_add<<<256, 256>>>(d_a, d_b, d_c, N);
    // ... 结果拷贝和清理 ...
}

这里有个坑，我当年踩过：核函数里不能调用printf（调试用除外），不能使用C++标准库，不能有递归。为什么？因为GPU的线程是轻量级的，没有完整的运行时环境支持这些操作。

1.4 线程层次结构：Grid、Block、Thread

这是CUDA编程模型里最核心的概念，也是初学者最容易懵的地方。

CUDA把线程组织成三层结构：

• Thread（线程）：最基础的执行单元
• Block（线程块）：一组线程，共享共享内存，可以同步
• Grid（线程网格）：一组线程块，构成整个核函数

我画了张图，帮你理解这个层次关系：

每个线程都有一个唯一的ID，通过内置变量获取：

• threadIdx.x：线程在块内的索引
• blockIdx.x：块在网格内的索引
• blockDim.x：每个块的线程数

你想想看，如果我们要处理一个长度为N的数组，每个线程处理一个元素，那么线程的全局ID就是：

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

这个公式，我建议你刻在脑子里。写CUDA代码，90%的时间都在和这个公式打交道。

1.5 实战：向量加法完整示例

光说不练假把式。我们写一个完整的向量加法程序，把上面讲的概念串起来：

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

// 核函数：向量加法
__global__ void vec_add(float *a, float *b, float *c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    int n = 1000000;
    size_t size = n * sizeof(float);
    
    // 1. 主机分配内存
    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;
    }
    
    // 2. 设备分配内存
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, size);
    cudaMalloc(&d_b, size);
    cudaMalloc(&d_c, size);
    
    // 3. 数据从主机拷贝到设备
    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);
    
    // 4. 配置线程并启动核函数
    int threads_per_block = 256;
    int blocks_per_grid = (n + threads_per_block - 1) / threads_per_block;
    vec_add<<<blocks_per_grid, threads_per_block>>>(d_a, d_b, d_c, n);
    
    // 5. 结果拷贝回主机
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
    
    // 验证结果
    for (int i = 0; i < 10; i++) {
        printf("c[%d] = %f\n", i, h_c[i]);
    }
    
    // 6. 清理
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);
    
    return 0;
}

这个例子虽然简单，但包含了CUDA编程的完整流程。我个人建议你亲手敲一遍，感受一下从CPU到GPU的数据流转。

好了，这一章的内容就到这里。CUDA编程模型的核心就是理解主机与设备的分工、核函数的定义与调用、以及线程的三层层次结构。把这些基础打牢，后面学各种并行模式就会轻松很多。

---

公众号:蓝海资料掘金营，微信deep3321