跳转至

GPU 架构与 CUDA

一句话总结: GPU 用数千核心的大规模并行改变了 AI，本节讲清 GPU 与 CPU 的设计哲学差异、显存层级、CUDA 编程模型、SIMT 执行、访存模式、同步、流、性能分析与 NVIDIA GPU 代际，让你能读懂并写出 GPU 内核。

GPU 用数千核心带来大规模并行，彻底改变了 AI. 本节涵盖 GPU 与 CPU 的设计哲学、GPU 显存层级、用 C++ 写 CUDA、SIMT 执行模型、访存模式、同步、流、性能分析，以及 NVIDIA GPU 各代际——这些是你写和读 GPU 内核所需的全部知识.

  • 想要可以直接跑起来的 CUDA 完整示例教程，见配套仓库：github.com/HenryNdubuaku/cuda-tutorials.

  • 一块现代 NVIDIA GPU (Graphics Processing Unit, GPU) 有超过 10,000 个 CUDA 核心。 一颗 CPU (Central Processing Unit, CPU) 只有 4-128 个核心。 这 100-1000 倍的核心数优势就是 GPU 在机器学习领域独霸一方的原因：训练一个 Transformer 要做数万亿次乘加运算，GPU 能并行处理的规模是 CPU 望尘莫及的。

  • 即使你自己永远不会写 CUDA 内核，理解 GPU 架构也能解释：为什么批大小 (batch size) 重要 (要有足够工作量把 GPU 喂饱)，为什么内存往往才是瓶颈 (而不是计算)，以及为什么某些操作 (散射、条件分支) 在 GPU 上慢。

GPU 与 CPU：根本不同的设计

  • CPU 是为延迟而设计的：让完成一个任务的时间最短。 它把大部分晶体管预算花在缓存、分支预测器、乱序执行上——所有这些都是为了让单线程更快。

  • GPU 是为吞吐而设计的：让每秒完成的任务数最大。 它把大部分晶体管留给执行单元 (算术逻辑单元，Arithmetic Logic Unit, ALU). 单个线程慢，但线程数以千计。

CPU GPU
核心数 4-128 (复杂，快) 1,000-20,000 (简单，慢)
时钟频率 3-5 GHz 1-2.5 GHz
缓存 大 (32 MB+ L3) 小 (每 SM 共享内存)
分支预测 精巧 无 (同一 warp 所有线程走同一路径)
擅长 低延迟，复杂控制流 高吞吐，数据并行
典型 FLOPS (FP32) 1-5 TFLOPS 30-80 TFLOPS
显存带宽 50-100 GB/s 1-3 TB/s
  • GPU 的显存带宽优势 (10-30 倍) 往往比它的计算优势更关键。 很多机器学习操作是内存受限的 (逐元素操作、归一化、注意力), GPU 的带宽让它有本事把数据以足够快的速度喂到核心里。

GPU 显存层级

  • 理解 GPU 显存是关键，因为内存访问才是主要瓶颈，不是计算。
存储 大小 延迟 带宽 作用域
寄存器 每 SM 约 256 KB 0 周期 最高 每线程
共享内存 每 SM 48-228 KB 约 5 周期 约 20 TB/s 每线程块
L1 缓存 每 SM 128-256 KB 约 30 周期 每 SM
L2 缓存 4-96 MB 约 200 周期 约 6 TB/s 全局
全局内存 (HBM) 24-192 GB 约 400 周期 1-3.3 TB/s 全局

  • 寄存器最快，但最稀缺。 每个线程有一组私有寄存器 (通常最多 255 个). 每线程用太多寄存器会降低占用率 (occupancy)——能同时跑的线程更少。

  • 共享内存是程序员可管理的缓存，由一个线程块内所有线程共用。 它是写快 CUDA 内核的关键：把一小块数据从慢的全局内存搬到快的共享内存里，然后在上面计算。 这就是主宰 GPU 编程的分块 (tiling) 模式。

  • 全局内存 (高带宽内存，High Bandwidth Memory, HBM): GPU 的主显存 (VRAM). 大但慢 (400 周期延迟). 所有数据的起点和终点都在这里。 内核优化的目标就是尽量减少对全局内存的访问。

CUDA 编程模型

  • CUDA (Compute Unified Device Architecture，统一计算设备架构) 是 NVIDIA 面向 GPU 的编程模型。 你写内核 (kernel)：跑在 GPU 上、由数千线程同时执行的函数。

层级：网格、块、线程

Grid (整次启动)
├── Block (0,0)
│   ├── Thread (0,0)
│   ├── Thread (1,0)
│   ├── Thread (2,0)
│   └── ... (每块最多 1024 个线程)
├── Block (1,0)
│   ├── Thread (0,0)
│   └── ...
└── ... (可以有上百万个块)
  • 线程 (Thread)：最小单位。 每个线程在所在块内有唯一 ID (threadIdx.x).
  • 块 (Block)：一组可以共享内存并互相同步的线程。 块 ID: blockIdx.x. 块大小：blockDim.x (最多 1024 个线程).

  • 网格 (Grid)：一次内核启动所有的块。 可以是 1D、2D 或 3D.

  • 每个线程算出自己的全局索引：int idx = blockIdx.x * blockDim.x + threadIdx.x;

你的第一个 CUDA 内核

// vector_add.cu — CUDA 源文件 (.cu 后缀)

#include <stdio.h>

// __global__ 标记这是一个 GPU 内核 (从 CPU 调用, 在 GPU 上运行)
__global__ void vector_add(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;  // 约 100 万个元素
    size_t bytes = n * sizeof(float);

    // 分配主机 (CPU) 内存
    float *h_a = new float[n];
    float *h_b = new float[n];
    float *h_c = new float[n];

    // 初始化
    for (int i = 0; i < n; i++) {
        h_a[i] = 1.0f;
        h_b[i] = 2.0f;
    }

    // 分配设备 (GPU) 内存
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);

    // 数据从 CPU 拷贝到 GPU
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

    // 启动内核: 每块 256 线程, 用足够多的块覆盖 n 个元素
    int block_size = 256;
    int grid_size = (n + block_size - 1) / block_size;  // 向上取整除
    vector_add<<<grid_size, block_size>>>(d_a, d_b, d_c, n);

    // 结果从 GPU 拷回 CPU
    cudaMemcpy(h_c, d_a, bytes, cudaMemcpyDeviceToHost);

    // 验证
    printf("c[0] = %f (expected 3.0)\n", h_c[0]);

    // 释放内存
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    delete[] h_a; delete[] h_b; delete[] h_c;

    return 0;
}

# 用 NVIDIA 编译器编译
nvcc -O3 -o vector_add vector_add.cu
./vector_add
  • CUDA 里的几个关键 C++ 概念:
    • __global__: CUDA 关键字，标记一个内核函数。 从 CPU (host) 调用，跑在 GPU (device) 上。
    • <<<grid_size, block_size>>>：内核启动语法。 指定用多少个块和线程。
    • cudaMalloc / cudaFree：分配/释放 GPU 内存 (类似 new/delete，但作用于 GPU).
    • cudaMemcpy：在 CPU 和 GPU 之间拷贝数据。 这常常是最大的瓶颈 (PCIe 带宽约 32 GB/s，而 GPU 显存带宽约 3 TB/s).

Warp 与 SIMT

  • GPU 以 32 个线程为一组来执行，这一组叫 warp. 一个 warp 内的 32 个线程在同一时刻执行同一条指令 (单指令多线程，Single Instruction Multiple Threads, SIMT). 这是 GPU 版的 SIMD，但作用在线程层级。

  • Warp 分歧 (warp divergence) 发生在同一 warp 内的线程在 if 语句上走了不同分支时。 GPU 一个 warp 里不能同时执行两条不同指令，于是它顺序执行两个分支，分别把不该参与的线程屏蔽掉。 这会让性能腰斩 (甚至更糟).

// 差的写法: warp 分歧 (同一 warp 内线程走不同路径)
if (threadIdx.x % 2 == 0) {
    c[idx] = a[idx] + b[idx];    // 偶数线程走这里
} else {
    c[idx] = a[idx] - b[idx];    // 奇数线程走这里 (同 warp, 被串行化)
}

// 更好的写法: 无分支 (无分歧)
float sign = (threadIdx.x % 2 == 0) ? 1.0f : -1.0f;
c[idx] = a[idx] + sign * b[idx];  // 所有线程执行同一条指令

访存合并

  • 合并访存 (coalesced access)：当连续的线程访问连续的内存地址时，GPU 会把它们合并成一次内存事务。 这对性能至关重要。
// 好: 合并访存 — 线程 0 读 a[0], 线程 1 读 a[1], ...
c[idx] = a[idx] + b[idx];

// 差: 跨步访存 — 线程 0 读 a[0], 线程 1 读 a[stride], ...
c[idx] = a[idx * stride] + b[idx * stride];  // stride > 1 浪费带宽
  • 对于一个 32 线程的 warp，合并访存在一次事务里加载 128 字节 (32 × 4 字节 float32). 跨步访存则需要多次事务，每次同样加载 128 字节，却只用其中一小部分。 步长为 32 是最糟情形：每次事务加载 128 字节，却只有一个线程用到 4 字节 (3% 利用率).

共享内存与分块

  • 分块模式 (tiling pattern) 是最重要的 GPU 优化技术。 思路：把一小块数据从慢的全局内存搬到快的共享内存里，在上面算完，再把结果写回去。
// 用共享内存做分块的矩阵乘 (简化版)
__global__ void matmul_tiled(const float* A, const float* B, float* C,
                              int M, int N, int K) {
    // 为 A 和 B 各准备一小块的共享内存
    __shared__ float tile_A[TILE_SIZE][TILE_SIZE];
    __shared__ float tile_B[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 < (K + TILE_SIZE - 1) / TILE_SIZE; t++) {
        // 把 A 和 B 各一块载入共享内存
        if (row < M && t * TILE_SIZE + threadIdx.x < K)
            tile_A[threadIdx.y][threadIdx.x] = A[row * K + t * TILE_SIZE + threadIdx.x];
        else
            tile_A[threadIdx.y][threadIdx.x] = 0.0f;

        if (col < N && t * TILE_SIZE + threadIdx.y < K)
            tile_B[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * N + col];
        else
            tile_B[threadIdx.y][threadIdx.x] = 0.0f;

        __syncthreads();  // 等所有线程加载完

        // 用这一块算局部点积
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += tile_A[threadIdx.y][k] * tile_B[k][threadIdx.x];
        }

        __syncthreads();  // 加载下一块前等一下
    }

    if (row < M && col < N)
        C[row * N + col] = sum;
}
  • __shared__：声明块内所有线程共享的内存 (快，位于片上).
  • __syncthreads()：一个屏障，等到块内所有线程都到这里。 写共享内存和读共享内存之间必须放一个 (否则有些线程会读到旧数据).
  • 分块为什么有效：不分块时，每个线程每次乘法都要从全局内存加载一次。 分块后，一个 TILE_SIZE × TILE_SIZE 的块只加载一次到共享内存，被块内所有线程复用。 复用倍数就是 TILE_SIZE，全局内存流量因此减小到原来的 1/TILE_SIZE.

流与并发

  • 默认情况下，CUDA 操作是串行的：CPU 启动一个内核，等它完成，再启动下一个。 流 (streams) 让操作可以重叠:
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// 下面这些操作可以重叠: 不同流并发执行
cudaMemcpyAsync(d_a, h_a, bytes, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(d_b, h_b, bytes, cudaMemcpyHostToDevice, stream2);

kernel1<<<grid, block, 0, stream1>>>(d_a, d_c);
kernel2<<<grid, block, 0, stream2>>>(d_b, d_d);
  • 流让数据传输与计算重叠：一个流的内核在跑时，另一个流在拷数据。 这样藏起 PCIe 传输的延迟，让 GPU 始终有活干。

CUDA 代码性能分析

# NVIDIA Nsight Compute: 内核级性能分析
ncu --set full ./my_program

# NVIDIA Nsight Systems: 系统级时间线
nsys profile ./my_program

# 快速看几个指标
ncu --metrics sm__throughput,dram__throughput ./my_program
  • 重点看什么:
    • 占用率 (Occupancy): SM (Streaming Multiprocessor，流式多处理器) 容量被利用的比例。 占用率低 (< 50%) 意味着线程太少，不足以隐藏内存延迟。 常见原因：每线程用太多寄存器，每块用太多共享内存。
    • 内存吞吐：跟峰值带宽比。 如果只达到峰值的 50% 以下，就说明访存模式不高效 (未合并、bank 冲突).
    • 计算吞吐：跟峰值 FLOPS 比。 如果内存吞吐和计算吞吐都低，内核就是延迟受限 (并行度不够).

进阶优化技术

  • 除了合并访存和共享内存分块这些基础，高性能 GPU (和 CPU) 代码还用一些进阶技巧:

数据布局：AoS 与 SoA

  • 结构体数组 (Array of Structures, AoS)：每个元素把所有字段挨在一起存。 [{x,y,z}, {x,y,z}, {x,y,z}].
  • 数组的结构体 (Structure of Arrays, SoA)：每个字段各存一个连续数组。 {[x,x,x], [y,y,y], [z,z,z]}.
// AoS: 对 SIMD/GPU 不友好 (访问所有 x 会碰不连续的内存)
struct Particle { float x, y, z, mass; };
Particle particles[N];
// particles[0].x 与 particles[1].x 相隔 16 字节

// SoA: 对 SIMD/GPU 友好 (所有 x 值连续)
struct Particles {
    float x[N], y[N], z[N], mass[N];
};
// x[0] 与 x[1] 相隔 4 字节 — 完美的合并访存与 SIMD 布局
  • 对数据并行的工作负载 (SIMD、GPU), SoA 几乎总是更快。 AoS 只在你总是把一个元素的所有字段一起访问时更好 (数值代码里少见). PyTorch 的张量本质就是 SoA：每个特征是一个连续的维度。

软件预取

  • 可以让 CPU 提前开始加载还没用到的数据，藏起内存延迟:
#include <xmmintrin.h>  // 为了 _mm_prefetch

for (int i = 0; i < n; i += 4) {
    _mm_prefetch((char*)(a + i + 64), _MM_HINT_T0);  // 提前 64 个元素预取
    // 用 SIMD 处理 a[i:i+4]
    __m128 va = _mm_load_ps(a + i);
    // ...
}
  • 预取指令只是提示：数据已在缓存中就是空操作；不在的话，CPU 会在后台开始加载，同时执行别的指令。 预取距离 (本例是提前 64 个元素) 要根据内存延迟和循环迭代时间来调。

算子融合

  • 算子融合 (kernel fusion) 把多个操作合并到一个内核里，避免把中间结果写回内存。 这是机器学习里最有威力的 GPU 优化:
// 未融合: 3 次内核启动, 3 次全局内存往返
y = matmul(x, W)     // 把 y 写入全局内存
z = y + bias          // 读 y, 写 z
out = relu(z)         // 读 z, 写 out

// 融合: 1 次内核启动, 1 次全局内存写入
out = fused_matmul_bias_relu(x, W, bias)  // y 和 z 从未离开 SRAM
  • 对内存受限的操作 (加偏置、ReLU、层归一化)，内存流量主宰执行时间。 融合把这些流量直接消掉。 PyTorch 的 torch.compile 和 Triton 让融合可以自动化，或几乎不费力地做到。

混合精度内核

  • 计算用低精度 (FP16、BF16、INT8)，累加用高精度 (FP32)，两全其美:
// Tensor Core: FP16 矩阵相乘, FP32 累加
// 每条 Tensor Core 指令: D (FP32) = A (FP16) × B (FP16) + C (FP32)
nvcuda::wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);
  • FP16 比 FP32 小一半，因此内存带宽 (通常的瓶颈) 翻倍，缓存里能塞下的数据也翻倍。 Tensor Core 处理 FP16 的速度是常规 CUDA 核心处理 FP32 的 8-16 倍。 这就是为什么混合精度训练 (第 6 章) 能带来 2-3 倍加速，而精度损失极小。

内存池分配器

  • cudaMalloc 慢 (每次调用约 1 ms)，因为它要跟 GPU 同步。 训练循环里每次迭代都分配临时缓冲区的话，这时间会累起来。

  • 内存池 (memory pool) (PyTorch 的缓存分配器、CUDA 内存池) 预先分配一大块 GPU 内存，然后从中做子分配，不走系统调用:

# PyTorch 自动做这件事 — 但明白背后原理很重要
# 每次 torch.empty() 复用内存池里的内存, 不调 cudaMalloc
temp = torch.empty(1024, 1024, device='cuda')  # 微秒级, 不是毫秒级
  • 这就是为什么 PyTorch 的 torch.cuda.memory_allocated() 和 torch.cuda.max_memory_allocated() 不一样：allocated 是当前在用的，max 是峰值 (内存池可能持有比当前更多的内存).

基于性能分析的优化

  • 别盲目优化。 先分析，定位瓶颈，针对性优化，再重新分析。 Roofline 模型 (01 节) 告诉你瓶颈在内存还是计算:

    • 内存受限 (算术强度低)：优化数据布局 (SoA)，融合内核，使用低精度，预取。
    • 计算受限 (算术强度高)：使用 Tensor Core，提高并行度，使用更快的指令 (融合乘加，Fused Multiply-Add, FMA).
    • 延迟受限 (并行度不足)：提高占用率，减少寄存器使用，启动更多线程。

  • 大多数机器学习工作负载是内存受限的. 意外的推论：更快的 GPU (更多 FLOPS) 常常帮不上忙。 更快的显存 (HBM3 vs HBM2e) 更有效。 这也是为什么 A100→H100 的升级不只在 FLOPS——H100 的显存带宽也翻了一番。

NVIDIA GPU 各代际

代际 年份 关键创新 AI 相关意义
Pascal (P100) 2016 HBM2, NVLink 第一代真正用于深度学习的 GPU
Volta (V100) 2017 Tensor Core (混合精度矩阵乘) 让 FP16 训练可行，125 TFLOPS FP16 (据 NVIDIA V100 datasheet; TF32 是 Ampere 才引入)
Ampere (A100) 2020 TF32，稀疏，第 3 代 Tensor Core 312 TFLOPS TF32，结构化 2:4 稀疏
Hopper (H100) 2022 Transformer Engine (FP8), HBM3 989 TFLOPS FP8，动态精度切换
Blackwell (B200) 2024 第 2 代 Transformer Engine, NVLink 5 2.5 PFLOPS FP4，多芯粒设计

  • Tensor Core 是专门的矩阵乘单元。 一条 Tensor Core 指令在一个周期内计算一次 4×4 矩阵乘法 (D = A×B + C). 普通 CUDA 核心要做 64 次 FMA 运算。 Tensor Core 就是混合精度训练 (FP16 计算、FP32 累加) 快的原因。

  • Transformer Engine (Hopper 及以后) 在一层之内动态在 FP8 和 FP16 精度之间切换，只在需要更高精度的地方用高精度。 这样在不牺牲模型质量的前提下把吞吐做到最大。 它专门为 Transformer 架构 (注意力 + 多层感知机，Multi-Layer Perceptron, MLP) 设计，而这正是现代 AI 的主流。

编码任务 (用 nvcc 编译)

  1. 写一个 CUDA 内核，对一个数组做 ReLU. 测量包含内存传输在内的耗时。 这个任务能让你学会写内核、cudaMalloc/cudaMemcpy，以及主机↔设备传输的瓶颈。 

    // task1_relu.cu
// 编译: nvcc -O3 -o task1_relu task1_relu.cu

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

__global__ void relu_kernel(const float* input, float* output, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        output[idx] = input[idx] > 0.0f ? input[idx] : 0.0f;
    }
}

int main() {
    const int N = 1 << 24;  // 约 1600 万个元素
    size_t bytes = N * sizeof(float);

    // 分配主机内存
    float* h_input  = (float*)malloc(bytes);
    float* h_output = (float*)malloc(bytes);
    for (int i = 0; i < N; i++) {
        h_input[i] = (float)(i % 100) - 50.0f;  // 正负混合
    }

    // 分配设备内存
    float *d_input, *d_output;
    cudaMalloc(&d_input, bytes);
    cudaMalloc(&d_output, bytes);

    // 对整条流水线计时: 拷到 GPU、计算、拷回
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    cudaEventRecord(start);
    cudaMemcpy(d_input, h_input, bytes, cudaMemcpyHostToDevice);

    int block_size = 256;
    int grid_size = (N + block_size - 1) / block_size;
    relu_kernel<<<grid_size, block_size>>>(d_input, d_output, N);

    cudaMemcpy(h_output, d_output, bytes, cudaMemcpyDeviceToHost);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);

    float ms = 0;
    cudaEventElapsedTime(&ms, start, stop);

    // 验证
    int errors = 0;
    for (int i = 0; i < N; i++) {
        float expected = h_input[i] > 0.0f ? h_input[i] : 0.0f;
        if (h_output[i] != expected) errors++;
    }

    printf("Time (including transfers): %.2f ms\n", ms);
    printf("Bandwidth: %.1f GB/s\n", 2.0 * bytes / ms / 1e6);  // 读 + 写
    printf("Errors: %d / %d\n", errors, N);

    cudaFree(d_input); cudaFree(d_output);
    free(h_input); free(h_output);
    return 0;
}

  2. 用共享内存写一个分块的 CUDA 矩阵乘。 与朴素 (未分块) 版本比性能。 这个任务教你共享内存、__syncthreads，以及为什么分块重要。 

    // task2_matmul.cu
// 编译: nvcc -O3 -o task2_matmul task2_matmul.cu

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

#define TILE 16

// 朴素矩阵乘: 每个线程算 C 的一个元素
__global__ void matmul_naive(const float* A, const float* B, float* C, int N) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < N && col < N) {
        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__ void matmul_tiled(const float* A, const float* B, float* C, int N) {
    __shared__ float sA[TILE][TILE];
    __shared__ float sB[TILE][TILE];

    int row = blockIdx.y * TILE + threadIdx.y;
    int col = blockIdx.x * TILE + threadIdx.x;
    float sum = 0.0f;

    for (int t = 0; t < (N + TILE - 1) / TILE; t++) {
        sA[threadIdx.y][threadIdx.x] = (row < N && t*TILE+threadIdx.x < N)
            ? A[row * N + t*TILE + threadIdx.x] : 0.0f;
        sB[threadIdx.y][threadIdx.x] = (t*TILE+threadIdx.y < N && col < N)
            ? B[(t*TILE + threadIdx.y) * N + col] : 0.0f;

        __syncthreads();
        for (int k = 0; k < TILE; k++)
            sum += sA[threadIdx.y][k] * sB[k][threadIdx.x];
        __syncthreads();
    }

    if (row < N && col < N)
        C[row * N + col] = sum;
}

int main() {
    const int N = 1024;
    size_t bytes = N * N * sizeof(float);

    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, bytes); cudaMalloc(&d_B, bytes); cudaMalloc(&d_C, bytes);

    // 全部初始化为 1 (易于验证: C 每个元素应为 N)
    float* h_A = new float[N*N];
    for (int i = 0; i < N*N; i++) h_A[i] = 1.0f;
    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_A, bytes, cudaMemcpyHostToDevice);

    dim3 block(TILE, TILE);
    dim3 grid((N+TILE-1)/TILE, (N+TILE-1)/TILE);

    // 朴素版基准
    cudaEvent_t start, stop;
    cudaEventCreate(&start); cudaEventCreate(&stop);

    cudaEventRecord(start);
    for (int i = 0; i < 10; i++)
        matmul_naive<<<grid, block>>>(d_A, d_B, d_C, N);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    float naive_ms; cudaEventElapsedTime(&naive_ms, start, stop);

    // 分块版基准
    cudaEventRecord(start);
    for (int i = 0; i < 10; i++)
        matmul_tiled<<<grid, block>>>(d_A, d_B, d_C, N);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    float tiled_ms; cudaEventElapsedTime(&tiled_ms, start, stop);

    double gflops_naive = 2.0 * N * N * N * 10 / naive_ms / 1e6;
    double gflops_tiled = 2.0 * N * N * N * 10 / tiled_ms / 1e6;

    printf("Naive:  %.2f ms, %.1f GFLOPS\n", naive_ms/10, gflops_naive);
    printf("Tiled:  %.2f ms, %.1f GFLOPS\n", tiled_ms/10, gflops_tiled);
    printf("Speedup: %.1fx\n", naive_ms / tiled_ms);

    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    delete[] h_A;
    return 0;
}

  3. 演示 warp 分歧。 写一个同一 warp 内线程走不同分支的内核，再和无分支版本比。 

    // task3_divergence.cu
// 编译: nvcc -O3 -o task3_diverge task3_divergence.cu

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

// 差: warp 分歧 — 偶数/奇数线程走不同路径
__global__ void divergent_kernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        if (idx % 2 == 0) {
            data[idx] = data[idx] * 2.0f + 1.0f;
        } else {
            data[idx] = data[idx] * 0.5f - 1.0f;
        }
    }
}

// 好: 无分支 — 所有线程执行同一条指令
__global__ void branchless_kernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float scale = (idx % 2 == 0) ? 2.0f : 0.5f;
        float offset = (idx % 2 == 0) ? 1.0f : -1.0f;
        data[idx] = data[idx] * scale + offset;
    }
}

int main() {
    const int N = 1 << 24;
    float* d_data;
    cudaMalloc(&d_data, N * sizeof(float));
    cudaMemset(d_data, 0, N * sizeof(float));

    int block = 256, grid = (N + block - 1) / block;

    cudaEvent_t start, stop;
    cudaEventCreate(&start); cudaEventCreate(&stop);

    // 分歧版
    cudaEventRecord(start);
    for (int i = 0; i < 100; i++)
        divergent_kernel<<<grid, block>>>(d_data, N);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    float div_ms; cudaEventElapsedTime(&div_ms, start, stop);

    // 无分支版
    cudaEventRecord(start);
    for (int i = 0; i < 100; i++)
        branchless_kernel<<<grid, block>>>(d_data, N);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    float nodiv_ms; cudaEventElapsedTime(&nodiv_ms, start, stop);

    printf("Divergent:  %.2f ms\n", div_ms / 100);
    printf("Branchless: %.2f ms\n", nodiv_ms / 100);
    printf("Speedup:    %.2fx\n", div_ms / nodiv_ms);

    cudaFree(d_data);
    return 0;
}

📌 本节要点

  • CPU vs GPU: CPU 为延迟优化 (少而强的核心), GPU 为吞吐优化 (数千个简单核心 + 10-30 倍显存带宽).
  • 显存层级：寄存器 → 共享内存 → L1/L2 → 全局 HBM，延迟差 400 倍，优化的目标是尽量减少全局内存访问。
  • CUDA 三级层次: Thread → Block → Grid；内核用 <<<grid, block>>> 启动，每线程通过 blockIdx * blockDim + threadIdx 算全局索引。
  • Warp 与 SIMT: GPU 以 32 线程为一组同步执行同一指令，同 warp 内走不同分支会被串行化 (warp 分歧).
  • 访存合并：连续线程访问连续地址合并成一次事务，跨步访问会浪费大部分带宽。
  • 共享内存分块：把数据从全局内存搬入片上共享内存，复用倍数等于 tile 大小，显著降低访存流量。
  • 流并发：用多个 CUDA 流让传输与计算重叠，藏起 PCIe 延迟。
  • 进阶优化：数据用 SoA 布局、软件预取、算子融合、混合精度 + Tensor Core、内存池。
  • 多数负载内存受限：先分析后优化；HBM 带宽提升往往比 FLOPS 提升更有价值。
  • NVIDIA 代际主线: Pascal → Volta (Tensor Core) → Ampere (TF32/稀疏) → Hopper (FP8/Transformer Engine) → Blackwell (FP4/多芯粒)，每代都在为 Transformer 加速。