CANN/cannbot-skills NPU编译参数速查

NPU 编译参数速查文档

【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体，本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills

触发条件：Agent 需要为 kernel 配置编译参数时查阅本文档

1. 编译参数分类速查表

NPU 编译参数通过kernel[grid](https://link.gitcode.com/i/9c429df1ca2567ee7e7f5434ee4fa5ba)调用时以关键字参数传入，底层映射到NPUOptions数据类的各字段。与 GPU 使用num_stages/num_warps控制流水线和并行度不同，NPU 使用一组专门的参数来控制 Cube-Vector 协同、多缓冲流水线和编译路径。

1.1 核心参数（按算子类型选择）

1.2 辅助参数

1.3 编译模式参数

compile_mode 详解：

1.4 精度参数

2. 不同算子类型的参数配置示例

2.1 纯 Vector 算子

纯 Vector 算子只使用 Vector 核心进行逐元素运算、归约等操作，不涉及 Cube（矩阵乘法）。

典型场景：向量加法、LayerNorm、Softmax、激活函数、逐元素运算

@triton.jit def vector_add_kernel(x_ptr, y_ptr, out_ptr, N, BLOCK_SIZE: tl.constexpr): pid = tl.program_id(0) offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE) mask = offsets < N x = tl.load(x_ptr + offsets, mask=mask) y = tl.load(y_ptr + offsets, mask=mask) tl.store(out_ptr + offsets, x + y, mask=mask) grid = (triton.cdiv(N, BLOCK_SIZE),) vector_add_kernel[grid](https://gitcode.com/cann/cannbot-skills/blob/eb07a8e06215aa74655bac6d14b86eadbafeb318/ops/triton-latency-optimizer/references/docs_triton_IR/docs_for_triton_agent/x, y, out, N, BLOCK_SIZE=1024, enable_flatten=True, multibuffer=True,?utm_source=gitcode_repo_files)

参数说明：

• enable_flatten=True：纯 Vector 算子无 Cube 参与，展平优化可提升向量化效率
• multibuffer=True：启用 ping-pong 流水线隐藏内存延迟
• 无需enable_auto_bind_sub_block、sync_solver、enable_mixed_cv等 CV 相关参数

2.2 CV 融合算子

CV 融合算子同时使用 Cube（矩阵计算）和 Vector（向量计算）核心，需要 Cube-Vector 协同调度和同步。

典型场景：Flash Attention、融合矩阵乘法+后处理、MatMul+ReLU+量化

@triton.jit def flash_attn_fwd_kernel(q_ptr, k_ptr, v_ptr, o_ptr, ..., BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, ...): ... for start_n in range(begin, end, BLOCK_N): k = tl.load(k_block_ptr, ...) s = tl.dot(q, k) s = s * scale + tl.where(mask, 0.0, -2.0**30) m_new = tl.maximum(m, tl.max(s, 1)) p = tl.math.exp(s - m_new[:, None]) v = tl.load(v_block_ptr, ...) pv = tl.dot(p.to(dtype), v) ... grid = (AICORE_NUM,) flash_attn_fwd_kernel[grid](https://gitcode.com/cann/cannbot-skills/blob/eb07a8e06215aa74655bac6d14b86eadbafeb318/ops/triton-latency-optimizer/references/docs_triton_IR/docs_for_triton_agent/q, k, v, o, ..., BLOCK_M=128, BLOCK_N=32, enable_auto_bind_sub_block=True, enable_flatten=False, set_workspace_multibuffer=2, sync_solver=True, limit_auto_multi_buffer_of_local_buffer="no-limit", multibuffer=True, enable_mixed_cv=True,?utm_source=gitcode_repo_files)

参数说明：

• enable_auto_bind_sub_block=True：将 Vector 核心划分为子块，与 Cube 并行执行
• enable_flatten=False：CV 融合算子不能展平，因为 Cube 和 Vector 需要各自独立的循环结构
• set_workspace_multibuffer=2：为中间结果分配双缓冲，配合 ping-pong 流水线
• sync_solver=True：自动求解 Cube-Vector 间的同步点，避免手动插入同步
• limit_auto_multi_buffer_of_local_buffer="no-limit"：不限制 UB 的多缓冲分配，避免 UB 不足导致编译失败
• multibuffer=True：启用多缓冲流水线
• enable_mixed_cv=True：启用混合 CV 模式，允许 Cube 和 Vector 交替执行

2.3 纯 Cube 算子（矩阵乘法前向）

@triton.jit def matmul_fwd_kernel(x_ptr, w_ptr, y_ptr, M, N, K, BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr, GROUP_SIZE_M: tl.constexpr): ... accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32) for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)): x = tl.load(x_ptrs, ...) w = tl.load(w_ptrs, ...) accumulator = tl.dot(x, w, accumulator) ... grid = lambda META: (triton.cdiv(M, META["BLOCK_SIZE_M"]) * triton.cdiv(N, META["BLOCK_SIZE_N"]),) matmul_fwd_kernel[grid](https://gitcode.com/cann/cannbot-skills/blob/eb07a8e06215aa74655bac6d14b86eadbafeb318/ops/triton-latency-optimizer/references/docs_triton_IR/docs_for_triton_agent/x, w, y, M, N, K, BLOCK_SIZE_M=128, BLOCK_SIZE_N=128, BLOCK_SIZE_K=128, GROUP_SIZE_M=8, enable_auto_bind_sub_block=True, set_workspace_multibuffer=2, sync_solver=True, limit_auto_multi_buffer_of_local_buffer="no-limit", multibuffer=True, enable_flatten=True,?utm_source=gitcode_repo_files)

参数说明：

• 矩阵乘法前向虽然主体是 Cube 操作，但 epilogue（类型转换、store）在 Vector 上执行
• enable_flatten=True：MatMul 前向的 Vector 后处理部分可以展平
• enable_auto_bind_sub_block=True：Cube 和 Vector 协同需要子块绑定
• sync_solver=True：自动求解 Cube→Vector 的同步点

2.4 纯 Vector 后处理算子（如 bias 梯度计算）

@triton.jit def bwd_b_kernel(dy_ptr, db_ptr, M, N, BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr): ... sum_b = tl.zeros((BLOCK_SIZE_N, BLOCK_SIZE_M), dtype=tl.float32) for row_idx in range(0, tl.cdiv(M, BLOCK_SIZE_M)): dy = tl.load(dy_ptrs, mask=mask, other=0.0).to(tl.float32) sum_b += dy dy_ptrs += BLOCK_SIZE_M * N tl.store(db_ptr + col_off, tl.sum(sum_b, 1), mask=col_off < N) grid = (triton.cdiv(N, BLOCK_SIZE_N),) bwd_b_kernel[grid](https://gitcode.com/cann/cannbot-skills/blob/eb07a8e06215aa74655bac6d14b86eadbafeb318/ops/triton-latency-optimizer/references/docs_triton_IR/docs_for_triton_agent/dy, db, M, N, enable_auto_bind_sub_block=False,?utm_source=gitcode_repo_files)

参数说明：

• enable_auto_bind_sub_block=False：纯 Vector 算子不需要子块绑定
• 无需 CV 相关参数

3. 参数传递方式

3.1 在 kernelgrid 调用时传入

这是最常用的方式，编译参数作为关键字参数直接传入 kernel 启动调用：

kernel[grid](https://gitcode.com/cann/cannbot-skills/blob/eb07a8e06215aa74655bac6d14b86eadbafeb318/ops/triton-latency-optimizer/references/docs_triton_IR/docs_for_triton_agent/arg1, arg2, ..., ?utm_source=gitcode_repo_files# kernel 的位置参数 BLOCK_SIZE=1024, # constexpr 参数 enable_flatten=True, # 编译参数 multibuffer=True, # 编译参数 sync_solver=True, # 编译参数)

编译参数与 kernel 参数混合传入，Triton 会自动区分：属于NPUOptions字段的被识别为编译参数，其余为 kernel 参数。

3.2 在 autotune Config 中传入

通过triton.Config的关键字参数传入，autotune 会为每个配置分别编译：

@triton.autotune( configs=[ triton.Config( {"BLOCK_M": 128, "BLOCK_N": 32}, multibuffer=True, enable_mixed_cv=True, enable_auto_bind_sub_block=True, sync_solver=True, enable_flatten=False, set_workspace_multibuffer=2, limit_auto_multi_buffer_of_local_buffer="no-limit", ), triton.Config( {"BLOCK_M": 64, "BLOCK_N": 64}, multibuffer=True, enable_mixed_cv=True, enable_auto_bind_sub_block=True, sync_solver=True, enable_flatten=False, set_workspace_multibuffer=2, limit_auto_multi_buffer_of_local_buffer="no-limit", ), ], key=["QK_DIM", "V_DIM"], ) @triton.jit def my_kernel(...): ...

3.3 通过 compile() 函数直接传入

compiled_kernel = triton.compile( kernel, options={ "compile_mode": "simt_only", "num_warps": 8, "enable_flatten": True, "multibuffer": True, } )

4. 精度回退策略

当 kernel 编译失败或运行时出现精度问题时，应按照以下策略逐步回退：

4.1 编译失败回退

1. 完整 CV 融合参数（最高性能） enable_auto_bind_sub_block=True, enable_flatten=False, set_workspace_multibuffer=2, sync_solver=True, limit_auto_multi_buffer_of_local_buffer="no-limit", multibuffer=True, enable_mixed_cv=True 2. 去掉 enable_mixed_cv（禁用混合 CV） enable_auto_bind_sub_block=True, enable_flatten=False, set_workspace_multibuffer=2, sync_solver=True, limit_auto_multi_buffer_of_local_buffer="no-limit", multibuffer=True 3. 去掉 enable_auto_bind_sub_block（禁用子块绑定） enable_flatten=False, multibuffer=True, sync_solver=True 4. 去掉 sync_solver（禁用同步求解器） enable_flatten=False, multibuffer=True 5. 最小参数集（最大兼容性） multibuffer=True

4.2 精度问题回退

1. 启用 FMA → disable_fma=True（禁用 FMA 提高精度） 2. enable_fp_fusion=True → enable_fp_fusion=False（禁用浮点融合） 3. default_dot_input_precision="ieee"（确保 IEEE 精度）

4.3 UB 溢出回退

当出现UB overflow错误时：

1. 减小 BLOCK_M / BLOCK_N（减少每个 block 的数据量） 2. set_workspace_multibuffer=2 → 1（减少 workspace 缓冲） 3. limit_auto_multi_buffer_of_local_buffer="no-limit" → 限制多缓冲 4. enable_flatten=False（展平可能增加 UB 占用） 5. enable_ubuf_saving=True（A2/A3 平台启用 UB 节省）

5. 910_95 特别注意

5.1 平台检测

910_95 平台通过is_compile_on_910_95自动检测，NPUOptions.compile_on_910_95默认值即为检测结果。

检测逻辑（get_ascend_devices.py）：

• PCI 设备 ID 包含0xd806
• 或npu-smi info输出包含ascend910_95/ascend950/910_958b

5.2 关键差异

5.3 910_95 上的参数调整建议

1. multibuffer 需显式开启：910_95 默认multibuffer=False，需要手动设为True以获得最佳性能
2. UB 空间更大：256KB UB 允许更大的 BLOCK_SIZE，但需配合limit_auto_multi_buffer_of_local_buffer="no-limit"使用
3. Vector 核心数翻倍：910_95 的 Vector 核心数是 Cube 的 2 倍，grid 可以使用更多核心
4. fixpipe 可直达 UB：利用fixpipe将 Cube 结果从 L0C 直接搬运到 UB，实现零拷贝融合
5. FFTS 自动禁用：910_95 不支持 FFTS，编译器会自动跳过 FFTS 相关 Pass

5.4 910_95 上的典型配置

# Flash Attention (910_95) fwd_kernel[grid](https://gitcode.com/cann/cannbot-skills/blob/eb07a8e06215aa74655bac6d14b86eadbafeb318/ops/triton-latency-optimizer/references/docs_triton_IR/docs_for_triton_agent/q, k, v, o, l, ..., BLOCK_M=128, BLOCK_N=32, multibuffer=True, ?utm_source=gitcode_repo_files# 910_95 默认 False，需显式开启 enable_mixed_cv=True, enable_auto_bind_sub_block=True, sync_solver=True, limit_auto_multi_buffer_of_local_buffer="no-limit", enable_flatten=False, set_workspace_multibuffer=2,)

6. 与 GPU 编译参数对比

关键区别：

• GPU 的num_stages控制软件流水线深度，NPU 的multibuffer控制 ping-pong 缓冲
• GPU 的num_warps直接控制线程并行度，NPU 的num_warps主要影响编译器的向量化策略
• NPU 特有的 Cube-Vector 分离架构需要enable_mixed_cv、sync_solver、enable_auto_bind_sub_block等参数来协调双核
• NPU 的 UB 空间有限，需要通过limit_auto_multi_buffer_of_local_buffer等参数精细控制

7. 实际代码中的参数使用参考

7.1 Flash Attention（CV 融合算子）

来源：flash_attention_npu_v8.py

前向 kernel（CV 融合）：

fwd_kernel[grid](https://gitcode.com/cann/cannbot-skills/blob/eb07a8e06215aa74655bac6d14b86eadbafeb318/ops/triton-latency-optimizer/references/docs_triton_IR/docs_for_triton_agent/q, k, v, o, l, ..., QK_DIM=qk_dim, V_DIM=v_dim, ..., multibuffer=True, enable_mixed_cv=True, enable_auto_bind_sub_block=True, sync_solver=True, limit_auto_multi_buffer_of_local_buffer="no-limit", enable_flatten=False, set_workspace_multibuffer=2,?utm_source=gitcode_repo_files)

反向预处理 kernel（纯 Vector）：

bwd_preprocess_ifmn[grid](https://gitcode.com/cann/cannbot-skills/blob/eb07a8e06215aa74655bac6d14b86eadbafeb318/ops/triton-latency-optimizer/references/docs_triton_IR/docs_for_triton_agent/o, do, d, ..., multibuffer=True, limit_auto_multi_buffer_of_local_buffer="no-limit",?utm_source=gitcode_repo_files)

反向 QKV kernel（CV 融合）：

bwd_qkv_kernel[grid](https://gitcode.com/cann/cannbot-skills/blob/eb07a8e06215aa74655bac6d14b86eadbafeb318/ops/triton-latency-optimizer/references/docs_triton_IR/docs_for_triton_agent/q, k, v, dq, dk, dv, ..., limit_auto_multi_buffer_of_local_buffer="no-limit", enable_flatten=False, sync_solver=True, enable_mixed_cv=True,?utm_source=gitcode_repo_files)

7.2 融合矩阵乘法（CV 融合算子）

来源：fused_matmul_npu_v3.py

前向 kernel（MatMul + Bias，CV 融合）：

fused_matmul_fwd_kernel[grid](https://gitcode.com/cann/cannbot-skills/blob/eb07a8e06215aa74655bac6d14b86eadbafeb318/ops/triton-latency-optimizer/references/docs_triton_IR/docs_for_triton_agent/x, w, b, y, M, N, K, HAS_BIAS=has_bias, enable_auto_bind_sub_block=True, set_workspace_multibuffer=2, sync_solver=True, limit_auto_multi_buffer_of_local_buffer="no-limit", multibuffer=True, enable_flatten=True,?utm_source=gitcode_repo_files)

反向 bias 梯度 kernel（纯 Vector）：

fused_matmul_bwd_b_kernel[grid](https://gitcode.com/cann/cannbot-skills/blob/eb07a8e06215aa74655bac6d14b86eadbafeb318/ops/triton-latency-optimizer/references/docs_triton_IR/docs_for_triton_agent/dy, db, M, N, enable_auto_bind_sub_block=False,?utm_source=gitcode_repo_files)

反向 x 梯度 kernel（纯 Vector 矩阵乘法）：

fused_matmul_bwd_x_kernel[grid](https://gitcode.com/cann/cannbot-skills/blob/eb07a8e06215aa74655bac6d14b86eadbafeb318/ops/triton-latency-optimizer/references/docs_triton_IR/docs_for_triton_agent/dy, w, dx, M, N, K, enable_auto_bind_sub_block=False,?utm_source=gitcode_repo_files)

反向 w 梯度 kernel（纯 Vector 矩阵乘法）：

fused_matmul_bwd_w_kernel[grid](https://gitcode.com/cann/cannbot-skills/blob/eb07a8e06215aa74655bac6d14b86eadbafeb318/ops/triton-latency-optimizer/references/docs_triton_IR/docs_for_triton_agent/dy, x, dw, lock_w, M, N, K, enable_auto_bind_sub_block=False,?utm_source=gitcode_repo_files)

8. 快速决策流程

kernel 是否包含 tl.dot()？ ├── 否 → 纯 Vector 算子 │ └── enable_flatten=True, multibuffer=True │ （无需 CV 相关参数） │ └── 是 → 是否有 Cube 后的 Vector 处理（如 softmax、mask、激活）？ ├── 否 → 纯 Cube 算子（MatMul 前向） │ └── enable_auto_bind_sub_block=True, enable_flatten=True, │ set_workspace_multibuffer=2, sync_solver=True, │ limit_auto_multi_buffer_of_local_buffer="no-limit", │ multibuffer=True │ └── 是 → CV 融合算子（Flash Attention、融合 MatMul+后处理） └── enable_auto_bind_sub_block=True, enable_flatten=False, set_workspace_multibuffer=2, sync_solver=True, limit_auto_multi_buffer_of_local_buffer="no-limit", multibuffer=True, enable_mixed_cv=True

9. 相关文档链接

• 01-extension-overview.md - Ascend 扩展 API 总览
• 02-pipe-and-core.md - PIPE/CORE 枚举详解
• 07-compile-options.md - 编译选项完整参考
• compiler.py - NPUOptions 数据类定义
• get_ascend_devices.py - 910_95 平台检测逻辑
• runtime/utils.py - NPU 运行时参数（核心数、UB 大小等）

【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体，本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills