第25章 GPU编程入门

课程目录第25章 GPU编程入门

第25章 GPU编程入门：CUDA基础、GPU内存模型、核函数编写、性能优化

说实话，我第一次接触CUDA的时候，心里是有点发怵的。那时候项目里要处理海量图像数据，CPU跑起来像老牛拉车。我记得当时组里一位老大哥说：“试试GPU吧，这东西天生就是干并行活的。” 嗯，从那以后，我就一头扎进了这个领域。今天咱们就来聊聊GPU编程的入门核心，我会尽量把那些踩过的坑、积累的经验都揉碎了讲给你听。

25.1 CUDA基础：从CPU到GPU的思维转变

你想想看，CPU是“精打细算”的管家，擅长处理复杂的逻辑和串行任务。而GPU呢？它更像一个“人海战术”的指挥官，手里握着成千上万个简单计算单元。说白了，GPU编程的核心就是并行化。

CUDA（Compute Unified Device Architecture）是NVIDIA推出的并行计算平台。我个人习惯把CUDA程序看作两部分：

Host（主机）：CPU及其内存，负责控制逻辑和数据搬运。
Device（设备）：GPU及其显存，负责大规模并行计算。

一个典型的CUDA程序流程是这样的：

在Host上分配内存，准备数据。
在Device上分配显存。
将数据从Host拷贝到Device。
在Device上启动核函数（Kernel）执行计算。
将结果从Device拷贝回Host。
释放资源。

核心概念： 核函数（Kernel）是运行在GPU上的函数，用 __global__ 修饰符声明。调用时使用 <<<grid, block>>> 语法指定线程组织方式。

25.2 GPU内存模型：你的数据住在哪？

GPU的内存模型，说白了就是一张“数据地图”。你写代码时，数据放在哪里，直接决定了程序的性能。我在项目中遇到过好几次，明明算法没问题，但就是跑不快，最后发现是内存访问模式不对。

来看这张图，它清晰地展示了CUDA的内存层次结构：

这张图我建议你多看几遍。为什么？因为很多性能问题，根源都在于数据放错了地方。比如，频繁访问的变量如果放在全局内存里，那性能肯定上不去。

我的经验： 在写核函数时，尽量把频繁使用的数据从全局内存搬到共享内存或寄存器中。我曾经优化过一个矩阵乘法，仅仅是把数据预取到共享内存，性能就提升了近10倍。

25.3 核函数编写：让GPU动起来

核函数是GPU编程的灵魂。它的写法跟普通C函数很像，但有几个关键点要注意。

先看一个最简单的例子——向量加法：

// 核函数定义
__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 threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    // ... 结果拷贝和清理 ...
}

这里有几个内置变量，你得记住：

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

为什么会这样设计？因为GPU的线程是分层组织的。一个Grid包含多个Block，一个Block包含多个Thread。这种层次结构，说白了就是为了方便管理和调度。

注意： 核函数内部不能调用 printf（调试时可以用 printf，但正式代码要删掉）。也不能使用C++标准库中的大部分函数。另外，核函数不能有返回值，必须是 void 类型。

25.4 性能优化：把GPU的潜力榨干

写一个能跑的CUDA程序不难，但写一个高性能的CUDA程序，就需要点功夫了。我总结了几个关键点，都是我在实际项目中反复验证过的。

25.4.1 内存访问合并

GPU访问全局内存时，是以“warp”（32个线程为一组）为单位进行的。如果同一个warp内的线程访问连续的内存地址，硬件会把多次访问合并成一次，效率极高。反之，如果访问是随机的，那性能就会大打折扣。

我曾经在一个图像处理项目中，因为数据排列方式不对，导致内存访问不连续，性能只有理论值的十分之一。后来调整了数据布局，问题就解决了。

25.4.2 共享内存的使用

共享内存是GPU上的“高速缓存”，由程序员手动管理。它的延迟比全局内存低得多。我建议你把那些会被多次访问的数据，比如卷积核的权重、矩阵的块数据，都放到共享内存里。

来看一个使用共享内存的矩阵乘法示例：

__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 bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    
    int row = by * TILE_SIZE + ty;
    int col = bx * TILE_SIZE + tx;
    
    float sum = 0.0f;
    for (int k = 0; k < N / TILE_SIZE; ++k) {
        As[ty][tx] = A[row * N + k * TILE_SIZE + tx];
        Bs[ty][tx] = B[(k * TILE_SIZE + ty) * N + col];
        __syncthreads();
        
        for (int i = 0; i < TILE_SIZE; ++i) {
            sum += As[ty][i] * Bs[i][tx];
        }
        __syncthreads();
    }
    C[row * N + col] = sum;
}

注意这里的 __syncthreads()，它的作用是同步同一个Block内的所有线程。为什么需要同步？因为我们要确保所有线程都把数据加载到共享内存后，才能开始计算。否则，有的线程可能还在读，有的已经开始算了，结果肯定不对。

25.4.3 减少分支发散

GPU的warp是“单指令多线程”（SIMT）执行的。如果warp内的线程走了不同的分支（比如if-else），那这些分支会串行执行，性能会下降。所以，写代码时要尽量避免warp内的分支发散。

优化原则： 尽量让同一个warp内的线程执行相同的指令。如果无法避免分支，尝试让分支条件与线程ID相关，这样至少能保证部分warp不发散。

25.4.4 最大化占用率

占用率（Occupancy）是指每个SM（流多处理器）上活跃的warp数量与最大warp数量的比值。占用率越高，说明GPU的计算资源利用得越充分。影响占用率的因素包括：每个线程使用的寄存器数量、共享内存大小、Block大小等。

你可以用NVIDIA提供的 cudaOccupancyMaxPotentialBlockSize 函数来计算最优的Block大小。我个人习惯是，先设一个合理的Block大小（比如256），然后根据性能测试结果微调。

优化方向	关键点	常见误区
内存访问	合并访问、使用共享内存	随机访问、未对齐访问
线程调度	减少分支发散、提高占用率	warp内分支过多、Block太小
数据传输	减少Host-Device拷贝、使用异步传输	频繁小数据拷贝、未使用流
指令优化	使用快速数学函数、减少除法和取模	过度优化、牺牲精度

避坑指南： 我曾经为了追求高占用率，把Block设得特别小，结果发现性能反而下降了。后来才明白，占用率不是越高越好，还要考虑每个线程的工作量。如果线程太“瘦”，调度开销会吃掉性能收益。

25.5 总结

GPU编程入门，说白了就是三件事：理解内存模型、写好核函数、做好性能优化。CUDA的学习曲线确实有点陡，但一旦你掌握了这些核心概念，就会发现它其实很有规律。

嗯，这里要提醒你一句：不要一开始就追求极致的性能。先把功能跑通，再用工具（比如NVIDIA Nsight）分析瓶颈，然后针对性地优化。我在项目中就是这么做的，效果还不错。

最后，送你一句话：GPU编程，慢就是快，快就是慢。多思考，多动手，你也能写出高性能的CUDA程序。

📚 课程目录