# Enable debug mode (default: False)
export CUTE_DSL_DEBUG=1

export CUTE_DSL_DEBUG=1
export CUTE_DSL_LINEINFO=0

Debug mode adds extra checks and diagnostics that increase compile time and
may affect the generated code (for example, by embedding line info). Enable
it while debugging, not for production runs.

The settings debug mode raises -- line info in particular -- change the
emitted IR/PTX, and every one of these settings is folded into the JIT
kernel cache key. A kernel compiled with debug mode on is therefore cached
separately from the same kernel compiled with it off: toggling
``CUTE_DSL_DEBUG`` forces a recompile instead of reusing a cached kernel,
and the kernel you inspect or profile under debug mode is not identical to
the one produced for a normal (debug-off) run. Validate performance and
generated-code conclusions with debug mode disabled. Because these settings
are part of the cache key, a debug-built kernel is never silently reused for
a production run. See :doc:`JIT caching <../cute_dsl_general/dsl_jit_caching>`
for how the cache key is formed.

export CUTE_DSL_COMPILER_OPT='warnings{nvvm},remarks{nvvm}'

CUTE_DSL_COMPILER_OPT='warnings{nvvm},remarks{nvvm}' \
python examples/python/CuTeDSL/experimental/compiler_diagnostic/prims_negative_cases.py

CUTE_DSL_COMPILER_OPT='warnings{nvvm},remarks{nvvm}' \
python examples/python/CuTeDSL/experimental/compiler_diagnostic/prims_negative_cases.py \
    --case expect_tx_without_complete_tx

import cutlass
import cutlass.cute as cute
from cutlass.cute.runtime import make_fake_stream
from cutlass.experimental import primitives as prims


@cute.kernel
def expect_tx_without_complete_tx_kernel() -> None:
    mbar = cutlass.Array(
        cutlass.Int64, 1, space=cutlass.AddressSpace.smem, alignment=8
    )

    if prims.elect_sync():
        prims.mbarrier_init(mbar, 1)
    prims.fence_mbarrier_init()
    prims.barrier_cta_sync(0)

    if prims.elect_sync():
        prims.mbarrier_arrive_expect_tx(mbar, 16_384)
    prims.mbarrier_try_wait_parity(mbar, 0, time_limit=10_000_000)


@cute.jit
def host_expect_tx_without_complete_tx(stream) -> None:
    expect_tx_without_complete_tx_kernel().launch(
        grid=(1, 1, 1), block=(32, 1, 1), stream=stream
    )


cute.compile(
    host_expect_tx_without_complete_tx,
    make_fake_stream(),
    options="warnings{nvvm},remarks{nvvm}",
)

===== expect_tx_without_complete_tx =====
error[E####]: mbarrier.arrive.expect_tx has no completion source for 16384 registered transaction bytes

  --> examples/python/CuTeDSL/experimental/compiler_diagnostic/prims_negative_cases.py:94:8
      in function `_expect_tx_without_complete_tx_kernel(...)`:
   |
  92 |
  93 |     if prims.elect_sync():
> 94 |         prims.mbarrier_arrive_expect_tx(mbar, 16_384)
     |        ^
  95 |     prims.mbarrier_try_wait_parity(mbar, 0, time_limit=_WAIT_TICKS)
  96 |
  error: arrive.expect_tx increments the barrier transaction count, but no explicit
         mbarrier.complete_tx or TMA complete_tx source targets this barrier. The transaction count
         can never be retired.
  suggestion: add a matching nvvm.mbarrier_complete_tx(...) on this barrier, or issue a TMA
              operation whose completion targets the same barrier
  note: PTX ISA mbarrier.arrive.expect_tx
        docs: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-expect-tx-operation
PASS expect_tx_without_complete_tx

cute.compile(my_host_fn, *args, options="remarks{ptx}")

cute.compile(
    my_host_fn,
    *args,
    options=(
        "remarks{ptx} "
        "--ptxas-options '--maxrregcount=128 --override-directive-values'"
    ),
)

python examples/python/CuTeDSL/experimental/compiler_diagnostic/ptxas_spill_cases.py

CuTe DSL provides environment variables to control logging level:

.. code:: bash

    # Enable console logging (default: False)
    export CUTE_DSL_LOG_TO_CONSOLE=1

    # Log to file instead of console (default: False).
    # Set to 1/True to enable; the log file path is chosen automatically by the DSL.
    export CUTE_DSL_LOG_TO_FILE=1

    # Control log verbosity (0=disabled, 1=all messages (debug and above), 10=debug, 20=info, 30=warning, 40=error, 50=critical; default: 1)
    export CUTE_DSL_LOG_LEVEL=20


Log Categories and Levels

CuTe DSL can save generated artifacts (IR, PTX, CUBIN, …) to files for offline inspection.
Use ``CUTE_DSL_KEEP`` with a comma-separated list of artifact tokens. Prefer
this consolidated option over deprecated per-artifact variables such as
``CUTE_DSL_KEEP_PTX=1``.

.. code:: bash

    # Save clean IR (after canonicalize+cse, human-readable) to a .mlir file
    export CUTE_DSL_KEEP=ir

    # Save raw IR (before any passes) to a .mlir file
    export CUTE_DSL_KEEP=ir-debug

    # Save PTX assembly to a .ptx file
    export CUTE_DSL_KEEP=ptx

    # Save CUBIN binary to a .cubin file
    export CUTE_DSL_KEEP=cubin

    # Save SASS disassembly to a .sass file
    export CUTE_DSL_KEEP=sass

    # Save multiple artifacts at once
    export CUTE_DSL_KEEP=ir,ptx,cubin,sass

    # Save all supported artifacts
    export CUTE_DSL_KEEP=all

Files are written to the current working directory by default. Use
``CUTE_DSL_DUMP_DIR`` to redirect them (see `Change the dump directory`_
below).

.. note::

    The ``sass`` token disassembles the CUBIN with ``nvdisasm``. Starting
    with CUTLASS 4.7, ``nvdisasm`` is no longer installed by default. To
    dump SASS, do one of the following:

    - install the ``sass`` extra (recommended — its ``nvdisasm`` version
      is guaranteed to match the toolchain the DSL ships):
      ``pip install nvidia-cutlass-dsl[sass]``
    - install a local CUDA Toolkit and expose it via ``CUDA_HOME`` /
      ``CUDA_PATH``

    A locally-provided ``nvdisasm`` must come from a CUDA Toolkit at least
    as new as the toolchain that produced the CUBIN; an older one (e.g.
    from CUDA 12.x) may fail to disassemble it.

    SASS dumping can also be controlled per compilation with ``KeepSASS``
    and ``NvdisasmOptions``; see :ref:`JIT_Compilation_Options`.

Print the generated IR to the console

# Print generated IR to stdout (default: False)
export CUTE_DSL_PRINT_IR=1

For compiled kernels, the generated PTX/CUBIN/IR can also be accessed
programmatically through the following attributes:

- ``__ptx__``: The generated PTX code of the compiled kernel.
- ``__cubin__``: The generated CUBIN data of the compiled kernel.
- ``__sass__``: The generated SASS disassembly of the compiled kernel, when
  SASS was requested.
- ``__mlir__``: The generated IR code of the compiled kernel.

.. code:: python

    compiled_foo = cute.compile(foo, ...)
    print(f"PTX: {compiled_foo.__ptx__}")
    with open("foo.cubin", "wb") as f:
        f.write(compiled_foo.__cubin__)


Change the dump directory
~~~~~~~~~~~~~~~~~~~~~~~~~

By default, all dumped files are saved in the current working directory. To specify a different directory for the dumped files, please set the environment variable CUTE_DSL_DUMP_DIR accordingly.


Kernel Functional Debugging
----------------------------

Using Python's ``print`` and CuTe's ``cute.printf``

When a kernel becomes unresponsive and ``SIGINT`` (``CTRL+C``) fails to terminate it,
you can follow these steps to forcefully terminate the process:

1. Use ``CTRL+Z`` to suspend the unresponsive kernel
2. Execute the following command to terminate the suspended process:

.. code:: bash

    # Terminate the most recently suspended process
    kill -9 $(jobs -p | tail -1)


CuTe DSL can also be debugged using standard NVIDIA CUDA tools.

Using Compute-Sanitizer
~~~~~~~~~~~~~~~~~~~~~~~

For detecting memory errors:

.. code:: bash

    compute-sanitizer --tool memcheck python your_dsl_code.py

For detecting race conditions:

.. code:: bash

    compute-sanitizer --tool racecheck python your_dsl_code.py

Please refer to the `compute-sanitizer documentation <https://developer.nvidia.com/compute-sanitizer>`_ for more details.

Set function name prefix
~~~~~~~~~~~~~~~~~~~~~~~~~

By default, the generated IR name of a host or kernel function is based on the
Python function name and its parameters. Call ``set_name_prefix`` on a JIT or
kernel function before invoking it to add runtime context or customize that IR
name.

``set_name_prefix`` also accepts two optional keyword-only parameters:

* ``remove_cutlass_symbol=False`` removes the ``cutlass`` component that CuTe
  DSL automatically inserts. It does not modify the user-provided prefix, the
  Python function name, or text derived from mangled arguments.
* ``keep_mangled_name=True`` retains the framework-generated function and
  argument components. When set to ``False``, those components are omitted,
  while the ``cutlass`` marker remains unless ``remove_cutlass_symbol=True``.
  The per-kernel numeric uniqueness suffix is always retained.

Calling ``set_name_prefix("")`` with the optional arguments left at their
defaults restores all default naming behavior. An empty prefix can also be
combined with non-default component options; for example,
``set_name_prefix("", remove_cutlass_symbol=True)`` removes the CuTe DSL marker
without adding a user prefix. A non-empty prefix is required when
``remove_cutlass_symbol=True`` and ``keep_mangled_name=False`` so that at least
one textual name component remains.

.. code:: python

    @cute.kernel
    def kernel(arg1, arg2, ...):
        ...

    @cute.jit
    def launch_kernel():
        kernel.set_name_prefix(
            "my_op",
            remove_cutlass_symbol=True,
            keep_mangled_name=False,
        )
        kernel(arg1, arg2, ...).launch(
            grid=[1, 1, 1], block=[1, 1, 1], ...
        )

For a first kernel trace whose default name resembles
``kernel_cutlass_kernel_<arguments>_0``, representative results are:

* ``set_name_prefix("my_op")``:
  ``my_op_kernel_cutlass_kernel_<arguments>_0``
* ``set_name_prefix("my_op", remove_cutlass_symbol=True)``:
  ``my_op_kernel_kernel_<arguments>_0``
* ``set_name_prefix("my_op", keep_mangled_name=False)``:
  ``my_op_cutlass_0``
* Enabling both options, as above: ``my_op_0``

The numeric suffix can differ when the same kernel is traced more than once.
Host JIT function names do not have the per-kernel numeric suffix.

To produce a generated IR name without ``cutlass``, use a prefix that does not
contain it. If ``keep_mangled_name=True``, also ensure that the Python function
name and mangled argument text do not contain it; setting
``keep_mangled_name=False`` omits those components.

The kernel suffix prevents collisions between traces in one generated module.
It does not make names unique across separately compiled modules, and host JIT
names have no such suffix. When ``keep_mangled_name=False``, use a prefix that
is unique in every final link or load scope where modules can be combined.

``set_name_prefix`` does not truncate the user prefix, so do not rely on
truncation to remove text from it. Keep the resulting name within the limits of
the tools that consume the generated artifact.

This API controls traced MLIR function names, not ABI wrappers added by later
export stages such as ``cutlass_call_<function-name>``.

Conclusion
----------

Effective CuTe DSL debugging typically combines source correlation, artifact
dumps, runtime prints, and CUDA tooling. When reporting an issue, include the
minimal reproducer, relevant generated artifacts, and logs collected with the
debugging options above, and share it with the CUTLASS team as a GitHub issue.