第三章 GPU架构与CUDA概述

CUDA之所以能够将GPU强大的并行计算能力释放出来，其根基在于对GPU硬件架构的深刻理解与适配。本章将首先剖析现代GPU的物理组成——从流式多处理器（SM）到CUDA核心，再到多层次显存体系；随后，我们将建立CUDA编程模型中的“主机-设备”协作视角，理解内核函数（Kernel）的执行逻辑；最后，完整梳理从驱动安装到编译器配置的CUDA开发环境搭建流程。这三部分构成了后续所有并行优化技巧的硬件与软件基石。

3.1 GPU硬件架构核心原理

现代GPU（如NVIDIA Ampere、Hopper架构）本质上是大规模并行吞吐处理器。与CPU追求低延迟、复杂控制不同，GPU将大量晶体管用于计算单元（ALU）和存储带宽，而非分支预测与乱序执行。其核心设计哲学是“用更多的线程掩盖延迟”——当一部分线程等待内存数据时，调度器立即切换到其他就绪线程，从而保持计算单元持续忙碌。

3.1.1 流式多处理器（SM）与SP核心

GPU的最小独立计算单元是流式多处理器（Streaming Multiprocessor, SM）。一个GPU由多个SM组成，例如NVIDIA A100拥有108个SM。每个SM内部包含：

• CUDA核心（SP, Streaming Processor）：执行整数和浮点运算的ALU单元。一个SM通常包含64~128个CUDA核心。注意，这些核心是“轻量级”的，它们共享控制逻辑，必须由调度器统一指挥。
• 特殊功能单元（SFU）：执行超越函数（如sin、cos、log）等复杂运算。
• 加载/存储单元（LD/ST）：负责在显存和寄存器之间搬运数据。
• Warp调度器：每个SM包含4个或更多调度器，每个调度器每时钟周期可发射一条指令给一个Warp（32个线程为一组）。这是GPU执行的基本单位。
• 寄存器文件（Register File）：容量极大（如256KB），用于存储线程的局部变量。寄存器分配直接决定每个SM能容纳多少线程。
• 共享内存（Shared Memory）：SM内部的高速可编程缓存，容量通常为96KB~228KB，供同一线程块内的所有线程访问，延迟极低（约20~30周期）。

3.1.2 显存层次结构

GPU的存储体系遵循“容量越大、速度越慢”的规律，理解这一层次对优化至关重要：

全局内存是GPU的主存储区，容量大但延迟高。访问全局内存时，硬件以32字节、64字节或128字节的段为单位进行传输。如果同一Warp的32个线程访问的地址能合并成连续的段，则只需一次事务；否则需要多次事务，造成带宽浪费。这就是内存合并（Memory Coalescing）技术。

3.2 CUDA编程模型：主机-设备协作

CUDA采用异构编程模型，将CPU视为主机（Host），GPU视为设备（Device）。主机负责串行逻辑、数据准备和内核启动；设备负责大规模并行计算。两者通过PCIe总线或NVLink互联。

3.2.1 内核函数（Kernel）

内核函数是运行在GPU上的函数，使用__global__修饰符声明。调用内核时，需要指定执行配置<<<gridDim, blockDim>>>，其中：

• gridDim：网格中线程块（Block）的数量。
• blockDim：每个线程块中线程（Thread）的数量。

这种两级并行结构直接映射到硬件：线程块被调度到SM上执行，块内的线程通过共享内存协作。一个SM可以同时驻留多个线程块，但受限于寄存器、共享内存等资源。

// 内核定义：每个线程计算一个元素
__global__ void vectorAdd(float *A, float *B, float *C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;  // 全局线程ID
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

// 主机端调用
int main() {
    int N = 1 << 20;
    size_t bytes = N * sizeof(float);
    // 1. 分配设备内存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, bytes);
    cudaMalloc(&d_B, bytes);
    cudaMalloc(&d_C, bytes);
    // 2. 将输入数据从主机拷贝到设备
    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);
    // 3. 启动内核：256线程/块，共N/256块
    int threads = 256;
    int blocks = (N + threads - 1) / threads;
    vectorAdd<<<blocks, threads>>>(d_A, d_B, d_C, N);
    // 4. 将结果拷贝回主机
    cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);
    // 5. 释放设备内存
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    return 0;
}

3.2.2 线程层次与索引计算

CUDA定义了清晰的线程层次：网格（Grid）→ 线程块（Block）→ 线程（Thread）。每个线程通过内置变量threadIdx（块内索引）、blockIdx（块索引）、blockDim（块维度）计算自己的全局唯一ID。这种设计使得程序可以灵活处理任意大小的数据，同时保持硬件利用率。

3.3 CUDA开发环境搭建

搭建一个可用的CUDA开发环境是进行并行编程的前提。环境包含三个层次：驱动（Driver）→ 工具包（Toolkit）→ 编译器（Compiler）。下面以Linux（Ubuntu 22.04）和Windows 11为例说明。

3.3.1 安装步骤与验证

1. 安装NVIDIA显卡驱动：驱动是操作系统与GPU通信的桥梁。推荐从NVIDIA官网下载对应型号的最新驱动，或使用系统包管理器安装。安装后运行nvidia-smi验证驱动版本和GPU信息。
2. 安装CUDA Toolkit：Toolkit包含编译器（nvcc）、运行时库、调试工具（cuda-gdb）、性能分析器（Nsight）等。下载对应操作系统的runfile或deb包。安装时注意选择与驱动兼容的版本。
3. 配置环境变量：将CUDA的bin和lib64目录添加到PATH和LD_LIBRARY_PATH中。
4. 验证安装：编写一个简单的“Hello World”内核并编译运行。

# 验证驱动
nvidia-smi
# 输出示例：
# +-----------------------------------------------------------------------------+
# | NVIDIA-SMI 525.85.12    Driver Version: 525.85.12    CUDA Version: 12.0      |
# +-------------------------------+----------------------+----------------------+

# 验证编译器
nvcc --version
# 输出示例：
# nvcc: NVIDIA (R) Cuda compiler driver
# Copyright (c) 2005-2023 NVIDIA Corporation
# Built on Wed_Nov_22_10:17:15_PST_2023
# Cuda compilation tools, release 12.3, V12.3.107

# 编译并运行示例
nvcc -o hello hello.cu
./hello

3.3.2 常见环境问题排查

3.4 实战案例：向量加法性能剖析

为了将上述概念串联起来，我们通过一个完整的向量加法案例，观察不同配置对性能的影响。实验环境：NVIDIA RTX 3090（GA102架构，82个SM，每个SM最大2048线程），CUDA 12.0。

3.4.1 实验设计与数据

我们固定数据规模N=1,000,000（约4MB），分别测试不同线程块大小（64、128、256、512、1024）下的内核执行时间。使用cudaEvent计时，并计算有效带宽。

// 计时代码片段
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
vectorAdd<<<blocks, threads>>>(d_A, d_B, d_C, N);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
float ms = 0;
cudaEventElapsedTime(&ms, start, stop);
printf("Time: %.3f ms, Bandwidth: %.2f GB/s\n", ms, 2*N*4/ms/1e6);

3.4.2 结果与复盘

分析过程：当块大小为64时，每个SM只能驻留32个块（2048/64），但每个块只有2个Warp，调度器无法有效隐藏内存延迟，导致带宽利用率低。块大小为256时，每个SM驻留8个块，共32个Warp，足以填满调度器流水线，带宽达到峰值。块大小继续增大到1024时，每个SM仅驻留2个块，虽然线程总数相同，但块间同步和调度灵活性下降，性能略有回落。

3.5 常见误区与总结

初学者在理解GPU架构和CUDA模型时，常陷入以下误区：

• 误区一：线程越多越快。实际上，当SM的线程容量饱和后，增加线程只会增加上下文切换开销。关键是提高占用率（Occupancy）与指令级并行。
• 误区二：共享内存是万能的。共享内存容量有限，且需要手动同步（__syncthreads()），使用不当会导致死锁或bank conflict。
• 误区三：忽略内存合并。如果线程访问全局内存的地址不连续，带宽会急剧下降。例如，按列访问二维数组时，应改为按行访问或使用共享内存转置。
• 误区四：主机与设备数据传输开销。PCIe传输（约16 GB/s）远低于显存带宽，频繁的小数据传输会严重拖慢性能。应尽量在设备端完成数据预处理。