fla-org
- 9 skills
- 0 followers
- 7 hours ago last updated
- ▌ Fla Kda · fla-orgFLA KDA kernel workflow and public technical notes. Use when modifying or reviewing fla/ops/kda/**, KDA gate modes, chunk intra/inter kernels, safe_gate behavior, KDA backends, or KDA-specific tests and benchmarks.
- ▌ Fla Mr Readiness · fla-orgChecklist and workflow for preparing an MR/PR in the FLA repo. Covers CONTRIBUTING.md compliance, test plan, benchmark evidence, and PR body structure.
- ▌ Fla Design Coverage · fla-orgContract-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.
- ▌ Fla Triton To Gluon · fla-orgWorkflow 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.
- ▌ Fla Dispatch Backends · fla-orgWorkflow 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 Optimization Loop · fla-org bundleDisciplined, 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.
- ▌ Fla Ascend Performance · fla-org bundleGuidelines 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.
- ▌ Fla Nvidia Performance · fla-orgGuidelines 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.
- ▌ Fla Correctness Coverage · fla-org bundleGuidelines 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.