{"owner":"fla-org","repo":"flash-linear-attention","hasSkills":true,"totalSkillsCount":23,"totalTokensCount":16710,"categories":["anthropic-skill","root-instruction"],"hasMcp":false,"mcpConfig":null,"found":[".agents/skills/README.md",".agents/skills/fla-ascend-performance/SKILL.md",".agents/skills/fla-ascend-performance/references/TRAPS.md",".agents/skills/fla-ascend-performance/references/cases.md",".agents/skills/fla-ascend-performance/references/g-contiguous-loading.md",".agents/skills/fla-ascend-performance/references/reference.md",".agents/skills/fla-correctness-coverage/SKILL.md",".agents/skills/fla-correctness-coverage/references/cp.md",".agents/skills/fla-correctness-coverage/references/delta-rule.md",".agents/skills/fla-correctness-coverage/references/generalized-delta-rule.md",".agents/skills/fla-correctness-coverage/references/simple-gla.md",".agents/skills/fla-design-coverage/SKILL.md",".agents/skills/fla-dispatch-backends/SKILL.md",".agents/skills/fla-kda/SKILL.md",".agents/skills/fla-nvidia-performance/SKILL.md",".agents/skills/fla-optimization-loop/SKILL.md",".agents/skills/fla-optimization-loop/references/TRAPS.md",".agents/skills/fla-optimization-loop/references/opt-log-template.md",".agents/skills/fla-pr-readiness/SKILL.md",".agents/skills/fla-pr-stacks/SKILL.md",".agents/skills/fla-triton-to-gluon/SKILL.md","AGENTS.md","evals/README.md"],"skills":{".agents/skills/README.md":"# Repository skills\n\n[CONTRIBUTING.md](../../CONTRIBUTING.md) defines contribution rules; [AGENTS.md](../../AGENTS.md) defines agent operations. Skills add task-specific guidance. Choose those relevant to the change.\n\nFor kernel work, define the contract and establish baseline coverage before optimization. Use the operator, dispatch, and hardware guides as needed, then prepare the PR with the collected evidence.\n\n| Skill                                                     | Use for                                                    |\n| --------------------------------------------------------- | ---------------------------------------------------------- |\n| [Design coverage](fla-design-coverage/SKILL.md)           | Supported inputs, numerical budgets, and routing contracts |\n| [Correctness coverage](fla-correctness-coverage/SKILL.md) | References, regression cases, and affected test selection  |\n| [KDA](fla-kda/SKILL.md)                                   | KDA gate modes and kernel invariants                       |\n| [Backend dispatch](fla-dispatch-backends/SKILL.md)        | Registration, verifiers, and fallback behavior             |\n| [Optimization loop](fla-optimization-loop/SKILL.md)       | A fixed correctness gate and measured candidate iterations |\n| [NVIDIA performance](fla-nvidia-performance/SKILL.md)     | GPU profiling and bottleneck diagnosis                     |\n| [Ascend performance](fla-ascend-performance/SKILL.md)     | NPU profiling and compiler constraints                     |\n| [Triton to Gluon](fla-triton-to-gluon/SKILL.md)           | Layouts, synchronization, and incremental kernel ports     |\n| [PR readiness](fla-pr-readiness/SKILL.md)                 | Scope, validation evidence, and submission                 |\n| [PR stacks](fla-pr-stacks/SKILL.md)                       | Splitting large changes, reusing PRs, and native stacks    |\n\n## Maintaining skills\n\nPut each workflow in `.agents/skills/<skill-name>/SKILL.md` with YAML `name` and `description` fields. Follow [documentation and skill style](../../CONTRIBUTING.md#prose-and-markdown) for scope, organization, wording, and shared policy.\n\nUse an existing `references/` directory for substantial task-specific detail. Add supporting files only when they make the workflow easier to use; `SKILL.md` is the entry point, so individual skill directories do not need a README. Symlinks may point to public documentation tracked in this repository.\n",".agents/skills/fla-ascend-performance/SKILL.md":"---\nname: fla-ascend-performance\ndescription: Profile and optimize FLA Triton-Ascend kernels using NPU traces, with guidance for UB capacity, memory movement, launch limits, and numerical correctness.\n---\n\n# Ascend kernel performance\n\nUse this skill for operator performance work in `triton_ascend.py` and `triton_ascend/` backend implementations. Follow [fla-optimization-loop](../fla-optimization-loop/SKILL.md) for sustained optimization and [CONTRIBUTING](../../../CONTRIBUTING.md#benchmarking) for validation and reporting.\n\nKeep collection, analysis, and benchmarks in the active Python/NPU environment. If it is not configured, activate the project's Ascend environment first.\n\n## 1. Establish the baseline\n\nLocate the public entry, dispatch route, and Ascend kernel. Record the supported inputs and production workloads, baseline commit, and synchronized latency. Preserve the public API, validated algorithm, reference, and tolerances. Confirm the intended NPU kernel runs; a Torch fallback must not hide an unsupported path or kernel bug.\n\n## 2. Collect a trace\n\nUse the existing collector instead of duplicating profiler setup. Run from the repository root; the workload file initializes its inputs and defines a repeatable `workload()` callable.\n\n```bash\npython .agents/skills/fla-ascend-performance/scripts/profile_npu.py \\\n  --name my_op --out-dir profile/my_op-npu \\\n  --metrics PipeUtilization --analyze \\\n  --kernel-filter my_kernel --exec-file path/to/workload.py\n```\n\nOne run collects one `aic_metrics` set. Start with `PipeUtilization`; collect `MemoryUB` separately when UB bandwidth is relevant. The default schedule is one warmup step and one active step. Warm the workload enough to exclude compilation from the measurement.\n\nThe collector prints the exact trace directory. To revisit it, pass that directory to [the analyzer](scripts/analyze_profile.py); see [collection and metrics](references/reference.md#collection-and-analysis) for the library interface and output layout.\n\n## 3. Diagnose bottlenecks\n\nUse `op_statistic.csv` to find the dominant operation, then `kernel_details.csv` to inspect its duration and pipe activity. Confirm the dispatch route before tuning a small or missing target kernel. Distinguish compilation failure, UB overflow, grid limits, numerical errors, and measured performance limits; they need different fixes.\n\nUse the [diagnosis table](references/reference.md#diagnosis) to choose the next experiment. Change one relevant factor, then measure whether the expected duration and pipe metrics moved.\n\n## 4. Optimize within the kernel contract\n\n- Estimate the peak live UB footprint and use `fla.utils.ascend_ub_manager` for device capacity and tiling. Model forward and backward separately; fuse only when the live set fits.\n- Match the launch to the work: Cube kernels use Cube cores, Vector kernels use Vector cores. Handle the grid limit through the existing helpers or a 1D task loop.\n- Preserve numerical precision, optional inputs, and fixed/variable-length behavior. Verify tail handling and pointer arithmetic before increasing tiles or fusing stages.\n- Omit `num_warps` and `num_stages` from Ascend launches and autotune settings; these are unsupported on this backend.\n- Check [compiler and memory traps](references/TRAPS.md) when changing DMA paths, pointer arithmetic, or reused `tl.dot` operands. Use [contiguous gate loading](references/g-contiguous-loading.md) when gate loads along time are strided.\n\n[Kernel cases](references/cases.md) explain existing convolution, recurrence, and solve implementations. Their tile choices are workload-specific examples.\n\n## 5. Validate and measure\n\nRun kernel comparisons and the affected operator/layer tests, including supported gradients, variable lengths, optional inputs, and boundary shapes. Keep NaN poisoning enabled and confirm that dispatch exercises the changed path. Run the full relevant correctness gate before reporting a speedup; a `--gate-k` subset is an iteration check.\n\nCompare synchronized forward and forward/backward latency on the same NPU, then repeat the relevant profiler collection. Report the timing change, the metrics that explain it, any compiler limitation encountered, and untested cases. If the expected metrics do not change, revisit the diagnosis before adding another optimization.\n",".agents/skills/fla-ascend-performance/references/TRAPS.md":"# Ascend compiler and memory traps\n\nRead the relevant section before changing DMA, launch indexing, or reused matrix operands. [Kernel cases](cases.md) show these constraints in existing implementations.\n\n## DMA paths and optional pointers\n\nA runtime branch between block-pointer DMA and masked loads can keep both paths live in UB. Larger tiles may then overflow even when measured UB bandwidth is low. Where the workload permits, split bulk and tail launches with a constexpr mode so each compilation removes the unused path. Include convolution halos when deciding whether a bulk load fits inside the packed allocation.\n\nAn optional pointer needs a compile-time guard before any runtime condition. `if USE_INITIAL_STATE or runtime_condition:` can still compile pointer arithmetic in the other branch when the pointer is `None`; nest the constexpr condition separately.\n\n## Address arithmetic\n\nCast indices to `tl.int64` before multiplying by strides or dimensions. Casting the product cannot repair overflow. On Ascend, specialized arguments and folded program IDs may be constexpr values without a `.to()` method; `tl.cast` works for both runtime and constexpr integers.\n\n```python\ntoken_offset = tl.cast(i_t, tl.int64) * BT\nbatch_offset = tl.cast(i_b, tl.int64) * T\nbos = tl.load(cu_seqlens + i_n).to(tl.int64)\n```\n\nLoad packed sequence offsets as int64 even when the input `cu_seqlens` uses int32. A small token index can overflow after multiplication by head count and feature size.\n\n`make_block_ptr` metadata has a separate constraint: its `offsets` and `block_shape` must remain int32. Keep large flattened base-address calculations in int64, then supply valid local block offsets. The Ascend backend is exempt from the mainline ban on block pointers.\n\n## Grid and task loops\n\nThe existing Ascend grid helpers enforce a grid-product limit of 65535. Use their host splitting or flatten independent tiles into a 1D loop over available cores. `get_multiprocessor_count()` selects Vector cores on NPU; `use_aicore=True` selects Cube cores.\n\nAfter slicing variable-length chunk metadata, apply the corresponding global offset only once. In a task loop, derive local pointers from the original base on each iteration; accumulating pointer updates across tasks can miscompile on Ascend.\n\n## Reused `tl.dot` operands\n\nTriton-Ascend can reuse the left operand's UB storage during `tl.dot`. A later read of that tile can therefore be wrong even without a compiler error. Audit all later uses, including another left operand, a right operand, arithmetic, and stores.\n\nReload the pristine tile from global memory when reuse is separated by stages. For tight reuse, create the needed `tile + 0.0` copies before the first dot that consumes the original as a left operand. Copying afterwards preserves the corrupted value. Account for the extra live copies in the UB budget and validate against the numerical reference.\n\n[The case index](cases.md#tldot-lhs-clobber--repo-wide-case-catalog) identifies the affected kernel patterns. A fresh load on each loop iteration or a tile used as a left operand only once needs no extra copy for this issue.\n\n## Numerical and tail behavior\n\nPreserve the validated accumulation precision and exponential base. Mask invalid values before exponentials and initialize every region consumed later. Factoring a gate difference into a ratio or reciprocal changes rounding and range behavior; check the existing reference and tolerances before using it.\n\nExercise gated and ungated paths when both are supported. A suite that always supplies a gate cannot catch `None` pointer handling or incorrect ungated residual updates.\n",".agents/skills/fla-ascend-performance/references/cases.md":"# Ascend kernel cases\n\nUse these examples to understand existing design choices. Read the current implementation before changing a path; tile settings depend on the workload and compiler. Paths are relative to the repository root.\n\n## Causal convolution: core grid and DMA paths\n\nImplementation: `fla/modules/causal_conv1d/backends/triton_ascend.py`.\n\nThe contiguous packed training path uses a 1D Vector-core task loop when channel tiling and state requirements permit. Other layouts, channel dimensions, and cache-state calls retain separate paths.\n\n- Transpose the public convolution weight to make channels contiguous for block loads. Preserve the public weight and output layouts.\n- Load one window including the convolution halo and use `extract_slice` for taps where supported. Keeping every tap live at once can exceed UB.\n- Preserve the saved pre-activation used by backward for silu/swish; removing it can trigger a second forward computation.\n- Separate bulk and tail DMA at compile time. A runtime choice keeps both paths live and can prevent larger tiles even when UB bandwidth is not saturated.\n\n| Path                               | DMA policy                                                     |\n| ---------------------------------- | -------------------------------------------------------------- |\n| Full chunks in packed input        | Block-pointer loads, with the tail path compiled out.          |\n| Final packed chunk                 | Masked loads/stores, including the halo bounds.                |\n| Variable lengths or a single chunk | Runtime bounds where a uniform bulk/tail split is unavailable. |\n\nUse `tests/modules/test_conv.py` for forward/backward, optional state, and layout validation. [TRAPS.md](TRAPS.md) covers optional pointers, packed bounds, and address types shared by these paths.\n\n## `chunk_delta_h.py` — bwd `dhu`\n\nImplementation: `fla/ops/common/backends/triton_ascend/chunk_delta_h.py`.\n\nGate loads along time use the [contiguous gate layout](g-contiguous-loading.md). The aligned fixed-length path can precompute gate terms, while variable-length or unaligned paths compute them inline. Model these paths separately: inline exponentials and temporary tiles can exceed the UB estimate for the precomputed path.\n\nChoose K and V tiles together. Reducing one dimension merely to maximize the other can increase repeated work. Rebuild task-local pointers from base addresses and return the corrected value gradient separately to keep the caller's `dv` unchanged. Backward also reuses dot operands across K slabs; apply the [operand lifetime rule](TRAPS.md#reused-tldot-operands).\n\n## `chunk_o.py` — fusion and gate loading\n\nImplementation: `fla/ops/common/backends/triton_ascend/chunk_o.py`.\n\nForward combines inter- and intra-chunk contributions when their live tiles fit UB, reducing intermediate output traffic. Backward gate loads use contiguous time storage while gradient outputs retain their public layout. Profile the combined launch and transpose cost, not just one kernel. See [gate loading](g-contiguous-loading.md) for pointer formulas and focused tests.\n\n## `chunk_bwd.py` — packed variable lengths\n\nImplementation: `fla/ops/kda/backends/triton_ascend/chunk_bwd.py`.\n\nPacked offsets multiplied by head and value dimensions can overflow int32 far before the token count reaches its limit. Load sequence boundaries as int64 and promote fixed-batch indices before multiplication. Only local lengths and block-pointer metadata should narrow to int32 when their bounds permit it.\n\n## `tl.dot` lhs clobber — repo-wide case catalog\n\nThe common rule is in [TRAPS.md](TRAPS.md#reused-tldot-operands). This index identifies where to inspect operand reuse without prescribing variable names or a fixed number of copies.\n\n| Implementation                                               | Reuse to inspect                                                                               |\n| ------------------------------------------------------------ | ---------------------------------------------------------------------------------------------- |\n| `fla/ops/common/backends/triton_ascend/chunk_o.py`           | Query/key tiles shared by multiple dots; attention tiles reused across value blocks.           |\n| `fla/ops/common/backends/triton_ascend/chunk_delta_h.py`     | Output-gradient and corrected residual tiles reused across K slabs.                            |\n| `fla/ops/gated_delta_rule/backends/triton_ascend/wy_fast.py` | Inverse tiles reused between stages; gradient tiles used first on the left, then on the right. |\n| `fla/ops/kda/backends/triton_ascend/wy_fast.py`              | Inverse tiles reused across the value and key loops.                                           |\n| `fla/ops/kda/backends/triton_ascend/chunk_intra.py`          | Triangular block merges that reuse operands for several products and final stores.             |\n| `fla/ops/utils/triton_ascend/solve_tril.py`                  | Block inverse operands used in multiple products or stored after a dot.                        |\n| `fla/ops/kda/backends/triton_ascend/chunk_bwd.py`            | Output gradients used as both left and right operands.                                         |\n\nAudit the changed kernel's first left-operand use and every subsequent read. Copies must precede that first use; reloads must come from pristine storage. Validate with the operator's kernel tests, including `tests/ops/test_gdn_kernels.py` and `tests/ops/test_solve_tril.py` where relevant.\n",".agents/skills/fla-ascend-performance/references/g-contiguous-loading.md":"# Contiguous gate loading\n\nUse this pattern when an Ascend kernel repeatedly loads gate values along time from `[B, T, HV]` storage. The time stride is `HV`, so these loads gather separated values. Transposing the input gate can reduce memory-transfer overhead; measure the transpose and kernel together.\n\nReference implementation: `fla/ops/common/backends/triton_ascend/chunk_o.py`.\n\n## Host layout\n\n```python\ng_t_contig = g is not None and HV != 1\ng_arg = g.transpose(1, 2).contiguous() if g_t_contig else g\n```\n\nPass the selected input and layout flag to the kernel. `HV == 1` is already contiguous along time. Keep gradient outputs in their original `[B, T, HV]` layout.\n\n## Kernel addressing\n\nSave the total packed length before computing a local sequence length. Promote batch, head, and packed sequence offsets before multiplying them; see [address arithmetic](TRAPS.md#address-arithmetic).\n\n| Input mode              | Contiguous gate base                    |\n| ----------------------- | --------------------------------------- |\n| Fixed batch             | `g + batch * HV * T_seq + head * T_seq` |\n| Packed variable lengths | `g + bos + head * T_seq`                |\n\nHere `batch`, `head`, and `bos` are int64 values, `T_seq` is the host-passed storage length, and the block-pointer bound uses the local sequence length. Block offsets are relative to that sequence's start.\n\nFor a contiguous gate base, use stride 1 along time. The original-layout branch instead uses `g + bos * HV + head` with time stride `HV`. Apply the base offset once and keep every load in a branch consistent with that layout. Reusing the original stride or offset after transposing produces wrong values or alignment faults.\n\n## Validation\n\nCompare both layout paths against the same reference, including head sharing, variable lengths, non-aligned tails, and long sequences. The focused tests are:\n\n- `tests/ops/test_gdn_kernels.py::test_chunk_bwd_dv_local`\n- `tests/ops/test_gdn_kernels.py::test_chunk_bwd_dqkwg`\n\nInspect wall-clock duration and MTE activity before and after, including transpose and gradient-finalization work. Sum-of-block profiler durations may differ from one launch's elapsed time.\n\nPreserve the exponential formula while changing layout. Replacing `exp2(a - b)` with a ratio can change overflow and rounding behavior. On Ascend, use the supported slicing/layout operations from the current implementation; a reshape or join that compiles does not prove the intended element ordering.\n",".agents/skills/fla-ascend-performance/references/reference.md":"# Ascend profiling reference\n\nUse this reference for collection options, CSV interpretation, and choosing the next optimization. Kernel correctness hazards are in [TRAPS.md](TRAPS.md); implementation examples are in [cases.md](cases.md).\n\n## Collection and analysis\n\nRun the scripts from the active NPU environment. `profile_npu.py` accepts a workload file or inline source defining `workload()`. It executes that callable repeatedly, so initialize reusable inputs outside the callable and make repeated forward/backward calls valid.\n\n```bash\npython .agents/skills/fla-ascend-performance/scripts/profile_npu.py \\\n  --name my_op --out-dir profile/my_op-npu \\\n  --metrics PipeUtilization --analyze --exec-file path/to/workload.py\n```\n\nFor library use, add the scripts directory to `PYTHONPATH` and call `profile_callable`:\n\n```bash\nexport PYTHONPATH=\"$PWD/.agents/skills/fla-ascend-performance/scripts${PYTHONPATH:+:$PYTHONPATH}\"\n```\n\n```python\nfrom profile_npu import profile_callable\n\ntrace_dir = profile_callable(fn=workload, name=\"my_op\", out_dir=\"profile/my_op-npu\", aic_metrics=\"PipeUtilization\")\n```\n\nThe returned directory contains the profiler output:\n\n```text\nmy_op_profiling_TIMESTAMP/\n└── localhost..._ascend_pt/\n    └── ASCEND_PROFILER_OUTPUT/\n        ├── op_statistic.csv\n        └── kernel_details.csv\n```\n\nPass that exact run directory when analyzing an existing trace. The analyzer takes the first matching CSV under the supplied directory, so a parent containing several runs can select an older result.\n\n```bash\npython .agents/skills/fla-ascend-performance/scripts/analyze_profile.py \\\n  profile/my_op-npu/my_op_profiling_TIMESTAMP --kernel-filter my_kernel --top-k 20\n```\n\n`--kernel-filter` applies to kernel details. If no name matches, the analyzer displays the top kernels instead; confirm the printed names before attributing metrics to the target.\n\n## Choosing metrics\n\nEach profiling run collects one `aic_metrics` set. Use separate runs for complementary metrics and keep the workload unchanged.\n\n| Metric set                                       | Use                                                                     |\n| ------------------------------------------------ | ----------------------------------------------------------------------- |\n| `PipeUtilization`                                | First pass: compare Cube, Vector, scalar, and memory-transfer activity. |\n| `MemoryUB`                                       | Check UB read/write bandwidth when memory traffic may limit the kernel. |\n| `Memory`, `MemoryAccess`, `MemoryL0`             | Inspect other memory paths.                                             |\n| `L2Cache`                                        | Investigate cache reuse.                                                |\n| `ArithmeticUtilization`, `ResourceConflictRatio` | Investigate arithmetic occupancy or resource conflicts.                 |\n\nAvailable metric names depend on the installed `torch_npu`; the collector reports available names when an option is unknown.\n\n## Diagnosis\n\nStart with total duration in `op_statistic.csv`, then inspect the dominant kernel in `kernel_details.csv`. Pipe activity alone does not establish the bottleneck.\n\n| Evidence                                               | Check next                                                                                                   |\n| ------------------------------------------------------ | ------------------------------------------------------------------------------------------------------------ |\n| Target kernel absent or many small fallback operations | Backend selection and fusion.                                                                                |\n| High Vector or Cube activity                           | Tile shape, useful work per tile, and scalar overhead around compute.                                        |\n| High MTE activity                                      | Reuse, intermediate writebacks, and strided loads; inspect [gate loading](g-contiguous-loading.md).          |\n| High scalar activity                                   | Address generation, branches, and repeated per-tile setup.                                                   |\n| High UB bandwidth with low compute activity            | Live buffers, layout conversions, and repeated UB reads/writes.                                              |\n| Low activity despite UB overflow at larger tiles       | Mutually exclusive DMA paths may remain live; see [DMA splitting](TRAPS.md#dma-paths-and-optional-pointers). |\n| Many host launches for grid chunks                     | Consider a 1D task loop with the appropriate Cube or Vector core count.                                      |\n\n## UB planning\n\nEstimate peak live storage from simultaneous buffers, their dtype, and any compiler-created copies. Use `fla.utils.ascend_ub_manager` for device capacity and tiling rather than hard-coding the UB size. The helper's safety margin reserves capacity; it is separate from the multiplier that estimates live buffers.\n\nIf more capacity is available, test a larger tile and verify compilation and numerics. If the live set already fills the budget, improve reuse or split stages. Model forward and backward separately; inline gate computation can need more UB than a precomputed-gate path.\n\nUseful code locations:\n\n- `fla/utils/ascend_ub_manager.py`: tiling and grid helpers.\n- `fla/utils/hardware.py`: `get_multiprocessor_count`, including `use_aicore=True` for Cube work.\n- `fla/modules/norm/triton_ascend/l2norm.py`: entry-point registration and the Ascend implementation.\n- `tests/ops/test_gdn_kernels.py`, `tests/ops/test_solve_tril.py`, and `tests/modules/test_conv.py`: relevant kernel comparisons.\n",".agents/skills/fla-correctness-coverage/SKILL.md":"---\nname: fla-correctness-coverage\ndescription: Select and run correctness coverage for FLA kernels and modules, including gradients, dispatch boundaries, and Triton addressing changes.\n---\n\n# FLA Correctness Coverage\n\nUse [fla-design-coverage](../fla-design-coverage/SKILL.md) when the change needs a new numerical or routing contract. This skill turns that contract into tests; [CONTRIBUTING.md](../../../CONTRIBUTING.md#testing) defines references, tolerances, test structure, and platform conventions.\n\n## Select coverage\n\nRead the affected tests and callers before adding cases. Follow [test organization and naming](../../../CONTRIBUTING.md#test-organization-and-naming): try the existing parameter matrix first, and add a function or file only when its setup or scope requires one. Cover reachable gaps across the applicable dimensions together, rather than testing each flag only in isolation:\n\n| Dimension   | Cases that can expose different behavior                                         |\n| ----------- | -------------------------------------------------------------------------------- |\n| Layout      | Dense and varlen, partial chunks, long sequences, asymmetric QK/value dimensions |\n| Features    | Gate bounds, raw/post-sigmoid beta, QK normalization, grouped value attention    |\n| State       | Initial state, final state, and state gradients where supported                  |\n| Numerics    | Supported dtypes, affected conversion/reduction paths, numerical boundaries      |\n| Routing     | Default and selected backends, verifier acceptance, rejection, and fallback      |\n\nCompare outputs and final states against the existing reference. Training APIs also need every supported gradient; forward-only APIs must state that backward is unsupported. Use `torch.autograd.gradcheck` where the implementation supports its required precision. Routing tests verify selection and argument forwarding; they do not replace numerical comparisons on the actual backend.\n\nFor addressing changes, inspect casts before multiplication and test shapes that expose large offsets, partial tiles, and varlen boundaries. Follow the `tl.int64` and platform requirements in [Triton Kernels](../../../CONTRIBUTING.md#triton-kernels). Exercise changed grid/address paths on affected supported platforms. Distinguish unsupported cases from tests skipped for unavailable hardware or dependencies; record the latter as validation gaps.\n\nThe NaN allocation guard covers eligible FLA allocations in operator/module tests, including the single-rank CP graph and intra-card cache cases moved from operator tests, but not layer/model tests, spawned CP workers, or every allocation. Add explicit finite-output and finite-gradient checks to adversarial cases, and investigate a poisoned-run failure even when an ad hoc run passes.\n\n## Run the affected paths\n\nDiscover dependent tests from the repository root:\n\n```bash\npython scripts/find_dependent_tests.py fla/ops/kda/chunk.py\nFLA_CI_ENV=0 pytest tests/ops/test_kda.py -v\n```\n\nThe helper lists files; it does not run them. Include reported callers and relevant `tests/modules/`, `tests/layers/`, `tests/models/`, or `tests/cp/` coverage when the change reaches those paths. Inspect skips and warning-only comparisons before calling the run a pass. Record commands, backend flags, environment, and results; reproduce failures on the unchanged baseline before labeling them pre-existing.\n\nRead operator math only when needed: [Delta Rule](../../../fla/ops/delta_rule/README.md), [Generalized Delta Rule](../../../fla/ops/generalized_delta_rule/README.md), [Simple GLA](../../../fla/ops/simple_gla/README.md), or [context parallelism](../../../fla/ops/cp/README.md).\n",".agents/skills/fla-correctness-coverage/references/cp.md":"../../../../fla/ops/cp/README.md",".agents/skills/fla-correctness-coverage/references/delta-rule.md":"../../../../fla/ops/delta_rule/README.md",".agents/skills/fla-correctness-coverage/references/generalized-delta-rule.md":"../../../../fla/ops/generalized_delta_rule/README.md",".agents/skills/fla-correctness-coverage/references/simple-gla.md":"../../../../fla/ops/simple_gla/README.md",".agents/skills/fla-design-coverage/SKILL.md":"---\nname: fla-design-coverage\ndescription: Define supported inputs, numerical budgets, routing, and validation before designing an FLA kernel or numerical change.\n---\n\n# FLA Design and Coverage\n\nUse this skill to establish the contract before implementing a kernel, backend route, or numerical change. Read [CONTRIBUTING.md](../../../CONTRIBUTING.md) for contribution rules, especially [public compatibility](../../../CONTRIBUTING.md#public-api-compatibility) and [design/RFC requirements](../../../CONTRIBUTING.md#submit-pull-requests).\n\n## Define the supported domain\n\nTrace public entries and defaults, backend verifiers, layer/model/CP callers, and existing tests. Record the merge-base SHA and the source of each supported case. Group cases only when they share the same numerical assumptions and execution route; omit combinations no supported caller can reach.\n\nDistinguish routes when kernel stages, chunk geometry or anchors, capability checks, dtype conversions, or input value ranges change. A backend here means an implementation selected by dispatch, including the default implementation, rather than a hardware product. For each group, identify its effective arguments, oracle, output/state and per-gradient tolerances, and test coverage. Assign one outcome:\n\n- The proposed implementation, with passing reference comparisons.\n- An existing fallback, with route-parity tests covering outputs, states, dtypes, and supported gradients.\n- A documented unsupported-input error raised before kernel execution and asserted by a test.\n\nPreserve previously supported cases under the public compatibility policy. A warning, partially initialized output, or fallback that changes semantics is not a valid outcome. When changing defaults, record the old and new effective values and selected routes.\n\n## Establish numerical safety\n\nFor each affected stage in forward and backward, record input-range assumptions, load/operand/accumulation/store dtypes, reduction behavior, overflow/underflow headroom, and error propagation to the end-to-end tolerance. Forward agreement alone does not establish backward safety, including inverse recomputation and state gradients.\n\nSafety flags and bounds apply to the stages and domains that justify them. Reassess them when chunk geometry, anchors, dtype staging, or input scaling changes; a safe exponent does not prove that a later inversion is safe. Document the algebra behind numerical acceptance thresholds, implement each shared predicate once, and test its acceptance boundary and the adjacent rejection or fallback. Enforce additional assumptions, such as input-norm bounds, at the public entry or every affected verifier; report the failed predicate and effective values.\n\nKeep the committed oracle, input distributions, assertions, and tolerances fixed while evaluating the candidate. A correction to that contract belongs in a separate change with before/after error distributions for the unchanged baseline and candidate; precision, tolerance, and algorithm changes follow the linked RFC policy. Keep the corrected gate green before the candidate lands. Every selectable autotune configuration must satisfy the same numerical contract.\n\n## Check routing\n\n- Share semantic argument normalization between the public entry, verifier, and executor. Test that omitted arguments match their documented explicit defaults in route, outputs, and gradients. `input_guard` handles contiguity and device context, not semantic defaults.\n- Use capability and availability helpers from `fla.utils`. Distinguish backend availability, policy enablement, and acceptance of an individual call; use [fla-dispatch-backends](../fla-dispatch-backends/SKILL.md) for implementation mechanics.\n- For context-parallel routes, every rank must make the same selection. If verifier inputs can differ by rank, combine acceptance with a collective and select the optimized route only if all ranks accept. Record whether forward commits backward to the same backend.\n- Test parity between routes claiming the same domain. Record precision settings such as `TRITON_F32_DEFAULT` and matmul flags, keep them identical across baseline, candidate, and tests, and restore global settings after tests.\n\n## Choose evidence\n\nApply the relevant axes from [fla-correctness-coverage](../fla-correctness-coverage/SKILL.md) to each supported group across these distinct layers:\n\n- **Production:** reproduce effective arguments and numerical input distributions from checked-in layer/model callers. A generic benchmark shape list alone does not establish production coverage.\n- **Public boundaries:** cover previously supported inputs, changed defaults, and verifier acceptance/rejection boundaries.\n- **Adversarial:** outside the contract, require documented rejection before kernel execution or explicit finite-output and finite-gradient checks. This does not expand the supported domain.\n\nBefore performance work, record baseline correctness, latency or throughput, commit, environment, workload, input distribution, and seed. Benchmark production cases first, then boundaries whose routing or expected performance changes. Unaffected boundaries still need correctness coverage. Use [fla-optimization-loop](../fla-optimization-loop/SKILL.md) to freeze the gate and compare candidates.\n",".agents/skills/fla-dispatch-backends/SKILL.md":"---\nname: fla-dispatch-backends\ndescription: Change FLA backend registration, dispatch, or verifiers while preserving routing and call contracts.\n---\n\n# Backend dispatch\n\nUse this skill for registration, selection, or verifier changes in `fla/backends.py` and operation backends. [CONTRIBUTING.md](../../../CONTRIBUTING.md) defines style and public compatibility; [fla/backends.py](../../../fla/backends.py) defines the routing behavior.\n\n## 1. Define the supported calls\n\nIdentify the default entry point and the calls the backend can handle: inputs, layouts, dtypes, gradient modes, and optional dependencies. Match the entry point's signature, parameter order, defaults, return structure, and supported gradients.\n\nA verifier returns `(True, None)` or `(False, reason)`. It checks support without changing inputs or global state. Keep it beside the implementation; put helpers shared by several implementations in a utility module.\n\n## 2. Register the implementation\n\nKeep the default entry point backend-independent with bare `@dispatch`. Choose the registration form that matches the implementation:\n\n| Implementation               | Registration                                                         |\n| ---------------------------- | -------------------------------------------------------------------- |\n| Function with the same API   | `@register(entry, backend=TritonAscendBackend, verifier=...)`        |\n| Adapter with backend policy  | `@register(entry)` on a `BaseBackend` subclass                       |\n\nAn adapter defines methods named after the entry point, such as `chunk_kda` and `chunk_kda_verifier`. Register it where it is defined and import optional kernels only when selected.\n\nImport the owning package's backends after its default entry points are defined. For implementations with platform-specific dependencies, guard the import with the shared backend's `is_available()`. Do not gate registration on enable switches; dispatch reads those at call time. Keep per-operation switches on `@register(..., env_var=...)` and implementation paths out of the shared dispatcher.\n\nExisting owners using `@dispatch('<operation>')` use `@register('<operation>')` on adapter classes and load implementations through the owner's backend package. Both forms use the same registration and selection logic.\n\n## 3. Preserve routing behavior\n\nFor each call, try backends in ascending priority order; ties retain registration order:\n\n1. Check availability and enablement.\n2. Run the verifier; rejection proceeds to the next candidate.\n3. Call the selected implementation. Its exceptions propagate.\n4. If no candidate accepts, call the default implementation.\n\nRegistration replaces an existing backend of the same type. Function registration keys the registry by the unwrapped entry point, so compiler decorators and public aliases share that registry. Legacy owner registries must load their backend package before lookup.\n\nKeep selection outside `torch.compile` graphs. Do not cache availability or enablement in the dispatch wrapper: registered backends must remain eligible when their switches change. Python's import cache handles modules already loaded, and unused optional dependencies must not prevent package import.\n\n## 4. Verify through the entry point\n\nRun the decorated entry point for accepted calls, verifier rejection, disabled/unavailable backends, and fallback. Check optional arguments omitted as well as supplied, changed support boundaries, outputs, and supported gradients.\n\nUse explicit backend flags in routing tests. For the default path, set `FLA_DISABLE_BACKEND_DISPATCH=1` before importing FLA. Use fresh processes for import-order checks, and verify public aliases and shared backend ownership after registration changes.\n\nFind affected consumers with `scripts/find_dependent_tests.py` and run the relevant operator, layer, and model tests. Host routing tests establish dispatch behavior; accelerator reference tests establish numerical correctness.\n",".agents/skills/fla-kda/SKILL.md":"---\nname: fla-kda\ndescription: Modify or review KDA gates, chunk and recurrent kernels, backend routes, and their correctness coverage.\n---\n\n# KDA\n\nUse this skill for `fla/ops/kda/` and its tests. Public entry points are `chunk_kda` and `fused_recurrent_kda`; inspect the affected entry point's signature before changing its implementation.\n\n## Code map\n\n- `gate.py` defines gate activations and cumulative decay.\n- `chunk_fwd.py` and `chunk_bwd.py` coordinate training; `chunk_intra.py`, `chunk_intra_token_parallel.py`, and `wy_fast.py` implement the main stages.\n- `backends/` contains alternative implementations and verifiers. Follow [backend dispatch](../fla-dispatch-backends/SKILL.md) when changing routing.\n- `tests/ops/test_kda.py` contains the reference comparisons and supported gate-mode cases.\n\n## Gate contracts\n\n`chunk_kda` takes log-space decay when `use_gate_in_kernel=False`. With it enabled, the caller supplies raw gates and optional bias:\n\n- Without a lower bound, activation is `-exp(A_log) * softplus(g + dt_bias)` and `A_log` is required.\n- With a lower bound, activation is `lower_bound * sigmoid(exp(A_log) * (g + dt_bias))`. Omitting `A_log` uses a multiplier of one.\n\n`safe_gate` selects the intra-chunk implementation independently of gate activation. Safe in-kernel activation requires `-5 <= lower_bound < 0`. Pre-gated safe mode requires bounded log-space gates supplied by the caller; tests cover `[-5, 0]`, but the wrapper does not validate tensor values.\n\n## Numerical invariants\n\nThe safe intra path uses 16-token diagonal blocks and midpoint offsets before exponentiation. At a lower bound of `-5`, a block can accumulate `-80` in natural-log units; local offsets limit each exponent's span to about half that range. Preserve these offsets rather than exponentiating the full cumulative gate. Inter-block decay uses paired differences whose exponents remain non-positive for monotonic decay.\n\nThe non-safe path uses token-parallel diagonal computation. Both paths share inter-block and triangular-solve work, so changes there require coverage of both modes. Keep precision and tolerance changes within the design process in [AGENTS.md](../../../AGENTS.md#scope-and-direction).\n\n## Validation\n\nUse [correctness coverage](../fla-correctness-coverage/SKILL.md) for the test strategy. Select the KDA cases affected by the change:\n\n- Dense and variable-length sequences; forward, final state, and supported gradients.\n- Pre-gated and in-kernel gates, safe and non-safe paths, and omitted optional arguments.\n- Beta activation and query/key normalization modes when touched.\n- Grouped value heads and differing key/value dimensions when indexing or shapes change.\n- Initial states, intermediate states, and context parallelism when their paths are touched.\n- Accepted and rejected backend routes when changing verifiers.\n\nFor gate or decay changes, include the lower-bound limits, saturated gate inputs, extreme gate scales, long cumulative decay, and chunk or ragged-sequence boundaries. Test through the public entry point so the selected route and effective arguments are exercised.\n",".agents/skills/fla-nvidia-performance/SKILL.md":"---\nname: fla-nvidia-performance\ndescription: Profile and optimize FLA kernels on NVIDIA GPUs, with same-hardware benchmarks and targeted Nsight Compute analysis.\n---\n\n# NVIDIA kernel performance\n\nUse this skill for NVIDIA kernel optimization and backend tuning. Follow [CONTRIBUTING](../../../CONTRIBUTING.md#benchmarking) for correctness and performance evidence, and [fla-optimization-loop](../fla-optimization-loop/SKILL.md) for sustained iteration.\n\n## Hardware baseline\n\nUse datacenter NVIDIA GPUs with sm_90 or newer for PR performance conclusions; H100 and H20 are accepted, and sm_100/sm_103 are preferred. Results from A100, older GPUs, or consumer cards are supplementary. Record the exact GPU and software versions with each comparison.\n\n## Workflow\n\n1. Identify the changed kernel, its callers, and the production workloads it affects. Record a baseline commit and confirm which backend runs.\n2. Pass the relevant correctness tests before timing a candidate. Preserve the reference and tolerances; include outputs, states, and supported gradients.\n3. Compare baseline and candidate on the same GPU with identical workloads and dispatch settings. Include dense and supported variable-length cases. Use an end-to-end benchmark for a training or generation throughput claim.\n4. Use Nsight Compute when a bottleneck or regression needs hardware metrics. Match the optimization to the measured limit, then rerun correctness and timing.\n5. Report latency or throughput, relevant memory changes, and any regressions. Keep raw profiler artifacts outside git, for example under `profile/<run_name>/`.\n\n## Benchmark commands\n\nRun from the repository root. Resolve the baseline to a commit SHA before comparing:\n\n```bash\nFLA_BENCH_BASE=$(git rev-parse origin/main)\nFLA_CI_ENV=0 python -m benchmarks.ops.verify --op chunk_kda --base \"$FLA_BENCH_BASE\"\n\npython benchmarks/benchmark_training_throughput.py --name kda --batch_size 2 --seq_len 8192\npython benchmarks/benchmark_training_throughput.py --name kda --batch_size 2 --seq_len 8192 --varlen\n```\n\nThe verification command runs the current tests; retain the validated baseline and tests yourself. A selected subset is useful during iteration but does not replace the full relevant gate before reporting a gain.\n\n## Nsight Compute\n\nCollect a representative changed kernel with the installed `ncu`. Check which sections and metrics that version supports.\n\nThe benchmark runner launches a child process, so use `--target-processes all`. Replace the kernel filter with the changed kernel's name. `--no-base` profiles the current implementation; collect the baseline separately with the same command and workload.\n\n```bash\nmkdir -p profile/kda\nncu --target-processes all --set full \\\n  -k 'regex:chunk_kda' -c 1 -o profile/kda/full \\\n  python -m benchmarks.ops.run --op chunk_kda --modes fwd --no-base\n\nncu --target-processes all --set source --section SourceCounters \\\n  -k 'regex:chunk_kda' -c 1 -o profile/kda/source \\\n  python -m benchmarks.ops.run --op chunk_kda --modes fwd --no-base\n```\n\nUse duration, memory throughput, occupancy, stalls, and hot instructions to explain the bottleneck. Profiling overhead makes these runs unsuitable as the latency comparison; use the synchronized benchmark for that measurement.\n",".agents/skills/fla-optimization-loop/SKILL.md":"---\nname: fla-optimization-loop\ndescription: Iterate on FLA kernel performance with a frozen correctness gate, a measured baseline, and reproducible candidate comparisons.\n---\n\n# FLA Optimization Loop\n\nUse this skill for repeated kernel performance work. Define supported workloads and numerical assumptions with [fla-design-coverage](../fla-design-coverage/SKILL.md); use the [NVIDIA](../fla-nvidia-performance/SKILL.md) or [Ascend](../fla-ascend-performance/SKILL.md) skill for profiling on that platform.\n\n## Freeze the gate and baseline\n\nAdd missing regression or new-path coverage before starting the loop. Then keep the relevant tests, reference implementation, input distributions, assertions, tolerances, and public signature fixed. Do not narrow parameter matrices, add skips, relax precision or thresholds, special-case test values, cache results across trials, or change timing behavior to make a candidate pass. Fully initialize exposed outputs. Use kernel implementations for the optimization; replacing the whole operator with a vendor call solely to win the benchmark does not meet this workflow.\n\nA suspected oracle or tolerance defect needs a separately justified correction, not an adjustment that makes the candidate pass. Follow [CONTRIBUTING.md](../../../CONTRIBUTING.md#submit-pull-requests) for precision, numerical-algorithm, and tolerance changes, and its [compatibility policy](../../../CONTRIBUTING.md#public-api-compatibility) for public behavior.\n\nRecord the entry point, target workloads, allowed implementation languages, validation/benchmark commands, baseline SHA, and success criterion in the local optimization log. Include production inputs from callers as well as applicable benchmark-registry shapes. Confirm the baseline passes the full gate and measure it before changing kernels. For a new kernel, establish correctness before tuning.\n\n```bash\nFLA_BENCH_BASE=$(git rev-parse origin/main)\nFLA_CI_ENV=0 python -m benchmarks.ops.verify --op chunk_gla --base \"$FLA_BENCH_BASE\"\n```\n\nUse a recorded SHA: the baseline worktree cannot reuse a checked-out branch, and slash-containing refs break its temporary directory naming. `--base` benchmarks that revision; it does not run its correctness tests. Establish baseline correctness separately if the current tree already differs.\n\nThe driver runs the current pytest file and stops before timing when pytest fails. It neither freezes files nor rejects an all-skipped run. Check that the expected cases and benchmark rows ran. `--test-file` overrides a mismatched derived test path; `--gate-k` selects a quick subset, while `--no-gate` produces unverified measurements. Neither subset nor skipped-gate results qualify for promotion.\n\n## Iterate from evidence\n\n1. Profile the current bottleneck and choose one reasoned change.\n2. Run the frozen gate, then benchmark only a passing candidate. A subset can provide an early signal; the final gate covers the full affected suite.\n3. Record the change, gate result, measurement, and keep/drop decision before the next experiment. Preserve a recoverable candidate under the repository's git rules and the user's authorization; failed and unchanged attempts need records, not empty commits.\n\nRead [measurement and correctness traps](references/TRAPS.md) before interpreting results. Keep seeds, numeric flags, hardware, and workload fixed. Candidate runtime helps rank local iterations; measure the final candidate directly against the frozen baseline in one session rather than adding speedups from separate runs.\n\nAdd specialized paths only when profiling shows different bottlenecks across shape regimes and measurements justify the extra routing. Record each condition, chosen path, and before/after result, and validate the full shape set including fallback boundaries.\n\nAfter repeated attempts without an improvement above noise, re-profile and review the failed directions before continuing. Respect the user's stopping condition. A no-win conclusion needs a measured baseline, a reasoned candidate attempt, gate status, benchmark evidence, and an identified bound or blocker. An unavailable profiler or one losing candidate alone does not establish a performance floor; report unresolved validation as a limitation.\n\n## Keep and report the result\n\nKeep iteration notes and raw traces under ignored `profile/<op>-opt/`; [opt-log-template.md](references/opt-log-template.md) provides a compact record format. This search history stays out of kernel comments and the final PR description.\n\nPromote only with a full green gate, a repeatable win on the claimed workloads, and profiling or a quantified bottleneck analysis that explains it. Report per-shape results and regressions, commands, baseline/candidate revisions, environment, and measurement variability using the relevant hardware skill. An unexplained large speedup needs investigation. Use [fla-pr-readiness](../fla-pr-readiness/SKILL.md) to prepare the final diff and evidence.\n",".agents/skills/fla-optimization-loop/references/TRAPS.md":"# Correctness and Measurement Traps\n\n## Partial writes\n\n`tests/conftest.py` poisons eligible FLA `empty` allocations in operator/module tests with NaNs. It excludes some allocations and does not cover layer/model/CP tests. A poisoned-run failure that disappears in an ad hoc script can reveal uninitialized reads; inspect masks and ensure every exposed output element is written. NaNs can also have numerical causes, so diagnose the failure rather than assuming its source.\n\n## Reference precision and tolerance\n\n`fla.utils.assert_close` checks relative RMS error with an absolute-error shortcut, not elementwise `rtol`. `FLA_CI_ENV=1` and explicit `warning=True` can downgrade tolerance failures to warnings. Use strict local validation and inspect warnings; keep the committed thresholds fixed.\n\nTF32 in a reference matmul can change an fp32 comparison. Inspect the reference's precision settings before blaming the kernel, but do not change the frozen oracle to pass. Baseline and candidate need identical precision and numeric flags. An intentional precision change follows the numerical-change process in [CONTRIBUTING.md](../../../../CONTRIBUTING.md#submit-pull-requests).\n\n## Address overflow\n\nCast program IDs and offsets to `tl.int64` before multiplication; widening an already-overflowed product cannot repair it. Small shapes may hide the error. Run large-offset and affected varlen/grid cases under the [Triton addressing rules](../../../../CONTRIBUTING.md#triton-kernels).\n\n## Missing work in a fast result\n\nCheck unexpectedly large speedups against the actual selected path, executed tests, and measured shapes. `verify.py` accepts pytest's exit status, including an all-skipped success. The benchmark runner filters constrained shapes and omits failed warmups or measurements. Compare expected and returned rows before making a claim; a passing process alone is insufficient.\n\n## Autotune and clocks\n\n`run.py`, also used by `verify.py`, warms all shapes before timing. Preserve that warmup. After changing the configuration space, check cached selections and rewarm or clear the relevant cache if results look stale.\n\nUnlocked clocks and changing device load can shift latency between runs. Rank local candidates by their own runtime, then remeasure the final candidate against the recorded baseline in one session. Report variability and the device state; speedups from separate experiments do not add up to a final result.\n",".agents/skills/fla-optimization-loop/references/opt-log-template.md":"# Optimization Record\n\nKeep one local log under ignored `profile/<op>-opt/`. Record the entry point, target workloads, baseline SHA and correctness, environment, commands, numeric settings, and success criterion once. Preserve enough candidate identity to reproduce each measurement.\n\n| Attempt  | Change or hypothesis | Gate | Latency by workload | Decision and reason |\n| -------- | -------------------- | ---- | ------------------- | ------------------- |\n| Baseline | Unchanged revision   | Pass | Measured baseline   | Comparison point    |\n\nAppend an entry after each attempt, including failures and unchanged results. Identify whether the gate was full or selected; record no performance result for a failed gate. Retained candidates need their revision or recoverable diff, the reason they helped, and any unresolved limitation. Keep failed directions here so later work does not repeat them.\n\nWhen specializing, add each routing condition, entry point, baseline/candidate latency, and measured reason for a separate path. Compare every bucket against the same baseline revision on that bucket's workloads, including fallback boundaries. Keep these details alongside the iteration results.\n",".agents/skills/fla-pr-readiness/SKILL.md":"---\nname: fla-pr-readiness\ndescription: Prepare or update an FLA pull request with focused scope, verified evidence, and the current repository template.\n---\n\n# FLA PR Readiness\n\nRead [CONTRIBUTING.md](../../../CONTRIBUTING.md) and [AGENTS.md](../../../AGENTS.md) first. Use the current [PR template](../../../.github/pull_request_template.md) for headings and checkboxes; contribution policy and template wording are not duplicated here.\n\n## Review scope and evidence\n\nRead the complete diff against the PR base and audit affected shared callers. State its concrete outcome in one sentence and separate independent work. Check existing issues/PRs and the applicable design discussion. Follow [scope and size](../../../CONTRIBUTING.md#pr-scope-and-size) for a cohesive change above the normal limit, and [public compatibility](../../../CONTRIBUTING.md#public-api-compatibility) when interfaces or defaults are affected.\n\nWhen a change needs splitting or dependent PRs need a native GitHub stack, use [PR stacks](../fla-pr-stacks/SKILL.md) to choose boundaries, reuse existing PRs, and verify the published stack.\n\nFor code changes, discover dependent tests and run the relevant targets. For example:\n\n```bash\npython scripts/find_dependent_tests.py fla/ops/kda/chunk.py\nFLA_CI_ENV=0 pytest tests/ops/test_kda.py\npre-commit run --files fla/ops/kda/chunk.py tests/ops/test_kda.py\n```\n\nFor documentation-only changes, validate links, examples, and formatting as described in [Evidence by Change Type](../../../CONTRIBUTING.md#evidence-by-change-type).\n\nThe dependency helper only lists tests. Check that regression coverage exercises the changed path and that required numerical comparisons passed, including supported gradients. Review new test functions and files against [test organization and naming](../../../CONTRIBUTING.md#test-organization-and-naming); use an existing test case where it fits. Record commands, tested revision, environment, and results; distinguish failures, skips, and unrun checks. Reproduce a failure on the baseline before calling it pre-existing.\n\nFor performance-relevant changes, collect same-hardware before/after measurements under [Benchmarking](../../../CONTRIBUTING.md#benchmarking) and the relevant hardware skill. A neutral result needs numbers; explain inapplicable checks. Inspect the final diff for stale comments and unrelated changes, leaving scratch logs and raw profiles outside it.\n\n## Prepare and publish\n\nWrite for a reviewer who has not seen the conversation. Follow [PR Description](../../../CONTRIBUTING.md#pr-description) and [Evidence by Change Type](../../../CONTRIBUTING.md#evidence-by-change-type): explain the problem, resulting behavior, affected users, and evidence. Include only useful rationale or limitations; omit file inventories and the history of attempts. Keep incomplete work in draft and leave claims and checklist items consistent with actual validation.\n\nThe `check-pr-title` workflow also checks the body and size exception on edits. Preserve template checklist wording and distinguish required items from conditional ones.\n\n`gh pr edit` fails with this repository's classic-Projects integration. Write the complete body to a temporary file and update through REST:\n\n```bash\ngh api -X PATCH \"repos/fla-org/flash-linear-attention/pulls/$pr_number\" -F \"body=@$pr_body_file\"\n```\n\nConfirm the published description, head revision, and checks. After review changes, rerun affected validation and update evidence that no longer describes the current code. Publishing and git operations remain subject to the user's authorization and repository rules.\n",".agents/skills/fla-pr-stacks/SKILL.md":"---\nname: fla-pr-stacks\ndescription: Split oversized FLA pull requests into focused changes, reuse existing PRs, and create or maintain verified GitHub native PR stacks.\n---\n\n# FLA PR Stacks\n\nUse for decomposing large changes or managing dependent PRs. Follow [CONTRIBUTING.md](../../../CONTRIBUTING.md#pr-scope-and-size) for scope and size, [AGENTS.md](../../../AGENTS.md) for git permissions, and [PR readiness](../fla-pr-readiness/SKILL.md) for each PR's evidence and description.\n\n## 1. Choose the split\n\nInspect the complete diff, commit history, and related open PRs, including their review discussions. Map each proposed layer to one outcome, its prerequisites, validation, and an existing PR to reuse where appropriate.\n\n| Relationship                            | Action                                                                   |\n| --------------------------------------- | ------------------------------------------------------------------------ |\n| Overlaps an existing PR's purpose       | Update that PR when ownership permits; preserve attribution and reviews. |\n| Independent of the other changes        | Submit separately against the trunk.                                     |\n| Depends on another layer                | Stack above that layer; include its own tests and documentation.         |\n| Cohesive but still above the size limit | Use the contribution guide's exception with a concrete justification.    |\n\nEach layer presents an increment over its parent and must work with its prerequisites, without needing later layers to fix it. Separate review does not imply dependency-free code. Do not split by arbitrary line counts, leave callers broken, or defer tests to a later PR.\n\nCheck head repository identity before arranging existing PRs. [GitHub stacks require all branches in the same repository](https://docs.github.com/en/pull-requests/reference/stacked-pull-requests); cross-fork stacks are unsupported. Keep incompatible PRs as explicit external prerequisites instead of duplicating them under new PR numbers.\n\n## 2. Build and validate\n\nPreserve the original head and uncommitted work before reconstructing commits in isolated branches or worktrees. Reuse suitable commits and preserve authorship. Check current remote heads before updating existing PRs.\n\nReview each layer against its parent, then compare the top branch with the preserved original result:\n\n```bash\ngit diff main...refactor/foundation\ngit diff refactor/foundation...refactor/consumer\ngit diff original-head refactor/consumer\n```\n\nSubstitute the recorded trunk and actual refs. Account for trunk changes and intentional omissions, including fixes retained in existing PRs. Run relevant tests at each layer's head and on the combined result; passing only the top branch does not validate intermediate PRs.\n\n## 3. Register and verify the native stack\n\nChanging PR bases alone does not create a native stack. Use the official [`gh stack` extension](https://docs.github.com/en/pull-requests/how-tos/create-pull-requests/creating-stacked-pull-requests), GitHub UI, or REST Stacks API. Check installed command help and existing publishing authorization first.\n\nFor prepared branches, adopt them bottom to top and submit:\n\n```bash\ngh stack init --base main refactor/foundation refactor/consumer\ngh stack submit --auto --remote origin\n```\n\nChoose the actual remote containing the branches. `submit` pushes, creates or updates PRs, and registers the stack. `--auto` creates new PRs as drafts; replace generated titles and bodies using PR readiness.\n\nFor existing PRs, use `gh stack link --base main --remote origin 123 124` with their real numbers in dependency order. This reuses PRs without requiring local stack tracking. Branch arguments also push and create missing PRs. After an error, inspect what succeeded before retrying.\n\nFor every PR, verify [remote stack membership](https://docs.github.com/en/pull-requests/reference/stacked-pull-requests-apis-and-webhooks):\n\n```bash\ngh api \"repos/fla-org/flash-linear-attention/pulls/$pr_number\" \\\n  --jq '{number, head_repo: .head.repo.full_name, head: .head.ref, base: .base.ref, stack}'\n```\n\nRequire non-null `stack` objects with the same stack identity, correct order, size, and trunk. Check each PR's incremental diff and report the verified links and remaining external prerequisites. A local branch chain or successful push is insufficient.\n\n## 4. Maintain the series\n\nAfter changing a lower layer, review affected upper diffs and rerun relevant tests. `gh stack rebase` rewrites history; `gh stack sync` also force-pushes with leases. Apply the repository's rewrite permissions and preserve concurrent remote work, then recheck published membership and evidence.\n\nCreating a series does not authorize merging or closing PRs. If merging is requested, establish the intended prefix: selecting an upper PR also merges those below it. Use a stack-aware merge operation; the legacy synchronous merge API does not support stacks.\n",".agents/skills/fla-triton-to-gluon/SKILL.md":"---\nname: fla-triton-to-gluon\ndescription: Port an FLA Triton kernel to Gluon when explicit layouts, asynchronous transfers, or scheduling can address a measured bottleneck.\n---\n\n# Triton to Gluon\n\nUse Gluon when profiling points to register pressure, layout conversion, or load/compute overlap that needs explicit control. A kernel already saturating its useful bandwidth or compute limit may not benefit. Keep the portable default path when adding hardware-specific implementations.\n\nFollow [fla-optimization-loop](../fla-optimization-loop/SKILL.md) for the baseline and correctness gate, and [fla-nvidia-performance](../fla-nvidia-performance/SKILL.md) for NVIDIA measurements.\n\n## Check the target environment\n\nGluon is experimental. Check the installed Triton API before copying [upstream examples](https://triton-lang.org/main/getting-started/tutorials/gluon/); their names and signatures can differ from the installed release.\n\nGate architecture-specific code appropriately: Ampere supports `cp.async`, Hopper adds TMA/WGMMA, and Blackwell adds TMEM/tcgen05. Code under Gluon's `nvidia` namespace requires NVIDIA hardware. Keep compile-time choices in explicit constexpr arguments so the kernel, launcher, and autotune pruning use the same values.\n\n```python\nfrom triton.experimental import gluon\nfrom triton.experimental.gluon import language as gl\n```\n\n## Port incrementally\n\n1. Record the baseline commit and retain the validated tests, reference, and tolerances. Keep the Triton fallback available.\n2. Translate loads, stores, and arithmetic with explicit layouts first. Match each tensor's contiguous dimension and preserve masking. Pass forward, state, and supported backward comparisons before adding asynchrony.\n3. Add async transfers only where load latency or addressing pressure limits the kernel. Define buffer ownership, completion, and reuse for the prologue, steady state, and drain.\n4. Add architecture-specific MMA or scheduling when profiling supports it. Recheck register/shared-memory use and retune after each change.\n5. Run the full relevant correctness gate and same-hardware benchmarks, including supported variable-length and boundary cases. Measure production entry points as well as individual kernels.\n\nA literal translation establishes parity; measure it before assuming a performance gain. Preserve the validated numerical algorithm and precision while changing execution layout or scheduling.\n\n## Layout and resource constraints\n\n- Pointer tensors and values need compatible layouts. A cross-warp `gl.convert_layout` may use shared memory; use `assert_trivial=True` when a conversion is intended to be free.\n- Broadcasts and oversized layout tiles can duplicate values and increase register use. Choose layouts from actual tensor strides rather than copying a tutorial's block sizes.\n- Account for all live shared-memory buffers and the target device's allocation limit. Prune infeasible autotune configurations; provide a streaming path when a resident design cannot fit supported shapes.\n- Large static unrolls multiplied by many autotune configurations can dominate compilation. Prune from the actual shape and live-memory budget, and reuse the worker and Triton cache during iteration.\n- Persistent schedules can trade launch overhead for worse L2 locality. Warp specialization also changes the total register budget; measure occupancy and stalls after changing either.\n\n## Async correctness\n\n### Transfers and buffer reuse\n\n`cp.async` groups are ordered as a queue. `wait_group(N)` bounds all outstanding groups, so it cannot identify a particular buffer after unrelated prefetches are interleaved. For per-buffer scheduling, use separate mbarriers and match their arrival counts to the participating threads. In the installed API, check whether `mbarrier_arrive` increments the count; a preinitialized thread count requires the non-incrementing form.\n\nWait for a load before consuming its buffer, and protect reuse until all readers finish. Even same-lane shared-memory staging needs protection against overwriting data still being read. Track each mbarrier's phase as its buffer is reused; do not run more than one phase ahead or share a completion barrier between TMA and tcgen05 without reinitializing it.\n\nTMA descriptors must satisfy the target's alignment and stride requirements. TMA stores may distinguish completion of the shared-memory read from completion of the global write. Check the installed wait API and require global completion before another operation reads the stored range.\n\n### MMA and memory ordering\n\nHopper WGMMA requires its B operand in shared memory and uses register accumulators; consume the values returned by its wait operation so compiler dependencies remain explicit. Blackwell tcgen05 uses TMEM accumulators and mbarrier completion; respect the participating warpgroup and TMEM layout requirements. Initialize accumulators explicitly, including `use_acc=False` where supported.\n\nGeneric shared-memory accesses and async operations use different memory proxies. Apply the required proxy fence when handing a buffer to an async consumer; an mbarrier alone does not replace that fence. A completed TMA load establishes the ordering needed to read its destination. Check ordering across warp-specialized partitions as well as within a single producer/consumer loop.\n\n### Tails and exact-zero cases\n\nMasked async copies can leave shared-memory elements uninitialized. Prevent invalid rows from entering reductions: zero-fill where supported, or use valid-row loads and explicitly remove every invalid contribution. Multiplying by zero does not neutralize a NaN. Keep output stores masked.\n\nDo not rely on separately compiled reductions cancelling bitwise. For a mathematically exact-zero special case, such as a single-source softmax gradient, preserve the exact-zero result explicitly and test it against the reference.\n\n## Verification commands\n\nRun from the repository root on the target GPU:\n\n```bash\nFLA_BENCH_BASE=$(git rev-parse origin/main)\nFLA_CI_ENV=0 python -m benchmarks.ops.verify --op chunk_kda --base \"$FLA_BENCH_BASE\"\n```\n\nReplace the operation with the one being ported. The command executes the current tests; it does not freeze them. Use `--gate-k` only for quick iteration, then run the full gate. Set backend environment flags before launching a fresh process so import-time choices cannot contaminate the comparison.\n","AGENTS.md":"# FLA Guidelines\n\nGuidance for AI coding agents (Claude Code, Codex, etc.) working in this repo.\n\n**Read `CONTRIBUTING.md` first.** It is the authoritative source for all code style, docstring, comment, commit, PR, and testing conventions, and applies to humans and agents alike. This file only covers agent-specific operational behavior that doesn't belong in a human contributor doc.\n\n## Git safety\n\n- **Never discard uncommitted work** with `git checkout HEAD -- <file>`, `git reset --hard`, or `git restore` to \"get a clean base\". Unstaged changes are unrecoverable (no blob, no reflog). Edit in place or `git stash` instead, and confirm with the user when in doubt.\n- **On `main`**: never commit or push without explicit per-action approval. Suggest a feature branch first.\n- Don't rewrite or amend already-pushed commits unless the user asks.\n- **Branch names and attribution**: name local and remote branches after the change. Do not include tool or assistant identities in branch names, commit messages, PR titles/bodies, or review comments. Do not add `Co-Authored-By` trailers or generated-by notices.\n\n## Scope and direction\n\n- **Check for duplicates first**: search open issues/PRs before starting so you don't redo in-flight work — `gh pr list --repo fla-org/flash-linear-attention --state open --search \"<keywords>\"`.\n- **Align before coding when a design decision is involved**: a new operator or model, a new public API, a kernel rewrite with a different algorithm, or a repo-wide policy change (e.g. adding a dtype everywhere). Open an issue or a draft PR first; if a matching issue/PR already exists, comment there instead of opening a duplicate. Judge design impact separately from the PR size check: tests, config boilerplate, and mechanical repetition of an existing pattern alone do not make a change a design decision.\n- **RFC first for precision or algorithm changes**: changing the precision a validated kernel computes or accumulates in (e.g. fp32 → bf16, bf16 → fp16), relaxing a tolerance, or replacing a numerical algorithm (e.g. a different inversion or solve scheme) needs an RFC issue before a PR, even if it would stay within test tolerance. FLA's kernels are widely validated and downstream users rely on their numerics, so these are design decisions to agree on, not fixes to ship. Ordinary reorderings within the same precision — tiling, accumulation order, scheduling — do not need an RFC; they are normal optimization PRs with before/after benchmarks as usual.\n- **Stop and ask before breaking changes**: renaming or removing public symbols or arguments in `fla/layers/` / `fla/models/` or documented config fields; changing checkpoint/state-dict compatibility; changing observable behavior, including numerics beyond the existing test tolerance. Bug fixes that restore intended behavior are not breaking. If the user doesn't answer — or you are running non-interactively — do the non-breaking parts, leave the breaking part out, and describe it in your final report. Silence is not consent.\n\n## Comment discipline\n\nThe style rules for comments and docstrings live in `CONTRIBUTING.md` (\"Docstrings and Comments\"), including the banned anti-patterns. What follows is agent-specific behavior:\n\n- **You own the comments your change invalidates.** When a change makes a comment or docstring in the code you touched factually wrong, fix it or delete it in the same commit — \"the code you touched\" means the same function, kernel, or class, not a repo-wide sweep for stale mentions. If you can't tell whether a comment is still true, keep it and flag it in the PR description; never delete a \"why\" comment just because you can't verify it.\n- **Comments are not your scratchpad.** State design conclusions, not the journey: \"not X, it does not compile\" is a welcome why comment; \"tried X first, then switched to Y\" is chain-of-thought noise. This covers docstrings too.\n- **Fix the words, not the style.** Correcting stale content is not a license to reformat the surrounding docstring or comment style.\n\n## PR workflow\n\n- **No busywork PRs**: don't open a one-off PR for a single typo or isolated style tweak; bundle trivial cleanups into substantive work.\n- **Minimal diff**: every changed line serves the PR's stated purpose. Drop drive-by reformats, unrelated cleanups, and speculative generality before pushing, not after a reviewer asks (see CONTRIBUTING \"Core Principles\" 5).\n- **Blast radius**: for every shared symbol you touched, all callsites still hold (Core Principle 4 — audit them in one pass, don't assume). For any behavior or numerics change, the PR description must state who is affected — which shapes, dtypes, backends, checkpoints, or downstream users — or state explicitly why no one is. \"Tests pass\" is not an impact analysis.\n- Before opening a PR, re-read your own diff once with fresh eyes: check every comment it adds against the anti-pattern list, and every comment it invalidates against [Comment discipline](#comment-discipline).\n- Check the diff for `tl.make_block_ptr` / `tl.advance`; these are banned from mainline Triton code. Follow the Ascend backend exemption in CONTRIBUTING \"Triton Kernels\".\n- `gh pr edit` fails on this repo (classic-Projects GraphQL error). Edit a PR title/body via the REST API instead: `gh api -X PATCH repos/fla-org/flash-linear-attention/pulls/<N> -f title='...' -F body=@file`.\n- Keep review/PR comments concise and natural — skip heavy `**1.** **2.**` scaffolding, write like a person.\n\n## Repo-local Skills\n\nUse the [skill index](.agents/skills/README.md) to select task-specific workflows and their execution order. Skills supplement `CONTRIBUTING.md` and this file; they do not define competing contribution policies.\n\nLoad a relevant skill when the task matches its scope. In particular, load [fla-pr-readiness](.agents/skills/fla-pr-readiness/SKILL.md) before opening any PR.\n","evals/README.md":"# Evaluation\n\nWe provide evaluation scripts for `fla` models with [lm-evaluation-harness](https://github.com/EleutherAI/lm-evaluation-harness), as well as a standalone perplexity evaluator.\nInstall `fla` following the [installation guide](../INSTALL.md), and run the commands below from the repository root.\n\n* [lm-evaluation-harness](#lm-evaluation-harness)\n  * [Single-GPU Evaluation](#single-gpu-evaluation)\n  * [Multi-GPU Evaluation](#multi-gpu-evaluation)\n  * [RULER Benchmarks](#ruler-benchmarks)\n* [Perplexity](#perplexity)\n  * [Quick Start](#quick-start)\n  * [Evaluation Settings](#evaluation-settings)\n  * [Scoring](#scoring)\n  * [Results](#results)\n\n## lm-evaluation-harness\n\nThe [lm-evaluation-harness](https://github.com/EleutherAI/lm-evaluation-harness) library allows you to easily perform (zero-shot) model evaluations.\nInstall `lm_eval` following [their instructions](https://github.com/EleutherAI/lm-evaluation-harness/blob/main/README.md).\n\n### Single-GPU Evaluation\n\nRun evaluation with:\n\n```sh\nMODEL='fla-hub/gla-1.3B-100B'\npython -m evals.harness --model hf \\\n    --model_args pretrained=$MODEL,dtype=bfloat16 \\\n    --tasks wikitext,lambada_openai,piqa,hellaswag,winogrande,arc_easy,arc_challenge,boolq,sciq,copa,openbookqa \\\n    --batch_size 64 \\\n    --num_fewshot 0 \\\n    --device cuda \\\n    --show_config\n```\n\nWe've made `fla` compatible with hf-style evaluations, you can call [evals.harness](harness.py) to finish the evaluations.\nThe command above evaluates the GLA checkpoint on the tasks used in the GLA paper.\n\n### Multi-GPU Evaluation\n\nUse the Hugging Face Accelerate launcher for data-parallel evaluation, where each GPU loads a full copy of the model:\n\n```sh\nMODEL='fla-hub/gla-1.3B-100B'\naccelerate launch -m evals.harness --model hf \\\n    --model_args pretrained=$MODEL,dtype=bfloat16,trust_remote_code=True \\\n    --tasks wikitext,lambada_openai,piqa,hellaswag,winogrande,arc_easy,arc_challenge,boolq,sciq,copa,openbookqa \\\n    --batch_size 64 \\\n    --num_fewshot 0 \\\n    --device cuda \\\n    --show_config \\\n    --trust_remote_code\n```\n\n### RULER Benchmarks\n\nThe RULER benchmarks are commonly used for evaluating model performance on long-context tasks.\nYou can evaluate `fla` models on RULER directly using `lm-evaluation-harness`. RULER is only available in a relatively recent version of `lm-evaluation-harness`, so make sure you have the latest version installed.\n\n```sh\ngit clone --depth 1 https://github.com/EleutherAI/lm-evaluation-harness\npip install -e './lm-evaluation-harness[ruler]'\n```\n\nFor example, to evaluate contexts up to 32K tokens:\n\n```sh\nMODEL='fla-hub/gla-1.3B-100B'\nOUTPUT='results/ruler'\naccelerate launch -m evals.harness --model hf \\\n    --output_path $OUTPUT \\\n    --tasks niah_single_1,niah_single_2,niah_single_3,niah_multikey_1,niah_multikey_2,niah_multikey_3,niah_multiquery,niah_multivalue,ruler_vt,ruler_cwe,ruler_fwe,ruler_qa_hotpot,ruler_qa_squad \\\n    --model_args pretrained=$MODEL,dtype=bfloat16,max_length=32768,trust_remote_code=True \\\n    --metadata='{\"max_seq_lengths\":[4096,8192,16384,32768]}' \\\n    --batch_size 2 \\\n    --show_config \\\n    --trust_remote_code\n```\n\nIf a GPU can't load a full copy of the model, please refer to [this link](https://github.com/EleutherAI/lm-evaluation-harness?tab=readme-ov-file#multi-gpu-evaluation-with-hugging-face-accelerate) for FSDP settings.\n\n> [!TIP]\n> If you are using `lm-evaluation-harness` as an external library and can't find (almost) any tasks available, before calling `lm_eval.evaluate()` or `lm_eval.simple_evaluate()`, simply run the following to load the library's stock tasks:\n> ```py\n> from lm_eval.tasks import TaskManager\n> TaskManager().initialize_tasks()\n> ```\n\n## Perplexity\n\n[evals.ppl](ppl.py) evaluates next-token perplexity over fixed-length contexts and reports perplexity by position within each context.\n\n### Quick Start\n\nInstall the dataset dependency:\n\n```sh\npip install datasets\n```\n\nFor example, to evaluate the GLA checkpoint on the PG19 test split and save the results:\n\n```sh\npython -m evals.ppl \\\n    --path fla-hub/gla-1.3B-100B \\\n    --data fla-hub/pg19 --split test \\\n    --block_size 28672 --bucket_size 2048 --batch_size 1 \\\n    --dtype bfloat16 --num_proc 4 \\\n    --output results/ppl.json\n```\n\n### Evaluation Settings\n\n| Argument          | Default                 | Description                                                 |\n|-------------------|-------------------------|-------------------------------------------------------------|\n| `--path`          | `fla-hub/gla-1.3B-100B` | Hugging Face model ID or local checkpoint directory.        |\n| `--data`          | `fla-hub/pg19`          | Dataset to evaluate.                                        |\n| `--split`         | `train`                 | Dataset split; set this explicitly for held-out evaluation. |\n| `--column_name`   | `text`                  | Dataset column containing text.                             |\n| `--block_size`    | `28672`                 | Context length in tokens; must be at least 2.               |\n| `--bucket_size`   | `2048`                  | Number of prediction positions per reported bucket.         |\n| `--batch_size`    | `1`                     | Number of full contexts evaluated together.                 |\n| `--device`        | Auto-detected           | Device on which to load and evaluate the model.             |\n| `--dtype`         | `bfloat16`              | Model dtype: `float32`, `float16`, or `bfloat16`.           |\n| `--num_proc`      | `32`                    | Number of dataset tokenization processes.                   |\n| `--revision`      | `main`                  | Model and tokenizer revision.                               |\n| `--data_revision` | `main`                  | Dataset revision.                                           |\n| `--output`        | None                    | Optional path for JSON results and evaluation settings.     |\n\nUse immutable commit hashes for `--revision` and `--data_revision` to reproduce a run.\nFor long contexts, reduce `--batch_size` to lower model memory use and `--bucket_size` to reduce the FP32 softmax workspace.\nThe model still materializes its full logits.\n\n### Scoring\n\nDocuments are tokenized with the tokenizer's default special-token policy, then concatenated in dataset order without extra separators.\nThe token stream is divided into independent contexts of `block_size` tokens, and the final partial context is evaluated separately without padding.\nContext is reset at every block boundary; there is no sliding window or recurrent state carried between blocks.\n\nEach context of length `L` contributes `L-1` next-token predictions.\nThe first token is unscored, and no synthetic EOS target is appended after the last token.\nEOS tokens present in the input are scored normally.\nA one-token tail contributes no predictions; a dataset with fewer than two tokens raises `ValueError`.\n\nOverall PPL and positional PPL use the same FP32 negative log-likelihoods, weighted by the number of scored predictions.\nBuckets group predictions by their zero-based logit position within each context: bucket 0 covers positions `0` through `bucket_size-1`.\nModels configured with fused linear cross entropy are supported, and evaluation runs with caching disabled.\n\n### Results\n\nThe script prints overall PPL, positional PPL, and token counts.\nWith `--output`, it also saves a JSON file containing:\n\n* `perplexity` and `block_perplexities`: overall and positional PPL; buckets with no predictions have `null` PPL.\n* `total_nll`, `block_nlls`, and `block_tokens`: summed NLLs and per-bucket prediction counts.\n* `total_tokens` and `total_input_tokens`: scored predictions and input tokens, respectively.\n* `total_sentences`: the number of evaluated contexts, including any one-token tail.\n* `config`, `model_commit`, `dataset_fingerprint`, and `versions`: evaluation settings, the resolved model commit when available, dataset fingerprint, and package versions.\n\n> [!NOTE]\n> Earlier versions appended a synthetic EOS only to positional statistics, used BF16 softmax for BF16 logits, and dropped the final partial context. Scores and token counts from those versions may differ from the corrected results.\n"},"files":{".agents/skills/README.md":"# Repository skills\n\n[CONTRIBUTING.md](../../CONTRIBUTING.md) defines contribution rules; [AGENTS.md](../../AGENTS.md) defines agent operations. Skills add task-specific guidance. Choose those relevant to the change.\n\nFor kernel work, define the contract and establish baseline coverage before optimization. Use the operator, dispatch, and hardware guides as needed, then prepare the PR with the collected evidence.\n\n| Skill                                                     | Use for                                                    |\n| --------------------------------------------------------- | ---------------------------------------------------------- |\n| [Design coverage](fla-design-coverage/SKILL.md)           | Supported inputs, numerical budgets, and routing contracts |\n| [Correctness coverage](fla-correctness-coverage/SKILL.md) | References, regression cases, and affected test selection  |\n| [KDA](fla-kda/SKILL.md)                                   | KDA gate modes and kernel invariants                       |\n| [Backend dispatch](fla-dispatch-backends/SKILL.md)        | Registration, verifiers, and fallback behavior             |\n| [Optimization loop](fla-optimization-loop/SKILL.md)       | A fixed correctness gate and measured candidate iterations |\n| [NVIDIA performance](fla-nvidia-performance/SKILL.md)     | GPU profiling and bottleneck diagnosis                     |\n| [Ascend performance](fla-ascend-performance/SKILL.md)     | NPU profiling and compiler constraints                     |\n| [Triton to Gluon](fla-triton-to-gluon/SKILL.md)           | Layouts, synchronization, and incremental kernel ports     |\n| [PR readiness](fla-pr-readiness/SKILL.md)                 | Scope, validation evidence, and submission                 |\n| [PR stacks](fla-pr-stacks/SKILL.md)                       | Splitting large changes, reusing PRs, and native stacks    |\n\n## Maintaining skills\n\nPut each workflow in `.agents/skills/<skill-name>/SKILL.md` with YAML `name` and `description` fields. Follow [documentation and skill style](../../CONTRIBUTING.md#prose-and-markdown) for scope, organization, wording, and shared policy.\n\nUse an existing `references/` directory for substantial task-specific detail. Add supporting files only when they make the workflow easier to use; `SKILL.md` is the entry point, so individual skill directories do not need a README. Symlinks may point to public documentation tracked in this repository.\n",".agents/skills/fla-ascend-performance/SKILL.md":"---\nname: fla-ascend-performance\ndescription: Profile and optimize FLA Triton-Ascend kernels using NPU traces, with guidance for UB capacity, memory movement, launch limits, and numerical correctness.\n---\n\n# Ascend kernel performance\n\nUse this skill for operator performance work in `triton_ascend.py` and `triton_ascend/` backend implementations. Follow [fla-optimization-loop](../fla-optimization-loop/SKILL.md) for sustained optimization and [CONTRIBUTING](../../../CONTRIBUTING.md#benchmarking) for validation and reporting.\n\nKeep collection, analysis, and benchmarks in the active Python/NPU environment. If it is not configured, activate the project's Ascend environment first.\n\n## 1. Establish the baseline\n\nLocate the public entry, dispatch route, and Ascend kernel. Record the supported inputs and production workloads, baseline commit, and synchronized latency. Preserve the public API, validated algorithm, reference, and tolerances. Confirm the intended NPU kernel runs; a Torch fallback must not hide an unsupported path or kernel bug.\n\n## 2. Collect a trace\n\nUse the existing collector instead of duplicating profiler setup. Run from the repository root; the workload file initializes its inputs and defines a repeatable `workload()` callable.\n\n```bash\npython .agents/skills/fla-ascend-performance/scripts/profile_npu.py \\\n  --name my_op --out-dir profile/my_op-npu \\\n  --metrics PipeUtilization --analyze \\\n  --kernel-filter my_kernel --exec-file path/to/workload.py\n```\n\nOne run collects one `aic_metrics` set. Start with `PipeUtilization`; collect `MemoryUB` separately when UB bandwidth is relevant. The default schedule is one warmup step and one active step. Warm the workload enough to exclude compilation from the measurement.\n\nThe collector prints the exact trace directory. To revisit it, pass that directory to [the analyzer](scripts/analyze_profile.py); see [collection and metrics](references/reference.md#collection-and-analysis) for the library interface and output layout.\n\n## 3. Diagnose bottlenecks\n\nUse `op_statistic.csv` to find the dominant operation, then `kernel_details.csv` to inspect its duration and pipe activity. Confirm the dispatch route before tuning a small or missing target kernel. Distinguish compilation failure, UB overflow, grid limits, numerical errors, and measured performance limits; they need different fixes.\n\nUse the [diagnosis table](references/reference.md#diagnosis) to choose the next experiment. Change one relevant factor, then measure whether the expected duration and pipe metrics moved.\n\n## 4. Optimize within the kernel contract\n\n- Estimate the peak live UB footprint and use `fla.utils.ascend_ub_manager` for device capacity and tiling. Model forward and backward separately; fuse only when the live set fits.\n- Match the launch to the work: Cube kernels use Cube cores, Vector kernels use Vector cores. Handle the grid limit through the existing helpers or a 1D task loop.\n- Preserve numerical precision, optional inputs, and fixed/variable-length behavior. Verify tail handling and pointer arithmetic before increasing tiles or fusing stages.\n- Omit `num_warps` and `num_stages` from Ascend launches and autotune settings; these are unsupported on this backend.\n- Check [compiler and memory traps](references/TRAPS.md) when changing DMA paths, pointer arithmetic, or reused `tl.dot` operands. Use [contiguous gate loading](references/g-contiguous-loading.md) when gate loads along time are strided.\n\n[Kernel cases](references/cases.md) explain existing convolution, recurrence, and solve implementations. Their tile choices are workload-specific examples.\n\n## 5. Validate and measure\n\nRun kernel comparisons and the affected operator/layer tests, including supported gradients, variable lengths, optional inputs, and boundary shapes. Keep NaN poisoning enabled and confirm that dispatch exercises the changed path. Run the full relevant correctness gate before reporting a speedup; a `--gate-k` subset is an iteration check.\n\nCompare synchronized forward and forward/backward latency on the same NPU, then repeat the relevant profiler collection. Report the timing change, the metrics that explain it, any compiler limitation encountered, and untested cases. If the expected metrics do not change, revisit the diagnosis before adding another optimization.\n",".agents/skills/fla-ascend-performance/references/TRAPS.md":"# Ascend compiler and memory traps\n\nRead the relevant section before changing DMA, launch indexing, or reused matrix operands. [Kernel cases](cases.md) show these constraints in existing implementations.\n\n## DMA paths and optional pointers\n\nA runtime branch between block-pointer DMA and masked loads can keep both paths live in UB. Larger tiles may then overflow even when measured UB bandwidth is low. Where the workload permits, split bulk and tail launches with a constexpr mode so each compilation removes the unused path. Include convolution halos when deciding whether a bulk load fits inside the packed allocation.\n\nAn optional pointer needs a compile-time guard before any runtime condition. `if USE_INITIAL_STATE or runtime_condition:` can still compile pointer arithmetic in the other branch when the pointer is `None`; nest the constexpr condition separately.\n\n## Address arithmetic\n\nCast indices to `tl.int64` before multiplying by strides or dimensions. Casting the product cannot repair overflow. On Ascend, specialized arguments and folded program IDs may be constexpr values without a `.to()` method; `tl.cast` works for both runtime and constexpr integers.\n\n```python\ntoken_offset = tl.cast(i_t, tl.int64) * BT\nbatch_offset = tl.cast(i_b, tl.int64) * T\nbos = tl.load(cu_seqlens + i_n).to(tl.int64)\n```\n\nLoad packed sequence offsets as int64 even when the input `cu_seqlens` uses int32. A small token index can overflow after multiplication by head count and feature size.\n\n`make_block_ptr` metadata has a separate constraint: its `offsets` and `block_shape` must remain int32. Keep large flattened base-address calculations in int64, then supply valid local block offsets. The Ascend backend is exempt from the mainline ban on block pointers.\n\n## Grid and task loops\n\nThe existing Ascend grid helpers enforce a grid-product limit of 65535. Use their host splitting or flatten independent tiles into a 1D loop over available cores. `get_multiprocessor_count()` selects Vector cores on NPU; `use_aicore=True` selects Cube cores.\n\nAfter slicing variable-length chunk metadata, apply the corresponding global offset only once. In a task loop, derive local pointers from the original base on each iteration; accumulating pointer updates across tasks can miscompile on Ascend.\n\n## Reused `tl.dot` operands\n\nTriton-Ascend can reuse the left operand's UB storage during `tl.dot`. A later read of that tile can therefore be wrong even without a compiler error. Audit all later uses, including another left operand, a right operand, arithmetic, and stores.\n\nReload the pristine tile from global memory when reuse is separated by stages. For tight reuse, create the needed `tile + 0.0` copies before the first dot that consumes the original as a left operand. Copying afterwards preserves the corrupted value. Account for the extra live copies in the UB budget and validate against the numerical reference.\n\n[The case index](cases.md#tldot-lhs-clobber--repo-wide-case-catalog) identifies the affected kernel patterns. A fresh load on each loop iteration or a tile used as a left operand only once needs no extra copy for this issue.\n\n## Numerical and tail behavior\n\nPreserve the validated accumulation precision and exponential base. Mask invalid values before exponentials and initialize every region consumed later. Factoring a gate difference into a ratio or reciprocal changes rounding and range behavior; check the existing reference and tolerances before using it.\n\nExercise gated and ungated paths when both are supported. A suite that always supplies a gate cannot catch `None` pointer handling or incorrect ungated residual updates.\n",".agents/skills/fla-ascend-performance/references/cases.md":"# Ascend kernel cases\n\nUse these examples to understand existing design choices. Read the current implementation before changing a path; tile settings depend on the workload and compiler. Paths are relative to the repository root.\n\n## Causal convolution: core grid and DMA paths\n\nImplementation: `fla/modules/causal_conv1d/backends/triton_ascend.py`.\n\nThe contiguous packed training path uses a 1D Vector-core task loop when channel tiling and state requirements permit. Other layouts, channel dimensions, and cache-state calls retain separate paths.\n\n- Transpose the public convolution weight to make channels contiguous for block loads. Preserve the public weight and output layouts.\n- Load one window including the convolution halo and use `extract_slice` for taps where supported. Keeping every tap live at once can exceed UB.\n- Preserve the saved pre-activation used by backward for silu/swish; removing it can trigger a second forward computation.\n- Separate bulk and tail DMA at compile time. A runtime choice keeps both paths live and can prevent larger tiles even when UB bandwidth is not saturated.\n\n| Path                               | DMA policy                                                     |\n| ---------------------------------- | -------------------------------------------------------------- |\n| Full chunks in packed input        | Block-pointer loads, with the tail path compiled out.          |\n| Final packed chunk                 | Masked loads/stores, including the halo bounds.                |\n| Variable lengths or a single chunk | Runtime bounds where a uniform bulk/tail split is unavailable. |\n\nUse `tests/modules/test_conv.py` for forward/backward, optional state, and layout validation. [TRAPS.md](TRAPS.md) covers optional pointers, packed bounds, and address types shared by these paths.\n\n## `chunk_delta_h.py` — bwd `dhu`\n\nImplementation: `fla/ops/common/backends/triton_ascend/chunk_delta_h.py`.\n\nGate loads along time use the [contiguous gate layout](g-contiguous-loading.md). The aligned fixed-length path can precompute gate terms, while variable-length or unaligned paths compute them inline. Model these paths separately: inline exponentials and temporary tiles can exceed the UB estimate for the precomputed path.\n\nChoose K and V tiles together. Reducing one dimension merely to maximize the other can increase repeated work. Rebuild task-local pointers from base addresses and return the corrected value gradient separately to keep the caller's `dv` unchanged. Backward also reuses dot operands across K slabs; apply the [operand lifetime rule](TRAPS.md#reused-tldot-operands).\n\n## `chunk_o.py` — fusion and gate loading\n\nImplementation: `fla/ops/common/backends/triton_ascend/chunk_o.py`.\n\nForward combines inter- and intra-chunk contributions when their live tiles fit UB, reducing intermediate output traffic. Backward gate loads use contiguous time storage while gradient outputs retain their public layout. Profile the combined launch and transpose cost, not just one kernel. See [gate loading](g-contiguous-loading.md) for pointer formulas and focused tests.\n\n## `chunk_bwd.py` — packed variable lengths\n\nImplementation: `fla/ops/kda/backends/triton_ascend/chunk_bwd.py`.\n\nPacked offsets multiplied by head and value dimensions can overflow int32 far before the token count reaches its limit. Load sequence boundaries as int64 and promote fixed-batch indices before multiplication. Only local lengths and block-pointer metadata should narrow to int32 when their bounds permit it.\n\n## `tl.dot` lhs clobber — repo-wide case catalog\n\nThe common rule is in [TRAPS.md](TRAPS.md#reused-tldot-operands). This index identifies where to inspect operand reuse without prescribing variable names or a fixed number of copies.\n\n| Implementation                                               | Reuse to inspect                                                                               |\n| ------------------------------------------------------------ | ---------------------------------------------------------------------------------------------- |\n| `fla/ops/common/backends/triton_ascend/chunk_o.py`           | Query/key tiles shared by multiple dots; attention tiles reused across value blocks.           |\n| `fla/ops/common/backends/triton_ascend/chunk_delta_h.py`     | Output-gradient and corrected residual tiles reused across K slabs.                            |\n| `fla/ops/gated_delta_rule/backends/triton_ascend/wy_fast.py` | Inverse tiles reused between stages; gradient tiles used first on the left, then on the right. |\n| `fla/ops/kda/backends/triton_ascend/wy_fast.py`              | Inverse tiles reused across the value and key loops.                                           |\n| `fla/ops/kda/backends/triton_ascend/chunk_intra.py`          | Triangular block merges that reuse operands for several products and final stores.             |\n| `fla/ops/utils/triton_ascend/solve_tril.py`                  | Block inverse operands used in multiple products or stored after a dot.                        |\n| `fla/ops/kda/backends/triton_ascend/chunk_bwd.py`            | Output gradients used as both left and right operands.                                         |\n\nAudit the changed kernel's first left-operand use and every subsequent read. Copies must precede that first use; reloads must come from pristine storage. Validate with the operator's kernel tests, including `tests/ops/test_gdn_kernels.py` and `tests/ops/test_solve_tril.py` where relevant.\n",".agents/skills/fla-ascend-performance/references/g-contiguous-loading.md":"# Contiguous gate loading\n\nUse this pattern when an Ascend kernel repeatedly loads gate values along time from `[B, T, HV]` storage. The time stride is `HV`, so these loads gather separated values. Transposing the input gate can reduce memory-transfer overhead; measure the transpose and kernel together.\n\nReference implementation: `fla/ops/common/backends/triton_ascend/chunk_o.py`.\n\n## Host layout\n\n```python\ng_t_contig = g is not None and HV != 1\ng_arg = g.transpose(1, 2).contiguous() if g_t_contig else g\n```\n\nPass the selected input and layout flag to the kernel. `HV == 1` is already contiguous along time. Keep gradient outputs in their original `[B, T, HV]` layout.\n\n## Kernel addressing\n\nSave the total packed length before computing a local sequence length. Promote batch, head, and packed sequence offsets before multiplying them; see [address arithmetic](TRAPS.md#address-arithmetic).\n\n| Input mode              | Contiguous gate base                    |\n| ----------------------- | --------------------------------------- |\n| Fixed batch             | `g + batch * HV * T_seq + head * T_seq` |\n| Packed variable lengths | `g + bos + head * T_seq`                |\n\nHere `batch`, `head`, and `bos` are int64 values, `T_seq` is the host-passed storage length, and the block-pointer bound uses the local sequence length. Block offsets are relative to that sequence's start.\n\nFor a contiguous gate base, use stride 1 along time. The original-layout branch instead uses `g + bos * HV + head` with time stride `HV`. Apply the base offset once and keep every load in a branch consistent with that layout. Reusing the original stride or offset after transposing produces wrong values or alignment faults.\n\n## Validation\n\nCompare both layout paths against the same reference, including head sharing, variable lengths, non-aligned tails, and long sequences. The focused tests are:\n\n- `tests/ops/test_gdn_kernels.py::test_chunk_bwd_dv_local`\n- `tests/ops/test_gdn_kernels.py::test_chunk_bwd_dqkwg`\n\nInspect wall-clock duration and MTE activity before and after, including transpose and gradient-finalization work. Sum-of-block profiler durations may differ from one launch's elapsed time.\n\nPreserve the exponential formula while changing layout. Replacing `exp2(a - b)` with a ratio can change overflow and rounding behavior. On Ascend, use the supported slicing/layout operations from the current implementation; a reshape or join that compiles does not prove the intended element ordering.\n",".agents/skills/fla-ascend-performance/references/reference.md":"# Ascend profiling reference\n\nUse this reference for collection options, CSV interpretation, and choosing the next optimization. Kernel correctness hazards are in [TRAPS.md](TRAPS.md); implementation examples are in [cases.md](cases.md).\n\n## Collection and analysis\n\nRun the scripts from the active NPU environment. `profile_npu.py` accepts a workload file or inline source defining `workload()`. It executes that callable repeatedly, so initialize reusable inputs outside the callable and make repeated forward/backward calls valid.\n\n```bash\npython .agents/skills/fla-ascend-performance/scripts/profile_npu.py \\\n  --name my_op --out-dir profile/my_op-npu \\\n  --metrics PipeUtilization --analyze --exec-file path/to/workload.py\n```\n\nFor library use, add the scripts directory to `PYTHONPATH` and call `profile_callable`:\n\n```bash\nexport PYTHONPATH=\"$PWD/.agents/skills/fla-ascend-performance/scripts${PYTHONPATH:+:$PYTHONPATH}\"\n```\n\n```python\nfrom profile_npu import profile_callable\n\ntrace_dir = profile_callable(fn=workload, name=\"my_op\", out_dir=\"profile/my_op-npu\", aic_metrics=\"PipeUtilization\")\n```\n\nThe returned directory contains the profiler output:\n\n```text\nmy_op_profiling_TIMESTAMP/\n└── localhost..._ascend_pt/\n    └── ASCEND_PROFILER_OUTPUT/\n        ├── op_statistic.csv\n        └── kernel_details.csv\n```\n\nPass that exact run directory when analyzing an existing trace. The analyzer takes the first matching CSV under the supplied directory, so a parent containing several runs can select an older result.\n\n```bash\npython .agents/skills/fla-ascend-performance/scripts/analyze_profile.py \\\n  profile/my_op-npu/my_op_profiling_TIMESTAMP --kernel-filter my_kernel --top-k 20\n```\n\n`--kernel-filter` applies to kernel details. If no name matches, the analyzer displays the top kernels instead; confirm the printed names before attributing metrics to the target.\n\n## Choosing metrics\n\nEach profiling run collects one `aic_metrics` set. Use separate runs for complementary metrics and keep the workload unchanged.\n\n| Metric set                                       | Use                                                                     |\n| ------------------------------------------------ | ----------------------------------------------------------------------- |\n| `PipeUtilization`                                | First pass: compare Cube, Vector, scalar, and memory-transfer activity. |\n| `MemoryUB`                                       | Check UB read/write bandwidth when memory traffic may limit the kernel. |\n| `Memory`, `MemoryAccess`, `MemoryL0`             | Inspect other memory paths.                                             |\n| `L2Cache`                                        | Investigate cache reuse.                                                |\n| `ArithmeticUtilization`, `ResourceConflictRatio` | Investigate arithmetic occupancy or resource conflicts.                 |\n\nAvailable metric names depend on the installed `torch_npu`; the collector reports available names when an option is unknown.\n\n## Diagnosis\n\nStart with total duration in `op_statistic.csv`, then inspect the dominant kernel in `kernel_details.csv`. Pipe activity alone does not establish the bottleneck.\n\n| Evidence                                               | Check next                                                                                                   |\n| ------------------------------------------------------ | ------------------------------------------------------------------------------------------------------------ |\n| Target kernel absent or many small fallback operations | Backend selection and fusion.                                                                                |\n| High Vector or Cube activity                           | Tile shape, useful work per tile, and scalar overhead around compute.                                        |\n| High MTE activity                                      | Reuse, intermediate writebacks, and strided loads; inspect [gate loading](g-contiguous-loading.md).          |\n| High scalar activity                                   | Address generation, branches, and repeated per-tile setup.                                                   |\n| High UB bandwidth with low compute activity            | Live buffers, layout conversions, and repeated UB reads/writes.                                              |\n| Low activity despite UB overflow at larger tiles       | Mutually exclusive DMA paths may remain live; see [DMA splitting](TRAPS.md#dma-paths-and-optional-pointers). |\n| Many host launches for grid chunks                     | Consider a 1D task loop with the appropriate Cube or Vector core count.                                      |\n\n## UB planning\n\nEstimate peak live storage from simultaneous buffers, their dtype, and any compiler-created copies. Use `fla.utils.ascend_ub_manager` for device capacity and tiling rather than hard-coding the UB size. The helper's safety margin reserves capacity; it is separate from the multiplier that estimates live buffers.\n\nIf more capacity is available, test a larger tile and verify compilation and numerics. If the live set already fills the budget, improve reuse or split stages. Model forward and backward separately; inline gate computation can need more UB than a precomputed-gate path.\n\nUseful code locations:\n\n- `fla/utils/ascend_ub_manager.py`: tiling and grid helpers.\n- `fla/utils/hardware.py`: `get_multiprocessor_count`, including `use_aicore=True` for Cube work.\n- `fla/modules/norm/triton_ascend/l2norm.py`: entry-point registration and the Ascend implementation.\n- `tests/ops/test_gdn_kernels.py`, `tests/ops/test_solve_tril.py`, and `tests/modules/test_conv.py`: relevant kernel comparisons.\n",".agents/skills/fla-correctness-coverage/SKILL.md":"---\nname: fla-correctness-coverage\ndescription: Select and run correctness coverage for FLA kernels and modules, including gradients, dispatch boundaries, and Triton addressing changes.\n---\n\n# FLA Correctness Coverage\n\nUse [fla-design-coverage](../fla-design-coverage/SKILL.md) when the change needs a new numerical or routing contract. This skill turns that contract into tests; [CONTRIBUTING.md](../../../CONTRIBUTING.md#testing) defines references, tolerances, test structure, and platform conventions.\n\n## Select coverage\n\nRead the affected tests and callers before adding cases. Follow [test organization and naming](../../../CONTRIBUTING.md#test-organization-and-naming): try the existing parameter matrix first, and add a function or file only when its setup or scope requires one. Cover reachable gaps across the applicable dimensions together, rather than testing each flag only in isolation:\n\n| Dimension   | Cases that can expose different behavior                                         |\n| ----------- | -------------------------------------------------------------------------------- |\n| Layout      | Dense and varlen, partial chunks, long sequences, asymmetric QK/value dimensions |\n| Features    | Gate bounds, raw/post-sigmoid beta, QK normalization, grouped value attention    |\n| State       | Initial state, final state, and state gradients where supported                  |\n| Numerics    | Supported dtypes, affected conversion/reduction paths, numerical boundaries      |\n| Routing     | Default and selected backends, verifier acceptance, rejection, and fallback      |\n\nCompare outputs and final states against the existing reference. Training APIs also need every supported gradient; forward-only APIs must state that backward is unsupported. Use `torch.autograd.gradcheck` where the implementation supports its required precision. Routing tests verify selection and argument forwarding; they do not replace numerical comparisons on the actual backend.\n\nFor addressing changes, inspect casts before multiplication and test shapes that expose large offsets, partial tiles, and varlen boundaries. Follow the `tl.int64` and platform requirements in [Triton Kernels](../../../CONTRIBUTING.md#triton-kernels). Exercise changed grid/address paths on affected supported platforms. Distinguish unsupported cases from tests skipped for unavailable hardware or dependencies; record the latter as validation gaps.\n\nThe NaN allocation guard covers eligible FLA allocations in operator/module tests, including the single-rank CP graph and intra-card cache cases moved from operator tests, but not layer/model tests, spawned CP workers, or every allocation. Add explicit finite-output and finite-gradient checks to adversarial cases, and investigate a poisoned-run failure even when an ad hoc run passes.\n\n## Run the affected paths\n\nDiscover dependent tests from the repository root:\n\n```bash\npython scripts/find_dependent_tests.py fla/ops/kda/chunk.py\nFLA_CI_ENV=0 pytest tests/ops/test_kda.py -v\n```\n\nThe helper lists files; it does not run them. Include reported callers and relevant `tests/modules/`, `tests/layers/`, `tests/models/`, or `tests/cp/` coverage when the change reaches those paths. Inspect skips and warning-only comparisons before calling the run a pass. Record commands, backend flags, environment, and results; reproduce failures on the unchanged baseline before labeling them pre-existing.\n\nRead operator math only when needed: [Delta Rule](../../../fla/ops/delta_rule/README.md), [Generalized Delta Rule](../../../fla/ops/generalized_delta_rule/README.md), [Simple GLA](../../../fla/ops/simple_gla/README.md), or [context parallelism](../../../fla/ops/cp/README.md).\n",".agents/skills/fla-correctness-coverage/references/cp.md":"../../../../fla/ops/cp/README.md",".agents/skills/fla-correctness-coverage/references/delta-rule.md":"../../../../fla/ops/delta_rule/README.md",".agents/skills/fla-correctness-coverage/references/generalized-delta-rule.md":"../../../../fla/ops/generalized_delta_rule/README.md",".agents/skills/fla-correctness-coverage/references/simple-gla.md":"../../../../fla/ops/simple_gla/README.md",".agents/skills/fla-design-coverage/SKILL.md":"---\nname: fla-design-coverage\ndescription: Define supported inputs, numerical budgets, routing, and validation before designing an FLA kernel or numerical change.\n---\n\n# FLA Design and Coverage\n\nUse this skill to establish the contract before implementing a kernel, backend route, or numerical change. Read [CONTRIBUTING.md](../../../CONTRIBUTING.md) for contribution rules, especially [public compatibility](../../../CONTRIBUTING.md#public-api-compatibility) and [design/RFC requirements](../../../CONTRIBUTING.md#submit-pull-requests).\n\n## Define the supported domain\n\nTrace public entries and defaults, backend verifiers, layer/model/CP callers, and existing tests. Record the merge-base SHA and the source of each supported case. Group cases only when they share the same numerical assumptions and execution route; omit combinations no supported caller can reach.\n\nDistinguish routes when kernel stages, chunk geometry or anchors, capability checks, dtype conversions, or input value ranges change. A backend here means an implementation selected by dispatch, including the default implementation, rather than a hardware product. For each group, identify its effective arguments, oracle, output/state and per-gradient tolerances, and test coverage. Assign one outcome:\n\n- The proposed implementation, with passing reference comparisons.\n- An existing fallback, with route-parity tests covering outputs, states, dtypes, and supported gradients.\n- A documented unsupported-input error raised before kernel execution and asserted by a test.\n\nPreserve previously supported cases under the public compatibility policy. A warning, partially initialized output, or fallback that changes semantics is not a valid outcome. When changing defaults, record the old and new effective values and selected routes.\n\n## Establish numerical safety\n\nFor each affected stage in forward and backward, record input-range assumptions, load/operand/accumulation/store dtypes, reduction behavior, overflow/underflow headroom, and error propagation to the end-to-end tolerance. Forward agreement alone does not establish backward safety, including inverse recomputation and state gradients.\n\nSafety flags and bounds apply to the stages and domains that justify them. Reassess them when chunk geometry, anchors, dtype staging, or input scaling changes; a safe exponent does not prove that a later inversion is safe. Document the algebra behind numerical acceptance thresholds, implement each shared predicate once, and test its acceptance boundary and the adjacent rejection or fallback. Enforce additional assumptions, such as input-norm bounds, at the public entry or every affected verifier; report the failed predicate and effective values.\n\nKeep the committed oracle, input distributions, assertions, and tolerances fixed while evaluating the candidate. A correction to that contract belongs in a separate change with before/after error distributions for the unchanged baseline and candidate; precision, tolerance, and algorithm changes follow the linked RFC policy. Keep the corrected gate green before the candidate lands. Every selectable autotune configuration must satisfy the same numerical contract.\n\n## Check routing\n\n- Share semantic argument normalization between the public entry, verifier, and executor. Test that omitted arguments match their documented explicit defaults in route, outputs, and gradients. `input_guard` handles contiguity and device context, not semantic defaults.\n- Use capability and availability helpers from `fla.utils`. Distinguish backend availability, policy enablement, and acceptance of an individual call; use [fla-dispatch-backends](../fla-dispatch-backends/SKILL.md) for implementation mechanics.\n- For context-parallel routes, every rank must make the same selection. If verifier inputs can differ by rank, combine acceptance with a collective and select the optimized route only if all ranks accept. Record whether forward commits backward to the same backend.\n- Test parity between routes claiming the same domain. Record precision settings such as `TRITON_F32_DEFAULT` and matmul flags, keep them identical across baseline, candidate, and tests, and restore global settings after tests.\n\n## Choose evidence\n\nApply the relevant axes from [fla-correctness-coverage](../fla-correctness-coverage/SKILL.md) to each supported group across these distinct layers:\n\n- **Production:** reproduce effective arguments and numerical input distributions from checked-in layer/model callers. A generic benchmark shape list alone does not establish production coverage.\n- **Public boundaries:** cover previously supported inputs, changed defaults, and verifier acceptance/rejection boundaries.\n- **Adversarial:** outside the contract, require documented rejection before kernel execution or explicit finite-output and finite-gradient checks. This does not expand the supported domain.\n\nBefore performance work, record baseline correctness, latency or throughput, commit, environment, workload, input distribution, and seed. Benchmark production cases first, then boundaries whose routing or expected performance changes. Unaffected boundaries still need correctness coverage. Use [fla-optimization-loop](../fla-optimization-loop/SKILL.md) to freeze the gate and compare candidates.\n",".agents/skills/fla-dispatch-backends/SKILL.md":"---\nname: fla-dispatch-backends\ndescription: Change FLA backend registration, dispatch, or verifiers while preserving routing and call contracts.\n---\n\n# Backend dispatch\n\nUse this skill for registration, selection, or verifier changes in `fla/backends.py` and operation backends. [CONTRIBUTING.md](../../../CONTRIBUTING.md) defines style and public compatibility; [fla/backends.py](../../../fla/backends.py) defines the routing behavior.\n\n## 1. Define the supported calls\n\nIdentify the default entry point and the calls the backend can handle: inputs, layouts, dtypes, gradient modes, and optional dependencies. Match the entry point's signature, parameter order, defaults, return structure, and supported gradients.\n\nA verifier returns `(True, None)` or `(False, reason)`. It checks support without changing inputs or global state. Keep it beside the implementation; put helpers shared by several implementations in a utility module.\n\n## 2. Register the implementation\n\nKeep the default entry point backend-independent with bare `@dispatch`. Choose the registration form that matches the implementation:\n\n| Implementation               | Registration                                                         |\n| ---------------------------- | -------------------------------------------------------------------- |\n| Function with the same API   | `@register(entry, backend=TritonAscendBackend, verifier=...)`        |\n| Adapter with backend policy  | `@register(entry)` on a `BaseBackend` subclass                       |\n\nAn adapter defines methods named after the entry point, such as `chunk_kda` and `chunk_kda_verifier`. Register it where it is defined and import optional kernels only when selected.\n\nImport the owning package's backends after its default entry points are defined. For implementations with platform-specific dependencies, guard the import with the shared backend's `is_available()`. Do not gate registration on enable switches; dispatch reads those at call time. Keep per-operation switches on `@register(..., env_var=...)` and implementation paths out of the shared dispatcher.\n\nExisting owners using `@dispatch('<operation>')` use `@register('<operation>')` on adapter classes and load implementations through the owner's backend package. Both forms use the same registration and selection logic.\n\n## 3. Preserve routing behavior\n\nFor each call, try backends in ascending priority order; ties retain registration order:\n\n1. Check availability and enablement.\n2. Run the verifier; rejection proceeds to the next candidate.\n3. Call the selected implementation. Its exceptions propagate.\n4. If no candidate accepts, call the default implementation.\n\nRegistration replaces an existing backend of the same type. Function registration keys the registry by the unwrapped entry point, so compiler decorators and public aliases share that registry. Legacy owner registries must load their backend package before lookup.\n\nKeep selection outside `torch.compile` graphs. Do not cache availability or enablement in the dispatch wrapper: registered backends must remain eligible when their switches change. Python's import cache handles modules already loaded, and unused optional dependencies must not prevent package import.\n\n## 4. Verify through the entry point\n\nRun the decorated entry point for accepted calls, verifier rejection, disabled/unavailable backends, and fallback. Check optional arguments omitted as well as supplied, changed support boundaries, outputs, and supported gradients.\n\nUse explicit backend flags in routing tests. For the default path, set `FLA_DISABLE_BACKEND_DISPATCH=1` before importing FLA. Use fresh processes for import-order checks, and verify public aliases and shared backend ownership after registration changes.\n\nFind affected consumers with `scripts/find_dependent_tests.py` and run the relevant operator, layer, and model tests. Host routing tests establish dispatch behavior; accelerator reference tests establish numerical correctness.\n",".agents/skills/fla-kda/SKILL.md":"---\nname: fla-kda\ndescription: Modify or review KDA gates, chunk and recurrent kernels, backend routes, and their correctness coverage.\n---\n\n# KDA\n\nUse this skill for `fla/ops/kda/` and its tests. Public entry points are `chunk_kda` and `fused_recurrent_kda`; inspect the affected entry point's signature before changing its implementation.\n\n## Code map\n\n- `gate.py` defines gate activations and cumulative decay.\n- `chunk_fwd.py` and `chunk_bwd.py` coordinate training; `chunk_intra.py`, `chunk_intra_token_parallel.py`, and `wy_fast.py` implement the main stages.\n- `backends/` contains alternative implementations and verifiers. Follow [backend dispatch](../fla-dispatch-backends/SKILL.md) when changing routing.\n- `tests/ops/test_kda.py` contains the reference comparisons and supported gate-mode cases.\n\n## Gate contracts\n\n`chunk_kda` takes log-space decay when `use_gate_in_kernel=False`. With it enabled, the caller supplies raw gates and optional bias:\n\n- Without a lower bound, activation is `-exp(A_log) * softplus(g + dt_bias)` and `A_log` is required.\n- With a lower bound, activation is `lower_bound * sigmoid(exp(A_log) * (g + dt_bias))`. Omitting `A_log` uses a multiplier of one.\n\n`safe_gate` selects the intra-chunk implementation independently of gate activation. Safe in-kernel activation requires `-5 <= lower_bound < 0`. Pre-gated safe mode requires bounded log-space gates supplied by the caller; tests cover `[-5, 0]`, but the wrapper does not validate tensor values.\n\n## Numerical invariants\n\nThe safe intra path uses 16-token diagonal blocks and midpoint offsets before exponentiation. At a lower bound of `-5`, a block can accumulate `-80` in natural-log units; local offsets limit each exponent's span to about half that range. Preserve these offsets rather than exponentiating the full cumulative gate. Inter-block decay uses paired differences whose exponents remain non-positive for monotonic decay.\n\nThe non-safe path uses token-parallel diagonal computation. Both paths share inter-block and triangular-solve work, so changes there require coverage of both modes. Keep precision and tolerance changes within the design process in [AGENTS.md](../../../AGENTS.md#scope-and-direction).\n\n## Validation\n\nUse [correctness coverage](../fla-correctness-coverage/SKILL.md) for the test strategy. Select the KDA cases affected by the change:\n\n- Dense and variable-length sequences; forward, final state, and supported gradients.\n- Pre-gated and in-kernel gates, safe and non-safe paths, and omitted optional arguments.\n- Beta activation and query/key normalization modes when touched.\n- Grouped value heads and differing key/value dimensions when indexing or shapes change.\n- Initial states, intermediate states, and context parallelism when their paths are touched.\n- Accepted and rejected backend routes when changing verifiers.\n\nFor gate or decay changes, include the lower-bound limits, saturated gate inputs, extreme gate scales, long cumulative decay, and chunk or ragged-sequence boundaries. Test through the public entry point so the selected route and effective arguments are exercised.\n",".agents/skills/fla-nvidia-performance/SKILL.md":"---\nname: fla-nvidia-performance\ndescription: Profile and optimize FLA kernels on NVIDIA GPUs, with same-hardware benchmarks and targeted Nsight Compute analysis.\n---\n\n# NVIDIA kernel performance\n\nUse this skill for NVIDIA kernel optimization and backend tuning. Follow [CONTRIBUTING](../../../CONTRIBUTING.md#benchmarking) for correctness and performance evidence, and [fla-optimization-loop](../fla-optimization-loop/SKILL.md) for sustained iteration.\n\n## Hardware baseline\n\nUse datacenter NVIDIA GPUs with sm_90 or newer for PR performance conclusions; H100 and H20 are accepted, and sm_100/sm_103 are preferred. Results from A100, older GPUs, or consumer cards are supplementary. Record the exact GPU and software versions with each comparison.\n\n## Workflow\n\n1. Identify the changed kernel, its callers, and the production workloads it affects. Record a baseline commit and confirm which backend runs.\n2. Pass the relevant correctness tests before timing a candidate. Preserve the reference and tolerances; include outputs, states, and supported gradients.\n3. Compare baseline and candidate on the same GPU with identical workloads and dispatch settings. Include dense and supported variable-length cases. Use an end-to-end benchmark for a training or generation throughput claim.\n4. Use Nsight Compute when a bottleneck or regression needs hardware metrics. Match the optimization to the measured limit, then rerun correctness and timing.\n5. Report latency or throughput, relevant memory changes, and any regressions. Keep raw profiler artifacts outside git, for example under `profile/<run_name>/`.\n\n## Benchmark commands\n\nRun from the repository root. Resolve the baseline to a commit SHA before comparing:\n\n```bash\nFLA_BENCH_BASE=$(git rev-parse origin/main)\nFLA_CI_ENV=0 python -m benchmarks.ops.verify --op chunk_kda --base \"$FLA_BENCH_BASE\"\n\npython benchmarks/benchmark_training_throughput.py --name kda --batch_size 2 --seq_len 8192\npython benchmarks/benchmark_training_throughput.py --name kda --batch_size 2 --seq_len 8192 --varlen\n```\n\nThe verification command runs the current tests; retain the validated baseline and tests yourself. A selected subset is useful during iteration but does not replace the full relevant gate before reporting a gain.\n\n## Nsight Compute\n\nCollect a representative changed kernel with the installed `ncu`. Check which sections and metrics that version supports.\n\nThe benchmark runner launches a child process, so use `--target-processes all`. Replace the kernel filter with the changed kernel's name. `--no-base` profiles the current implementation; collect the baseline separately with the same command and workload.\n\n```bash\nmkdir -p profile/kda\nncu --target-processes all --set full \\\n  -k 'regex:chunk_kda' -c 1 -o profile/kda/full \\\n  python -m benchmarks.ops.run --op chunk_kda --modes fwd --no-base\n\nncu --target-processes all --set source --section SourceCounters \\\n  -k 'regex:chunk_kda' -c 1 -o profile/kda/source \\\n  python -m benchmarks.ops.run --op chunk_kda --modes fwd --no-base\n```\n\nUse duration, memory throughput, occupancy, stalls, and hot instructions to explain the bottleneck. Profiling overhead makes these runs unsuitable as the latency comparison; use the synchronized benchmark for that measurement.\n",".agents/skills/fla-optimization-loop/SKILL.md":"---\nname: fla-optimization-loop\ndescription: Iterate on FLA kernel performance with a frozen correctness gate, a measured baseline, and reproducible candidate comparisons.\n---\n\n# FLA Optimization Loop\n\nUse this skill for repeated kernel performance work. Define supported workloads and numerical assumptions with [fla-design-coverage](../fla-design-coverage/SKILL.md); use the [NVIDIA](../fla-nvidia-performance/SKILL.md) or [Ascend](../fla-ascend-performance/SKILL.md) skill for profiling on that platform.\n\n## Freeze the gate and baseline\n\nAdd missing regression or new-path coverage before starting the loop. Then keep the relevant tests, reference implementation, input distributions, assertions, tolerances, and public signature fixed. Do not narrow parameter matrices, add skips, relax precision or thresholds, special-case test values, cache results across trials, or change timing behavior to make a candidate pass. Fully initialize exposed outputs. Use kernel implementations for the optimization; replacing the whole operator with a vendor call solely to win the benchmark does not meet this workflow.\n\nA suspected oracle or tolerance defect needs a separately justified correction, not an adjustment that makes the candidate pass. Follow [CONTRIBUTING.md](../../../CONTRIBUTING.md#submit-pull-requests) for precision, numerical-algorithm, and tolerance changes, and its [compatibility policy](../../../CONTRIBUTING.md#public-api-compatibility) for public behavior.\n\nRecord the entry point, target workloads, allowed implementation languages, validation/benchmark commands, baseline SHA, and success criterion in the local optimization log. Include production inputs from callers as well as applicable benchmark-registry shapes. Confirm the baseline passes the full gate and measure it before changing kernels. For a new kernel, establish correctness before tuning.\n\n```bash\nFLA_BENCH_BASE=$(git rev-parse origin/main)\nFLA_CI_ENV=0 python -m benchmarks.ops.verify --op chunk_gla --base \"$FLA_BENCH_BASE\"\n```\n\nUse a recorded SHA: the baseline worktree cannot reuse a checked-out branch, and slash-containing refs break its temporary directory naming. `--base` benchmarks that revision; it does not run its correctness tests. Establish baseline correctness separately if the current tree already differs.\n\nThe driver runs the current pytest file and stops before timing when pytest fails. It neither freezes files nor rejects an all-skipped run. Check that the expected cases and benchmark rows ran. `--test-file` overrides a mismatched derived test path; `--gate-k` selects a quick subset, while `--no-gate` produces unverified measurements. Neither subset nor skipped-gate results qualify for promotion.\n\n## Iterate from evidence\n\n1. Profile the current bottleneck and choose one reasoned change.\n2. Run the frozen gate, then benchmark only a passing candidate. A subset can provide an early signal; the final gate covers the full affected suite.\n3. Record the change, gate result, measurement, and keep/drop decision before the next experiment. Preserve a recoverable candidate under the repository's git rules and the user's authorization; failed and unchanged attempts need records, not empty commits.\n\nRead [measurement and correctness traps](references/TRAPS.md) before interpreting results. Keep seeds, numeric flags, hardware, and workload fixed. Candidate runtime helps rank local iterations; measure the final candidate directly against the frozen baseline in one session rather than adding speedups from separate runs.\n\nAdd specialized paths only when profiling shows different bottlenecks across shape regimes and measurements justify the extra routing. Record each condition, chosen path, and before/after result, and validate the full shape set including fallback boundaries.\n\nAfter repeated attempts without an improvement above noise, re-profile and review the failed directions before continuing. Respect the user's stopping condition. A no-win conclusion needs a measured baseline, a reasoned candidate attempt, gate status, benchmark evidence, and an identified bound or blocker. An unavailable profiler or one losing candidate alone does not establish a performance floor; report unresolved validation as a limitation.\n\n## Keep and report the result\n\nKeep iteration notes and raw traces under ignored `profile/<op>-opt/`; [opt-log-template.md](references/opt-log-template.md) provides a compact record format. This search history stays out of kernel comments and the final PR description.\n\nPromote only with a full green gate, a repeatable win on the claimed workloads, and profiling or a quantified bottleneck analysis that explains it. Report per-shape results and regressions, commands, baseline/candidate revisions, environment, and measurement variability using the relevant hardware skill. An unexplained large speedup needs investigation. Use [fla-pr-readiness](../fla-pr-readiness/SKILL.md) to prepare the final diff and evidence.\n",".agents/skills/fla-optimization-loop/references/TRAPS.md":"# Correctness and Measurement Traps\n\n## Partial writes\n\n`tests/conftest.py` poisons eligible FLA `empty` allocations in operator/module tests with NaNs. It excludes some allocations and does not cover layer/model/CP tests. A poisoned-run failure that disappears in an ad hoc script can reveal uninitialized reads; inspect masks and ensure every exposed output element is written. NaNs can also have numerical causes, so diagnose the failure rather than assuming its source.\n\n## Reference precision and tolerance\n\n`fla.utils.assert_close` checks relative RMS error with an absolute-error shortcut, not elementwise `rtol`. `FLA_CI_ENV=1` and explicit `warning=True` can downgrade tolerance failures to warnings. Use strict local validation and inspect warnings; keep the committed thresholds fixed.\n\nTF32 in a reference matmul can change an fp32 comparison. Inspect the reference's precision settings before blaming the kernel, but do not change the frozen oracle to pass. Baseline and candidate need identical precision and numeric flags. An intentional precision change follows the numerical-change process in [CONTRIBUTING.md](../../../../CONTRIBUTING.md#submit-pull-requests).\n\n## Address overflow\n\nCast program IDs and offsets to `tl.int64` before multiplication; widening an already-overflowed product cannot repair it. Small shapes may hide the error. Run large-offset and affected varlen/grid cases under the [Triton addressing rules](../../../../CONTRIBUTING.md#triton-kernels).\n\n## Missing work in a fast result\n\nCheck unexpectedly large speedups against the actual selected path, executed tests, and measured shapes. `verify.py` accepts pytest's exit status, including an all-skipped success. The benchmark runner filters constrained shapes and omits failed warmups or measurements. Compare expected and returned rows before making a claim; a passing process alone is insufficient.\n\n## Autotune and clocks\n\n`run.py`, also used by `verify.py`, warms all shapes before timing. Preserve that warmup. After changing the configuration space, check cached selections and rewarm or clear the relevant cache if results look stale.\n\nUnlocked clocks and changing device load can shift latency between runs. Rank local candidates by their own runtime, then remeasure the final candidate against the recorded baseline in one session. Report variability and the device state; speedups from separate experiments do not add up to a final result.\n",".agents/skills/fla-optimization-loop/references/opt-log-template.md":"# Optimization Record\n\nKeep one local log under ignored `profile/<op>-opt/`. Record the entry point, target workloads, baseline SHA and correctness, environment, commands, numeric settings, and success criterion once. Preserve enough candidate identity to reproduce each measurement.\n\n| Attempt  | Change or hypothesis | Gate | Latency by workload | Decision and reason |\n| -------- | -------------------- | ---- | ------------------- | ------------------- |\n| Baseline | Unchanged revision   | Pass | Measured baseline   | Comparison point    |\n\nAppend an entry after each attempt, including failures and unchanged results. Identify whether the gate was full or selected; record no performance result for a failed gate. Retained candidates need their revision or recoverable diff, the reason they helped, and any unresolved limitation. Keep failed directions here so later work does not repeat them.\n\nWhen specializing, add each routing condition, entry point, baseline/candidate latency, and measured reason for a separate path. Compare every bucket against the same baseline revision on that bucket's workloads, including fallback boundaries. Keep these details alongside the iteration results.\n",".agents/skills/fla-pr-readiness/SKILL.md":"---\nname: fla-pr-readiness\ndescription: Prepare or update an FLA pull request with focused scope, verified evidence, and the current repository template.\n---\n\n# FLA PR Readiness\n\nRead [CONTRIBUTING.md](../../../CONTRIBUTING.md) and [AGENTS.md](../../../AGENTS.md) first. Use the current [PR template](../../../.github/pull_request_template.md) for headings and checkboxes; contribution policy and template wording are not duplicated here.\n\n## Review scope and evidence\n\nRead the complete diff against the PR base and audit affected shared callers. State its concrete outcome in one sentence and separate independent work. Check existing issues/PRs and the applicable design discussion. Follow [scope and size](../../../CONTRIBUTING.md#pr-scope-and-size) for a cohesive change above the normal limit, and [public compatibility](../../../CONTRIBUTING.md#public-api-compatibility) when interfaces or defaults are affected.\n\nWhen a change needs splitting or dependent PRs need a native GitHub stack, use [PR stacks](../fla-pr-stacks/SKILL.md) to choose boundaries, reuse existing PRs, and verify the published stack.\n\nFor code changes, discover dependent tests and run the relevant targets. For example:\n\n```bash\npython scripts/find_dependent_tests.py fla/ops/kda/chunk.py\nFLA_CI_ENV=0 pytest tests/ops/test_kda.py\npre-commit run --files fla/ops/kda/chunk.py tests/ops/test_kda.py\n```\n\nFor documentation-only changes, validate links, examples, and formatting as described in [Evidence by Change Type](../../../CONTRIBUTING.md#evidence-by-change-type).\n\nThe dependency helper only lists tests. Check that regression coverage exercises the changed path and that required numerical comparisons passed, including supported gradients. Review new test functions and files against [test organization and naming](../../../CONTRIBUTING.md#test-organization-and-naming); use an existing test case where it fits. Record commands, tested revision, environment, and results; distinguish failures, skips, and unrun checks. Reproduce a failure on the baseline before calling it pre-existing.\n\nFor performance-relevant changes, collect same-hardware before/after measurements under [Benchmarking](../../../CONTRIBUTING.md#benchmarking) and the relevant hardware skill. A neutral result needs numbers; explain inapplicable checks. Inspect the final diff for stale comments and unrelated changes, leaving scratch logs and raw profiles outside it.\n\n## Prepare and publish\n\nWrite for a reviewer who has not seen the conversation. Follow [PR Description](../../../CONTRIBUTING.md#pr-description) and [Evidence by Change Type](../../../CONTRIBUTING.md#evidence-by-change-type): explain the problem, resulting behavior, affected users, and evidence. Include only useful rationale or limitations; omit file inventories and the history of attempts. Keep incomplete work in draft and leave claims and checklist items consistent with actual validation.\n\nThe `check-pr-title` workflow also checks the body and size exception on edits. Preserve template checklist wording and distinguish required items from conditional ones.\n\n`gh pr edit` fails with this repository's classic-Projects integration. Write the complete body to a temporary file and update through REST:\n\n```bash\ngh api -X PATCH \"repos/fla-org/flash-linear-attention/pulls/$pr_number\" -F \"body=@$pr_body_file\"\n```\n\nConfirm the published description, head revision, and checks. After review changes, rerun affected validation and update evidence that no longer describes the current code. Publishing and git operations remain subject to the user's authorization and repository rules.\n",".agents/skills/fla-pr-stacks/SKILL.md":"---\nname: fla-pr-stacks\ndescription: Split oversized FLA pull requests into focused changes, reuse existing PRs, and create or maintain verified GitHub native PR stacks.\n---\n\n# FLA PR Stacks\n\nUse for decomposing large changes or managing dependent PRs. Follow [CONTRIBUTING.md](../../../CONTRIBUTING.md#pr-scope-and-size) for scope and size, [AGENTS.md](../../../AGENTS.md) for git permissions, and [PR readiness](../fla-pr-readiness/SKILL.md) for each PR's evidence and description.\n\n## 1. Choose the split\n\nInspect the complete diff, commit history, and related open PRs, including their review discussions. Map each proposed layer to one outcome, its prerequisites, validation, and an existing PR to reuse where appropriate.\n\n| Relationship                            | Action                                                                   |\n| --------------------------------------- | ------------------------------------------------------------------------ |\n| Overlaps an existing PR's purpose       | Update that PR when ownership permits; preserve attribution and reviews. |\n| Independent of the other changes        | Submit separately against the trunk.                                     |\n| Depends on another layer                | Stack above that layer; include its own tests and documentation.         |\n| Cohesive but still above the size limit | Use the contribution guide's exception with a concrete justification.    |\n\nEach layer presents an increment over its parent and must work with its prerequisites, without needing later layers to fix it. Separate review does not imply dependency-free code. Do not split by arbitrary line counts, leave callers broken, or defer tests to a later PR.\n\nCheck head repository identity before arranging existing PRs. [GitHub stacks require all branches in the same repository](https://docs.github.com/en/pull-requests/reference/stacked-pull-requests); cross-fork stacks are unsupported. Keep incompatible PRs as explicit external prerequisites instead of duplicating them under new PR numbers.\n\n## 2. Build and validate\n\nPreserve the original head and uncommitted work before reconstructing commits in isolated branches or worktrees. Reuse suitable commits and preserve authorship. Check current remote heads before updating existing PRs.\n\nReview each layer against its parent, then compare the top branch with the preserved original result:\n\n```bash\ngit diff main...refactor/foundation\ngit diff refactor/foundation...refactor/consumer\ngit diff original-head refactor/consumer\n```\n\nSubstitute the recorded trunk and actual refs. Account for trunk changes and intentional omissions, including fixes retained in existing PRs. Run relevant tests at each layer's head and on the combined result; passing only the top branch does not validate intermediate PRs.\n\n## 3. Register and verify the native stack\n\nChanging PR bases alone does not create a native stack. Use the official [`gh stack` extension](https://docs.github.com/en/pull-requests/how-tos/create-pull-requests/creating-stacked-pull-requests), GitHub UI, or REST Stacks API. Check installed command help and existing publishing authorization first.\n\nFor prepared branches, adopt them bottom to top and submit:\n\n```bash\ngh stack init --base main refactor/foundation refactor/consumer\ngh stack submit --auto --remote origin\n```\n\nChoose the actual remote containing the branches. `submit` pushes, creates or updates PRs, and registers the stack. `--auto` creates new PRs as drafts; replace generated titles and bodies using PR readiness.\n\nFor existing PRs, use `gh stack link --base main --remote origin 123 124` with their real numbers in dependency order. This reuses PRs without requiring local stack tracking. Branch arguments also push and create missing PRs. After an error, inspect what succeeded before retrying.\n\nFor every PR, verify [remote stack membership](https://docs.github.com/en/pull-requests/reference/stacked-pull-requests-apis-and-webhooks):\n\n```bash\ngh api \"repos/fla-org/flash-linear-attention/pulls/$pr_number\" \\\n  --jq '{number, head_repo: .head.repo.full_name, head: .head.ref, base: .base.ref, stack}'\n```\n\nRequire non-null `stack` objects with the same stack identity, correct order, size, and trunk. Check each PR's incremental diff and report the verified links and remaining external prerequisites. A local branch chain or successful push is insufficient.\n\n## 4. Maintain the series\n\nAfter changing a lower layer, review affected upper diffs and rerun relevant tests. `gh stack rebase` rewrites history; `gh stack sync` also force-pushes with leases. Apply the repository's rewrite permissions and preserve concurrent remote work, then recheck published membership and evidence.\n\nCreating a series does not authorize merging or closing PRs. If merging is requested, establish the intended prefix: selecting an upper PR also merges those below it. Use a stack-aware merge operation; the legacy synchronous merge API does not support stacks.\n",".agents/skills/fla-triton-to-gluon/SKILL.md":"---\nname: fla-triton-to-gluon\ndescription: Port an FLA Triton kernel to Gluon when explicit layouts, asynchronous transfers, or scheduling can address a measured bottleneck.\n---\n\n# Triton to Gluon\n\nUse Gluon when profiling points to register pressure, layout conversion, or load/compute overlap that needs explicit control. A kernel already saturating its useful bandwidth or compute limit may not benefit. Keep the portable default path when adding hardware-specific implementations.\n\nFollow [fla-optimization-loop](../fla-optimization-loop/SKILL.md) for the baseline and correctness gate, and [fla-nvidia-performance](../fla-nvidia-performance/SKILL.md) for NVIDIA measurements.\n\n## Check the target environment\n\nGluon is experimental. Check the installed Triton API before copying [upstream examples](https://triton-lang.org/main/getting-started/tutorials/gluon/); their names and signatures can differ from the installed release.\n\nGate architecture-specific code appropriately: Ampere supports `cp.async`, Hopper adds TMA/WGMMA, and Blackwell adds TMEM/tcgen05. Code under Gluon's `nvidia` namespace requires NVIDIA hardware. Keep compile-time choices in explicit constexpr arguments so the kernel, launcher, and autotune pruning use the same values.\n\n```python\nfrom triton.experimental import gluon\nfrom triton.experimental.gluon import language as gl\n```\n\n## Port incrementally\n\n1. Record the baseline commit and retain the validated tests, reference, and tolerances. Keep the Triton fallback available.\n2. Translate loads, stores, and arithmetic with explicit layouts first. Match each tensor's contiguous dimension and preserve masking. Pass forward, state, and supported backward comparisons before adding asynchrony.\n3. Add async transfers only where load latency or addressing pressure limits the kernel. Define buffer ownership, completion, and reuse for the prologue, steady state, and drain.\n4. Add architecture-specific MMA or scheduling when profiling supports it. Recheck register/shared-memory use and retune after each change.\n5. Run the full relevant correctness gate and same-hardware benchmarks, including supported variable-length and boundary cases. Measure production entry points as well as individual kernels.\n\nA literal translation establishes parity; measure it before assuming a performance gain. Preserve the validated numerical algorithm and precision while changing execution layout or scheduling.\n\n## Layout and resource constraints\n\n- Pointer tensors and values need compatible layouts. A cross-warp `gl.convert_layout` may use shared memory; use `assert_trivial=True` when a conversion is intended to be free.\n- Broadcasts and oversized layout tiles can duplicate values and increase register use. Choose layouts from actual tensor strides rather than copying a tutorial's block sizes.\n- Account for all live shared-memory buffers and the target device's allocation limit. Prune infeasible autotune configurations; provide a streaming path when a resident design cannot fit supported shapes.\n- Large static unrolls multiplied by many autotune configurations can dominate compilation. Prune from the actual shape and live-memory budget, and reuse the worker and Triton cache during iteration.\n- Persistent schedules can trade launch overhead for worse L2 locality. Warp specialization also changes the total register budget; measure occupancy and stalls after changing either.\n\n## Async correctness\n\n### Transfers and buffer reuse\n\n`cp.async` groups are ordered as a queue. `wait_group(N)` bounds all outstanding groups, so it cannot identify a particular buffer after unrelated prefetches are interleaved. For per-buffer scheduling, use separate mbarriers and match their arrival counts to the participating threads. In the installed API, check whether `mbarrier_arrive` increments the count; a preinitialized thread count requires the non-incrementing form.\n\nWait for a load before consuming its buffer, and protect reuse until all readers finish. Even same-lane shared-memory staging needs protection against overwriting data still being read. Track each mbarrier's phase as its buffer is reused; do not run more than one phase ahead or share a completion barrier between TMA and tcgen05 without reinitializing it.\n\nTMA descriptors must satisfy the target's alignment and stride requirements. TMA stores may distinguish completion of the shared-memory read from completion of the global write. Check the installed wait API and require global completion before another operation reads the stored range.\n\n### MMA and memory ordering\n\nHopper WGMMA requires its B operand in shared memory and uses register accumulators; consume the values returned by its wait operation so compiler dependencies remain explicit. Blackwell tcgen05 uses TMEM accumulators and mbarrier completion; respect the participating warpgroup and TMEM layout requirements. Initialize accumulators explicitly, including `use_acc=False` where supported.\n\nGeneric shared-memory accesses and async operations use different memory proxies. Apply the required proxy fence when handing a buffer to an async consumer; an mbarrier alone does not replace that fence. A completed TMA load establishes the ordering needed to read its destination. Check ordering across warp-specialized partitions as well as within a single producer/consumer loop.\n\n### Tails and exact-zero cases\n\nMasked async copies can leave shared-memory elements uninitialized. Prevent invalid rows from entering reductions: zero-fill where supported, or use valid-row loads and explicitly remove every invalid contribution. Multiplying by zero does not neutralize a NaN. Keep output stores masked.\n\nDo not rely on separately compiled reductions cancelling bitwise. For a mathematically exact-zero special case, such as a single-source softmax gradient, preserve the exact-zero result explicitly and test it against the reference.\n\n## Verification commands\n\nRun from the repository root on the target GPU:\n\n```bash\nFLA_BENCH_BASE=$(git rev-parse origin/main)\nFLA_CI_ENV=0 python -m benchmarks.ops.verify --op chunk_kda --base \"$FLA_BENCH_BASE\"\n```\n\nReplace the operation with the one being ported. The command executes the current tests; it does not freeze them. Use `--gate-k` only for quick iteration, then run the full gate. Set backend environment flags before launching a fresh process so import-time choices cannot contaminate the comparison.\n","AGENTS.md":"# FLA Guidelines\n\nGuidance for AI coding agents (Claude Code, Codex, etc.) working in this repo.\n\n**Read `CONTRIBUTING.md` first.** It is the authoritative source for all code style, docstring, comment, commit, PR, and testing conventions, and applies to humans and agents alike. This file only covers agent-specific operational behavior that doesn't belong in a human contributor doc.\n\n## Git safety\n\n- **Never discard uncommitted work** with `git checkout HEAD -- <file>`, `git reset --hard`, or `git restore` to \"get a clean base\". Unstaged changes are unrecoverable (no blob, no reflog). Edit in place or `git stash` instead, and confirm with the user when in doubt.\n- **On `main`**: never commit or push without explicit per-action approval. Suggest a feature branch first.\n- Don't rewrite or amend already-pushed commits unless the user asks.\n- **Branch names and attribution**: name local and remote branches after the change. Do not include tool or assistant identities in branch names, commit messages, PR titles/bodies, or review comments. Do not add `Co-Authored-By` trailers or generated-by notices.\n\n## Scope and direction\n\n- **Check for duplicates first**: search open issues/PRs before starting so you don't redo in-flight work — `gh pr list --repo fla-org/flash-linear-attention --state open --search \"<keywords>\"`.\n- **Align before coding when a design decision is involved**: a new operator or model, a new public API, a kernel rewrite with a different algorithm, or a repo-wide policy change (e.g. adding a dtype everywhere). Open an issue or a draft PR first; if a matching issue/PR already exists, comment there instead of opening a duplicate. Judge design impact separately from the PR size check: tests, config boilerplate, and mechanical repetition of an existing pattern alone do not make a change a design decision.\n- **RFC first for precision or algorithm changes**: changing the precision a validated kernel computes or accumulates in (e.g. fp32 → bf16, bf16 → fp16), relaxing a tolerance, or replacing a numerical algorithm (e.g. a different inversion or solve scheme) needs an RFC issue before a PR, even if it would stay within test tolerance. FLA's kernels are widely validated and downstream users rely on their numerics, so these are design decisions to agree on, not fixes to ship. Ordinary reorderings within the same precision — tiling, accumulation order, scheduling — do not need an RFC; they are normal optimization PRs with before/after benchmarks as usual.\n- **Stop and ask before breaking changes**: renaming or removing public symbols or arguments in `fla/layers/` / `fla/models/` or documented config fields; changing checkpoint/state-dict compatibility; changing observable behavior, including numerics beyond the existing test tolerance. Bug fixes that restore intended behavior are not breaking. If the user doesn't answer — or you are running non-interactively — do the non-breaking parts, leave the breaking part out, and describe it in your final report. Silence is not consent.\n\n## Comment discipline\n\nThe style rules for comments and docstrings live in `CONTRIBUTING.md` (\"Docstrings and Comments\"), including the banned anti-patterns. What follows is agent-specific behavior:\n\n- **You own the comments your change invalidates.** When a change makes a comment or docstring in the code you touched factually wrong, fix it or delete it in the same commit — \"the code you touched\" means the same function, kernel, or class, not a repo-wide sweep for stale mentions. If you can't tell whether a comment is still true, keep it and flag it in the PR description; never delete a \"why\" comment just because you can't verify it.\n- **Comments are not your scratchpad.** State design conclusions, not the journey: \"not X, it does not compile\" is a welcome why comment; \"tried X first, then switched to Y\" is chain-of-thought noise. This covers docstrings too.\n- **Fix the words, not the style.** Correcting stale content is not a license to reformat the surrounding docstring or comment style.\n\n## PR workflow\n\n- **No busywork PRs**: don't open a one-off PR for a single typo or isolated style tweak; bundle trivial cleanups into substantive work.\n- **Minimal diff**: every changed line serves the PR's stated purpose. Drop drive-by reformats, unrelated cleanups, and speculative generality before pushing, not after a reviewer asks (see CONTRIBUTING \"Core Principles\" 5).\n- **Blast radius**: for every shared symbol you touched, all callsites still hold (Core Principle 4 — audit them in one pass, don't assume). For any behavior or numerics change, the PR description must state who is affected — which shapes, dtypes, backends, checkpoints, or downstream users — or state explicitly why no one is. \"Tests pass\" is not an impact analysis.\n- Before opening a PR, re-read your own diff once with fresh eyes: check every comment it adds against the anti-pattern list, and every comment it invalidates against [Comment discipline](#comment-discipline).\n- Check the diff for `tl.make_block_ptr` / `tl.advance`; these are banned from mainline Triton code. Follow the Ascend backend exemption in CONTRIBUTING \"Triton Kernels\".\n- `gh pr edit` fails on this repo (classic-Projects GraphQL error). Edit a PR title/body via the REST API instead: `gh api -X PATCH repos/fla-org/flash-linear-attention/pulls/<N> -f title='...' -F body=@file`.\n- Keep review/PR comments concise and natural — skip heavy `**1.** **2.**` scaffolding, write like a person.\n\n## Repo-local Skills\n\nUse the [skill index](.agents/skills/README.md) to select task-specific workflows and their execution order. Skills supplement `CONTRIBUTING.md` and this file; they do not define competing contribution policies.\n\nLoad a relevant skill when the task matches its scope. In particular, load [fla-pr-readiness](.agents/skills/fla-pr-readiness/SKILL.md) before opening any PR.\n","evals/README.md":"# Evaluation\n\nWe provide evaluation scripts for `fla` models with [lm-evaluation-harness](https://github.com/EleutherAI/lm-evaluation-harness), as well as a standalone perplexity evaluator.\nInstall `fla` following the [installation guide](../INSTALL.md), and run the commands below from the repository root.\n\n* [lm-evaluation-harness](#lm-evaluation-harness)\n  * [Single-GPU Evaluation](#single-gpu-evaluation)\n  * [Multi-GPU Evaluation](#multi-gpu-evaluation)\n  * [RULER Benchmarks](#ruler-benchmarks)\n* [Perplexity](#perplexity)\n  * [Quick Start](#quick-start)\n  * [Evaluation Settings](#evaluation-settings)\n  * [Scoring](#scoring)\n  * [Results](#results)\n\n## lm-evaluation-harness\n\nThe [lm-evaluation-harness](https://github.com/EleutherAI/lm-evaluation-harness) library allows you to easily perform (zero-shot) model evaluations.\nInstall `lm_eval` following [their instructions](https://github.com/EleutherAI/lm-evaluation-harness/blob/main/README.md).\n\n### Single-GPU Evaluation\n\nRun evaluation with:\n\n```sh\nMODEL='fla-hub/gla-1.3B-100B'\npython -m evals.harness --model hf \\\n    --model_args pretrained=$MODEL,dtype=bfloat16 \\\n    --tasks wikitext,lambada_openai,piqa,hellaswag,winogrande,arc_easy,arc_challenge,boolq,sciq,copa,openbookqa \\\n    --batch_size 64 \\\n    --num_fewshot 0 \\\n    --device cuda \\\n    --show_config\n```\n\nWe've made `fla` compatible with hf-style evaluations, you can call [evals.harness](harness.py) to finish the evaluations.\nThe command above evaluates the GLA checkpoint on the tasks used in the GLA paper.\n\n### Multi-GPU Evaluation\n\nUse the Hugging Face Accelerate launcher for data-parallel evaluation, where each GPU loads a full copy of the model:\n\n```sh\nMODEL='fla-hub/gla-1.3B-100B'\naccelerate launch -m evals.harness --model hf \\\n    --model_args pretrained=$MODEL,dtype=bfloat16,trust_remote_code=True \\\n    --tasks wikitext,lambada_openai,piqa,hellaswag,winogrande,arc_easy,arc_challenge,boolq,sciq,copa,openbookqa \\\n    --batch_size 64 \\\n    --num_fewshot 0 \\\n    --device cuda \\\n    --show_config \\\n    --trust_remote_code\n```\n\n### RULER Benchmarks\n\nThe RULER benchmarks are commonly used for evaluating model performance on long-context tasks.\nYou can evaluate `fla` models on RULER directly using `lm-evaluation-harness`. RULER is only available in a relatively recent version of `lm-evaluation-harness`, so make sure you have the latest version installed.\n\n```sh\ngit clone --depth 1 https://github.com/EleutherAI/lm-evaluation-harness\npip install -e './lm-evaluation-harness[ruler]'\n```\n\nFor example, to evaluate contexts up to 32K tokens:\n\n```sh\nMODEL='fla-hub/gla-1.3B-100B'\nOUTPUT='results/ruler'\naccelerate launch -m evals.harness --model hf \\\n    --output_path $OUTPUT \\\n    --tasks niah_single_1,niah_single_2,niah_single_3,niah_multikey_1,niah_multikey_2,niah_multikey_3,niah_multiquery,niah_multivalue,ruler_vt,ruler_cwe,ruler_fwe,ruler_qa_hotpot,ruler_qa_squad \\\n    --model_args pretrained=$MODEL,dtype=bfloat16,max_length=32768,trust_remote_code=True \\\n    --metadata='{\"max_seq_lengths\":[4096,8192,16384,32768]}' \\\n    --batch_size 2 \\\n    --show_config \\\n    --trust_remote_code\n```\n\nIf a GPU can't load a full copy of the model, please refer to [this link](https://github.com/EleutherAI/lm-evaluation-harness?tab=readme-ov-file#multi-gpu-evaluation-with-hugging-face-accelerate) for FSDP settings.\n\n> [!TIP]\n> If you are using `lm-evaluation-harness` as an external library and can't find (almost) any tasks available, before calling `lm_eval.evaluate()` or `lm_eval.simple_evaluate()`, simply run the following to load the library's stock tasks:\n> ```py\n> from lm_eval.tasks import TaskManager\n> TaskManager().initialize_tasks()\n> ```\n\n## Perplexity\n\n[evals.ppl](ppl.py) evaluates next-token perplexity over fixed-length contexts and reports perplexity by position within each context.\n\n### Quick Start\n\nInstall the dataset dependency:\n\n```sh\npip install datasets\n```\n\nFor example, to evaluate the GLA checkpoint on the PG19 test split and save the results:\n\n```sh\npython -m evals.ppl \\\n    --path fla-hub/gla-1.3B-100B \\\n    --data fla-hub/pg19 --split test \\\n    --block_size 28672 --bucket_size 2048 --batch_size 1 \\\n    --dtype bfloat16 --num_proc 4 \\\n    --output results/ppl.json\n```\n\n### Evaluation Settings\n\n| Argument          | Default                 | Description                                                 |\n|-------------------|-------------------------|-------------------------------------------------------------|\n| `--path`          | `fla-hub/gla-1.3B-100B` | Hugging Face model ID or local checkpoint directory.        |\n| `--data`          | `fla-hub/pg19`          | Dataset to evaluate.                                        |\n| `--split`         | `train`                 | Dataset split; set this explicitly for held-out evaluation. |\n| `--column_name`   | `text`                  | Dataset column containing text.                             |\n| `--block_size`    | `28672`                 | Context length in tokens; must be at least 2.               |\n| `--bucket_size`   | `2048`                  | Number of prediction positions per reported bucket.         |\n| `--batch_size`    | `1`                     | Number of full contexts evaluated together.                 |\n| `--device`        | Auto-detected           | Device on which to load and evaluate the model.             |\n| `--dtype`         | `bfloat16`              | Model dtype: `float32`, `float16`, or `bfloat16`.           |\n| `--num_proc`      | `32`                    | Number of dataset tokenization processes.                   |\n| `--revision`      | `main`                  | Model and tokenizer revision.                               |\n| `--data_revision` | `main`                  | Dataset revision.                                           |\n| `--output`        | None                    | Optional path for JSON results and evaluation settings.     |\n\nUse immutable commit hashes for `--revision` and `--data_revision` to reproduce a run.\nFor long contexts, reduce `--batch_size` to lower model memory use and `--bucket_size` to reduce the FP32 softmax workspace.\nThe model still materializes its full logits.\n\n### Scoring\n\nDocuments are tokenized with the tokenizer's default special-token policy, then concatenated in dataset order without extra separators.\nThe token stream is divided into independent contexts of `block_size` tokens, and the final partial context is evaluated separately without padding.\nContext is reset at every block boundary; there is no sliding window or recurrent state carried between blocks.\n\nEach context of length `L` contributes `L-1` next-token predictions.\nThe first token is unscored, and no synthetic EOS target is appended after the last token.\nEOS tokens present in the input are scored normally.\nA one-token tail contributes no predictions; a dataset with fewer than two tokens raises `ValueError`.\n\nOverall PPL and positional PPL use the same FP32 negative log-likelihoods, weighted by the number of scored predictions.\nBuckets group predictions by their zero-based logit position within each context: bucket 0 covers positions `0` through `bucket_size-1`.\nModels configured with fused linear cross entropy are supported, and evaluation runs with caching disabled.\n\n### Results\n\nThe script prints overall PPL, positional PPL, and token counts.\nWith `--output`, it also saves a JSON file containing:\n\n* `perplexity` and `block_perplexities`: overall and positional PPL; buckets with no predictions have `null` PPL.\n* `total_nll`, `block_nlls`, and `block_tokens`: summed NLLs and per-bucket prediction counts.\n* `total_tokens` and `total_input_tokens`: scored predictions and input tokens, respectively.\n* `total_sentences`: the number of evaluated contexts, including any one-token tail.\n* `config`, `model_commit`, `dataset_fingerprint`, and `versions`: evaluation settings, the resolved model commit when available, dataset fingerprint, and package versions.\n\n> [!NOTE]\n> Earlier versions appended a synthetic EOS only to positional statistics, used BF16 softmax for BF16 logits, and dropped the final partial context. Scores and token counts from those versions may differ from the corrected results.\n"},"items":[{"name":"AGENTS.md","path":"AGENTS.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/AGENTS.md","title":"AI Agent Protocol & Instructions","category":"root-instruction","format":"markdown","content":"# FLA Guidelines\n\nGuidance for AI coding agents (Claude Code, Codex, etc.) working in this repo.\n\n**Read `CONTRIBUTING.md` first.** It is the authoritative source for all code style, docstring, comment, commit, PR, and testing conventions, and applies to humans and agents alike. This file only covers agent-specific operational behavior that doesn't belong in a human contributor doc.\n\n## Git safety\n\n- **Never discard uncommitted work** with `git checkout HEAD -- <file>`, `git reset --hard`, or `git restore` to \"get a clean base\". Unstaged changes are unrecoverable (no blob, no reflog). Edit in place or `git stash` instead, and confirm with the user when in doubt.\n- **On `main`**: never commit or push without explicit per-action approval. Suggest a feature branch first.\n- Don't rewrite or amend already-pushed commits unless the user asks.\n- **Branch names and attribution**: name local and remote branches after the change. Do not include tool or assistant identities in branch names, commit messages, PR titles/bodies, or review comments. Do not add `Co-Authored-By` trailers or generated-by notices.\n\n## Scope and direction\n\n- **Check for duplicates first**: search open issues/PRs before starting so you don't redo in-flight work — `gh pr list --repo fla-org/flash-linear-attention --state open --search \"<keywords>\"`.\n- **Align before coding when a design decision is involved**: a new operator or model, a new public API, a kernel rewrite with a different algorithm, or a repo-wide policy change (e.g. adding a dtype everywhere). Open an issue or a draft PR first; if a matching issue/PR already exists, comment there instead of opening a duplicate. Judge design impact separately from the PR size check: tests, config boilerplate, and mechanical repetition of an existing pattern alone do not make a change a design decision.\n- **RFC first for precision or algorithm changes**: changing the precision a validated kernel computes or accumulates in (e.g. fp32 → bf16, bf16 → fp16), relaxing a tolerance, or replacing a numerical algorithm (e.g. a different inversion or solve scheme) needs an RFC issue before a PR, even if it would stay within test tolerance. FLA's kernels are widely validated and downstream users rely on their numerics, so these are design decisions to agree on, not fixes to ship. Ordinary reorderings within the same precision — tiling, accumulation order, scheduling — do not need an RFC; they are normal optimization PRs with before/after benchmarks as usual.\n- **Stop and ask before breaking changes**: renaming or removing public symbols or arguments in `fla/layers/` / `fla/models/` or documented config fields; changing checkpoint/state-dict compatibility; changing observable behavior, including numerics beyond the existing test tolerance. Bug fixes that restore intended behavior are not breaking. If the user doesn't answer — or you are running non-interactively — do the non-breaking parts, leave the breaking part out, and describe it in your final report. Silence is not consent.\n\n## Comment discipline\n\nThe style rules for comments and docstrings live in `CONTRIBUTING.md` (\"Docstrings and Comments\"), including the banned anti-patterns. What follows is agent-specific behavior:\n\n- **You own the comments your change invalidates.** When a change makes a comment or docstring in the code you touched factually wrong, fix it or delete it in the same commit — \"the code you touched\" means the same function, kernel, or class, not a repo-wide sweep for stale mentions. If you can't tell whether a comment is still true, keep it and flag it in the PR description; never delete a \"why\" comment just because you can't verify it.\n- **Comments are not your scratchpad.** State design conclusions, not the journey: \"not X, it does not compile\" is a welcome why comment; \"tried X first, then switched to Y\" is chain-of-thought noise. This covers docstrings too.\n- **Fix the words, not the style.** Correcting stale content is not a license to reformat the surrounding docstring or comment style.\n\n## PR workflow\n\n- **No busywork PRs**: don't open a one-off PR for a single typo or isolated style tweak; bundle trivial cleanups into substantive work.\n- **Minimal diff**: every changed line serves the PR's stated purpose. Drop drive-by reformats, unrelated cleanups, and speculative generality before pushing, not after a reviewer asks (see CONTRIBUTING \"Core Principles\" 5).\n- **Blast radius**: for every shared symbol you touched, all callsites still hold (Core Principle 4 — audit them in one pass, don't assume). For any behavior or numerics change, the PR description must state who is affected — which shapes, dtypes, backends, checkpoints, or downstream users — or state explicitly why no one is. \"Tests pass\" is not an impact analysis.\n- Before opening a PR, re-read your own diff once with fresh eyes: check every comment it adds against the anti-pattern list, and every comment it invalidates against [Comment discipline](#comment-discipline).\n- Check the diff for `tl.make_block_ptr` / `tl.advance`; these are banned from mainline Triton code. Follow the Ascend backend exemption in CONTRIBUTING \"Triton Kernels\".\n- `gh pr edit` fails on this repo (classic-Projects GraphQL error). Edit a PR title/body via the REST API instead: `gh api -X PATCH repos/fla-org/flash-linear-attention/pulls/<N> -f title='...' -F body=@file`.\n- Keep review/PR comments concise and natural — skip heavy `**1.** **2.**` scaffolding, write like a person.\n\n## Repo-local Skills\n\nUse the [skill index](.agents/skills/README.md) to select task-specific workflows and their execution order. Skills supplement `CONTRIBUTING.md` and this file; they do not define competing contribution policies.\n\nLoad a relevant skill when the task matches its scope. In particular, load [fla-pr-readiness](.agents/skills/fla-pr-readiness/SKILL.md) before opening any PR.\n","isInternal":false,"tokens":1322,"sizeBytes":5902},{"name":"cases.md","path":".agents/skills/fla-ascend-performance/references/cases.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/.agents/skills/fla-ascend-performance/references/cases.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# Ascend kernel cases\n\nUse these examples to understand existing design choices. Read the current implementation before changing a path; tile settings depend on the workload and compiler. Paths are relative to the repository root.\n\n## Causal convolution: core grid and DMA paths\n\nImplementation: `fla/modules/causal_conv1d/backends/triton_ascend.py`.\n\nThe contiguous packed training path uses a 1D Vector-core task loop when channel tiling and state requirements permit. Other layouts, channel dimensions, and cache-state calls retain separate paths.\n\n- Transpose the public convolution weight to make channels contiguous for block loads. Preserve the public weight and output layouts.\n- Load one window including the convolution halo and use `extract_slice` for taps where supported. Keeping every tap live at once can exceed UB.\n- Preserve the saved pre-activation used by backward for silu/swish; removing it can trigger a second forward computation.\n- Separate bulk and tail DMA at compile time. A runtime choice keeps both paths live and can prevent larger tiles even when UB bandwidth is not saturated.\n\n| Path                               | DMA policy                                                     |\n| ---------------------------------- | -------------------------------------------------------------- |\n| Full chunks in packed input        | Block-pointer loads, with the tail path compiled out.          |\n| Final packed chunk                 | Masked loads/stores, including the halo bounds.                |\n| Variable lengths or a single chunk | Runtime bounds where a uniform bulk/tail split is unavailable. |\n\nUse `tests/modules/test_conv.py` for forward/backward, optional state, and layout validation. [TRAPS.md](TRAPS.md) covers optional pointers, packed bounds, and address types shared by these paths.\n\n## `chunk_delta_h.py` — bwd `dhu`\n\nImplementation: `fla/ops/common/backends/triton_ascend/chunk_delta_h.py`.\n\nGate loads along time use the [contiguous gate layout](g-contiguous-loading.md). The aligned fixed-length path can precompute gate terms, while variable-length or unaligned paths compute them inline. Model these paths separately: inline exponentials and temporary tiles can exceed the UB estimate for the precomputed path.\n\nChoose K and V tiles together. Reducing one dimension merely to maximize the other can increase repeated work. Rebuild task-local pointers from base addresses and return the corrected value gradient separately to keep the caller's `dv` unchanged. Backward also reuses dot operands across K slabs; apply the [operand lifetime rule](TRAPS.md#reused-tldot-operands).\n\n## `chunk_o.py` — fusion and gate loading\n\nImplementation: `fla/ops/common/backends/triton_ascend/chunk_o.py`.\n\nForward combines inter- and intra-chunk contributions when their live tiles fit UB, reducing intermediate output traffic. Backward gate loads use contiguous time storage while gradient outputs retain their public layout. Profile the combined launch and transpose cost, not just one kernel. See [gate loading](g-contiguous-loading.md) for pointer formulas and focused tests.\n\n## `chunk_bwd.py` — packed variable lengths\n\nImplementation: `fla/ops/kda/backends/triton_ascend/chunk_bwd.py`.\n\nPacked offsets multiplied by head and value dimensions can overflow int32 far before the token count reaches its limit. Load sequence boundaries as int64 and promote fixed-batch indices before multiplication. Only local lengths and block-pointer metadata should narrow to int32 when their bounds permit it.\n\n## `tl.dot` lhs clobber — repo-wide case catalog\n\nThe common rule is in [TRAPS.md](TRAPS.md#reused-tldot-operands). This index identifies where to inspect operand reuse without prescribing variable names or a fixed number of copies.\n\n| Implementation                                               | Reuse to inspect                                                                               |\n| ------------------------------------------------------------ | ---------------------------------------------------------------------------------------------- |\n| `fla/ops/common/backends/triton_ascend/chunk_o.py`           | Query/key tiles shared by multiple dots; attention tiles reused across value blocks.           |\n| `fla/ops/common/backends/triton_ascend/chunk_delta_h.py`     | Output-gradient and corrected residual tiles reused across K slabs.                            |\n| `fla/ops/gated_delta_rule/backends/triton_ascend/wy_fast.py` | Inverse tiles reused between stages; gradient tiles used first on the left, then on the right. |\n| `fla/ops/kda/backends/triton_ascend/wy_fast.py`              | Inverse tiles reused across the value and key loops.                                           |\n| `fla/ops/kda/backends/triton_ascend/chunk_intra.py`          | Triangular block merges that reuse operands for several products and final stores.             |\n| `fla/ops/utils/triton_ascend/solve_tril.py`                  | Block inverse operands used in multiple products or stored after a dot.                        |\n| `fla/ops/kda/backends/triton_ascend/chunk_bwd.py`            | Output gradients used as both left and right operands.                                         |\n\nAudit the changed kernel's first left-operand use and every subsequent read. Copies must precede that first use; reloads must come from pristine storage. Validate with the operator's kernel tests, including `tests/ops/test_gdn_kernels.py` and `tests/ops/test_solve_tril.py` where relevant.\n","isInternal":false,"tokens":1086,"sizeBytes":5527},{"name":"g-contiguous-loading.md","path":".agents/skills/fla-ascend-performance/references/g-contiguous-loading.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/.agents/skills/fla-ascend-performance/references/g-contiguous-loading.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# Contiguous gate loading\n\nUse this pattern when an Ascend kernel repeatedly loads gate values along time from `[B, T, HV]` storage. The time stride is `HV`, so these loads gather separated values. Transposing the input gate can reduce memory-transfer overhead; measure the transpose and kernel together.\n\nReference implementation: `fla/ops/common/backends/triton_ascend/chunk_o.py`.\n\n## Host layout\n\n```python\ng_t_contig = g is not None and HV != 1\ng_arg = g.transpose(1, 2).contiguous() if g_t_contig else g\n```\n\nPass the selected input and layout flag to the kernel. `HV == 1` is already contiguous along time. Keep gradient outputs in their original `[B, T, HV]` layout.\n\n## Kernel addressing\n\nSave the total packed length before computing a local sequence length. Promote batch, head, and packed sequence offsets before multiplying them; see [address arithmetic](TRAPS.md#address-arithmetic).\n\n| Input mode              | Contiguous gate base                    |\n| ----------------------- | --------------------------------------- |\n| Fixed batch             | `g + batch * HV * T_seq + head * T_seq` |\n| Packed variable lengths | `g + bos + head * T_seq`                |\n\nHere `batch`, `head`, and `bos` are int64 values, `T_seq` is the host-passed storage length, and the block-pointer bound uses the local sequence length. Block offsets are relative to that sequence's start.\n\nFor a contiguous gate base, use stride 1 along time. The original-layout branch instead uses `g + bos * HV + head` with time stride `HV`. Apply the base offset once and keep every load in a branch consistent with that layout. Reusing the original stride or offset after transposing produces wrong values or alignment faults.\n\n## Validation\n\nCompare both layout paths against the same reference, including head sharing, variable lengths, non-aligned tails, and long sequences. The focused tests are:\n\n- `tests/ops/test_gdn_kernels.py::test_chunk_bwd_dv_local`\n- `tests/ops/test_gdn_kernels.py::test_chunk_bwd_dqkwg`\n\nInspect wall-clock duration and MTE activity before and after, including transpose and gradient-finalization work. Sum-of-block profiler durations may differ from one launch's elapsed time.\n\nPreserve the exponential formula while changing layout. Replacing `exp2(a - b)` with a ratio can change overflow and rounding behavior. On Ascend, use the supported slicing/layout operations from the current implementation; a reshape or join that compiles does not prove the intended element ordering.\n","isInternal":false,"tokens":550,"sizeBytes":2496},{"name":"reference.md","path":".agents/skills/fla-ascend-performance/references/reference.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/.agents/skills/fla-ascend-performance/references/reference.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# Ascend profiling reference\n\nUse this reference for collection options, CSV interpretation, and choosing the next optimization. Kernel correctness hazards are in [TRAPS.md](TRAPS.md); implementation examples are in [cases.md](cases.md).\n\n## Collection and analysis\n\nRun the scripts from the active NPU environment. `profile_npu.py` accepts a workload file or inline source defining `workload()`. It executes that callable repeatedly, so initialize reusable inputs outside the callable and make repeated forward/backward calls valid.\n\n```bash\npython .agents/skills/fla-ascend-performance/scripts/profile_npu.py \\\n  --name my_op --out-dir profile/my_op-npu \\\n  --metrics PipeUtilization --analyze --exec-file path/to/workload.py\n```\n\nFor library use, add the scripts directory to `PYTHONPATH` and call `profile_callable`:\n\n```bash\nexport PYTHONPATH=\"$PWD/.agents/skills/fla-ascend-performance/scripts${PYTHONPATH:+:$PYTHONPATH}\"\n```\n\n```python\nfrom profile_npu import profile_callable\n\ntrace_dir = profile_callable(fn=workload, name=\"my_op\", out_dir=\"profile/my_op-npu\", aic_metrics=\"PipeUtilization\")\n```\n\nThe returned directory contains the profiler output:\n\n```text\nmy_op_profiling_TIMESTAMP/\n└── localhost..._ascend_pt/\n    └── ASCEND_PROFILER_OUTPUT/\n        ├── op_statistic.csv\n        └── kernel_details.csv\n```\n\nPass that exact run directory when analyzing an existing trace. The analyzer takes the first matching CSV under the supplied directory, so a parent containing several runs can select an older result.\n\n```bash\npython .agents/skills/fla-ascend-performance/scripts/analyze_profile.py \\\n  profile/my_op-npu/my_op_profiling_TIMESTAMP --kernel-filter my_kernel --top-k 20\n```\n\n`--kernel-filter` applies to kernel details. If no name matches, the analyzer displays the top kernels instead; confirm the printed names before attributing metrics to the target.\n\n## Choosing metrics\n\nEach profiling run collects one `aic_metrics` set. Use separate runs for complementary metrics and keep the workload unchanged.\n\n| Metric set                                       | Use                                                                     |\n| ------------------------------------------------ | ----------------------------------------------------------------------- |\n| `PipeUtilization`                                | First pass: compare Cube, Vector, scalar, and memory-transfer activity. |\n| `MemoryUB`                                       | Check UB read/write bandwidth when memory traffic may limit the kernel. |\n| `Memory`, `MemoryAccess`, `MemoryL0`             | Inspect other memory paths.                                             |\n| `L2Cache`                                        | Investigate cache reuse.                                                |\n| `ArithmeticUtilization`, `ResourceConflictRatio` | Investigate arithmetic occupancy or resource conflicts.                 |\n\nAvailable metric names depend on the installed `torch_npu`; the collector reports available names when an option is unknown.\n\n## Diagnosis\n\nStart with total duration in `op_statistic.csv`, then inspect the dominant kernel in `kernel_details.csv`. Pipe activity alone does not establish the bottleneck.\n\n| Evidence                                               | Check next                                                                                                   |\n| ------------------------------------------------------ | ------------------------------------------------------------------------------------------------------------ |\n| Target kernel absent or many small fallback operations | Backend selection and fusion.                                                                                |\n| High Vector or Cube activity                           | Tile shape, useful work per tile, and scalar overhead around compute.                                        |\n| High MTE activity                                      | Reuse, intermediate writebacks, and strided loads; inspect [gate loading](g-contiguous-loading.md).          |\n| High scalar activity                                   | Address generation, branches, and repeated per-tile setup.                                                   |\n| High UB bandwidth with low compute activity            | Live buffers, layout conversions, and repeated UB reads/writes.                                              |\n| Low activity despite UB overflow at larger tiles       | Mutually exclusive DMA paths may remain live; see [DMA splitting](TRAPS.md#dma-paths-and-optional-pointers). |\n| Many host launches for grid chunks                     | Consider a 1D task loop with the appropriate Cube or Vector core count.                                      |\n\n## UB planning\n\nEstimate peak live storage from simultaneous buffers, their dtype, and any compiler-created copies. Use `fla.utils.ascend_ub_manager` for device capacity and tiling rather than hard-coding the UB size. The helper's safety margin reserves capacity; it is separate from the multiplier that estimates live buffers.\n\nIf more capacity is available, test a larger tile and verify compilation and numerics. If the live set already fills the budget, improve reuse or split stages. Model forward and backward separately; inline gate computation can need more UB than a precomputed-gate path.\n\nUseful code locations:\n\n- `fla/utils/ascend_ub_manager.py`: tiling and grid helpers.\n- `fla/utils/hardware.py`: `get_multiprocessor_count`, including `use_aicore=True` for Cube work.\n- `fla/modules/norm/triton_ascend/l2norm.py`: entry-point registration and the Ascend implementation.\n- `tests/ops/test_gdn_kernels.py`, `tests/ops/test_solve_tril.py`, and `tests/modules/test_conv.py`: relevant kernel comparisons.\n","isInternal":false,"tokens":1051,"sizeBytes":5785},{"name":"TRAPS.md","path":".agents/skills/fla-ascend-performance/references/TRAPS.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/.agents/skills/fla-ascend-performance/references/TRAPS.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# Ascend compiler and memory traps\n\nRead the relevant section before changing DMA, launch indexing, or reused matrix operands. [Kernel cases](cases.md) show these constraints in existing implementations.\n\n## DMA paths and optional pointers\n\nA runtime branch between block-pointer DMA and masked loads can keep both paths live in UB. Larger tiles may then overflow even when measured UB bandwidth is low. Where the workload permits, split bulk and tail launches with a constexpr mode so each compilation removes the unused path. Include convolution halos when deciding whether a bulk load fits inside the packed allocation.\n\nAn optional pointer needs a compile-time guard before any runtime condition. `if USE_INITIAL_STATE or runtime_condition:` can still compile pointer arithmetic in the other branch when the pointer is `None`; nest the constexpr condition separately.\n\n## Address arithmetic\n\nCast indices to `tl.int64` before multiplying by strides or dimensions. Casting the product cannot repair overflow. On Ascend, specialized arguments and folded program IDs may be constexpr values without a `.to()` method; `tl.cast` works for both runtime and constexpr integers.\n\n```python\ntoken_offset = tl.cast(i_t, tl.int64) * BT\nbatch_offset = tl.cast(i_b, tl.int64) * T\nbos = tl.load(cu_seqlens + i_n).to(tl.int64)\n```\n\nLoad packed sequence offsets as int64 even when the input `cu_seqlens` uses int32. A small token index can overflow after multiplication by head count and feature size.\n\n`make_block_ptr` metadata has a separate constraint: its `offsets` and `block_shape` must remain int32. Keep large flattened base-address calculations in int64, then supply valid local block offsets. The Ascend backend is exempt from the mainline ban on block pointers.\n\n## Grid and task loops\n\nThe existing Ascend grid helpers enforce a grid-product limit of 65535. Use their host splitting or flatten independent tiles into a 1D loop over available cores. `get_multiprocessor_count()` selects Vector cores on NPU; `use_aicore=True` selects Cube cores.\n\nAfter slicing variable-length chunk metadata, apply the corresponding global offset only once. In a task loop, derive local pointers from the original base on each iteration; accumulating pointer updates across tasks can miscompile on Ascend.\n\n## Reused `tl.dot` operands\n\nTriton-Ascend can reuse the left operand's UB storage during `tl.dot`. A later read of that tile can therefore be wrong even without a compiler error. Audit all later uses, including another left operand, a right operand, arithmetic, and stores.\n\nReload the pristine tile from global memory when reuse is separated by stages. For tight reuse, create the needed `tile + 0.0` copies before the first dot that consumes the original as a left operand. Copying afterwards preserves the corrupted value. Account for the extra live copies in the UB budget and validate against the numerical reference.\n\n[The case index](cases.md#tldot-lhs-clobber--repo-wide-case-catalog) identifies the affected kernel patterns. A fresh load on each loop iteration or a tile used as a left operand only once needs no extra copy for this issue.\n\n## Numerical and tail behavior\n\nPreserve the validated accumulation precision and exponential base. Mask invalid values before exponentials and initialize every region consumed later. Factoring a gate difference into a ratio or reciprocal changes rounding and range behavior; check the existing reference and tolerances before using it.\n\nExercise gated and ungated paths when both are supported. A suite that always supplies a gate cannot catch `None` pointer handling or incorrect ungated residual updates.\n","isInternal":false,"tokens":746,"sizeBytes":3649},{"name":"SKILL.md","path":".agents/skills/fla-ascend-performance/SKILL.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/.agents/skills/fla-ascend-performance/SKILL.md","title":"fla-ascend-performance","category":"anthropic-skill","format":"markdown","content":"---\nname: fla-ascend-performance\ndescription: Profile and optimize FLA Triton-Ascend kernels using NPU traces, with guidance for UB capacity, memory movement, launch limits, and numerical correctness.\n---\n\n# Ascend kernel performance\n\nUse this skill for operator performance work in `triton_ascend.py` and `triton_ascend/` backend implementations. Follow [fla-optimization-loop](../fla-optimization-loop/SKILL.md) for sustained optimization and [CONTRIBUTING](../../../CONTRIBUTING.md#benchmarking) for validation and reporting.\n\nKeep collection, analysis, and benchmarks in the active Python/NPU environment. If it is not configured, activate the project's Ascend environment first.\n\n## 1. Establish the baseline\n\nLocate the public entry, dispatch route, and Ascend kernel. Record the supported inputs and production workloads, baseline commit, and synchronized latency. Preserve the public API, validated algorithm, reference, and tolerances. Confirm the intended NPU kernel runs; a Torch fallback must not hide an unsupported path or kernel bug.\n\n## 2. Collect a trace\n\nUse the existing collector instead of duplicating profiler setup. Run from the repository root; the workload file initializes its inputs and defines a repeatable `workload()` callable.\n\n```bash\npython .agents/skills/fla-ascend-performance/scripts/profile_npu.py \\\n  --name my_op --out-dir profile/my_op-npu \\\n  --metrics PipeUtilization --analyze \\\n  --kernel-filter my_kernel --exec-file path/to/workload.py\n```\n\nOne run collects one `aic_metrics` set. Start with `PipeUtilization`; collect `MemoryUB` separately when UB bandwidth is relevant. The default schedule is one warmup step and one active step. Warm the workload enough to exclude compilation from the measurement.\n\nThe collector prints the exact trace directory. To revisit it, pass that directory to [the analyzer](scripts/analyze_profile.py); see [collection and metrics](references/reference.md#collection-and-analysis) for the library interface and output layout.\n\n## 3. Diagnose bottlenecks\n\nUse `op_statistic.csv` to find the dominant operation, then `kernel_details.csv` to inspect its duration and pipe activity. Confirm the dispatch route before tuning a small or missing target kernel. Distinguish compilation failure, UB overflow, grid limits, numerical errors, and measured performance limits; they need different fixes.\n\nUse the [diagnosis table](references/reference.md#diagnosis) to choose the next experiment. Change one relevant factor, then measure whether the expected duration and pipe metrics moved.\n\n## 4. Optimize within the kernel contract\n\n- Estimate the peak live UB footprint and use `fla.utils.ascend_ub_manager` for device capacity and tiling. Model forward and backward separately; fuse only when the live set fits.\n- Match the launch to the work: Cube kernels use Cube cores, Vector kernels use Vector cores. Handle the grid limit through the existing helpers or a 1D task loop.\n- Preserve numerical precision, optional inputs, and fixed/variable-length behavior. Verify tail handling and pointer arithmetic before increasing tiles or fusing stages.\n- Omit `num_warps` and `num_stages` from Ascend launches and autotune settings; these are unsupported on this backend.\n- Check [compiler and memory traps](references/TRAPS.md) when changing DMA paths, pointer arithmetic, or reused `tl.dot` operands. Use [contiguous gate loading](references/g-contiguous-loading.md) when gate loads along time are strided.\n\n[Kernel cases](references/cases.md) explain existing convolution, recurrence, and solve implementations. Their tile choices are workload-specific examples.\n\n## 5. Validate and measure\n\nRun kernel comparisons and the affected operator/layer tests, including supported gradients, variable lengths, optional inputs, and boundary shapes. Keep NaN poisoning enabled and confirm that dispatch exercises the changed path. Run the full relevant correctness gate before reporting a speedup; a `--gate-k` subset is an iteration check.\n\nCompare synchronized forward and forward/backward latency on the same NPU, then repeat the relevant profiler collection. Report the timing change, the metrics that explain it, any compiler limitation encountered, and untested cases. If the expected metrics do not change, revisit the diagnosis before adding another optimization.\n","frontmatter":{"name":"fla-ascend-performance","description":"Profile and optimize FLA Triton-Ascend kernels using NPU traces, with guidance for UB capacity, memory movement, launch limits, and numerical correctness."},"isInternal":false,"tokens":892,"sizeBytes":4329},{"name":"cp.md","path":".agents/skills/fla-correctness-coverage/references/cp.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/.agents/skills/fla-correctness-coverage/references/cp.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"../../../../fla/ops/cp/README.md","isInternal":false,"tokens":10,"sizeBytes":32},{"name":"delta-rule.md","path":".agents/skills/fla-correctness-coverage/references/delta-rule.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/.agents/skills/fla-correctness-coverage/references/delta-rule.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"../../../../fla/ops/delta_rule/README.md","isInternal":false,"tokens":11,"sizeBytes":40},{"name":"generalized-delta-rule.md","path":".agents/skills/fla-correctness-coverage/references/generalized-delta-rule.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/.agents/skills/fla-correctness-coverage/references/generalized-delta-rule.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"../../../../fla/ops/generalized_delta_rule/README.md","isInternal":false,"tokens":12,"sizeBytes":52},{"name":"simple-gla.md","path":".agents/skills/fla-correctness-coverage/references/simple-gla.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/.agents/skills/fla-correctness-coverage/references/simple-gla.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"../../../../fla/ops/simple_gla/README.md","isInternal":false,"tokens":11,"sizeBytes":40},{"name":"SKILL.md","path":".agents/skills/fla-correctness-coverage/SKILL.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/.agents/skills/fla-correctness-coverage/SKILL.md","title":"fla-correctness-coverage","category":"anthropic-skill","format":"markdown","content":"---\nname: fla-correctness-coverage\ndescription: Select and run correctness coverage for FLA kernels and modules, including gradients, dispatch boundaries, and Triton addressing changes.\n---\n\n# FLA Correctness Coverage\n\nUse [fla-design-coverage](../fla-design-coverage/SKILL.md) when the change needs a new numerical or routing contract. This skill turns that contract into tests; [CONTRIBUTING.md](../../../CONTRIBUTING.md#testing) defines references, tolerances, test structure, and platform conventions.\n\n## Select coverage\n\nRead the affected tests and callers before adding cases. Follow [test organization and naming](../../../CONTRIBUTING.md#test-organization-and-naming): try the existing parameter matrix first, and add a function or file only when its setup or scope requires one. Cover reachable gaps across the applicable dimensions together, rather than testing each flag only in isolation:\n\n| Dimension   | Cases that can expose different behavior                                         |\n| ----------- | -------------------------------------------------------------------------------- |\n| Layout      | Dense and varlen, partial chunks, long sequences, asymmetric QK/value dimensions |\n| Features    | Gate bounds, raw/post-sigmoid beta, QK normalization, grouped value attention    |\n| State       | Initial state, final state, and state gradients where supported                  |\n| Numerics    | Supported dtypes, affected conversion/reduction paths, numerical boundaries      |\n| Routing     | Default and selected backends, verifier acceptance, rejection, and fallback      |\n\nCompare outputs and final states against the existing reference. Training APIs also need every supported gradient; forward-only APIs must state that backward is unsupported. Use `torch.autograd.gradcheck` where the implementation supports its required precision. Routing tests verify selection and argument forwarding; they do not replace numerical comparisons on the actual backend.\n\nFor addressing changes, inspect casts before multiplication and test shapes that expose large offsets, partial tiles, and varlen boundaries. Follow the `tl.int64` and platform requirements in [Triton Kernels](../../../CONTRIBUTING.md#triton-kernels). Exercise changed grid/address paths on affected supported platforms. Distinguish unsupported cases from tests skipped for unavailable hardware or dependencies; record the latter as validation gaps.\n\nThe NaN allocation guard covers eligible FLA allocations in operator/module tests, including the single-rank CP graph and intra-card cache cases moved from operator tests, but not layer/model tests, spawned CP workers, or every allocation. Add explicit finite-output and finite-gradient checks to adversarial cases, and investigate a poisoned-run failure even when an ad hoc run passes.\n\n## Run the affected paths\n\nDiscover dependent tests from the repository root:\n\n```bash\npython scripts/find_dependent_tests.py fla/ops/kda/chunk.py\nFLA_CI_ENV=0 pytest tests/ops/test_kda.py -v\n```\n\nThe helper lists files; it does not run them. Include reported callers and relevant `tests/modules/`, `tests/layers/`, `tests/models/`, or `tests/cp/` coverage when the change reaches those paths. Inspect skips and warning-only comparisons before calling the run a pass. Record commands, backend flags, environment, and results; reproduce failures on the unchanged baseline before labeling them pre-existing.\n\nRead operator math only when needed: [Delta Rule](../../../fla/ops/delta_rule/README.md), [Generalized Delta Rule](../../../fla/ops/generalized_delta_rule/README.md), [Simple GLA](../../../fla/ops/simple_gla/README.md), or [context parallelism](../../../fla/ops/cp/README.md).\n","frontmatter":{"name":"fla-correctness-coverage","description":"Select and run correctness coverage for FLA kernels and modules, including gradients, dispatch boundaries, and Triton addressing changes."},"isInternal":false,"tokens":732,"sizeBytes":3704},{"name":"SKILL.md","path":".agents/skills/fla-design-coverage/SKILL.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/.agents/skills/fla-design-coverage/SKILL.md","title":"fla-design-coverage","category":"anthropic-skill","format":"markdown","content":"---\nname: fla-design-coverage\ndescription: Define supported inputs, numerical budgets, routing, and validation before designing an FLA kernel or numerical change.\n---\n\n# FLA Design and Coverage\n\nUse this skill to establish the contract before implementing a kernel, backend route, or numerical change. Read [CONTRIBUTING.md](../../../CONTRIBUTING.md) for contribution rules, especially [public compatibility](../../../CONTRIBUTING.md#public-api-compatibility) and [design/RFC requirements](../../../CONTRIBUTING.md#submit-pull-requests).\n\n## Define the supported domain\n\nTrace public entries and defaults, backend verifiers, layer/model/CP callers, and existing tests. Record the merge-base SHA and the source of each supported case. Group cases only when they share the same numerical assumptions and execution route; omit combinations no supported caller can reach.\n\nDistinguish routes when kernel stages, chunk geometry or anchors, capability checks, dtype conversions, or input value ranges change. A backend here means an implementation selected by dispatch, including the default implementation, rather than a hardware product. For each group, identify its effective arguments, oracle, output/state and per-gradient tolerances, and test coverage. Assign one outcome:\n\n- The proposed implementation, with passing reference comparisons.\n- An existing fallback, with route-parity tests covering outputs, states, dtypes, and supported gradients.\n- A documented unsupported-input error raised before kernel execution and asserted by a test.\n\nPreserve previously supported cases under the public compatibility policy. A warning, partially initialized output, or fallback that changes semantics is not a valid outcome. When changing defaults, record the old and new effective values and selected routes.\n\n## Establish numerical safety\n\nFor each affected stage in forward and backward, record input-range assumptions, load/operand/accumulation/store dtypes, reduction behavior, overflow/underflow headroom, and error propagation to the end-to-end tolerance. Forward agreement alone does not establish backward safety, including inverse recomputation and state gradients.\n\nSafety flags and bounds apply to the stages and domains that justify them. Reassess them when chunk geometry, anchors, dtype staging, or input scaling changes; a safe exponent does not prove that a later inversion is safe. Document the algebra behind numerical acceptance thresholds, implement each shared predicate once, and test its acceptance boundary and the adjacent rejection or fallback. Enforce additional assumptions, such as input-norm bounds, at the public entry or every affected verifier; report the failed predicate and effective values.\n\nKeep the committed oracle, input distributions, assertions, and tolerances fixed while evaluating the candidate. A correction to that contract belongs in a separate change with before/after error distributions for the unchanged baseline and candidate; precision, tolerance, and algorithm changes follow the linked RFC policy. Keep the corrected gate green before the candidate lands. Every selectable autotune configuration must satisfy the same numerical contract.\n\n## Check routing\n\n- Share semantic argument normalization between the public entry, verifier, and executor. Test that omitted arguments match their documented explicit defaults in route, outputs, and gradients. `input_guard` handles contiguity and device context, not semantic defaults.\n- Use capability and availability helpers from `fla.utils`. Distinguish backend availability, policy enablement, and acceptance of an individual call; use [fla-dispatch-backends](../fla-dispatch-backends/SKILL.md) for implementation mechanics.\n- For context-parallel routes, every rank must make the same selection. If verifier inputs can differ by rank, combine acceptance with a collective and select the optimized route only if all ranks accept. Record whether forward commits backward to the same backend.\n- Test parity between routes claiming the same domain. Record precision settings such as `TRITON_F32_DEFAULT` and matmul flags, keep them identical across baseline, candidate, and tests, and restore global settings after tests.\n\n## Choose evidence\n\nApply the relevant axes from [fla-correctness-coverage](../fla-correctness-coverage/SKILL.md) to each supported group across these distinct layers:\n\n- **Production:** reproduce effective arguments and numerical input distributions from checked-in layer/model callers. A generic benchmark shape list alone does not establish production coverage.\n- **Public boundaries:** cover previously supported inputs, changed defaults, and verifier acceptance/rejection boundaries.\n- **Adversarial:** outside the contract, require documented rejection before kernel execution or explicit finite-output and finite-gradient checks. This does not expand the supported domain.\n\nBefore performance work, record baseline correctness, latency or throughput, commit, environment, workload, input distribution, and seed. Benchmark production cases first, then boundaries whose routing or expected performance changes. Unaffected boundaries still need correctness coverage. Use [fla-optimization-loop](../fla-optimization-loop/SKILL.md) to freeze the gate and compare candidates.\n","frontmatter":{"name":"fla-design-coverage","description":"Define supported inputs, numerical budgets, routing, and validation before designing an FLA kernel or numerical change."},"isInternal":false,"tokens":981,"sizeBytes":5293},{"name":"SKILL.md","path":".agents/skills/fla-dispatch-backends/SKILL.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/.agents/skills/fla-dispatch-backends/SKILL.md","title":"fla-dispatch-backends","category":"anthropic-skill","format":"markdown","content":"---\nname: fla-dispatch-backends\ndescription: Change FLA backend registration, dispatch, or verifiers while preserving routing and call contracts.\n---\n\n# Backend dispatch\n\nUse this skill for registration, selection, or verifier changes in `fla/backends.py` and operation backends. [CONTRIBUTING.md](../../../CONTRIBUTING.md) defines style and public compatibility; [fla/backends.py](../../../fla/backends.py) defines the routing behavior.\n\n## 1. Define the supported calls\n\nIdentify the default entry point and the calls the backend can handle: inputs, layouts, dtypes, gradient modes, and optional dependencies. Match the entry point's signature, parameter order, defaults, return structure, and supported gradients.\n\nA verifier returns `(True, None)` or `(False, reason)`. It checks support without changing inputs or global state. Keep it beside the implementation; put helpers shared by several implementations in a utility module.\n\n## 2. Register the implementation\n\nKeep the default entry point backend-independent with bare `@dispatch`. Choose the registration form that matches the implementation:\n\n| Implementation               | Registration                                                         |\n| ---------------------------- | -------------------------------------------------------------------- |\n| Function with the same API   | `@register(entry, backend=TritonAscendBackend, verifier=...)`        |\n| Adapter with backend policy  | `@register(entry)` on a `BaseBackend` subclass                       |\n\nAn adapter defines methods named after the entry point, such as `chunk_kda` and `chunk_kda_verifier`. Register it where it is defined and import optional kernels only when selected.\n\nImport the owning package's backends after its default entry points are defined. For implementations with platform-specific dependencies, guard the import with the shared backend's `is_available()`. Do not gate registration on enable switches; dispatch reads those at call time. Keep per-operation switches on `@register(..., env_var=...)` and implementation paths out of the shared dispatcher.\n\nExisting owners using `@dispatch('<operation>')` use `@register('<operation>')` on adapter classes and load implementations through the owner's backend package. Both forms use the same registration and selection logic.\n\n## 3. Preserve routing behavior\n\nFor each call, try backends in ascending priority order; ties retain registration order:\n\n1. Check availability and enablement.\n2. Run the verifier; rejection proceeds to the next candidate.\n3. Call the selected implementation. Its exceptions propagate.\n4. If no candidate accepts, call the default implementation.\n\nRegistration replaces an existing backend of the same type. Function registration keys the registry by the unwrapped entry point, so compiler decorators and public aliases share that registry. Legacy owner registries must load their backend package before lookup.\n\nKeep selection outside `torch.compile` graphs. Do not cache availability or enablement in the dispatch wrapper: registered backends must remain eligible when their switches change. Python's import cache handles modules already loaded, and unused optional dependencies must not prevent package import.\n\n## 4. Verify through the entry point\n\nRun the decorated entry point for accepted calls, verifier rejection, disabled/unavailable backends, and fallback. Check optional arguments omitted as well as supplied, changed support boundaries, outputs, and supported gradients.\n\nUse explicit backend flags in routing tests. For the default path, set `FLA_DISABLE_BACKEND_DISPATCH=1` before importing FLA. Use fresh processes for import-order checks, and verify public aliases and shared backend ownership after registration changes.\n\nFind affected consumers with `scripts/find_dependent_tests.py` and run the relevant operator, layer, and model tests. Host routing tests establish dispatch behavior; accelerator reference tests establish numerical correctness.\n","frontmatter":{"name":"fla-dispatch-backends","description":"Change FLA backend registration, dispatch, or verifiers while preserving routing and call contracts."},"isInternal":false,"tokens":741,"sizeBytes":3990},{"name":"SKILL.md","path":".agents/skills/fla-kda/SKILL.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/.agents/skills/fla-kda/SKILL.md","title":"fla-kda","category":"anthropic-skill","format":"markdown","content":"---\nname: fla-kda\ndescription: Modify or review KDA gates, chunk and recurrent kernels, backend routes, and their correctness coverage.\n---\n\n# KDA\n\nUse this skill for `fla/ops/kda/` and its tests. Public entry points are `chunk_kda` and `fused_recurrent_kda`; inspect the affected entry point's signature before changing its implementation.\n\n## Code map\n\n- `gate.py` defines gate activations and cumulative decay.\n- `chunk_fwd.py` and `chunk_bwd.py` coordinate training; `chunk_intra.py`, `chunk_intra_token_parallel.py`, and `wy_fast.py` implement the main stages.\n- `backends/` contains alternative implementations and verifiers. Follow [backend dispatch](../fla-dispatch-backends/SKILL.md) when changing routing.\n- `tests/ops/test_kda.py` contains the reference comparisons and supported gate-mode cases.\n\n## Gate contracts\n\n`chunk_kda` takes log-space decay when `use_gate_in_kernel=False`. With it enabled, the caller supplies raw gates and optional bias:\n\n- Without a lower bound, activation is `-exp(A_log) * softplus(g + dt_bias)` and `A_log` is required.\n- With a lower bound, activation is `lower_bound * sigmoid(exp(A_log) * (g + dt_bias))`. Omitting `A_log` uses a multiplier of one.\n\n`safe_gate` selects the intra-chunk implementation independently of gate activation. Safe in-kernel activation requires `-5 <= lower_bound < 0`. Pre-gated safe mode requires bounded log-space gates supplied by the caller; tests cover `[-5, 0]`, but the wrapper does not validate tensor values.\n\n## Numerical invariants\n\nThe safe intra path uses 16-token diagonal blocks and midpoint offsets before exponentiation. At a lower bound of `-5`, a block can accumulate `-80` in natural-log units; local offsets limit each exponent's span to about half that range. Preserve these offsets rather than exponentiating the full cumulative gate. Inter-block decay uses paired differences whose exponents remain non-positive for monotonic decay.\n\nThe non-safe path uses token-parallel diagonal computation. Both paths share inter-block and triangular-solve work, so changes there require coverage of both modes. Keep precision and tolerance changes within the design process in [AGENTS.md](../../../AGENTS.md#scope-and-direction).\n\n## Validation\n\nUse [correctness coverage](../fla-correctness-coverage/SKILL.md) for the test strategy. Select the KDA cases affected by the change:\n\n- Dense and variable-length sequences; forward, final state, and supported gradients.\n- Pre-gated and in-kernel gates, safe and non-safe paths, and omitted optional arguments.\n- Beta activation and query/key normalization modes when touched.\n- Grouped value heads and differing key/value dimensions when indexing or shapes change.\n- Initial states, intermediate states, and context parallelism when their paths are touched.\n- Accepted and rejected backend routes when changing verifiers.\n\nFor gate or decay changes, include the lower-bound limits, saturated gate inputs, extreme gate scales, long cumulative decay, and chunk or ragged-sequence boundaries. Test through the public entry point so the selected route and effective arguments are exercised.\n","frontmatter":{"name":"fla-kda","description":"Modify or review KDA gates, chunk and recurrent kernels, backend routes, and their correctness coverage."},"isInternal":false,"tokens":675,"sizeBytes":3118},{"name":"SKILL.md","path":".agents/skills/fla-nvidia-performance/SKILL.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/.agents/skills/fla-nvidia-performance/SKILL.md","title":"fla-nvidia-performance","category":"anthropic-skill","format":"markdown","content":"---\nname: fla-nvidia-performance\ndescription: Profile and optimize FLA kernels on NVIDIA GPUs, with same-hardware benchmarks and targeted Nsight Compute analysis.\n---\n\n# NVIDIA kernel performance\n\nUse this skill for NVIDIA kernel optimization and backend tuning. Follow [CONTRIBUTING](../../../CONTRIBUTING.md#benchmarking) for correctness and performance evidence, and [fla-optimization-loop](../fla-optimization-loop/SKILL.md) for sustained iteration.\n\n## Hardware baseline\n\nUse datacenter NVIDIA GPUs with sm_90 or newer for PR performance conclusions; H100 and H20 are accepted, and sm_100/sm_103 are preferred. Results from A100, older GPUs, or consumer cards are supplementary. Record the exact GPU and software versions with each comparison.\n\n## Workflow\n\n1. Identify the changed kernel, its callers, and the production workloads it affects. Record a baseline commit and confirm which backend runs.\n2. Pass the relevant correctness tests before timing a candidate. Preserve the reference and tolerances; include outputs, states, and supported gradients.\n3. Compare baseline and candidate on the same GPU with identical workloads and dispatch settings. Include dense and supported variable-length cases. Use an end-to-end benchmark for a training or generation throughput claim.\n4. Use Nsight Compute when a bottleneck or regression needs hardware metrics. Match the optimization to the measured limit, then rerun correctness and timing.\n5. Report latency or throughput, relevant memory changes, and any regressions. Keep raw profiler artifacts outside git, for example under `profile/<run_name>/`.\n\n## Benchmark commands\n\nRun from the repository root. Resolve the baseline to a commit SHA before comparing:\n\n```bash\nFLA_BENCH_BASE=$(git rev-parse origin/main)\nFLA_CI_ENV=0 python -m benchmarks.ops.verify --op chunk_kda --base \"$FLA_BENCH_BASE\"\n\npython benchmarks/benchmark_training_throughput.py --name kda --batch_size 2 --seq_len 8192\npython benchmarks/benchmark_training_throughput.py --name kda --batch_size 2 --seq_len 8192 --varlen\n```\n\nThe verification command runs the current tests; retain the validated baseline and tests yourself. A selected subset is useful during iteration but does not replace the full relevant gate before reporting a gain.\n\n## Nsight Compute\n\nCollect a representative changed kernel with the installed `ncu`. Check which sections and metrics that version supports.\n\nThe benchmark runner launches a child process, so use `--target-processes all`. Replace the kernel filter with the changed kernel's name. `--no-base` profiles the current implementation; collect the baseline separately with the same command and workload.\n\n```bash\nmkdir -p profile/kda\nncu --target-processes all --set full \\\n  -k 'regex:chunk_kda' -c 1 -o profile/kda/full \\\n  python -m benchmarks.ops.run --op chunk_kda --modes fwd --no-base\n\nncu --target-processes all --set source --section SourceCounters \\\n  -k 'regex:chunk_kda' -c 1 -o profile/kda/source \\\n  python -m benchmarks.ops.run --op chunk_kda --modes fwd --no-base\n```\n\nUse duration, memory throughput, occupancy, stalls, and hot instructions to explain the bottleneck. Profiling overhead makes these runs unsuitable as the latency comparison; use the synchronized benchmark for that measurement.\n","frontmatter":{"name":"fla-nvidia-performance","description":"Profile and optimize FLA kernels on NVIDIA GPUs, with same-hardware benchmarks and targeted Nsight Compute analysis."},"isInternal":false,"tokens":708,"sizeBytes":3270},{"name":"opt-log-template.md","path":".agents/skills/fla-optimization-loop/references/opt-log-template.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/.agents/skills/fla-optimization-loop/references/opt-log-template.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# Optimization Record\n\nKeep one local log under ignored `profile/<op>-opt/`. Record the entry point, target workloads, baseline SHA and correctness, environment, commands, numeric settings, and success criterion once. Preserve enough candidate identity to reproduce each measurement.\n\n| Attempt  | Change or hypothesis | Gate | Latency by workload | Decision and reason |\n| -------- | -------------------- | ---- | ------------------- | ------------------- |\n| Baseline | Unchanged revision   | Pass | Measured baseline   | Comparison point    |\n\nAppend an entry after each attempt, including failures and unchanged results. Identify whether the gate was full or selected; record no performance result for a failed gate. Retained candidates need their revision or recoverable diff, the reason they helped, and any unresolved limitation. Keep failed directions here so later work does not repeat them.\n\nWhen specializing, add each routing condition, entry point, baseline/candidate latency, and measured reason for a separate path. Compare every bucket against the same baseline revision on that bucket's workloads, including fallback boundaries. Keep these details alongside the iteration results.\n","isInternal":false,"tokens":220,"sizeBytes":1198},{"name":"TRAPS.md","path":".agents/skills/fla-optimization-loop/references/TRAPS.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/.agents/skills/fla-optimization-loop/references/TRAPS.md","title":"References Skill","category":"anthropic-skill","format":"markdown","content":"# Correctness and Measurement Traps\n\n## Partial writes\n\n`tests/conftest.py` poisons eligible FLA `empty` allocations in operator/module tests with NaNs. It excludes some allocations and does not cover layer/model/CP tests. A poisoned-run failure that disappears in an ad hoc script can reveal uninitialized reads; inspect masks and ensure every exposed output element is written. NaNs can also have numerical causes, so diagnose the failure rather than assuming its source.\n\n## Reference precision and tolerance\n\n`fla.utils.assert_close` checks relative RMS error with an absolute-error shortcut, not elementwise `rtol`. `FLA_CI_ENV=1` and explicit `warning=True` can downgrade tolerance failures to warnings. Use strict local validation and inspect warnings; keep the committed thresholds fixed.\n\nTF32 in a reference matmul can change an fp32 comparison. Inspect the reference's precision settings before blaming the kernel, but do not change the frozen oracle to pass. Baseline and candidate need identical precision and numeric flags. An intentional precision change follows the numerical-change process in [CONTRIBUTING.md](../../../../CONTRIBUTING.md#submit-pull-requests).\n\n## Address overflow\n\nCast program IDs and offsets to `tl.int64` before multiplication; widening an already-overflowed product cannot repair it. Small shapes may hide the error. Run large-offset and affected varlen/grid cases under the [Triton addressing rules](../../../../CONTRIBUTING.md#triton-kernels).\n\n## Missing work in a fast result\n\nCheck unexpectedly large speedups against the actual selected path, executed tests, and measured shapes. `verify.py` accepts pytest's exit status, including an all-skipped success. The benchmark runner filters constrained shapes and omits failed warmups or measurements. Compare expected and returned rows before making a claim; a passing process alone is insufficient.\n\n## Autotune and clocks\n\n`run.py`, also used by `verify.py`, warms all shapes before timing. Preserve that warmup. After changing the configuration space, check cached selections and rewarm or clear the relevant cache if results look stale.\n\nUnlocked clocks and changing device load can shift latency between runs. Rank local candidates by their own runtime, then remeasure the final candidate against the recorded baseline in one session. Report variability and the device state; speedups from separate experiments do not add up to a final result.\n","isInternal":false,"tokens":487,"sizeBytes":2440},{"name":"SKILL.md","path":".agents/skills/fla-optimization-loop/SKILL.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/.agents/skills/fla-optimization-loop/SKILL.md","title":"fla-optimization-loop","category":"anthropic-skill","format":"markdown","content":"---\nname: fla-optimization-loop\ndescription: Iterate on FLA kernel performance with a frozen correctness gate, a measured baseline, and reproducible candidate comparisons.\n---\n\n# FLA Optimization Loop\n\nUse this skill for repeated kernel performance work. Define supported workloads and numerical assumptions with [fla-design-coverage](../fla-design-coverage/SKILL.md); use the [NVIDIA](../fla-nvidia-performance/SKILL.md) or [Ascend](../fla-ascend-performance/SKILL.md) skill for profiling on that platform.\n\n## Freeze the gate and baseline\n\nAdd missing regression or new-path coverage before starting the loop. Then keep the relevant tests, reference implementation, input distributions, assertions, tolerances, and public signature fixed. Do not narrow parameter matrices, add skips, relax precision or thresholds, special-case test values, cache results across trials, or change timing behavior to make a candidate pass. Fully initialize exposed outputs. Use kernel implementations for the optimization; replacing the whole operator with a vendor call solely to win the benchmark does not meet this workflow.\n\nA suspected oracle or tolerance defect needs a separately justified correction, not an adjustment that makes the candidate pass. Follow [CONTRIBUTING.md](../../../CONTRIBUTING.md#submit-pull-requests) for precision, numerical-algorithm, and tolerance changes, and its [compatibility policy](../../../CONTRIBUTING.md#public-api-compatibility) for public behavior.\n\nRecord the entry point, target workloads, allowed implementation languages, validation/benchmark commands, baseline SHA, and success criterion in the local optimization log. Include production inputs from callers as well as applicable benchmark-registry shapes. Confirm the baseline passes the full gate and measure it before changing kernels. For a new kernel, establish correctness before tuning.\n\n```bash\nFLA_BENCH_BASE=$(git rev-parse origin/main)\nFLA_CI_ENV=0 python -m benchmarks.ops.verify --op chunk_gla --base \"$FLA_BENCH_BASE\"\n```\n\nUse a recorded SHA: the baseline worktree cannot reuse a checked-out branch, and slash-containing refs break its temporary directory naming. `--base` benchmarks that revision; it does not run its correctness tests. Establish baseline correctness separately if the current tree already differs.\n\nThe driver runs the current pytest file and stops before timing when pytest fails. It neither freezes files nor rejects an all-skipped run. Check that the expected cases and benchmark rows ran. `--test-file` overrides a mismatched derived test path; `--gate-k` selects a quick subset, while `--no-gate` produces unverified measurements. Neither subset nor skipped-gate results qualify for promotion.\n\n## Iterate from evidence\n\n1. Profile the current bottleneck and choose one reasoned change.\n2. Run the frozen gate, then benchmark only a passing candidate. A subset can provide an early signal; the final gate covers the full affected suite.\n3. Record the change, gate result, measurement, and keep/drop decision before the next experiment. Preserve a recoverable candidate under the repository's git rules and the user's authorization; failed and unchanged attempts need records, not empty commits.\n\nRead [measurement and correctness traps](references/TRAPS.md) before interpreting results. Keep seeds, numeric flags, hardware, and workload fixed. Candidate runtime helps rank local iterations; measure the final candidate directly against the frozen baseline in one session rather than adding speedups from separate runs.\n\nAdd specialized paths only when profiling shows different bottlenecks across shape regimes and measurements justify the extra routing. Record each condition, chosen path, and before/after result, and validate the full shape set including fallback boundaries.\n\nAfter repeated attempts without an improvement above noise, re-profile and review the failed directions before continuing. Respect the user's stopping condition. A no-win conclusion needs a measured baseline, a reasoned candidate attempt, gate status, benchmark evidence, and an identified bound or blocker. An unavailable profiler or one losing candidate alone does not establish a performance floor; report unresolved validation as a limitation.\n\n## Keep and report the result\n\nKeep iteration notes and raw traces under ignored `profile/<op>-opt/`; [opt-log-template.md](references/opt-log-template.md) provides a compact record format. This search history stays out of kernel comments and the final PR description.\n\nPromote only with a full green gate, a repeatable win on the claimed workloads, and profiling or a quantified bottleneck analysis that explains it. Report per-shape results and regressions, commands, baseline/candidate revisions, environment, and measurement variability using the relevant hardware skill. An unexplained large speedup needs investigation. Use [fla-pr-readiness](../fla-pr-readiness/SKILL.md) to prepare the final diff and evidence.\n","frontmatter":{"name":"fla-optimization-loop","description":"Iterate on FLA kernel performance with a frozen correctness gate, a measured baseline, and reproducible candidate comparisons."},"isInternal":false,"tokens":966,"sizeBytes":4968},{"name":"SKILL.md","path":".agents/skills/fla-pr-readiness/SKILL.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/.agents/skills/fla-pr-readiness/SKILL.md","title":"fla-pr-readiness","category":"anthropic-skill","format":"markdown","content":"---\nname: fla-pr-readiness\ndescription: Prepare or update an FLA pull request with focused scope, verified evidence, and the current repository template.\n---\n\n# FLA PR Readiness\n\nRead [CONTRIBUTING.md](../../../CONTRIBUTING.md) and [AGENTS.md](../../../AGENTS.md) first. Use the current [PR template](../../../.github/pull_request_template.md) for headings and checkboxes; contribution policy and template wording are not duplicated here.\n\n## Review scope and evidence\n\nRead the complete diff against the PR base and audit affected shared callers. State its concrete outcome in one sentence and separate independent work. Check existing issues/PRs and the applicable design discussion. Follow [scope and size](../../../CONTRIBUTING.md#pr-scope-and-size) for a cohesive change above the normal limit, and [public compatibility](../../../CONTRIBUTING.md#public-api-compatibility) when interfaces or defaults are affected.\n\nWhen a change needs splitting or dependent PRs need a native GitHub stack, use [PR stacks](../fla-pr-stacks/SKILL.md) to choose boundaries, reuse existing PRs, and verify the published stack.\n\nFor code changes, discover dependent tests and run the relevant targets. For example:\n\n```bash\npython scripts/find_dependent_tests.py fla/ops/kda/chunk.py\nFLA_CI_ENV=0 pytest tests/ops/test_kda.py\npre-commit run --files fla/ops/kda/chunk.py tests/ops/test_kda.py\n```\n\nFor documentation-only changes, validate links, examples, and formatting as described in [Evidence by Change Type](../../../CONTRIBUTING.md#evidence-by-change-type).\n\nThe dependency helper only lists tests. Check that regression coverage exercises the changed path and that required numerical comparisons passed, including supported gradients. Review new test functions and files against [test organization and naming](../../../CONTRIBUTING.md#test-organization-and-naming); use an existing test case where it fits. Record commands, tested revision, environment, and results; distinguish failures, skips, and unrun checks. Reproduce a failure on the baseline before calling it pre-existing.\n\nFor performance-relevant changes, collect same-hardware before/after measurements under [Benchmarking](../../../CONTRIBUTING.md#benchmarking) and the relevant hardware skill. A neutral result needs numbers; explain inapplicable checks. Inspect the final diff for stale comments and unrelated changes, leaving scratch logs and raw profiles outside it.\n\n## Prepare and publish\n\nWrite for a reviewer who has not seen the conversation. Follow [PR Description](../../../CONTRIBUTING.md#pr-description) and [Evidence by Change Type](../../../CONTRIBUTING.md#evidence-by-change-type): explain the problem, resulting behavior, affected users, and evidence. Include only useful rationale or limitations; omit file inventories and the history of attempts. Keep incomplete work in draft and leave claims and checklist items consistent with actual validation.\n\nThe `check-pr-title` workflow also checks the body and size exception on edits. Preserve template checklist wording and distinguish required items from conditional ones.\n\n`gh pr edit` fails with this repository's classic-Projects integration. Write the complete body to a temporary file and update through REST:\n\n```bash\ngh api -X PATCH \"repos/fla-org/flash-linear-attention/pulls/$pr_number\" -F \"body=@$pr_body_file\"\n```\n\nConfirm the published description, head revision, and checks. After review changes, rerun affected validation and update evidence that no longer describes the current code. Publishing and git operations remain subject to the user's authorization and repository rules.\n","frontmatter":{"name":"fla-pr-readiness","description":"Prepare or update an FLA pull request with focused scope, verified evidence, and the current repository template."},"isInternal":false,"tokens":746,"sizeBytes":3614},{"name":"SKILL.md","path":".agents/skills/fla-pr-stacks/SKILL.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/.agents/skills/fla-pr-stacks/SKILL.md","title":"fla-pr-stacks","category":"anthropic-skill","format":"markdown","content":"---\nname: fla-pr-stacks\ndescription: Split oversized FLA pull requests into focused changes, reuse existing PRs, and create or maintain verified GitHub native PR stacks.\n---\n\n# FLA PR Stacks\n\nUse for decomposing large changes or managing dependent PRs. Follow [CONTRIBUTING.md](../../../CONTRIBUTING.md#pr-scope-and-size) for scope and size, [AGENTS.md](../../../AGENTS.md) for git permissions, and [PR readiness](../fla-pr-readiness/SKILL.md) for each PR's evidence and description.\n\n## 1. Choose the split\n\nInspect the complete diff, commit history, and related open PRs, including their review discussions. Map each proposed layer to one outcome, its prerequisites, validation, and an existing PR to reuse where appropriate.\n\n| Relationship                            | Action                                                                   |\n| --------------------------------------- | ------------------------------------------------------------------------ |\n| Overlaps an existing PR's purpose       | Update that PR when ownership permits; preserve attribution and reviews. |\n| Independent of the other changes        | Submit separately against the trunk.                                     |\n| Depends on another layer                | Stack above that layer; include its own tests and documentation.         |\n| Cohesive but still above the size limit | Use the contribution guide's exception with a concrete justification.    |\n\nEach layer presents an increment over its parent and must work with its prerequisites, without needing later layers to fix it. Separate review does not imply dependency-free code. Do not split by arbitrary line counts, leave callers broken, or defer tests to a later PR.\n\nCheck head repository identity before arranging existing PRs. [GitHub stacks require all branches in the same repository](https://docs.github.com/en/pull-requests/reference/stacked-pull-requests); cross-fork stacks are unsupported. Keep incompatible PRs as explicit external prerequisites instead of duplicating them under new PR numbers.\n\n## 2. Build and validate\n\nPreserve the original head and uncommitted work before reconstructing commits in isolated branches or worktrees. Reuse suitable commits and preserve authorship. Check current remote heads before updating existing PRs.\n\nReview each layer against its parent, then compare the top branch with the preserved original result:\n\n```bash\ngit diff main...refactor/foundation\ngit diff refactor/foundation...refactor/consumer\ngit diff original-head refactor/consumer\n```\n\nSubstitute the recorded trunk and actual refs. Account for trunk changes and intentional omissions, including fixes retained in existing PRs. Run relevant tests at each layer's head and on the combined result; passing only the top branch does not validate intermediate PRs.\n\n## 3. Register and verify the native stack\n\nChanging PR bases alone does not create a native stack. Use the official [`gh stack` extension](https://docs.github.com/en/pull-requests/how-tos/create-pull-requests/creating-stacked-pull-requests), GitHub UI, or REST Stacks API. Check installed command help and existing publishing authorization first.\n\nFor prepared branches, adopt them bottom to top and submit:\n\n```bash\ngh stack init --base main refactor/foundation refactor/consumer\ngh stack submit --auto --remote origin\n```\n\nChoose the actual remote containing the branches. `submit` pushes, creates or updates PRs, and registers the stack. `--auto` creates new PRs as drafts; replace generated titles and bodies using PR readiness.\n\nFor existing PRs, use `gh stack link --base main --remote origin 123 124` with their real numbers in dependency order. This reuses PRs without requiring local stack tracking. Branch arguments also push and create missing PRs. After an error, inspect what succeeded before retrying.\n\nFor every PR, verify [remote stack membership](https://docs.github.com/en/pull-requests/reference/stacked-pull-requests-apis-and-webhooks):\n\n```bash\ngh api \"repos/fla-org/flash-linear-attention/pulls/$pr_number\" \\\n  --jq '{number, head_repo: .head.repo.full_name, head: .head.ref, base: .base.ref, stack}'\n```\n\nRequire non-null `stack` objects with the same stack identity, correct order, size, and trunk. Check each PR's incremental diff and report the verified links and remaining external prerequisites. A local branch chain or successful push is insufficient.\n\n## 4. Maintain the series\n\nAfter changing a lower layer, review affected upper diffs and rerun relevant tests. `gh stack rebase` rewrites history; `gh stack sync` also force-pushes with leases. Apply the repository's rewrite permissions and preserve concurrent remote work, then recheck published membership and evidence.\n\nCreating a series does not authorize merging or closing PRs. If merging is requested, establish the intended prefix: selecting an upper PR also merges those below it. Use a stack-aware merge operation; the legacy synchronous merge API does not support stacks.\n","frontmatter":{"name":"fla-pr-stacks","description":"Split oversized FLA pull requests into focused changes, reuse existing PRs, and create or maintain verified GitHub native PR stacks."},"isInternal":false,"tokens":1003,"sizeBytes":4984},{"name":"SKILL.md","path":".agents/skills/fla-triton-to-gluon/SKILL.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/.agents/skills/fla-triton-to-gluon/SKILL.md","title":"fla-triton-to-gluon","category":"anthropic-skill","format":"markdown","content":"---\nname: fla-triton-to-gluon\ndescription: Port an FLA Triton kernel to Gluon when explicit layouts, asynchronous transfers, or scheduling can address a measured bottleneck.\n---\n\n# Triton to Gluon\n\nUse Gluon when profiling points to register pressure, layout conversion, or load/compute overlap that needs explicit control. A kernel already saturating its useful bandwidth or compute limit may not benefit. Keep the portable default path when adding hardware-specific implementations.\n\nFollow [fla-optimization-loop](../fla-optimization-loop/SKILL.md) for the baseline and correctness gate, and [fla-nvidia-performance](../fla-nvidia-performance/SKILL.md) for NVIDIA measurements.\n\n## Check the target environment\n\nGluon is experimental. Check the installed Triton API before copying [upstream examples](https://triton-lang.org/main/getting-started/tutorials/gluon/); their names and signatures can differ from the installed release.\n\nGate architecture-specific code appropriately: Ampere supports `cp.async`, Hopper adds TMA/WGMMA, and Blackwell adds TMEM/tcgen05. Code under Gluon's `nvidia` namespace requires NVIDIA hardware. Keep compile-time choices in explicit constexpr arguments so the kernel, launcher, and autotune pruning use the same values.\n\n```python\nfrom triton.experimental import gluon\nfrom triton.experimental.gluon import language as gl\n```\n\n## Port incrementally\n\n1. Record the baseline commit and retain the validated tests, reference, and tolerances. Keep the Triton fallback available.\n2. Translate loads, stores, and arithmetic with explicit layouts first. Match each tensor's contiguous dimension and preserve masking. Pass forward, state, and supported backward comparisons before adding asynchrony.\n3. Add async transfers only where load latency or addressing pressure limits the kernel. Define buffer ownership, completion, and reuse for the prologue, steady state, and drain.\n4. Add architecture-specific MMA or scheduling when profiling supports it. Recheck register/shared-memory use and retune after each change.\n5. Run the full relevant correctness gate and same-hardware benchmarks, including supported variable-length and boundary cases. Measure production entry points as well as individual kernels.\n\nA literal translation establishes parity; measure it before assuming a performance gain. Preserve the validated numerical algorithm and precision while changing execution layout or scheduling.\n\n## Layout and resource constraints\n\n- Pointer tensors and values need compatible layouts. A cross-warp `gl.convert_layout` may use shared memory; use `assert_trivial=True` when a conversion is intended to be free.\n- Broadcasts and oversized layout tiles can duplicate values and increase register use. Choose layouts from actual tensor strides rather than copying a tutorial's block sizes.\n- Account for all live shared-memory buffers and the target device's allocation limit. Prune infeasible autotune configurations; provide a streaming path when a resident design cannot fit supported shapes.\n- Large static unrolls multiplied by many autotune configurations can dominate compilation. Prune from the actual shape and live-memory budget, and reuse the worker and Triton cache during iteration.\n- Persistent schedules can trade launch overhead for worse L2 locality. Warp specialization also changes the total register budget; measure occupancy and stalls after changing either.\n\n## Async correctness\n\n### Transfers and buffer reuse\n\n`cp.async` groups are ordered as a queue. `wait_group(N)` bounds all outstanding groups, so it cannot identify a particular buffer after unrelated prefetches are interleaved. For per-buffer scheduling, use separate mbarriers and match their arrival counts to the participating threads. In the installed API, check whether `mbarrier_arrive` increments the count; a preinitialized thread count requires the non-incrementing form.\n\nWait for a load before consuming its buffer, and protect reuse until all readers finish. Even same-lane shared-memory staging needs protection against overwriting data still being read. Track each mbarrier's phase as its buffer is reused; do not run more than one phase ahead or share a completion barrier between TMA and tcgen05 without reinitializing it.\n\nTMA descriptors must satisfy the target's alignment and stride requirements. TMA stores may distinguish completion of the shared-memory read from completion of the global write. Check the installed wait API and require global completion before another operation reads the stored range.\n\n### MMA and memory ordering\n\nHopper WGMMA requires its B operand in shared memory and uses register accumulators; consume the values returned by its wait operation so compiler dependencies remain explicit. Blackwell tcgen05 uses TMEM accumulators and mbarrier completion; respect the participating warpgroup and TMEM layout requirements. Initialize accumulators explicitly, including `use_acc=False` where supported.\n\nGeneric shared-memory accesses and async operations use different memory proxies. Apply the required proxy fence when handing a buffer to an async consumer; an mbarrier alone does not replace that fence. A completed TMA load establishes the ordering needed to read its destination. Check ordering across warp-specialized partitions as well as within a single producer/consumer loop.\n\n### Tails and exact-zero cases\n\nMasked async copies can leave shared-memory elements uninitialized. Prevent invalid rows from entering reductions: zero-fill where supported, or use valid-row loads and explicitly remove every invalid contribution. Multiplying by zero does not neutralize a NaN. Keep output stores masked.\n\nDo not rely on separately compiled reductions cancelling bitwise. For a mathematically exact-zero special case, such as a single-source softmax gradient, preserve the exact-zero result explicitly and test it against the reference.\n\n## Verification commands\n\nRun from the repository root on the target GPU:\n\n```bash\nFLA_BENCH_BASE=$(git rev-parse origin/main)\nFLA_CI_ENV=0 python -m benchmarks.ops.verify --op chunk_kda --base \"$FLA_BENCH_BASE\"\n```\n\nReplace the operation with the one being ported. The command executes the current tests; it does not freeze them. Use `--gate-k` only for quick iteration, then run the full gate. Set backend environment flags before launching a fresh process so import-time choices cannot contaminate the comparison.\n","frontmatter":{"name":"fla-triton-to-gluon","description":"Port an FLA Triton kernel to Gluon when explicit layouts, asynchronous transfers, or scheduling can address a measured bottleneck."},"isInternal":false,"tokens":1257,"sizeBytes":6414},{"name":"README.md","path":".agents/skills/README.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/.agents/skills/README.md","title":"Skills Skill","category":"anthropic-skill","format":"markdown","content":"# Repository skills\n\n[CONTRIBUTING.md](../../CONTRIBUTING.md) defines contribution rules; [AGENTS.md](../../AGENTS.md) defines agent operations. Skills add task-specific guidance. Choose those relevant to the change.\n\nFor kernel work, define the contract and establish baseline coverage before optimization. Use the operator, dispatch, and hardware guides as needed, then prepare the PR with the collected evidence.\n\n| Skill                                                     | Use for                                                    |\n| --------------------------------------------------------- | ---------------------------------------------------------- |\n| [Design coverage](fla-design-coverage/SKILL.md)           | Supported inputs, numerical budgets, and routing contracts |\n| [Correctness coverage](fla-correctness-coverage/SKILL.md) | References, regression cases, and affected test selection  |\n| [KDA](fla-kda/SKILL.md)                                   | KDA gate modes and kernel invariants                       |\n| [Backend dispatch](fla-dispatch-backends/SKILL.md)        | Registration, verifiers, and fallback behavior             |\n| [Optimization loop](fla-optimization-loop/SKILL.md)       | A fixed correctness gate and measured candidate iterations |\n| [NVIDIA performance](fla-nvidia-performance/SKILL.md)     | GPU profiling and bottleneck diagnosis                     |\n| [Ascend performance](fla-ascend-performance/SKILL.md)     | NPU profiling and compiler constraints                     |\n| [Triton to Gluon](fla-triton-to-gluon/SKILL.md)           | Layouts, synchronization, and incremental kernel ports     |\n| [PR readiness](fla-pr-readiness/SKILL.md)                 | Scope, validation evidence, and submission                 |\n| [PR stacks](fla-pr-stacks/SKILL.md)                       | Splitting large changes, reusing PRs, and native stacks    |\n\n## Maintaining skills\n\nPut each workflow in `.agents/skills/<skill-name>/SKILL.md` with YAML `name` and `description` fields. Follow [documentation and skill style](../../CONTRIBUTING.md#prose-and-markdown) for scope, organization, wording, and shared policy.\n\nUse an existing `references/` directory for substantial task-specific detail. Add supporting files only when they make the workflow easier to use; `SKILL.md` is the entry point, so individual skill directories do not need a README. Symlinks may point to public documentation tracked in this repository.\n","isInternal":false,"tokens":507,"sizeBytes":2458},{"name":"README.md","path":"evals/README.md","rawUrl":"https://raw.githubusercontent.com/fla-org/flash-linear-attention/HEAD/evals/README.md","title":"Evals Skill","category":"anthropic-skill","format":"markdown","content":"# Evaluation\n\nWe provide evaluation scripts for `fla` models with [lm-evaluation-harness](https://github.com/EleutherAI/lm-evaluation-harness), as well as a standalone perplexity evaluator.\nInstall `fla` following the [installation guide](../INSTALL.md), and run the commands below from the repository root.\n\n* [lm-evaluation-harness](#lm-evaluation-harness)\n  * [Single-GPU Evaluation](#single-gpu-evaluation)\n  * [Multi-GPU Evaluation](#multi-gpu-evaluation)\n  * [RULER Benchmarks](#ruler-benchmarks)\n* [Perplexity](#perplexity)\n  * [Quick Start](#quick-start)\n  * [Evaluation Settings](#evaluation-settings)\n  * [Scoring](#scoring)\n  * [Results](#results)\n\n## lm-evaluation-harness\n\nThe [lm-evaluation-harness](https://github.com/EleutherAI/lm-evaluation-harness) library allows you to easily perform (zero-shot) model evaluations.\nInstall `lm_eval` following [their instructions](https://github.com/EleutherAI/lm-evaluation-harness/blob/main/README.md).\n\n### Single-GPU Evaluation\n\nRun evaluation with:\n\n```sh\nMODEL='fla-hub/gla-1.3B-100B'\npython -m evals.harness --model hf \\\n    --model_args pretrained=$MODEL,dtype=bfloat16 \\\n    --tasks wikitext,lambada_openai,piqa,hellaswag,winogrande,arc_easy,arc_challenge,boolq,sciq,copa,openbookqa \\\n    --batch_size 64 \\\n    --num_fewshot 0 \\\n    --device cuda \\\n    --show_config\n```\n\nWe've made `fla` compatible with hf-style evaluations, you can call [evals.harness](harness.py) to finish the evaluations.\nThe command above evaluates the GLA checkpoint on the tasks used in the GLA paper.\n\n### Multi-GPU Evaluation\n\nUse the Hugging Face Accelerate launcher for data-parallel evaluation, where each GPU loads a full copy of the model:\n\n```sh\nMODEL='fla-hub/gla-1.3B-100B'\naccelerate launch -m evals.harness --model hf \\\n    --model_args pretrained=$MODEL,dtype=bfloat16,trust_remote_code=True \\\n    --tasks wikitext,lambada_openai,piqa,hellaswag,winogrande,arc_easy,arc_challenge,boolq,sciq,copa,openbookqa \\\n    --batch_size 64 \\\n    --num_fewshot 0 \\\n    --device cuda \\\n    --show_config \\\n    --trust_remote_code\n```\n\n### RULER Benchmarks\n\nThe RULER benchmarks are commonly used for evaluating model performance on long-context tasks.\nYou can evaluate `fla` models on RULER directly using `lm-evaluation-harness`. RULER is only available in a relatively recent version of `lm-evaluation-harness`, so make sure you have the latest version installed.\n\n```sh\ngit clone --depth 1 https://github.com/EleutherAI/lm-evaluation-harness\npip install -e './lm-evaluation-harness[ruler]'\n```\n\nFor example, to evaluate contexts up to 32K tokens:\n\n```sh\nMODEL='fla-hub/gla-1.3B-100B'\nOUTPUT='results/ruler'\naccelerate launch -m evals.harness --model hf \\\n    --output_path $OUTPUT \\\n    --tasks niah_single_1,niah_single_2,niah_single_3,niah_multikey_1,niah_multikey_2,niah_multikey_3,niah_multiquery,niah_multivalue,ruler_vt,ruler_cwe,ruler_fwe,ruler_qa_hotpot,ruler_qa_squad \\\n    --model_args pretrained=$MODEL,dtype=bfloat16,max_length=32768,trust_remote_code=True \\\n    --metadata='{\"max_seq_lengths\":[4096,8192,16384,32768]}' \\\n    --batch_size 2 \\\n    --show_config \\\n    --trust_remote_code\n```\n\nIf a GPU can't load a full copy of the model, please refer to [this link](https://github.com/EleutherAI/lm-evaluation-harness?tab=readme-ov-file#multi-gpu-evaluation-with-hugging-face-accelerate) for FSDP settings.\n\n> [!TIP]\n> If you are using `lm-evaluation-harness` as an external library and can't find (almost) any tasks available, before calling `lm_eval.evaluate()` or `lm_eval.simple_evaluate()`, simply run the following to load the library's stock tasks:\n> ```py\n> from lm_eval.tasks import TaskManager\n> TaskManager().initialize_tasks()\n> ```\n\n## Perplexity\n\n[evals.ppl](ppl.py) evaluates next-token perplexity over fixed-length contexts and reports perplexity by position within each context.\n\n### Quick Start\n\nInstall the dataset dependency:\n\n```sh\npip install datasets\n```\n\nFor example, to evaluate the GLA checkpoint on the PG19 test split and save the results:\n\n```sh\npython -m evals.ppl \\\n    --path fla-hub/gla-1.3B-100B \\\n    --data fla-hub/pg19 --split test \\\n    --block_size 28672 --bucket_size 2048 --batch_size 1 \\\n    --dtype bfloat16 --num_proc 4 \\\n    --output results/ppl.json\n```\n\n### Evaluation Settings\n\n| Argument          | Default                 | Description                                                 |\n|-------------------|-------------------------|-------------------------------------------------------------|\n| `--path`          | `fla-hub/gla-1.3B-100B` | Hugging Face model ID or local checkpoint directory.        |\n| `--data`          | `fla-hub/pg19`          | Dataset to evaluate.                                        |\n| `--split`         | `train`                 | Dataset split; set this explicitly for held-out evaluation. |\n| `--column_name`   | `text`                  | Dataset column containing text.                             |\n| `--block_size`    | `28672`                 | Context length in tokens; must be at least 2.               |\n| `--bucket_size`   | `2048`                  | Number of prediction positions per reported bucket.         |\n| `--batch_size`    | `1`                     | Number of full contexts evaluated together.                 |\n| `--device`        | Auto-detected           | Device on which to load and evaluate the model.             |\n| `--dtype`         | `bfloat16`              | Model dtype: `float32`, `float16`, or `bfloat16`.           |\n| `--num_proc`      | `32`                    | Number of dataset tokenization processes.                   |\n| `--revision`      | `main`                  | Model and tokenizer revision.                               |\n| `--data_revision` | `main`                  | Dataset revision.                                           |\n| `--output`        | None                    | Optional path for JSON results and evaluation settings.     |\n\nUse immutable commit hashes for `--revision` and `--data_revision` to reproduce a run.\nFor long contexts, reduce `--batch_size` to lower model memory use and `--bucket_size` to reduce the FP32 softmax workspace.\nThe model still materializes its full logits.\n\n### Scoring\n\nDocuments are tokenized with the tokenizer's default special-token policy, then concatenated in dataset order without extra separators.\nThe token stream is divided into independent contexts of `block_size` tokens, and the final partial context is evaluated separately without padding.\nContext is reset at every block boundary; there is no sliding window or recurrent state carried between blocks.\n\nEach context of length `L` contributes `L-1` next-token predictions.\nThe first token is unscored, and no synthetic EOS target is appended after the last token.\nEOS tokens present in the input are scored normally.\nA one-token tail contributes no predictions; a dataset with fewer than two tokens raises `ValueError`.\n\nOverall PPL and positional PPL use the same FP32 negative log-likelihoods, weighted by the number of scored predictions.\nBuckets group predictions by their zero-based logit position within each context: bucket 0 covers positions `0` through `bucket_size-1`.\nModels configured with fused linear cross entropy are supported, and evaluation runs with caching disabled.\n\n### Results\n\nThe script prints overall PPL, positional PPL, and token counts.\nWith `--output`, it also saves a JSON file containing:\n\n* `perplexity` and `block_perplexities`: overall and positional PPL; buckets with no predictions have `null` PPL.\n* `total_nll`, `block_nlls`, and `block_tokens`: summed NLLs and per-bucket prediction counts.\n* `total_tokens` and `total_input_tokens`: scored predictions and input tokens, respectively.\n* `total_sentences`: the number of evaluated contexts, including any one-token tail.\n* `config`, `model_commit`, `dataset_fingerprint`, and `versions`: evaluation settings, the resolved model commit when available, dataset fingerprint, and package versions.\n\n> [!NOTE]\n> Earlier versions appended a synthetic EOS only to positional statistics, used BF16 softmax for BF16 logits, and dropped the final partial context. Scores and token counts from those versions may differ from the corrected results.\n","isInternal":false,"tokens":1996,"sizeBytes":8221}],"systemPromptSnippet":"<yakaai_skills repo=\"fla-org/flash-linear-attention\">\n<!-- File: AGENTS.md (Tokens: ~1322 | Category: root-instruction) -->\n# FLA Guidelines\n\nGuidance for AI coding agents (Claude Code, Codex, etc.) working in this repo.\n\n**Read `CONTRIBUTING.md` first.** It is the authoritative source for all code style, docstring, comment, commit, PR, and testing conventions, and applies to humans and agents alike. This file only covers agent-specific operational behavior that doesn't belong in a human contributor doc.\n\n## Git safety\n\n- **Never discard uncommitted work** with `git checkout HEAD -- <file>`, `git reset --hard`, or `git restore` to \"get a clean base\". Unstaged changes are unrecoverable (no blob, no reflog). Edit in place or `git stash` instead, and confirm with the user when in doubt.\n- **On `main`**: never commit or push without explicit per-action approval. Suggest a feature branch first.\n- Don't rewrite or amend already-pushed commits unless the user asks.\n- **Branch names and attribution**: name local and remote branches after the change. Do not include tool or assistant identities in branch names, commit messages, PR titles/bodies, or review comments. Do not add `Co-Authored-By` trailers or generated-by notices.\n\n## Scope and direction\n\n- **Check for duplicates first**: search open issues/PRs before starting so you don't redo in-flight work — `gh pr list --repo fla-org/flash-linear-attention --state open --search \"<keywords>\"`.\n- **Align before coding when a design decision is involved**: a new operator or model, a new public API, a kernel rewrite with a different algorithm, or a repo-wide policy change (e.g. adding a dtype everywhere). Open an issue or a draft PR first; if a matching issue/PR already exists, comment there instead of opening a duplicate. Judge design impact separately from the PR size check: tests, config boilerplate, and mechanical repetition of an existing pattern alone do not make a change a design decision.\n- **RFC first for precision or algorithm changes**: changing the precision a validated kernel computes or accumulates in (e.g. fp32 → bf16, bf16 → fp16), relaxing a tolerance, or replacing a numerical algorithm (e.g. a different inversion or solve scheme) needs an RFC issue before a PR, even if it would stay within test tolerance. FLA's kernels are widely validated and downstream users rely on their numerics, so these are design decisions to agree on, not fixes to ship. Ordinary reorderings within the same precision — tiling, accumulation order, scheduling — do not need an RFC; they are normal optimization PRs with before/after benchmarks as usual.\n- **Stop and ask before breaking changes**: renaming or removing public symbols or arguments in `fla/layers/` / `fla/models/` or documented config fields; changing checkpoint/state-dict compatibility; changing observable behavior, including numerics beyond the existing test tolerance. Bug fixes that restore intended behavior are not breaking. If the user doesn't answer — or you are running non-interactively — do the non-breaking parts, leave the breaking part out, and describe it in your final report. Silence is not consent.\n\n## Comment discipline\n\nThe style rules for comments and docstrings live in `CONTRIBUTING.md` (\"Docstrings and Comments\"), including the banned anti-patterns. What follows is agent-specific behavior:\n\n- **You own the comments your change invalidates.** When a change makes a comment or docstring in the code you touched factually wrong, fix it or delete it in the same commit — \"the code you touched\" means the same function, kernel, or class, not a repo-wide sweep for stale mentions. If you can't tell whether a comment is still true, keep it and flag it in the PR description; never delete a \"why\" comment just because you can't verify it.\n- **Comments are not your scratchpad.** State design conclusions, not the journey: \"not X, it does not compile\" is a welcome why comment; \"tried X first, then switched to Y\" is chain-of-thought noise. This covers docstrings too.\n- **Fix the words, not the style.** Correcting stale content is not a license to reformat the surrounding docstring or comment style.\n\n## PR workflow\n\n- **No busywork PRs**: don't open a one-off PR for a single typo or isolated style tweak; bundle trivial cleanups into substantive work.\n- **Minimal diff**: every changed line serves the PR's stated purpose. Drop drive-by reformats, unrelated cleanups, and speculative generality before pushing, not after a reviewer asks (see CONTRIBUTING \"Core Principles\" 5).\n- **Blast radius**: for every shared symbol you touched, all callsites still hold (Core Principle 4 — audit them in one pass, don't assume). For any behavior or numerics change, the PR description must state who is affected — which shapes, dtypes, backends, checkpoints, or downstream users — or state explicitly why no one is. \"Tests pass\" is not an impact analysis.\n- Before opening a PR, re-read your own diff once with fresh eyes: check every comment it adds against the anti-pattern list, and every comment it invalidates against [Comment discipline](#comment-discipline).\n- Check the diff for `tl.make_block_ptr` / `tl.advance`; these are banned from mainline Triton code. Follow the Ascend backend exemption in CONTRIBUTING \"Triton Kernels\".\n- `gh pr edit` fails on this repo (classic-Projects GraphQL error). Edit a PR title/body via the REST API instead: `gh api -X PATCH repos/fla-org/flash-linear-attention/pulls/<N> -f title='...' -F body=@file`.\n- Keep review/PR comments concise and natural — skip heavy `**1.** **2.**` scaffolding, write like a person.\n\n## Repo-local Skills\n\nUse the [skill index](.agents/skills/README.md) to select task-specific workflows and their execution order. Skills supplement `CONTRIBUTING.md` and this file; they do not define competing contribution policies.\n\nLoad a relevant skill when the task matches its scope. In particular, load [fla-pr-readiness](.agents/skills/fla-pr-readiness/SKILL.md) before opening any PR.\n\n<!-- File: .agents/skills/fla-ascend-performance/references/cases.md (Tokens: ~1086 | Category: anthropic-skill) -->\n# Ascend kernel cases\n\nUse these examples to understand existing design choices. Read the current implementation before changing a path; tile settings depend on the workload and compiler. Paths are relative to the repository root.\n\n## Causal convolution: core grid and DMA paths\n\nImplementation: `fla/modules/causal_conv1d/backends/triton_ascend.py`.\n\nThe contiguous packed training path uses a 1D Vector-core task loop when channel tiling and state requirements permit. Other layouts, channel dimensions, and cache-state calls retain separate paths.\n\n- Transpose the public convolution weight to make channels contiguous for block loads. Preserve the public weight and output layouts.\n- Load one window including the convolution halo and use `extract_slice` for taps where supported. Keeping every tap live at once can exceed UB.\n- Preserve the saved pre-activation used by backward for silu/swish; removing it can trigger a second forward computation.\n- Separate bulk and tail DMA at compile time. A runtime choice keeps both paths live and can prevent larger tiles even when UB bandwidth is not saturated.\n\n| Path                               | DMA policy                                                     |\n| ---------------------------------- | -------------------------------------------------------------- |\n| Full chunks in packed input        | Block-pointer loads, with the tail path compiled out.          |\n| Final packed chunk                 | Masked loads/stores, including the halo bounds.                |\n| Variable lengths or a single chunk | Runtime bounds where a uniform bulk/tail split is unavailable. |\n\nUse `tests/modules/test_conv.py` for forward/backward, optional state, and layout validation. [TRAPS.md](TRAPS.md) covers optional pointers, packed bounds, and address types shared by these paths.\n\n## `chunk_delta_h.py` — bwd `dhu`\n\nImplementation: `fla/ops/common/backends/triton_ascend/chunk_delta_h.py`.\n\nGate loads along time use the [contiguous gate layout](g-contiguous-loading.md). The aligned fixed-length path can precompute gate terms, while variable-length or unaligned paths compute them inline. Model these paths separately: inline exponentials and temporary tiles can exceed the UB estimate for the precomputed path.\n\nChoose K and V tiles together. Reducing one dimension merely to maximize the other can increase repeated work. Rebuild task-local pointers from base addresses and return the corrected value gradient separately to keep the caller's `dv` unchanged. Backward also reuses dot operands across K slabs; apply the [operand lifetime rule](TRAPS.md#reused-tldot-operands).\n\n## `chunk_o.py` — fusion and gate loading\n\nImplementation: `fla/ops/common/backends/triton_ascend/chunk_o.py`.\n\nForward combines inter- and intra-chunk contributions when their live tiles fit UB, reducing intermediate output traffic. Backward gate loads use contiguous time storage while gradient outputs retain their public layout. Profile the combined launch and transpose cost, not just one kernel. See [gate loading](g-contiguous-loading.md) for pointer formulas and focused tests.\n\n## `chunk_bwd.py` — packed variable lengths\n\nImplementation: `fla/ops/kda/backends/triton_ascend/chunk_bwd.py`.\n\nPacked offsets multiplied by head and value dimensions can overflow int32 far before the token count reaches its limit. Load sequence boundaries as int64 and promote fixed-batch indices before multiplication. Only local lengths and block-pointer metadata should narrow to int32 when their bounds permit it.\n\n## `tl.dot` lhs clobber — repo-wide case catalog\n\nThe common rule is in [TRAPS.md](TRAPS.md#reused-tldot-operands). This index identifies where to inspect operand reuse without prescribing variable names or a fixed number of copies.\n\n| Implementation                                               | Reuse to inspect                                                                               |\n| ------------------------------------------------------------ | ---------------------------------------------------------------------------------------------- |\n| `fla/ops/common/backends/triton_ascend/chunk_o.py`           | Query/key tiles shared by multiple dots; attention tiles reused across value blocks.           |\n| `fla/ops/common/backends/triton_ascend/chunk_delta_h.py`     | Output-gradient and corrected residual tiles reused across K slabs.                            |\n| `fla/ops/gated_delta_rule/backends/triton_ascend/wy_fast.py` | Inverse tiles reused between stages; gradient tiles used first on the left, then on the right. |\n| `fla/ops/kda/backends/triton_ascend/wy_fast.py`              | Inverse tiles reused across the value and key loops.                                           |\n| `fla/ops/kda/backends/triton_ascend/chunk_intra.py`          | Triangular block merges that reuse operands for several products and final stores.             |\n| `fla/ops/utils/triton_ascend/solve_tril.py`                  | Block inverse operands used in multiple products or stored after a dot.                        |\n| `fla/ops/kda/backends/triton_ascend/chunk_bwd.py`            | Output gradients used as both left and right operands.                                         |\n\nAudit the changed kernel's first left-operand use and every subsequent read. Copies must precede that first use; reloads must come from pristine storage. Validate with the operator's kernel tests, including `tests/ops/test_gdn_kernels.py` and `tests/ops/test_solve_tril.py` where relevant.\n\n<!-- File: .agents/skills/fla-ascend-performance/references/g-contiguous-loading.md (Tokens: ~550 | Category: anthropic-skill) -->\n# Contiguous gate loading\n\nUse this pattern when an Ascend kernel repeatedly loads gate values along time from `[B, T, HV]` storage. The time stride is `HV`, so these loads gather separated values. Transposing the input gate can reduce memory-transfer overhead; measure the transpose and kernel together.\n\nReference implementation: `fla/ops/common/backends/triton_ascend/chunk_o.py`.\n\n## Host layout\n\n```python\ng_t_contig = g is not None and HV != 1\ng_arg = g.transpose(1, 2).contiguous() if g_t_contig else g\n```\n\nPass the selected input and layout flag to the kernel. `HV == 1` is already contiguous along time. Keep gradient outputs in their original `[B, T, HV]` layout.\n\n## Kernel addressing\n\nSave the total packed length before computing a local sequence length. Promote batch, head, and packed sequence offsets before multiplying them; see [address arithmetic](TRAPS.md#address-arithmetic).\n\n| Input mode              | Contiguous gate base                    |\n| ----------------------- | --------------------------------------- |\n| Fixed batch             | `g + batch * HV * T_seq + head * T_seq` |\n| Packed variable lengths | `g + bos + head * T_seq`                |\n\nHere `batch`, `head`, and `bos` are int64 values, `T_seq` is the host-passed storage length, and the block-pointer bound uses the local sequence length. Block offsets are relative to that sequence's start.\n\nFor a contiguous gate base, use stride 1 along time. The original-layout branch instead uses `g + bos * HV + head` with time stride `HV`. Apply the base offset once and keep every load in a branch consistent with that layout. Reusing the original stride or offset after transposing produces wrong values or alignment faults.\n\n## Validation\n\nCompare both layout paths against the same reference, including head sharing, variable lengths, non-aligned tails, and long sequences. The focused tests are:\n\n- `tests/ops/test_gdn_kernels.py::test_chunk_bwd_dv_local`\n- `tests/ops/test_gdn_kernels.py::test_chunk_bwd_dqkwg`\n\nInspect wall-clock duration and MTE activity before and after, including transpose and gradient-finalization work. Sum-of-block profiler durations may differ from one launch's elapsed time.\n\nPreserve the exponential formula while changing layout. Replacing `exp2(a - b)` with a ratio can change overflow and rounding behavior. On Ascend, use the supported slicing/layout operations from the current implementation; a reshape or join that compiles does not prove the intended element ordering.\n\n<!-- File: .agents/skills/fla-ascend-performance/references/reference.md (Tokens: ~1051 | Category: anthropic-skill) -->\n# Ascend profiling reference\n\nUse this reference for collection options, CSV interpretation, and choosing the next optimization. Kernel correctness hazards are in [TRAPS.md](TRAPS.md); implementation examples are in [cases.md](cases.md).\n\n## Collection and analysis\n\nRun the scripts from the active NPU environment. `profile_npu.py` accepts a workload file or inline source defining `workload()`. It executes that callable repeatedly, so initialize reusable inputs outside the callable and make repeated forward/backward calls valid.\n\n```bash\npython .agents/skills/fla-ascend-performance/scripts/profile_npu.py \\\n  --name my_op --out-dir profile/my_op-npu \\\n  --metrics PipeUtilization --analyze --exec-file path/to/workload.py\n```\n\nFor library use, add the scripts directory to `PYTHONPATH` and call `profile_callable`:\n\n```bash\nexport PYTHONPATH=\"$PWD/.agents/skills/fla-ascend-performance/scripts${PYTHONPATH:+:$PYTHONPATH}\"\n```\n\n```python\nfrom profile_npu import profile_callable\n\ntrace_dir = profile_callable(fn=workload, name=\"my_op\", out_dir=\"profile/my_op-npu\", aic_metrics=\"PipeUtilization\")\n```\n\nThe returned directory contains the profiler output:\n\n```text\nmy_op_profiling_TIMESTAMP/\n└── localhost..._ascend_pt/\n    └── ASCEND_PROFILER_OUTPUT/\n        ├── op_statistic.csv\n        └── kernel_details.csv\n```\n\nPass that exact run directory when analyzing an existing trace. The analyzer takes the first matching CSV under the supplied directory, so a parent containing several runs can select an older result.\n\n```bash\npython .agents/skills/fla-ascend-performance/scripts/analyze_profile.py \\\n  profile/my_op-npu/my_op_profiling_TIMESTAMP --kernel-filter my_kernel --top-k 20\n```\n\n`--kernel-filter` applies to kernel details. If no name matches, the analyzer displays the top kernels instead; confirm the printed names before attributing metrics to the target.\n\n## Choosing metrics\n\nEach profiling run collects one `aic_metrics` set. Use separate runs for complementary metrics and keep the workload unchanged.\n\n| Metric set                                       | Use                                                                     |\n| ------------------------------------------------ | ----------------------------------------------------------------------- |\n| `PipeUtilization`                                | First pass: compare Cube, Vector, scalar, and memory-transfer activity. |\n| `MemoryUB`                                       | Check UB read/write bandwidth when memory traffic may limit the kernel. |\n| `Memory`, `MemoryAccess`, `MemoryL0`             | Inspect other memory paths.                                             |\n| `L2Cache`                                        | Investigate cache reuse.                                                |\n| `ArithmeticUtilization`, `ResourceConflictRatio` | Investigate arithmetic occupancy or resource conflicts.                 |\n\nAvailable metric names depend on the installed `torch_npu`; the collector reports available names when an option is unknown.\n\n## Diagnosis\n\nStart with total duration in `op_statistic.csv`, then inspect the dominant kernel in `kernel_details.csv`. Pipe activity alone does not establish the bottleneck.\n\n| Evidence                                               | Check next                                                                                                   |\n| ------------------------------------------------------ | ------------------------------------------------------------------------------------------------------------ |\n| Target kernel absent or many small fallback operations | Backend selection and fusion.                                                                                |\n| High Vector or Cube activity                           | Tile shape, useful work per tile, and scalar overhead around compute.                                        |\n| High MTE activity                                      | Reuse, intermediate writebacks, and strided loads; inspect [gate loading](g-contiguous-loading.md).          |\n| High scalar activity                                   | Address generation, branches, and repeated per-tile setup.                                                   |\n| High UB bandwidth with low compute activity            | Live buffers, layout conversions, and repeated UB reads/writes.                                              |\n| Low activity despite UB overflow at larger tiles       | Mutually exclusive DMA paths may remain live; see [DMA splitting](TRAPS.md#dma-paths-and-optional-pointers). |\n| Many host launches for grid chunks                     | Consider a 1D task loop with the appropriate Cube or Vector core count.                                      |\n\n## UB planning\n\nEstimate peak live storage from simultaneous buffers, their dtype, and any compiler-created copies. Use `fla.utils.ascend_ub_manager` for device capacity and tiling rather than hard-coding the UB size. The helper's safety margin reserves capacity; it is separate from the multiplier that estimates live buffers.\n\nIf more capacity is available, test a larger tile and verify compilation and numerics. If the live set already fills the budget, improve reuse or split stages. Model forward and backward separately; inline gate computation can need more UB than a precomputed-gate path.\n\nUseful code locations:\n\n- `fla/utils/ascend_ub_manager.py`: tiling and grid helpers.\n- `fla/utils/hardware.py`: `get_multiprocessor_count`, including `use_aicore=True` for Cube work.\n- `fla/modules/norm/triton_ascend/l2norm.py`: entry-point registration and the Ascend implementation.\n- `tests/ops/test_gdn_kernels.py`, `tests/ops/test_solve_tril.py`, and `tests/modules/test_conv.py`: relevant kernel comparisons.\n\n<!-- File: .agents/skills/fla-ascend-performance/references/TRAPS.md (Tokens: ~746 | Category: anthropic-skill) -->\n# Ascend compiler and memory traps\n\nRead the relevant section before changing DMA, launch indexing, or reused matrix operands. [Kernel cases](cases.md) show these constraints in existing implementations.\n\n## DMA paths and optional pointers\n\nA runtime branch between block-pointer DMA and masked loads can keep both paths live in UB. Larger tiles may then overflow even when measured UB bandwidth is low. Where the workload permits, split bulk and tail launches with a constexpr mode so each compilation removes the unused path. Include convolution halos when deciding whether a bulk load fits inside the packed allocation.\n\nAn optional pointer needs a compile-time guard before any runtime condition. `if USE_INITIAL_STATE or runtime_condition:` can still compile pointer arithmetic in the other branch when the pointer is `None`; nest the constexpr condition separately.\n\n## Address arithmetic\n\nCast indices to `tl.int64` before multiplying by strides or dimensions. Casting the product cannot repair overflow. On Ascend, specialized arguments and folded program IDs may be constexpr values without a `.to()` method; `tl.cast` works for both runtime and constexpr integers.\n\n```python\ntoken_offset = tl.cast(i_t, tl.int64) * BT\nbatch_offset = tl.cast(i_b, tl.int64) * T\nbos = tl.load(cu_seqlens + i_n).to(tl.int64)\n```\n\nLoad packed sequence offsets as int64 even when the input `cu_seqlens` uses int32. A small token index can overflow after multiplication by head count and feature size.\n\n`make_block_ptr` metadata has a separate constraint: its `offsets` and `block_shape` must remain int32. Keep large flattened base-address calculations in int64, then supply valid local block offsets. The Ascend backend is exempt from the mainline ban on block pointers.\n\n## Grid and task loops\n\nThe existing Ascend grid helpers enforce a grid-product limit of 65535. Use their host splitting or flatten independent tiles into a 1D loop over available cores. `get_multiprocessor_count()` selects Vector cores on NPU; `use_aicore=True` selects Cube cores.\n\nAfter slicing variable-length chunk metadata, apply the corresponding global offset only once. In a task loop, derive local pointers from the original base on each iteration; accumulating pointer updates across tasks can miscompile on Ascend.\n\n## Reused `tl.dot` operands\n\nTriton-Ascend can reuse the left operand's UB storage during `tl.dot`. A later read of that tile can therefore be wrong even without a compiler error. Audit all later uses, including another left operand, a right operand, arithmetic, and stores.\n\nReload the pristine tile from global memory when reuse is separated by stages. For tight reuse, create the needed `tile + 0.0` copies before the first dot that consumes the original as a left operand. Copying afterwards preserves the corrupted value. Account for the extra live copies in the UB budget and validate against the numerical reference.\n\n[The case index](cases.md#tldot-lhs-clobber--repo-wide-case-catalog) identifies the affected kernel patterns. A fresh load on each loop iteration or a tile used as a left operand only once needs no extra copy for this issue.\n\n## Numerical and tail behavior\n\nPreserve the validated accumulation precision and exponential base. Mask invalid values before exponentials and initialize every region consumed later. Factoring a gate difference into a ratio or reciprocal changes rounding and range behavior; check the existing reference and tolerances before using it.\n\nExercise gated and ungated paths when both are supported. A suite that always supplies a gate cannot catch `None` pointer handling or incorrect ungated residual updates.\n\n<!-- File: .agents/skills/fla-ascend-performance/SKILL.md (Tokens: ~892 | Category: anthropic-skill) -->\n---\nname: fla-ascend-performance\ndescription: Profile and optimize FLA Triton-Ascend kernels using NPU traces, with guidance for UB capacity, memory movement, launch limits, and numerical correctness.\n---\n\n# Ascend kernel performance\n\nUse this skill for operator performance work in `triton_ascend.py` and `triton_ascend/` backend implementations. Follow [fla-optimization-loop](../fla-optimization-loop/SKILL.md) for sustained optimization and [CONTRIBUTING](../../../CONTRIBUTING.md#benchmarking) for validation and reporting.\n\nKeep collection, analysis, and benchmarks in the active Python/NPU environment. If it is not configured, activate the project's Ascend environment first.\n\n## 1. Establish the baseline\n\nLocate the public entry, dispatch route, and Ascend kernel. Record the supported inputs and production workloads, baseline commit, and synchronized latency. Preserve the public API, validated algorithm, reference, and tolerances. Confirm the intended NPU kernel runs; a Torch fallback must not hide an unsupported path or kernel bug.\n\n## 2. Collect a trace\n\nUse the existing collector instead of duplicating profiler setup. Run from the repository root; the workload file initializes its inputs and defines a repeatable `workload()` callable.\n\n```bash\npython .agents/skills/fla-ascend-performance/scripts/profile_npu.py \\\n  --name my_op --out-dir profile/my_op-npu \\\n  --metrics PipeUtilization --analyze \\\n  --kernel-filter my_kernel --exec-file path/to/workload.py\n```\n\nOne run collects one `aic_metrics` set. Start with `PipeUtilization`; collect `MemoryUB` separately when UB bandwidth is relevant. The default schedule is one warmup step and one active step. Warm the workload enough to exclude compilation from the measurement.\n\nThe collector prints the exact trace directory. To revisit it, pass that directory to [the analyzer](scripts/analyze_profile.py); see [collection and metrics](references/reference.md#collection-and-analysis) for the library interface and output layout.\n\n## 3. Diagnose bottlenecks\n\nUse `op_statistic.csv` to find the dominant operation, then `kernel_details.csv` to inspect its duration and pipe activity. Confirm the dispatch route before tuning a small or missing target kernel. Distinguish compilation failure, UB overflow, grid limits, numerical errors, and measured performance limits; they need different fixes.\n\nUse the [diagnosis table](references/reference.md#diagnosis) to choose the next experiment. Change one relevant factor, then measure whether the expected duration and pipe metrics moved.\n\n## 4. Optimize within the kernel contract\n\n- Estimate the peak live UB footprint and use `fla.utils.ascend_ub_manager` for device capacity and tiling. Model forward and backward separately; fuse only when the live set fits.\n- Match the launch to the work: Cube kernels use Cube cores, Vector kernels use Vector cores. Handle the grid limit through the existing helpers or a 1D task loop.\n- Preserve numerical precision, optional inputs, and fixed/variable-length behavior. Verify tail handling and pointer arithmetic before increasing tiles or fusing stages.\n- Omit `num_warps` and `num_stages` from Ascend launches and autotune settings; these are unsupported on this backend.\n- Check [compiler and memory traps](references/TRAPS.md) when changing DMA paths, pointer arithmetic, or reused `tl.dot` operands. Use [contiguous gate loading](references/g-contiguous-loading.md) when gate loads along time are strided.\n\n[Kernel cases](references/cases.md) explain existing convolution, recurrence, and solve implementations. Their tile choices are workload-specific examples.\n\n## 5. Validate and measure\n\nRun kernel comparisons and the affected operator/layer tests, including supported gradients, variable lengths, optional inputs, and boundary shapes. Keep NaN poisoning enabled and confirm that dispatch exercises the changed path. Run the full relevant correctness gate before reporting a speedup; a `--gate-k` subset is an iteration check.\n\nCompare synchronized forward and forward/backward latency on the same NPU, then repeat the relevant profiler collection. Report the timing change, the metrics that explain it, any compiler limitation encountered, and untested cases. If the expected metrics do not change, revisit the diagnosis before adding another optimization.\n\n<!-- File: .agents/skills/fla-correctness-coverage/references/cp.md (Tokens: ~10 | Category: anthropic-skill) -->\n../../../../fla/ops/cp/README.md\n\n<!-- File: .agents/skills/fla-correctness-coverage/references/delta-rule.md (Tokens: ~11 | Category: anthropic-skill) -->\n../../../../fla/ops/delta_rule/README.md\n\n<!-- File: .agents/skills/fla-correctness-coverage/references/generalized-delta-rule.md (Tokens: ~12 | Category: anthropic-skill) -->\n../../../../fla/ops/generalized_delta_rule/README.md\n\n<!-- File: .agents/skills/fla-correctness-coverage/references/simple-gla.md (Tokens: ~11 | Category: anthropic-skill) -->\n../../../../fla/ops/simple_gla/README.md\n\n<!-- File: .agents/skills/fla-correctness-coverage/SKILL.md (Tokens: ~732 | Category: anthropic-skill) -->\n---\nname: fla-correctness-coverage\ndescription: Select and run correctness coverage for FLA kernels and modules, including gradients, dispatch boundaries, and Triton addressing changes.\n---\n\n# FLA Correctness Coverage\n\nUse [fla-design-coverage](../fla-design-coverage/SKILL.md) when the change needs a new numerical or routing contract. This skill turns that contract into tests; [CONTRIBUTING.md](../../../CONTRIBUTING.md#testing) defines references, tolerances, test structure, and platform conventions.\n\n## Select coverage\n\nRead the affected tests and callers before adding cases. Follow [test organization and naming](../../../CONTRIBUTING.md#test-organization-and-naming): try the existing parameter matrix first, and add a function or file only when its setup or scope requires one. Cover reachable gaps across the applicable dimensions together, rather than testing each flag only in isolation:\n\n| Dimension   | Cases that can expose different behavior                                         |\n| ----------- | -------------------------------------------------------------------------------- |\n| Layout      | Dense and varlen, partial chunks, long sequences, asymmetric QK/value dimensions |\n| Features    | Gate bounds, raw/post-sigmoid beta, QK normalization, grouped value attention    |\n| State       | Initial state, final state, and state gradients where supported                  |\n| Numerics    | Supported dtypes, affected conversion/reduction paths, numerical boundaries      |\n| Routing     | Default and selected backends, verifier acceptance, rejection, and fallback      |\n\nCompare outputs and final states against the existing reference. Training APIs also need every supported gradient; forward-only APIs must state that backward is unsupported. Use `torch.autograd.gradcheck` where the implementation supports its required precision. Routing tests verify selection and argument forwarding; they do not replace numerical comparisons on the actual backend.\n\nFor addressing changes, inspect casts before multiplication and test shapes that expose large offsets, partial tiles, and varlen boundaries. Follow the `tl.int64` and platform requirements in [Triton Kernels](../../../CONTRIBUTING.md#triton-kernels). Exercise changed grid/address paths on affected supported platforms. Distinguish unsupported cases from tests skipped for unavailable hardware or dependencies; record the latter as validation gaps.\n\nThe NaN allocation guard covers eligible FLA allocations in operator/module tests, including the single-rank CP graph and intra-card cache cases moved from operator tests, but not layer/model tests, spawned CP workers, or every allocation. Add explicit finite-output and finite-gradient checks to adversarial cases, and investigate a poisoned-run failure even when an ad hoc run passes.\n\n## Run the affected paths\n\nDiscover dependent tests from the repository root:\n\n```bash\npython scripts/find_dependent_tests.py fla/ops/kda/chunk.py\nFLA_CI_ENV=0 pytest tests/ops/test_kda.py -v\n```\n\nThe helper lists files; it does not run them. Include reported callers and relevant `tests/modules/`, `tests/layers/`, `tests/models/`, or `tests/cp/` coverage when the change reaches those paths. Inspect skips and warning-only comparisons before calling the run a pass. Record commands, backend flags, environment, and results; reproduce failures on the unchanged baseline before labeling them pre-existing.\n\nRead operator math only when needed: [Delta Rule](../../../fla/ops/delta_rule/README.md), [Generalized Delta Rule](../../../fla/ops/generalized_delta_rule/README.md), [Simple GLA](../../../fla/ops/simple_gla/README.md), or [context parallelism](../../../fla/ops/cp/README.md).\n\n<!-- File: .agents/skills/fla-design-coverage/SKILL.md (Tokens: ~981 | Category: anthropic-skill) -->\n---\nname: fla-design-coverage\ndescription: Define supported inputs, numerical budgets, routing, and validation before designing an FLA kernel or numerical change.\n---\n\n# FLA Design and Coverage\n\nUse this skill to establish the contract before implementing a kernel, backend route, or numerical change. Read [CONTRIBUTING.md](../../../CONTRIBUTING.md) for contribution rules, especially [public compatibility](../../../CONTRIBUTING.md#public-api-compatibility) and [design/RFC requirements](../../../CONTRIBUTING.md#submit-pull-requests).\n\n## Define the supported domain\n\nTrace public entries and defaults, backend verifiers, layer/model/CP callers, and existing tests. Record the merge-base SHA and the source of each supported case. Group cases only when they share the same numerical assumptions and execution route; omit combinations no supported caller can reach.\n\nDistinguish routes when kernel stages, chunk geometry or anchors, capability checks, dtype conversions, or input value ranges change. A backend here means an implementation selected by dispatch, including the default implementation, rather than a hardware product. For each group, identify its effective arguments, oracle, output/state and per-gradient tolerances, and test coverage. Assign one outcome:\n\n- The proposed implementation, with passing reference comparisons.\n- An existing fallback, with route-parity tests covering outputs, states, dtypes, and supported gradients.\n- A documented unsupported-input error raised before kernel execution and asserted by a test.\n\nPreserve previously supported cases under the public compatibility policy. A warning, partially initialized output, or fallback that changes semantics is not a valid outcome. When changing defaults, record the old and new effective values and selected routes.\n\n## Establish numerical safety\n\nFor each affected stage in forward and backward, record input-range assumptions, load/operand/accumulation/store dtypes, reduction behavior, overflow/underflow headroom, and error propagation to the end-to-end tolerance. Forward agreement alone does not establish backward safety, including inverse recomputation and state gradients.\n\nSafety flags and bounds apply to the stages and domains that justify them. Reassess them when chunk geometry, anchors, dtype staging, or input scaling changes; a safe exponent does not prove that a later inversion is safe. Document the algebra behind numerical acceptance thresholds, implement each shared predicate once, and test its acceptance boundary and the adjacent rejection or fallback. Enforce additional assumptions, such as input-norm bounds, at the public entry or every affected verifier; report the failed predicate and effective values.\n\nKeep the committed oracle, input distributions, assertions, and tolerances fixed while evaluating the candidate. A correction to that contract belongs in a separate change with before/after error distributions for the unchanged baseline and candidate; precision, tolerance, and algorithm changes follow the linked RFC policy. Keep the corrected gate green before the candidate lands. Every selectable autotune configuration must satisfy the same numerical contract.\n\n## Check routing\n\n- Share semantic argument normalization between the public entry, verifier, and executor. Test that omitted arguments match their documented explicit defaults in route, outputs, and gradients. `input_guard` handles contiguity and device context, not semantic defaults.\n- Use capability and availability helpers from `fla.utils`. Distinguish backend availability, policy enablement, and acceptance of an individual call; use [fla-dispatch-backends](../fla-dispatch-backends/SKILL.md) for implementation mechanics.\n- For context-parallel routes, every rank must make the same selection. If verifier inputs can differ by rank, combine acceptance with a collective and select the optimized route only if all ranks accept. Record whether forward commits backward to the same backend.\n- Test parity between routes claiming the same domain. Record precision settings such as `TRITON_F32_DEFAULT` and matmul flags, keep them identical across baseline, candidate, and tests, and restore global settings after tests.\n\n## Choose evidence\n\nApply the relevant axes from [fla-correctness-coverage](../fla-correctness-coverage/SKILL.md) to each supported group across these distinct layers:\n\n- **Production:** reproduce effective arguments and numerical input distributions from checked-in layer/model callers. A generic benchmark shape list alone does not establish production coverage.\n- **Public boundaries:** cover previously supported inputs, changed defaults, and verifier acceptance/rejection boundaries.\n- **Adversarial:** outside the contract, require documented rejection before kernel execution or explicit finite-output and finite-gradient checks. This does not expand the supported domain.\n\nBefore performance work, record baseline correctness, latency or throughput, commit, environment, workload, input distribution, and seed. Benchmark production cases first, then boundaries whose routing or expected performance changes. Unaffected boundaries still need correctness coverage. Use [fla-optimization-loop](../fla-optimization-loop/SKILL.md) to freeze the gate and compare candidates.\n\n<!-- File: .agents/skills/fla-dispatch-backends/SKILL.md (Tokens: ~741 | Category: anthropic-skill) -->\n---\nname: fla-dispatch-backends\ndescription: Change FLA backend registration, dispatch, or verifiers while preserving routing and call contracts.\n---\n\n# Backend dispatch\n\nUse this skill for registration, selection, or verifier changes in `fla/backends.py` and operation backends. [CONTRIBUTING.md](../../../CONTRIBUTING.md) defines style and public compatibility; [fla/backends.py](../../../fla/backends.py) defines the routing behavior.\n\n## 1. Define the supported calls\n\nIdentify the default entry point and the calls the backend can handle: inputs, layouts, dtypes, gradient modes, and optional dependencies. Match the entry point's signature, parameter order, defaults, return structure, and supported gradients.\n\nA verifier returns `(True, None)` or `(False, reason)`. It checks support without changing inputs or global state. Keep it beside the implementation; put helpers shared by several implementations in a utility module.\n\n## 2. Register the implementation\n\nKeep the default entry point backend-independent with bare `@dispatch`. Choose the registration form that matches the implementation:\n\n| Implementation               | Registration                                                         |\n| ---------------------------- | -------------------------------------------------------------------- |\n| Function with the same API   | `@register(entry, backend=TritonAscendBackend, verifier=...)`        |\n| Adapter with backend policy  | `@register(entry)` on a `BaseBackend` subclass                       |\n\nAn adapter defines methods named after the entry point, such as `chunk_kda` and `chunk_kda_verifier`. Register it where it is defined and import optional kernels only when selected.\n\nImport the owning package's backends after its default entry points are defined. For implementations with platform-specific dependencies, guard the import with the shared backend's `is_available()`. Do not gate registration on enable switches; dispatch reads those at call time. Keep per-operation switches on `@register(..., env_var=...)` and implementation paths out of the shared dispatcher.\n\nExisting owners using `@dispatch('<operation>')` use `@register('<operation>')` on adapter classes and load implementations through the owner's backend package. Both forms use the same registration and selection logic.\n\n## 3. Preserve routing behavior\n\nFor each call, try backends in ascending priority order; ties retain registration order:\n\n1. Check availability and enablement.\n2. Run the verifier; rejection proceeds to the next candidate.\n3. Call the selected implementation. Its exceptions propagate.\n4. If no candidate accepts, call the default implementation.\n\nRegistration replaces an existing backend of the same type. Function registration keys the registry by the unwrapped entry point, so compiler decorators and public aliases share that registry. Legacy owner registries must load their backend package before lookup.\n\nKeep selection outside `torch.compile` graphs. Do not cache availability or enablement in the dispatch wrapper: registered backends must remain eligible when their switches change. Python's import cache handles modules already loaded, and unused optional dependencies must not prevent package import.\n\n## 4. Verify through the entry point\n\nRun the decorated entry point for accepted calls, verifier rejection, disabled/unavailable backends, and fallback. Check optional arguments omitted as well as supplied, changed support boundaries, outputs, and supported gradients.\n\nUse explicit backend flags in routing tests. For the default path, set `FLA_DISABLE_BACKEND_DISPATCH=1` before importing FLA. Use fresh processes for import-order checks, and verify public aliases and shared backend ownership after registration changes.\n\nFind affected consumers with `scripts/find_dependent_tests.py` and run the relevant operator, layer, and model tests. Host routing tests establish dispatch behavior; accelerator reference tests establish numerical correctness.\n\n<!-- File: .agents/skills/fla-kda/SKILL.md (Tokens: ~675 | Category: anthropic-skill) -->\n---\nname: fla-kda\ndescription: Modify or review KDA gates, chunk and recurrent kernels, backend routes, and their correctness coverage.\n---\n\n# KDA\n\nUse this skill for `fla/ops/kda/` and its tests. Public entry points are `chunk_kda` and `fused_recurrent_kda`; inspect the affected entry point's signature before changing its implementation.\n\n## Code map\n\n- `gate.py` defines gate activations and cumulative decay.\n- `chunk_fwd.py` and `chunk_bwd.py` coordinate training; `chunk_intra.py`, `chunk_intra_token_parallel.py`, and `wy_fast.py` implement the main stages.\n- `backends/` contains alternative implementations and verifiers. Follow [backend dispatch](../fla-dispatch-backends/SKILL.md) when changing routing.\n- `tests/ops/test_kda.py` contains the reference comparisons and supported gate-mode cases.\n\n## Gate contracts\n\n`chunk_kda` takes log-space decay when `use_gate_in_kernel=False`. With it enabled, the caller supplies raw gates and optional bias:\n\n- Without a lower bound, activation is `-exp(A_log) * softplus(g + dt_bias)` and `A_log` is required.\n- With a lower bound, activation is `lower_bound * sigmoid(exp(A_log) * (g + dt_bias))`. Omitting `A_log` uses a multiplier of one.\n\n`safe_gate` selects the intra-chunk implementation independently of gate activation. Safe in-kernel activation requires `-5 <= lower_bound < 0`. Pre-gated safe mode requires bounded log-space gates supplied by the caller; tests cover `[-5, 0]`, but the wrapper does not validate tensor values.\n\n## Numerical invariants\n\nThe safe intra path uses 16-token diagonal blocks and midpoint offsets before exponentiation. At a lower bound of `-5`, a block can accumulate `-80` in natural-log units; local offsets limit each exponent's span to about half that range. Preserve these offsets rather than exponentiating the full cumulative gate. Inter-block decay uses paired differences whose exponents remain non-positive for monotonic decay.\n\nThe non-safe path uses token-parallel diagonal computation. Both paths share inter-block and triangular-solve work, so changes there require coverage of both modes. Keep precision and tolerance changes within the design process in [AGENTS.md](../../../AGENTS.md#scope-and-direction).\n\n## Validation\n\nUse [correctness coverage](../fla-correctness-coverage/SKILL.md) for the test strategy. Select the KDA cases affected by the change:\n\n- Dense and variable-length sequences; forward, final state, and supported gradients.\n- Pre-gated and in-kernel gates, safe and non-safe paths, and omitted optional arguments.\n- Beta activation and query/key normalization modes when touched.\n- Grouped value heads and differing key/value dimensions when indexing or shapes change.\n- Initial states, intermediate states, and context parallelism when their paths are touched.\n- Accepted and rejected backend routes when changing verifiers.\n\nFor gate or decay changes, include the lower-bound limits, saturated gate inputs, extreme gate scales, long cumulative decay, and chunk or ragged-sequence boundaries. Test through the public entry point so the selected route and effective arguments are exercised.\n\n<!-- File: .agents/skills/fla-nvidia-performance/SKILL.md (Tokens: ~708 | Category: anthropic-skill) -->\n---\nname: fla-nvidia-performance\ndescription: Profile and optimize FLA kernels on NVIDIA GPUs, with same-hardware benchmarks and targeted Nsight Compute analysis.\n---\n\n# NVIDIA kernel performance\n\nUse this skill for NVIDIA kernel optimization and backend tuning. Follow [CONTRIBUTING](../../../CONTRIBUTING.md#benchmarking) for correctness and performance evidence, and [fla-optimization-loop](../fla-optimization-loop/SKILL.md) for sustained iteration.\n\n## Hardware baseline\n\nUse datacenter NVIDIA GPUs with sm_90 or newer for PR performance conclusions; H100 and H20 are accepted, and sm_100/sm_103 are preferred. Results from A100, older GPUs, or consumer cards are supplementary. Record the exact GPU and software versions with each comparison.\n\n## Workflow\n\n1. Identify the changed kernel, its callers, and the production workloads it affects. Record a baseline commit and confirm which backend runs.\n2. Pass the relevant correctness tests before timing a candidate. Preserve the reference and tolerances; include outputs, states, and supported gradients.\n3. Compare baseline and candidate on the same GPU with identical workloads and dispatch settings. Include dense and supported variable-length cases. Use an end-to-end benchmark for a training or generation throughput claim.\n4. Use Nsight Compute when a bottleneck or regression needs hardware metrics. Match the optimization to the measured limit, then rerun correctness and timing.\n5. Report latency or throughput, relevant memory changes, and any regressions. Keep raw profiler artifacts outside git, for example under `profile/<run_name>/`.\n\n## Benchmark commands\n\nRun from the repository root. Resolve the baseline to a commit SHA before comparing:\n\n```bash\nFLA_BENCH_BASE=$(git rev-parse origin/main)\nFLA_CI_ENV=0 python -m benchmarks.ops.verify --op chunk_kda --base \"$FLA_BENCH_BASE\"\n\npython benchmarks/benchmark_training_throughput.py --name kda --batch_size 2 --seq_len 8192\npython benchmarks/benchmark_training_throughput.py --name kda --batch_size 2 --seq_len 8192 --varlen\n```\n\nThe verification command runs the current tests; retain the validated baseline and tests yourself. A selected subset is useful during iteration but does not replace the full relevant gate before reporting a gain.\n\n## Nsight Compute\n\nCollect a representative changed kernel with the installed `ncu`. Check which sections and metrics that version supports.\n\nThe benchmark runner launches a child process, so use `--target-processes all`. Replace the kernel filter with the changed kernel's name. `--no-base` profiles the current implementation; collect the baseline separately with the same command and workload.\n\n```bash\nmkdir -p profile/kda\nncu --target-processes all --set full \\\n  -k 'regex:chunk_kda' -c 1 -o profile/kda/full \\\n  python -m benchmarks.ops.run --op chunk_kda --modes fwd --no-base\n\nncu --target-processes all --set source --section SourceCounters \\\n  -k 'regex:chunk_kda' -c 1 -o profile/kda/source \\\n  python -m benchmarks.ops.run --op chunk_kda --modes fwd --no-base\n```\n\nUse duration, memory throughput, occupancy, stalls, and hot instructions to explain the bottleneck. Profiling overhead makes these runs unsuitable as the latency comparison; use the synchronized benchmark for that measurement.\n\n<!-- File: .agents/skills/fla-optimization-loop/references/opt-log-template.md (Tokens: ~220 | Category: anthropic-skill) -->\n# Optimization Record\n\nKeep one local log under ignored `profile/<op>-opt/`. Record the entry point, target workloads, baseline SHA and correctness, environment, commands, numeric settings, and success criterion once. Preserve enough candidate identity to reproduce each measurement.\n\n| Attempt  | Change or hypothesis | Gate | Latency by workload | Decision and reason |\n| -------- | -------------------- | ---- | ------------------- | ------------------- |\n| Baseline | Unchanged revision   | Pass | Measured baseline   | Comparison point    |\n\nAppend an entry after each attempt, including failures and unchanged results. Identify whether the gate was full or selected; record no performance result for a failed gate. Retained candidates need their revision or recoverable diff, the reason they helped, and any unresolved limitation. Keep failed directions here so later work does not repeat them.\n\nWhen specializing, add each routing condition, entry point, baseline/candidate latency, and measured reason for a separate path. Compare every bucket against the same baseline revision on that bucket's workloads, including fallback boundaries. Keep these details alongside the iteration results.\n\n<!-- File: .agents/skills/fla-optimization-loop/references/TRAPS.md (Tokens: ~487 | Category: anthropic-skill) -->\n# Correctness and Measurement Traps\n\n## Partial writes\n\n`tests/conftest.py` poisons eligible FLA `empty` allocations in operator/module tests with NaNs. It excludes some allocations and does not cover layer/model/CP tests. A poisoned-run failure that disappears in an ad hoc script can reveal uninitialized reads; inspect masks and ensure every exposed output element is written. NaNs can also have numerical causes, so diagnose the failure rather than assuming its source.\n\n## Reference precision and tolerance\n\n`fla.utils.assert_close` checks relative RMS error with an absolute-error shortcut, not elementwise `rtol`. `FLA_CI_ENV=1` and explicit `warning=True` can downgrade tolerance failures to warnings. Use strict local validation and inspect warnings; keep the committed thresholds fixed.\n\nTF32 in a reference matmul can change an fp32 comparison. Inspect the reference's precision settings before blaming the kernel, but do not change the frozen oracle to pass. Baseline and candidate need identical precision and numeric flags. An intentional precision change follows the numerical-change process in [CONTRIBUTING.md](../../../../CONTRIBUTING.md#submit-pull-requests).\n\n## Address overflow\n\nCast program IDs and offsets to `tl.int64` before multiplication; widening an already-overflowed product cannot repair it. Small shapes may hide the error. Run large-offset and affected varlen/grid cases under the [Triton addressing rules](../../../../CONTRIBUTING.md#triton-kernels).\n\n## Missing work in a fast result\n\nCheck unexpectedly large speedups against the actual selected path, executed tests, and measured shapes. `verify.py` accepts pytest's exit status, including an all-skipped success. The benchmark runner filters constrained shapes and omits failed warmups or measurements. Compare expected and returned rows before making a claim; a passing process alone is insufficient.\n\n## Autotune and clocks\n\n`run.py`, also used by `verify.py`, warms all shapes before timing. Preserve that warmup. After changing the configuration space, check cached selections and rewarm or clear the relevant cache if results look stale.\n\nUnlocked clocks and changing device load can shift latency between runs. Rank local candidates by their own runtime, then remeasure the final candidate against the recorded baseline in one session. Report variability and the device state; speedups from separate experiments do not add up to a final result.\n\n<!-- File: .agents/skills/fla-optimization-loop/SKILL.md (Tokens: ~966 | Category: anthropic-skill) -->\n---\nname: fla-optimization-loop\ndescription: Iterate on FLA kernel performance with a frozen correctness gate, a measured baseline, and reproducible candidate comparisons.\n---\n\n# FLA Optimization Loop\n\nUse this skill for repeated kernel performance work. Define supported workloads and numerical assumptions with [fla-design-coverage](../fla-design-coverage/SKILL.md); use the [NVIDIA](../fla-nvidia-performance/SKILL.md) or [Ascend](../fla-ascend-performance/SKILL.md) skill for profiling on that platform.\n\n## Freeze the gate and baseline\n\nAdd missing regression or new-path coverage before starting the loop. Then keep the relevant tests, reference implementation, input distributions, assertions, tolerances, and public signature fixed. Do not narrow parameter matrices, add skips, relax precision or thresholds, special-case test values, cache results across trials, or change timing behavior to make a candidate pass. Fully initialize exposed outputs. Use kernel implementations for the optimization; replacing the whole operator with a vendor call solely to win the benchmark does not meet this workflow.\n\nA suspected oracle or tolerance defect needs a separately justified correction, not an adjustment that makes the candidate pass. Follow [CONTRIBUTING.md](../../../CONTRIBUTING.md#submit-pull-requests) for precision, numerical-algorithm, and tolerance changes, and its [compatibility policy](../../../CONTRIBUTING.md#public-api-compatibility) for public behavior.\n\nRecord the entry point, target workloads, allowed implementation languages, validation/benchmark commands, baseline SHA, and success criterion in the local optimization log. Include production inputs from callers as well as applicable benchmark-registry shapes. Confirm the baseline passes the full gate and measure it before changing kernels. For a new kernel, establish correctness before tuning.\n\n```bash\nFLA_BENCH_BASE=$(git rev-parse origin/main)\nFLA_CI_ENV=0 python -m benchmarks.ops.verify --op chunk_gla --base \"$FLA_BENCH_BASE\"\n```\n\nUse a recorded SHA: the baseline worktree cannot reuse a checked-out branch, and slash-containing refs break its temporary directory naming. `--base` benchmarks that revision; it does not run its correctness tests. Establish baseline correctness separately if the current tree already differs.\n\nThe driver runs the current pytest file and stops before timing when pytest fails. It neither freezes files nor rejects an all-skipped run. Check that the expected cases and benchmark rows ran. `--test-file` overrides a mismatched derived test path; `--gate-k` selects a quick subset, while `--no-gate` produces unverified measurements. Neither subset nor skipped-gate results qualify for promotion.\n\n## Iterate from evidence\n\n1. Profile the current bottleneck and choose one reasoned change.\n2. Run the frozen gate, then benchmark only a passing candidate. A subset can provide an early signal; the final gate covers the full affected suite.\n3. Record the change, gate result, measurement, and keep/drop decision before the next experiment. Preserve a recoverable candidate under the repository's git rules and the user's authorization; failed and unchanged attempts need records, not empty commits.\n\nRead [measurement and correctness traps](references/TRAPS.md) before interpreting results. Keep seeds, numeric flags, hardware, and workload fixed. Candidate runtime helps rank local iterations; measure the final candidate directly against the frozen baseline in one session rather than adding speedups from separate runs.\n\nAdd specialized paths only when profiling shows different bottlenecks across shape regimes and measurements justify the extra routing. Record each condition, chosen path, and before/after result, and validate the full shape set including fallback boundaries.\n\nAfter repeated attempts without an improvement above noise, re-profile and review the failed directions before continuing. Respect the user's stopping condition. A no-win conclusion needs a measured baseline, a reasoned candidate attempt, gate status, benchmark evidence, and an identified bound or blocker. An unavailable profiler or one losing candidate alone does not establish a performance floor; report unresolved validation as a limitation.\n\n## Keep and report the result\n\nKeep iteration notes and raw traces under ignored `profile/<op>-opt/`; [opt-log-template.md](references/opt-log-template.md) provides a compact record format. This search history stays out of kernel comments and the final PR description.\n\nPromote only with a full green gate, a repeatable win on the claimed workloads, and profiling or a quantified bottleneck analysis that explains it. Report per-shape results and regressions, commands, baseline/candidate revisions, environment, and measurement variability using the relevant hardware skill. An unexplained large speedup needs investigation. Use [fla-pr-readiness](../fla-pr-readiness/SKILL.md) to prepare the final diff and evidence.\n\n<!-- File: .agents/skills/fla-pr-readiness/SKILL.md (Tokens: ~746 | Category: anthropic-skill) -->\n---\nname: fla-pr-readiness\ndescription: Prepare or update an FLA pull request with focused scope, verified evidence, and the current repository template.\n---\n\n# FLA PR Readiness\n\nRead [CONTRIBUTING.md](../../../CONTRIBUTING.md) and [AGENTS.md](../../../AGENTS.md) first. Use the current [PR template](../../../.github/pull_request_template.md) for headings and checkboxes; contribution policy and template wording are not duplicated here.\n\n## Review scope and evidence\n\nRead the complete diff against the PR base and audit affected shared callers. State its concrete outcome in one sentence and separate independent work. Check existing issues/PRs and the applicable design discussion. Follow [scope and size](../../../CONTRIBUTING.md#pr-scope-and-size) for a cohesive change above the normal limit, and [public compatibility](../../../CONTRIBUTING.md#public-api-compatibility) when interfaces or defaults are affected.\n\nWhen a change needs splitting or dependent PRs need a native GitHub stack, use [PR stacks](../fla-pr-stacks/SKILL.md) to choose boundaries, reuse existing PRs, and verify the published stack.\n\nFor code changes, discover dependent tests and run the relevant targets. For example:\n\n```bash\npython scripts/find_dependent_tests.py fla/ops/kda/chunk.py\nFLA_CI_ENV=0 pytest tests/ops/test_kda.py\npre-commit run --files fla/ops/kda/chunk.py tests/ops/test_kda.py\n```\n\nFor documentation-only changes, validate links, examples, and formatting as described in [Evidence by Change Type](../../../CONTRIBUTING.md#evidence-by-change-type).\n\nThe dependency helper only lists tests. Check that regression coverage exercises the changed path and that required numerical comparisons passed, including supported gradients. Review new test functions and files against [test organization and naming](../../../CONTRIBUTING.md#test-organization-and-naming); use an existing test case where it fits. Record commands, tested revision, environment, and results; distinguish failures, skips, and unrun checks. Reproduce a failure on the baseline before calling it pre-existing.\n\nFor performance-relevant changes, collect same-hardware before/after measurements under [Benchmarking](../../../CONTRIBUTING.md#benchmarking) and the relevant hardware skill. A neutral result needs numbers; explain inapplicable checks. Inspect the final diff for stale comments and unrelated changes, leaving scratch logs and raw profiles outside it.\n\n## Prepare and publish\n\nWrite for a reviewer who has not seen the conversation. Follow [PR Description](../../../CONTRIBUTING.md#pr-description) and [Evidence by Change Type](../../../CONTRIBUTING.md#evidence-by-change-type): explain the problem, resulting behavior, affected users, and evidence. Include only useful rationale or limitations; omit file inventories and the history of attempts. Keep incomplete work in draft and leave claims and checklist items consistent with actual validation.\n\nThe `check-pr-title` workflow also checks the body and size exception on edits. Preserve template checklist wording and distinguish required items from conditional ones.\n\n`gh pr edit` fails with this repository's classic-Projects integration. Write the complete body to a temporary file and update through REST:\n\n```bash\ngh api -X PATCH \"repos/fla-org/flash-linear-attention/pulls/$pr_number\" -F \"body=@$pr_body_file\"\n```\n\nConfirm the published description, head revision, and checks. After review changes, rerun affected validation and update evidence that no longer describes the current code. Publishing and git operations remain subject to the user's authorization and repository rules.\n\n<!-- File: .agents/skills/fla-pr-stacks/SKILL.md (Tokens: ~1003 | Category: anthropic-skill) -->\n---\nname: fla-pr-stacks\ndescription: Split oversized FLA pull requests into focused changes, reuse existing PRs, and create or maintain verified GitHub native PR stacks.\n---\n\n# FLA PR Stacks\n\nUse for decomposing large changes or managing dependent PRs. Follow [CONTRIBUTING.md](../../../CONTRIBUTING.md#pr-scope-and-size) for scope and size, [AGENTS.md](../../../AGENTS.md) for git permissions, and [PR readiness](../fla-pr-readiness/SKILL.md) for each PR's evidence and description.\n\n## 1. Choose the split\n\nInspect the complete diff, commit history, and related open PRs, including their review discussions. Map each proposed layer to one outcome, its prerequisites, validation, and an existing PR to reuse where appropriate.\n\n| Relationship                            | Action                                                                   |\n| --------------------------------------- | ------------------------------------------------------------------------ |\n| Overlaps an existing PR's purpose       | Update that PR when ownership permits; preserve attribution and reviews. |\n| Independent of the other changes        | Submit separately against the trunk.                                     |\n| Depends on another layer                | Stack above that layer; include its own tests and documentation.         |\n| Cohesive but still above the size limit | Use the contribution guide's exception with a concrete justification.    |\n\nEach layer presents an increment over its parent and must work with its prerequisites, without needing later layers to fix it. Separate review does not imply dependency-free code. Do not split by arbitrary line counts, leave callers broken, or defer tests to a later PR.\n\nCheck head repository identity before arranging existing PRs. [GitHub stacks require all branches in the same repository](https://docs.github.com/en/pull-requests/reference/stacked-pull-requests); cross-fork stacks are unsupported. Keep incompatible PRs as explicit external prerequisites instead of duplicating them under new PR numbers.\n\n## 2. Build and validate\n\nPreserve the original head and uncommitted work before reconstructing commits in isolated branches or worktrees. Reuse suitable commits and preserve authorship. Check current remote heads before updating existing PRs.\n\nReview each layer against its parent, then compare the top branch with the preserved original result:\n\n```bash\ngit diff main...refactor/foundation\ngit diff refactor/foundation...refactor/consumer\ngit diff original-head refactor/consumer\n```\n\nSubstitute the recorded trunk and actual refs. Account for trunk changes and intentional omissions, including fixes retained in existing PRs. Run relevant tests at each layer's head and on the combined result; passing only the top branch does not validate intermediate PRs.\n\n## 3. Register and verify the native stack\n\nChanging PR bases alone does not create a native stack. Use the official [`gh stack` extension](https://docs.github.com/en/pull-requests/how-tos/create-pull-requests/creating-stacked-pull-requests), GitHub UI, or REST Stacks API. Check installed command help and existing publishing authorization first.\n\nFor prepared branches, adopt them bottom to top and submit:\n\n```bash\ngh stack init --base main refactor/foundation refactor/consumer\ngh stack submit --auto --remote origin\n```\n\nChoose the actual remote containing the branches. `submit` pushes, creates or updates PRs, and registers the stack. `--auto` creates new PRs as drafts; replace generated titles and bodies using PR readiness.\n\nFor existing PRs, use `gh stack link --base main --remote origin 123 124` with their real numbers in dependency order. This reuses PRs without requiring local stack tracking. Branch arguments also push and create missing PRs. After an error, inspect what succeeded before retrying.\n\nFor every PR, verify [remote stack membership](https://docs.github.com/en/pull-requests/reference/stacked-pull-requests-apis-and-webhooks):\n\n```bash\ngh api \"repos/fla-org/flash-linear-attention/pulls/$pr_number\" \\\n  --jq '{number, head_repo: .head.repo.full_name, head: .head.ref, base: .base.ref, stack}'\n```\n\nRequire non-null `stack` objects with the same stack identity, correct order, size, and trunk. Check each PR's incremental diff and report the verified links and remaining external prerequisites. A local branch chain or successful push is insufficient.\n\n## 4. Maintain the series\n\nAfter changing a lower layer, review affected upper diffs and rerun relevant tests. `gh stack rebase` rewrites history; `gh stack sync` also force-pushes with leases. Apply the repository's rewrite permissions and preserve concurrent remote work, then recheck published membership and evidence.\n\nCreating a series does not authorize merging or closing PRs. If merging is requested, establish the intended prefix: selecting an upper PR also merges those below it. Use a stack-aware merge operation; the legacy synchronous merge API does not support stacks.\n\n<!-- File: .agents/skills/fla-triton-to-gluon/SKILL.md (Tokens: ~1257 | Category: anthropic-skill) -->\n---\nname: fla-triton-to-gluon\ndescription: Port an FLA Triton kernel to Gluon when explicit layouts, asynchronous transfers, or scheduling can address a measured bottleneck.\n---\n\n# Triton to Gluon\n\nUse Gluon when profiling points to register pressure, layout conversion, or load/compute overlap that needs explicit control. A kernel already saturating its useful bandwidth or compute limit may not benefit. Keep the portable default path when adding hardware-specific implementations.\n\nFollow [fla-optimization-loop](../fla-optimization-loop/SKILL.md) for the baseline and correctness gate, and [fla-nvidia-performance](../fla-nvidia-performance/SKILL.md) for NVIDIA measurements.\n\n## Check the target environment\n\nGluon is experimental. Check the installed Triton API before copying [upstream examples](https://triton-lang.org/main/getting-started/tutorials/gluon/); their names and signatures can differ from the installed release.\n\nGate architecture-specific code appropriately: Ampere supports `cp.async`, Hopper adds TMA/WGMMA, and Blackwell adds TMEM/tcgen05. Code under Gluon's `nvidia` namespace requires NVIDIA hardware. Keep compile-time choices in explicit constexpr arguments so the kernel, launcher, and autotune pruning use the same values.\n\n```python\nfrom triton.experimental import gluon\nfrom triton.experimental.gluon import language as gl\n```\n\n## Port incrementally\n\n1. Record the baseline commit and retain the validated tests, reference, and tolerances. Keep the Triton fallback available.\n2. Translate loads, stores, and arithmetic with explicit layouts first. Match each tensor's contiguous dimension and preserve masking. Pass forward, state, and supported backward comparisons before adding asynchrony.\n3. Add async transfers only where load latency or addressing pressure limits the kernel. Define buffer ownership, completion, and reuse for the prologue, steady state, and drain.\n4. Add architecture-specific MMA or scheduling when profiling supports it. Recheck register/shared-memory use and retune after each change.\n5. Run the full relevant correctness gate and same-hardware benchmarks, including supported variable-length and boundary cases. Measure production entry points as well as individual kernels.\n\nA literal translation establishes parity; measure it before assuming a performance gain. Preserve the validated numerical algorithm and precision while changing execution layout or scheduling.\n\n## Layout and resource constraints\n\n- Pointer tensors and values need compatible layouts. A cross-warp `gl.convert_layout` may use shared memory; use `assert_trivial=True` when a conversion is intended to be free.\n- Broadcasts and oversized layout tiles can duplicate values and increase register use. Choose layouts from actual tensor strides rather than copying a tutorial's block sizes.\n- Account for all live shared-memory buffers and the target device's allocation limit. Prune infeasible autotune configurations; provide a streaming path when a resident design cannot fit supported shapes.\n- Large static unrolls multiplied by many autotune configurations can dominate compilation. Prune from the actual shape and live-memory budget, and reuse the worker and Triton cache during iteration.\n- Persistent schedules can trade launch overhead for worse L2 locality. Warp specialization also changes the total register budget; measure occupancy and stalls after changing either.\n\n## Async correctness\n\n### Transfers and buffer reuse\n\n`cp.async` groups are ordered as a queue. `wait_group(N)` bounds all outstanding groups, so it cannot identify a particular buffer after unrelated prefetches are interleaved. For per-buffer scheduling, use separate mbarriers and match their arrival counts to the participating threads. In the installed API, check whether `mbarrier_arrive` increments the count; a preinitialized thread count requires the non-incrementing form.\n\nWait for a load before consuming its buffer, and protect reuse until all readers finish. Even same-lane shared-memory staging needs protection against overwriting data still being read. Track each mbarrier's phase as its buffer is reused; do not run more than one phase ahead or share a completion barrier between TMA and tcgen05 without reinitializing it.\n\nTMA descriptors must satisfy the target's alignment and stride requirements. TMA stores may distinguish completion of the shared-memory read from completion of the global write. Check the installed wait API and require global completion before another operation reads the stored range.\n\n### MMA and memory ordering\n\nHopper WGMMA requires its B operand in shared memory and uses register accumulators; consume the values returned by its wait operation so compiler dependencies remain explicit. Blackwell tcgen05 uses TMEM accumulators and mbarrier completion; respect the participating warpgroup and TMEM layout requirements. Initialize accumulators explicitly, including `use_acc=False` where supported.\n\nGeneric shared-memory accesses and async operations use different memory proxies. Apply the required proxy fence when handing a buffer to an async consumer; an mbarrier alone does not replace that fence. A completed TMA load establishes the ordering needed to read its destination. Check ordering across warp-specialized partitions as well as within a single producer/consumer loop.\n\n### Tails and exact-zero cases\n\nMasked async copies can leave shared-memory elements uninitialized. Prevent invalid rows from entering reductions: zero-fill where supported, or use valid-row loads and explicitly remove every invalid contribution. Multiplying by zero does not neutralize a NaN. Keep output stores masked.\n\nDo not rely on separately compiled reductions cancelling bitwise. For a mathematically exact-zero special case, such as a single-source softmax gradient, preserve the exact-zero result explicitly and test it against the reference.\n\n## Verification commands\n\nRun from the repository root on the target GPU:\n\n```bash\nFLA_BENCH_BASE=$(git rev-parse origin/main)\nFLA_CI_ENV=0 python -m benchmarks.ops.verify --op chunk_kda --base \"$FLA_BENCH_BASE\"\n```\n\nReplace the operation with the one being ported. The command executes the current tests; it does not freeze them. Use `--gate-k` only for quick iteration, then run the full gate. Set backend environment flags before launching a fresh process so import-time choices cannot contaminate the comparison.\n\n<!-- File: .agents/skills/README.md (Tokens: ~507 | Category: anthropic-skill) -->\n# Repository skills\n\n[CONTRIBUTING.md](../../CONTRIBUTING.md) defines contribution rules; [AGENTS.md](../../AGENTS.md) defines agent operations. Skills add task-specific guidance. Choose those relevant to the change.\n\nFor kernel work, define the contract and establish baseline coverage before optimization. Use the operator, dispatch, and hardware guides as needed, then prepare the PR with the collected evidence.\n\n| Skill                                                     | Use for                                                    |\n| --------------------------------------------------------- | ---------------------------------------------------------- |\n| [Design coverage](fla-design-coverage/SKILL.md)           | Supported inputs, numerical budgets, and routing contracts |\n| [Correctness coverage](fla-correctness-coverage/SKILL.md) | References, regression cases, and affected test selection  |\n| [KDA](fla-kda/SKILL.md)                                   | KDA gate modes and kernel invariants                       |\n| [Backend dispatch](fla-dispatch-backends/SKILL.md)        | Registration, verifiers, and fallback behavior             |\n| [Optimization loop](fla-optimization-loop/SKILL.md)       | A fixed correctness gate and measured candidate iterations |\n| [NVIDIA performance](fla-nvidia-performance/SKILL.md)     | GPU profiling and bottleneck diagnosis                     |\n| [Ascend performance](fla-ascend-performance/SKILL.md)     | NPU profiling and compiler constraints                     |\n| [Triton to Gluon](fla-triton-to-gluon/SKILL.md)           | Layouts, synchronization, and incremental kernel ports     |\n| [PR readiness](fla-pr-readiness/SKILL.md)                 | Scope, validation evidence, and submission                 |\n| [PR stacks](fla-pr-stacks/SKILL.md)                       | Splitting large changes, reusing PRs, and native stacks    |\n\n## Maintaining skills\n\nPut each workflow in `.agents/skills/<skill-name>/SKILL.md` with YAML `name` and `description` fields. Follow [documentation and skill style](../../CONTRIBUTING.md#prose-and-markdown) for scope, organization, wording, and shared policy.\n\nUse an existing `references/` directory for substantial task-specific detail. Add supporting files only when they make the workflow easier to use; `SKILL.md` is the entry point, so individual skill directories do not need a README. Symlinks may point to public documentation tracked in this repository.\n\n<!-- File: evals/README.md (Tokens: ~1996 | Category: anthropic-skill) -->\n# Evaluation\n\nWe provide evaluation scripts for `fla` models with [lm-evaluation-harness](https://github.com/EleutherAI/lm-evaluation-harness), as well as a standalone perplexity evaluator.\nInstall `fla` following the [installation guide](../INSTALL.md), and run the commands below from the repository root.\n\n* [lm-evaluation-harness](#lm-evaluation-harness)\n  * [Single-GPU Evaluation](#single-gpu-evaluation)\n  * [Multi-GPU Evaluation](#multi-gpu-evaluation)\n  * [RULER Benchmarks](#ruler-benchmarks)\n* [Perplexity](#perplexity)\n  * [Quick Start](#quick-start)\n  * [Evaluation Settings](#evaluation-settings)\n  * [Scoring](#scoring)\n  * [Results](#results)\n\n## lm-evaluation-harness\n\nThe [lm-evaluation-harness](https://github.com/EleutherAI/lm-evaluation-harness) library allows you to easily perform (zero-shot) model evaluations.\nInstall `lm_eval` following [their instructions](https://github.com/EleutherAI/lm-evaluation-harness/blob/main/README.md).\n\n### Single-GPU Evaluation\n\nRun evaluation with:\n\n```sh\nMODEL='fla-hub/gla-1.3B-100B'\npython -m evals.harness --model hf \\\n    --model_args pretrained=$MODEL,dtype=bfloat16 \\\n    --tasks wikitext,lambada_openai,piqa,hellaswag,winogrande,arc_easy,arc_challenge,boolq,sciq,copa,openbookqa \\\n    --batch_size 64 \\\n    --num_fewshot 0 \\\n    --device cuda \\\n    --show_config\n```\n\nWe've made `fla` compatible with hf-style evaluations, you can call [evals.harness](harness.py) to finish the evaluations.\nThe command above evaluates the GLA checkpoint on the tasks used in the GLA paper.\n\n### Multi-GPU Evaluation\n\nUse the Hugging Face Accelerate launcher for data-parallel evaluation, where each GPU loads a full copy of the model:\n\n```sh\nMODEL='fla-hub/gla-1.3B-100B'\naccelerate launch -m evals.harness --model hf \\\n    --model_args pretrained=$MODEL,dtype=bfloat16,trust_remote_code=True \\\n    --tasks wikitext,lambada_openai,piqa,hellaswag,winogrande,arc_easy,arc_challenge,boolq,sciq,copa,openbookqa \\\n    --batch_size 64 \\\n    --num_fewshot 0 \\\n    --device cuda \\\n    --show_config \\\n    --trust_remote_code\n```\n\n### RULER Benchmarks\n\nThe RULER benchmarks are commonly used for evaluating model performance on long-context tasks.\nYou can evaluate `fla` models on RULER directly using `lm-evaluation-harness`. RULER is only available in a relatively recent version of `lm-evaluation-harness`, so make sure you have the latest version installed.\n\n```sh\ngit clone --depth 1 https://github.com/EleutherAI/lm-evaluation-harness\npip install -e './lm-evaluation-harness[ruler]'\n```\n\nFor example, to evaluate contexts up to 32K tokens:\n\n```sh\nMODEL='fla-hub/gla-1.3B-100B'\nOUTPUT='results/ruler'\naccelerate launch -m evals.harness --model hf \\\n    --output_path $OUTPUT \\\n    --tasks niah_single_1,niah_single_2,niah_single_3,niah_multikey_1,niah_multikey_2,niah_multikey_3,niah_multiquery,niah_multivalue,ruler_vt,ruler_cwe,ruler_fwe,ruler_qa_hotpot,ruler_qa_squad \\\n    --model_args pretrained=$MODEL,dtype=bfloat16,max_length=32768,trust_remote_code=True \\\n    --metadata='{\"max_seq_lengths\":[4096,8192,16384,32768]}' \\\n    --batch_size 2 \\\n    --show_config \\\n    --trust_remote_code\n```\n\nIf a GPU can't load a full copy of the model, please refer to [this link](https://github.com/EleutherAI/lm-evaluation-harness?tab=readme-ov-file#multi-gpu-evaluation-with-hugging-face-accelerate) for FSDP settings.\n\n> [!TIP]\n> If you are using `lm-evaluation-harness` as an external library and can't find (almost) any tasks available, before calling `lm_eval.evaluate()` or `lm_eval.simple_evaluate()`, simply run the following to load the library's stock tasks:\n> ```py\n> from lm_eval.tasks import TaskManager\n> TaskManager().initialize_tasks()\n> ```\n\n## Perplexity\n\n[evals.ppl](ppl.py) evaluates next-token perplexity over fixed-length contexts and reports perplexity by position within each context.\n\n### Quick Start\n\nInstall the dataset dependency:\n\n```sh\npip install datasets\n```\n\nFor example, to evaluate the GLA checkpoint on the PG19 test split and save the results:\n\n```sh\npython -m evals.ppl \\\n    --path fla-hub/gla-1.3B-100B \\\n    --data fla-hub/pg19 --split test \\\n    --block_size 28672 --bucket_size 2048 --batch_size 1 \\\n    --dtype bfloat16 --num_proc 4 \\\n    --output results/ppl.json\n```\n\n### Evaluation Settings\n\n| Argument          | Default                 | Description                                                 |\n|-------------------|-------------------------|-------------------------------------------------------------|\n| `--path`          | `fla-hub/gla-1.3B-100B` | Hugging Face model ID or local checkpoint directory.        |\n| `--data`          | `fla-hub/pg19`          | Dataset to evaluate.                                        |\n| `--split`         | `train`                 | Dataset split; set this explicitly for held-out evaluation. |\n| `--column_name`   | `text`                  | Dataset column containing text.                             |\n| `--block_size`    | `28672`                 | Context length in tokens; must be at least 2.               |\n| `--bucket_size`   | `2048`                  | Number of prediction positions per reported bucket.         |\n| `--batch_size`    | `1`                     | Number of full contexts evaluated together.                 |\n| `--device`        | Auto-detected           | Device on which to load and evaluate the model.             |\n| `--dtype`         | `bfloat16`              | Model dtype: `float32`, `float16`, or `bfloat16`.           |\n| `--num_proc`      | `32`                    | Number of dataset tokenization processes.                   |\n| `--revision`      | `main`                  | Model and tokenizer revision.                               |\n| `--data_revision` | `main`                  | Dataset revision.                                           |\n| `--output`        | None                    | Optional path for JSON results and evaluation settings.     |\n\nUse immutable commit hashes for `--revision` and `--data_revision` to reproduce a run.\nFor long contexts, reduce `--batch_size` to lower model memory use and `--bucket_size` to reduce the FP32 softmax workspace.\nThe model still materializes its full logits.\n\n### Scoring\n\nDocuments are tokenized with the tokenizer's default special-token policy, then concatenated in dataset order without extra separators.\nThe token stream is divided into independent contexts of `block_size` tokens, and the final partial context is evaluated separately without padding.\nContext is reset at every block boundary; there is no sliding window or recurrent state carried between blocks.\n\nEach context of length `L` contributes `L-1` next-token predictions.\nThe first token is unscored, and no synthetic EOS target is appended after the last token.\nEOS tokens present in the input are scored normally.\nA one-token tail contributes no predictions; a dataset with fewer than two tokens raises `ValueError`.\n\nOverall PPL and positional PPL use the same FP32 negative log-likelihoods, weighted by the number of scored predictions.\nBuckets group predictions by their zero-based logit position within each context: bucket 0 covers positions `0` through `bucket_size-1`.\nModels configured with fused linear cross entropy are supported, and evaluation runs with caching disabled.\n\n### Results\n\nThe script prints overall PPL, positional PPL, and token counts.\nWith `--output`, it also saves a JSON file containing:\n\n* `perplexity` and `block_perplexities`: overall and positional PPL; buckets with no predictions have `null` PPL.\n* `total_nll`, `block_nlls`, and `block_tokens`: summed NLLs and per-bucket prediction counts.\n* `total_tokens` and `total_input_tokens`: scored predictions and input tokens, respectively.\n* `total_sentences`: the number of evaluated contexts, including any one-token tail.\n* `config`, `model_commit`, `dataset_fingerprint`, and `versions`: evaluation settings, the resolved model commit when available, dataset fingerprint, and package versions.\n\n> [!NOTE]\n> Earlier versions appended a synthetic EOS only to positional statistics, used BF16 softmax for BF16 logits, and dropped the final partial context. Scores and token counts from those versions may differ from the corrected results.\n</yakaai_skills>"}