3、CUDA编程模型入门

好，咱们今天聊聊CUDA编程模型。说实话，我第一次接触CUDA的时候，也被它那套Host/Device、Kernel、Grid/Block/Thread的概念搞得有点晕。但别担心，这些东西说白了就是一套规则——告诉GPU怎么干活、谁来干活、活怎么分。

我个人习惯把CUDA编程模型理解成「老板-工人」模式。CPU是老板，负责发号施令；GPU是工人，负责埋头苦干。老板把任务拆成小块，扔给工人去并行处理。嗯，就这么简单。

3.1 Host与Device：老板和工人的分工

在CUDA里，Host指的是CPU及其内存系统，Device指的是GPU及其显存。这两者各有各的职责：

• Host（CPU端）：负责逻辑控制、串行任务、数据准备、启动Kernel。
• Device（GPU端）：负责大规模并行计算、数据吞吐。

我在项目中遇到过一个问题：有人把大量数据频繁地在Host和Device之间来回拷贝，结果性能还不如纯CPU跑。你想想看，PCIe带宽再快，也架不住你来回折腾啊。所以记住一个原则：数据尽量留在Device上，减少Host-Device传输。

3.2 核函数（Kernel）：给GPU下的指令

Kernel是什么？说白了，就是一段在GPU上执行的函数。你把它定义好，然后告诉CUDA「去，把这个函数在GPU上跑N遍」。每个线程执行同一个Kernel，但处理不同的数据。

定义Kernel的语法很简单，用__global__修饰符就行：

// 定义一个Kernel函数
__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];
    }
}

调用Kernel的方式也很特别，用三括号语法<<<grid, block>>>：

// 调用Kernel
int N = 1024;
vecAdd<<<1, 256>>>(d_A, d_B, d_C, N);

这里<<<1, 256>>>的意思是：启动1个Block，每个Block里有256个线程。总共256个线程并行执行vecAdd函数。

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

这是CUDA最核心的概念。你想想看，GPU里有几千个核心，怎么管理这么多线程？CUDA的做法是分层管理：

• Thread（线程）：最基础的执行单元，每个线程执行Kernel的一份拷贝。
• Block（线程块）：一组线程的集合，Block内的线程可以协作（共享内存、同步）。
• Grid（线程网格）：一组Block的集合，一个Kernel对应一个Grid。

用一张图来理解：

每个线程都有一个唯一的ID，通过threadIdx、blockIdx、blockDim这三个内置变量来定位：

变量	含义	示例值
threadIdx	线程在Block内的索引	threadIdx.x = 0, 1, ..., 255
blockIdx	Block在Grid内的索引	blockIdx.x = 0, 1, ..., 15
blockDim	每个Block包含的线程数	blockDim.x = 256
gridDim	每个Grid包含的Block数	gridDim.x = 16

计算全局线程ID的公式：int tid = threadIdx.x + blockIdx.x * blockDim.x;

3.4 一个完整的例子：向量加法

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

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

// Kernel定义
__global__ void vecAddKernel(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 = 1024;
    size_t size = N * sizeof(float);
    
    // 1. 在Host上分配内存
    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. 在Device上分配显存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);
    
    // 3. 把数据从Host拷贝到Device
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
    
    // 4. 配置Kernel启动参数
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    
    // 5. 启动Kernel
    vecAddKernel<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    
    // 6. 把结果从Device拷贝回Host
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
    
    // 7. 验证结果
    for (int i = 0; i < 10; i++) {
        printf("C[%d] = %f\n", i, h_C[i]);
    }
    
    // 8. 清理
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    free(h_A);
    free(h_B);
    free(h_C);
    
    return 0;
}

这个流程你记住：分配内存 → 拷贝数据 → 启动Kernel → 拷回结果 → 清理。我刚开始学的时候，老是忘记cudaFree，结果显存泄漏，跑几次程序就卡死了。嗯，血的教训。

3.5 线程层次结构的实际意义

你可能会问：为什么非要搞Grid、Block、Thread三层？直接搞一堆线程不行吗？

原因有两点：

1. 硬件限制：GPU的SM（流多处理器）是以Block为单位调度任务的。一个SM可以同时处理多个Block，但Block内的线程必须在一个SM上执行。所以Block大小直接影响资源利用率。
2. 协作需求：Block内的线程可以通过共享内存（Shared Memory）通信，也可以同步（__syncthreads()）。但不同Block之间的线程无法直接通信。这种设计让并行更高效。

好了，CUDA编程模型的核心概念就这些。说白了就是：CPU发号施令，GPU埋头干活；干活的时候，Grid分给Block，Block分给Thread，各司其职。把这个模型刻在脑子里，后面学内存管理、优化技巧就顺了。

---

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