FAQ

This document collects common questions about using and developing AscendNPU IR, grouped by topic. For build details, see Build and Installation; for the contribution flow, see Contributing Guide.

Build and Installation

Q1.1 What should I do if the error message “ninja: error: loading ‘build.ninja’: No such file or directory” is displayed when I run the build.sh script?

Add the -r option to build-tools/build.sh to return CMake and regenerate build.ninja. For example:

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

Q1.2 What should I do if the error message “Too many open files” is displayed and the build fails?

The system limits open files per process. Raise the limit and rebuild the project. For example:

ulimit -n 65535

Q1.3 Why is --apply-patches required on the first build?

--apply-patches enables AscendNPU IR extensions (patches) for LLVM/MLIR and other third-party repos; it is required on the first build. You can omit it for later incremental builds.

Running and Debugging

Q2.1 How do I run a test?

In the build directory:

  • BishengIR test: ninja check-bishengir or cmake --build . --target check-bishengir

  • LIT test suite: ./bin/llvm-lit ../bishengir/test (Use the actual paths to your repository and build directory.)

For details, see “Running Tests” in Build and Installation.

Q2.2 What is needed to run on devices?

You need CANN (installed and set_env.sh sourced), a device binary built with bishengir-compile (e.g. kernel.o), and a Host program that uses CANN runtime to register and launch the kernel. See Compile and Run Example and Quick Start.

Q2.3 How do I obtain the intermediate MLIR (e.g. HFusion and HIVM)?

  1. At build time: set ENABLE_IR_PRINT and BISHENGIR_PUBLISH to ON in the build script (see build-tools/build.sh and the docuentation).

  2. At runtime: use bishengir-compile print options to dump MLIR before/after a pass, e.g.:

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

You can use other pass names. See Compile Options and Debug Options.

Q2.4 How do I compile MLIR to a device binary with bishengir-compile?

Use options such as -enable-hivm-compile to compile high-level MLIR to an NPU binary, e.g.:

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

See Compile Options and Architecture Design.

Q2.5 How do I debug LIT or check-bishengir failures?

Use the failing test name to find the test file and assertions; check whether the issue is IR transformation, numerical result, or environment (CANN, paths, etc.). Use “How do I obtain the intermediate MLIR” (Q2.3) to inspect intermediate state. See Debug Options.

Performance Tuning

Q3.1 How do I locate operator performance bottlenecks?

MindStudio

Link: https://www.hiascend.com/developer/software/mindstudio On Ascend, MindStudio’s Profiler collects runtime metrics to help locate kernel bottlenecks when debugging Triton kernels.

Torch NPU profiler

torch_npu.profiler.profile It is the main API for profiling PyTorch training/inference on Ascend. It collects and parses runtime performance data to help locate and fix bottlenecks.

Core functions and positioning: It injects instrumentation to collect CPU and NPU data, including: PyTorch-layer info (ops, memory, and call stacks) CANN-layer scheduling and execution Hardware-layer info (operator time, AI core metrics such as pipeline utilization, cache hit rate)

It bridges your training script and visualization tools (e.g. MindStudio Insight or TensorBoard).

Example:

@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_example[6,1,1](input0, input1, output, 86, 64, XBLOCK=16, XBLOCK_SUB=16)

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, # Controls whether to collect the function call stack of the torch operator. Optional. Default: disabled.
    record_shapes=False, # Controls whether to collect the input shape and input type of the torch operator. Optional. Default: disabled.
    profile_memory=False, # Controls whether to collect memory-related data. Optional. Default: disabled.
    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),
    # The default value of warmup is 0. Required for the early version of the torch_npu package.
    experimental_config=experimental_config, # Optional. Default: Level0.
    # Location of the generated profile file.
    on_trace_ready=torch_npu.profiler.tensorboard_trace_handler("./result_dir")
    # Format of data exported to TensorBoard. You can specify worker_name. Default: {host name}_{process ID}.
) as prof:
    for i in range(20):
        triton_example[6,1,1](input0, input1, output, 86, 64, XBLOCK=16, XBLOCK_SUB=16)
        prof.step()

Accuracy and Debugging

Q4.1 How do I debug when the operator result is different from the reference (such as CPU/GPU or reference implementation)? When debugging precision issues in Triton kernels, tl.device_print is an indispensable tool. It allows you to directly print intermediate values of tensors or scalars at NPU runtime, thereby pinpointing where the error occurs. Usage guide:

# Before using this function, set the environment variable TRITON_DEVICE_PRINT=1.
tl.device_print("prefix string", value)

Precision troubleshooting strategies

  1. Segmented printing: Insert tl.device_print before and after key computation steps (e.g., matrix multiply-add, reduction, activation functions) to observe numerical changes.

  2. Compare with expected values: After printing intermediate results, compare them with manual calculations or the CPU reference implementation to quickly locate the source of error.

  3. Watch for abnormal values: If values suddenly become NaN or Inf, print more context around the corresponding positions.

Example:

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)
        # Print tmp0 data directly at NPU runtime
        tl.device_print("tmp0",  tmp0)
        tmp1 = tl.load(in_ptr1 + (x0), None)
        tmp2 = tmp0 + tmp1
        tl.store(out_ptr0 + (x0), tmp2, None)

Q4.2 How do I compare numerical results across MLIR layers? bishengir-opt is a tool similar to mlir-opt, primarily used for loading, optimizing, and transforming (lowering) MLIR code. You can think of it as a “Swiss Army knife” testing and debugging tool: it reads an .mlir file, applies a series of user-specified compilation passes, and outputs the result. Hence, it can be used for independent pass debugging of AscendNPU IR. Through it, developers can apply a specific pass individually and compare the IR differences before and after application, thereby verifying whether the pass achieves the intended functionality.

Basic syntax: bishengir-opt xx.mlir –{pass name}

Usage example:

bishengir-opt test.mlir –hfusion-normalize-ops

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>
}

After executing the single hfusion-normalize-ops pass:

// 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>
  }
}

Q4.3 What are common accuracy issues (e.g. BF16/FP16 loss, reduction order)?

How to determine whether precision loss meets standards: Use a three-way comparison scheme to verify precision loss (NPU, GPU, CPU).

This is a classic and necessary verification process, especially when porting algorithms from CPU to NPU or GPU, to ensure that hardware acceleration does not introduce unacceptable precision loss. Taking CPU float64 as the “ground truth” benchmark and comparing the float32 outputs of all three is the gold standard for measuring precision loss.

Why does precision loss occur? Computers use binary to represent decimal numbers; many decimal fractions (e.g., 0.1) cannot be exactly represented by finite binary length and can only be approximated. The precision difference between float32 and float64 is huge: float32 (single precision): approximately 7 significant digits. Memory footprint: 4 bytes. float64 (double precision): approximately 15–16 significant digits. Memory footprint: 8 bytes.

Why is a three-way comparison needed? CPU (float64): Serves as the reference benchmark, providing the highest precision computation results. CPU (float32): Isolates the source of “precision loss.” Comparing float32 CPU results with float64 results reveals the theoretical loss caused purely by “single precision.” GPU/NPU (float32): Observes additional errors introduced by specific hardware accelerators due to differences in instruction sets, operator implementation algorithms, intermediate result retention precision ( e.g., some NPUs may use float16 for accumulation), or driver/library optimization strategies.

Core comparison logic Since floating-point numbers cannot be directly compared with ==, tolerance-based comparison must be used. Common methods are: Absolute Error: |a - b| Relative Error: |a - b| / max(|a|, |b|), which is suitable for comparing large numbers. Mixed tolerance: Combines both, e.g., the implementation of np.isclose().

Contributions and Community

Q5.1 How can I contribute?

You must sign the Ascend Community CLA and follow the ascend-community code of conduct. Flow: open or claim an Issue, fork and develop, self-test (e.g. ninja check-bishengir), open a PR, and pass CI (build, static check, tests). Merge requires 2 Reviewers’ /lgtm and 1 Approver’s /approve. See Contributing Guide.

Q5.2 How do I debug when a PR fails CI (build, static check, or tests)?

Follow the CI log: fix build failures (errors and environment), static check (style/logic as indicated), and failing tests (fix and re-run CI). See “Dealing with CI failures” in Contributing Guide.

Q5.3 What should I do before opening a PR?

Avoid unrelated changes; keep history clear (squash/rebase if needed); rebase your branch onto the latest upstream master; for bug-fix PRs, reference related Issues and PRs in the description. See “Notes” in Contributing Guide.