fla-triton-to-gluon

Ports Triton kernels in fla/ops to Gluon for explicit control over layouts, shared memory, and MMA scheduling.

Updated Sep 11, 2026
One-click install
npx skills add https://github.com/swiss-ai/flash-linear-attention --skill fla-triton-to-gluon-swiss-ai
Or copy as Structured Prompt for Agent
Please help me install this Agent Skill.
Skill: fla-triton-to-gluon
Source: https://github.com/swiss-ai/flash-linear-attention/tree/main/.agents/skills/fla-triton-to-gluon
Command: npx skills add https://github.com/swiss-ai/flash-linear-attention --skill fla-triton-to-gluon-swiss-ai

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 unreachable. This Skill provides a disciplined workflow for porting existing Triton kernels in fla/ops/** to Gluon (triton.experimental.gluon) while preserving 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 async data movement, WGMMA/tcgen05 MMA, and persistent/warp-specialized scheduling, re-autotuning after each layer. - Triton-to-Gluon API mapping: a reference table covering @gluon.jit, gl.load/gl.store, BlockedLayout, NVMMASharedLayout, mbarrier pipelining, and manual multi-buffering replacing num_stages. - Compile-time and smem-budget management: autotune config pruning for heavily unrolled kernels, constexpr-switched resident vs streaming designs under the 228KB shared memory cap, and a pitfall checklist covering proxy fences, mbarrier phases, bitwise-cancellation traps, and NaN-poisoned OOB handling. - Use Case: A Gated DeltaNet kernel is register-bound and cannot grow its block size. Use this Skill to port it to Gluon with TMA loads tracked by mbarriers, freeing registers so blocks can grow, then verify forward and backward parity with the frozen pytest before benchmarking against the Triton baseline. ## Quick Start Port the Triton kernel in fla/ops/<op> to Gluon following the incremental sequence, keeping the frozen pytest passing at every 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?

Port incrementally: freeze the existing pytest as a parity contract, write a literal translation with explicit layouts using gl.load/gl.store, then add cp.async or TMA async data movement, WGMMA or tcgen05 MMA, and scheduling layers. Re-autotune after every layer and keep numerical parity at 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 when the kernel already saturates bandwidth or the bottleneck is algorithmic.

Does Gluon work on AMD GPUs?

Gluon's nvidia modules (ampere, hopper, blackwell) are NVIDIA-only; AMD is a separate submodule. cp.async needs Ampere+, TMA and warp specialization need Hopper+, and TMEM, tcgen05, and CLC require Blackwell.

Why does my Gluon kernel produce wrong values without crashing?

Wrong values without a crash usually come from layout broadcast or conversion misunderstandings. Debug with gl.static_print on layouts and convert_layout(..., assert_trivial=True) to prove a conversion is free, and check proxy fences between generic and async proxy accesses.

Why is my Gluon kernel slower after adding optimizations?

Common causes are a blown register budget in warp-specialized kernels, 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 stalled_long_scoreboard in NCU to diagnose.

How do I handle shared memory limits in Gluon kernels?

Shared memory is capped at 228KB per SM on Hopper/Blackwell; budget around 192KB. Put resident and streaming double-buffer designs in one kernel behind a constexpr switch, and prune autotune configs whose minimal footprint exceeds the budget.