第一章：GPU计算基础与CUDA入门

各位同学好，我是你们这门课的主讲。今天咱们聊聊GPU计算的基础，还有CUDA怎么入门。

说实话，我第一次接触GPU编程的时候，也是一头雾水。那时候我还在做图像处理，CPU算得实在太慢了，老板催得紧。后来有人跟我说：「试试CUDA吧。」结果一试，真香。

1.1 GPU架构概览：它到底强在哪？

先说说GPU和CPU的区别。CPU像是一个全能型选手，什么活都能干，但一次只能干几件事。GPU呢？它更像是一群流水线上的工人，每个人只会干一件事，但人多力量大。

我给大家打个比方。CPU就像是一个米其林大厨，能做满汉全席，但一次只能做一桌。GPU呢，就像是一百个麦当劳员工，每个人只会做汉堡，但一分钟能做出几百个。

为什么会这样？因为GPU的设计哲学就是「以量取胜」。它把大量的晶体管用来做计算单元，而不是做复杂的控制逻辑和缓存。你想想看，一个GPU里有几千个核心，每个核心虽然跑得不如CPU快，但几千个一起上，那计算能力就恐怖了。

具体到英伟达的GPU架构，我简单说几个关键点：

• SM（流式多处理器）：这是GPU的基本计算单元。每个SM里包含多个CUDA核心、共享内存、寄存器等资源。
• 全局内存：所有线程都能访问，但速度慢。我习惯叫它「显存」。
• 共享内存：同一个SM内的线程可以共享，速度很快。我在项目中经常用它来做数据缓存，能大幅提升性能。
• 寄存器：每个线程私有的，速度最快，但数量有限。

1.2 CUDA编程模型：怎么指挥这群工人？

CUDA是英伟达推出的并行计算平台和编程模型。说白了，它就是一套让你能用C/C++语言指挥GPU干活的工具。

我个人觉得，理解CUDA编程模型，最关键的是搞懂三个概念：

• 线程层次结构：线程、线程块、网格
• 内存层次结构：全局内存、共享内存、寄存器
• 主机与设备：CPU是主机，GPU是设备

先说说线程层次。在CUDA里，你写的函数叫「核函数」，它会在GPU上被成千上万个线程同时执行。这些线程是怎么组织的呢？

• 一个线程是最小的执行单元
• 多个线程组成一个线程块（Block）
• 多个线程块组成一个网格（Grid）

嗯，这里要注意：同一个线程块内的线程可以协作，通过共享内存交换数据。不同线程块之间的线程，基本上老死不相往来。

我在做项目时，经常需要权衡线程块的大小。太小了，SM的资源用不满；太大了，又可能超出硬件限制。一般来说，线程块的大小取32的倍数比较好，因为GPU的warp（线程束）大小就是32。

1.3 第一个CUDA程序：Hello, World!

光说不练假把式。咱们直接上手写一个CUDA程序。这个程序很简单：在GPU上计算两个向量的加法。

#include <stdio.h>

// 核函数：在GPU上执行
__global__ void vectorAdd(const float *a, const float *b, float *c, int n) {
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

int main() {
    int n = 1024;
    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;
    vectorAdd<<<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]);
    }

    // 释放内存
    free(h_a); free(h_b); free(h_c);
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);

    return 0;
}

这段代码看着挺长，其实核心就几件事：

1. 写一个核函数（前面加__global__）
2. 在GPU上分配内存（cudaMalloc）
3. 把数据拷到GPU（cudaMemcpy）
4. 启动核函数（<<<grid, block>>>）
5. 把结果拷回来
6. 释放内存

1.4 避坑指南：新手最容易犯的错

我曾经在项目里犯过一个低级错误：核函数里忘记加__global__修饰符。结果程序编译通过了，但运行起来什么也没发生。查了半天才发现，核函数根本没在GPU上执行，而是在CPU上跑了一遍。

还有一次，我分配了设备内存，但忘了释放。程序跑了几百次后，显存爆了，整个系统都卡死了。嗯，从那以后，我每次写cudaMalloc，都会立刻写上对应的cudaFree。

1.5 总结与思考

好了，第一章的内容就到这里。咱们回顾一下：

• GPU架构的核心是「以量取胜」，用大量简单计算单元并行工作
• CUDA编程模型围绕线程层次和内存层次展开
• 写CUDA程序就是：分配内存、拷贝数据、启动核函数、拷回结果

最后留个思考题：为什么线程块的大小取32的倍数比较好？这个问题咱们下一章会详细讲。如果你等不及，可以先查查「warp」这个概念。

下一章，咱们聊聊CUDA的内存管理，这可是性能优化的关键。到时候我会分享一些我在实际项目中踩过的坑，保证让你少走弯路。