第3章 CUDA编程模型入门

第3章：CUDA编程模型入门：CUDA平台概述、主机与设备、核函数（Kernel）概念

好，我们开始进入CUDA编程模型。说实话，这是整个GPU编程最核心的一块。很多初学者一上来就盯着语法看，结果写出来的代码跑得比CPU还慢。我当年也犯过这个错，后来才明白——理解编程模型，比记住几个API重要得多。

3.1 CUDA平台：不只是个编译器

CUDA全称是Compute Unified Device Architecture，统一计算设备架构。嗯，名字挺长，但你只需要记住一点：它不是一个单纯的编程语言，而是一个完整的软硬件协同平台。

我个人习惯把CUDA平台拆成三层来看：

硬件层：GPU本身，包含SM（流多处理器）、显存、缓存等
驱动层：CUDA Driver API，负责硬件交互
运行时层：CUDA Runtime API，我们平时用的最多的那一层

为什么这么分？因为我在项目中遇到过一个问题：直接用Driver API写代码，代码量翻了三倍，但性能提升不到5%。后来我学乖了，除非做底层驱动开发，否则老老实实用Runtime API。

关键认知：CUDA不是C语言的扩展，而是一个异构计算平台。你的代码会被分成两部分——一部分跑在CPU上，一部分跑在GPU上。

3.2 主机与设备：两个世界

在CUDA的世界里，有两个角色：

主机（Host）：CPU及其内存系统
设备（Device）：GPU及其显存系统

说白了，主机负责控制逻辑、I/O操作、串行任务。设备负责大规模并行计算。两者通过PCIe总线通信，但要注意——这个通信带宽远低于显存带宽。

我曾经帮一个团队优化代码，发现他们把数据在主机和设备之间来回拷贝了十几次。你想想看，每次拷贝都要经过PCIe，这能不慢吗？后来我们改成一次拷贝、多次计算，性能直接提升了8倍。

特性	主机（CPU）	设备（GPU）
内存	系统内存（DDR）	显存（GDDR/HBM）
计算单元	几个到几十个核心	数千个CUDA核心
适合任务	串行、逻辑控制	数据并行、计算密集
内存延迟	较低	较高（但带宽极高）

避坑指南：我曾经见过有人把malloc和cudaMalloc混用，结果程序崩溃了一整天。记住：主机内存用malloc，设备内存用cudaMalloc，别搞混了。

3.3 核函数（Kernel）：GPU上的执行单元

核函数是什么？简单说，就是跑在GPU上的函数。用__global__修饰符声明，调用时用<<<grid, block>>>语法指定执行配置。

来看一个最基础的例子：

// 核函数定义
__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() {
    // ... 分配内存、拷贝数据 ...
    
    // 执行配置：1个grid，256个线程/block
    vecAdd<<<1, 256>>>(d_A, d_B, d_C, N);
    
    // ... 拷贝结果回主机 ...
}

这里有几个关键点：

__global__：告诉编译器这个函数跑在设备上，但从主机调用
blockIdx.x：当前线程块在grid中的索引
threadIdx.x：当前线程在线程块中的索引
blockDim.x：每个线程块包含的线程数

为什么会这样设计？因为GPU需要把成千上万个线程组织起来。你想想看，如果每个线程都独立管理，调度器会疯掉的。所以CUDA采用了层次结构：Grid → Block → Thread。

个人经验：我刚开始写核函数时，总喜欢把if (i < N)这个边界检查去掉，觉得多余。结果有一次数据量不是256的整数倍，程序直接访问了非法内存。从那以后，我再也不敢省这个检查了。

3.4 执行配置：<<>>

这个语法看起来奇怪，但背后有深意。<<<grid, block>>>告诉运行时系统：

启动多少个线程块（gridDim）
每个线程块有多少个线程（blockDim）

举个例子：

// 处理1024个元素
int N = 1024;
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

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

这里用了一个经典技巧：(N + threadsPerBlock - 1) / threadsPerBlock。这是向上取整的整数除法，确保即使N不是256的倍数，也能覆盖所有元素。

重要原则：线程块的数量通常设置为SM数量的倍数。比如你的GPU有20个SM，那就设成40或80个block。这样能充分利用硬件资源。

3.5 一个完整的例子

把上面所有概念串起来，写一个完整的向量加法：

#include <cuda_runtime.h>
#include <stdio.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 = 1 << 20;  // 约100万个元素
    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);
    
    // 拷贝数据到设备
    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;
}

这个例子包含了CUDA编程的完整流程：分配内存 → 拷贝数据 → 执行核函数 → 拷贝结果 → 清理资源。嗯，每一步都不能少。

曾经踩过的坑：我刚开始写CUDA时，经常忘记调用cudaFree。结果程序跑着跑着显存就爆了。记住：显存泄漏比内存泄漏更难排查，因为GPU不会像CPU那样给你报段错误。

3.6 本章小结

这一章我们讲了三个核心概念：

CUDA平台：一个完整的软硬件协同系统，不是单纯的编程语言
主机与设备：CPU和GPU各司其职，通过PCIe通信
核函数：用__global__声明，用<<<grid, block>>>调用

下一章我们会深入线程层次结构，讲清楚Grid、Block、Thread到底是怎么映射到硬件上的。到时候你会发现，理解了这些，你就能写出真正高效的CUDA代码了。

📚 课程目录