企业官网建设流程全解析

• 编译器
• 人工智能
• Ascend
• CANN

本文围绕 AscendNPU-IR 官方 FAQ 展开，系统梳理构建安装、运行调试、性能调优、精度调试与社区贡献五大场景中开发者最常遇到的问题及解决方案，并结合仓库源码（如 build.sh、PassManagerOptions.cpp、Passes.td）给出底层原理佐证。读完本文，你将掌握 build 脚本关键参数、LIT 测试与单 Pass 调试方法、bishengir-compile设备侧编译流程、tl.device_print精度定位手段以及 PR 合入规范，能够在昇腾 NPU 算子开发与调优中快速排障。

构建与安装阶段的常见问题

问题一：ninja: error: loading 'build.ninja': No such file or directory

在首次构建或增量构建失败后，如果直接再次执行build-tools/build.sh，可能会遇到上述错误。原因是构建目录中没有可用的build.ninja文件（例如 CMake 配置未生成或缓存被清理）。

解决方案：调用build-tools/build.sh时加上-r（rebuild）选项，强制重新执行 CMake 生成并刷新build.ninja：

./build-tools/build.sh -o ./build -r --build-type Debug

从 build.sh 源码可以看到该选项的底层行为：parse_arguments中-r|--rebuild会将REBUILD置为1（见 build.sh），随后在main()中执行clean_build_dir（清理CMake*缓存文件）并重新调用cmake_generate（见 build.sh），从而保证重新走一遍完整的 CMake 配置流程。

该脚本还支持一系列其他常用选项，便于按需定制构建：

问题二：构建报错 "Too many open files"

该错误表示进程同时打开的文件句柄数超过了系统限制。MLIR 项目构建涉及大量头文件与编译单元，容易触发该上限。

解决方案：通过ulimit -n提高单进程可同时打开的文件数量上限：

ulimit -n 65535

建议在构建前于当前 shell 会话中执行，然后再运行构建脚本。

问题三：The CMAKE_CXX_COMPILER: clang++ is not a full path and was not found in the PATH

该报错说明 CMake 未找到指定的 C++ 编译器，通常有两种原因：未显式指定编译器路径，或编译器二进制损坏/缺失。

解决方案（按优先级尝试）：

1. 显式指定编译器路径：

./build-tools/build.sh --cxx-compiler=/usr/bin/clang++ --c-compiler=/usr/bin/clang

2. 若指定后仍报错，则重装编译器或换用其他版本，例如推荐使用clang++-15。

从 build.sh 源码看，Linux 下默认编译器即为clang/clang++（macOS 优先使用/usr/bin/clang），因此若 PATH 中缺少 clang 可执行文件，就会命中该错误分支。

运行与调试阶段的常见问题

如何运行测试？

项目测试主要分两类入口，均在构建目录中执行：

• bishengir 测试：

ninja check-bishengir # 或等价的 CMake 形式 cmake --build . --target check-bishengir

• LIT 测试套件：

./bin/llvm-lit ../bishengir/test

其中 LIT 测试路径取决于实际仓库与构建目录的相对位置。在 build.sh 的cmake_build中，--build-test模式下会以--target "check-mlir;check-bishengir"并行构建测试目标，并设置了LIT_OPTS="--timeout=30"的用例超时保护。

更详细的构建与测试步骤可参考 安装与构建指南。

板端部署需要哪些环境依赖？

要在 NPU 上端到端运行算子，需要以下三部分配合：

1. CANN：完成安装并执行source set_env.sh配置环境变量；
2. bishengir-compile：负责编译生成设备侧二进制文件（如kernel.o）；
3. CANN Runtime 的 Host 程序：用于算子注册与调用。

完整工作流与动手示例见 编译与执行示例 与 快速开始。

如何获取各层 MLIR 的中间编译状态（如 HFusion、HIVM）？

分为构建阶段与运行阶段两条路径：

• 构建阶段：修改构建脚本配置，将ENABLE_IR_PRINT与BISHENGIR_PUBLISH置为 ON。具体配置方式参考 build.sh 及其配套文档。构建脚本中通过-DBISHENGIR_PUBLISH="${BISHENGIR_PUBLISH}"将该开关透传给 CMake（见 build.sh），它控制尚未公开的特性是否参与编译。

• 运行阶段：使用bishengir-compile提供的打印参数，导出指定 Pass 执行前/后的 MLIR：

bishengir-compile your.mlir --bishengir-print-ir-before=hivm-inject-block-sync --bishengir-print-ir-after=hivm-inject-block-sync

其中 Pass 名称可按需替换。这两个选项的底层实现在 PassManagerOptions.cpp：bishengir-print-ir-before与bishengir-print-ir-after通过PassNameCLParser解析，随后通过pm.enableIRPrinting向 PassManager 注入 IR 打印探针（见 PassManagerOptions.cpp）。需要说明的是，该机制维护了一个白名单（kAllowedPassList，见 PassManagerOptions.cpp），目前允许打印的 Pass 包括hivm-inject-sync、hivm-inject-block-sync、hivm-graph-sync-solver。这也解释了 FAQ 示例中为何选用hivm-inject-block-sync——它本身就是源码中定义并被白名单放行的 Pass（定义见 Passes.td）。

完整参数说明见 编译选项 与 调试与调优。

如何使用 bishengir-compile 将 MLIR 编译为设备侧二进制？

使用-enable-hivm-compile等选项，即可将高层 MLIR 编译为可在 NPU 上执行的二进制：

bishengir-compile input.mlir -enable-hivm-compile -o kernel.o

该选项在源码中的定义为 "Enable BiShengHIR HIVM compile"（见 HIVMCConfig.cpp），是 HIVM 设备侧编译的总开关。仓库测试用例大量使用这一编译路径，例如：

• 集成测试 VecAdd README 展示了bishengir-compile add.mlir -enable-hivm-compile -o kernel.o的完整用法；
• 端到端测试 compile-return-arg.mlir 中可见-enable-hfusion-compile=true -enable-hivm-compile=true -enable-lir-compile=false的组合使用模式。

实际场景中-enable-hivm-compile常与-enable-hfusion-compile（HFusion 融合编译）、-enable-triton-kernel-compile（Triton kernel 编译）、-enable-lir-compile（LIR 编译开关）等配合使用，形成一个"高层 IR → 融合优化 → 设备侧代码生成"的完整流水线。具体选项与流水线细节见 编译选项 与 架构设计。

如何排查 LIT 或 check-bishengir 测试失败？

1. 根据失败用例名称定位对应测试文件与断言信息，区分失败类型：IR 变换问题、数值结果问题、运行环境异常（CANN 版本、文件路径配置等）。
2. 对 IR 变换问题，参考上文"获取各层 MLIR 中间编译状态"的方法查看中间态。

调试选项详见 调试与调优。仓库中大量.mlir测试文件都使用 LIT 的// RUN:行声明编译命令与FileCheck断言，例如 commandline.mlir 中通过bishengir-compile --help | FileCheck %s校验命令行输出，这些文件本身就是定位失败用例时最直接的参考样本。

性能调优：定位算子性能瓶颈

使用 MindStudio 调试 Triton Kernel 性能

MindStudio 内置 Profiler 性能分析组件，可在硬件运行期间采集关键指标，帮助开发者定位 kernel 执行瓶颈。

使用torch_npu.profiler.profile做 PyTorch 训练/推理性能分析

torch_npu.profiler.profile是昇腾 AI 处理器上 PyTorch 训练/推理任务性能分析的核心 API。它通过代码注入，在模型运行期间全面采集 CPU 与 NPU（昇腾 AI 处理器）的性能数据，主要包含多维数据：

• PyTorch 层信息：框架侧算子调用、内存使用、调用栈等；
• CANN 层信息：昇腾计算语言（Ascend Computing Language）接口层的调度与执行情况；
• 硬件层信息：NPU 上的算子执行时间、AI Core 性能指标（如 pipeline 利用率）、缓存命中率等。

它充当 PyTorch 训练脚本与可视化分析工具（如 MindStudio Insight 或 TensorBoard 插件）之间的桥梁。

示例（在 Triton kernel 中采集 NPU 侧性能数据）：

@triton.jit def triton_example(in_ptr0, in_ptr1, out_ptr0, x0_numel, r1_numel, XBLOCK: tl.constexpr, XBLOCK_SUB: tl.constexpr): ... dtype = torch.float16 torch.manual_seed(0) input0 = rand_strided((86, 64, 130), (8320, 130, 1), device='npu:0', dtype=dtype) input1 = rand_strided((1, 64, 1), (64, 1, 1), device='npu:0', dtype=dtype) output = empty_strided((86, 1), (1, 86), device='npu', dtype=dtype) triton_example6,1,1 experimental_config = torch_npu.profiler._ExperimentalConfig( aic_metrics=torch_npu.profiler.AiCMetrics.PipeUtilization, profiler_level=torch_npu.profiler.ProfilerLevel.Level1, l2_cache=False ) with torch_npu.profiler.profile( activities=[ # torch_npu.profiler.ProfilerActivity.CPU, torch_npu.profiler.ProfilerActivity.NPU], with_stack=False, # Switch for collecting the function call stack of torch operators. This parameter is optional and disabled by default. record_shapes=False, # Switch for collecting the input shape and input type of torch operators. This parameter is optional and disabled by default. profile_memory=False, # Switch for collecting memory-related data. This parameter is optional and disabled by default. schedule=torch_npu.profiler.schedule(wait=1, warmup=1, active=10, repeat=1, skip_first=1), # schedule=torch_npu.profiler.schedule(wait=1, warmup=1, active=1, skip_first=6), # warmup defaults to 0. In older versions of the torch_npu package, this parameter is required. experimental_config=experimental_config, # This parameter is optional and defaults to Level0. # Location of the generated profiling files. on_trace_ready=torch_npu.profiler.tensorboard_trace_handler("./result_dir") # Export data in a format that TensorBoard can present. You can specify worker_name, which defaults to {host name}_{process ID}. ) as prof: for i in range(20): triton_example6,1,1 prof.step()

示例中几个关键参数的作用：

• activities=[ProfilerActivity.NPU]：只采集 NPU 活动（CPU 活动可选）；with_stack控制 torch 算子函数调用栈采集（默认关闭）；record_shapes控制输入 shape 与类型采集（默认关闭）；profile_memory控制内存数据采集（默认关闭）；
• schedule：定义wait/warmup/active/repeat/skip_first的采集节奏，用于跳过预热迭代、只采集稳定运行段的性能数据；
• experimental_config：通过aic_metrics=PipeUtilization采集 AI Core 流水线利用率，通过profiler_level控制采样详细等级（默认 Level0）；
• on_trace_ready：将采集结果以 TensorBoard 可展示的格式导出到./result_dir，worker_name默认为{host name}_{process ID}。

精度调试

算子结果与参考实现不一致时如何排查？

调试 Triton kernel 精度问题时，tl.device_print是定位错误不可或缺的工具，它允许在 NPU 运行时直接打印张量或标量的中间值，从而精确定位出错位置。

# 使用前先开启 TRITON_DEVICE_PRINT=1 环境变量 tl.device_print("prefix string", value)

从源码层面看，tl.device_print在 DebugOpToLLVM.cpp 中完成到 LLVM 的降级：其可打印的运行时类型包括整型i1/i8/i16/i32/i64以及fp16/bf16/fp32浮点类型（见 DebugOpToLLVM.cpp），支持直接打印张量元素或标量。下游还有专门为该类打印算子做内存布局适配的变换，例如 InsertL12UBForDebug.cpp 与 InsertNZ2NDForDebug.cpp 会为打印对象调整 L1/L2 数据排布，保证打印出的数值与语义一致。

精度排查策略：

1. 分段打印：在关键计算步骤（如矩阵乘加、归约、激活函数）前后插入tl.device_print，观察数值变化；
2. 对比期望值：打印中间结果后，与手工计算或 CPU 参考实现对比，快速定位误差来源；
3. 关注异常值：如果某处数值突然变成NaN或Inf，在该位置前后打印更多上下文。

示例：

import triton import triton.language as tl @triton.jit def triton_add(in_ptr0, in_ptr1, out_ptr0, XBLOCK: tl.constexpr, XBLOCK_SUB: tl.constexpr): offset = tl.program_id(0) * XBLOCK base1 = tl.arange(0, XBLOCK_SUB) loops1: tl.constexpr = (XBLOCK + XBLOCK_SUB - 1) // XBLOCK_SUB for loop1 in range(loops1): x0_prime = offset + (loop1 * XBLOCK_SUB) + base1 x0 = offset + (loop1 * XBLOCK_SUB) + base1 tmp0 = tl.load(in_ptr0 + (x0), None) # 在 NPU 运行时直接打印 tmp0 的数据 tl.device_print("tmp0", tmp0) tmp1 = tl.load(in_ptr1 + (x0), None) tmp2 = tmp0 + tmp1 tl.store(out_ptr0 + (x0), tmp2, None)

如何使用 bishengir-opt 对比各层 MLIR？

bishengir-opt是类似mlir-opt的工具，用于加载、优化和降级 MLIR 代码的综合测试调试工具。它读取.mlir文件，执行用户指定的编译 Pass，并输出变换后的 IR，支持 AscendNPU-IR 的独立 Pass 调试——开发者可单独执行某个 Pass，对比变换前后的 IR 差异，验证 Pass 是否达到预期功能。

基本语法：

bishengir-opt xx.mlir --{Pass 名称}

示例：输入变换前的 IR（test.mlir）：

// before hfusion-normalize-ops func.func @test_normalize_rec_i32_to_f32(%arg0 : tensor<1x2xi32>) -> tensor<1x2xi32> { %0 = tensor.empty() : tensor<1x2xi32> %1 = hfusion.elemwise_unary {fun = #hfusion.unary_fn<rec>, rec} ins(%arg0 : tensor<1x2xi32>) outs(%0 : tensor<1x2xi32>) -> tensor<1x2xi32> return %1 : tensor<1x2xi32> }

执行bishengir-opt test.mlir --hfusion-normalize-ops后输出：

// after hfusion-normalize-ops module { func.func @test_normalize_rec_i32_to_f32(%arg0: tensor<1x2xi32>) -> tensor<1x2xi32> { %cst = arith.constant 1.000000e+00 : f32 %0 = tensor.empty() : tensor<1x2xf32> %1 = hfusion.cast {cast = #hfusion.type_fn<cast_signed>, enable_overflow = true, round_mode = #hfusion.round_mode<rint>} ins(%arg0 : tensor<1x2xi32>) outs(%0 : tensor<1x2xf32>) -> tensor<1x2xf32> %2 = tensor.empty() : tensor<1x2xf32> %3 = hfusion.elemwise_unary {fun = #hfusion.unary_fn<rec>} ins(%1 : tensor<1x2xf32>) outs(%2 : tensor<1x2xf32>) -> tensor<1x2xf32> %4 = tensor.empty() : tensor<1x2xi32> %5 = hfusion.cast {cast = #hfusion.type_fn<cast_signed>, enable_overflow = true, round_mode = #hfusion.round_mode<trunc>} ins(%3 : tensor<1x2xf32>) outs(%4 : tensor<1x2xi32>) -> tensor<1x2xi32> return %5 : tensor<1x2xi32> } }

hfusion-normalize-opsPass 在源码中定义为 "Normalize Hfusion"，作用是把非规范的高层算子改写为语义等价的规范化算子组合（见 Passes.td）。上述示例展示了该 Pass 的关键行为：i32 -> i32的rec（取倒数）一元算子被规范化成 "i32有符号转f32（cast_signed+rint舍入）→f32上做rec→ 有符号转回i32（trunc截断舍入）" 三步组合，从而保证在昇腾向量指令集上可正确实现。仓库中该 Pass 的 LIT 测试可参考 hfusion-normalize-ops.mlir 与 RegBase 版本用例。

常见精度问题（如 BF16/FP16 精度损失、累加顺序）

典型精度偏差场景：BF16/FP16 低精度计算引起的数值损失，以及不同张量累加顺序造成的误差累积。

浮点精度损失原理：计算机用二进制表示小数，多数十进制小数（如 0.1）无法用有限长度的二进制精确表示，只能近似。不同浮点位宽在数值表示能力上差异显著：

• float32（单精度）：约 7 位有效数字，占 4 字节存储；
• float64（双精度）：约 15～16 位有效数字，占 8 字节存储。

精度验证方案：采用 CPU、GPU、NPU 三方结果对比的验证方案来度量精度偏差。将算法从 CPU 迁移到 NPU/GPU 时，这是标准的强制验证流程，确保硬件加速不会引入不可接受的精度损失。验证基线规则为：以 CPU 上高精度 float64 计算结果作为 ground-truth 基线，再将三类硬件的 float32 输出与之对比，量化整体精度误差。

三方对比的目的：

1. CPU（float64）：作为参考基线，提供最高精度的计算结果；
2. CPU（float32）：用于隔离"精度损失"来源。对比 float32 的 CPU 结果与 float64 结果，可以观察单纯由"单精度"引起的理论损失；
3. GPU/NPU（float32）：用于定位额外的硬件误差。误差来源包括硬件指令集、算子实现逻辑、中间计算存储位宽（例如部分 NPU 使用 FP16 累加）、驱动/库的优化策略等。

浮点比较的核心逻辑：浮点数不能直接用==比较，必须基于容差阈值判断。常见判断方法分两类：

• 绝对误差：|a - b|
• 相对误差：|a - b| / max(|a|, |b|)，适合比较数值较大的情况

混合容差（mixed tolerance）综合了绝对误差与相对误差两种判断逻辑，典型实现如np.isclose()。

贡献与社区

如何贡献？

1. 前置条件：参与前必须签署昇腾社区贡献者许可协议（CLA），并遵循 ascend-community 的行为准则。

2. 标准贡献流程：通过 Issues 反馈问题或认领任务；fork 目标仓库并在本地完成特性开发；本地自测（例如运行ninja check-bishengir）；提交 PR；通过门禁检查（包括编译、静态检查与 CI）。

3. PR 合入条件：PR 必须获得两位评审者的/lgtm以及至少一位 approver 的/approve，方可合入主干分支。

完整贡献规范见 贡献指南。

PR 门禁失败（编译失败、静态检查失败或 CI 失败）如何排查？

根据 CI 输出提示逐项处理门禁异常：

• 编译失败：检查构建错误日志，排查代码与构建环境；
• 静态检查失败：按工具提示修正代码格式、编码规范或逻辑问题；
• CI 测试失败：定位失败测试用例，修复代码后重新触发 CI。

详细解决方案见 贡献指南 - 门禁异常处理。

提交 PR 前有哪些注意事项？

• 避免在 PR 中引入与本次修改无关的变更；
• 保持提交历史精简（适当 squash/rebase）；
• 创建 PR 前将分支 rebase 到最新的上游 master；
• 对于 bug 修复类 PR，在描述中关联相关 Issue 与 PR。

更多细节见 贡献指南 - 注意事项。

小结

AscendNPU-IR 的 FAQ 覆盖了从"构建环境搭建"到"板上精度对齐"再到"社区合码"的全生命周期问题。结合源码可以更深刻地理解每个方案背后的机制：-r参数触发 CMake 重新生成、bishengir-print-ir-*基于 PassManager 白名单注入打印探针、-enable-hivm-compile是 HIVM 设备侧编译总开关、hfusion-normalize-ops负责把高层算子改写为规范化组合、tl.device_print经由 DebugOp 降级链输出运行时中间值。将这些经验与 编译选项、调试与调优、安装与构建指南 配合使用，即可形成一套完整的昇腾 NPU 算子开发、调试与调优闭环。

• 编译器
• 人工智能
• Ascend
• CANN