Agent skill

Fla Kda

by fla-org in fla-org/flash-linear-attention

FLA KDA kernel workflow and public technical notes. An agent skill from fla-org/flash-linear-attention.

MITAuto-check passed

Install Fla Kda

skills CLI
$ npx skills add fla-org/flash-linear-attention --skill fla-kda -a claude-code

Project install by default; add -g for ~/.claude/skills/.

GitHub CLI
$ gh skill install fla-org/flash-linear-attention fla-kda --agent claude-code

Project scope by default; add --scope user for a personal install. Needs GitHub CLI 2.90.0 or later (public preview).

Manual copy
$ git clone --depth 1 https://github.com/fla-org/flash-linear-attention.git skills-src && mkdir -p .claude/skills && cp -r skills-src/.agents/skills/fla-kda .claude/skills/fla-kda && rm -rf skills-src

Use ~/.claude/skills/ instead of .claude/skills for a personal install. The folder must contain SKILL.md.

Claude Code skills documentation · loads skills from .claude/skills/

Facts

Skill name
fla-kda
GitHub stars
5.8k
Token cost
~1.3k tokens
SKILL.md length
563 words
Files
1
Skills in repo
9
Repo updated
First seen
Licence
MIT

At a glance

FLA KDA kernel workflow and public technical notes. An agent skill from fla-org/flash-linear-attention.

  • Works in 2 steps: Pre-gated mode: use_gate_in_kernel=False. → In-kernel mode: use_gate_in_kernel=True.
  • Reviewing fla/ops/kda/
  • SKILL.md covers Public code map, Gate modes, Safe gate numerical note and Safe vs non-safe intra path, plus 2 more sections
  • Instructions only: no scripts, shell commands, URLs or credentials in SKILL.md

What it does

Fla Kda is an agent skill from fla-org/flash-linear-attention. FLA KDA kernel workflow and public technical notes. Use when modifying or reviewing fla/ops/kda/, KDA gate modes, chunk intra/inter kernels, safegate behavior, KDA backends, or KDA-specific tests and benchmarks.

Its SKILL.md is about 1.3k tokens, which your agent loads only when the skill is triggered. It is a single SKILL.md file with no bundled scripts.

The repository describes itself as: 🚀 Efficient implementations for emerging model architectures. The licence is MIT.

When your agent uses it

  • Reviewing fla/ops/kda/
  • Chunk intra/inter kernels
  • Safegate behavior
  • KDA-specific tests and benchmarks

Example prompts

  • “/fla-kda”

Workflow steps

2 steps, taken from the first numbered list in SKILL.md.

  1. Pre-gated mode: use_gate_in_kernel=False.
  2. In-kernel mode: use_gate_in_kernel=True.

What it can do on your machine

Read from SKILL.md and the folder at commit b8ff848. It shows what the files ask for, not the result of running them.

  • Tool permissions

    Pre-approves nothing: there is no allowed-tools line, so your agent's usual permission prompts apply.

    From allowed-tools in the SKILL.md frontmatter.

  • Runs code

    No scripts in the folder and no shell commands in SKILL.md.

    From the folder's file list and the shell code blocks in SKILL.md.

  • Network

    No URLs in SKILL.md.

    From URLs in SKILL.md, links to its own repository left out.

  • Credentials

    Names no API keys, tokens, secrets or passwords.

    From names ending in _API_KEY, _TOKEN, _SECRET, _KEY or _PASSWORD in SKILL.md.

Context cost

Fla Kda loads about 1.3k tokens when it runs. Until then it costs about 56 tokens; SKILL.md has 563 words of instructions outside code blocks.

Always · name and description, kept in context so the agent knows when to use it
~56
When it runs · the whole SKILL.md, loaded when a task matches
~1.3k

Estimates: characters ÷ 4, the usual rule of thumb; real counts depend on the model's tokenizer. Scripts and assets cost tokens only if the agent reads them.

Safety

Auto-check passed

The automated check found no risky patterns in SKILL.md.

Automated static check — not a guarantee. Review scripts before installing. It scans the text of SKILL.md for risky patterns (piping downloads into a shell, reading credential files, hidden Unicode, destructive commands); files beside SKILL.md are not scanned.

SKILL.md

The full file from fla-org/flash-linear-attention at commit b8ff848, republished under its MIT licence (© fla-org). 563 words, ~1,297 tokens.

Download SKILL.mdSave it as .claude/skills/fla-kda/SKILL.md (or your agent's skills folder).
name
fla-kda
description
FLA KDA kernel workflow and public technical notes. Use when modifying or reviewing fla/ops/kda/**, KDA gate modes, chunk intra/inter kernels, safe_gate behavior, KDA backends, or KDA-specific tests and benchmarks.

FLA KDA Skill

Use this skill for KDA-specific work under fla/ops/kda/** and tests that exercise KDA behavior.

Public code map

  • Public API: fla.ops.kda.chunk_kda, fla.ops.kda.fused_recurrent_kda.
  • Gate helpers: naive_kda_gate, naive_kda_lowerbound_gate, kda_gate_fwd, kda_gate_bwd, fused_kda_gate, kda_gate_chunk_cumsum in fla/ops/kda/gate.py.
  • Chunk forward: chunk_kda_fwd in chunk_fwd.py.
  • Intra/inter forward: chunk_kda_fwd_intra, chunk_kda_fwd_kernel_intra_sub_chunk, chunk_kda_fwd_kernel_inter_solve_fused in chunk_intra.py.
  • Token-parallel non-safe path: chunk_kda_fwd_intra_token_parallel in chunk_intra_token_parallel.py.
  • WY recompute: recompute_w_u_fwd and recompute_w_u_fwd_kda_kernel in wy_fast.py.
  • Backward: chunk_kda_bwd, chunk_kda_bwd_intra, chunk_kda_bwd_wy_dqkg_fused.
  • Backends: FlashKDABackend, KDATileLangBackend.

Gate modes

chunk_kda has two gate input contracts:

  1. Pre-gated mode: use_gate_in_kernel=False.
    • g is already the log-space decay tensor.
    • A_log, dt_bias, and lower_bound are not part of the gate activation.
  2. In-kernel mode: use_gate_in_kernel=True.
    • g is raw gate input.
    • A_log is required and dt_bias is optional.
    • Without safe_gate, activation is -exp(A_log) * softplus(g + dt_bias).
    • With safe_gate, activation is lower_bound * sigmoid(exp(A_log) * (g + dt_bias)).

safe_gate=True requires use_gate_in_kernel=True, lower_bound is not None, and -5 <= lower_bound < 0.

Safe gate numerical note

With lower_bound=-5, every per-token gate value is in [-5, 0) before the RCP_LN2 conversion used by chunk_kda_fwd. A 16-token sub-chunk can therefore accumulate -80 in natural-log units. Directly feeding the full span to exp2 would be larger in base-2 units, so the safe intra path relies on offsetting.

chunk_kda_fwd_kernel_intra_sub_chunk uses a midpoint offset before exponentiation:

  • b_gm = b_g - b_gn;
  • exp2(b_gm) and exp2(-b_gm).

With the midpoint offset, each exponent operand covers at most about half of the 16-token sub-chunk. Under lower_bound=-5, this is about 40 / ln(2), which is below the kernel's exp2 safety comment threshold. The important invariant is not the raw cumulative value alone; it is that each exponentiation uses a local offset rather than the full chunk cumsum directly.

For inter-subchunk work, chunk_kda_fwd_kernel_inter_solve_fused computes decay ratios with paired offsets such as:

  • exp2(b_g1 - b_gn1) and exp2(b_gn1 - b_g0);
  • exp2(b_g2 - b_gn2) and exp2(b_gn2 - b_g1).

Both terms are non-positive under monotonic accumulated decay, so the off-diagonal inter path avoids positive exponent growth. The triangular solve operates on masked lower-triangular blocks, so it does not introduce an unbounded exponent path.

Show full SKILL.md (246 more words)Show less

Safe vs non-safe intra path

  • Safe path: chunk_kda_fwd_intra(..., safe_gate=True) calls chunk_kda_fwd_kernel_intra_sub_chunk for 16-token diagonal blocks, then calls chunk_kda_fwd_kernel_inter_solve_fused with USE_SAFE_GATE=True.
  • Non-safe path: safe_gate=False calls chunk_kda_fwd_intra_token_parallel for diagonal blocks, then calls the same inter/solve kernel with USE_SAFE_GATE=False.
  • Do not change one path without checking the other path unless the contract is explicitly safe-only or non-safe-only.

Correctness checklist

Before finishing a KDA behavior change, use fla-correctness-coverage and cover only axes affected by the change:

  • dense and varlen sequence layout;
  • forward and backward if training path is touched;
  • pre-gated, non-safe in-kernel, and safe in-kernel gate modes where supported;
  • raw beta logits and post-sigmoid beta where supported;
  • use_qk_l2norm_in_kernel=True/False where relevant;
  • MHA and GVA (HV > H);
  • D != Dv when value dimension is involved;
  • initial/final state, return_intermediate_states, and CP paths when touched;
  • backend verifier behavior for FlashKDA / TileLang changes.
  • gate numerical extremes when gate math or intra/inter decay is touched: lower_bound=-5, a lower bound close to 0, large positive and negative g + dt_bias, extreme A_log, long-sequence cumulative decay, chunk boundaries, and ragged varlen boundaries.

Style constraints

  • Use platform helpers from fla.utils (device, device_platform, IS_NVIDIA, IS_NVIDIA_HOPPER, IS_NVIDIA_BLACKWELL, IS_AMD, IS_INTEL) instead of adding new direct torch.cuda platform checks in tests or public code. If no helper covers the condition, add one in fla.utils first.
  • Keep math derivations in operator docs or PR text; in Triton kernels, prefer compact shape comments and one-line rationale comments.
  • Do not include internal-only paths, private model names, local machine paths, or private workload identifiers in public tests or skills.

© fla-org, MIT. Rendered from Markdown: HTML in the file is shown as text, images as links, and headings moved down two levels. Raw file

Files

Just SKILL.md in .agents/skills/fla-kda of fla-org/flash-linear-attention.

Open the folder on GitHubat commit b8ff848

Compare with similar skills

Fla Kda next to the 5 skills that share the most tags, products or categories with it. Stars are the repository's; “used in” counts other GitHub owners with a copy.

Fla Kda compared with similar skills
SkillStarsUsed inTokensAuto-checkLicenceRepo updated
Fla Kda this skillfla-org/flash-linear-attention5.8k—~1.3kAutomated safety check: PassMIT
Kernel Organizationsgl-project/sglang37k—~1.3kAutomated safety check: PassApache-2.0
Metal Kernelpytorch/pytorch104k—~4.9kAutomated safety check: PassCustom licence
Semantic Kernelgithub/awesome-copilot40k2 repos~756Automated safety check: PassMIT
Add Sgl Kernelsgl-project/sglang37k2 repos~3.4kAutomated safety check: PassApache-2.0
Add Jit Kernelsgl-project/sglang37k—~13kAutomated safety check: PassApache-2.0

Similar skills

  • Kernel Organization

    sgl-project/sglang

    Apply the SGLang kernels RFC when adding, moving, splitting, or reviewing kernel APIs, registry metadata, kernel tests, benchmarks, and model-specific implementations.

    37k GitHub stars~1.3k tokensUpdated today
    AI & LLM EngineeringAuto-check passed
  • Metal Kernel

    pytorch/pytorch

    Write Metal/MPS kernels for PyTorch operators. An agent skill from pytorch/pytorch.

    104k GitHub stars~4.9k tokensUpdated today
    AI & LLM EngineeringAuto-check passed
  • Semantic Kernel

    github/awesome-copilot

    Official

    Create, update, refactor, explain, or review Semantic Kernel solutions using shared guidance plus language-specific references for .NET and Python.

    40k GitHub starsUsed in 2 repos~756 tokens
    DevelopmentAuto-check passed
  • Add Sgl Kernel

    sgl-project/sglang

    Step-by-step tutorial for adding a heavyweight AOT CUDA/C++ kernel to sgl-kernel (including tests & benchmarks)

    37k GitHub starsUsed in 2 repos~3.4k tokens
    AI & LLM EngineeringAuto-check passed
  • Add Jit Kernel

    sgl-project/sglang

    Step-by-step tutorial for adding a new lightweight JIT CUDA kernel to sglang.kernels JIT infrastructure and public operator groups

    37k GitHub stars~13k tokensUpdated today
    AI & LLM EngineeringAuto-check passed
  • Paddle Phi Kernel

    PaddlePaddle/Paddle

    A skill your agent uses when working with Paddle's PHI kernel system: registering new kernels, debugging kernel selection/dispatch, understanding code auto-generation from YAML, or implementing…

    24k GitHub stars~656 tokensUpdated 7 days ago
    DevelopmentAuto-check passed

More from fla-org/flash-linear-attention

All 9 skills in this repo
  • Fla Ascend Performance

    fla-org/flash-linear-attention

    Guidelines for Ascend NPU kernel / Triton-Ascend backend performance work in the FLA repo.

    5.8k GitHub stars~5.6k tokensUpdated yesterday
    Auto-check passed
  • Fla Optimization Loop

    fla-org/flash-linear-attention

    Disciplined, reproducible loop for making an FLA kernel faster (Triton, Gluon, TileLang, CuTe) without ever breaking or gaming correctness.

    5.8k GitHub stars~2.6k tokensUpdated yesterday
    Auto-check passed
  • Fla Triton To Gluon

    fla-org/flash-linear-attention

    Workflow for porting an existing Triton kernel in fla/ops/ to Gluon (triton.experimental.gluon) to gain explicit control over tensor layouts, shared memory, async data movement (cp.async / TMA), MMA…

    5.8k GitHub stars~4.2k tokensUpdated yesterday
    Auto-check passed
  • Fla Correctness Coverage

    fla-org/flash-linear-attention

    Guidelines for kernel correctness testing and coverage in fla/ops/ and related modules, including common Triton grid/addressing pitfalls.

    5.8k GitHub stars~1.2k tokensUpdated yesterday
    Auto-check passed
  • Fla Design Coverage

    fla-org/flash-linear-attention

    Contract-first design and coverage discipline for FLA kernel and numerical changes.

    5.8k GitHub stars~3.3k tokensUpdated yesterday
    Auto-check passed
  • Fla Dispatch Backends

    fla-org/flash-linear-attention

    Workflow for FLA backend dispatch decorators and backend implementations.

    5.8k GitHub stars~1.1k tokensUpdated yesterday
    Auto-check passed

Questions about Fla Kda

What does Fla Kda do?

FLA KDA kernel workflow and public technical notes. An agent skill from fla-org/flash-linear-attention. Fla Kda is an agent skill from fla-org/flash-linear-attention. FLA KDA kernel workflow and public technical notes.

When should I use Fla Kda?

Fla Kda fits situations like: reviewing fla/ops/kda/; chunk intra/inter kernels; safegate behavior; KDA-specific tests and benchmarks.

How do I install Fla Kda in Claude Code?

Run `npx skills add fla-org/flash-linear-attention --skill fla-kda -a claude-code`. Or copy the skill folder (.agents/skills/fla-kda in fla-org/flash-linear-attention) into .claude/skills/fla-kda in your project. Claude Code loads it when a task matches its description.

How do I install Fla Kda in Codex?

Run `npx skills add fla-org/flash-linear-attention --skill fla-kda -a codex`. Or copy the skill folder (.agents/skills/fla-kda in fla-org/flash-linear-attention) into .agents/skills/fla-kda in your project. Codex loads it when a task matches its description.

Can I use Fla Kda in Cursor, Gemini CLI or GitHub Copilot?

Cursor, Gemini CLI, GitHub Copilot and OpenCode also load SKILL.md folders. With the skills CLI, run `npx skills add fla-org/flash-linear-attention --skill fla-kda -a cursor` (or -a gemini-cli, github-copilot or opencode for the others). To copy it by hand, put the folder in .cursor/skills/fla-kda, .gemini/skills/fla-kda, .github/skills/fla-kda and .opencode/skills/fla-kda in your project.

What does Fla Kda need to run?

SKILL.md names no scripts, command-line tools or credentials: Fla Kda is instructions for the agent only.

Does Fla Kda access the network?

SKILL.md contains no URLs. Any network use would come from the scripts or tools the agent runs. This is read from the text; nothing was executed.

Is Fla Kda safe to install?

Our automated static check of SKILL.md found no risky patterns, such as piping downloads into a shell, reading credential files or hidden Unicode. It is not a guarantee. Review the folder before installing.

What licence does Fla Kda use?

Fla Kda is published under the MIT licence (the repository's licence). It allows redistribution, so the full SKILL.md is shown on this page.

How many tokens does Fla Kda use?

About 1.3k tokens (SKILL.md is roughly 5.2k characters). Agents keep only the skill's name and description in context until a task matches; then they load SKILL.md in full.

What are the alternatives to Fla Kda?

Skills that share tags, products or a category with Fla Kda: Kernel Organization (sgl-project/sglang, 37k stars), Metal Kernel (pytorch/pytorch, 104k stars), Semantic Kernel (github/awesome-copilot, 40k stars) and Add Sgl Kernel (sgl-project/sglang, 37k stars). The comparison table on this page puts their stars, adoption, token cost, safety result and licence side by side.

Who maintains Fla Kda?

fla-org (a GitHub organization) maintains it in fla-org/flash-linear-attention, which has 5,828 GitHub stars. The repository holds 9 skills in this directory. The repository was last updated on October 6, 2026.

Source: fla-org/flash-linear-attention on GitHub. Facts on this page come from the repository at the commit we read; the author's words are quoted as theirs.