메뉴
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.
Codex 또는 Claude로 설치 이 Prompt를 복사해 Codex, Claude 또는 다른 어시스턴트에 붙여 넣으면 Skill 페이지를 검토하고 설치를 진행할 수 있습니다.
SOC 직업 분류 기준
| 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)Author Triton kernels with automatic warp specialization (AutoWS). Use when writing new AutoWS kernels, adding warp_specialize=True to tl.range loops, choosing tl.range kwargs and JIT options, debugging why WS was not applied, or structuring a kernel to work with both Meta WS and upstream OAI Triton. Covers GEMM and Flash Attention patterns on Hopper and Blackwell.
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 find the barrier that actually deadlocks -- reasoning with the mbarrier phase model (NOT raw arrive/wait counts, which give false positives), plus missing backward barriers and other synchronization issues. Covers mbarriers, named barriers, tcgen05 commit, TMA-implicit arrives, Aref-based synchronization, and producer/consumer barrier patterns.
Design and run Triton TTGIR debugging ablations using ir_override. Use when reducing a provided or dumped TTGIR, trying user-provided or agent-generated ablation/oblation ideas, updating a test harness around ir_override, or preserving a compile/runtime failure while simplifying IR to expose a fundamental compiler or lowering gap.
How to build and run GPU targets under Buck in fbcode. Use when invoking buck2 run / buck2 build for any GPU benchmark, test, or kernel — selecting the GPU architecture and CUDA version, using @mode/opt and the beta Triton modifier, passing environment variables through, and running from the right directory. Covers the general requirements plus the B200/GB200 (b200a, CUDA >= 12.8) and GB300 (b300a, CUDA >= 13.0) hardware requirements.
Recover from GPU-busy / GPU-unavailable failures. Use when a command (pytest, python, a TLX/Triton kernel run, a benchmark) fails with errors indicating the GPU is busy, out of memory, or unavailable — e.g. "CUDA error: out of memory", "all CUDA-capable devices are busy or unavailable", "CUDA-capable device(s) is/are busy or unavailable", "RuntimeError: No CUDA GPUs are available", "device-side assert", or a hang on the first CUDA call. Runs find_working_gpu.sh to locate a healthy GPU and re-runs the failed command pinned to it via CUDA_VISIBLE_DEVICES.
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.