Transformer 的一个 Block 包含 12+ 个独立算子：LayerNorm → QKV Linear → Reshape → Transpose → Attention → Concat → Linear → LayerNorm → FFN Up → Gelu → FFN Down → Residual Add。每个独立算子的 launch 开销 ~50μs——12 个算子 × 50μs = 600μs 的 launch 总开销。这个 Block 的计算只需 2ms → launch 占比 30%。

graph-autofusion 的自动融合引擎把这 12 个算子合成了 4 个融合 kernel——launch 开销从 600μs 降到 200μs。

融合引擎的图分析流程

图分析流水线
输入：PyTorch 计算图（torch.fx 或 torch.jit）
  ↓
步骤 1：子图匹配（Pattern Matching）
  扫描计算图，找可融合的子图模式
  ↓
步骤 2：依赖分析（Dependency Analysis）
  检查数据依赖和内存依赖——确保融合后语义不变
  ↓
步骤 3：代价估计（Cost Estimation）
  评估融合后的性能增益——不值得的融合跳过
  ↓
步骤 4：代码生成（Code Generation）
  生成融合后的 Ascend C kernel 代码
  ↓
输出：优化后的计算图（算子数减少 60-80%）

步骤 1：子图匹配

graph-autofusion 内置了上百个融合模式（fuse patterns），用图匹配算法扫描计算图：

# graph-autofusion/tools/fusion_patterns.py

FUSION_PATTERNS = {
    # 模式 1：LayerNorm + Dropout + Linear
    "layernorm_dropout_linear": {
        "nodes": [
            {"op": "layer_norm", "inputs": ["x", "gamma", "beta"]},
            {"op": "dropout", "inputs": ["layernorm_out"], "attrs": {"p": 0.1}},
            {"op": "linear", "inputs": ["dropout_out", "weight"], "attrs": {"bias": True}}
        ],
        "conditions": [
            "layernorm_out.shape == dropout_out.shape",
            "dropout_out.shape[-1] == linear_weight.shape[0]"
        ]
    },

    # 模式 2：Gelu + Linear + Residual Add
    "gelu_linear_residual": {
        "nodes": [
            {"op": "gelu", "inputs": ["x"]},
            {"op": "linear", "inputs": ["gelu_out", "weight"], "attrs": {"bias": True}},
            {"op": "add", "inputs": ["linear_out", "residual"]}
        ],
        "conditions": [
            "gelu_out.shape == linear_out.shape",
            "linear_out.shape == residual.shape"
        ]
    },

    # 模式 3：MatMul + Scale + Softmax + MatMul（Attention 核心）
    "attention_core": {
        "nodes": [
            {"op": "matmul", "inputs": ["Q", "K.T"]},
            {"op": "div", "inputs": ["matmul_out", "scale"]},
            {"op": "softmax", "inputs": ["scale_out"]},
            {"op": "matmul", "inputs": ["softmax_out", "V"]}
        ],
        "conditions": [
            "Q.shape[0] == K.shape[0]",
            "Q.shape[-1] == K.shape[-1]",
            "matmul_out.shape[-1] == V.shape[0]"
        ]
    },

    # 模式 4：LayerNorm + QKV Linear（前向融合到 QKV 投影）
    "layernorm_qkv_linear": {
        "nodes": [
            {"op": "layer_norm", "inputs": ["x", "gamma", "beta"]},
            {"op": "linear", "inputs": ["layernorm_out", "W_qkv"]},
            {"op": "split", "inputs": ["qkv_out"], "attrs": {"splits": [3, "hidden"]}}
        ],
    },
}

步骤 2：依赖分析

融合不只是操作串联——必须保证数据依赖正确：

# graph-autofusion/tools/dependency_analysis.py

def analyze_dependencies(fusion_candidate):
    """检查融合候选：确保融合后语义不变"""

    # 检查 1：没有外部消费者
    for node in fusion_candidate.nodes[:-1]:  # 除最后一个外的所有
        if has_external_consumer(node.output):
            raise FusionError(f"{node.name} 有外部消费者，不能融合")

    # 检查 2：没有内部依赖冲突
    for node in fusion_candidate.nodes:
        for dep in node.dependencies:
            if dep in fusion_candidate.nodes:
                if dep != node.prev:
                    raise FusionError(f"{node.name} 依赖 {dep}，但 {dep} 不在前面")

    # 检查 3：内存别名冲突
    for node in fusion_candidate.nodes:
        if node.output == fusion_candidate.nodes[0].input:
            raise FusionError(f"{node.name} 的输出和输入共享内存，不能融合")

    # 检查 4：动态 shape 冲突
    for node in fusion_candidate.nodes:
        if node.has_dynamic_shape:
            raise FusionError(f"{node.name} 有动态 shape，不能融合")

    return True  # 通过所有检查

步骤 3：代价估计

不是所有融合都有收益——代价估计决定是否融合：

# graph-autofusion/tools/cost_estimation.py

def estimate_fusion_benefit(fusion_candidate):
    """估算融合的收益"""

    # 原始代价（融合前）
    original_launch_cost = len(fusion_candidate.nodes) * 50e-6  # 50μs per launch
    original_mem_read = sum(node.input_size for node in fusion_candidate.nodes)
    original_mem_write = sum(node.output_size for node in fusion_candidate.nodes)

    # 融合后代价
    fused_launch_cost = 50e-6  # 1 次 launch
    fused_mem_read = fusion_candidate.nodes[0].input_size  # 只读一次
    fused_mem_write = fusion_candidate.nodes[-1].output_size  # 只写一次

    # 计算 HBM 带宽节省
    hbm_bandwidth = 900e9  # 900 GB/s
    read_time_original = original_mem_read / hbm_bandwidth
    write_time_original = original_mem_write / hbm_bandwidth
    read_time_fused = fused_mem_read / hbm_bandwidth
    write_time_fused = fused_mem_write / hbm_bandwidth

    hbm_saving = (read_time_original + write_time_original) - (read_time_fused + write_time_fused)
    launch_saving = original_launch_cost - fused_launch_cost

    total_saving = hbm_saving + launch_saving

    # 阈值：收益 > 10μs 才融合（避免无意义的融合）
    if total_saving < 10e-6:
        return None  # 收益太小，不融合

    return {
        "hbm_saving_seconds": hbm_saving,
        "launch_saving_seconds": launch_saving,
        "total_saving_seconds": total_saving,
        "hbm_read_reduction": f"{100 * (1 - fused_mem_read / original_mem_read):.1f}%",
        "hbm_write_reduction": f"{100 * (1 - fused_mem_write / original_mem_write):.1f}%",
    }

步骤 4：代码生成

从融合模式生成 Ascend C kernel：

# graph-autofusion/tools/code_generator.py

def generate_fused_kernel(pattern_name, nodes):
    """从融合模式生成 Ascend C kernel 代码"""

    if pattern_name == "layernorm_qkv_linear":
        return generate_layernorm_qkv_linear(nodes)
    elif pattern_name == "attention_core":
        return generate_attention_core_kernel(nodes)
    # ...

def generate_layernorm_qkv_linear(nodes):
    """生成 LayerNorm + QKV Linear 融合 kernel"""

    kernel_code = """
__aicore__ void LayerNormQKVLinearFused(
    GlobalTensor<float16>& input,      // [batch, seq, hidden]
    GlobalTensor<float16>& gamma,      // [hidden]
    GlobalTensor<float16>& beta,       // [hidden]
    GlobalTensor<float16>& W_qkv,      // [3*hidden, hidden]
    GlobalTensor<float16>& output,     // [batch, seq, 3*hidden]
    int batch, int seq_len, int hidden
) {
    for (int b = blockIdx.x; b < batch * seq_len; b += gridDim.x) {
        // ===== 阶段 1：LayerNorm（仅 L1 计算，不写 HBM）=====
        float mean = 0.0f;
        float M2 = 0.0f;

        // Welford 算法
        for (int h = 0; h < hidden; h += 256) {
            LocalTensor<float16> x_block(256);
            DataCopy(x_block, input[b * hidden + h], 256);

            for (int i = 0; i < 256; i++) {
                float x = float(x_block[i]);
                float delta = x - mean;
                mean += delta / float(h + i + 1);
                float delta2 = x - mean;
                M2 += delta * delta2;
            }
        }

        float inv_std = rsqrtf(M2 / hidden + 1e-5f);

        // 归一化（输出在 L1 中，不写 HBM）
        LocalTensor<float16> normalized(hidden);
        for (int h = 0; h < hidden; h += 256) {
            // ... 归一化：normalized = (x - mean) * inv_std * gamma + beta
        }

        // ===== 阶段 2：QKV Linear（直接在 L1 中的 normalized 上算）=====
        // QKV 投影：output = normalized @ W_qkv^T
        // W_qkv shape: [3*hidden, hidden]
        // output shape: [3*hidden]（每个 token 的 QKV）

        for (int qkv = 0; qkv < 3; qkv++) {
            int offset = qkv * hidden;

            for (int o = 0; o < hidden; o += 64) {
                float accum = 0.0f;
                for (int i = 0; i < hidden; i++) {
                    accum += float(normalized[i]) * float(W_qkv[offset + o]);
                }
                output[b * 3 * hidden + offset + o] = float16(accum);
            }
        }
    }
}
"""
    return kernel_code

生成的 kernel 会编译成 NPU 可执行的代码——图层面自动完成，开发者无需手动写融合。

Transformer Block 的完整融合案例

一个标准 Transformer Block 的 12 个算子，经过 graph-autofusion 自动融合后：

原始图（12 个算子）
LayerNorm → QKV_Linear → Reshape → Transpose → Attn_MatMul →
Scale → Softmax → Attn_MatMul2 → Concat → Out_Linear →
LayerNorm2 → Gelu → FFN_Up → FFN_Down → Residual_Add

融合后图（4 个融合 kernel + 2 个独立算子）
┌─ Fused_1：LayerNorm + QKV_Linear + Reshape + Transpose
├─ Fused_2：Attn_MatMul + Scale + Softmax + Attn_MatMul2 + Concat
├─ Fused_3：LayerNorm2 + Gelu + FFN_Up + FFN_Down
├─ Fused_4：Out_Linear + Residual_Add

独立算子：
├─ TokenEmbedding（图开头，不属于 Block）
└─ LM_Head（图末尾）

Launch 开销：12 × 50μs = 600μs → 4 × 50μs + 2 × 50μs = 300μs
HBM 读写：12 次输入 + 12 次输出 ≈ 24MB → 4 次输入 + 4 次输出 ≈ 8MB

踩坑一：融合导致中间结果不可调试

12 个算子融合成 4 个 kernel → 中间结果 “消失” 了。调试时看不到 Reshape 后的形状、Softmax 前的值、Gelu 的输出——这些都只在生成的 kernel 里出现。

缓解：添加 debug 模式

# 设置环境变量启用 debug
# export GF_AUTOFUSION_DEBUG=1  # 关闭融合（所有算子独立运行）

# 或选择性禁用个别融合
# export GF_AUTOFUSION_DISABLE="layernorm_qkv_linear,attention_core"

踩坑二：融合过度导致 L1 溢出

12 个算子融合成 4 个 kernel——每个 kernel 要在 L1 中存更多的中间变量。当 hidden=8192（LLaMA 3.1 的 hidden 维度），LayerNorm 的 normalized 中间量是 8192×4 = 32KB → 刚好填满 L1。加上 QKV Linear 的中间结果，L1 溢出到 HBM → 性能不升反降。

修复：代价估计中检查 L1 使用情况

def estimate_l1_usage(fusion_candidate):
    """估计融合后的 L1 使用量"""
    total_l1 = sum(node.l1_footprint for node in fusion_candidate.nodes)
    l1_capacity = 32 * 1024  # 32KB

    if total_l1 > l1_capacity * 0.8:  # 80% 阈值
        raise FusionError(f"融合后 L1 使用 {total_l1} > {l1_capacity*0.8}，跳过融合")
    return total_l1

踩坑三：融合改变计算顺序导致精度差异

标准 Attn Scale + Softmax：先 Div（除以 scale），再 Softmax。融合后：先 Softmax（带 temperature），隐含 Scale。多个中间结果经过 FP16 截断——融合后精度可能下降。

问题：Div + Softmax 的 FP32 精度 vs Fused Softmax 的 FP16 精度。

缓解：融合后内部用 FP32 计算，只在最后一步转 FP16

// 融合 kernel 内用 FP32
float scale = 1.0f / sqrtf(head_dim);
for (int i = 0; i < seq_len; i++) {
    float x = float(scores[i]) * scale;  // FP32 scale（内部用 FP32）
    float exp_val = expf(x - max_val_fp32);
    output[i] = float16(exp_val / sum_exp_fp32);  // 最后才转 FP16
}

---

graph-autofusion 的价值在于"自动"——不需要开发者手动写融合 kernel。图的 Pattern Matching → Dependency Analysis → Cost Estimation → Code Generation 四个步骤全自动完成。12 个算子的 Transformer Block → 4 个融合 kernel，HBM 读写减少 67%，launch 开销减半。代价是中间结果的不可见性和 L1 溢出的风险——付出这些代价，换回了更低延迟。