📅 创建时间：2026-06-03 🏷️ 标签：#算子融合 #Fusion #ElementwiseFusion #MatmulFusion #HorizontalFusion #Transformer优化 #MemoryBound 📚 前置知识：[[/04-ai/01-llm-engineering/07-llm-evolution]]（LLM 发展脉络） 📚 前置知识：[[07-graph-optimization-passes]]（图优化 Pass） 📚 相关知识：[[09-memory-planning]]（内存规划） [[15-backend-cuda]]（CUDA 后端）

┌──────────────────────────────────────────────────────────────────────────────┐ │ 🔍 场景：GPU 利用率只有 15%？ │ ├──────────────────────────────────────────────────────────────────────────────┤ │ 你的 transformer 推理速度上不去。用 nsys profiler 分析后发现： │ │ - GPU Utilization 只有 15% │ │ - 大部分时间花在 Memory Copy 和 Kernel Launch overhead │ │ - 每个单独的 GEMM 和 Activation 都单独 launch 一个 CUDA kernel │ │ - 100 个算子 = 100 次 kernel launch，GPU 大部分时间在等指令调度 │ │ │ │ 融合这些算子后，GPU 利用率提升到 55%，推理速度翻倍。 │ └──────────────────────────────────────────────────────────────────────────────┘

算子融合——编译器最重要的性能优化 / Operator Fusion as a Core Compiler Optimization

第1节：为什么融合是最重要的优化

1.1 Kernel Launch Overhead

在 CUDA 编程中，每次调用 GPU kernel 都有固定开销：

cuda

// 每个 kernel launch 都有以下开销：
// 1. 主机到设备的命令传输 (~1-5 微秒)
// 2. 设备端命令队列排队 (~0.1-1 微秒)
// 3. 核函数启动开销 (~0.5-2 微秒)
// 4. 寄存器分配和线程块调度

// 如果我们有 100 个小算子：
for (int i = 0; i < 100; i++) {
    cudaMemcpyAsync(...);    // 同步开销
    my_kernel<<<...>>>(...); // kernel launch
    cudaStreamSynchronize();  // 等待完成
}
// 总开销：100 × (5 + 1 + 2 + 1) = 900 微秒 = 0.9ms

// 融合后：
fused_kernel<<<...>>>(...);  // 单次 launch
// 总开销：5 微秒 = 0.005ms

// 性能提升：180 倍！

1.2 内存带宽浪费

未融合的算子之间需要读写显存：

python

# 未融合的情况
def unfused_attention(Q, K, V):
    # 每个步骤都是独立的 kernel
    QK = matmul(Q, K.transpose(-2, -1))  # kernel 1: 读 Q, K，写 QK
    QK_scaled = QK * scale               # kernel 2: 读 QK，写 scaled
    attn_weights = softmax(QK_scaled)    # kernel 3: 读 scaled，写 attn
    output = matmul(attn_weights, V)    # kernel 4: 读 attn, V，写 output
    
    # 中间结果全部写回显存再读出
    # Memory bandwidth: (Q + K + QK + scaled + attn + V + output) 全部经过显存
    
    return output


def fused_attention(Q, K, V):
    # 所有操作在一个 kernel 中完成
    # 只需要读写 Q, K, V, output
    # Memory bandwidth: (Q + K + V + output)
    
    # 节省的内存带宽：
    # - QK: 不需要
    # - scaled: 不需要
    # - attn: 不需要
    return fused_kernel(Q, K, V)

1.3 融合的收益分析

指标	未融合	融合后	改善
Kernel Launch 次数	100	5-10	10-20×
中间 Activation 显存	500 MB	50 MB	10×
内存带宽占用	800 GB/s	400 GB/s	2×
GPU Utilization	15%	55%	3.7×
端到端延迟	100ms	45ms	2.2×

第2节：融合的类型详解

2.1 Element-wise Fusion（逐元素融合）

最常见的融合类型：将 element-wise 操作与主要算子融合。

python

# 原始图结构
# x → matmul → add(bias) → relu → dropout → output

# Element-wise 融合后
# x → fused_linear_relu_dropout → output

class ElementwiseFusion:
    """
    逐元素融合：将 element-wise 操作（ReLU, Sigmoid, Add, Mul 等）
    融合到其输入的 producers 中
    """
    
    FUSION_PATTERNS = [
        # matmul + bias
        ('matmul', 'add'),           # → linear
        # conv + bias + activation
        ('conv2d', 'add', 'relu'),   # → conv2d_relu (有 bias)
        ('conv2d', 'relu'),          # → conv2d_relu (无 bias)
        # matmul + add + activation
        ('matmul', 'add', 'gelu'),  # → fused_gemm_gelu
    ]
    
    def fuse_elementwise(self, graph):
        """扫描图中所有可融合的 element-wise 模式"""
        
        for pattern in self.FUSION_PATTERNS:
            matches = self._find_matches(graph, pattern)
            
            for match in matches:
                fused_op = self._create_fused_op(match)
                graph.replace(match, fused_op)

ReLU(GEMM(x)) 融合示例：

python

# 未融合的 PyTorch 实现
class UnfusedLayer(torch.nn.Module):
    def forward(self, x):
        x = torch.matmul(x, self.weight)  # kernel 1: GEMM
        x = x + self.bias                  # kernel 2: Element-wise add
        x = torch.relu(x)                  # kernel 3: ReLU
        return x

# 等价于：
# 3 个 kernel
# 3 次显存读写（中间结果）
# 3 次 kernel launch overhead

# 融合后的实现（cuBLAS 风格）
class FusedLayer(torch.nn.Module):
    def forward(self, x):
        # 使用 cuBLAS 的 bias ReLU GEMM
        # 单一 kernel 调用
        return torch.ops.aten._cuBLAS_linear_relu(x, self.weight, self.bias)

2.2 Consumer-Producer Fusion（生产者-消费者融合）

将数据依赖链上的多个算子融合成一个：

python

# Attention 机制是典型的 producer-consumer 链
# Q = W_q @ x
# K = W_k @ x  
# V = W_v @ x
# scores = Q @ K.T
# attn = softmax(scores)
# output = attn @ V

# Consumer-Producer 融合后：
# Q, K, V = fused_qkv_proj(x)              # 一次性计算 QKV
# output = fused_attention(Q, K, V)        # 完整的 attention 计算

python

# Flash Attention 的融合策略
def flash_attention_fused(Q, K, V, scale, causal_mask=None):
    """
    Flash Attention 融合 kernel
    
    包含的操作：
    1. Scaled dot-product: Q @ K.T * scale
    2. Masking (causal or attention mask)
    3. Softmax normalization
    4. Dropout (training)
    5. Matmul with V
    6. Residue connection (Fused with output projection)
    
    所有操作在一个 kernel 中完成
    """
    
    # 块级计算，减少 HBM 访问
    # 每个 block 只需要加载 Q, K, V 的一个块
    # 中间结果保存在 SM 寄存器/SRAM 中
    
    return output


# 融合带来的收益：
# - HBM 访问量从 O(N^2) 减少到 O(N^2 / block_size)
# - 不需要存储完整的 attention matrix
# - Memory complexity 从 O(N^2) 降到 O(N)

2.3 Horizontal Fusion（水平融合）

将多个独立的、可以并行执行的小操作合并：

python

# Horizontal Fusion 示例：批量 LayerNorm

# 原始：4 个独立的 LayerNorm
ln1 = LayerNorm(x1)
ln2 = LayerNorm(x2)
ln3 = LayerNorm(x3)
ln4 = LayerNorm(x4)

# 融合后：将多个 LayerNorm 合并成一个大 kernel
# 输入：张量堆叠 (4, seq_len, hidden)
# 输出：堆叠的归一化结果 (4, seq_len, hidden)

def fused_multi_layernorm(tensors, eps=1e-5):
    """
    水平融合多个 LayerNorm
    
    优点：
    - 一次 kernel launch
    - 更好的并行度（更多线程）
    - 减少 kernel 间同步开销
    """
    # Concatenate along batch dimension
    x = torch.cat(tensors, dim=0)
    
    # Fused LayerNorm kernel
    return fused_layernorm_kernel(x, eps)

python

# ViT 中的 Horizontal Fusion 示例
class FusedViTBlock(torch.nn.Module):
    """
    ViT Block 包含：
    1. Attention: QKV projection + attention + output projection
    2. MLP: 2-3 层全连接 + GELU
    
    可以水平融合多个 ViT Block 的 attention 计算
    """
    
    def forward(self, x_list):
        # 输入：多个 token 的序列
        # 每个序列独立计算 attention
        
        # 水平融合：一次性计算所有序列的 QKV
        # 利用 GPU 的并行性，一次处理更多数据
        
        all_qkv = self.compute_qkv_fused(x_list)
        all_attn_out = self.compute_attention_fused(all_qkv)
        
        return all_attn_out

第3节：融合的决策边界

3.1 融合过多的问题

python

# 融合过度的反模式
class OverFusedModel(torch.nn.Module):
    """
    问题：过度融合导致寄存器溢出
    
    假设我们将整个 Transformer 融合成一个 kernel：
    - 输入: [batch, seq_len, hidden]
    - 输出: [batch, seq_len, vocab_size]
    
    问题：
    1. 寄存器压力：需要同时持有 Q, K, V, attn, mlp 等中间结果
    2. Shared Memory 不足：无法存储所有中间结果
    3. Occupancy 降低：每个线程做太多工作
    4. 编译时间爆炸：融合 kernel 很大，优化时间长
    """
    
    def forward(self, x):
        # 这不是一个好主意
        return mega_transformer_kernel(x)

融合过度的症状：

问题	症状	原因
寄存器溢出	性能下降 30-50%	寄存器溢出到 local memory
Shared Memory 不足	融合失败	中间结果太大
Occupancy 低	GPU 利用率低	每个 SM 运行的 block 太少
编译时间过长	编译需要 10+ 分钟	优化搜索空间太大

3.2 不融合的问题

python

# 不融合的反模式
class UnfusedModel(torch.nn.Module):
    """
    问题：过多的 kernel launch overhead
    
    每个操作都独立 kernel，导致：
    - 100 个算子 = 100 次 launch overhead
    - 大量中间结果显存读写
    - GPU 等待同步
    """
    
    def __init__(self):
        super().__init__()
        self.layers = torch.nn.ModuleList([
            torch.nn.Linear(512, 2048),
            torch.nn.GELU(),
            torch.nn.Linear(2048, 512),
        ]) * 24  # 24 层
    
    def forward(self, x):
        for layer in self.layers:
            x = layer(x)  # 每个 layer 都是独立 kernel
        return x
    # 72 个独立的 kernel！

3.3 融合决策指南

python

# 融合决策算法
def should_fuse(pattern, graph, hardware):
    """
    判断是否应该融合
    
    考虑因素：
    1. Kernel launch overhead vs 融合收益
    2. 中间结果显存 vs 融合后计算量
    3. 硬件限制（寄存器、shared memory）
    """
    
    # 经验法则：
    # 1. Element-wise 操作几乎总是值得融合
    # 2. Memory-bound 算子融合收益大
    # 3. Compute-bound 大算子融合要谨慎
    
    launch_overhead = estimate_launch_overhead(pattern)
    memory_saving = estimate_memory_saving(pattern)
    compute_change = estimate_compute_change(pattern)
    
    # 决策条件
    if is_elementwise(pattern[-1]):
        return True  # element-wise 操作几乎总是值得融合
    
    if memory_saving > compute_change * hardware.memory_bandwidth_ratio:
        return True  # 内存节省大于计算变化
    
    if pattern_size(pattern) > hardware.max_fusable_ops:
        return False  # 超过硬件限制
    
    if register_pressure(pattern) > hardware.max_registers_per_block:
        return False  # 寄存器压力太大
    
    return launch_overhead > memory_saving

第4节：Fused Kernel 生成

4.1 从融合图生成 Kernel

python

class FusedKernelGenerator:
    """
    如何从融合图中生成一个融合 kernel
    """
    
    def generate(self, fused_op):
        """
        融合 kernel 生成流程：
        
        1. 分析数据流：确定输入、输出、中间结果
        2. 确定循环结构：外层 batch/sequence 循环，内层 feature 循环
        3. 生成计算代码：按照数据流拼接各个算子
        4. 寄存器分配：尽量将中间结果保存在寄存器
        5. Shared Memory 规划：如果中间结果太大，使用 shared memory
        """
        
        # Step 1: 数据流分析
        dataflow = self._analyze_dataflow(fused_op)
        # 输出: {
        #   'inputs': ['x'],
        #   'computations': ['matmul', 'add', 'relu'],
        #   'outputs': ['y'],
        #   'intermediates': ['matmul_out', 'relu_out']
        # }
        
        # Step 2: 确定循环结构
        loop_structure = self._determine_loops(dataflow)
        
        # Step 3: 生成 CUDA 代码
        kernel_code = self._generate_cuda_code(dataflow, loop_structure)
        
        return kernel_code


# 生成的融合 kernel 示例
CUDA_KERNEL_TEMPLATE = """
__global__ void fused_linear_relu_kernel(
    const float* __restrict__ input,   // 输入 (N, K)
    const float* __restrict__ weight,  // 权重 (M, K)
    const float* __restrict__ bias,    // 偏置 (M,)
    float* __restrict__ output,        // 输出 (N, M)
    int N, int K, int M
) {{
    // 外层循环：batch
    for (int n = blockIdx.x * blockDim.x + threadIdx.x; 
         n < N; 
         n += blockDim.x * gridDim.x) {
        
        // 临时寄存器
        float sum[M];  // 存储一行的结果
        
        // 清零
        #pragma unroll
        for (int m = 0; m < M; m++) {
            sum[m] = 0.0f;
        }
        
        // 内层循环：累加
        // 注意：K 可能很大，不能在寄存器中累加
        for (int k = 0; k < K; k++) {
            float x_val = input[n * K + k];
            #pragma unroll
            for (int m = 0; m < M; m++) {
                sum[m] += x_val * weight[m * K + k];
            }
        }
        
        // 加上 bias
        #pragma unroll
        for (int m = 0; m < M; m++) {
            sum[m] += bias[m];
        }
        
        // ReLU 激活（融合在此）
        #pragma unroll
        for (int m = 0; m < M; m++) {
            output[n * M + m] = sum[m] > 0 ? sum[m] : 0;
        }
    }}
}}
"""

4.2 使用 Triton 生成融合 Kernel

python

import triton
import triton.language as tl

@triton.jit
def fused_linear_relu_kernel(
    # Pointers to matrices
    x_ptr, W_ptr, B_ptr, Y_ptr,
    # Matrix dimensions
    M, N, K,
    # Strides
    stride_xm, stride_xk,
    stride_wn, stride_wk,
    stride_ym, stride_yn,
    # Meta-parameters
    BLOCK_SIZE_M: tl.constexpr,
    BLOCK_SIZE_N: tl.constexpr,
    BLOCK_SIZE_K: tl.constexpr,
):
    """
    Triton 融合 kernel：Linear + ReLU
    
    x: (M, K) input
    W: (N, K) weights (row-major)
    B: (N,) bias
    Y: (M, N) output
    """
    
    # 程序 ID
    pid_m = tl.program_id(0)
    pid_n = tl.program_id(1)
    
    # 初始化累加器
    offs_m = (pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)) % M
    offs_n = (pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) % N
    offs_k = tl.arange(0, BLOCK_SIZE_K)
    
    # 加载权重到 shared memory
    w_ptrs = W_ptr + offs_n[:, None] * stride_wn + offs_k[None, :] * stride_wk
    w_mask = (offs_n[:, None] < N) & (offs_k[None, :] < K)
    w = tl.load(w_ptrs, mask=w_mask, other=0.0)
    
    # 累加循环
    acc = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
    
    for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)):
        # 加载输入
        x_offs_m = offs_m[:, None]
        x_offs_k = (k * BLOCK_SIZE_K + offs_k)[None, :]
        x_ptrs = x_ptr + x_offs_m * stride_xm + x_offs_k * stride_xk
        x_mask = (offs_m[:, None] < M) & ((k * BLOCK_SIZE_K + offs_k)[None, :] < K)
        x = tl.load(x_ptrs, mask=x_mask, other=0.0)
        
        # 矩阵乘法累加
        acc += tl.dot(x, w)
        
        # 更新指针
        offs_k += BLOCK_SIZE_K
    
    # 加上 bias
    b_ptrs = B_ptr + offs_n
    b = tl.load(b_ptrs, mask=offs_n < N, other=0.0)
    acc = acc + b[None, :]
    
    # ReLU 激活
    acc = tl.where(acc > 0, acc, 0.0)
    
    # 写回
    y_ptrs = Y_ptr + offs_m[:, None] * stride_ym + offs_n[None, :] * stride_yn
    y_mask = (offs_m[:, None] < M) & (offs_n[None, :] < N)
    tl.store(y_ptrs, acc, mask=y_mask)


def fused_linear_relu(x, weight, bias):
    """
    调用 Triton 融合 kernel
    """
    M, K = x.shape
    N, K2 = weight.shape
    assert K == K2
    
    Y = torch.empty((M, N), device=x.device, dtype=x.dtype)
    
    # 配置 grid
    BLOCK_SIZE_M = 16
    BLOCK_SIZE_N = 16
    BLOCK_SIZE_K = 32
    
    grid = (triton.cdiv(M, BLOCK_SIZE_M), triton.cdiv(N, BLOCK_SIZE_N))
    
    fused_linear_relu_kernel[grid](
        x, weight, bias, Y,
        M, N, K,
        x.stride(0), x.stride(1),
        weight.stride(0), weight.stride(1),
        Y.stride(0), Y.stride(1),
        BLOCK_SIZE_M, BLOCK_SIZE_N, BLOCK_SIZE_K,
    )
    
    return Y

第5节：Transformer 中的关键融合模式

5.1 Attention Pattern 融合

python

# Attention 计算的完整融合模式

# 原始分解（每个操作独立 kernel）：
# 1. QKV = W_qkv @ x              # 矩阵乘法
# 2. Q = QKV[:, :, :hidden]       # 切片
# 3. K = QKV[:, :, hidden:2*hidden]  # 切片
# 4. V = QKV[:, :, 2*hidden:]     # 切片
# 5. QK = Q @ K.T                 # 矩阵乘法
# 6. scaled = QK / sqrt(d)        # 缩放
# 7. masked = scaled + mask       # 掩码
# 8. attn = softmax(masked)       # Softmax
# 9. attn_drop = dropout(attn)     # Dropout (训练)
# 10. out = attn_drop @ V         # 矩阵乘法
# 11. out = out + residue         # 残差

# 融合后（多个 kernel）：
# Kernel 1: fused_qkv_proj(x) → {Q, K, V}
# Kernel 2: fused_attention_core(Q, K, V, mask, dropout_mask) → output
# Kernel 3: fused_residue_norm(output, residue) → final_output

python

# Flash Attention 的融合策略
class FlashAttentionKernel:
    """
    Flash Attention 2.0 融合策略
    
    核心思想：
    1. 不存储完整的 N×N attention matrix
    2. 分块计算，逐步归约
    3. 所有操作在一个 kernel 中完成
    """
    
    def forward_kernel(self, Q, K, V, scale, causal_mask=None):
        """
        融合的操作序列：
        
        1. 加载 Q, K, V 的一个块
        2. 计算 Q @ K.T * scale
        3. 应用 causal mask（如果需要）
        4. 计算 softmax（在线算法）
        5. 计算 softmax(QK) @ V
        6. 更新累加器
        
        所有这些操作在一个 kernel 中完成
        只需要 O(N) 的额外显存，而不是 O(N^2)
        """
        
        # 块级 Flash Attention
        # 每块大小通常为 32x64 或 64x64
        pass

5.2 MLP 融合模式

python

# Transformer MLP 层融合

# 原始：
# h = x @ W_up          # Up-projection (N, hidden) @ (hidden, 4*hidden) = (N, 4*hidden)
# h = gelu(h)            # Element-wise
# h = h @ W_down        # Down-projection (N, 4*hidden) @ (4*hidden, hidden) = (N, hidden)

# 融合后：fused_mlp_kernel(x) → output
# 单一 kernel 完成所有计算

# 融合的额外收益：
# - GELU 的非线性可以融合到矩阵乘法中
# - 中间结果 h 不需要写回显存

5.3 完整 Transformer Block 融合

python

# 极致融合：将多个 Transformer Block 融合

# 融合范围：
# - QKV Projection（3 个 MatMul 融合）
# - Attention 计算（Flash Attention）
# - Output Projection
# - MLP（Up + GELU + Down）
# - 残差连接
# - LayerNorm

# 分层融合策略：
# 1. Intra-layer fusion：单个 Block 内的算子融合
# 2. Inter-layer fusion：多个 Block 共享某些计算（如共享 K, V）
# 3. Inter-model fusion：多个模型实例融合（如 batch 内多个样本）

def fused_transformer_block(x, W_qkv, W_o, W_up, W_down, 
                            norm1_weight, norm2_weight, 
                            scale, mask):
    """
    完整 Transformer Block 融合
    
    包含：
    1. LayerNorm
    2. QKV Projection
    3. Self-Attention
    4. Attention Output Projection
    5. 残差连接
    6. LayerNorm
    7. MLP
    8. 残差连接
    """
    
    # LayerNorm + QKV
    normed = layer_norm_fused(x, norm1_weight)
    qkv = fused_qkv_matmul(normed, W_qkv)
    
    # Split QKV
    Q, K, V = split_qkv(qkv)
    
    # Flash Attention
    attn_out = flash_attention_fused(Q, K, V, scale, mask)
    
    # Output Projection + 残差
    out1 = fused_matmul_add(attn_out, W_o, x)
    
    # LayerNorm + MLP + 残差
    normed2 = layer_norm_fused(out1, norm2_weight)
    mlp_out = fused_mlp(normed2, W_up, W_down)
    out2 = out1 + mlp_out
    
    return out2

第6节：融合的调试——使用 Profiler 找到未融合的 Pattern

6.1 PyTorch Profiler 使用

python

import torch
from torch.profiler import profile, ProfilerActivity

# 使用 PyTorch Profiler 分析融合情况
def profile_fusion(model, input_tensor):
    """
    分析模型的融合情况
    """
    
    with profile(
        activities=[ProfilerActivity.CPU, ProfilerActivity.CUDA],
        record_shapes=True,
        profile_memory=True,
        with_stack=True,
    ) as prof:
        # Warmup
        for _ in range(3):
            _ = model(input_tensor)
        
        # Profiling
        for _ in range(10):
            _ = model(input_tensor)
    
    # 打印 CUDA kernel 统计
    print(prof.key_averages().table(
        sort_by="cuda_time_total", 
        row_limit=20
    ))
    
    # 查找未融合的操作
    # 如果看到太多单独的 element-wise kernel，说明融合不够
    
    return prof


# 示例输出分析
def analyze_profiler_output(profiler_output):
    """
    分析 Profiler 输出，识别融合问题
    """
    
    key_averages = profiler_output.key_averages()
    
    # 统计各类 kernel 的调用次数
    kernel_counts = {}
    for item in key_averages:
        name = item.key
        if "cuda" in str(type(item)).lower():
            kernel_counts[name] = item.count
    
    # 识别过多 kernel launch 的模式
    suspicious_patterns = []
    
    # 1. 大量独立的 element-wise kernels
    elementwise_kernels = [k for k in kernel_counts if 
                          any(op in k.lower() for op in ['relu', 'sigmoid', 'add', 'mul'])]
    
    if sum(kernel_counts[k] for k in elementwise_kernels) > 100:
        suspicious_patterns.append(
            f"太多独立的 element-wise kernels: {len(elementwise_kernels)} 种, "
            f"共 {sum(kernel_counts[k] for k in elementwise_kernels)} 次调用"
        )
    
    # 2. 大量小的 MatMul kernels
    matmul_kernels = [k for k in kernel_counts if 'gemm' in k.lower() or 'matmul' in k.lower()]
    
    if len(matmul_kernels) > 20:
        suspicious_patterns.append(
            f"太多独立的 MatMul kernels: {len(matmul_kernels)} 种"
        )
    
    return suspicious_patterns

6.2 使用 torch.compile 调试融合

python

# 检查 torch.compile 的融合情况
import torch._inductor.config as inductor_config

# 开启融合调试
inductor_config.debug_fusion = True

# 编译模型
model = create_transformer_model()
compiled = torch.compile(model, backend='inductor')

# 第一次运行
input_tensor = torch.randn(1, 512, 768).cuda()
output = compiled(input_tensor)

# 查看生成的代码
inductor_output_dir = "torchinductor_output"
# 会在这个目录下生成 CUDA 代码

python

# 分析 torch.compile 生成的代码
def analyze_generated_kernel(code_path):
    """
    分析 torch.compile 生成的融合 kernel
    """
    
    import os
    
    kernel_files = [f for f in os.listdir(code_path) if f.endswith('.cu')]
    
    print(f"生成了 {len(kernel_files)} 个 kernel 文件：\n")
    
    for kernel_file in kernel_files[:10]:  # 只显示前 10 个
        with open(os.path.join(code_path, kernel_file), 'r') as f:
            content = f.read()
        
        # 统计 kernel 大小
        lines = content.count('\n')
        
        # 查找融合的操作
        fused_ops = []
        if 'linear' in content.lower():
            fused_ops.append('linear')
        if 'relu' in content.lower():
            fused_ops.append('relu')
        if 'gelu' in content.lower():
            fused_ops.append('gelu')
        
        print(f"{kernel_file}:")
        print(f"  - 行数: {lines}")
        print(f"  - 融合的操作: {fused_ops}")
        print()

6.3 Nsight Systems 分析

bash

# 使用 nsys 分析 CUDA kernel 融合情况
# nsys profile -o profile_result -c cudaMemoryWorkloadView_Default \
#     python your_script.py

# 分析结果中的关键指标：
# - GPU Time: GPU 实际执行时间
# - Kernel Statistics: kernel 调用次数和时间
# - Memory Transfer: 显存读写量

# 融合良好的标志：
# - Kernel 数量少（每个融合操作对应一个 kernel）
# - 大部分时间是少量大型 kernel
# - 显存读写主要是输入输出，中间结果很少

升华

┌─────────────────────────────────────────────────────────────────────────────┐
│                       算子融合核心原则                                       │
├─────────────────────────────────────────────────────────────────────────────┤
│  1. 融合是性能的关键杠杆                                                     │
│     → GPU 利用率低通常是因为 kernel launch overhead 或内存带宽瓶颈           │
│  2. Element-wise 操作几乎总是值得融合                                        │
│     → ReLU, Sigmoid, Add, Mul 与上游算子融合                                │
│  3. 融合有边界，不是越多越好                                                  │
│     → 寄存器压力、shared memory 限制、编译时间是制约因素                      │
│  4. 使用 Profiler 驱动融合优化                                               │
│     → 先测量，找到瓶颈，再决定融合策略                                        │
│  5. Flash Attention 是融合的最佳范例                                        │
│     → O(N²) → O(N) 显存，完整的计算融合                                       │
└─────────────────────────────────────────────────────────────────────────────┘

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

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

🔴 为什么 kernel launch overhead 是问题：每次 launch 都有 ~5-10μs 的固定开销
🔴 Element-wise fusion 的原理：ReLU(GEMM(x)) 如何在单 kernel 中完成
🔴 融合过度的后果：寄存器溢出、shared memory 不足、occupancy 降低
🔴 Flash Attention 的融合策略：如何用 O(N) 显存替代 O(N²) 的 attention matrix
🔴 如何使用 Profiler 判断融合是否充分：kernel 数量、调用时间分布

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

✅ 特定硬件的 shared memory 大小和寄存器数量限制
✅ Triton 或 CUTLASS 的具体 API 用法
✅ Flash Attention 的内部算法细节（在线 softmax）
✅ cuBLAS 的融合 GEMM 接口（如 cublasGemmEx）
✅ Nsight Systems 的具体使用命令

学习状态：🟡 开始学习

算子融合——编译器最重要的性能优化 / Operator Fusion as a Core Compiler Optimization ​

第1节：为什么融合是最重要的优化 ​

1.1 Kernel Launch Overhead ​

1.2 内存带宽浪费 ​

1.3 融合的收益分析 ​

第2节：融合的类型详解 ​

2.1 Element-wise Fusion（逐元素融合） ​

2.2 Consumer-Producer Fusion（生产者-消费者融合） ​

2.3 Horizontal Fusion（水平融合） ​

第3节：融合的决策边界 ​

3.1 融合过多的问题 ​

3.2 不融合的问题 ​

3.3 融合决策指南 ​

第4节：Fused Kernel 生成 ​

4.1 从融合图生成 Kernel ​

4.2 使用 Triton 生成融合 Kernel ​

第5节：Transformer 中的关键融合模式 ​

5.1 Attention Pattern 融合 ​

5.2 MLP 融合模式 ​

5.3 完整 Transformer Block 融合 ​

第6节：融合的调试——使用 Profiler 找到未融合的 Pattern ​

6.1 PyTorch Profiler 使用 ​

6.2 使用 torch.compile 调试融合 ​

6.3 Nsight Systems 分析 ​

升华 ​

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

算子融合——编译器最重要的性能优化 / Operator Fusion as a Core Compiler Optimization

第1节：为什么融合是最重要的优化

1.1 Kernel Launch Overhead

1.2 内存带宽浪费

1.3 融合的收益分析

第2节：融合的类型详解

2.1 Element-wise Fusion（逐元素融合）

2.2 Consumer-Producer Fusion（生产者-消费者融合）

2.3 Horizontal Fusion（水平融合）

第3节：融合的决策边界

3.1 融合过多的问题

3.2 不融合的问题

3.3 融合决策指南

第4节：Fused Kernel 生成

4.1 从融合图生成 Kernel

4.2 使用 Triton 生成融合 Kernel

第5节：Transformer 中的关键融合模式

5.1 Attention Pattern 融合

5.2 MLP 融合模式

5.3 完整 Transformer Block 融合

第6节：融合的调试——使用 Profiler 找到未融合的 Pattern

6.1 PyTorch Profiler 使用

6.2 使用 torch.compile 调试融合

6.3 Nsight Systems 分析

升华

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