// Use when working on fence-related compiler passes, TMA store lowering, proxy fence insertion, investigating missing or spurious fences, or debugging correctness issues in TLX kernels that use tlx.async_descriptor_store or MMA operations.
| name | proxy-fence-insertion |
| description | Use when working on fence-related compiler passes, TMA store lowering, proxy fence insertion, investigating missing or spurious fences, or debugging correctness issues in TLX kernels that use tlx.async_descriptor_store or MMA operations. |
Use when working on fence-related compiler passes, TMA store lowering, proxy fence insertion, investigating missing or spurious fences, or debugging correctness issue in TLX kernels that use tlx.async_descriptor_store or MMA operations.
Hopper+ (sm90+) has separate generic and async memory proxies. Writes
through one proxy are not visible to reads through the other without an explicit
proxy fence (fence.proxy.async.shared::cta). For example, a register→SMEM
store (generic proxy) followed by a TMA store from SMEM (async proxy) requires
a fence between the two.
Source: third_party/tlx/language/tlx/mem_ops.py
tlx.fence(scope)Unified fence entry point.
scope | PTX emitted | Use case |
|---|---|---|
"async_shared" | fence.proxy.async.shared::cta | Bridge generic↔async proxy (e.g. between local_store and TMA store) |
"gpu" | fence.acq_rel.gpu | Device-scope ordering of global/shared memory |
"sys" | fence.acq_rel.sys | System-scope ordering (visible to host CPU) |
tlx.fence_async_shared()Deprecated alias for tlx.fence("async_shared").
tlx.local_store(smem, data)
tlx.fence("async_shared") # proxy fence
tlx.async_descriptor_store(desc, smem)
tlx.async_descriptor_store_wait(0)
local_store (generic proxy write) followed by async_descriptor_store (async
proxy read). The TMA hardware reads SMEM via the async proxy, so a fence is
needed after the generic-proxy store. This is handled by TMALowering and
covered by the canonical TMA store pattern above.
When MMA operands are populated by writing registers to SMEM (via LocalAllocOp
with a source or LocalStoreOp), the write goes through the generic proxy.
wgmma and tcgen5 MMA instructions read their SMEM operands through the async
proxy. A proxy fence is required between the register→SMEM copy and the MMA.
This is handled automatically by FenceInsertionPass.
In TLX kernels this shows up when, for example, scales or other data are
written to SMEM from registers and then consumed by a wgmma — the compiler
inserts the fence, but understanding the pattern helps when debugging
correctness issues where the fence might be missing.
Three passes insert proxy fences at different stages of the compilation pipeline. They are listed in the order they run.
File: lib/Dialect/TritonNvidiaGPU/Transforms/FenceInsertion.cpp
Walks every DotOpInterface op (wgmma / tcgen5 MMA). If an operand traces
back to a register→SMEM copy (generic proxy write feeding an async proxy read),
inserts a FenceAsyncSharedOp before the dot. Can hoist the fence out of loops
when safe. Only runs on sm90+.
File: lib/Dialect/TritonNvidiaGPU/Transforms/TMALowering.cpp
Rewrites high-level TMA store ops. Unconditionally inserts a
FenceAsyncSharedOp between the LocalAllocOp (register→SMEM) and the
lowered TMA store:
LocalAllocOp → FenceAsyncSharedOp → TMA store → TMAStoreWaitOp
File: lib/Dialect/TritonNvidiaGPU/Transforms/ProxFenceInsertion.cpp
Runs after shared memory allocation. Uses alias analysis over allocated
buffers to find remaining generic↔async proxy conflicts not caught by earlier
passes. Conservatively inserts fences to avoid races. Only runs on sm90+
(computeCapability >= 90).
FenceAsyncSharedOp (TritonNvidiaGPU dialect)
→ NVVM::FenceProxyOp (NVVM dialect)
→ fence.proxy.async.shared::cta (PTX)
Lowering lives in
third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/BarrierOpToLLVM.cpp
(FenceAsyncSharedOpConversion). The bCluster attribute selects
shared::cluster vs shared::cta scope.
FenceAsyncSharedOp)TLX DSL API reference for low-level GPU primitives. Use when writing or modifying TLX kernel code that uses barriers (mbarrier, named barriers), memory allocation (local_alloc, SMEM, TMEM), TMA operations, warp specialization (async_tasks, async_task), CLC (cluster launch control), or wgmma instructions. Covers Hopper and Blackwell hardware differences.
Run autoWS (automatic warp specialization) correctness tests. Use when working on autoWS compiler code — files under WarpSpecialization/, partition scheduling, warp_specialize ops, WSCodePartition, WSDataPartition, WSTaskPartition, WSMemoryPlanner, or related passes. Do NOT use TLX correctness tests (third_party/tlx/tutorials/testing/test_correctness.py) for autoWS work — those test manual warp specialization via TLX, not the automatic compiler pipeline.
Consult and maintain AutoWS documentation. Use BEFORE exploring AutoWS source code — when investigating, planning, or modifying files under WarpSpecialization/, partition scheduling, warp_specialize ops, WSCodePartition, WSDataPartition, WSTaskPartition, WSMemoryPlanner, or related passes. Also use AFTER making non-trivial changes to AutoWS code to keep docs in sync.
Diagnose CUDA "illegal instruction" / kernel crashes on Triton kernels that reference to TMA loads or stores (`make_tensor_descriptor`, `TensorDescriptor`, `descriptor.load`, `descriptor.store`, `tl.async_descriptor_load`, async TMA copies) as the source code line. Use when the user reports CUDA error 716, "an illegal instruction was encountered", segfault inside a TMA op, kernel hang followed by an illegal instruction trap, or a crash that only fires on the first or last tile of a launch. Covers the pattern where a TMA store/load is issued at an offset entirely past a tensor's shape — TMA does NOT silently mask out-of-bounds tile accesses; it traps. The root cause is almost never "missing in-kernel mask" — it is commonly a structural launcher / tile-mapping bug.
Produce a structured barrier report for AutoWS (automatic warp specialization) IR. Use when the user wants to visualize, audit, or debug barrier usage across warp-specialized partitions, or when debugging a GPU kernel hang (deadlock). For hangs, first dump IR using the ir-debugging skill, then run this barrier analysis to identify mismatched arrive/wait counts, missing backward barriers, or other synchronization issues that cause deadlocks. Covers mbarriers, named barriers, tcgen05 commit, TMA-implicit arrives, Aref-based synchronization, and producer/consumer barrier patterns.
Debug Triton compilation by dumping IR at each stage (TTIR, TTGIR, LLVM, PTX). Use when investigating compilation failures, kernel performance, register spills, or when user asks to inspect IR output. Covers TRITON_KERNEL_DUMP, MLIR_ENABLE_DUMP, LLVM_IR_ENABLE_DUMP, TRITON_DUMP_PTXAS_LOG, and related env vars.