Kernel Debugging

This guide covers techniques for debugging Triton kernels compiled for Ascend NPUs, including extracting intermediate representations (IR), analyzing compiler transformations, and diagnosing synchronization issues.

Extracting Intermediate Representations

Triton caches compiled kernels and intermediate artifacts. By inspecting the cache directory, you can examine the IR at various compilation stages.

Default Cache Location

By default, Triton caches compiled kernels in ~/.triton/cache/. Each kernel is stored in a subdirectory named with an MD5 hash based on the kernel signature and compilation options.

# List cached kernels
ls ~/.triton/cache/

For Ascend NPU compilation, the cache directory for each kernel contains:

• kernel.ttir — High-level Triton IR with distribution primitives

• kernel.ttadapter (release/3.2.2) or kernel.bcmlir (master branch) — Adapter IR used as input to bishengir-compile

• Compiled binaries — NPU executable files (.o)

The compilation pipeline:

[Triton Python Kernel]
     ↓ (triton.compile)
[kernel.ttir]
     ↓ (Ascend backend adaptation)
[kernel.ttadapter]  (release/3.2.2)
or
[kernel.bcmlir]     (master branch)
     ↓ (bishengir-compile)
[kernel.o]  (NPU executable)

Enabling Debug Dumps

To dump IR files during compilation (useful when cache is hit), enable debug mode:

export TRITON_DEBUG=1
python your_script.py

This creates IR dumps in ~/.triton/dump/ with files like:

• kernel.ttir.mlir

• kernel.ttadapter.mlir or kernel.bcmlir.mlir

To force recompilation (bypassing cache):

export TRITON_ALWAYS_COMPILE=1
export TRITON_DEBUG=1
python your_script.py

Inspecting Cached IR

# Find the latest kernel hash
ls -lt ~/.triton/cache/ | head

# Examine a specific kernel's IR
cd ~/.triton/cache/<hash>
cat kernel.ttir        # High-level Triton IR
cat kernel.ttadapter   # Adapter IR (release/3.2.2)
# or
cat kernel.bcmlir      # BCM IR (master branch)

Use this to verify:

• Correct kernel parameters and tensor shapes

• Expected distribution primitives (wait, notify, symm_at)

• Memory allocations and layout transformations

Printing IR During Compilation

The Ascend backend uses bishengir-compile to lower adapter IR to NPU binaries. You can pass MLIR debugging flags to bishengir-compile to print IR after each compiler pass.

Manual Invocation with Debug Flags

Extract the adapter IR from the cache, then manually invoke bishengir-compile with debug flags:

# First, run your script to generate cached IR
export TRITON_DEBUG=1
python your_script.py

# Find the cached adapter IR
cd ~/.triton/cache/<hash>

# Manually compile with IR printing
bishengir-compile kernel.ttadapter --target=Ascend910B3 \
  --mlir-print-ir-after-all \
  --enable-auto-multi-buffer=True \
  --enable-hfusion-compile=true \
  --enable-hivm-compile=true \
  --enable-triton-kernel-compile=true

This produces verbose output showing IR transformations after each pass:

// -----// IR Dump After SomePass //----- //
module {
  func.func @kernel(...) {
    ...
  }
}

// -----// IR Dump After AnotherPass //----- //
module {
  func.func @kernel(...) {
    ...
  }
}

Useful MLIR Debug Flags

Common flags to pass to bishengir-compile:

# Print IR after all passes
--mlir-print-ir-after-all

# Print IR only after specific passes
--mlir-print-ir-after=pass-name

# Print IR before all passes
--mlir-print-ir-before-all

# Print only when IR changes (reduces noise)
--mlir-print-ir-after-change

# Disable multithreading for deterministic output
--mlir-disable-threading

Additional bishengir-compile flags for debugging:

# Enable debug info generation
--enable-debug-info=true

# Print IR after specific HIVM passes
--hivm-compile-args=bishengir-print-ir-after=hivm-inject-sync

View all available bishengir-compile options:

bishengir-compile --help

Environment Variables for IR Dumping

Triton-Ascend provides environment variables to control IR dumping during compilation:

# Dump MLIR IR before each optimization pass
export MLIR_ENABLE_DUMP=1

# Dump LLVM IR before each LLVM optimization
export LLVM_IR_ENABLE_DUMP=1

# Generate MLIR reproducer files at each compilation phase
export TRITON_REPRODUCER_PATH=/tmp/reproducer

# Enable verbose debugging output
export TRITON_DEBUG=1

# Force recompilation (bypass cache)
export TRITON_ALWAYS_COMPILE=1

python your_script.py

Note: MLIR_ENABLE_DUMP may not work if the cache is hit. Use TRITON_ALWAYS_COMPILE=1 to force recompilation, or clear the cache with rm -rf ~/.triton/cache/.

Debugging Synchronization Issues

Distributed kernels use fine-grained synchronization based on data dependencies. The compiler calculates synchronization points, but these calculations may have bugs or edge cases. To isolate compiler synchronization issues, you can force the use of global barriers.

Using --enable-hivm-inject-barrier-all-sync

This flag replaces fine-grained synchronization with global barriers at every synchronization point.

To use it, manually invoke bishengir-compile:

# Extract cached adapter IR
export TRITON_DEBUG=1
python your_script.py

cd ~/.triton/cache/<hash>

# Recompile with global barriers
bishengir-compile kernel.ttadapter --target=Ascend910B3 \
  --enable-hivm-inject-barrier-all-sync=true \
  --enable-auto-multi-buffer=True \
  --enable-hfusion-compile=true \
  --enable-hivm-compile=true \
  --enable-triton-kernel-compile=true

Effect:

• Every synchronization operation is replaced with barrier_all()

• All PEs wait at every synchronization point

• Performance degrades, but correctness improves if the issue is synchronization-related

When to Use:

• Suspect data races or incorrect synchronization

• Kernel produces non-deterministic or intermittent incorrect results

• Want to isolate whether the issue is in synchronization logic or computation logic

Workflow:

1. Run kernel with fine-grained sync (default) — observe the issue

2. Manually recompile with --enable-hivm-inject-barrier-all-sync=true and replace the cached binary

3. If global barriers fix the issue, the compiler’s synchronization analysis likely has a bug — report with a minimal reproducer

4. If the issue persists, the problem is elsewhere (computation, memory access, etc.)

Example Debugging Session

Here’s a complete debugging workflow for a kernel producing incorrect results:

# Step 1: Enable dumps and debug mode
export TRITON_DEBUG=1
export TRITON_ALWAYS_COMPILE=1
python failing_kernel.py

# Step 2: Examine the dumped IR
ls ~/.triton/dump/
cat ~/.triton/dump/kernel.ttir.mlir
cat ~/.triton/dump/kernel.ttadapter.mlir  # or kernel.bcmlir.mlir

# Step 3: Find the cached kernel hash
ls -lt ~/.triton/cache/ | head
cd ~/.triton/cache/<latest_hash>

# Step 4: Manually recompile with verbose IR printing
bishengir-compile kernel.ttadapter --target=Ascend910B3 \
  --mlir-print-ir-after-all \
  --enable-auto-multi-buffer=True \
  --enable-hfusion-compile=true \
  --enable-hivm-compile=true \
  --enable-triton-kernel-compile=true \
  2>&1 | tee compile_log.txt

# Step 5: Search for specific operations in the log
grep -A 10 "notify" compile_log.txt
grep -A 10 "symm_at" compile_log.txt
grep -A 10 "inject-sync" compile_log.txt

# Step 6: Try with global barriers
bishengir-compile kernel.ttadapter --target=Ascend910B3 \
  --enable-hivm-inject-barrier-all-sync=true \
  --enable-auto-multi-buffer=True \
  --enable-hfusion-compile=true \
  --enable-hivm-compile=true \
  --enable-triton-kernel-compile=true \
  -o kernel_fixed.o

# Step 7: Replace the cached binary and test
cp kernel_fixed.o <original_binary_name>
python failing_kernel.py

Advanced Techniques

Dumping IR at Specific Passes

To reduce output volume, dump IR only at specific passes:

export MLIR_ENABLE_DUMP=kernel_name  # Dump only specific kernel
python your_script.py

Or when manually invoking bishengir-compile:

bishengir-compile kernel.ttadapter --target=Ascend910B3 \
  --mlir-print-ir-after=hivm-inject-sync

Kernel Parameter Validation

Print runtime kernel parameters for debugging:

import triton
import triton_dist.language as tdl

@triton.jit
def my_kernel(ptr, size, ...):
    rank = tdl.rank()
    if rank == 0:
        # Use device_print for runtime debugging
        pass  # Insert debugging logic here

Using TRITON_INTERPRET

For CPU-based debugging without NPU execution:

export TRITON_INTERPRET=1
python your_script.py

This runs the kernel in interpreted mode on the CPU, helping isolate kernel logic issues from NPU-specific behavior.

Enabling Device Print

To enable tl.device_print() and tl.static_print() for runtime debugging:

export TRITON_DEVICE_PRINT=1
python your_script.py

Note: Each thread has a 16 KB GM buffer limit for device print output.

Common Pitfalls

Cached Kernels

Triton caches compiled kernels. After changing compiler flags or debugging settings, clear the cache:

rm -rf ~/.triton/cache/*

Or force recompilation:

export TRITON_ALWAYS_COMPILE=1

Non-deterministic Output

Multithreaded compilation may produce non-deterministic IR ordering:

bishengir-compile kernel.ttadapter --mlir-disable-threading ...

Verbose Logging

Enable verbose Triton logging:

export TRITON_DEBUG=1
python your_script.py

Additional environment variables for detailed logging:

export TRITON_ENABLE_LLVM_DEBUG=1  # Extensive LLVM CodeGen logs (very large)
export MLIR_ENABLE_TIMING=1        # Timing statistics for MLIR passes
export LLVM_ENABLE_TIMING=1        # Timing statistics for LLVM passes

Reporting Issues

When reporting compiler bugs, please include:

1. Minimal reproducer — Smallest kernel that triggers the issue

2. Cached IR — Attach ~/.triton/cache/<hash>/ contents for the failing kernel

3. Compiler output — Full logs with --mlir-print-ir-after-all or relevant environment variables

4. Environment info — CANN version, Triton-Ascend version, NPU model, branch (release/3.2.2 vs master)

5. Workaround status — Does --enable-hivm-inject-barrier-all-sync=true fix it?

Submit issues to:

• Triton-distributed-ascend: https://gitcode.com/Ascend/Triton-distributed-ascend/issues

• AscendNPU-IR: https://gitcode.com/Ascend/AscendNPU-IR/issues

Summary

Technique	Command
Enable IR dumps	export TRITON_DEBUG=1
Force recompilation	export TRITON_ALWAYS_COMPILE=1
Dump MLIR passes	export MLIR_ENABLE_DUMP=1
Dump LLVM IR	export LLVM_IR_ENABLE_DUMP=1
Print all IR passes (manual)	bishengir-compile ... --mlir-print-ir-after-all
Print only changes (manual)	bishengir-compile ... --mlir-print-ir-after-change
Force global barriers (manual)	bishengir-compile ... --enable-hivm-inject-barrier-all-sync=true
Disable multithreading (manual)	bishengir-compile ... --mlir-disable-threading
CPU interpretation	export TRITON_INTERPRET=1
Enable device print	export TRITON_DEVICE_PRINT=1
Verbose LLVM debug	export TRITON_ENABLE_LLVM_DEBUG=1
Clear cache	rm -rf ~/.triton/cache/*

Key Points:

• IR files: kernel.ttir → kernel.ttadapter (release/3.2.2) or kernel.bcmlir (master) → kernel.o

• Cache location: ~/.triton/cache/

• Dump location: ~/.triton/dump/ (when TRITON_DEBUG=1)

• Manual compilation needed for bishengir-compile debug flags

• Use --enable-hivm-inject-barrier-all-sync=true to diagnose synchronization bugs