2. CUDA编程模型入门

说实话，很多初学者第一次接触CUDA时，都会被一堆概念搞懵。什么Host、Device、Grid、Block、Thread……我当年也一样。但别急，咱们一步步来拆解。

1.1 Host与Device：CPU和GPU的分工

在CUDA的世界里，Host指的是CPU及其内存，Device指的是GPU及其显存。说白了，Host是发号施令的老板，Device是干活的工人。

我习惯把Host比作项目经理，Device比作施工队。项目经理（CPU）负责逻辑控制、串行任务，比如读取数据、分配任务、检查结果。施工队（GPU）负责大规模并行计算，比如矩阵乘法、图像处理。

举个例子，你在CPU上有一个数组，想用GPU加速计算，流程是这样的：

1. 在GPU上分配显存
2. 把数据从CPU拷贝到GPU
3. 在GPU上执行Kernel函数
4. 把结果从GPU拷贝回CPU
5. 释放GPU显存

嗯，这里要注意：数据拷贝是开销很大的操作。我在项目中遇到过，有人为了图省事，每次计算都来回拷贝数据，结果性能还不如纯CPU。所以，尽量减少Host和Device之间的数据传输。

1.2 Kernel函数定义与启动：让GPU干活

Kernel函数，就是运行在GPU上的函数。它的定义方式很特别，需要在函数前面加上 __global__ 修饰符。

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

你可能会问：为什么用 __global__？因为CUDA有三种函数修饰符：

修饰符	调用位置	执行位置
__global__	Host端调用	Device端执行
__device__	Device端调用	Device端执行
__host__	Host端调用	Host端执行

启动Kernel的语法也很特别，用三括号 <<<grid, block>>> 来指定线程层次结构：

// 启动Kernel
int n = 1000000;
int threadsPerBlock = 256;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;

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

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

这是CUDA最核心的概念。你想想看，GPU有成千上万个核心，怎么组织它们？CUDA用了三层结构：

• Thread（线程）：最小的执行单元，每个线程执行相同的Kernel代码，但处理不同的数据
• Block（线程块）：一组线程的集合，Block内的线程可以共享内存和同步
• Grid（网格）：一组Block的集合，一个Kernel对应一个Grid

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

• threadIdx.x：线程在Block内的索引
• blockIdx.x：Block在Grid内的索引
• blockDim.x：每个Block包含的线程数
• gridDim.x：每个Grid包含的Block数

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

我习惯把Grid想象成一个城市，Block是小区，Thread是住户。每个住户（Thread）有自己门牌号（threadIdx），每个小区（Block）有小区编号（blockIdx），每个小区有固定户数（blockDim）。

1.4 第一个向量加法Kernel：完整示例

好了，理论说完了，咱们写一个完整的向量加法程序。这是CUDA的"Hello World"。

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

// Kernel定义
__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);
    
    // 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. 拷贝数据到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;
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, n);
    
    // 5. 拷贝结果回Host
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
    
    // 6. 验证结果
    for (int i = 0; i < 10; i++) {
        printf("C[%d] = %f\n", i, h_C[i]);
    }
    
    // 7. 清理
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    free(h_A);
    free(h_B);
    free(h_C);
    
    return 0;
}

1.5 知识体系结构图

下面这张图总结了本章的核心内容，帮你理清思路：

这张图展示了CUDA编程模型的完整流程。从Host发起，到Device执行，再到线程层次的组织方式。你想想看，是不是一下子就清晰了？

1.6 本章小结

咱们这一章讲了CUDA编程模型的基础：

• Host与Device：CPU和GPU的分工协作，数据需要显式拷贝
• Kernel函数：用 __global__ 定义，用 <<<grid, block>>> 启动
• 线程层次：Grid→Block→Thread，每个线程有唯一ID
• 向量加法：完整的代码示例，从分配到清理

我个人觉得，向量加法虽然简单，但它包含了CUDA编程的所有核心要素。把这个搞懂了，后面的矩阵乘法、卷积优化什么的，都是在这个基础上加东西。

好了，这一章就到这里。记住：多写代码，多调试，多思考。CUDA没那么神秘，就是个工具，用多了就熟了。