第一章：CUDA编程模型与GPU架构基础

大家好，我是你们的老朋友。今天咱们聊聊CUDA编程最核心的东西——GPU架构和编程模型。说实话，我刚开始接触CUDA时，也被那一堆术语搞得头晕：SM、Warp、Thread、Grid、Block……但别担心，我会用最直白的方式讲清楚。

1.1 GPU为什么这么快？

先问个问题：为什么GPU能比CPU快那么多？

CPU像是一个全能型选手，什么活都能干，但一次只能干几件事。GPU呢？它更像是一支军队，成千上万个士兵同时干活。你想想看，处理一张图片的每个像素，CPU要一个个来，GPU可以同时处理几万个像素——这就是并行计算的力量。

我个人习惯把GPU架构理解成三个层次：

SM（Streaming Multiprocessor）——流式多处理器，相当于GPU里的“计算单元”。每个SM有自己的寄存器、共享内存和调度器。
Warp——线程束，32个线程组成一个warp。这是GPU执行指令的最小单位。记住这个数字：32。
Thread——线程，最基础的执行单元。

我在项目中遇到过一个问题：写了个核函数，性能死活上不去。后来发现是warp发散导致的——同一个warp里的线程走了不同的分支。嗯，这里要注意，warp里的线程最好走同一条路。

1.2 CUDA编程模型：Grid、Block、Thread

CUDA的编程模型其实很直观。你把任务想象成一个三维网格：

Grid——整个任务，由多个Block组成
Block——任务的一个子块，由多个Thread组成
Thread——最细粒度的执行单元

为什么搞这么复杂？说白了，是为了灵活性和硬件适配。不同的GPU有不同数量的SM，每个SM能跑的block数量也有限制。通过调整block大小，你可以让程序在不同GPU上都跑得不错。

关键点：一个block里的线程可以共享数据（通过共享内存），也可以同步（通过__syncthreads()）。不同block之间不能直接通信。

1.3 第一个CUDA程序：向量加法

光说不练假把式。咱们写个最简单的向量加法程序。这个例子我用了不下百次，每次给新人讲都从它开始。

// 核函数：向量加法
__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);

    // 拷贝数据到GPU
    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;
}

小技巧：threadsPerBlock选256是个不错的起点。为什么？因为它是32的倍数（warp大小），而且大多数GPU都能同时跑多个block。我曾经试过128和512，256在大多数场景下表现最均衡。

1.4 避坑指南

我曾经犯过一个低级错误：忘记检查cudaMalloc的返回值。结果程序跑着跑着就崩了，查了半天才发现是显存不够。所以，我建议你养成好习惯：

每次调用CUDA API后检查返回值
用cudaGetLastError()捕获核函数错误
记得cudaFree，不然显存泄漏很麻烦

警告：核函数里不要用printf太多！GPU上printf是异步的，大量输出会严重影响性能。调试时用用就好，正式代码里删掉。

1.5 知识体系总览

下面这张图是我自己画的，把本章的核心知识点串起来了。你看一眼就能明白GPU架构和CUDA编程模型的关系：

这张图你看懂了吗？从上到下是硬件层次：GPU里有多个SM，每个SM管理多个warp，每个warp包含32个线程。编程时我们用的是Grid-Block-Thread这个逻辑模型，硬件会自动把block分配到SM上执行。

1.6 总结

好了，第一章就讲这么多。记住三个数字：32（warp大小）、256（推荐block线程数）、还有——记得检查错误。下一章咱们聊聊内存管理，那可是性能优化的重头戏。

想要更多资料，联系可可

📚 课程目录

第一章：CUDA编程模型与GPU架构基础

1.1 GPU为什么这么快？

1.2 CUDA编程模型：Grid、Block、Thread

1.3 第一个CUDA程序：向量加法

1.4 避坑指南

1.5 知识体系总览

1.6 总结