Triton codegen inspection

Reading Triton's TTGIR and PTX shows which tiling, swizzling, and pipelining decisions the compiler made for you.

Where the IR comes from

TRITON_KERNEL_DUMP=1 with TRITON_DUMP_DIR=<dir> writes the pipeline’s artifacts for each compiled kernel: TTIR, TTGIR, LLIR and PTX. From Python, a compiled kernel exposes the same text as kernel.asm["ttgir"] and kernel.asm["ptx"], with kernel.n_regs and kernel.n_spills beside it. Add TRITON_ALWAYS_COMPILE=1 when a cache hit would otherwise hand you an artifact from an earlier build.

TTGIR: the decisions made for you

TTIR is generic; TTGIR is the tiled GPU IR where every value carries a layout. The operations to look for, abridged:

%a_smem = ttg.local_alloc   %a      -> a shared-memory buffer for the A tile
%a_frag = ttg.local_load    %a_smem -> registers in the dot operand layout
%acc_mma = ttg.convert_layout %acc  -> the layout the matrix instruction wants
ttg.async_copy_global_to_local %k_ptrs, %k_smem
ttg.async_wait

local_alloc and local_load are the shared-memory staging the compiler inserted so one tile feeds many matrix instructions. convert_layout is where data changes distribution across lanes, and it is the first thing to count, because each one becomes shuffles or a round trip through shared memory. The async copy pair is num_stages made visible: tiles are fetched ahead while the dot consumes the current one. Printed type and encoding names change between releases, so read what your build emits rather than what a blog post shows.

PTX: the instruction selected

In PTX you see which matrix instruction the backend chose on your target: mma.sync.aligned.m16n8k16 for Ampere-class MMA, wgmma.mma_async on Hopper, tcgen05.mma on Blackwell. You also see cp.async for the pipelined loads, the swizzle arithmetic folded into shared-memory addresses, and the register count that decides occupancy.

Things that break

Inspecting a cached kernel built with different flags and drawing conclusions about the code you just wrote. Grepping for a layout or operation name remembered from an older release: tt.make_tensor_ptr was removed in 3.8, and the shared-memory encoding attributes have been renamed more than once. And treating tidy PTX as proof of speed, when n_spills and achieved occupancy say otherwise. The dump reflects Triton 3.9; confirm the knob names and the asm keys against your installed version.

Sources