01-GPU架构概述¶
重要性：⭐⭐⭐⭐⭐ 实用度：⭐⭐⭐⭐⭐ 学习时间：2天 必须掌握：是
为什么学这一章？¶
GPU已经从单纯的图形处理器发展成为通用并行计算设备。理解GPU架构能帮助你： - 理解为什么GPU适合AI和深度学习 - 编写高效的并行程序 - 优化CUDA程序性能 - 理解现代AI框架的底层实现
学完这一章，你将能够： - ✅ 解释GPU与CPU的设计差异 - ✅ 理解GPU的核心架构组件 - ✅ 掌握GPU编程的基本概念 - ✅ 了解GPU的应用场景
📖 核心概念¶
1. GPU与CPU的设计哲学¶
Text Only
┌─────────────────────────────────────────────────────────────────────┐
│ CPU vs GPU 设计哲学对比 │
├─────────────────────────────────────────────────────────────────────┤
│ │
│ CPU（中央处理器） │
│ ┌───────────────────────────────────────────────────────────────┐ │
│ │ 设计目标：最小化单个任务的延迟（Latency） │ │
│ │ │ │
│ │ 特点： │ │
│ │ • 强大的单核性能 │ │
│ │ • 复杂的控制逻辑（分支预测、乱序执行） │ │
│ │ • 大容量缓存（减少内存访问延迟） │ │
│ │ • 适合串行任务 │ │
│ │ │ │
│ │ 比喻：F1赛车 - 速度快但昂贵，适合复杂赛道 │ │
│ │ │ │
│ │ ┌─────┐ ┌─────┐ ┌─────┐ ┌─────┐ │ │
│ │ │核心1│ │核心2│ │核心3│ │核心4│ ← 少量强大核心 │ │
│ │ └──┬──┘ └──┬──┘ └──┬──┘ └──┬──┘ │ │
│ │ └───────┴───────┴───────┘ │ │
│ │ 共享L3缓存 │ │
│ └───────────────────────────────────────────────────────────────┘ │
│ │
│ GPU（图形处理器） │
│ ┌───────────────────────────────────────────────────────────────┐ │
│ │ 设计目标：最大化吞吐量（Throughput） │ │
│ │ │ │
│ │ 特点： │ │
│ │ • 大量简单核心（数千个） │ │
│ │ • 简单的控制逻辑（无分支预测、SIMT顺序流水线） │ │
│ │ • 小容量缓存（依赖高带宽内存） │ │
│ │ • 适合并行任务 │ │
│ │ │ │
│ │ 比喻：货运船队 - 单船慢但数量多，适合运输大量货物 │ │
│ │ │ │
│ │ ┌──┐ ┌──┐ ┌──┐ ┌──┐ ┌──┐ ┌──┐ ┌──┐ ┌──┐ │ │
│ │ │核│ │核│ │核│ │核│ │核│ │核│ │核│ │核│ ← 大量简单核心 │ │
│ │ └──┘ └──┘ └──┘ └──┘ └──┘ └──┘ └──┘ └──┘ │ │
│ │ ┌──┐ ┌──┐ ┌──┐ ┌──┐ ┌──┐ ┌──┐ ┌──┐ ┌──┐ │ │
│ │ │核│ │核│ │核│ │核│ │核│ │核│ │核│ │核│ │ │
│ │ └──┘ └──┘ └──┘ └──┘ └──┘ └──┘ └──┘ └──┘ │ │
│ │ ... 数千个核心 ... │ │
│ └───────────────────────────────────────────────────────────────┘ │
│ │
└─────────────────────────────────────────────────────────────────────┘
核心差异对比¶
| 特性 | CPU | GPU |
|---|---|---|
| 核心数量 | 4-64个 | 数千个 |
| 核心复杂度 | 复杂（乱序执行、分支预测） | 简单（SIMT模型，同一Warp内线程同步执行相同指令） |
| 时钟频率 | 3-5 GHz | 1-2 GHz |
| 缓存大小 | 大（MB级） | 小（KB级） |
| 内存带宽 | 50-100 GB/s | 500-1000+ GB/s |
| 设计目标 | 低延迟 | 高吞吐量 |
| 适用任务 | 串行、复杂逻辑 | 并行、计算密集 |
2. GPU架构详解¶
NVIDIA GPU架构（以Ampere为例）¶
Text Only
┌─────────────────────────────────────────────────────────────────────┐
│ NVIDIA GPU架构（Ampere） │
├─────────────────────────────────────────────────────────────────────┤
│ │
│ GPU（图形处理器） │
│ ┌───────────────────────────────────────────────────────────────┐ │
│ │ │ │
│ │ ┌─────────────┐ ┌─────────────┐ ┌─────────────┐ │ │
│ │ │ GPC 0 │ │ GPC 1 │ │ GPC 2... │ ← 图形处理集群│ │
│ │ │ (Graphics │ │ │ │ │ │ │
│ │ │ Processing │ │ │ │ │ │ │
│ │ │ Cluster) │ │ │ │ │ │ │
│ │ └──────┬──────┘ └──────┬──────┘ └──────┬──────┘ │ │
│ │ │ │ │ │ │
│ │ ┌──────┴──────┐ ┌──────┴──────┐ ┌──────┴──────┐ │ │
│ │ │ TPC 0 │ │ TPC 0 │ │ TPC 0 │ ← 纹理处理集群│ │
│ │ │ TPC 1 │ │ TPC 1 │ │ TPC 1 │ │ │
│ │ └──────┬──────┘ └──────┬──────┘ └──────┬──────┘ │ │
│ │ │ │ │ │ │
│ │ ┌──────┴──────┐ ┌──────┴──────┐ ┌──────┴──────┐ │ │
│ │ │ SM 0 │ │ SM 0 │ │ SM 0 │ ← 流式多处理器│ │
│ │ │ SM 1 │ │ SM 1 │ │ SM 1 │ (Streaming │ │
│ │ │ ... │ │ ... │ │ ... │ Multiprocessor)│
│ │ └──────┬──────┘ └──────┬──────┘ └──────┬──────┘ │ │
│ │ │ │ │ │ │
│ │ ┌──────┴──────┐ ┌──────┴──────┐ ┌──────┴──────┐ │ │
│ │ │ CUDA Core │ │ CUDA Core │ │ CUDA Core │ ← CUDA核心 │ │
│ │ │ Tensor Core │ │ Tensor Core │ │ Tensor Core │ ← Tensor核心 │ │
│ │ │ RT Core │ │ RT Core │ │ RT Core │ ← 光追核心 │ │
│ │ │ ... │ │ ... │ │ ... │ │ │
│ │ └─────────────┘ └─────────────┘ └─────────────┘ │ │
│ │ │ │
│ │ ┌─────────────────────────────────────────────────────────┐ │ │
│ │ │ 全局内存（Global Memory） │ │ │
│ │ │ 容量：数GB到数十GB │ │ │
│ │ │ 带宽：数百GB/s到TB/s │ │ │
│ │ └─────────────────────────────────────────────────────────┘ │ │
│ │ │ │
│ └───────────────────────────────────────────────────────────────┘ │
│ │
└─────────────────────────────────────────────────────────────────────┘
流式多处理器（SM）详解¶
Text Only
┌─────────────────────────────────────────────────────────────────────┐
│ 流式多处理器（SM）内部结构 │
├─────────────────────────────────────────────────────────────────────┤
│ │
│ SM（Streaming Multiprocessor） │
│ ┌───────────────────────────────────────────────────────────────┐ │
│ │ │ │
│ │ ┌─────────────┐ ┌─────────────┐ ┌─────────────┐ │ │
│ │ │ Warp调度器1 │ │ Warp调度器2 │ │ Warp调度器3 │ ← 调度器 │ │
│ │ │ Warp调度器4 │ │ │ │ │ │ │
│ │ └──────┬──────┘ └──────┬──────┘ └──────┬──────┘ │ │
│ │ │ │ │ │ │
│ │ ┌──────┴──────┐ ┌──────┴──────┐ ┌──────┴──────┐ │ │
│ │ │ CUDA Core │ │ CUDA Core │ │ CUDA Core │ ← FP32单元│ │
│ │ │ (FP32) │ │ (FP32) │ │ (FP32) │ │ │
│ │ │ CUDA Core │ │ CUDA Core │ │ CUDA Core │ │ │
│ │ │ (FP32) │ │ (FP32) │ │ (FP32) │ │ │
│ │ │ ... │ │ ... │ │ ... │ │ │
│ │ │ (64-128个) │ │ │ │ │ │ │
│ │ └─────────────┘ └─────────────┘ └─────────────┘ │ │
│ │ │ │
│ │ ┌─────────────┐ ┌─────────────┐ ┌─────────────┐ │ │
│ │ │ Tensor Core │ │ Tensor Core │ │ Tensor Core │ ← Tensor核心│ │
│ │ │ (4-8个) │ │ │ │ │ (AI加速) │ │
│ │ └─────────────┘ └─────────────┘ └─────────────┘ │ │
│ │ │ │
│ │ ┌─────────────────────────────────────────────────────────┐ │ │
│ │ │ 共享内存（Shared Memory） │ │ │
│ │ │ 容量：几十KB到几百KB │ │ │
│ │ │ 速度：比全局内存快100倍 │ │ │
│ │ └─────────────────────────────────────────────────────────┘ │ │
│ │ │ │
│ │ ┌─────────────────────────────────────────────────────────┐ │ │
│ │ │ 寄存器文件（Register File） │ │ │
│ │ │ 容量：数百KB │ │ │
│ │ │ 每个线程有私有寄存器 │ │ │
│ │ └─────────────────────────────────────────────────────────┘ │ │
│ │ │ │
│ └───────────────────────────────────────────────────────────────┘ │
│ │
└─────────────────────────────────────────────────────────────────────┘
3. GPU编程模型¶
CUDA线程层次结构¶
Text Only
┌─────────────────────────────────────────────────────────────────────┐
│ CUDA线程层次结构 │
├─────────────────────────────────────────────────────────────────────┤
│ │
│ Grid（网格）- 整个GPU程序 │
│ ┌───────────────────────────────────────────────────────────────┐ │
│ │ │ │
│ │ Block(0,0) Block(1,0) Block(2,0) │ │
│ │ ┌─────────┐ ┌─────────┐ ┌─────────┐ │ │
│ │ │Thread │ │Thread │ │Thread │ │ │
│ │ │(0,0) │ │(0,0) │ │(0,0) │ │ │
│ │ │(1,0) │ │(1,0) │ │(1,0) │ │ │
│ │ │(2,0) │ │(2,0) │ │(2,0) │ │ │
│ │ │... │ │... │ │... │ │ │
│ │ │(0,1) │ │(0,1) │ │(0,1) │ │ │
│ │ │... │ │... │ │... │ │ │
│ │ └─────────┘ └─────────┘ └─────────┘ │ │
│ │ │ │
│ │ Block(0,1) Block(1,1) Block(2,1) │ │
│ │ ┌─────────┐ ┌─────────┐ ┌─────────┐ │ │
│ │ │Thread │ │Thread │ │Thread │ │ │
│ │ │... │ │... │ │... │ │ │
│ │ └─────────┘ └─────────┘ └─────────┘ │ │
│ │ │ │
│ │ ... │ │
│ │ │ │
│ └───────────────────────────────────────────────────────────────┘ │
│ │
│ 层次关系： │
│ ├── Grid：包含多个Block，可以是一维、二维或三维 │
│ ├── Block：包含多个Thread，可以是一维、二维或三维 │
│ ├── Thread：最基本的执行单元 │
│ └── Warp：32个Thread组成一个Warp，是调度的基本单位 │
│ │
│ 示例： │
│ dim3 gridDim(3, 2); // 3x2 = 6个Block │
│ dim3 blockDim(4, 4); // 4x4 = 16个Thread/Block │
│ 总线程数 = 6 * 16 = 96个线程 │
│ │
└─────────────────────────────────────────────────────────────────────┘
线程索引计算¶
C++
// CUDA内核函数示例
__global__ void vectorAdd(float* A, float* B, float* C, int n) {
// 计算全局线程ID
int i = blockIdx.x * blockDim.x + threadIdx.x;
// 每个线程处理一个元素
if (i < n) {
C[i] = A[i] + B[i];
}
}
// 调用内核
int main() {
int n = 1000000;
// 配置：256线程/Block，足够多的Block覆盖所有数据
int threadsPerBlock = 256;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, n);
return 0;
}
4. GPU内存层次结构¶
Text Only
┌─────────────────────────────────────────────────────────────────────┐
│ GPU内存层次结构 │
├─────────────────────────────────────────────────────────────────────┤
│ │
│ 速度 ↑ │
│ 快 │ ┌───────────────────────────────────────────────────┐ │
│ │ │ 寄存器（Register） │ │
│ │ │ • 每个线程私有 │ │
│ │ │ • 容量：几十KB/SM │ │
│ │ │ • 延迟：1周期 │ │
│ │ │ • 速度最快 │ │
│ │ └───────────────────────────────────────────────────┘ │
│ │ ↓ │
│ │ ┌───────────────────────────────────────────────────┐ │
│ │ │ 共享内存（Shared Memory / L1 Cache） │ │
│ │ │ • Block内所有线程共享 │ │
│ │ │ • 容量：几十KB到几百KB │ │
│ │ │ • 延迟：20-30周期 │ │
│ │ │ • 可编程管理 │ │
│ │ └───────────────────────────────────────────────────┘ │
│ │ ↓ │
│ │ ┌───────────────────────────────────────────────────┐ │
│ │ │ L2缓存 │ │
│ │ │ • 所有SM共享 │ │
│ │ │ • 容量：几MB │ │
│ │ │ • 延迟：几百周期 │ │
│ │ └───────────────────────────────────────────────────┘ │
│ │ ↓ │
│ 慢 │ ┌───────────────────────────────────────────────────┐ │
│ │ │ 全局内存（Global Memory / 显存） │ │
│ │ │ • 所有线程可访问 │ │
│ │ │ • 容量：几GB到几十GB │ │
│ │ │ • 延迟：几百到几千周期 │ │
│ │ │ • 带宽：数百GB/s │ │
│ │ └───────────────────────────────────────────────────┘ │
│ └ │
│ 容量 → │
│ 小 大 │
│ │
└─────────────────────────────────────────────────────────────────────┘
内存类型对比¶
| 内存类型 | 作用域 | 生命周期 | 速度 | 容量 | 用途 |
|---|---|---|---|---|---|
| 寄存器 | 线程 | 线程 | 最快 | 有限 | 局部变量 |
| 共享内存 | Block | Block | 很快 | 几十KB | 线程间通信 |
| L2缓存 | 所有线程 | 程序 | 快 | 几MB | 自动缓存 |
| 全局内存 | 所有线程 | 程序 | 慢 | 几GB | 主数据存储 |
| 常量内存 | 所有线程 | 程序 | 快（缓存） | 64KB | 只读常量 |
| 纹理内存 | 所有线程 | 程序 | 快（缓存） | 大 | 图像处理 |
🧪 动手实验¶
实验1：查看GPU信息¶
目的：了解你的GPU硬件信息
步骤：
-
安装CUDA工具包（Linux）
-
使用nvidia-smi查看GPU信息
-
使用deviceQuery查看详细信息
观察输出： - GPU型号和架构 - CUDA核心数量 - 显存大小和带宽 - 计算能力（Compute Capability）
实验2：第一个CUDA程序¶
目的：编写并运行简单的CUDA程序
步骤：
-
创建CUDA文件
C++// first_cuda.cu #include <iostream> #include <cuda_runtime.h> // CUDA内核函数 __global__ void helloFromGPU() { int tid = blockIdx.x * blockDim.x + threadIdx.x; printf("Hello from thread %d (Block %d, Thread %d)\n", tid, blockIdx.x, threadIdx.x); } int main() { std::cout << "Hello from CPU!" << std::endl; // 启动内核：2个Block，每个Block 4个线程 helloFromGPU<<<2, 4>>>(); // 等待GPU完成 cudaDeviceSynchronize(); std::cout << "Done!" << std::endl; return 0; } -
编译运行
-
观察输出
- 总共8个线程（2 Block × 4 Thread）
- 每个线程打印自己的ID
实验3：向量加法¶
目的：实现并优化向量加法
步骤：
-
创建CUDA文件
C++// vector_add.cu #include <iostream> // 引入头文件 #include <cuda_runtime.h> #include <chrono> // CUDA内核：向量加法 __global__ void vectorAdd(const float* A, const float* B, float* C, int n) { // 指针：存储变量的内存地址 int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) { C[i] = A[i] + B[i]; } } // CPU版本：用于对比 void vectorAddCPU(const float* A, const float* B, float* C, int n) { for (int i = 0; i < n; i++) { C[i] = A[i] + B[i]; } } int main() { int n = 10000000; // 1000万元素 size_t size = n * sizeof(float); // 分配主机内存 float* h_A = new float[n]; float* h_B = new float[n]; float* h_C = new float[n]; float* h_C_CPU = new float[n]; // 初始化数据 for (int i = 0; i < n; i++) { h_A[i] = rand() / (float)RAND_MAX; h_B[i] = rand() / (float)RAND_MAX; } // CPU计算 auto start = std::chrono::high_resolution_clock::now(); vectorAddCPU(h_A, h_B, h_C_CPU, n); auto end = std::chrono::high_resolution_clock::now(); auto cpu_time = std::chrono::duration_cast<std::chrono::microseconds>(end - start).count(); // 分配设备内存 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; cudaEvent_t start_gpu, stop_gpu; cudaEventCreate(&start_gpu); cudaEventCreate(&stop_gpu); cudaEventRecord(start_gpu); vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, n); cudaEventRecord(stop_gpu); cudaEventSynchronize(stop_gpu); float gpu_time = 0; cudaEventElapsedTime(&gpu_time, start_gpu, stop_gpu); // 拷贝结果回主机 cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost); // 验证结果 bool correct = true; for (int i = 0; i < n; i++) { if (fabs(h_C[i] - h_C_CPU[i]) > 1e-5) { correct = false; break; } } std::cout << "CPU Time: " << cpu_time / 1000.0 << " ms" << std::endl; std::cout << "GPU Time: " << gpu_time << " ms" << std::endl; std::cout << "Speedup: " << cpu_time / 1000.0 / gpu_time << "x" << std::endl; std::cout << "Result: " << (correct ? "CORRECT" : "INCORRECT") << std::endl; // 清理 delete[] h_A; delete[] h_B; delete[] h_C; delete[] h_C_CPU; cudaFree(d_A); cudaFree(d_B); cudaFree(d_C); return 0; } -
编译运行
-
观察加速比
- 通常GPU比CPU快10-100倍（取决于GPU型号和数据大小）
💡 核心要点总结¶
GPU vs CPU¶
| 特性 | CPU | GPU |
|---|---|---|
| 核心数 | 少（4-64） | 多（数千） |
| 核心复杂度 | 高 | 低 |
| 设计目标 | 低延迟 | 高吞吐量 |
| 适用场景 | 串行任务 | 并行任务 |
GPU架构层次¶
Text Only
GPU
├── GPC（图形处理集群）
│ └── TPC（纹理处理集群）
│ └── SM（流式多处理器）
│ ├── CUDA Core（计算核心）
│ ├── Tensor Core（AI加速）
│ ├── 共享内存
│ └── 寄存器
└── 全局内存
CUDA线程层次¶
GPU内存层次¶
❓ 常见问题¶
Q1：为什么GPU适合深度学习？
A：深度学习主要是矩阵运算，具有： - 高并行性：矩阵乘法可以分解成大量独立的乘加运算 - 计算密集：需要大量浮点运算 - 数据局部性：可以充分利用缓存和共享内存
Q2：所有程序都适合用GPU加速吗？
A：不是。适合GPU的程序特征： - 大规模数据并行 - 计算密集（计算/访存比高） - 控制流简单（少分支）
不适合GPU的程序： - 串行任务 - 大量分支判断 - 需要复杂同步
Q3：CUDA和OpenCL有什么区别？
A： - CUDA：NVIDIA专有，功能丰富，生态完善 - OpenCL：开放标准，跨平台（NVIDIA/AMD/Intel），但功能相对简单
Q4：如何选择Block大小？
A：一般原则： - 使用32的倍数（Warp大小） - 常见选择：128, 256, 512 - 考虑寄存器使用量（太多线程会导致寄存器溢出）
📚 扩展阅读¶
- 《CUDA并行程序设计》 - 相关书籍
- NVIDIA CUDA文档：docs.nvidia.com/cuda/
- 《Programming Massively Parallel Processors》 - David Kirk
🎯 下一步¶
学习 02-GPU编程模型，深入理解CUDA线程层次和编程模型。