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

大家好，我是你们的老朋友。今天咱们聊聊CUDA编程最核心的东西——GPU到底是怎么并行计算的。说实话，我刚开始接触CUDA时，也被那一堆线程、内存层次搞得晕头转向。但后来我发现，只要理解了三个关键概念，一切就豁然开朗了。

1.1 GPU并行计算原理：为什么GPU这么快？

先问个问题：为什么CPU和GPU长得不一样？

CPU像是一个全能型选手——什么都能干，但一次只能干几件事。GPU呢？它更像是一支军队——每个士兵只会简单的操作，但成千上万个士兵一起上，威力惊人。

我个人习惯把GPU比作「人海战术」。你想想看，一个复杂的计算任务，比如矩阵乘法，CPU可能派4个核心去算，每个核心都很聪明，但数量少。GPU直接派出几千个核心，每个核心只做「乘加」这种简单操作，但架不住人多啊！

核心思想：GPU通过大量轻量级线程并行执行，用数量弥补单线程性能的不足。适合数据并行、计算密集型的任务。

我在项目中遇到过这样一个场景：训练一个BERT模型，CPU跑一次前向传播要十几秒，换成GPU后，同样的计算量，几百毫秒就搞定了。这就是并行计算的威力。

1.2 CUDA线程层次结构：Grid-Block-Thread

好了，现在我们知道GPU有很多线程。但这些线程怎么组织呢？CUDA给出了一个三层结构：

Thread（线程）：最基础的计算单元
Block（线程块）：一组线程，可以协作
Grid（网格）：一组Block，构成整个任务

说白了，这就是一个「公司-部门-员工」的关系。Grid是公司，Block是部门，Thread是员工。公司（Grid）由多个部门（Block）组成，每个部门里有若干员工（Thread）。

为什么会这样设计？嗯，这里要注意：GPU硬件上，线程是以「warp」（32个线程一组）为单位调度的。Block内的线程可以共享数据，但Block之间是独立的。这个设计直接影响了我们写代码的方式。

避坑指南：我曾经犯过一个错误——把Block设得太大，结果每个Block里的线程数超过了硬件限制（通常是1024），程序直接崩溃。后来我养成了习惯：先查硬件规格，再设参数。

来看一个简单的代码示例：

// 定义Grid和Block的维度
dim3 gridDim(2, 2);      // 2x2的网格，共4个Block
dim3 blockDim(16, 16);   // 16x16的Block，共256个线程

// 启动内核
myKernel<<<gridDim, blockDim>>>(d_data);

// 内核函数
__global__ void myKernel(float* data) {
    // 获取当前线程的全局索引
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int idy = blockIdx.y * blockDim.y + threadIdx.y;
    int globalIdx = idy * gridDim.x * blockDim.x + idx;
    
    // 执行计算
    data[globalIdx] = data[globalIdx] * 2.0f;
}

这段代码里，每个线程都知道了自己的「坐标」，然后去处理对应的数据。这就是CUDA编程的基本模式。

1.3 内存层次结构：Global/Shared/Register/Local

讲完线程，咱们聊聊数据。GPU的内存就像一座金字塔，从下到上，速度越来越快，容量越来越小。

内存类型	位置	速度	容量	作用域
Global Memory	显存	慢（~200周期）	大（GB级）	所有线程
Shared Memory	芯片内	快（~5周期）	小（KB级）	Block内
Register	寄存器	极快（1周期）	极小（每个线程几十个）	单个线程
Local Memory	显存（寄存器溢出）	慢	取决于寄存器使用	单个线程

我个人习惯把Global Memory比作「仓库」——什么都能放，但取东西要走很远。Shared Memory是「工位上的抽屉」——同一个小组的人可以快速共享。Register是「手里的工具」——最快，但只能自己用。

注意：Local Memory其实是个「陷阱」。当你声明太多局部变量，寄存器放不下时，编译器会自动把多余的变量放到Local Memory里。但Local Memory物理上还是在Global Memory里，速度很慢。我曾经调试一个内核，发现性能比预期差很多，一查——寄存器溢出导致大量Local Memory访问。解决办法：减少每个线程的局部变量数量。

1.4 实战：矩阵乘法中的内存优化

光说不练假把式。咱们拿矩阵乘法来演示一下内存层次怎么用。

最朴素的写法：

__global__ void matMulNaive(float* A, float* B, float* C, int N) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    
    float sum = 0.0f;
    for (int k = 0; k < N; k++) {
        sum += A[row * N + k] * B[k * N + col];
    }
    C[row * N + col] = sum;
}

这段代码能跑，但效率很低。为什么？因为每个线程都在反复读取Global Memory，而A和B的数据其实可以被多个线程复用。

优化版本——使用Shared Memory：

__global__ void matMulShared(float* A, float* B, float* C, int N) {
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];
    
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    
    float sum = 0.0f;
    for (int tile = 0; tile < N / TILE_SIZE; tile++) {
        // 协作加载数据到Shared Memory
        As[threadIdx.y][threadIdx.x] = A[row * N + tile * TILE_SIZE + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(tile * TILE_SIZE + threadIdx.y) * N + col];
        __syncthreads();  // 同步：确保所有线程都加载完毕
        
        // 从Shared Memory计算
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        __syncthreads();  // 同步：确保所有线程都计算完毕，再加载下一块
    }
    C[row * N + col] = sum;
}

这个优化版本，说白了就是把数据先搬到Shared Memory里，然后反复使用。我实测过，对于1024x1024的矩阵，优化后的版本比朴素版本快了将近10倍。

小技巧：使用Shared Memory时，记得加__syncthreads()同步。我曾经漏掉一个同步点，结果线程A还在读数据，线程B已经把数据覆盖了——算出来的结果全是错的。调试了整整一个下午才找到问题。

1.5 知识体系总览

为了让大家更直观地理解本章的知识结构，我画了一张图：

这张图把本章的核心内容串起来了。从上到下，先理解GPU为什么快，再学会怎么组织线程，最后掌握怎么用好内存。这三块搞懂了，CUDA编程就算入门了。

好了，第一章就到这里。记住：线程层次决定「谁干活」，内存层次决定「数据怎么传」。这两件事想明白了，后面优化性能就顺理成章了。

本章要点回顾：

GPU通过大量线程并行执行，适合数据密集型任务
线程层次：Grid → Block → Thread，三层结构
内存层次：Global（慢但大）→ Shared（快但小）→ Register（最快）
Shared Memory是性能优化的关键，但要注意同步
避免寄存器溢出导致Local Memory访问

想要更多资料，联系可可

📚 课程目录

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

1.1 GPU并行计算原理：为什么GPU这么快？

1.2 CUDA线程层次结构：Grid-Block-Thread

1.3 内存层次结构：Global/Shared/Register/Local

1.4 实战：矩阵乘法中的内存优化

1.5 知识体系总览