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

大家好，我是你们的老朋友。今天咱们聊聊GPU编程的根基——硬件架构和编程模型。说实话，我见过太多人一上来就写核函数，结果性能惨不忍睹。为什么？因为你不了解GPU到底是怎么工作的。就像开车，你总得知道油门刹车在哪吧？

GPU硬件架构概览

先说说GPU的硬件结构。你想想看，一块GPU里面，最核心的单元叫SM（Streaming Multiprocessor，流式多处理器）。SM就是GPU的计算核心，类似于CPU里的核，但又不完全一样。

SM内部长什么样？

CUDA Core：每个SM里有几十到几百个CUDA核心。它们负责执行算术运算。
Warp调度器：负责把指令分发给CUDA核心。这个后面会细说。
共享内存（Shared Memory）：SM内部的高速缓存，所有线程都能访问。我在项目中经常用它来减少全局内存访问，效果立竿见影。
寄存器文件（Register File）：每个线程私有的存储空间。嗯，这里要注意，寄存器是有限的，线程太多会导致寄存器溢出，性能直接崩掉。

Warp是什么？

Warp是GPU执行指令的最小单位。一个Warp包含32个线程。这32个线程必须执行相同的指令，但可以处理不同的数据。这就是所谓的SIMT（单指令多线程）模型。

我曾经踩过一个坑：写了一个条件分支很多的核函数，结果Warp里的线程走了不同的分支路径，导致部分线程被阻塞，性能直接腰斩。所以，写核函数时尽量让Warp内的线程走同一条路。

内存层次结构

GPU的内存体系，说白了就是一个金字塔结构。从上到下，速度越来越慢，容量越来越大。

内存类型	位置	速度	容量	作用域
寄存器	SM内部	最快	极小（每个线程几十个）	线程私有
共享内存	SM内部	快	几十KB	Block内共享
L1缓存	SM内部	快	几十KB	SM内共享
L2缓存	芯片级	中等	几MB	全局共享
全局内存	显存	慢（几百个周期）	几GB到几十GB	所有线程

我个人的习惯是：能放寄存器的绝不放共享内存，能放共享内存的绝不放全局内存。这个原则帮我优化了不少算子。

核心要点：GPU性能优化的本质，就是尽量减少慢速内存的访问次数，把数据尽量往金字塔上层搬。

CUDA编程模型：Grid-Block-Thread层次结构

CUDA的编程模型，说白了就是怎么组织你的计算任务。它分三层：

Grid（网格）：一个核函数启动后，对应一个Grid。Grid包含多个Block。
Block（线程块）：Block是线程的集合。同一个Block内的线程可以共享共享内存，也可以同步。
Thread（线程）：最小的执行单元。每个线程执行核函数中的代码。

你可能会问：为什么要分这么多层？

原因很简单：为了灵活性和可扩展性。Grid对应整个问题域，Block对应子问题，Thread对应最小计算单元。这样，同一个程序可以在不同规格的GPU上运行，只要调整Grid和Block的维度就行。

线程索引怎么用？

CUDA提供了几个内置变量来定位线程：

threadIdx.x：线程在Block内的索引
blockIdx.x：Block在Grid内的索引
blockDim.x：Block的维度（包含多少个线程）
gridDim.x：Grid的维度（包含多少个Block）

计算全局线程ID的公式：int tid = blockIdx.x * blockDim.x + threadIdx.x;。这个公式我闭着眼睛都能写出来，因为每个核函数几乎都要用。

小技巧：Block的大小通常是32的倍数（因为Warp是32线程）。我一般用128或256，这两个值在大多数GPU上表现都不错。

Host与Device内存管理

Host是CPU端，Device是GPU端。两者有独立的内存空间。数据要来回拷贝，这是CUDA编程中最基础也最容易出错的地方。

常用内存操作函数：

cudaMalloc()：在Device上分配内存
cudaFree()：释放Device内存
cudaMemcpy()：在Host和Device之间拷贝数据
cudaMallocHost()：分配页锁定内存（Host端），可以加速拷贝

我曾经犯过一个低级错误：在核函数里直接访问Host指针，结果程序崩溃了。记住，GPU不能直接访问CPU内存（除非用统一内存，但那又是另一回事了）。

典型的数据流：

在Host上准备数据
用cudaMalloc在Device上分配空间
用cudaMemcpy把数据从Host拷贝到Device
启动核函数
用cudaMemcpy把结果从Device拷回Host
释放内存

注意：cudaMemcpy是同步操作，会阻塞CPU直到拷贝完成。如果你需要异步操作，可以用cudaMemcpyAsync配合流（Stream）使用。

第一个CUDA程序：向量加法

好了，理论说完了，咱们来写个实际的程序。向量加法是CUDA的"Hello World"，虽然简单，但五脏俱全。

#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);

    // 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] = 1.0f;
        h_B[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. 配置核函数启动参数
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

    // 5. 启动核函数
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

    // 6. 拷贝结果回Host
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    // 验证结果
    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;
}

这个程序虽然简单，但包含了CUDA编程的所有关键步骤。你想想看，如果连这个都跑不通，后面的优化就更别谈了。

调试建议：每次调用CUDA API后，最好检查一下返回值。可以用cudaGetLastError()来捕获错误。我刚开始写CUDA时，经常因为忘记检查错误而浪费大量时间。

知识体系总览

下面这张图，是我自己总结的本章知识体系。你可以把它当作一个地图，随时回来查阅。

好了，第一章的内容就到这里。记住，GPU编程不是玄学，它是有章可循的。理解了硬件结构，你写的每个核函数都会更有底气。

想要更多资料，联系可可

📚 课程目录

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

GPU硬件架构概览

CUDA编程模型：Grid-Block-Thread层次结构

Host与Device内存管理

第一个CUDA程序：向量加法

知识体系总览