CUDA编程模型：从Host到Device的旅程

大家好，我是你们的讲师。今天我们来聊聊CUDA编程模型的核心概念。说实话，这部分内容是我当年入门CUDA时最头疼的，但也是理解GPU计算的关键。咱们一步步来，别急。

Host与Device：CPU和GPU的分工

先问个问题：你的电脑里，CPU和GPU谁说了算？

答案是CPU。在CUDA编程模型里，CPU被称为Host（主机），GPU被称为Device（设备）。Host负责控制流程、逻辑判断、I/O操作，而Device负责大规模并行计算。

我刚开始做项目时，总想把所有计算都扔给GPU。结果发现，数据传输的开销比计算本身还大。嗯，这里要注意：Host和Device之间的数据传输是昂贵的。

举个例子，你在CPU上分配内存，把数据拷贝到GPU，让GPU计算，再把结果拷回来。这个过程就像：你（Host）把材料送到工厂（Device），工厂加工完，你再把成品取回来。来回跑路的时间，有时候比加工时间还长。

Kernel函数：GPU上的“小任务”

Kernel函数，说白了就是在GPU上执行的函数。它的定义方式很特别——前面要加一个__global__修饰符。

__global__ void myKernel(float *data) {
    // 这里的代码会在GPU上执行
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    data[idx] = data[idx] * 2.0f;
}

调用方式也跟普通函数不一样：

myKernel<<<gridDim, blockDim>>>(d_data);

看到那对尖括号了吗？那是CUDA特有的语法，用来指定线程的组织结构。我刚开始写的时候，经常忘记加这玩意儿，结果函数在CPU上执行了，GPU纹丝不动。你说尴尬不尴尬？

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

这是CUDA最核心的概念，也是很多人绕不过去的坎。咱们用一张图来理解：

你看，一个Grid包含多个Block，每个Block包含多个Thread。这就像：

• Grid = 整个工厂
• Block = 车间
• Thread = 工人

每个工人（Thread）干自己的活，车间（Block）里的工人可以互相交流（共享内存），不同车间的工人不能直接交流。

内置变量：找到你的位置

CUDA提供了几个内置变量，帮你确定当前线程在全局中的位置：

变量名	含义	取值范围
threadIdx	当前线程在Block内的索引	0 到 blockDim - 1
blockIdx	当前Block在Grid内的索引	0 到 gridDim - 1
blockDim	Block的维度（包含多少个Thread）	由用户指定
gridDim	Grid的维度（包含多少个Block）	由用户指定

这些变量可以是1维、2维或3维的。比如threadIdx.x、threadIdx.y、threadIdx.z。

计算全局索引的公式很简单：

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

这个公式我闭着眼睛都能写出来。你想想看，每个Block有blockDim.x个线程，当前是第blockIdx.x个Block，所以偏移量就是blockIdx.x * blockDim.x，再加上线程在Block内的位置threadIdx.x，就是全局索引了。

向量加法实战：Hello World of CUDA

好了，理论说完了，咱们来写个真正的CUDA程序。向量加法就是CUDA界的"Hello World"。

先看完整代码：

#include <cuda_runtime.h>
#include <stdio.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 = 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. 把数据从Host拷贝到Device
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
    
    // 4. 配置线程层次
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    
    // 5. 启动Kernel
    vecAddKernel<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    
    // 6. 把结果拷回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;
}

这段代码有8个步骤，我一个个说：

1. Host分配内存：用malloc在CPU上分配空间
2. Device分配内存：用cudaMalloc在GPU上分配空间
3. 数据拷贝：把数据从CPU搬到GPU
4. 配置线程：决定用多少Block、每个Block多少Thread
5. 启动Kernel：GPU开始干活
6. 结果拷回：把计算结果从GPU搬回CPU
7. 验证：检查结果对不对
8. 清理：释放内存

关于线程配置，有个小技巧：

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

这个公式保证即使N不能被threadsPerBlock整除，也能覆盖所有数据。比如N=1000，threadsPerBlock=256，那么blocksPerGrid=4。最后一个Block只有232个线程干活，剩下的24个线程闲着。但总比漏掉数据强。

Kernel函数里的if (i < N)检查也很重要。为什么？因为多余的线程会访问越界内存。我刚开始写的时候没加这个判断，结果程序偶尔崩溃，偶尔正常，查了好久才发现是越界访问的问题。

总结一下

今天咱们聊了CUDA编程模型的几个核心概念：

• Host与Device：CPU指挥，GPU干活
• Kernel函数：用__global__修饰，用<<<grid, block>>>调用
• 线程层次：Grid → Block → Thread，三层结构
• 内置变量：threadIdx、blockIdx、blockDim、gridDim
• 向量加法：完整的CUDA程序流程

这些东西看着简单，但真正用好需要大量实践。我建议你亲手敲一遍代码，跑一跑，看看结果。遇到问题别怕，GPU编程就是这样，踩坑多了自然就熟了。

想要更多资料，联系可可