flash-linear-attention

Guidelines for Ascend NPU kernel / Triton-Ascend backend performance work in the FLA repo. Covers profiling with torch_npu, PipeUtilization/MemoryUB CSV analysis, Cube/Vector/MTE/UB bottleneck diagnosis, and kernel optimization (UB tiling, grid splits, fusion/split, varlen, G_T_CONTIG gate loading, constexpr DMA-path split / TAIL_MODE, extract_slice, MTE OOB, int32 address overflow, tl.cast vs constexpr .to, make_block_ptr int32 offsets, correctness gates, tl.dot left-operand clobber). NPU kernels must not use num_warps/num_stages. Per-kernel catalog: references/cases.md (incl. causal_conv1d core-grid). Use when working on NPU profiling, kernel_details/op_statistic, aic_metrics, fla triton_ascend backends (ops or modules), g transpose stride-1, UB overflow, dual-path DCE, grid limits, int64 pointer math, tl.dot reuse, or Ascend performance.

Guidelines for kernel correctness testing and coverage in fla/ops/** and related modules, including common Triton grid/addressing pitfalls. Helps decide what tests to add or run before an MR.

Contract-first design and coverage discipline for FLA kernel and numerical changes. Use before implementation to define supported cells, numerical budgets, dispatch semantics, compatibility, tests, and benchmarks.

Workflow for FLA backend dispatch decorators and backend implementations. Use when touching fla.ops.backends, @dispatch-decorated functions, BaseBackend subclasses, backend verifier methods, backend env vars, or backend tests.

FLA KDA kernel workflow and public technical notes. Use when modifying or reviewing fla/ops/kda/**, KDA gate modes, chunk intra/inter kernels, safe_gate behavior, KDA backends, or KDA-specific tests and benchmarks.

Checklist and workflow for preparing an MR/PR in the FLA repo. Covers CONTRIBUTING.md compliance, test plan, benchmark evidence, and PR body structure.

Guidelines for NVIDIA GPU kernel / Triton / Gluon / TileLang / CUDA backend performance work in the FLA repo. Covers profiling workflow, hardware baselines, and MR-ready performance evidence requirements. Uses an installed ncu-report-skill when a task needs detailed Nsight Compute collection and diagnosis.

Disciplined, reproducible loop for making an FLA kernel faster (Triton, Gluon, TileLang, CuTe) without ever breaking or gaming correctness. Synthesizes the task-contract / three-phase / iteration-protocol / silent-bug-catalog discipline of agent kernel-optimization frameworks (KDA, the MLSys FlashInfer contest workflow, AKO4ALL/AKO4X), and anchors all of it on FLA's frozen pytest (forward AND backward, under NaN poisoning) as the immutable correctness gate. Use when iterating on `fla/ops/**` performance over multiple rounds.

Workflow for porting an existing Triton kernel in `fla/ops/**` to Gluon (`triton.experimental.gluon`) to gain explicit control over tensor layouts, shared memory, async data movement (cp.async / TMA), MMA (WGMMA / tcgen05), and scheduling (persistent kernels, warp specialization). Covers when a port is worth it, an incremental porting sequence that keeps numerical parity at every step, a Triton-to-Gluon API mapping, compile-time / autotune / smem-budget management for heavily unrolled kernels, and a pitfall checklist (proxy fences, mbarrier semantics, layout costs, bitwise-cancellation traps, NaN-poisoned OOB handling). Use when a Triton kernel is register-bound, when `num_stages` pipelining underperforms, or when Hopper/Blackwell features (TMA, TMEM, tcgen05) are needed.