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 核心参数按算子类型选择参数类型默认值含义适用算子类型enable_flattenboolNone启用 IR 展平优化将多维循环展平为一维提升向量化效率纯 Vector: True; CV 融合: Falsemultibufferbool非910_95时 True启用 ping-pong 多缓冲流水线在计算当前数据的同时预取下一批数据隐藏内存延迟通用enable_auto_bind_sub_blockboolNone启用自动绑定子块sub-block将 Vector 核心划分为多个子块并行执行CV 融合: True; 纯 Vector: Falsesync_solverboolNone启用同步求解器自动求解 Cube-Vector 间的同步点位置CV 融合: Trueset_workspace_multibufferintNone设置 workspace 多缓冲数量为中间计算结果分配多份缓冲区CV 融合: 2limit_auto_multi_buffer_of_local_bufferstrNone限制本地缓冲UB的多缓冲策略。no-limit表示不限制CV 融合: no-limitenable_mixed_cvboolNone启用混合 CV 模式允许 Cube 和 Vector 在同一 kernel 中混合执行CV 融合: True1.2 辅助参数参数类型默认值含义副作用enable_ubuf_savingboolNone启用 UB 节省优化A2/A3 平台减少 UB 占用可能降低计算并行度enable_hivm_auto_cv_balanceboolNone启用 HIVM 自动 CV 负载均衡编译时间增加inject_barrier_allboolNone在所有操作间注入屏障同步性能可能下降用于调试inject_block_allboolNone在所有块间注入同步性能可能下降用于调试disable_auto_inject_block_syncboolNone禁用自动注入块同步可能导致数据竞争enable_vf_fusionboolFalse启用 VFVector Function融合UB 占用增加limit_auto_multi_buffer_only_for_local_bufferboolNone限制多缓冲仅用于本地缓冲UB减少全局缓冲开销enable_cce_vf_auto_syncboolNone启用 CCE VF 自动同步编译时间增加enable_cce_vf_remove_membarboolNone启用 CCE VF 移除内存屏障可能导致同步问题disable_size_align_for_castboolNone禁用类型转换的大小对齐要求可能产生非对齐访问tile_mix_vector_loopintNone混合 Vector 循环 tiling 大小A2/A3影响 UB 占用tile_mix_cube_loopintNone混合 Cube 循环 tiling 大小A2/A3影响 L1 占用unit_flagboolNone启用单元标志同步调试用1.3 编译模式参数参数类型默认值含义编译路径compile_modestrsimd编译模式选择见下表num_warpsint4Warp 数量SIMD 模式影响 HFusion 向量化策略SIMT 模式控制线程数通用num_stagesint1流水线阶段数NPU 上通常保持 1通用auto_blockify_sizeint1AutoBlockify 分块大小配合TRITON_ALL_BLOCKS_PARALLEL使用通用add_auto_schedulingboolFalse启用自动调度DAG 亲和性优化SIMT/混合模式enable_bishengir_simt_optimizationint000SIMT 优化控制位位模式SIMT 模式compile_mode 详解值行为编译路径simd默认模式parallel_mode设为simdLinalg → HFusion → HIVM → Binaryunstructured_in_simt非结构化转 SIMT自动设force_simt_templateTrueSIMD SIMT 混合路径simt_only纯 SIMT自动设force_simt_onlyTrue、parallel_modesimtTTIR → TTGIR → LLVM → Binary1.4 精度参数参数类型默认值含义default_dot_input_precisionstrieeeDot 操作默认输入精度allowed_dot_input_precisionstuple(ieee, hf32)允许的 Dot 输入精度enable_fp_fusionboolTrue启用浮点融合FMAdisable_fmaboolFalse禁用 FMA提高精度但降低性能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, maskmask) y tl.load(y_ptr offsets, maskmask) tl.store(out_ptr offsets, x y, maskmask) 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_SIZE1024, enable_flattenTrue, multibufferTrue,?utm_sourcegitcode_repo_files)参数说明enable_flattenTrue纯 Vector 算子无 Cube 参与展平优化可提升向量化效率multibufferTrue启用 ping-pong 流水线隐藏内存延迟无需enable_auto_bind_sub_block、sync_solver、enable_mixed_cv等 CV 相关参数2.2 CV 融合算子CV 融合算子同时使用 Cube矩阵计算和 Vector向量计算核心需要 Cube-Vector 协同调度和同步。典型场景Flash Attention、融合矩阵乘法后处理、MatMulReLU量化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_M128, BLOCK_N32, enable_auto_bind_sub_blockTrue, enable_flattenFalse, set_workspace_multibuffer2, sync_solverTrue, limit_auto_multi_buffer_of_local_bufferno-limit, multibufferTrue, enable_mixed_cvTrue,?utm_sourcegitcode_repo_files)参数说明enable_auto_bind_sub_blockTrue将 Vector 核心划分为子块与 Cube 并行执行enable_flattenFalseCV 融合算子不能展平因为 Cube 和 Vector 需要各自独立的循环结构set_workspace_multibuffer2为中间结果分配双缓冲配合 ping-pong 流水线sync_solverTrue自动求解 Cube-Vector 间的同步点避免手动插入同步limit_auto_multi_buffer_of_local_bufferno-limit不限制 UB 的多缓冲分配避免 UB 不足导致编译失败multibufferTrue启用多缓冲流水线enable_mixed_cvTrue启用混合 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), dtypetl.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_M128, BLOCK_SIZE_N128, BLOCK_SIZE_K128, GROUP_SIZE_M8, enable_auto_bind_sub_blockTrue, set_workspace_multibuffer2, sync_solverTrue, limit_auto_multi_buffer_of_local_bufferno-limit, multibufferTrue, enable_flattenTrue,?utm_sourcegitcode_repo_files)参数说明矩阵乘法前向虽然主体是 Cube 操作但 epilogue类型转换、store在 Vector 上执行enable_flattenTrueMatMul 前向的 Vector 后处理部分可以展平enable_auto_bind_sub_blockTrueCube 和 Vector 协同需要子块绑定sync_solverTrue自动求解 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), dtypetl.float32) for row_idx in range(0, tl.cdiv(M, BLOCK_SIZE_M)): dy tl.load(dy_ptrs, maskmask, other0.0).to(tl.float32) sum_b dy dy_ptrs BLOCK_SIZE_M * N tl.store(db_ptr col_off, tl.sum(sum_b, 1), maskcol_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_blockFalse,?utm_sourcegitcode_repo_files)参数说明enable_auto_bind_sub_blockFalse纯 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_sourcegitcode_repo_files# kernel 的位置参数 BLOCK_SIZE1024, # constexpr 参数 enable_flattenTrue, # 编译参数 multibufferTrue, # 编译参数 sync_solverTrue, # 编译参数)编译参数与 kernel 参数混合传入Triton 会自动区分属于NPUOptions字段的被识别为编译参数其余为 kernel 参数。3.2 在 autotune Config 中传入通过triton.Config的关键字参数传入autotune 会为每个配置分别编译triton.autotune( configs[ triton.Config( {BLOCK_M: 128, BLOCK_N: 32}, multibufferTrue, enable_mixed_cvTrue, enable_auto_bind_sub_blockTrue, sync_solverTrue, enable_flattenFalse, set_workspace_multibuffer2, limit_auto_multi_buffer_of_local_bufferno-limit, ), triton.Config( {BLOCK_M: 64, BLOCK_N: 64}, multibufferTrue, enable_mixed_cvTrue, enable_auto_bind_sub_blockTrue, sync_solverTrue, enable_flattenFalse, set_workspace_multibuffer2, limit_auto_multi_buffer_of_local_bufferno-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_blockTrue, enable_flattenFalse, set_workspace_multibuffer2, sync_solverTrue, limit_auto_multi_buffer_of_local_bufferno-limit, multibufferTrue, enable_mixed_cvTrue 2. 去掉 enable_mixed_cv禁用混合 CV enable_auto_bind_sub_blockTrue, enable_flattenFalse, set_workspace_multibuffer2, sync_solverTrue, limit_auto_multi_buffer_of_local_bufferno-limit, multibufferTrue 3. 去掉 enable_auto_bind_sub_block禁用子块绑定 enable_flattenFalse, multibufferTrue, sync_solverTrue 4. 去掉 sync_solver禁用同步求解器 enable_flattenFalse, multibufferTrue 5. 最小参数集最大兼容性 multibufferTrue4.2 精度问题回退1. 启用 FMA → disable_fmaTrue禁用 FMA 提高精度 2. enable_fp_fusionTrue → enable_fp_fusionFalse禁用浮点融合 3. default_dot_input_precisionieee确保 IEEE 精度4.3 UB 溢出回退当出现UB overflow错误时1. 减小 BLOCK_M / BLOCK_N减少每个 block 的数据量 2. set_workspace_multibuffer2 → 1减少 workspace 缓冲 3. limit_auto_multi_buffer_of_local_bufferno-limit → 限制多缓冲 4. enable_flattenFalse展平可能增加 UB 占用 5. enable_ubuf_savingTrueA2/A3 平台启用 UB 节省5. 910_95 特别注意5.1 平台检测910_95 平台通过is_compile_on_910_95自动检测NPUOptions.compile_on_910_95默认值即为检测结果。检测逻辑get_ascend_devices.pyPCI 设备 ID 包含0xd806或npu-smi info输出包含ascend910_95/ascend950/910_958b5.2 关键差异特性A2/A3 (910B)910_95multibuffer默认值TrueFalseUB 大小192 KB256 KBRF 大小无128 KBFFTS 支持支持不支持自动禁用fixpipe L0C→UB不支持支持copy(UB→UB/L1)不支持支持Vector 核心数 Cube 核心数 Cube 核心数 x 2shared_mem_dynamic_sizeSIMT2211841228805.3 910_95 上的参数调整建议multibuffer 需显式开启910_95 默认multibufferFalse需要手动设为True以获得最佳性能UB 空间更大256KB UB 允许更大的 BLOCK_SIZE但需配合limit_auto_multi_buffer_of_local_bufferno-limit使用Vector 核心数翻倍910_95 的 Vector 核心数是 Cube 的 2 倍grid 可以使用更多核心fixpipe 可直达 UB利用fixpipe将 Cube 结果从 L0C 直接搬运到 UB实现零拷贝融合FFTS 自动禁用910_95 不支持 FFTS编译器会自动跳过 FFTS 相关 Pass5.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_M128, BLOCK_N32, multibufferTrue, ?utm_sourcegitcode_repo_files# 910_95 默认 False需显式开启 enable_mixed_cvTrue, enable_auto_bind_sub_blockTrue, sync_solverTrue, limit_auto_multi_buffer_of_local_bufferno-limit, enable_flattenFalse, set_workspace_multibuffer2,)6. 与 GPU 编译参数对比维度GPU (CUDA)NPU (Ascend)流水线控制num_stages2~4软件流水线阶段数multibufferTrueping-pong 多缓冲并行度控制num_warps4/8/16warp 数量num_warps4影响向量化策略非线程数概念线程级并行GPU 天然 SIMTcompile_modesimt_onlynum_warps8/16Cube-Vector 协同无GPU 无 Cube/Vector 分离enable_mixed_cvTruesync_solverTrue子块绑定无enable_auto_bind_sub_blockTrue展平优化无enable_flattenTrue/Falseworkspace 缓冲无set_workspace_multibuffer2UB 限制无使用 Shared Memorylimit_auto_multi_buffer_of_local_bufferno-limit精度控制tf32/ieee/tf32x3ieee/hf32关键区别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 AttentionCV 融合算子来源flash_attention_npu_v8.py前向 kernelCV 融合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_DIMqk_dim, V_DIMv_dim, ..., multibufferTrue, enable_mixed_cvTrue, enable_auto_bind_sub_blockTrue, sync_solverTrue, limit_auto_multi_buffer_of_local_bufferno-limit, enable_flattenFalse, set_workspace_multibuffer2,?utm_sourcegitcode_repo_files)反向预处理 kernel纯 Vectorbwd_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, ..., multibufferTrue, limit_auto_multi_buffer_of_local_bufferno-limit,?utm_sourcegitcode_repo_files)反向 QKV kernelCV 融合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_bufferno-limit, enable_flattenFalse, sync_solverTrue, enable_mixed_cvTrue,?utm_sourcegitcode_repo_files)7.2 融合矩阵乘法CV 融合算子来源fused_matmul_npu_v3.py前向 kernelMatMul BiasCV 融合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_BIAShas_bias, enable_auto_bind_sub_blockTrue, set_workspace_multibuffer2, sync_solverTrue, limit_auto_multi_buffer_of_local_bufferno-limit, multibufferTrue, enable_flattenTrue,?utm_sourcegitcode_repo_files)反向 bias 梯度 kernel纯 Vectorfused_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_blockFalse,?utm_sourcegitcode_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_blockFalse,?utm_sourcegitcode_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_blockFalse,?utm_sourcegitcode_repo_files)8. 快速决策流程kernel 是否包含 tl.dot() ├── 否 → 纯 Vector 算子 │ └── enable_flattenTrue, multibufferTrue │ 无需 CV 相关参数 │ └── 是 → 是否有 Cube 后的 Vector 处理如 softmax、mask、激活 ├── 否 → 纯 Cube 算子MatMul 前向 │ └── enable_auto_bind_sub_blockTrue, enable_flattenTrue, │ set_workspace_multibuffer2, sync_solverTrue, │ limit_auto_multi_buffer_of_local_bufferno-limit, │ multibufferTrue │ └── 是 → CV 融合算子Flash Attention、融合 MatMul后处理 └── enable_auto_bind_sub_blockTrue, enable_flattenFalse, set_workspace_multibuffer2, sync_solverTrue, limit_auto_multi_buffer_of_local_bufferno-limit, multibufferTrue, enable_mixed_cvTrue9. 相关文档链接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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
