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.

.. code-block:: bash

   # 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:

.. code-block:: text

   [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:

.. code-block:: bash

   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):

.. code-block:: bash

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

Inspecting Cached IR
^^^^^^^^^^^^^^^^^^^^^

.. code-block:: bash

   # 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:

.. code-block:: bash

   # 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:

.. code-block:: text

   // -----// 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``:

.. code-block:: bash

   # 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:

.. code-block:: bash

   # 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:

.. code-block:: bash

   bishengir-compile --help

Environment Variables for IR Dumping
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^

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

.. code-block:: bash

   # 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``:

.. code-block:: bash

   # 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:

.. code-block:: bash

   # 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:

.. code-block:: bash

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

Or when manually invoking ``bishengir-compile``:

.. code-block:: bash

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

Kernel Parameter Validation
^^^^^^^^^^^^^^^^^^^^^^^^^^^^

Print runtime kernel parameters for debugging:

.. code-block:: python

   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:

.. code-block:: bash

   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:

.. code-block:: bash

   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:

.. code-block:: bash

   rm -rf ~/.triton/cache/*

Or force recompilation:

.. code-block:: bash

   export TRITON_ALWAYS_COMPILE=1

Non-deterministic Output
^^^^^^^^^^^^^^^^^^^^^^^^

Multithreaded compilation may produce non-deterministic IR ordering:

.. code-block:: bash

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

Verbose Logging
^^^^^^^^^^^^^^^

Enable verbose Triton logging:

.. code-block:: bash

   export TRITON_DEBUG=1
   python your_script.py

Additional environment variables for detailed logging:

.. code-block:: bash

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

.. list-table::
   :header-rows: 1
   :widths: 30 70

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