Debugging

Read compilation errors

When the compiler can't compile a kernel, it raises enceladus.CompilationError with the file, line, and column, a description, and a suggested fix. For example, a block size that isn't annotated as tl.constexpr fails as follows:

enceladus.CompilationError: add.py:9:12: the end of tl.arange must be a compile-time
integer, but got a runtime value `BLOCK` of type tl.int32. If `BLOCK` comes from a kernel
parameter, annotate the parameter as tl.constexpr, for example `BLOCK: tl.constexpr`.
offs = tl.arange(0, BLOCK)
       ^

Run kernels in the interpreter

The interpreter runs a kernel on the CPU with NumPy, one program at a time. You can use print(), pdb, and breakpoints inside the kernel and see real tile values. To run every kernel in the interpreter, set ENCELADUS_INTERPRET=1. To interpret one kernel, pass interpret=True to the decorator.

ENCELADUS_INTERPRET=1 uv run python my_script.py

The interpreter matches the GPU with a few exceptions: the GPU flushes float32 denormals to zero, math functions such as exp differ in the last bits, and colliding atomics can apply in a different order. The test suite uses the interpreter as the reference for every compiled kernel.

Inspect the generated IR and MSL

To see what the compiler generates, set ENCELADUS_DUMP=1. Enceladus writes ir.txt, kernel.metal, and meta.json to the kernel's cache entry and prints its path to stderr. To get the MSL from Python without launching, call warmup:

compiled = add_kernel.warmup(x, y, out, n, BLOCK=1024)
print(compiled.msl)
print(compiled.threadgroup_memory_bytes)

To test a hand edit, copy kernel.metal to DIR/KERNEL_NAME.metal and set ENCELADUS_OVERRIDE_DIR=DIR. Keep the entry point and the [[buffer(N)]] arguments unchanged.

Explain the compiler's decisions

kernel.explain(*args, grid=..., **constexprs) prints and returns a report for one specialization without launching it. It lists the tl.dot backend and why a dot fell back, each tile's layout and registers per thread, layout conversions, and peak threadgroup memory:

Kernel `rowsum_matmul`: num_warps=4 (128 threads per program), grid (4, 1, 1)
MSL language version 3.2
dot backend: simdgroup
dbg.py:30  64x64x32 (MxNxK), SIMD-group grid 4x1  `acc = tl.dot(...)`
uses simdgroup instead of mpp: dbg.py:30:15: the tl.dot result feeds `reduce`,
which isn't an elementwise op in the loop's block

Tiles (layout; registers per thread):
dbg.py:29  loop-carried `acc` f32[64, 64]: simdgroup_matrix fragments: regs 2x16,
lanes 8x4, warps 4x1; 32 registers  `for k in range(0, K, BK):`
dbg.py:30  desc_load f16[64, 32]: read by tl.dot straight from device memory;
0 registers  `acc = tl.dot(...)`

Layout conversions:
(none)

Threadgroup memory: 0 bytes of 32768 (one arena that each operation reuses)

Print and assert on the GPU

tl.device_print(prefix, *args) prints one line per element to stderr when the stream synchronizes. tl.device_assert(cond, msg, mask=None) compiles only when ENCELADUS_DEBUG=1 and raises enceladus.DeviceAssertionError at the next sync:

gather.py:12:5: device assertion failed in program (1, 0, 0): index out of range
tl.device_assert((idx >= 0) & (idx < n), "index out of range", mask=mask)
^

The kernel keeps running after a failed assert, so guard the access with a mask as well. For compile-time values, use tl.static_print and tl.static_assert.

Capture a GPU trace

enceladus.capture(path) records GPU work to a .gputrace document that you can open in Xcode. Start the process with MTL_CAPTURE_ENABLED=1. The path must end in .gputrace and must not exist. Capture and tl.device_print can't be used in the same process.

with enceladus.capture("add.gputrace"):
add_kernel[(4,)](x, y, out, 4096, BLOCK=1024)