fla-triton-to-gluon

Ports Triton kernels in fla/ops to Gluon with explicit layouts, async copies, and MMA control.

Updated Aug 18, 2026
One-click install
npx skills add https://github.com/weichengz0616/fla --skill fla-triton-to-gluon-weichengz0616
Or copy as Structured Prompt for Agent
Please help me install this Agent Skill.
Skill: fla-triton-to-gluon
Source: https://github.com/weichengz0616/fla/tree/main/.agents/skills/fla-triton-to-gluon
Command: npx skills add https://github.com/weichengz0616/fla --skill fla-triton-to-gluon-weichengz0616

SYSTEM DOCUMENTATION & REQUIREMENTS

What problem does it solve? Triton kernels can hit performance ceilings when register spills cap block sizes, num_stages pipelining underperforms, or Hopper/Blackwell hardware features like TMA, TMEM, and tcgen05 are needed but not exposed. This Skill provides a disciplined workflow for porting those kernels to Gluon while keeping numerical parity at every step. ## Core Features & Use Cases - Incremental Porting Sequence: Freeze the baseline pytest contract, do a literal translation with explicit layouts, then layer in cp.async/TMA data movement, WGMMA/tcgen05 MMA, and scheduling (persistent kernels, warp specialization) driven by profiling. - Triton-to-Gluon API Mapping: A concrete table mapping tl.load/tl.dot/num_stages to their Gluon equivalents, plus guidance on layouts, shared memory budgets, autotune pruning, and compile-time management for heavily unrolled kernels. - Pitfall Checklist: Covers proxy fences, mbarrier phase semantics, tma.store_wait behavior, bitwise-cancellation traps, and NaN-poisoned out-of-bounds handling. - Use Case: A bandwidth-bound attention kernel in fla/ops is register-bound on an H100. Use this Skill to port it to Gluon with TMA loads and per-slot mbarrier pipelining, verifying forward and backward parity against the frozen pytest after each layer. ## Quick Start Port the Triton kernel in fla/ops/attnres to Gluon following the incremental sequence, keeping the existing pytest passing after each step.

Frequently Asked Questions about fla-triton-to-gluon

High-intent search queries and answers about installing and using this skill.

FAQPage Schema
How do I port a Triton kernel to Gluon?

Start with a literal translation using gl.load/gl.store and explicit layouts, keeping the op's frozen pytest passing. Then add async data movement (cp.async or TMA), MMA via warpgroup_mma or tcgen05_mma, and scheduling layers, re-autotuning and verifying parity after each step.

When is porting a Triton kernel to Gluon worth it?

Port when register spills cap block sizes, num_stages pipelining fails to overlap load and compute, or you need Hopper/Blackwell features like TMA, TMEM, tcgen05, or warp specialization. Skip it if the kernel already saturates bandwidth or the bottleneck is algorithmic.

What is the difference between Triton and Gluon?

Gluon shares Triton's compiler stack, JIT, and SPMD tile model, and host-side launch code is unchanged. The difference is that layouts, shared memory allocation, asynchrony, and synchronization are all explicit rather than compiler-managed.

Does Gluon work on AMD GPUs?

Gluon's nvidia modules (ampere, hopper, blackwell) are NVIDIA-only; AMD is a separate submodule. Hardware gating applies: cp.async needs Ampere+, TMA and WGMMA need Hopper+, and TMEM and tcgen05 need Blackwell.

Why does my Gluon kernel produce wrong values without crashing?

Wrong values without a crash usually indicate a layout broadcast or conversion misunderstanding. Debug with gl.static_print on layouts and convert_layout(..., assert_trivial=True) to prove a conversion is actually free.

Why is my Gluon kernel slower after adding optimizations?

Common causes are a blown register budget in warp-specialized partitions, a cross-warp convert_layout silently routing through shared memory, or a persistent schedule tanking L2 hit rate. Check lts__t_sector_hit_rate and stall metrics in NCU.