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.