GPU 计算基础

GPU（图形处理器）是现代 AI 计算的核心硬件。理解 GPU 的工作原理和编程模型，对于优化 AI 应用的性能至关重要。本章将从硬件架构、编程模型到性能优化，全面介绍 GPU 计算的核心知识。

GPU 与 CPU 的区别

GPU 和 CPU 在设计理念上有根本性的不同：

特性	CPU	GPU
核心数量	少（几十个）	多（数千个）
单核性能	强	弱
内存带宽	较低	极高
适用场景	串行任务、复杂逻辑	并行任务、简单计算
设计目标	低延迟	高吞吐量

CPU 适合处理复杂的逻辑控制和串行任务，而 GPU 则擅长处理大规模并行计算。这种差异源于它们的设计目标：CPU 追求单线程的低延迟，GPU 追求整体的高吞吐量。

一个形象的比喻

如果把计算任务比作运送货物：

CPU 就像一辆法拉利，速度快但载货量小，适合运送少量急需的货物
GPU 就像一支由小型卡车组成的车队，每辆车速度不快但数量庞大，适合运送大量货物

为什么 GPU 适合深度学习

深度学习的核心计算是矩阵乘法，这正是 GPU 擅长的并行计算任务。以一个 $1024 \times 1024$ 的矩阵乘法为例：

矩阵中的每个元素计算是独立的，可以并行执行
GPU 可以同时启动数百个线程计算不同的元素
单个元素的计算相对简单，不需要复杂的控制逻辑

这就是为什么一个拥有 16896 个 CUDA 核心的 H100 GPU，在矩阵运算上比传统 CPU 快数百倍的原因。

GPU 架构详解

流式多处理器（SM）

GPU 的核心计算单元是流式多处理器（Streaming Multiprocessor，SM）。每个 SM 包含：

CUDA 核心：执行浮点和整数运算
张量核心：专门用于矩阵运算，对深度学习至关重要
共享内存：SM 内部的高速缓存
寄存器文件：存储线程的局部变量
调度单元：管理线程的执行

以 NVIDIA H100 为例，单个 GPU 包含 132 个 SM，每个 SM 有 128 个 CUDA 核心，总计 16896 个核心。H200 基于 Hopper 架构，SM 数量与 H100 相同，但配备更大的 141GB HBM3e 显存。

SM 的执行模型：

SM 采用 SIMT（单指令多线程）执行模型。每个 SM 可以同时执行多个线程块（Block），每个线程块包含多个线程。所有线程执行相同的代码，但可以处理不同的数据。

理解 SM 的工作方式对于优化 GPU 程序至关重要：

线程束（Warp）：每 32 个线程组成一个线程束，是 SM 的基本调度单位
线程束分歧：同一线程束内的线程如果执行不同的分支路径，会串行执行，降低效率
占用率：SM 上活跃的线程束数量与最大容量的比值，影响隐藏延迟的能力

内存层次结构

GPU 的内存层次结构对性能有重大影响：

┌─────────────────────────────────────────┐
│              HBM（高带宽内存）            │
│           容量：80GB，带宽：3.35TB/s      │
├─────────────────────────────────────────┤
│              L2 缓存                     │
│           容量：50MB，延迟：低            │
├─────────────────────────────────────────┤
│    共享内存 / L1 缓存（每个 SM）          │
│           容量：228KB，延迟：极低          │
├─────────────────────────────────────────┤
│              寄存器文件                   │
│           容量：256KB，延迟：最低          │
└─────────────────────────────────────────┘

关键原则：数据越靠近计算单元，访问速度越快。优化 GPU 程序的核心就是最大化数据在高速存储中的复用。

各级内存访问延迟对比：

存储类型	延迟（时钟周期）	带宽	容量
寄存器	1	极高	最小
共享内存	~20	高	有限
L1/L2 缓存	~30-50	高	有限
HBM	~400-600	中	较大

优化策略：将频繁访问的数据放入共享内存或寄存器，减少对 HBM 的访问次数。

张量核心

张量核心（Tensor Core）是 NVIDIA GPU 中专门用于矩阵乘法加速的硬件单元。它可以在一个时钟周期内完成一个 4×4 矩阵乘法运算。

对于深度学习中最常见的矩阵乘法 C = A × B：

传统 CUDA 核心：需要 64 次乘法和 48 次加法
张量核心：单次操作完成

H100 的张量核心在 FP16 精度下可提供近 2000 TFLOPS 的算力。

张量核心支持的数据类型：

数据类型	说明	适用场景
FP16	半精度浮点	训练、推理
BF16	脑浮点	训练（更好的数值稳定性）
TF32	Tensor Float 32	A100+ 默认精度
INT8	8 位整数	推理量化
FP8	8 位浮点	H100+ 推理加速

如何利用张量核心：

在实际开发中，不需要手动编写张量核心代码。深度学习框架和库（如 cuBLAS、cuDNN）会自动使用张量核心。要充分利用张量核心，需要注意：

使用 FP16/BF16 混合精度训练
确保矩阵维度是 8 或 16 的倍数
使用高度优化的库函数

CUDA 编程模型

CUDA（Compute Unified Device Architecture）是 NVIDIA 推出的并行计算平台和编程模型。它允许开发者使用类似 C 的语法编写 GPU 程序，无需深入了解底层硬件细节。

线程层次结构

CUDA 的线程组织分为三个层次：

Grid（网格）
  └── Block（线程块）
        └── Thread（线程）

Thread（线程）：最小的执行单元，执行 kernel 中的代码
Block（线程块）：一组线程，可以协作执行，共享内存
Warp（线程束）：32 个线程的集合，是 SM 的调度单位
Grid（网格）：一组线程块，组成完整的 kernel 执行

理解线程层次结构的意义：

层次	特性	限制
Thread	最小执行单元	拥有私有寄存器和局部内存
Block	线程可同步、共享内存	最多 1024 个线程
Warp	32 线程同步执行	分歧会降低效率
Grid	完整的 kernel 执行	线程块相互独立

线程索引

每个线程都有唯一的索引，用于确定它应该处理的数据：

// 一维索引
int idx = blockIdx.x * blockDim.x + threadIdx.x;

// 二维索引
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;

// 三维索引（用于处理立体数据）
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;

索引计算示例：

假设有 1024 个数据元素，每个 Block 有 256 个线程：

需要 4 个 Block（gridDim = 4）
Block 0 的线程索引范围：0-255
Block 1 的线程索引范围：256-511
Block 2 的线程索引范围：512-767
Block 3 的线程索引范围：768-1023

Kernel 函数

Kernel 是在 GPU 上执行的函数，使用 __global__ 关键字声明：

// 向量加法 kernel
__global__ void vectorAdd(float *a, float *b, float *c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

// 主机端调用
int blockSize = 256;
int numBlocks = (n + blockSize - 1) / blockSize;
vectorAdd<<<numBlocks, blockSize>>>(d_a, d_b, d_c, n);

Kernel 函数的修饰符：

修饰符	执行位置	调用者
`__global__`	GPU	CPU 或 GPU
`__device__`	GPU	GPU
`__host__`	CPU	CPU

一个完整的 CUDA 程序示例：

#include <cuda_runtime.h>
#include <stdio.h>

// Kernel 函数：向量加法
__global__ void vectorAdd(const float *a, const float *b, float *c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    int n = 1 << 20;  // 1M 元素
    size_t size = n * sizeof(float);
    
    // 分配主机内存
    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;
    }
    
    // 分配设备内存
    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);
    
    // 启动 kernel
    int blockSize = 256;
    int numBlocks = (n + blockSize - 1) / blockSize;
    vectorAdd<<<numBlocks, blockSize>>>(d_a, d_b, d_c, n);
    
    // 数据传输：设备到主机
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
    
    // 验证结果
    for (int i = 0; i < 5; i++) {
        printf("c[%d] = %f\n", i, h_c[i]);
    }
    
    // 释放内存
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);
    
    return 0;
}

内存管理

CUDA 程序需要显式管理主机（CPU）和设备（GPU）之间的数据传输：

// 分配设备内存
float *d_a;
cudaMalloc(&d_a, size);

// 数据传输：主机到设备
cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);

// 数据传输：设备到主机
cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);

// 释放内存
cudaFree(d_a);

CUDA 内存类型：

内存类型	位置	访问者	速度
全局内存	HBM	所有线程	较慢
共享内存	SM 内	同 Block 线程	快
寄存器	SM 内	单线程	最快
常量内存	HBM+缓存	所有线程	快（命中时）
纹理内存	HBM+缓存	所有线程	快（空间局部性）

使用共享内存优化

共享内存是 SM 内部的高速存储，用于同一线程块内的线程协作：

#define TILE_SIZE 16

// 矩阵乘法：使用共享内存优化
__global__ void matrixMulShared(float *A, float *B, float *C, int width) {
    __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 t = 0; t < width / TILE_SIZE; t++) {
        // 加载数据到共享内存
        As[threadIdx.y][threadIdx.x] = A[row * width + t * TILE_SIZE + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * width + col];
        __syncthreads();  // 同步：确保数据加载完成
        
        // 计算部分结果
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        __syncthreads();  // 同步：确保计算完成
    }
    
    C[row * width + col] = sum;
}

优化效果：使用共享内存后，矩阵乘法的全局内存访问量减少到原来的 1/TILE_SIZE，显著提升性能。

性能优化原则

GPU 程序优化的核心目标是最大化硬件利用率。以下是几个关键的优化原则：

最大化并行度

GPU 的性能来自于大规模并行。要充分利用 GPU，需要：

足够的线程数量：通常需要数千到数万个线程来隐藏内存延迟
合理的线程块大小：通常是 128、256 或 512，需要是 32 的倍数
避免线程束分歧：同一线程束内的线程应尽量执行相同的代码路径

线程块大小选择指南：

线程块大小	适用场景	注意事项
32	简单操作、高寄存器需求	占用率可能较低
64	平衡选择	-
128	常用选择	较好的占用率
256	内存带宽敏感	共享内存使用较多
512	特定场景	可能超出寄存器限制

优化内存访问

内存带宽通常是 GPU 程序的瓶颈。优化内存访问是提升性能的关键。

合并访问（Coalesced Access）：相邻线程访问相邻的内存地址，可以合并为一次内存事务

// 好的访问模式（合并访问）
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[idx];

// 差的访问模式（跨步访问）
float val = data[threadIdx.x * stride];

内存访问模式对比：

访问模式	带宽利用率	说明
合并访问	100%	相邻线程访问相邻地址
跨步访问	降低	跨度为 2 降低 50%，跨度为 32 降低 97%
随机访问	极低	每次访问单独事务

使用共享内存：共享内存比全局内存快很多，适合存储需要重复访问的数据

__shared__ float sharedData[BLOCK_SIZE];
sharedData[threadIdx.x] = globalData[idx];
__syncthreads();  // 同步线程块内所有线程

减少数据传输

主机和设备之间的数据传输是性能瓶颈：

尽量减少传输次数
使用 pinned memory 加速传输
考虑使用统一内存（Unified Memory）

Pinned Memory 使用示例：

// 分配 pinned 内存（页锁定内存）
float *h_a;
cudaMallocHost(&h_a, size);  // 比 malloc 更快传输

// 异步传输
cudaStream_t stream;
cudaStreamCreate(&stream);
cudaMemcpyAsync(d_a, h_a, size, cudaMemcpyHostToDevice, stream);

// 释放
cudaFreeHost(h_a);

统一内存示例：

// 使用统一内存，自动管理数据传输
float *data;
cudaMallocManaged(&data, size);

// CPU 访问
data[0] = 1.0f;

// GPU 访问（自动传输）
kernel<<<blocks, threads>>>(data);

// 释放
cudaFree(data);

利用张量核心

对于矩阵乘法等操作，使用张量核心可以大幅提升性能：

// 使用 cublas 库调用张量核心
cublasGemmEx(handle, 
             CUBLAS_OP_N, CUBLAS_OP_N,
             m, n, k,
             &alpha, A, CUDA_R_16F, lda,
                     B, CUDA_R_16F, ldb,
             &beta,  C, CUDA_R_16F, ldc,
             CUDA_R_32F, CUBLAS_GEMM_DEFAULT_TENSOR_OP);

张量核心使用条件：

使用 FP16、BF16 或 INT8 数据类型
矩阵维度最好是 8 或 16 的倍数
使用高度优化的库（cuBLAS、cuDNN）

性能分析工具

使用工具定位瓶颈：

# Nsight Systems：整体性能分析
nsys profile --stats=true ./my_program

# Nsight Compute：kernel 级别分析
ncu ./my_program

# nvprof：命令行性能分析（旧版）
nvprof ./my_program

常见性能指标：

指标	理想值	说明
SM 占用率	> 50%	活跃线程束占总容量的比例
内存带宽利用率	> 70%	实际带宽 / 峰值带宽
IPC	接近峰值	每周期执行的指令数
Warp 执行效率	> 90%	非分歧线程比例

常用工具

NVIDIA Nsight

Nsight 是 NVIDIA 提供的性能分析工具套件：

Nsight Systems：系统级性能分析，查看整体执行时间线
Nsight Compute：kernel 级性能分析，分析内存访问效率等

CUDA-MEMCHECK

检测内存错误：

cuda-memcheck ./my_program

nvprof

命令行性能分析工具：

nvprof ./my_program

多 GPU 编程

现代 AI 计算通常需要多个 GPU 协同工作。多 GPU 编程主要涉及数据并行和点对点通信两种模式。

数据并行

每个 GPU 持有完整模型，处理不同数据：

int deviceCount;
cudaGetDeviceCount(&deviceCount);

for (int i = 0; i < deviceCount; i++) {
    cudaSetDevice(i);
    // 在设备 i 上执行 kernel
    myKernel<<<blocks, threads>>>(...);
}

多流执行：

使用 CUDA 流可以在单个 GPU 上并发执行多个操作：

cudaStream_t streams[2];
cudaStreamCreate(&streams[0]);
cudaStreamCreate(&streams[1]);

// 并发执行两个 kernel
kernel1<<<blocks, threads, 0, streams[0]>>>(...);
kernel2<<<blocks, threads, 0, streams[1]>>>(...);

// 等待所有流完成
cudaDeviceSynchronize();

// 清理
cudaStreamDestroy(streams[0]);
cudaStreamDestroy(streams[1]);

点对点通信

GPU 之间直接通信：

// 启用 P2P 访问
cudaDeviceEnablePeerAccess(peerDevice, 0);

// 直接从另一个 GPU 复制数据
cudaMemcpyPeer(dst, dstDevice, src, srcDevice, size);

P2P 访问检查：

int canAccess;
cudaDeviceCanAccessPeer(&canAccess, device0, device1);
if (canAccess) {
    cudaDeviceEnablePeerAccess(device1, 0);
    // 现在 device0 可以直接访问 device1 的内存
}

NCCL 集合通信

对于分布式训练，NCCL（NVIDIA Collective Communications Library）提供了高效的集合通信原语：

#include <nccl.h>

// 初始化 NCCL
ncclComm_t comm;
ncclCommInitAll(&comm, nDevices, devices);

// All-Reduce：所有 GPU 同步数据
ncclAllReduce(sendbuff, recvbuff, count, ncclFloat, ncclSum, 
              comm, stream);

// All-Reduce 示例：梯度同步
ncclGroupStart();
for (int i = 0; i < nGpus; i++) {
    cudaSetDevice(i);
    ncclAllReduce(gradients[i], gradients[i], paramCount,
                  ncclFloat, ncclSum, comms[i], streams[i]);
}
ncclGroupEnd();

NCCL 支持的集合操作：

操作	说明	典型用途
All-Reduce	所有 GPU 得到相同结果	梯度同步
Reduce-Scatter	分片聚合	数据并行
All-Gather	收集所有数据	模型并行
Broadcast	一对多广播	参数初始化
Reduce	聚合到一个 GPU	汇总结果

实践建议

开发流程建议

先保证正确性，再优化性能：使用 CUDA-MEMCHECK 检查内存错误
使用性能分析工具：找到真正的瓶颈再优化
参考最佳实践：NVIDIA 提供了许多优化示例
利用现有库：cuBLAS、cuDNN 等库已经高度优化

常见问题与解决

问题	症状	解决方案
显存不足	cudaErrorMemoryAllocation	减小批次大小、使用梯度检查点
内存越界	未定义行为	使用 cuda-memcheck 检查
性能低下	GPU 利用率低	检查内存访问模式、增加并行度
同步错误	结果不一致	检查 __syncthreads() 位置
P2P 不工作	传输慢	检查 NVLink 连接、使用 NCCL

调试技巧

# 内存检查
cuda-memcheck ./my_program

# 详细内存检查
cuda-memcheck --tool memcheck --leak-check full ./my_program

# 性能分析
nsys profile --stats=true ./my_program

# Kernel 分析
ncu --set full ./my_program

学习路径

入门阶段：学习 CUDA 基本概念，理解线程层次和内存模型
进阶阶段：掌握共享内存优化、流和事件
高级阶段：深入理解架构特性、使用性能分析工具
实战阶段：使用 cuBLAS、cuDNN 等库构建实际应用

小结

GPU 是 AI 计算的核心硬件，理解其架构和编程模型对于构建高效的 AI 系统至关重要。本章介绍了：

GPU 架构：SM、内存层次、张量核心等核心组件
CUDA 编程模型：线程层次、Kernel 函数、内存管理
性能优化：内存访问优化、并行度最大化、张量核心利用
多 GPU 编程：数据并行、P2P 通信、NCCL 集合操作

在实际开发中，建议优先使用高度优化的库（如 cuBLAS、cuDNN），避免重复造轮子。当需要自定义 CUDA kernel 时，遵循先正确后优化的原则，善用性能分析工具定位瓶颈。

GPU 计算基础

GPU 与 CPU 的区别

一个形象的比喻

为什么 GPU 适合深度学习

GPU 架构详解

流式多处理器（SM）

内存层次结构

张量核心

CUDA 编程模型

线程层次结构

线程索引

Kernel 函数

内存管理

使用共享内存优化

性能优化原则

最大化并行度

优化内存访问

减少数据传输

利用张量核心

性能分析工具

常用工具

NVIDIA Nsight

CUDA-MEMCHECK

nvprof

多 GPU 编程

数据并行

点对点通信

NCCL 集合通信

实践建议

开发流程建议

常见问题与解决

调试技巧

学习路径

小结

参考资料

官方文档

学习资源

GPU 与 CPU 的区别​

一个形象的比喻​

为什么 GPU 适合深度学习​

GPU 架构详解​

流式多处理器（SM）​

内存层次结构​

张量核心​

CUDA 编程模型​

线程层次结构​

线程索引​

Kernel 函数​

内存管理​

使用共享内存优化​

性能优化原则​

最大化并行度​

优化内存访问​

减少数据传输​

利用张量核心​

性能分析工具​

常用工具​

NVIDIA Nsight​

CUDA-MEMCHECK​

nvprof​

多 GPU 编程​

数据并行​

点对点通信​

NCCL 集合通信​

实践建议​

开发流程建议​

常见问题与解决​

调试技巧​

学习路径​

小结​

参考资料​

官方文档​

学习资源​

GPU 与 CPU 的区别

一个形象的比喻

为什么 GPU 适合深度学习

GPU 架构详解

流式多处理器（SM）

内存层次结构

张量核心

CUDA 编程模型

线程层次结构

线程索引

Kernel 函数

内存管理

使用共享内存优化

性能优化原则

最大化并行度

优化内存访问

减少数据传输

利用张量核心

性能分析工具

常用工具

NVIDIA Nsight

CUDA-MEMCHECK

nvprof

多 GPU 编程

数据并行

点对点通信

NCCL 集合通信

实践建议

开发流程建议

常见问题与解决

调试技巧

学习路径

小结

参考资料

官方文档

学习资源