📅 创建时间：2026-06-03 🏷️ 标签：#Kernel #Roofline #Occupancy #ArithmeticIntensity #内存带宽 #计算密度 #性能优化 📚 前置知识：[[/04-ai/01-llm-engineering/07-llm-evolution]]（LLM 发展脉络） 📚 相关知识：[[15-backend-cuda]]（CUDA 后端）[[14-backend-cpu]]（CPU 后端）[[18-cutlass]]（CUTLASS）[[19-tvm-te]]（TVM TE）[[20-triton]]（Triton）

Kernel 性能基础：Roofline 与 Occupancy / Kernel Performance Fundamentals with Roofline and Occupancy

┌──────────────────────────────────────────────────────────────────────────────┐ │ 📖 场景：你写了一个矩阵乘法 kernel │ ├──────────────────────────────────────────────────────────────────────────────┤ │ 你写了一个矩阵乘法 kernel，跑 benchmark 比 cuBLAS 慢 10 倍。 │ │ 你知道逻辑是对的，问题在硬件层面。但你不知道具体哪里出了问题。 │ └──────────────────────────────────────────────────────────────────────────────┘

第1节 Roofline Model——理解计算上限

1.1 什么是 Roofline Model

Roofline Model 是一个性能分析可视化工具，它告诉你一个 kernel 的理论上限取决于什么。在 GPU 上，内存带宽和计算峰值是两个硬上限——你的 kernel 性能不会超过其中任何一个。

         A100 FP32 Peak: 19.5 TFLOPS
        /
       /  ← 计算峰值上限 (Compute-Bound)
      /
     /
    /
   /
  / ← 斜率 = 内存带宽 / 计算峰值
 /
/
+--------------------------------------------------------→ 算术强度 (Flops/Byte)
   ↑
   内存带宽上限 (Memory-Bound)
   
   A100 HBM Bandwidth = 2039 GB/s

关键公式：

算术强度 (Arithmetic Intensity) = Flops / Bytes — 每次内存访问做多少计算
计算峰值上限 = min(Peak FLOPS, Bandwidth × Arithmetic Intensity)
** Roofline 曲线** = min(AithmeticIntensity × Bandwidth, PeakFLOPS)

1.2 两种瓶颈类型

瓶颈类型	英文	特征	优化方向
内存瓶颈	Memory-Bound	算术强度低，等待内存	减少内存访问、增加复用
计算瓶颈	Compute-Bound	算术强度高，CPU/ALU 饱和	提升计算利用率、增加并行度

1.3 A100 具体数据

硬件参数 (A100 SXM):
  - FP32 Peak FLOPS:    19.5 TFLOPS    (单精度计算上限)
  - FP64 Peak FLOPS:    9.7 TFLOPS    (双精度)
  - Tensor FP16 Peak:   312 TFLOPS    (Tensor Core)
  - HBM Bandwidth:      2039 GB/s      (显存带宽)
  - L2 Cache Bandwidth: ~2.5 TB/s     (约 1.2x HBM)
  - L1 Cache Bandwidth: ~13 TB/s       (约 6x HBM)

算术强度临界点 (FP32):
  - 临界算术强度 = PeakFLOPS / Bandwidth
                   = 19.5e12 / 2039e9
                   ≈ 9.56 Flops/Byte
                   
  - 如果 Flops/Byte < 9.56  → Memory-Bound
  - 如果 Flops/Byte > 9.56  → Compute-Bound

1.4 用 Roofline 诊断你的 Matmul Kernel

回头看你的 naive matmul kernel：

cuda

__global__ void matmul(float* C, float* A, float* B, int M, int N, int K) {
    // 每次外层循环迭代:
    //   读取: A[row*K+k]  → 1 float = 4 bytes
    //   读取: B[k*N+col]  → 1 float = 4 bytes
    //   计算: 1次乘 + 1次加 = 2 Flops
    // 
    // 总共读取: 2 * K * 4 bytes per element
    // 总共计算: 2 * K Flops per element
    //
    // 算术强度 = 2K / (2K * 4) = 2/4 = 0.5 Flops/Byte
    // 远低于 A100 临界点 9.56 → 严重 Memory-Bound！
    
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    int col = blockIdx.y * blockDim.y + threadIdx.y;
    if (row < M && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < K; k++)
            sum += A[row * K + k] * B[k * N + col];  // 每步读 8 bytes，算 2 Flops
        C[row * N + col] = sum;
    }
}

问题诊断：

每个线程每次迭代只处理 1 个输出元素，但每个元素需要读取 K 行 + K 列的数据
K=1024 时，每个线程要读 2048 个 float = 8KB 数据，但只做了 2048 次乘加
算术强度 = 2K / (2K × 4) = 0.5 Flops/Byte，比临界点低 19 倍
GPU 大部分时间在等待内存，而不是计算

1.5 Roofline 分析 Python 脚本

python

#!/usr/bin/env python3
"""
Roofline Model 分析工具
分析给定 kernel 的算术强度，判断瓶颈类型
"""

# ======================== 硬件参数 ========================
# A100 SXM4 80GB
HARDWARE = {
    "A100": {
        "fp32_peak_tflops": 19.5,        # TFLOPS
        "fp16_tensor_tflops": 312.0,      # TFLOPS (Tensor Core)
        "hbm_bandwidth_gbps": 2039.0,    # GB/s
        "l2_bandwidth_gbps": 2500.0,     # GB/s (估算)
        "l1_bandwidth_gbps": 13000.0,    # GB/s (估算)
    },
    "H100": {
        "fp32_peak_tflops": 67.0,        # TFLOPS
        "fp8_tensor_tflops": 3958.0,      # TFLOPS (Hopper FP8)
        "hbm_bandwidth_gbps": 3350.0,    # GB/s (H100 SXM5)
        "hbm3_bandwidth_gbps": 5100.0,   # GB/s (H200)
    },
}

def compute_arithmetic_intensity(flops_per_element: float, 
                                  bytes_per_element: float) -> float:
    """
    计算算术强度：Flops/Bytes
    
    Args:
        flops_per_element: 每个输出元素的浮点运算次数
        bytes_per_element:  每个输出元素需要读取的总字节数
    Returns:
        算术强度 (Flops/Byte)
    """
    return flops_per_element / bytes_per_element

def roofline_attainment(arithmetic_intensity: float,
                        hardware: dict,
                        dtype: str = "fp32") -> dict:
    """
    计算 Roofline 上限和实际可达性能
    
    Args:
        arithmetic_intensity: 算术强度 (Flops/Byte)
        hardware: 硬件参数字典
        dtype: 数据类型 ("fp32", "fp16", "fp8")
    Returns:
        包含分析结果的字典
    """
    # 根据 dtype 选择峰值 FLOPS
    if dtype == "fp32":
        peak_tflops = hardware["fp32_peak_tflops"]
    elif dtype == "fp16":
        peak_tflops = hardware.get("fp16_tensor_tflops", hardware["fp32_peak_tflops"] * 2)
    elif dtype == "fp8":
        peak_tflops = hardware.get("fp8_tensor_tflops", 0)
    else:
        peak_tflops = hardware["fp32_peak_tflops"]
    
    bandwidth = hardware["hbm_bandwidth_gbps"]  # GB/s
    peak_gflops = peak_tflops * 1000            # GFLOPS
    
    # Roofline 临界点: PeakFLOPS / Bandwidth
    # 在这个点上，内存带宽和计算峰值相等
    # 低于这个点 → memory-bound；高于这个点 → compute-bound
    critical_ai = peak_gflops / bandwidth       # Flops/Byte
    memory_bound_roof = bandwidth * arithmetic_intensity  # GFLOPS
    compute_bound_roof = peak_gflops            # GFLOPS
    
    # 实际可达性能上限
    attainable = min(memory_bound_roof, compute_bound_roof)
    
    # 判断瓶颈类型
    if arithmetic_intensity < critical_ai:
        bottleneck = "Memory-Bound"
        theoretical_tflops = memory_bound_roof / 1000  # TFLOPS
    else:
        bottleneck = "Compute-Bound"
        theoretical_tflops = peak_tflops
    
    # 计算利用率（相对于峰值）
    # 实际性能 / Roofline 上限 = 利用率百分比
    # 这个比例告诉你 kernel 有多少优化空间
    roofline_limit = min(bandwidth * arithmetic_intensity, peak_gflops)
    roofline_utilization = min(arithmetic_intensity / critical_ai, 1.0) if critical_ai > 0 else 0
    
    return {
        "arithmetic_intensity": arithmetic_intensity,
        "critical_ai": critical_ai,
        "bottleneck": bottleneck,
        "peak_tflops": peak_tflops,
        "memory_bandwidth_gbps": bandwidth,
        "roofline_limit_gflops": roofline_limit,
        "attainable_tflops": theoretical_tflops,
        "roofline_utilization": roofline_utilization,
        "efficiency": (theoretical_tflops / peak_tflops * 100) if peak_tflops > 0 else 0,
    }

def analyze_naive_matmul(M, N, K, dtype="fp32"):
    """分析 naive matmul kernel 的算术强度"""
    # 每个输出元素 C[i,j] = sum_k A[i,k] * B[k,j]
    # 读取: A 需要 M*K 个 float，B 需要 K*N 个 float
    #       分摊到每个输出: (M*K + K*N) / (M*N) = K*(M+N)/(M*N) ≈ 2K/MN * M*N = 2K bytes
    # 等等，这个算法每个线程读一行 A 和一列 B
    # 实际每个输出元素读取: 2*K float = 8K bytes
    bytes_per_output = 2 * K * 4  # K 行 A + K 列 B，每 float 4 bytes
    flops_per_output = 2 * K      # K 次乘 + K 次加 = 2K Flops
    
    ai = compute_arithmetic_intensity(flops_per_output, bytes_per_output)
    print(f"\n{'='*60}")
    print(f"Naive Matmul Analysis: M={M}, N={N}, K={K}, dtype={dtype}")
    print(f"{'='*60}")
    print(f"  每元素读取字节数: {bytes_per_output} bytes ({2*K} floats)")
    print(f"  每元素浮点运算数: {flops_per_output} Flops ({K}乘 + {K}加)")
    print(f"  算术强度: {ai:.2f} Flops/Byte")
    
    result = roofline_attainment(ai, HARDWARE["A100"], dtype)
    
    print(f"\n  【瓶颈诊断】")
    print(f"  瓶颈类型: {result['bottleneck']}")
    print(f"  Roofline 临界点: {result['critical_ai']:.2f} Flops/Byte")
    print(f"  理论上限: {result['attainable_tflops']:.2f} TFLOPS")
    print(f"  计算峰值: {result['peak_tflops']:.2f} TFLOPS")
    print(f"  理论效率: {result['efficiency']:.1f}%")
    
    if result['bottleneck'] == "Memory-Bound":
        print(f"\n  ⚠️  Memory-Bound: 大部分时间在等内存")
        print(f"     优化方向: 增加数据复用 (tiling)、使用 shared memory")
        print(f"     理论加速: {result['peak_tflops'] / result['attainable_tflops']:.1f}x (如果变成 Compute-Bound)")
    
    return result

# 分析不同 K 值
if __name__ == "__main__":
    print("=" * 60)
    print("Roofline Model 分析 - A100 GPU")
    print("=" * 60)
    
    # K 值越小，算术强度越低（每个元素读 2K floats，但 K 小）
    # 注意：算术强度 = 2K / (8K) = 0.25，和 K 无关！
    # 真正的问题是：每个线程独立读一行 A 和一列 B，没有数据复用
    
    # 但如果我们用 tiling，每次把一个 tile 加载到 shared memory:
    # tile 大小 TILE_M × TILE_K
    # 每个 tile 加载一次 A[TILE_M, TILE_K]，供 TILE_M 个线程复用
    # 每个元素读取: TILE_K * 4 bytes (A) + TILE_K * 4 bytes (B) / TILE_M 线程分摊
    # 实际上 tiling 后算术强度 = (2 * TILE_K * TILE_M) / (TILE_K * 4 + TILE_K * 4) = TILE_M / 4
    
    TILE_M = 64  # 假设每个 tile 64 行
    print(f"\n使用 TILE_M={TILE_M} 的 tiling 后:")
    print(f"  每个输出元素的算术强度提升到: {TILE_M / 4:.1f} Flops/Byte")
    
    result_tiled = roofline_attainment(TILE_M / 4, HARDWARE["A100"], "fp32")
    print(f"  瓶颈类型: {result_tiled['bottleneck']}")
    print(f"  理论上限: {result_tiled['attainable_tflops']:.2f} TFLOPS")
    print(f"  理论效率: {result_tiled['efficiency']:.1f}%")
    
    # 当 TILE_M >= 38 时 (临界算术强度 ~9.56)，算术强度 > 9.56
    # kernel 从 memory-bound 变为 compute-bound
    print(f"\n{'='*60}")
    print("关键发现:")
    print("  - Naive kernel: 算术强度 0.25 Flops/Byte，效率 <3%")
    print("  - Tiling 后: 算术强度 = TILE_M/4，提升到 16+ 时接近峰值")
    print("  - 优化空间: ~7x 加速 (从 2.7 TFLOPS 到 19.5 TFLOPS)")

第2节 GPU Occupancy——每 SM 活跃线程数

2.1 什么是 Occupancy

GPU Occupancy（占用率） = 实际使用的线程数 / 最大可支持线程数。它衡量 GPU SM（流多处理器）的利用程度。

A100 SM 资源限制:
┌─────────────────────────────────────────────────────────┐
│  每个 SM 最多 2048 线程 (128 threads/block × 16 blocks)  │
│  每个 SM 最多 32 个 block                               │
│  每个 SM 最多 64 warps (32 threads/warp × 64 warps)     │
│  每个 SM 最多 16 个 CTA (cooperative thread array)      │
│                                                         │
│  Register 限制: 每 SM 最多 65536 registers              │
│  Shared Memory 限制: 每 SM 最多 164 KB                  │
└─────────────────────────────────────────────────────────┘

2.2 寄存器压力和 Shared Memory 压力

Occupancy 下降的两个主要原因：

压力类型	原因	影响
Register 压力	每个线程使用过多寄存器	SM 能放的线程数减少
Shared Memory 压力	每个 block 使用过多 shared memory	SM 能放的 block 数减少

Occupancy 计算公式：

Occupancy = min(
    ThreadsPerSM / MaxThreadsPerSM,          # 线程数限制
    RegistersPerSM / (RegistersPerThread × ThreadsPerSM),  # 寄存器限制
    SharedMemPerSM / (SharedMemPerBlock × BlocksPerSM)     # Shared Memory 限制
)

2.3 计算 Naive Matmul 的 Occupancy

cuda

__global__ void matmul(float* C, float* A, float* B, int M, int N, int K) {
    // 分析每个线程的寄存器使用:
    // - row, col, k: int 类型 → 2 registers (或 1)
    // - sum: float → 1 register
    // - A[row*K+k], B[k*N+col]: 寄存器间接寻址，不占寄存器
    // 总计约 3-4 registers per thread
    
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    int col = blockIdx.y * blockDim.y + threadIdx.y;
    if (row < M && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < K; k++)
            sum += A[row * K + k] * B[k * N + col];
        C[row * N + col] = sum;
    }
}

用 nvcc --ptxas-info 分析：

ptxas info: 0 bytes gmem, 0 bytes cmem[3]
ptxas info: Compiling entry function 'matmul' for sm_80
ptxas info: Function properties for matmul
    0 bytes stack frame, 0 bytes spilled stores, 0 bytes spilled loads
ptxas info: Used 4 registers, 0 bytes smem, 0 bytes cmem

每个线程 4 registers:
  - A100 每 SM 65536 registers / 4 = 16384 threads per SM
  - 实际限制: 16384 < 2048 (最大值) → Occupancy 由寄存器限制
  - 但 16384/2048 = 100% → 寄存器够用
  
Block 大小 256 (8×32):
  - 2048 threads / SM ÷ 256 threads/block = 8 blocks per SM
  - 8 blocks × 0 smem = 0 smem < 164KB → shared memory 足够
  - 实际 Occupancy ≈ 100%

但注意：高 Occupancy 不等于高性能。Naive matmul 虽然 Occupancy 接近 100%，但因为算术强度只有 0.25 Flops/Byte，它仍然是 Memory-Bound，大量时间在等待内存读取。

2.4 为什么高 Occupancy 可能还是慢

指标	Naive Matmul	优化后 Matmul
Occupancy	~100%	~75%
算术强度	0.25 Flops/Byte	16 Flops/Byte
实际性能	~2.7 TFLOPS	~16 TFLOPS
瓶颈	Memory-Bound	Compute-Bound

结论： Occupancy 告诉你"GPU 有多忙"，但 Roofline 告诉你忙在哪里。一个 100% Occupancy 的 Memory-Bound kernel，可能比 75% Occupancy 的 Compute-Bound kernel 慢 5 倍。

第3节 Arithmetic Intensity 深入分析

3.1 GEMM 的计算密度随矩阵大小变化

GEMM 的算术强度公式：

C = A × B, 其中 A ∈ R^(M×K), B ∈ R^(K×N), C ∈ R^(M×N)

对于标准 GEMM:
  - 总 Flops = 2 × M × N × K (乘加)
  - 总 Bytes = M×K×4 + K×N×4 + M×N×4 (A+B+C 各一次读)
  - 每输出元素 Bytes = (M×K + K×N) / (M×N) × 4 = 4K × (1/M + 1/N)
  - 每输出元素 Flops = 2K
  - 算术强度 = 2K / (4K × (1/M + 1/N)) = 1 / (2 × (1/M + 1/N))
                                        = MN / (2(M+N))

简化近似：对于 M, N >> K 的情况（常见于深度学习 batch matmul）：

算术强度 ≈ MN / (2(M+N)) ≈ min(M, N) / 2

当 M = N = 2048, K = 64:
  算术强度 = 2048 × 2048 / (2 × (2048 + 2048))
           = 4,194,304 / 8192
           ≈ 512 Flops/Byte  → Compute-Bound!

3.2 为什么大矩阵快，小矩阵慢

算术强度对比 (FP32, A100):

矩阵大小          K值   算术强度    瓶颈类型    理论性能
──────────────────────────────────────────────────────
64×64             64    ~16        Memory     ~32 GFLOPS
256×256           256   ~64        Memory     ~130 GFLOPS  
512×512           512   ~128       临界附近    ~250 GFLOPS
1024×1024         1024  ~256       Compute    ~2 TFLOPS
2048×2048         2048  ~512       Compute    ~10 TFLOPS
4096×4096         4096  ~1024      Compute    ~19 TFLOPS (接近峰值)

关键规律： 当矩阵很大时，每个元素需要读取的数据量相对减少（数据复用率提高），算术强度上升，变成 Compute-Bound，性能接近峰值。

3.3 深度学习中 Batch Matmul 的特殊之处

Transformer 中的 Attention 计算：

Q ∈ R^(B, H, S, D), K ∈ R^(B, H, S, D), V ∈ R^(B, H, S, D)
Attention Score = Q × K^T ∈ R^(B, H, S, S)

对于 S=512, D=64:
  - M = B×H = 12×16 = 192
  - N = S = 512
  - K = D = 64
  
  算术强度 = MN / (2(M+N)) = 192×512 / (2×704) ≈ 70 Flops/Byte
  
  → Memory-Bound 但接近临界点
  → Flash Attention 通过 IO-Aware 算法优化这个问题

第4节 Memory Access Pattern 分析

4.1 Global Memory 访问模式

Coalesced vs Uncoalesced 访问：

cuda

// ❌ Uncoalesced 访问 (每个线程访问不连续地址)
__global__ void bad_access(float* A, float* B, int N) {
    int tid = threadIdx.x;
    int lane_id = tid % 32;
    // 线程束中每个线程访问 A[lane_id * stride + offset]
    // stride = N/32，导致 32 个线程访问 32 个不相邻的内存行
    float val = A[lane_id * (N/32)];  // 内存访问不合并
    B[tid] = val;
}

// ✅ Coalesced 访问 (连续线程访问连续地址)
__global__ void good_access(float* A, float* B, int N) {
    int tid = threadIdx.x;
    int block_start = blockIdx.x * blockDim.x;
    // 线程 0 访问 A[block_start + 0]
    // 线程 1 访问 A[block_start + 1]
    // ... 线程 31 访问 A[block_start + 31]
    // 合并成一次内存事务！
    float val = A[block_start + tid];
    B[tid] = val * 2.0f;
}

4.2 Shared Memory Bandwidth vs Global Memory

内存层级带宽对比 (A100):
┌─────────────────────────────────────────────┐
│  Register     → ~100,000 GB/s (理论无限)     │
│  L1 Cache     →    13,000 GB/s              │
│  L2 Cache     →     2,500 GB/s              │
│  Shared Mem   →     13,000 GB/s              │
│  HBM (Global) →     2,039 GB/s              │
└─────────────────────────────────────────────┘

Shared Memory 带宽是 Global Memory 的 ~6.4 倍
这就是 tiling 优化的核心：通过 shared memory 复用数据，减少 global memory 访问

4.3 矩阵乘法的分块策略

Naive Matmul 的内存访问 (K=1024):
  
  每个线程计算 C[i,j] = sum_k A[i,k] * B[k,j]
  
  对于线程 (i,j):
    - 读取 A[i,*]: K 个 float (行 i 的全部元素)
    - 读取 B[*,j]: K 个 float (列 j 的全部元素)
    - 写 C[i,j]: 1 个 float
  
  问题：同一行 A[i,*] 被 N 个线程重复读取 N 次！
        同一列 B[*,j] 被 M 个线程重复读取 M 次！

使用 Tiling 后的内存访问 (Tile=32):
  
  Block (bi, bj) 处理子矩阵 C[bi:bi+32, bj:bj+32]
  
  Step 1: 把 A[bi:bi+32, :] 和 B[: , bj:bj+32] 加载到 shared memory
    - 一次读取，供 32×32 = 1024 个线程复用
    - 相比 naive，减少了 32 倍的 global memory 访问
  
  Step 2: 每个线程从 shared memory 读取 32 个 A 元素 + 32 个 B 元素
    - Shared memory 带宽 13 TB/s vs Global memory 2 TB/s
    - 加速 6.4 倍

第5节 Kernel 性能瓶颈诊断流程

5.1 诊断步骤

┌──────────────────────────────────────────────────────────────┐
│  步骤 1: 测量实际性能                                         │
│    → 运行 benchmark，记录 GFLOPS / 延迟                       │
├──────────────────────────────────────────────────────────────┤
│  步骤 2: 计算算术强度                                         │
│    → Flops / Bytes = ? Flops/Byte                           │
├──────────────────────────────────────────────────────────────┤
│  步骤 3: 对比 Roofline                                        │
│    → Memory-Bound 还是 Compute-Bound？                       │
│    → 距离理论上限有多远？                                     │
├──────────────────────────────────────────────────────────────┤
│  步骤 4: 内存访问分析                                         │
│    → 访问模式是否 coalesced？                                │
│    → shared memory 复用率如何？                               │
│    → 使用 Nsight Compute 分析内存效率                         │
├──────────────────────────────────────────────────────────────┤
│  步骤 5: Occupancy 分析                                       │
│    → 使用 Nsight Compute 或 nvprof                          │
│    → 是寄存器压力还是 shared memory 压力？                   │
├──────────────────────────────────────────────────────────────┤
│  步骤 6: 针对性优化                                          │
│    → Memory-Bound: tiling, shared memory, 数据预取          │
│    → Compute-Bound: 减少分支，提升 ILP, 使用 Tensor Core     │
└──────────────────────────────────────────────────────────────┘

5.2 Nsight Compute 分析命令

bash

# 编译 kernel
nvcc -O3 -use_fast_norm -arch=sm_80 -lineinfo matmul.cu -o matmul

# 运行 Nsight Compute 分析
ncu --set full ./matmul

# 关键指标解读:
#  achieved_occupancy: 实际 SM 占用率
#  sm_efficiency: SM 利用率
#  warp_execution_efficiency: Warp 执行效率
#  gld_throughput: 全局加载吞吐
#  gst_throughput: 全局存储吞吐
#  l1_texture_cache_hit_rate: L1 缓存命中率
#  shared_efficiency: Shared memory 效率

# 详细内存分析
ncu --set full --metrics l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,
                        l1tex__t_sectors_pipe_lsu_mem_global_op_st.sum,
                        sm__throughput.avg.pct_of_peak_sustained_elapsed ./matmul

第6节 CPU vs GPU Kernel 设计差异

6.1 核心设计哲学差异

维度	CPU Kernel	GPU Kernel
并行粒度	线程级并行（几个到几十个线程）	数千个线程并行执行
内存层级	L1/L2/L3 缓存层次清晰	Shared Memory + Global Memory
瓶颈	通常是计算（ALU 饱和）	通常是内存带宽
优化重点	缓存局部性、分支预测	内存合并、Occupancy、算术强度
编程模型	多线程 + SIMD 向量化	SIMT（单指令多线程）

6.2 CPU 上的 GEMM 优化策略

// CPU GEMM 优化：分块 + 缓存友好
void gemm_blocked(float* C, float* A, float* B, int M, int N, int K) {
    // 分块大小：32×32 (能在 L1 cache 中放下)
    const int BLOCK = 32;
    
    // 外层循环：遍历 C 的块
    for (int i = 0; i < M; i += BLOCK) {
        for (int j = 0; j < N; j += BLOCK) {
            // 中层循环：累加 K 方向
            for (int k = 0; k < K; k += BLOCK) {
                // 内层核心：2 层嵌套的小矩阵乘
                // BLOCK×BLOCK 块能完全放入 L1 cache (32×32×4×3 ≈ 12KB < 48KB L1)
                for (int ii = i; ii < min(i+BLOCK, M); ii++) {
                    for (int jj = j; jj < min(j+BLOCK, N); jj++) {
                        float sum = 0.0f;
                        for (int kk = k; kk < min(k+BLOCK, K); kk++) {
                            sum += A[ii * K + kk] * B[kk * N + jj];
                        }
                        C[ii * N + jj] += sum;
                    }
                }
            }
        }
    }
}

CPU vs GPU 优化对比：

CPU GEMM 优化:
  - 分块大小由 L1/L2/L3 cache 大小决定（通常 32-64）
  - 关注 cache line 复用（一次加载后用多次）
  - 使用 SIMD 指令（AVX-512: 512bits = 16×float）
  - 分支预测优化（循环展开）

GPU GEMM 优化:
  - 分块大小由 shared memory 和 register 决定（通常 16-256）
  - 关注 global memory 合并访问
  - 使用 Tensor Core（16×16 矩阵乘）
  - Warp 协同调度（__syncwarp()）

升华

┌────────────────────────────────────────────────────────────────────────┐
│                     🚀 Kernel 性能优化核心原则                          │
├────────────────────────────────────────────────────────────────────────┤
│                                                                        │
│  ① 先测后优化：永远用 Roofline Model 确定瓶颈类型再动手                  │
│     盲目优化 Memory-Bound kernel 的计算部分 = 浪费时间                  │
│                                                                        │
│  ② 算术强度是根本：提高数据复用率才是 Memory-Bound 的正解                │
│     Tiling、Shared Memory、数据布局转换都是手段                          │
│                                                                        │
│  ③ Occupancy 高 ≠ 性能好：高 Occupancy + Memory-Bound = 假忙            │
│     要同时关注 SM Efficiency 和 Memory Efficiency                      │
│                                                                        │
│  ④ 测量驱动：nsys/ncu 分析结果是事实，猜测不是                         │
│     优化前后对比，目标是什么指标提升了多少？                               │
│                                                                        │
└────────────────────────────────────────────────────────────────────────┘

"AI 可查 vs 必须理解"清单

必须理解（不理解就等于不会）：

🔴 Roofline Model 的物理意义：为什么算术强度决定瓶颈类型，临界点怎么算
🔴 Memory-Bound vs Compute-Bound 的区别：分别意味着什么，如何针对性优化
🔴 算术强度公式：Flops/Bytes 怎么推导，GEMM 的算术强度怎么算
🔴 Tiling 的本质：为什么能提升算术强度，shared memory 复用数据的原理
🔴 Occupancy 局限性：Occupancy 高但还是慢的原因分析

AI 可查（知道去哪查就行）：

✅ A100/H100 具体带宽和算力数据（硬件参数随时变化）
✅ Nsight Compute 具体 metrics 含义（官方文档）
✅ nvcc 编译选项和 PTXAS 寄存器分配策略（官方文档）
✅ 特定 GPU architecture 的 shared memory 大小限制（硬件规格）
✅ 具体的 cuBLAS/cuDNN kernel 性能数据（NVIDIA 官方 benchmark）

学习状态：🟡 开始学习

Kernel 性能基础：Roofline 与 Occupancy / Kernel Performance Fundamentals with Roofline and Occupancy ​

第1节 Roofline Model——理解计算上限 ​

1.1 什么是 Roofline Model ​

1.2 两种瓶颈类型 ​

1.3 A100 具体数据 ​

1.4 用 Roofline 诊断你的 Matmul Kernel ​

1.5 Roofline 分析 Python 脚本 ​

第2节 GPU Occupancy——每 SM 活跃线程数 ​

2.1 什么是 Occupancy ​

2.2 寄存器压力和 Shared Memory 压力 ​

2.3 计算 Naive Matmul 的 Occupancy ​

2.4 为什么高 Occupancy 可能还是慢 ​

第3节 Arithmetic Intensity 深入分析 ​

3.1 GEMM 的计算密度随矩阵大小变化 ​

3.2 为什么大矩阵快，小矩阵慢 ​

3.3 深度学习中 Batch Matmul 的特殊之处 ​

第4节 Memory Access Pattern 分析 ​

4.1 Global Memory 访问模式 ​

4.2 Shared Memory Bandwidth vs Global Memory ​

4.3 矩阵乘法的分块策略 ​

第5节 Kernel 性能瓶颈诊断流程 ​

5.1 诊断步骤 ​

5.2 Nsight Compute 分析命令 ​

第6节 CPU vs GPU Kernel 设计差异 ​

6.1 核心设计哲学差异 ​

6.2 CPU 上的 GEMM 优化策略 ​

升华 ​

"AI 可查 vs 必须理解"清单 ​

Kernel 性能基础：Roofline 与 Occupancy / Kernel Performance Fundamentals with Roofline and Occupancy

第1节 Roofline Model——理解计算上限

1.1 什么是 Roofline Model

1.2 两种瓶颈类型

1.3 A100 具体数据

1.4 用 Roofline 诊断你的 Matmul Kernel

1.5 Roofline 分析 Python 脚本

第2节 GPU Occupancy——每 SM 活跃线程数

2.1 什么是 Occupancy

2.2 寄存器压力和 Shared Memory 压力

2.3 计算 Naive Matmul 的 Occupancy

2.4 为什么高 Occupancy 可能还是慢

第3节 Arithmetic Intensity 深入分析

3.1 GEMM 的计算密度随矩阵大小变化

3.2 为什么大矩阵快，小矩阵慢

3.3 深度学习中 Batch Matmul 的特殊之处

第4节 Memory Access Pattern 分析

4.1 Global Memory 访问模式

4.2 Shared Memory Bandwidth vs Global Memory

4.3 矩阵乘法的分块策略

第5节 Kernel 性能瓶颈诊断流程

5.1 诊断步骤

5.2 Nsight Compute 分析命令

第6节 CPU vs GPU Kernel 设计差异

6.1 核心设计哲学差异

6.2 CPU 上的 GEMM 优化策略

升华

"AI 可查 vs 必须理解"清单