From 42d70599a8724fca63685ea11aa1040fd6dc18d6 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Mon, 10 Aug 2026 21:53:00 +0000 Subject: [PATCH 01/10] FlashKDA: keep this branch's project state across the trunk merge --- projects/FlashKDA/notes.md | 1 + projects/FlashKDA/status.json | 17 +++++++++++++++++ 2 files changed, 18 insertions(+) create mode 100644 projects/FlashKDA/notes.md create mode 100644 projects/FlashKDA/status.json diff --git a/projects/FlashKDA/notes.md b/projects/FlashKDA/notes.md new file mode 100644 index 00000000..01adf092 --- /dev/null +++ b/projects/FlashKDA/notes.md @@ -0,0 +1 @@ +# FlashKDA notes diff --git a/projects/FlashKDA/status.json b/projects/FlashKDA/status.json new file mode 100644 index 00000000..9ab6e59e --- /dev/null +++ b/projects/FlashKDA/status.json @@ -0,0 +1,17 @@ +{ + "schema_version": 3, + "name": "FlashKDA", + "upstream_url": "https://github.com/MoonshotAI/FlashKDA", + "fork_url": null, + "fork_default_branch": "main", + "priority": 0.0, + "ext_type": "torch-extension", + "adopted_at": "2026-08-07T17:05:59Z", + "updated_at": "2026-08-07T17:05:59Z", + "head_sha": null, + "depends_on": [], + "porting": null, + "waivers": {}, + "stage": "unclaimed", + "platforms": {} +} From 34f6a0ec65c83375368a774f90abe1b3d11bc8ea Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Thu, 13 Aug 2026 03:18:38 +0000 Subject: [PATCH 02/10] intake: FlashKDA screened -- MIT tier 1, recommend decline (cant-port) --- projects/FlashKDA/notes.md | 123 ++++++++++++++++++++++++++++++++++ projects/FlashKDA/stats.jsonl | 1 + projects/FlashKDA/status.json | 13 +++- 3 files changed, 135 insertions(+), 2 deletions(-) create mode 100644 projects/FlashKDA/stats.jsonl diff --git a/projects/FlashKDA/notes.md b/projects/FlashKDA/notes.md index 01adf092..0c8702aa 100644 --- a/projects/FlashKDA/notes.md +++ b/projects/FlashKDA/notes.md @@ -1 +1,124 @@ # FlashKDA notes + +## Intake screen (2026-08-13, linux-gfx1100) + +Upstream: https://github.com/MoonshotAI/FlashKDA — "FlashKDA: Flash Kimi Delta +Attention", high-performance KDA attention kernels built on CUTLASS/CuTe. Screened +from a shallow clone at `agent_space/FlashKDA-screen`; no fork exists and none was +requested. + +**Recommendation: decline, reason `cant-port`.** The argument is below. It is a +recommendation, not a decision — a person answers it through the intake queue. + +### Licence + +`license_spdx = MIT` (tier 1, cleared to contribute). Confirmed by reading `LICENSE` +in the tree, not just GitHub's field: verbatim MIT, "Copyright (c) 2026 MoonshotAI". +GitHub's API agrees (`MIT`). + +`scan-nvidia` over the main tree is clean. + +**One finding a person must rule on if this is ever adopted.** The repo carries a +submodule, `cutlass` -> https://github.com/NVIDIA/cutlass.git, pinned at +`5c149f52a436782210263fb2f19b354443a61c6a`. That tree is dual-licensed per-part: + +- `LICENSE.txt` is BSD-3-Clause and covers the C++ headers, and it says explicitly + that "the files located in the `python/CuTeDSL` directory are licensed under the + NVIDIA End User License Agreement (EULA)". +- `EULA.txt` is that agreement. `scan-nvidia` over the submodule flags ~40 files, all + of them under `python/CuTeDSL/`, and none anywhere else. + +The EULA restricts the Software to "systems with NVIDIA GPUs" (1.1) and clause 2.11 +forbids reverse-engineering output artifacts "for the purpose of translating such +output artifacts to target a non-NVIDIA platform" — which names the activity a +CUDA-to-ROCm port performs. It applies only to the CuTeDSL Python package. + +FlashKDA does not use that package. Its build pulls in only the BSD-3-Clause C++ +headers: `setup.py` adds `cutlass/include`, `cutlass/examples/common`, and +`cutlass/tools/util/include` to the include path, and every source include is a +`cute/...` or `cutlass/...` C++ header. Nothing imports CuTeDSL, and the port would +neither modify nor redistribute it. + +So the finding is almost certainly benign — but per the intake role an NVIDIA +proprietary licence on any file needs a person's decision before proceeding, and I do +not clear it myself. If the decline is upheld the question is moot; if someone +overrides to fork, this must be answered first and not skipped. + +### Duplicate effort + +No AMD or ROCm effort on FlashKDA itself: + +- No `FlashKDA` repository in AMD-Ecosystem (checked the org listing; the near-name + matches are `flashinfer`, `flashinfer-bench`, `ffpa-attn` — different projects). +- Nothing matching in the ROCm org. +- GitHub repo search for `flashkda` returns upstream plus `vllm-project/FlashKDA` + (12 stars, branches `master`/`dev` plus two feature branches, no ROCm branch), + two empty personal forks, and `popfido/FlashKDA-mlx` — an Apple MLX port, so there + is porting interest, but none toward AMD. +- `grep -rniE 'amd|rocm|hip|gfx[0-9]' README* docs/ BENCHMARK*.md` yields one hit, + and it is the substring in "on-chip". No notable-forks section, no platform ports. +- No MOAT disposition, no opt-out, no other `port/` branch for it or for a related + project. + +**The capability, though, already reaches AMD by another route.** FlashKDA is not a +standalone library — it is an optional backend for `flash-linear-attention`'s +`chunk_kda`, auto-dispatched when installed and disabled with `FLA_FLASH_KDA=0`, at +which point FLA falls back to its own Triton kernels. Those Triton kernels +(`fla/ops/kda/`) run on ROCm, and fla-org tests them on AMD in CI — +`.github/workflows/amd-mi300.yml`. So an AMD user wanting Kimi Delta Attention has a +supported, upstream-CI-tested path today. This is not `already-supported` for +FlashKDA, which has no AMD port at all, but it changes what a FlashKDA port would +buy: speed on hardware-specific code, not access to a missing capability. + +### Viability + +Genuine CUDA, and unusually concentrated: a torch extension of ~2,300 lines across +six files, all of the GPU work in `csrc/smxx/`. + +The problem is what it is built on. + +- **Every kernel is CuTe/CUTLASS.** `csrc/smxx/utils.cuh` alone pulls 19 CUTLASS + headers, including `cute/arch/copy_sm90_tma.hpp`, `cute/arch/cluster_sm90.hpp`, + `cutlass/cluster_launch.hpp`, `cutlass/arch/barrier.h`, and + `cutlass/pipeline/sm90_pipeline.hpp`. Layouts, tensors, and the GEMMs are all CuTe + types. NVIDIA CUTLASS has no ROCm/HIP backend; AMD's analogue is Composable Kernel, + a different library with a different API. Porting therefore is not translating a + CUDA path — it is reimplementing the kernels against a different tile library. +- **TMA is pervasive, not incidental.** 25 `SM90_TMA` references and 24 `make_tma` + call sites, plus `CUTE_GRID_CONSTANT` TMA descriptor parameters threaded through + every kernel entry point (`tma_load_q`, `tma_load_k`, `tma_store_ws_*`, ...), and + `ClusterTransactionBarrier` in shared memory. The most recent upstream commit is + literally "fix missing proxy fences around TMA accesses". Neither CDNA3 nor RDNA3 + has a TMA equivalent; the data movement would have to be rewritten, and with it the + pipeline and barrier structure that is the design. +- **Upstream targets Hopper and Blackwell only.** `SUPPORTED_CUDA_ARCHS = + ["90a", "100a", "103a", "120a"]`, README requirement "SM90 and above", benchmark + files for GB200 and H20. +- **This host's platform is the worst fit of the family.** gfx1100 is RDNA3: wave32, + no MFMA, no async-copy analogue of any of this. The `wave64` gate would need CDNA + and would still need the same full rewrite; there is no cheap gate to satisfy. +- Three inline PTX asm uses (`ex2.approx.ftz.f32`, `tanh.approx.f32`, + `cvt.f32.bf16`) — trivial next to the above, noted only for completeness. + +That combination puts the work far outside MOAT's standing rule to build the smallest +complete port preserving upstream structure and the CUDA path. A from-scratch +Composable-Kernel or raw-HIP reimplementation of a Hopper-tuned attention kernel is a +new kernel project, and it would arrive upstream as a large unsolicited parallel +backend in a repo whose reason for existing is Hopper/Blackwell tuning — a poor +prospect for acceptance even if it were built and validated. + +**Dependencies.** None on any MOAT project, so `depends_on` stays empty. External +build dependencies are PyTorch (>= 2.4), CUDA 12.9+, and the vendored CUTLASS +submodule. The `flash-linear-attention` relationship is integration, not a build +dependency. + +**Upstream health.** Healthy, and not a factor in the decline: not archived, 1208 +stars, 114 forks, last push 2026-07-30, active commits through late July 2026. If the +recommendation is overridden, upstream is at least alive enough to receive a PR. + +### If a person overrides to fork + +Two things must happen before any porting work, in this order: rule on the CuTeDSL +EULA finding above, and accept that the first real task is a kernel reimplementation +rather than a translation — so the planner should scope it as such, and on a CDNA +host, not on gfx1100. diff --git a/projects/FlashKDA/stats.jsonl b/projects/FlashKDA/stats.jsonl new file mode 100644 index 00000000..3b5e4a1c --- /dev/null +++ b/projects/FlashKDA/stats.jsonl @@ -0,0 +1 @@ +{"kind":"session","ts":"2026-08-13T03:12:45Z","epoch":1786590765.138518808,"event":"start","platform":"linux-gfx1100"} diff --git a/projects/FlashKDA/status.json b/projects/FlashKDA/status.json index 9ab6e59e..bce4eabc 100644 --- a/projects/FlashKDA/status.json +++ b/projects/FlashKDA/status.json @@ -7,11 +7,20 @@ "priority": 0.0, "ext_type": "torch-extension", "adopted_at": "2026-08-07T17:05:59Z", - "updated_at": "2026-08-07T17:05:59Z", + "updated_at": "2026-08-13T03:18:25Z", "head_sha": null, "depends_on": [], "porting": null, "waivers": {}, "stage": "unclaimed", - "platforms": {} + "platforms": {}, + "license_spdx": "MIT", + "intake": { + "verdict": "decline", + "reason": "cant-port", + "duplicate_effort": "none for FlashKDA itself (no AMD-Ecosystem or ROCm fork, no AMD mention in upstream docs); the KDA capability reaches AMD via fla-org/flash-linear-attention's Triton fla/ops/kda, CI-tested on MI300", + "viable": false, + "summary": "MIT (tier 1) Hopper/Blackwell-only KDA attention kernels written entirely in CUTLASS/CuTe -- 25 SM90_TMA sites, sm90 cluster pipelines, archs 90a/100a/103a/120a; CuTe has no ROCm backend and TMA no AMD equivalent, so this is a from-scratch CK/HIP kernel rewrite rather than a port, and AMD users already get KDA through flash-linear-attention's Triton path (CI-tested on MI300)", + "at": "2026-08-13T03:18:25Z" + } } From 9540a68fd055c9453392f6e2366d6874ee454f4e Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Thu, 13 Aug 2026 03:19:17 +0000 Subject: [PATCH 03/10] intake: FlashKDA session telemetry --- projects/FlashKDA/stats.jsonl | 1 + 1 file changed, 1 insertion(+) diff --git a/projects/FlashKDA/stats.jsonl b/projects/FlashKDA/stats.jsonl index 3b5e4a1c..368ccac0 100644 --- a/projects/FlashKDA/stats.jsonl +++ b/projects/FlashKDA/stats.jsonl @@ -1 +1,2 @@ {"kind":"session","ts":"2026-08-13T03:12:45Z","epoch":1786590765.138518808,"event":"start","platform":"linux-gfx1100"} +{"kind":"session","ts":"2026-08-13T03:19:12Z","epoch":1786591152.066405121,"event":"end","platform":"linux-gfx1100"} From 0a6931c425d326e979a0f12b8507672bcade7f73 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Fri, 14 Aug 2026 07:01:45 +0000 Subject: [PATCH 04/10] FlashKDA: intake re-screen on linux-gfx90a confirms cant-port decline --- projects/FlashKDA/notes.md | 73 +++++++++++++++++++++++++++++++++++ projects/FlashKDA/stats.jsonl | 1 + projects/FlashKDA/status.json | 6 +-- 3 files changed, 77 insertions(+), 3 deletions(-) diff --git a/projects/FlashKDA/notes.md b/projects/FlashKDA/notes.md index 0c8702aa..99c056dd 100644 --- a/projects/FlashKDA/notes.md +++ b/projects/FlashKDA/notes.md @@ -122,3 +122,76 @@ Two things must happen before any porting work, in this order: rule on the CuTeD EULA finding above, and accept that the first real task is a kernel reimplementation rather than a translation — so the planner should scope it as such, and on a CDNA host, not on gfx1100. + +## Intake re-screen (2026-08-14, linux-gfx90a) + +Independent second-host confirmation of the 2026-08-13 screen. **The decline +recommendation stands, unchanged: `cant-port`.** Verified against a fresh shallow +clone at upstream head `1ce47ea` ("optimize kda prepare cu_seqlens scan with +prefix-sum and binary search (#13)"), one commit ahead of the first screen. Nothing +material changed. + +Re-verified, each independently rather than taken on trust: + +- Licence MIT, tier 1 (`licenses.py check` → `license=MIT tier=1`). `scan-nvidia` + over the main tree: clean. +- The `cutlass` submodule is still pinned at `5c149f52`, and that pin's `LICENSE.txt` + still carries the per-part clause ("The files located in the `python/CuTeDSL` + directory are licensed under the NVIDIA End User License Agreement"), with + `EULA.txt` alongside it. FlashKDA still uses none of it: no `cutedsl`, + `cutlass.cute`, or `import cutlass` anywhere in `csrc/`, `flash_kda/`, `setup.py`, + `tests/`, or `benchmarks/`, and `setup.py` adds only the BSD-3-Clause C++ include + paths. **Still a person's ruling, still not cleared here.** +- No AMD-Ecosystem or ROCm FlashKDA repo (both 404). Repo search adds nothing new + toward AMD. Note for whoever re-runs this: a fork-list regex for `amd` matches + `Shamdon/FlashKDA` on the substring in the owner's name — a false positive, not an + AMD fork. +- `grep -rniE 'amd|rocm|hip|gfx[0-9]' README* docs/ BENCHMARK*.md` still yields + exactly one hit, still the substring in "on-chip". +- The AMD route to the capability is real and still live: `fla-org/flash-linear-attention` + has `fla/ops/kda/` (Triton, with a `backends/` subpackage) and + `.github/workflows/amd-mi300.yml`. +- `SUPPORTED_CUDA_ARCHS = ["90a", "100a", "103a", "120a"]`; README requirements + "SM90 and above", CUDA 12.9+, PyTorch 2.4+. Source is 2,346 lines across six files. +- CuTe/TMA density confirmed: 24 `SM90_TMA` references, 24 `make_tma` sites, 8 + cluster/sm90-pipeline references. The first screen said 25 TMA sites; the correct + count is 24, so the queue summary is corrected. The difference changes nothing. + +**One refinement this host adds, and it is the reason the re-screen was worth +running.** The first screen was performed on gfx1100 and noted that RDNA3 is "the +worst fit of the family" — wave32, no MFMA. That invites the objection that the +decline is an artifact of screening on unsuitable hardware. It is not. This host is +gfx90a: CDNA2, wave64, MFMA present — the `wave64` gate architecture, and a far +better-suited target. The decline is unchanged from here, because it never rested on +the host's wavefront or matrix-core support. It rests on two things neither CDNA2 nor +CDNA3 changes: + +1. NVIDIA CUTLASS/CuTe has no ROCm backend, and every layout, tensor, and GEMM in + this codebase is a CuTe type. AMD's Composable Kernel is a different library with + a different API, not a drop-in. +2. TMA has no AMD equivalent on any current architecture, and TMA is the design here, + not an optimization layered on top — descriptors are threaded through every kernel + entry point as `CUTE_GRID_CONSTANT` parameters and the pipeline and barrier + structure is built around them. + +So a CDNA host would face the same from-scratch kernel reimplementation that gfx1100 +would. There is no platform in the fleet on which this becomes a translation rather +than a rewrite, which is precisely what `cant-port` means. + +**Dependencies re-checked:** `depends_on` stays empty. No MOAT project provides +CUTLASS or flash-linear-attention (checked against `moatlib.py projects` and +`DEPENDENCIES.md`), so there is no unknown hard dependency needing an intake request. +PyTorch, CUDA 12.9+, and the vendored CUTLASS submodule are external build +dependencies only. + +**Upstream health:** not archived, not disabled, 1211 stars, 115 forks, last push +2026-07-30. A PR would have a live destination if the recommendation were overridden. + +**Unrelated data defect noticed while screening, for whoever maintains discovery.** +`data/candidates.json` carries two entries named `FlashKDA`: +`MoonshotAI/FlashKDA` (the one screened here) and `0xwilliamortiz/FlashKDA` +(195 stars, MIT, "memory-efficient KDA kernels for training and decode"). The second +is a **404 on the GitHub API** — deleted or never public. Beyond being stale, it is a +name collision: adopting it would scaffold to the same `projects/FlashKDA` and the +same `port/FlashKDA` claim as this project. Not fixed here, since candidate curation +is not the intake role's to edit and the decline makes it non-urgent. diff --git a/projects/FlashKDA/stats.jsonl b/projects/FlashKDA/stats.jsonl index 368ccac0..e985ea4e 100644 --- a/projects/FlashKDA/stats.jsonl +++ b/projects/FlashKDA/stats.jsonl @@ -1,2 +1,3 @@ {"kind":"session","ts":"2026-08-13T03:12:45Z","epoch":1786590765.138518808,"event":"start","platform":"linux-gfx1100"} {"kind":"session","ts":"2026-08-13T03:19:12Z","epoch":1786591152.066405121,"event":"end","platform":"linux-gfx1100"} +{"kind":"session","ts":"2026-08-14T06:59:31Z","epoch":1786690771.859204086,"event":"start","platform":"linux-gfx90a"} diff --git a/projects/FlashKDA/status.json b/projects/FlashKDA/status.json index bce4eabc..e91f5436 100644 --- a/projects/FlashKDA/status.json +++ b/projects/FlashKDA/status.json @@ -7,7 +7,7 @@ "priority": 0.0, "ext_type": "torch-extension", "adopted_at": "2026-08-07T17:05:59Z", - "updated_at": "2026-08-13T03:18:25Z", + "updated_at": "2026-08-14T07:01:39Z", "head_sha": null, "depends_on": [], "porting": null, @@ -20,7 +20,7 @@ "reason": "cant-port", "duplicate_effort": "none for FlashKDA itself (no AMD-Ecosystem or ROCm fork, no AMD mention in upstream docs); the KDA capability reaches AMD via fla-org/flash-linear-attention's Triton fla/ops/kda, CI-tested on MI300", "viable": false, - "summary": "MIT (tier 1) Hopper/Blackwell-only KDA attention kernels written entirely in CUTLASS/CuTe -- 25 SM90_TMA sites, sm90 cluster pipelines, archs 90a/100a/103a/120a; CuTe has no ROCm backend and TMA no AMD equivalent, so this is a from-scratch CK/HIP kernel rewrite rather than a port, and AMD users already get KDA through flash-linear-attention's Triton path (CI-tested on MI300)", - "at": "2026-08-13T03:18:25Z" + "summary": "MIT (tier 1) Hopper/Blackwell-only KDA attention kernels written entirely in CUTLASS/CuTe -- 24 SM90_TMA sites, sm90 cluster pipelines, archs 90a/100a/103a/120a; CuTe has no ROCm backend and TMA no AMD equivalent, so this is a from-scratch CK/HIP kernel rewrite rather than a port on CDNA as much as RDNA, and AMD users already get KDA through flash-linear-attention's Triton path (CI-tested on MI300)", + "at": "2026-08-14T07:01:39Z" } } From d4c98a5482600a82e20bc54c9af745e6f9f8ca44 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Fri, 14 Aug 2026 07:02:00 +0000 Subject: [PATCH 05/10] FlashKDA: record intake session telemetry --- projects/FlashKDA/stats.jsonl | 1 + 1 file changed, 1 insertion(+) diff --git a/projects/FlashKDA/stats.jsonl b/projects/FlashKDA/stats.jsonl index e985ea4e..c7515114 100644 --- a/projects/FlashKDA/stats.jsonl +++ b/projects/FlashKDA/stats.jsonl @@ -1,3 +1,4 @@ {"kind":"session","ts":"2026-08-13T03:12:45Z","epoch":1786590765.138518808,"event":"start","platform":"linux-gfx1100"} {"kind":"session","ts":"2026-08-13T03:19:12Z","epoch":1786591152.066405121,"event":"end","platform":"linux-gfx1100"} {"kind":"session","ts":"2026-08-14T06:59:31Z","epoch":1786690771.859204086,"event":"start","platform":"linux-gfx90a"} +{"kind":"session","ts":"2026-08-14T07:01:55Z","epoch":1786690915.683138792,"event":"end","platform":"linux-gfx90a"} From e8bb8da76f13aa90c577550fcedf9fa0c8b29bf6 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Fri, 14 Aug 2026 14:25:25 +0000 Subject: [PATCH 06/10] FlashKDA: third screen on gfx942 -- decline stands, MI300 CI claim corrected --- projects/FlashKDA/notes.md | 114 +++++++++++++++++++++++++++++++++- projects/FlashKDA/stats.jsonl | 2 + projects/FlashKDA/status.json | 8 +-- 3 files changed, 117 insertions(+), 7 deletions(-) diff --git a/projects/FlashKDA/notes.md b/projects/FlashKDA/notes.md index 99c056dd..b24ba151 100644 --- a/projects/FlashKDA/notes.md +++ b/projects/FlashKDA/notes.md @@ -66,7 +66,11 @@ standalone library — it is an optional backend for `flash-linear-attention`'s which point FLA falls back to its own Triton kernels. Those Triton kernels (`fla/ops/kda/`) run on ROCm, and fla-org tests them on AMD in CI — `.github/workflows/amd-mi300.yml`. So an AMD user wanting Kimi Delta Attention has a -supported, upstream-CI-tested path today. This is not `already-supported` for +supported, upstream-CI-tested path today. +[**CORRECTED 2026-08-14 by the gfx942 screen: the CI half of this is wrong.** That +workflow is disabled (`if: false`) and has never run — 0 runs. The Triton path is +installable on ROCm but upstream-unvalidated on AMD hardware. See the gfx942 section.] +This is not `already-supported` for FlashKDA, which has no AMD port at all, but it changes what a FlashKDA port would buy: speed on hardware-specific code, not access to a missing capability. @@ -150,7 +154,9 @@ Re-verified, each independently rather than taken on trust: exactly one hit, still the substring in "on-chip". - The AMD route to the capability is real and still live: `fla-org/flash-linear-attention` has `fla/ops/kda/` (Triton, with a `backends/` subpackage) and - `.github/workflows/amd-mi300.yml`. + `.github/workflows/amd-mi300.yml`. [**CORRECTED by the gfx942 screen:** confirming + the workflow file exists is not confirming CI coverage. It is disabled and has + never run.] - `SUPPORTED_CUDA_ARCHS = ["90a", "100a", "103a", "120a"]`; README requirements "SM90 and above", CUDA 12.9+, PyTorch 2.4+. Source is 2,346 lines across six files. - CuTe/TMA density confirmed: 24 `SM90_TMA` references, 24 `make_tma` sites, 8 @@ -187,7 +193,7 @@ dependencies only. **Upstream health:** not archived, not disabled, 1211 stars, 115 forks, last push 2026-07-30. A PR would have a live destination if the recommendation were overridden. -**Unrelated data defect noticed while screening, for whoever maintains discovery.** +**Unrelated data defect noticed while screening (first noted 2026-08-14, gfx90a).** `data/candidates.json` carries two entries named `FlashKDA`: `MoonshotAI/FlashKDA` (the one screened here) and `0xwilliamortiz/FlashKDA` (195 stars, MIT, "memory-efficient KDA kernels for training and decode"). The second @@ -195,3 +201,105 @@ is a **404 on the GitHub API** — deleted or never public. Beyond being stale, name collision: adopting it would scaffold to the same `projects/FlashKDA` and the same `port/FlashKDA` claim as this project. Not fixed here, since candidate curation is not the intake role's to edit and the decline makes it non-urgent. + +## Intake re-screen (2026-08-14, linux-gfx942) + +Third-host screen, on CDNA3/MI300 — the architecture the previous two screens +pointed at as the best-case target and the one the claimed AMD fallback route names. +**The decline recommendation stands: `cant-port`.** But this screen corrects a +supporting fact that both earlier screens recorded as true, and that correction is +the reason this run was worth doing. + +Verified against a fresh shallow clone at upstream head `1ce47ea` — the same commit +the gfx90a screen saw, so the code has not moved. + +### Correction: the "CI-tested on MI300" claim is false + +Both prior screens, and the recorded `intake.duplicate_effort` and `summary` fields, +stated that `fla-org/flash-linear-attention` tests its Triton KDA kernels on AMD in +CI, citing `.github/workflows/amd-mi300.yml`. The file exists. It has never run. + +- The job is guarded `if: false`, and the file's own header comment says: "Disabled + by default: fla does not yet operate an amd-mi300 runner. Flip the `if:` guard + below (or replace with a real condition) once a runner is wired in. Workflow lives + here so the install / sanity-check pattern is documented." +- Its trigger is `workflow_dispatch` only — no push or PR trigger. +- Run count via the Actions API: **0**. For contrast, `nvidia-h100.yml` has 3028. + +The existence of a workflow file is not evidence of CI coverage. Checking the run +count is the cheap disambiguation, and it should be the habit whenever an +existing-support claim rests on a CI file. + +**What is actually true about the AMD route**, stated at the strength the evidence +supports: fla ships KDA as Triton kernels (`fla/ops/kda/`, the default path) and +offers a documented ROCm install extra (`pip install -e ".[rocm]"`, which +deliberately does not pin a Triton flavor so the ROCm wheel index supplies +`pytorch-triton-rocm`). So the AMD path is intended and plausible — Triton targets +ROCm — but it is **unvalidated by upstream on AMD hardware**. Nobody should record +that AMD users have a tested KDA path today. + +This weakens the "capability already reaches AMD" argument rather than the decline. +It also points at the tractable work, which is worth saying plainly in the queue: if +someone wants Kimi Delta Attention working on MI300, the cheap, high-value task is +validating fla's existing Triton `fla/ops/kda` on gfx942 — a different project, and +one whose upstream has already built the scaffolding and is visibly waiting for a +runner. Rewriting FlashKDA's CuTe kernels is the expensive way to the same capability. + +### Re-verified independently on this host + +- Licence MIT, tier 1 (`licenses.py check` → `license=MIT tier=1`); `LICENSE` read + directly, verbatim MIT, "Copyright (c) 2026 MoonshotAI". `scan-nvidia` over the + main tree: clean. +- `cutlass` submodule still pinned at `5c149f52a436782210263fb2f19b354443a61c6a` + (`git ls-tree HEAD` gitlink), url `https://github.com/NVIDIA/cutlass.git`. The + per-part CuTeDSL EULA finding from the first screen is unchanged and **still not + cleared here — it remains a person's ruling if this is ever adopted.** FlashKDA + uses none of it: a case-insensitive search for `cutedsl`, `cutlass.cute`, + `import cutlass`, `from cutlass` across the whole tree returns nothing, and + `setup.py` adds only the BSD-3-Clause C++ include paths (`cutlass/include`, + `cutlass/examples/common`, `cutlass/tools/util/include`). +- No `FlashKDA` repo in AMD-Ecosystem or ROCm (both 404), nor `AMD-Ecosystem/flash-kda` + or `ROCm/flash-linear-attention`. Repo search returns upstream, `vllm-project/FlashKDA` + (12 stars), two empty personal repos, and `popfido/FlashKDA-mlx`. Nothing toward AMD. +- No MOAT disposition for FlashKDA in `data/dispositions.json` (282 entries), no + opt-out (`optout.py list` → nobody), no other `port/` branch. +- `SUPPORTED_CUDA_ARCHS = ["90a", "100a", "103a", "120a"]`; 2346 lines across six + source files; three inline PTX uses. +- CuTe/TMA density: 24 `SM90_TMA` uses, 24 `make_tma` sites, 22 `CUTE_GRID_CONSTANT` + parameters, 5 `ClusterTransactionBarrier`, plus `cutlass/pipeline/sm90_pipeline.hpp` + and `cutlass/cluster_launch.hpp`. + +**The 24-vs-25 TMA discrepancy between the first two screens is settled**, since it +cost a correction once already: case-sensitive `SM90_TMA` in `csrc/` is 24; +case-insensitive is 25, the extra match being the include filename +`cute/arch/copy_sm90_tma.hpp` on line 14 of `utils.cuh`. Both counts were right about +different questions. 24 is the number of code uses. + +### Why CDNA3 does not rescue it + +gfx942 is the strongest case AMD can make here — CDNA3, wave64, MFMA, and the +`wave64` gate architecture. The decline is unchanged from this host for the reason +the gfx90a screen already gave, which this host confirms rather than repeats: the +blocker is not wavefront width or matrix-core availability. It is that every layout, +tensor, and GEMM is a CuTe type and NVIDIA CUTLASS has no ROCm backend, and that TMA +is the design rather than an optimization — descriptors threaded through every kernel +entry point as `CUTE_GRID_CONSTANT` parameters, with the pipeline and barrier +structure built around them. CDNA3 has no TMA equivalent. Three hosts spanning RDNA3, +CDNA2 and CDNA3 now agree, which is as much as screening can establish: there is no +platform in the fleet on which this becomes a translation rather than a from-scratch +kernel reimplementation. That is what `cant-port` means. + +**Dependencies:** `depends_on` stays empty; no MOAT project provides CUTLASS or +flash-linear-attention. PyTorch, CUDA 12.9+, and the vendored CUTLASS submodule are +external build dependencies only. + +**Upstream health:** not archived, not disabled, 1211 stars, 115 forks, last push +2026-07-30, default branch `master`. A PR would have a live destination if the +recommendation were overridden. + +**Minor data observation, not load-bearing.** `data/retired_stats.jsonl` line 83 +carries a single token record for project `FlashKDA` dated 2026-06-04 with source +`porter` — two months before this project's `adopted_at` (2026-08-07) and before +upstream's own recorded activity. There is no matching disposition and no other +trace of a prior lifecycle. Flagged for whoever maintains telemetry; it does not +affect the recommendation and I did not edit it. diff --git a/projects/FlashKDA/stats.jsonl b/projects/FlashKDA/stats.jsonl index c7515114..2fc88399 100644 --- a/projects/FlashKDA/stats.jsonl +++ b/projects/FlashKDA/stats.jsonl @@ -2,3 +2,5 @@ {"kind":"session","ts":"2026-08-13T03:19:12Z","epoch":1786591152.066405121,"event":"end","platform":"linux-gfx1100"} {"kind":"session","ts":"2026-08-14T06:59:31Z","epoch":1786690771.859204086,"event":"start","platform":"linux-gfx90a"} {"kind":"session","ts":"2026-08-14T07:01:55Z","epoch":1786690915.683138792,"event":"end","platform":"linux-gfx90a"} +{"kind":"session","ts":"2026-08-14T14:22:06Z","epoch":1786717326.587487960,"event":"start","platform":"linux-gfx942"} +{"kind":"session","ts":"2026-08-14T14:25:21Z","epoch":1786717521.427508328,"event":"end","platform":"linux-gfx942"} diff --git a/projects/FlashKDA/status.json b/projects/FlashKDA/status.json index e91f5436..efe82f2c 100644 --- a/projects/FlashKDA/status.json +++ b/projects/FlashKDA/status.json @@ -7,7 +7,7 @@ "priority": 0.0, "ext_type": "torch-extension", "adopted_at": "2026-08-07T17:05:59Z", - "updated_at": "2026-08-14T07:01:39Z", + "updated_at": "2026-08-14T14:25:14Z", "head_sha": null, "depends_on": [], "porting": null, @@ -18,9 +18,9 @@ "intake": { "verdict": "decline", "reason": "cant-port", - "duplicate_effort": "none for FlashKDA itself (no AMD-Ecosystem or ROCm fork, no AMD mention in upstream docs); the KDA capability reaches AMD via fla-org/flash-linear-attention's Triton fla/ops/kda, CI-tested on MI300", + "duplicate_effort": "none: no AMD-Ecosystem or ROCm fork, no AMD mention in upstream docs. KDA reaches AMD only via flash-linear-attention's Triton fla/ops/kda, which installs on ROCm but is upstream-unvalidated on AMD hardware -- its amd-mi300 workflow is disabled (if: false) with 0 runs ever vs 3028 for nvidia-h100, correcting the earlier screens' 'CI-tested on MI300'", "viable": false, - "summary": "MIT (tier 1) Hopper/Blackwell-only KDA attention kernels written entirely in CUTLASS/CuTe -- 24 SM90_TMA sites, sm90 cluster pipelines, archs 90a/100a/103a/120a; CuTe has no ROCm backend and TMA no AMD equivalent, so this is a from-scratch CK/HIP kernel rewrite rather than a port on CDNA as much as RDNA, and AMD users already get KDA through flash-linear-attention's Triton path (CI-tested on MI300)", - "at": "2026-08-14T07:01:39Z" + "summary": "MIT (tier 1) Hopper/Blackwell-only KDA attention kernels written entirely in CUTLASS/CuTe -- 24 SM90_TMA uses, 22 CUTE_GRID_CONSTANT TMA descriptors, archs 90a/100a/103a/120a; CuTe has no ROCm backend and TMA no AMD equivalent, so this is a from-scratch CK/HIP kernel rewrite on CDNA3 exactly as on RDNA3 (three hosts agree), and the cheap route to KDA on MI300 is validating flash-linear-attention's existing Triton fla/ops/kda rather than porting this", + "at": "2026-08-14T14:25:14Z" } } From bd1306450ea723931c623da09fb036778f9b62bc Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Wed, 19 Aug 2026 05:04:05 +0000 Subject: [PATCH 07/10] FlashKDA: delta-check confirms the decline; queue answer is the blocker --- projects/FlashKDA/notes.md | 68 ++++++++++++++++++++++++++++++++++- projects/FlashKDA/stats.jsonl | 2 ++ projects/FlashKDA/status.json | 8 ++--- 3 files changed, 73 insertions(+), 5 deletions(-) diff --git a/projects/FlashKDA/notes.md b/projects/FlashKDA/notes.md index b24ba151..8e0ede9c 100644 --- a/projects/FlashKDA/notes.md +++ b/projects/FlashKDA/notes.md @@ -297,9 +297,75 @@ external build dependencies only. 2026-07-30, default branch `master`. A PR would have a live destination if the recommendation were overridden. -**Minor data observation, not load-bearing.** `data/retired_stats.jsonl` line 83 +**Minor data observation, not load-bearing (gfx942 screen).** `data/retired_stats.jsonl` line 83 carries a single token record for project `FlashKDA` dated 2026-06-04 with source `porter` — two months before this project's `adopted_at` (2026-08-07) and before upstream's own recorded activity. There is no matching disposition and no other trace of a prior lifecycle. Flagged for whoever maintains telemetry; it does not affect the recommendation and I did not edit it. + +## Intake delta-check (2026-08-19, linux-gfx1100) — fourth dispatch, NOT a fourth screen + +This host was dispatched FlashKDA at `stage: unclaimed` for the fourth time. Three +screens (2026-08-13 gfx1100, 2026-08-14 gfx90a, 2026-08-14 gfx942) already agree on +`decline` / `cant-port`, and that recommendation is already recorded in +`status.json.intake`. So this run deliberately did **not** re-derive the screen. It +checked only what could have changed since 2026-08-14, and stopped. + +**Nothing changed. The recommendation stands, unchanged: `decline`, `cant-port`.** + +Delta checks, all read-only, no clone: + +- **Upstream head is still `1ce47ea`** ("optimize kda prepare cu_seqlens scan with + prefix-sum and binary search (#13)", authored 2026-07-29), the same commit the + gfx90a and gfx942 screens verified. `master` is the only branch. The tree is + therefore byte-identical to what was screened three times, so every code-level + finding — 24 `SM90_TMA` uses, 24 `make_tma` sites, 22 `CUTE_GRID_CONSTANT` + descriptors, `SUPPORTED_CUDA_ARCHS = ["90a","100a","103a","120a"]`, 2346 lines + across six files, the `cutlass` submodule pinned at `5c149f52` — holds by + construction. Re-cloning to recount them would produce the same numbers from the + same bytes, which is why it was skipped rather than repeated. +- **Licence unchanged: MIT, tier 1.** GitHub API `license.spdx_id = MIT`; the + `LICENSE` file was read directly by all three prior screens at this same SHA. + `license_spdx` was already recorded and stays `MIT`. +- **The CuTeDSL EULA finding is unchanged and still uncleared.** Same submodule pin, + so same per-part licensing. It remains a person's ruling if this is ever adopted, + and this host does not clear it either. +- **No AMD effort appeared.** `AMD-Ecosystem/FlashKDA`, `ROCm/FlashKDA`, and + `AMD-Ecosystem/flash-kda` are all still 404. A fresh repo search surfaces one name + the earlier screens did not see, `Unitflexmed1821/FlashKDA` (0 stars, not a fork, + pushed 2026-08-19, head "Update README.md", description advertising CUTLASS + kernels). Its README has zero matches for `amd|rocm|hip|gfx[0-9]` — an unrelated + re-upload, not an AMD port. Also new: `atomicmilkshake/godzilla-llama.cpp`, an + MSVC+CUDA Windows llama.cpp fork carrying KDA, likewise nothing toward AMD. +- **Upstream still healthy**, still a live PR destination if the recommendation is + overridden: not archived, not disabled, 1218 stars (was 1211), 117 forks (was 115), + last push 2026-07-30, default branch `master`. +- `depends_on` stays empty. + +### Why this project keeps coming back, which is the finding worth acting on + +The screen is not stuck — the *decision* is. FlashKDA is row 1 of the open intake +queue, issue AMD-Ecosystem/moat#8, opened 2026-08-07 and carrying **zero comments +after 12 days**. Its recommendation has been correct and complete since 2026-08-13. + +An agent may not record a decline, so `stage` correctly stays `unclaimed` — and the +selector offers `unclaimed` projects to whatever host asks next. A decline +recommendation awaiting a person is therefore indistinguishable, to the selector, +from a project nobody has looked at. That is the whole mechanism, and it is the same +one that had spconv screened six times on four platforms; spconv is row 2 of the same +unanswered issue. Four sessions have now been spent on FlashKDA to reach the answer +the first one reached. + +Two things would stop it, and both belong to a person, not to this role: + +1. **Answer issue #8.** One reply decides this batch, and `intake_queue.py apply` + records it. This is the real fix; everything else is a workaround. +2. **`moatlib.py set-hold FlashKDA`** if the answer will be a while. AGENTS.md + reserves `set-hold` to a person, so this host did not set it — but a held project + is skipped by the selector on every platform with no state touched, which is + exactly the behaviour wanted for a screened project queued behind a pending + decision. + +Neither was done here. The only write this run made was refreshing the queue summary +so the row itself reports that four hosts agree and the screen is not the bottleneck. diff --git a/projects/FlashKDA/stats.jsonl b/projects/FlashKDA/stats.jsonl index 2fc88399..8030a517 100644 --- a/projects/FlashKDA/stats.jsonl +++ b/projects/FlashKDA/stats.jsonl @@ -4,3 +4,5 @@ {"kind":"session","ts":"2026-08-14T07:01:55Z","epoch":1786690915.683138792,"event":"end","platform":"linux-gfx90a"} {"kind":"session","ts":"2026-08-14T14:22:06Z","epoch":1786717326.587487960,"event":"start","platform":"linux-gfx942"} {"kind":"session","ts":"2026-08-14T14:25:21Z","epoch":1786717521.427508328,"event":"end","platform":"linux-gfx942"} +{"kind":"session","ts":"2026-08-19T05:01:27Z","epoch":1787115687.749938491,"event":"start","platform":"linux-gfx1100"} +{"kind":"session","ts":"2026-08-19T05:04:05Z","epoch":1787115845.057749640,"event":"end","platform":"linux-gfx1100"} diff --git a/projects/FlashKDA/status.json b/projects/FlashKDA/status.json index efe82f2c..509a13ed 100644 --- a/projects/FlashKDA/status.json +++ b/projects/FlashKDA/status.json @@ -7,7 +7,7 @@ "priority": 0.0, "ext_type": "torch-extension", "adopted_at": "2026-08-07T17:05:59Z", - "updated_at": "2026-08-14T14:25:14Z", + "updated_at": "2026-08-19T05:03:51Z", "head_sha": null, "depends_on": [], "porting": null, @@ -18,9 +18,9 @@ "intake": { "verdict": "decline", "reason": "cant-port", - "duplicate_effort": "none: no AMD-Ecosystem or ROCm fork, no AMD mention in upstream docs. KDA reaches AMD only via flash-linear-attention's Triton fla/ops/kda, which installs on ROCm but is upstream-unvalidated on AMD hardware -- its amd-mi300 workflow is disabled (if: false) with 0 runs ever vs 3028 for nvidia-h100, correcting the earlier screens' 'CI-tested on MI300'", + "duplicate_effort": "none: no AMD-Ecosystem or ROCm fork (all 404 again on 2026-08-19), no AMD mention in upstream docs, and no AMD-directed fork among the copies that have appeared since. KDA reaches AMD only via flash-linear-attention's Triton fla/ops/kda, which installs on ROCm but is upstream-unvalidated on AMD hardware -- its amd-mi300 workflow is disabled (if: false) with 0 runs ever vs 3028 for nvidia-h100, correcting the earliest screens' 'CI-tested on MI300'", "viable": false, - "summary": "MIT (tier 1) Hopper/Blackwell-only KDA attention kernels written entirely in CUTLASS/CuTe -- 24 SM90_TMA uses, 22 CUTE_GRID_CONSTANT TMA descriptors, archs 90a/100a/103a/120a; CuTe has no ROCm backend and TMA no AMD equivalent, so this is a from-scratch CK/HIP kernel rewrite on CDNA3 exactly as on RDNA3 (three hosts agree), and the cheap route to KDA on MI300 is validating flash-linear-attention's existing Triton fla/ops/kda rather than porting this", - "at": "2026-08-14T14:25:14Z" + "summary": "MIT (tier 1) Hopper/Blackwell-only KDA attention kernels written entirely in CUTLASS/CuTe -- 24 SM90_TMA uses, 22 CUTE_GRID_CONSTANT TMA descriptors, archs 90a/100a/103a/120a; CuTe has no ROCm backend and TMA no AMD equivalent, so this is a from-scratch CK/HIP kernel rewrite on CDNA3 exactly as on RDNA3, and the cheap route to KDA on MI300 is validating flash-linear-attention's existing Triton fla/ops/kda rather than porting this -- screened four times on three platforms (RDNA3, CDNA2, CDNA3), all agreeing, the latest a delta-check confirming upstream has not moved off 1ce47ea since 2026-07-29, so the screen is settled and only this queue answer is outstanding", + "at": "2026-08-19T05:03:51Z" } } From 16a012e3cbc1edabbfa4066b90ba833bde40cc56 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Thu, 20 Aug 2026 06:33:56 +0000 Subject: [PATCH 08/10] FlashKDA: fifth screen finds ROCm/aiter already ships KDA -- reason moves to ported-elsewhere --- projects/FlashKDA/notes.md | 175 ++++++++++++++++++++++++++++++++++ projects/FlashKDA/stats.jsonl | 2 + projects/FlashKDA/status.json | 10 +- 3 files changed, 182 insertions(+), 5 deletions(-) diff --git a/projects/FlashKDA/notes.md b/projects/FlashKDA/notes.md index 8e0ede9c..0015eec0 100644 --- a/projects/FlashKDA/notes.md +++ b/projects/FlashKDA/notes.md @@ -369,3 +369,178 @@ Two things would stop it, and both belong to a person, not to this role: Neither was done here. The only write this run made was refreshing the queue summary so the row itself reports that four hosts agree and the screen is not the bottleneck. + +## Intake re-screen (2026-08-20, linux-gfx90a) — fifth dispatch, and the reason changes + +Fifth dispatch at `stage: unclaimed`. The decline is unchanged, but **the recorded +reason moves from `cant-port` to `ported-elsewhere`**, because a check no earlier +screen performed found that AMD itself now ships this capability on ROCm. The +technical `cant-port` argument still holds independently — see below — so the decline +is robust under either reason; `ported-elsewhere` is recorded because it is the +stronger and less arguable of the two, and because it names the standing rule that +applies: we do not duplicate AMD's own work. + +### The new finding: ROCm/aiter ships Kimi Delta Attention + +`ROCm/aiter` — AMD's own kernel library, MIT, "Copyright (C) 2024-2026, Advanced Micro +Devices, Inc." — contains a Triton KDA implementation, including a path named after +this very project: + +- `aiter/ops/triton/kimi_delta_attn/` — public entry point `chunk_kimi_delta_attn`, + documented as "Kimi Delta Attention Operations (Forward Only) ... for the KDA + linear-attention mixer used by Kimi-Linear / Kimi-K3. The chunked prefill op + mirrors `fla.ops.kda.chunk_kda`." +- `aiter/ops/triton/_triton_kernels/chunk_delta_attn/flash_kda.py` — the fused + variant, exporting `flash_kda_fwd` and `flash_kda_supported`. +- `op_tests/triton_tests/chunk_delta_attn/test_flash_kda.py` — tests the fused + FlashKDA path against the default chunk_delta_attn pipeline as reference, on + relative error because the two differ at bf16 level by design. +- `op_tests/op_benchmarks/triton/bench_flash_kda.py` — a benchmark. +- Listed in `aiter/ops/triton/README.md`: "kimi_delta_attn/ # Kimi Delta Attention + (chunked delta attention)". + +It is live work, not a stub, and it runs on real AMD hardware: + + 01ddb01 2026-08-10 [chunk_delta_attn] Add Triton kernels for chunk delta attention forward pass (#4568) + 0f6766f 2026-08-12 [Triton] Fix LDS OOM issue on MI300 (#4671) + 3679a25 2026-08-13 [Triton] Optimize chunk_delta_attn performance. (#4683) + +An LDS-OOM fix specific to MI300 is only written by someone running the kernel on +MI300. That is the hardware evidence the `fla` route never had. + +**Why the earlier screens missed it, which is the transferable lesson.** All four +prior screens searched for the project NAME (`FlashKDA` in AMD-Ecosystem, in ROCm, in +repo search) and for forks of the upstream. AMD's support is neither: it is a +differently-named op inside a large general-purpose kernel library, added on +2026-08-10 — after the three real screens (08-13, 08-14, 08-14) had already run, and +the 08-19 delta-check looked only at the project name again. Searching the CAPABILITY +inside `ROCm/aiter` (`repo:ROCm/aiter kda in:path`, `repo:ROCm/aiter "kimi delta"`) +is two API calls and is what found it. For any attention/GEMM/quantization kernel +project, `ROCm/aiter` should be searched by capability, not just the org by name. + +### Scope parity: it is the same op, not a loose equivalent + +Upstream FlashKDA's entire public API is one function. `flash_kda/__init__.py` exports +`fwd(q, k, v, g, beta, scale, out, A_log, dt_bias, lower_bound, initial_state, +final_state, cu_seqlens)` — forward only, bf16, K = V = 128 — backed by four CUDA +files in `csrc/smxx/` (`fwd_kernel1.cuh`, `fwd_kernel2.cuh`, `fwd_launch.cu`, +`utils.cuh`). The README's stated purpose is to be auto-dispatched as +`flash-linear-attention`'s `chunk_kda` backend (fla-org/flash-linear-attention#852). + +aiter's `chunk_kimi_delta_attn` is forward-only and explicitly "mirrors +`fla.ops.kda.chunk_kda`" — the same op, in the same role, on ROCm. Upstream itself is +forward-only too: training kernels are not in the tree, they are an unmerged +third-party PR (#28, "Add CUDA training kernels (fwd+bwd) for KDA"). So there is no +"but AMD only has part of it" gap to point at. + +Honest limits on the claim: aiter's implementation is Triton rather than CUTLASS-class +tuned, it is 10 days old, and forward/prefill only. It is authoritative (AMD-official, +actively maintained, MI300-fixed) rather than mature-by-age. Per +`assess-existing-support.md`, authoritativeness is the deciding axis, not age — an +AMD-official effort means reuse-and-improve or skip, never re-port from scratch. + +**Where the work goes if someone wants more KDA on AMD:** contribute to +`ROCm/aiter/aiter/ops/triton/kimi_delta_attn/` — backward/training, decode, or +performance. That is a different project with a different intake, and it is a far +cheaper path than reimplementing 2,346 lines of CuTe/TMA kernels against Composable +Kernel. + +### Delta-checks, all re-verified this run + +- **Licence: MIT, tier 1** (`licenses.py check MoonshotAI/FlashKDA` → + `license=MIT tier=1`, cleared to contribute). GitHub API `license.spdx_id = MIT`. + `license_spdx` stays `MIT`. Read directly from the `LICENSE` file by all three full + screens at this same SHA. +- **Upstream head is still `1ce47ea`** ("optimize kda prepare cu_seqlens scan with + prefix-sum and binary search (#13)", 2026-07-29), and `master` is still the only + branch. The tree is byte-identical to what was screened three times, so the + code-level counts (24 `SM90_TMA`, 24 `make_tma`, 22 `CUTE_GRID_CONSTANT`, + `SUPPORTED_CUDA_ARCHS = ["90a","100a","103a","120a"]`, 2346 lines / six files) hold + by construction and were not recounted. No shallow clone was taken for the same + reason; re-running `scan-nvidia` over identical bytes would reproduce the identical + result recorded on 08-13, 08-14 and 08-14. +- **The CuTeDSL EULA finding is unchanged and still uncleared.** The `cutlass` + submodule is still pinned at `5c149f52a436782210263fb2f19b354443a61c6a`, url + `https://github.com/NVIDIA/cutlass.git` (verified through the contents API at + `ref=1ce47ea`, no clone needed). Same pin means the same per-part licensing: the + ~40 `python/CuTeDSL/` files under the NVIDIA EULA, everything FlashKDA actually + builds against under BSD-3-Clause. Still a person's ruling if this is ever adopted; + this host does not clear it either. Moot if the decline is upheld. + +### Duplicate effort, including the open-PR/branch check no earlier screen ran + +This screen adds the upstream open-PR and branch check that the role now requires +(added after FLAMEGPU2 was adopted while the maintainer's own HIP draft PR sat open). + +- **Upstream open PRs: 10, none AMD-related.** `gh pr list --state open` returns #16, + #17, #18, #19, #27, #28, #29, #30, #31, #32 — numerics tolerance, store + parallelization, FLA validation, CUTLASS error reporting, a CPU training backend, a + CUDA training-kernel PR, tile-prefix reuse, CP transition, reusable workspace, and + empty varlen batches. `--search "hip OR rocm OR amd"` over open PRs returns **empty**, + over ALL states returns **empty**, and the same search over issues returns **empty**. + No maintainer HIP draft, no competing AMD PR. +- **Upstream branches: `master` only.** No rocm/hip/amd branch. +- `AMD-Ecosystem/FlashKDA`, `ROCm/FlashKDA`, `AMD-Ecosystem/flash-kda`, + `ROCm/flash-kda`, `AMD-Ecosystem/flash-linear-attention`, + `ROCm/flash-linear-attention`: all still 404. Org-scoped repo search for `kda` in + AMD-Ecosystem and in ROCm: 0 each — which is exactly why the name search keeps + missing aiter, whose repo name contains no "kda". +- **Forks: 118, none AMD-directed.** Scanning all fork full names for + `amd|rocm|hip|gfx[0-9]` yields two substring false positives and nothing else: + `Shamdon/FlashKDA` (known, "amd" inside the owner name) and, new this run, + `woshipapa/FlashKDA` ("hip" inside the owner name). Neither is an AMD port. +- `vllm-project/FlashKDA` (12 stars, pushed 2026-08-19) re-checked directly since vLLM + does care about ROCm: branches are `master` (at upstream's `1ce47ea`), `dev`, + `agent/export-intermediate-recurrent-checkpoints`, + `codex/eagle-dual-checkpoint-pr7`, `thien-codex/state-store-refactor`; open PRs #7 + and #8 are smem layout and recurrent-state export. Nothing toward AMD. +- Repo search for `flashkda` adds nothing new: upstream, vllm-project's copy, + `Unitflexmed1821/FlashKDA` and `allanleewh/FlashKDA` (re-uploads), + `popfido/FlashKDA-mlx` (Apple MLX), `atomicmilkshake/godzilla-llama.cpp` (Windows + MSVC+CUDA). +- No MOAT disposition for FlashKDA (282 entries, no `kda` key), no opt-out + (`optout.py list` → nobody), no other `port/` branch. + +**The `fla` amd-mi300 CI correction still stands, and has not moved.** Workflow run +counts re-read this run: `amd-mi300.yml` = **0 runs**, `nvidia-h100.yml` = **3051** +(was 3028 on 08-14). Both workflows report `state=active`, which is the trap — the job +body is guarded `if: false`, so "active" describes the file, not any execution. Read +the run count, never the state field. The `fla` Triton path remains +installable-but-unvalidated on AMD; aiter, not fla, is now the credible AMD route. + +### The `cant-port` argument, unchanged and still independently sufficient + +Recorded here so that changing the reason does not lose it. Every layout, tensor and +GEMM is a CuTe type and NVIDIA CUTLASS has no ROCm backend; TMA is the design rather +than an optimization, with descriptors threaded through every kernel entry point as +`CUTE_GRID_CONSTANT` parameters and the pipeline and barrier structure built around +them; no AMD architecture has a TMA equivalent. Three hosts spanning RDNA3 (gfx1100), +CDNA2 (gfx90a) and CDNA3 (gfx942) agreed there is no fleet platform on which this +becomes a translation rather than a from-scratch kernel reimplementation. If a person +prefers to record the decline on technical grounds instead, `cant-port` is equally +defensible — the answer is the same either way. + +### Viability and dependencies + +Genuine CUDA, confirmed again: four CUDA sources in `csrc/smxx/`, a `flash_kda_C` +compiled torch extension, `SUPPORTED_CUDA_ARCHS = ["90a","100a","103a","120a"]`, +CUDA 12.9+ / PyTorch 2.4+ requirements. `viable = no` for MOAT's purposes. + +`depends_on` stays empty. No MOAT project provides CUTLASS, aiter, or +flash-linear-attention, so there is no unknown hard dependency needing an intake +request. PyTorch, CUDA 12.9+ and the vendored CUTLASS submodule are external build +dependencies only. + +**Upstream health:** not archived, not disabled, 1221 stars (was 1218), 118 forks (was +117), last push 2026-07-30, default branch `master`. A PR would still have a live +destination if the recommendation were overridden. + +### Queue status + +Still row 1 of intake queue issue AMD-Ecosystem/moat#8, opened 2026-08-07, unanswered +after 13 days. Five sessions have now been spent reaching the answer the first one +reached — though this one earned its keep by finding the aiter implementation and by +running the open-PR check for the first time. The two things that would stop the +re-dispatch loop remain a person's: answer issue #8, or `moatlib.py set-hold FlashKDA` +while it waits. Neither was done here. No fork was requested, no disposition was +written, no PR was opened. diff --git a/projects/FlashKDA/stats.jsonl b/projects/FlashKDA/stats.jsonl index 8030a517..4836b202 100644 --- a/projects/FlashKDA/stats.jsonl +++ b/projects/FlashKDA/stats.jsonl @@ -6,3 +6,5 @@ {"kind":"session","ts":"2026-08-14T14:25:21Z","epoch":1786717521.427508328,"event":"end","platform":"linux-gfx942"} {"kind":"session","ts":"2026-08-19T05:01:27Z","epoch":1787115687.749938491,"event":"start","platform":"linux-gfx1100"} {"kind":"session","ts":"2026-08-19T05:04:05Z","epoch":1787115845.057749640,"event":"end","platform":"linux-gfx1100"} +{"kind":"session","ts":"2026-08-20T06:29:26Z","epoch":1787207366.578456516,"event":"start","platform":"linux-gfx90a"} +{"kind":"session","ts":"2026-08-20T06:33:56Z","epoch":1787207636.778818635,"event":"end","platform":"linux-gfx90a"} diff --git a/projects/FlashKDA/status.json b/projects/FlashKDA/status.json index 509a13ed..3ffbab42 100644 --- a/projects/FlashKDA/status.json +++ b/projects/FlashKDA/status.json @@ -7,7 +7,7 @@ "priority": 0.0, "ext_type": "torch-extension", "adopted_at": "2026-08-07T17:05:59Z", - "updated_at": "2026-08-19T05:03:51Z", + "updated_at": "2026-08-20T06:33:47Z", "head_sha": null, "depends_on": [], "porting": null, @@ -17,10 +17,10 @@ "license_spdx": "MIT", "intake": { "verdict": "decline", - "reason": "cant-port", - "duplicate_effort": "none: no AMD-Ecosystem or ROCm fork (all 404 again on 2026-08-19), no AMD mention in upstream docs, and no AMD-directed fork among the copies that have appeared since. KDA reaches AMD only via flash-linear-attention's Triton fla/ops/kda, which installs on ROCm but is upstream-unvalidated on AMD hardware -- its amd-mi300 workflow is disabled (if: false) with 0 runs ever vs 3028 for nvidia-h100, correcting the earliest screens' 'CI-tested on MI300'", + "reason": "ported-elsewhere", + "duplicate_effort": "ROCm/aiter ships Kimi Delta Attention on ROCm: aiter/ops/triton/kimi_delta_attn/ (chunk_kimi_delta_attn, forward-only, mirrors fla.ops.kda.chunk_kda) plus _triton_kernels/chunk_delta_attn/flash_kda.py (flash_kda_fwd), with tests and a benchmark; added 2026-08-10 (#4568), MI300 LDS OOM fix 2026-08-12 (#4671), perf work 2026-08-13 (#4683) -- AMD-official and running on MI300. No FlashKDA repo in AMD-Ecosystem or ROCm (404), no AMD-directed fork among 118, no rocm/hip/amd branch or PR upstream (10 open PRs, hip/rocm/amd search empty in all states), vllm-project/FlashKDA has 5 branches none AMD. fla's amd-mi300 CI is still 0 runs vs 3051 h100, so aiter not fla is the credible AMD route", "viable": false, - "summary": "MIT (tier 1) Hopper/Blackwell-only KDA attention kernels written entirely in CUTLASS/CuTe -- 24 SM90_TMA uses, 22 CUTE_GRID_CONSTANT TMA descriptors, archs 90a/100a/103a/120a; CuTe has no ROCm backend and TMA no AMD equivalent, so this is a from-scratch CK/HIP kernel rewrite on CDNA3 exactly as on RDNA3, and the cheap route to KDA on MI300 is validating flash-linear-attention's existing Triton fla/ops/kda rather than porting this -- screened four times on three platforms (RDNA3, CDNA2, CDNA3), all agreeing, the latest a delta-check confirming upstream has not moved off 1ce47ea since 2026-07-29, so the screen is settled and only this queue answer is outstanding", - "at": "2026-08-19T05:03:51Z" + "summary": "MIT (tier 1) Hopper/Blackwell-only KDA forward kernel, entirely CUTLASS/CuTe with TMA as the design -- decline as ported-elsewhere, NEW since the earlier screens: ROCm/aiter (AMD's own library) already ships the same op on ROCm as Triton kimi_delta_attn/chunk_kimi_delta_attn with a flash_kda fused path, tests, benchmark and an MI300 LDS fix, forward-only exactly like upstream's single fwd() export, so a port would duplicate AMD's own work; the older cant-port argument (no ROCm CuTe backend, no AMD TMA, from-scratch CK rewrite on any fleet platform) still holds independently, and if more KDA on AMD is wanted the cheap path is contributing backward/decode to aiter", + "at": "2026-08-20T06:33:47Z" } } From d8aa2e237b4a13cd56d1779a0cd03f2345ffe57d Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Thu, 20 Aug 2026 06:42:29 +0000 Subject: [PATCH 09/10] FlashKDA: fix ROCm casing in intake duplicate-effort field --- projects/FlashKDA/status.json | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/FlashKDA/status.json b/projects/FlashKDA/status.json index 3ffbab42..62058fbd 100644 --- a/projects/FlashKDA/status.json +++ b/projects/FlashKDA/status.json @@ -7,7 +7,7 @@ "priority": 0.0, "ext_type": "torch-extension", "adopted_at": "2026-08-07T17:05:59Z", - "updated_at": "2026-08-20T06:33:47Z", + "updated_at": "2026-08-20T06:42:24Z", "head_sha": null, "depends_on": [], "porting": null, @@ -18,9 +18,9 @@ "intake": { "verdict": "decline", "reason": "ported-elsewhere", - "duplicate_effort": "ROCm/aiter ships Kimi Delta Attention on ROCm: aiter/ops/triton/kimi_delta_attn/ (chunk_kimi_delta_attn, forward-only, mirrors fla.ops.kda.chunk_kda) plus _triton_kernels/chunk_delta_attn/flash_kda.py (flash_kda_fwd), with tests and a benchmark; added 2026-08-10 (#4568), MI300 LDS OOM fix 2026-08-12 (#4671), perf work 2026-08-13 (#4683) -- AMD-official and running on MI300. No FlashKDA repo in AMD-Ecosystem or ROCm (404), no AMD-directed fork among 118, no rocm/hip/amd branch or PR upstream (10 open PRs, hip/rocm/amd search empty in all states), vllm-project/FlashKDA has 5 branches none AMD. fla's amd-mi300 CI is still 0 runs vs 3051 h100, so aiter not fla is the credible AMD route", + "duplicate_effort": "ROCm/aiter ships Kimi Delta Attention on ROCm: aiter/ops/triton/kimi_delta_attn/ (chunk_kimi_delta_attn, forward-only, mirrors fla.ops.kda.chunk_kda) plus _triton_kernels/chunk_delta_attn/flash_kda.py (flash_kda_fwd), with tests and a benchmark; added 2026-08-10 (#4568), MI300 LDS OOM fix 2026-08-12 (#4671), perf work 2026-08-13 (#4683) -- AMD-official and running on MI300. No FlashKDA repo in AMD-Ecosystem or ROCm (404), no AMD-directed fork among 118, no ROCm/HIP/AMD branch or PR upstream (10 open PRs, HIP/ROCm/AMD search empty in all states), vllm-project/FlashKDA has 5 branches none AMD. fla's amd-mi300 CI is still 0 runs vs 3051 h100, so aiter not fla is the credible AMD route", "viable": false, "summary": "MIT (tier 1) Hopper/Blackwell-only KDA forward kernel, entirely CUTLASS/CuTe with TMA as the design -- decline as ported-elsewhere, NEW since the earlier screens: ROCm/aiter (AMD's own library) already ships the same op on ROCm as Triton kimi_delta_attn/chunk_kimi_delta_attn with a flash_kda fused path, tests, benchmark and an MI300 LDS fix, forward-only exactly like upstream's single fwd() export, so a port would duplicate AMD's own work; the older cant-port argument (no ROCm CuTe backend, no AMD TMA, from-scratch CK rewrite on any fleet platform) still holds independently, and if more KDA on AMD is wanted the cheap path is contributing backward/decode to aiter", - "at": "2026-08-20T06:33:47Z" + "at": "2026-08-20T06:42:24Z" } } From 1e27445ac061e4b2b0ee1eba312a02a318b8c4de Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Thu, 20 Aug 2026 19:14:49 +0000 Subject: [PATCH 10/10] FlashKDA: retire after skip decision -- ported-elsewhere (ROCm/aiter ships KDA) Jeff Daily decided in-session on 2026-08-20 to record the intake recommendation as the disposition: skip, ported-elsewhere. ROCm/aiter already ships Kimi Delta Attention on ROCm (kimi_delta_attn Triton ops with a flash_kda fused path, tests, benchmark, MI300 fixes), and the CUTLASS/CuTe+TMA cant-port argument holds independently. This ends the re-dispatch cycle five intake sessions circled. Following the 27f7646 precedent: the decision lives in data/dispositions.json (with by/reason/note and the retrieval command), the 10 telemetry records move to data/retired_stats.jsonl with a project field, the folder is removed, and the README table is regenerated. Notes retrievable: git show 5729ef615ea8:projects/FlashKDA/notes.md --- README.md | 28 +- data/dispositions.json | 9 + data/retired_stats.jsonl | 10 + projects/FlashKDA/notes.md | 546 ---------------------------------- projects/FlashKDA/stats.jsonl | 10 - projects/FlashKDA/status.json | 26 -- 6 files changed, 33 insertions(+), 596 deletions(-) delete mode 100644 projects/FlashKDA/notes.md delete mode 100644 projects/FlashKDA/stats.jsonl delete mode 100644 projects/FlashKDA/status.json diff --git a/README.md b/README.md index e9d6d117..432354ae 100644 --- a/README.md +++ b/README.md @@ -147,11 +147,11 @@ The project name links upstream. | [3P-ADMM-PC2](https://github.com/Samarvivian/3P-ADMM-PC2) ([fork](https://github.com/AMD-Ecosystem/3P-ADMM-PC2/tree/moat-port)) | ✅ | ✅ | ✅ | 🟣 [#10](https://github.com/Samarvivian/3P-ADMM-PC2/pull/10) | | [accelerated-scan](https://github.com/proger/accelerated-scan) ([fork](https://github.com/AMD-Ecosystem/accelerated-scan/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#17](https://github.com/proger/accelerated-scan/pull/17) | | [aihwkit](https://github.com/IBM/aihwkit) ([fork](https://github.com/AMD-Ecosystem/aihwkit/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#770](https://github.com/IBM/aihwkit/pull/770) | -| [alien](https://github.com/chrxh/alien) ([fork](https://github.com/AMD-Ecosystem/alien/tree/moat-port)) | ✅ | ✅ | 🔄 | 🟣 [#710](https://github.com/chrxh/alien/pull/710) | +| [alien](https://github.com/chrxh/alien) ([fork](https://github.com/AMD-Ecosystem/alien/tree/moat-port)) | ✅ | ✅ | ✅ | 🟣 [#710](https://github.com/chrxh/alien/pull/710) | | [amgcl](https://github.com/ddemidov/amgcl) ([fork](https://github.com/AMD-Ecosystem/amgcl/tree/moat-port)) | ✅ | ✅ | ✅ | 🟣 [#315](https://github.com/ddemidov/amgcl/pull/315) | | [anari-visionaray](https://github.com/szellmann/anari-visionaray) ([fork](https://github.com/AMD-Ecosystem/anari-visionaray/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#10](https://github.com/szellmann/anari-visionaray/pull/10) | | [arbor](https://github.com/arbor-sim/arbor) ([fork](https://github.com/AMD-Ecosystem/arbor/tree/moat-port)) | ✅ | ✅ | ✅ | 🟣 [#2512](https://github.com/arbor-sim/arbor/pull/2512) | -| [arrayfire](https://github.com/arrayfire/arrayfire) ([fork](https://github.com/AMD-Ecosystem/arrayfire/tree/moat-port)) | ✅ | ✅ | 🔄 | 🟢 [#3708](https://github.com/arrayfire/arrayfire/pull/3708) | +| [arrayfire](https://github.com/arrayfire/arrayfire) ([fork](https://github.com/AMD-Ecosystem/arrayfire/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#3708](https://github.com/arrayfire/arrayfire/pull/3708) | | [AutoDock-GPU](https://github.com/ccsb-scripps/AutoDock-GPU) ([fork](https://github.com/AMD-Ecosystem/AutoDock-GPU/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#320](https://github.com/ccsb-scripps/AutoDock-GPU/pull/320) | | [bam](https://github.com/ZaidQureshi/bam) ([fork](https://github.com/AMD-Ecosystem/bam/tree/moat-port)) | 🚫 | 🚫 | 🚫 | ⏸ on hold | | [barney](https://github.com/NVIDIA/barney) ([fork](https://github.com/AMD-Ecosystem/barney/tree/moat-port)) | ✅ | ✅ | ✅ | 🟣 [#46](https://github.com/NVIDIA/barney/pull/46) | @@ -161,7 +161,7 @@ The project name links upstream. | [catboost](https://github.com/catboost/catboost) ([fork](https://github.com/AMD-Ecosystem/catboost-moat/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#3111](https://github.com/catboost/catboost/pull/3111) | | [colmap](https://github.com/colmap/colmap) ([fork](https://github.com/AMD-Ecosystem/colmap/tree/moat-port)) | ✅ | ✅ | 🎫 | 🟢 [#4635](https://github.com/colmap/colmap/pull/4635) | | [CPM.cu](https://github.com/OpenBMB/CPM.cu) | 🚫 | — | — | ⚪ not-portable | -| [CubbyFlow](https://github.com/utilForever/CubbyFlow) ([fork](https://github.com/AMD-Ecosystem/CubbyFlow/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#145](https://github.com/utilForever/CubbyFlow/pull/145) | +| [CubbyFlow](https://github.com/utilForever/CubbyFlow) ([fork](https://github.com/AMD-Ecosystem/CubbyFlow/tree/moat-port)) | 🔄 | ✅ | 🔄 | 🟢 [#145](https://github.com/utilForever/CubbyFlow/pull/145) | | [cuBQL](https://github.com/NVIDIA/cuBQL) ([fork](https://github.com/AMD-Ecosystem/cuBQL/tree/moat-port)) | ✅ | ✅ | ✅ | 🟣 [#35](https://github.com/NVIDIA/cuBQL/pull/35) | | [cubvh](https://github.com/ashawkey/cubvh) ([fork](https://github.com/AMD-Ecosystem/cubvh/tree/moat-port)) | ✅ | ✅ | ✅ | 🟣 [#33](https://github.com/ashawkey/cubvh/pull/33) | | [cuda-efficient-features](https://github.com/fixstars/cuda-efficient-features) ([fork](https://github.com/AMD-Ecosystem/cuda-efficient-features/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#3](https://github.com/fixstars/cuda-efficient-features/pull/3) | @@ -174,14 +174,14 @@ The project name links upstream. | [cuPDLP-C](https://github.com/COPT-Public/cuPDLP-C) ([fork](https://github.com/AMD-Ecosystem/cuPDLP-C/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#41](https://github.com/COPT-Public/cuPDLP-C/pull/41) | | [cuPDLPx](https://github.com/MIT-Lu-Lab/cuPDLPx) ([fork](https://github.com/AMD-Ecosystem/cuPDLPx/tree/moat-port)) | ✅ | ✅ | ✅ | 🟣 [#94](https://github.com/MIT-Lu-Lab/cuPDLPx/pull/94) | | [cupoch](https://github.com/neka-nat/cupoch) ([fork](https://github.com/AMD-Ecosystem/cupoch/tree/moat-port)) | ✅ | ✅ | ✅ | 🟣 [#143](https://github.com/neka-nat/cupoch/pull/143) | -| [CuRast](https://github.com/m-schuetz/CuRast) ([fork](https://github.com/AMD-Ecosystem/CuRast/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#2](https://github.com/m-schuetz/CuRast/pull/2) | +| [CuRast](https://github.com/m-schuetz/CuRast) ([fork](https://github.com/AMD-Ecosystem/CuRast/tree/moat-port)) | 🔄 | ✅ | 🔄 | 🟢 [#2](https://github.com/m-schuetz/CuRast/pull/2) | | [cuSZ](https://github.com/szcompressor/cuSZ) ([fork](https://github.com/AMD-Ecosystem/cuSZ/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#91](https://github.com/szcompressor/cuSZ/pull/91) | | [CV-CUDA](https://github.com/CVCUDA/CV-CUDA) ([fork](https://github.com/AMD-Ecosystem/CV-CUDA/tree/moat-port)) | ✅ | ✅ | 🚫 | 🟢 [#293](https://github.com/CVCUDA/CV-CUDA/pull/293) | | [DEM-Engine](https://github.com/projectchrono/DEM-Engine) ([fork](https://github.com/AMD-Ecosystem/DEM-Engine/tree/moat-port)) | ✅ | ✅ | ✅ | 🟣 [#69](https://github.com/projectchrono/DEM-Engine/pull/69) | | [dgSPARSE-Lib](https://github.com/dgSPARSE/dgSPARSE-Lib) ([fork](https://github.com/AMD-Ecosystem/dgSPARSE-Lib/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#39](https://github.com/dgSPARSE/dgSPARSE-Lib/pull/39) | | [dietgpu](https://github.com/facebookresearch/dietgpu) ([fork](https://github.com/AMD-Ecosystem/dietgpu/tree/moat-port)) | ✅ | ✅ | ✅ | — | | [diff-surfel-rasterizations](https://github.com/xbillowy/diff-surfel-rasterizations) ([fork](https://github.com/AMD-Ecosystem/diff-surfel-rasterizations/tree/moat-port)) | ✅ | ✅ | ✅ | — | -| [diff-surfel-tracing](https://github.com/xbillowy/diff-surfel-tracing) ([fork](https://github.com/AMD-Ecosystem/diff-surfel-tracing/tree/moat-port)) | ✅ | ✅ | ⬜ | — | +| [diff-surfel-tracing](https://github.com/xbillowy/diff-surfel-tracing) ([fork](https://github.com/AMD-Ecosystem/diff-surfel-tracing/tree/moat-port)) | ✅ | ✅ | 🔧 | — | | [DiffPhysDrone](https://github.com/HenryHuYu/DiffPhysDrone) ([fork](https://github.com/AMD-Ecosystem/DiffPhysDrone/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#45](https://github.com/HenryHuYu/DiffPhysDrone/pull/45) | | [DynOSAM](https://github.com/ACFR-RPG/DynOSAM) | ⬜ | ⬜ | ⬜ | — | | [egg.c](https://github.com/d0rc/egg.c) ([fork](https://github.com/AMD-Ecosystem/egg.c/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#8](https://github.com/d0rc/egg.c/pull/8) | @@ -189,7 +189,7 @@ The project name links upstream. | [EnvGS](https://github.com/zju3dv/EnvGS) ([fork](https://github.com/AMD-Ecosystem/EnvGS/tree/moat-port)) | ✅ | ✅ | ✅ | — | | [espresso](https://github.com/espressomd/espresso) ([fork](https://github.com/AMD-Ecosystem/espresso/tree/moat-port)) | ✅ | ✅ | 🚫 | — | | [evogp](https://github.com/EMI-Group/evogp) ([fork](https://github.com/AMD-Ecosystem/evogp/tree/moat-port)) | ✅ | ✅ | ✅ | 🟣 [#12](https://github.com/EMI-Group/evogp/pull/12) | -| [faiss](https://github.com/facebookresearch/faiss) ([fork](https://github.com/AMD-Ecosystem/faiss/tree/moat-port)) | 🔄 | 🔄 | 🔄 | — | +| [faiss](https://github.com/facebookresearch/faiss) ([fork](https://github.com/AMD-Ecosystem/faiss/tree/moat-port)) | ✅ | ✅ | ✅ | — | | [FaithC](https://github.com/Luo-Yihao/FaithC) ([fork](https://github.com/AMD-Ecosystem/FaithC/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#12](https://github.com/Luo-Yihao/FaithC/pull/12) | | [Fast-Poisson-Image-Editing](https://github.com/Trinkle23897/Fast-Poisson-Image-Editing) ([fork](https://github.com/AMD-Ecosystem/Fast-Poisson-Image-Editing/tree/moat-port)) | ✅ | ✅ | ✅ | 🟣 [#25](https://github.com/Trinkle23897/Fast-Poisson-Image-Editing/pull/25) | | [faster-gaussian-splatting](https://github.com/nerficg-project/faster-gaussian-splatting) ([fork](https://github.com/AMD-Ecosystem/faster-gaussian-splatting/tree/moat-port)) | ✅ | ✅ | ✅ | — | @@ -197,7 +197,7 @@ The project name links upstream. | [fdtd3d](https://github.com/zer011b/fdtd3d) ([fork](https://github.com/AMD-Ecosystem/fdtd3d/tree/moat-port)) | ✅ | ✅ | ✅ | — | | [ffpa-attn](https://github.com/xlite-dev/ffpa-attn) ([fork](https://github.com/AMD-Ecosystem/ffpa-attn/tree/moat-port)) | ✅ | ✅ | ✅ | 🟣 [#268](https://github.com/xlite-dev/ffpa-attn/pull/268) | | [FLAMEGPU2](https://github.com/FLAMEGPU/FLAMEGPU2) ([fork](https://github.com/AMD-Ecosystem/FLAMEGPU2/tree/moat-port)) | ✅ | ✅ | ✅ | ⏸ on hold | -| [FlashKDA](https://github.com/MoonshotAI/FlashKDA) | ⬜ | ⬜ | ⬜ | — | +| [FlashKDA](https://github.com/MoonshotAI/FlashKDA) | ⬜ | ⬜ | ⬜ | ⚪ ported-elsewhere | | [FlashMoE](https://github.com/osayamenja/FlashMoE) | 🚫 | — | — | ⚪ not-portable | | [FlashRT](https://github.com/flashrt-project/FlashRT) | 🚫 | — | — | ⚪ not-portable | | [foldmason](https://github.com/steineggerlab/foldmason) ([fork](https://github.com/AMD-Ecosystem/foldmason/tree/moat-port)) | ✅ | ✅ | ✅ | — | @@ -205,13 +205,13 @@ The project name links upstream. | [gaussian_splatting](https://github.com/joeyan/gaussian_splatting) ([fork](https://github.com/AMD-Ecosystem/gaussian_splatting/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#37](https://github.com/joeyan/gaussian_splatting/pull/37) | | [gdtk](https://github.com/gdtk-uq/gdtk) ([fork](https://github.com/AMD-Ecosystem/gdtk/tree/moat-port)) | ✅ | ✅ | ✅ | — | | [GOMC](https://github.com/GOMC-WSU/GOMC) ([fork](https://github.com/AMD-Ecosystem/GOMC/tree/moat-port)) | ✅ | ✅ | ✅ | — | -| [GooFit](https://github.com/GooFit/GooFit) ([fork](https://github.com/AMD-Ecosystem/GooFit/tree/moat-port)) | ✅ | ✅ | ⬜ | — | +| [GooFit](https://github.com/GooFit/GooFit) ([fork](https://github.com/AMD-Ecosystem/GooFit/tree/moat-port)) | ✅ | ✅ | 🚫 | — | | [gpu4pyscf](https://github.com/pyscf/gpu4pyscf) ([fork](https://github.com/AMD-Ecosystem/gpu4pyscf/tree/moat-port)) | ✅ | ✅ | 🚫 | — | | [GPU_IPC](https://github.com/KemengHuang/GPU_IPC) ([fork](https://github.com/AMD-Ecosystem/GPU_IPC/tree/moat-port)) | ✅ | ✅ | 🚫 | — | | [Gpufit](https://github.com/gpufit/Gpufit) ([fork](https://github.com/AMD-Ecosystem/Gpufit/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#149](https://github.com/gpufit/Gpufit/pull/149) | | [GPUMD](https://github.com/brucefan1983/GPUMD) ([fork](https://github.com/AMD-Ecosystem/GPUMD/tree/moat-port)) | ✅ | ✅ | ✅ | 🟣 [#1538](https://github.com/brucefan1983/GPUMD/pull/1538) | | [gpuRIR](https://github.com/DavidDiazGuerra/gpuRIR) ([fork](https://github.com/AMD-Ecosystem/gpuRIR/tree/moat-port)) | ✅ | ✅ | ✅ | — | -| [gRASPA](https://github.com/snurr-group/gRASPA) ([fork](https://github.com/AMD-Ecosystem/gRASPA/tree/moat-port)) | ✅ | ✅ | 🔄 | — | +| [gRASPA](https://github.com/snurr-group/gRASPA) ([fork](https://github.com/AMD-Ecosystem/gRASPA/tree/moat-port)) | ✅ | ✅ | ✅ | — | | [gtsam_points](https://github.com/koide3/gtsam_points) ([fork](https://github.com/AMD-Ecosystem/gtsam_points/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#99](https://github.com/koide3/gtsam_points/pull/99) | | [h2o4gpu](https://github.com/h2oai/h2o4gpu) | ⬜ | ⬜ | ⬜ | — | | [HEonGPU](https://github.com/Alisah-Ozcan/HEonGPU) ([fork](https://github.com/AMD-Ecosystem/HEonGPU/tree/moat-port)) | ✅ | ✅ | ✅ | — | @@ -242,7 +242,7 @@ The project name links upstream. | [OCTproZ](https://github.com/spectralcode/OCTproZ) ([fork](https://github.com/AMD-Ecosystem/OCTproZ/tree/moat-port)) | ✅ | ✅ | ✅ | — | | [ohm](https://github.com/csiro-robotics/ohm) ([fork](https://github.com/AMD-Ecosystem/ohm/tree/moat-port)) | ✅ | ✅ | ✅ | — | | [oneflow](https://github.com/Oneflow-Inc/oneflow) ([fork](https://github.com/AMD-Ecosystem/oneflow/tree/moat-port)) | ✅ | ✅ | 🚫 | — | -| [op43dgs](https://github.com/LetianHuang/op43dgs) ([fork](https://github.com/AMD-Ecosystem/op43dgs/tree/moat-port)) | ✅ | ✅ | ✅ | — | +| [op43dgs](https://github.com/LetianHuang/op43dgs) ([fork](https://github.com/AMD-Ecosystem/op43dgs/tree/moat-port)) | 🔄 | ✅ | 🔄 | — | | [Open3D](https://github.com/isl-org/Open3D) ([fork](https://github.com/AMD-Ecosystem/Open3D/tree/moat-port)) | ✅ | ✅ | 🔄 | 🟢 [#7509](https://github.com/isl-org/Open3D/pull/7509) | | [opencv](https://github.com/opencv/opencv) ([fork](https://github.com/AMD-Ecosystem/opencv/tree/moat-port)) | 🚫 | 🚫 | 🚫 | 🟢 [#29285](https://github.com/opencv/opencv/pull/29285) | | [opencv_contrib](https://github.com/opencv/opencv_contrib) ([fork](https://github.com/AMD-Ecosystem/opencv_contrib/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#4147](https://github.com/opencv/opencv_contrib/pull/4147) | @@ -250,13 +250,13 @@ The project name links upstream. | [plumed2](https://github.com/plumed/plumed2) ([fork](https://github.com/AMD-Ecosystem/plumed2/tree/moat-port)) | ✅ | ✅ | 🚫 | — | | [plvs](https://github.com/luigifreda/plvs) ([fork](https://github.com/AMD-Ecosystem/plvs/tree/moat-port)) | ✅ | ✅ | 🚫 | — | | [Pointcept](https://github.com/Pointcept/Pointcept) ([fork](https://github.com/AMD-Ecosystem/Pointcept/tree/moat-port)) | ✅ | ✅ | ✅ | 🟣 [#604](https://github.com/Pointcept/Pointcept/pull/604) | -| [popsift](https://github.com/alicevision/popsift) ([fork](https://github.com/AMD-Ecosystem/popsift/tree/moat-port)) | ✅ | ✅ | 🔄 | 🟢 [#186](https://github.com/alicevision/popsift/pull/186) | +| [popsift](https://github.com/alicevision/popsift) ([fork](https://github.com/AMD-Ecosystem/popsift/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#186](https://github.com/alicevision/popsift/pull/186) | | [prismatic](https://github.com/prism-em/prismatic) ([fork](https://github.com/AMD-Ecosystem/prismatic/tree/moat-port)) | ✅ | ✅ | ✅ | — | | [pytorch3d](https://github.com/facebookresearch/pytorch3d) | ✅ | ✅ | ✅ | 🟣 [#2039](https://github.com/facebookresearch/pytorch3d/pull/2039) | | [qrack](https://github.com/unitaryfoundation/qrack) ([fork](https://github.com/AMD-Ecosystem/qrack/tree/moat-port)) | ✅ | ✅ | ✅ | — | -| [Quest](https://github.com/mit-han-lab/Quest) ([fork](https://github.com/AMD-Ecosystem/Quest/tree/moat-port)) | 🔄 | ✅ | 🔄 | — | +| [Quest](https://github.com/mit-han-lab/Quest) ([fork](https://github.com/AMD-Ecosystem/Quest/tree/moat-port)) | ✅ | ✅ | ✅ | — | | [QUICK](https://github.com/merzlab/QUICK) ([fork](https://github.com/AMD-Ecosystem/QUICK/tree/moat-port)) | ✅ | ✅ | 🚫 | — | -| [rmagine](https://github.com/uos/rmagine) ([fork](https://github.com/AMD-Ecosystem/rmagine/tree/moat-port)) | ✅ | ✅ | 🔄 | — | +| [rmagine](https://github.com/uos/rmagine) ([fork](https://github.com/AMD-Ecosystem/rmagine/tree/moat-port)) | ✅ | ✅ | ✅ | — | | [rmcl](https://github.com/uos/rmcl) ([fork](https://github.com/AMD-Ecosystem/rmcl/tree/moat-port)) | ✅ | ✅ | ⬜ | — | | [RWKV-CUDA](https://github.com/BlinkDL/RWKV-CUDA) ([fork](https://github.com/AMD-Ecosystem/RWKV-CUDA/tree/moat-port)) | ✅ | ✅ | ✅ | — | | [RXMesh](https://github.com/owensgroup/RXMesh) ([fork](https://github.com/AMD-Ecosystem/RXMesh/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#73](https://github.com/owensgroup/RXMesh/pull/73) | @@ -266,7 +266,7 @@ The project name links upstream. | [splatad](https://github.com/carlinds/splatad) ([fork](https://github.com/AMD-Ecosystem/splatad/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#24](https://github.com/carlinds/splatad/pull/24) | | [sppark](https://github.com/supranational/sppark) ([fork](https://github.com/AMD-Ecosystem/sppark/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#82](https://github.com/supranational/sppark/pull/82) | | [stdgpu](https://github.com/stotko/stdgpu) ([fork](https://github.com/AMD-Ecosystem/stdgpu/tree/moat-port)) | ✅ | ✅ | ✅ | 🟣 [#484](https://github.com/stotko/stdgpu/pull/484) | -| [symforce](https://github.com/symforce-org/symforce) ([fork](https://github.com/AMD-Ecosystem/symforce/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#465](https://github.com/symforce-org/symforce/pull/465) | +| [symforce](https://github.com/symforce-org/symforce) ([fork](https://github.com/AMD-Ecosystem/symforce/tree/moat-port)) | 🔄 | ✅ | 🔄 | 🟢 [#465](https://github.com/symforce-org/symforce/pull/465) | | [TIGRE](https://github.com/CERN/TIGRE) ([fork](https://github.com/AMD-Ecosystem/TIGRE/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#747](https://github.com/CERN/TIGRE/pull/747) | | [tiny-vllm](https://github.com/jmaczan/tiny-vllm) ([fork](https://github.com/AMD-Ecosystem/tiny-vllm/tree/moat-port)) | ✅ | ✅ | ✅ | 🟣 [#2](https://github.com/jmaczan/tiny-vllm/pull/2) | | [torch-linear-assignment](https://github.com/ivan-chai/torch-linear-assignment) ([fork](https://github.com/AMD-Ecosystem/torch-linear-assignment/tree/moat-port)) | ✅ | ✅ | ✅ | 🟢 [#31](https://github.com/ivan-chai/torch-linear-assignment/pull/31) | diff --git a/data/dispositions.json b/data/dispositions.json index 958ac79c..78a9e46d 100644 --- a/data/dispositions.json +++ b/data/dispositions.json @@ -1222,6 +1222,15 @@ "reason": "already-supported", "repo_id": 780323643 }, + "moonshotai/flashkda": { + "by": "Jeff Daily", + "decided": "2026-08-20T19:12:56Z", + "disposition": "skip", + "full_name": "MoonshotAI/FlashKDA", + "note": "Decided in-session 2026-08-20: skip as ported-elsewhere per the intake recommendation -- ROCm/aiter already ships Kimi Delta Attention on ROCm (kimi_delta_attn Triton ops with a flash_kda fused path, tests, benchmark, MI300 fixes), so a port would duplicate AMD's own work; the CUTLASS/CuTe+TMA cant-port argument holds independently. Analysis retained in history: `git show 5729ef615ea8:projects/FlashKDA/notes.md` (notes.md).", + "reason": "ported-elsewhere", + "repo_id": 1215986180 + }, "mrshaw01/software-engineer": { "decided": "2026-05-30T00:45:17Z", "disposition": "skip", diff --git a/data/retired_stats.jsonl b/data/retired_stats.jsonl index f1e15710..dbb1b889 100644 --- a/data/retired_stats.jsonl +++ b/data/retired_stats.jsonl @@ -86,3 +86,13 @@ {"kind": "tokens", "ts": "2026-06-16T06:29:17Z", "tokens": 41627, "source": "planner", "project": "mirage"} {"kind": "tokens", "ts": "2026-06-04T22:08:22Z", "tokens": 106800, "source": "planner", "project": "sparser-faster-llms"} {"kind": "tokens", "ts": "2026-06-11T04:17:11Z", "tokens": 83059, "source": "planner", "project": "spconv"} +{"kind": "session", "ts": "2026-08-13T03:12:45Z", "epoch": 1786590765.1385188, "event": "start", "platform": "linux-gfx1100", "project": "FlashKDA"} +{"kind": "session", "ts": "2026-08-13T03:19:12Z", "epoch": 1786591152.066405, "event": "end", "platform": "linux-gfx1100", "project": "FlashKDA"} +{"kind": "session", "ts": "2026-08-14T06:59:31Z", "epoch": 1786690771.859204, "event": "start", "platform": "linux-gfx90a", "project": "FlashKDA"} +{"kind": "session", "ts": "2026-08-14T07:01:55Z", "epoch": 1786690915.6831388, "event": "end", "platform": "linux-gfx90a", "project": "FlashKDA"} +{"kind": "session", "ts": "2026-08-14T14:22:06Z", "epoch": 1786717326.587488, "event": "start", "platform": "linux-gfx942", "project": "FlashKDA"} +{"kind": "session", "ts": "2026-08-14T14:25:21Z", "epoch": 1786717521.4275084, "event": "end", "platform": "linux-gfx942", "project": "FlashKDA"} +{"kind": "session", "ts": "2026-08-19T05:01:27Z", "epoch": 1787115687.7499385, "event": "start", "platform": "linux-gfx1100", "project": "FlashKDA"} +{"kind": "session", "ts": "2026-08-19T05:04:05Z", "epoch": 1787115845.0577497, "event": "end", "platform": "linux-gfx1100", "project": "FlashKDA"} +{"kind": "session", "ts": "2026-08-20T06:29:26Z", "epoch": 1787207366.5784564, "event": "start", "platform": "linux-gfx90a", "project": "FlashKDA"} +{"kind": "session", "ts": "2026-08-20T06:33:56Z", "epoch": 1787207636.7788186, "event": "end", "platform": "linux-gfx90a", "project": "FlashKDA"} diff --git a/projects/FlashKDA/notes.md b/projects/FlashKDA/notes.md deleted file mode 100644 index 0015eec0..00000000 --- a/projects/FlashKDA/notes.md +++ /dev/null @@ -1,546 +0,0 @@ -# FlashKDA notes - -## Intake screen (2026-08-13, linux-gfx1100) - -Upstream: https://github.com/MoonshotAI/FlashKDA — "FlashKDA: Flash Kimi Delta -Attention", high-performance KDA attention kernels built on CUTLASS/CuTe. Screened -from a shallow clone at `agent_space/FlashKDA-screen`; no fork exists and none was -requested. - -**Recommendation: decline, reason `cant-port`.** The argument is below. It is a -recommendation, not a decision — a person answers it through the intake queue. - -### Licence - -`license_spdx = MIT` (tier 1, cleared to contribute). Confirmed by reading `LICENSE` -in the tree, not just GitHub's field: verbatim MIT, "Copyright (c) 2026 MoonshotAI". -GitHub's API agrees (`MIT`). - -`scan-nvidia` over the main tree is clean. - -**One finding a person must rule on if this is ever adopted.** The repo carries a -submodule, `cutlass` -> https://github.com/NVIDIA/cutlass.git, pinned at -`5c149f52a436782210263fb2f19b354443a61c6a`. That tree is dual-licensed per-part: - -- `LICENSE.txt` is BSD-3-Clause and covers the C++ headers, and it says explicitly - that "the files located in the `python/CuTeDSL` directory are licensed under the - NVIDIA End User License Agreement (EULA)". -- `EULA.txt` is that agreement. `scan-nvidia` over the submodule flags ~40 files, all - of them under `python/CuTeDSL/`, and none anywhere else. - -The EULA restricts the Software to "systems with NVIDIA GPUs" (1.1) and clause 2.11 -forbids reverse-engineering output artifacts "for the purpose of translating such -output artifacts to target a non-NVIDIA platform" — which names the activity a -CUDA-to-ROCm port performs. It applies only to the CuTeDSL Python package. - -FlashKDA does not use that package. Its build pulls in only the BSD-3-Clause C++ -headers: `setup.py` adds `cutlass/include`, `cutlass/examples/common`, and -`cutlass/tools/util/include` to the include path, and every source include is a -`cute/...` or `cutlass/...` C++ header. Nothing imports CuTeDSL, and the port would -neither modify nor redistribute it. - -So the finding is almost certainly benign — but per the intake role an NVIDIA -proprietary licence on any file needs a person's decision before proceeding, and I do -not clear it myself. If the decline is upheld the question is moot; if someone -overrides to fork, this must be answered first and not skipped. - -### Duplicate effort - -No AMD or ROCm effort on FlashKDA itself: - -- No `FlashKDA` repository in AMD-Ecosystem (checked the org listing; the near-name - matches are `flashinfer`, `flashinfer-bench`, `ffpa-attn` — different projects). -- Nothing matching in the ROCm org. -- GitHub repo search for `flashkda` returns upstream plus `vllm-project/FlashKDA` - (12 stars, branches `master`/`dev` plus two feature branches, no ROCm branch), - two empty personal forks, and `popfido/FlashKDA-mlx` — an Apple MLX port, so there - is porting interest, but none toward AMD. -- `grep -rniE 'amd|rocm|hip|gfx[0-9]' README* docs/ BENCHMARK*.md` yields one hit, - and it is the substring in "on-chip". No notable-forks section, no platform ports. -- No MOAT disposition, no opt-out, no other `port/` branch for it or for a related - project. - -**The capability, though, already reaches AMD by another route.** FlashKDA is not a -standalone library — it is an optional backend for `flash-linear-attention`'s -`chunk_kda`, auto-dispatched when installed and disabled with `FLA_FLASH_KDA=0`, at -which point FLA falls back to its own Triton kernels. Those Triton kernels -(`fla/ops/kda/`) run on ROCm, and fla-org tests them on AMD in CI — -`.github/workflows/amd-mi300.yml`. So an AMD user wanting Kimi Delta Attention has a -supported, upstream-CI-tested path today. -[**CORRECTED 2026-08-14 by the gfx942 screen: the CI half of this is wrong.** That -workflow is disabled (`if: false`) and has never run — 0 runs. The Triton path is -installable on ROCm but upstream-unvalidated on AMD hardware. See the gfx942 section.] -This is not `already-supported` for -FlashKDA, which has no AMD port at all, but it changes what a FlashKDA port would -buy: speed on hardware-specific code, not access to a missing capability. - -### Viability - -Genuine CUDA, and unusually concentrated: a torch extension of ~2,300 lines across -six files, all of the GPU work in `csrc/smxx/`. - -The problem is what it is built on. - -- **Every kernel is CuTe/CUTLASS.** `csrc/smxx/utils.cuh` alone pulls 19 CUTLASS - headers, including `cute/arch/copy_sm90_tma.hpp`, `cute/arch/cluster_sm90.hpp`, - `cutlass/cluster_launch.hpp`, `cutlass/arch/barrier.h`, and - `cutlass/pipeline/sm90_pipeline.hpp`. Layouts, tensors, and the GEMMs are all CuTe - types. NVIDIA CUTLASS has no ROCm/HIP backend; AMD's analogue is Composable Kernel, - a different library with a different API. Porting therefore is not translating a - CUDA path — it is reimplementing the kernels against a different tile library. -- **TMA is pervasive, not incidental.** 25 `SM90_TMA` references and 24 `make_tma` - call sites, plus `CUTE_GRID_CONSTANT` TMA descriptor parameters threaded through - every kernel entry point (`tma_load_q`, `tma_load_k`, `tma_store_ws_*`, ...), and - `ClusterTransactionBarrier` in shared memory. The most recent upstream commit is - literally "fix missing proxy fences around TMA accesses". Neither CDNA3 nor RDNA3 - has a TMA equivalent; the data movement would have to be rewritten, and with it the - pipeline and barrier structure that is the design. -- **Upstream targets Hopper and Blackwell only.** `SUPPORTED_CUDA_ARCHS = - ["90a", "100a", "103a", "120a"]`, README requirement "SM90 and above", benchmark - files for GB200 and H20. -- **This host's platform is the worst fit of the family.** gfx1100 is RDNA3: wave32, - no MFMA, no async-copy analogue of any of this. The `wave64` gate would need CDNA - and would still need the same full rewrite; there is no cheap gate to satisfy. -- Three inline PTX asm uses (`ex2.approx.ftz.f32`, `tanh.approx.f32`, - `cvt.f32.bf16`) — trivial next to the above, noted only for completeness. - -That combination puts the work far outside MOAT's standing rule to build the smallest -complete port preserving upstream structure and the CUDA path. A from-scratch -Composable-Kernel or raw-HIP reimplementation of a Hopper-tuned attention kernel is a -new kernel project, and it would arrive upstream as a large unsolicited parallel -backend in a repo whose reason for existing is Hopper/Blackwell tuning — a poor -prospect for acceptance even if it were built and validated. - -**Dependencies.** None on any MOAT project, so `depends_on` stays empty. External -build dependencies are PyTorch (>= 2.4), CUDA 12.9+, and the vendored CUTLASS -submodule. The `flash-linear-attention` relationship is integration, not a build -dependency. - -**Upstream health.** Healthy, and not a factor in the decline: not archived, 1208 -stars, 114 forks, last push 2026-07-30, active commits through late July 2026. If the -recommendation is overridden, upstream is at least alive enough to receive a PR. - -### If a person overrides to fork - -Two things must happen before any porting work, in this order: rule on the CuTeDSL -EULA finding above, and accept that the first real task is a kernel reimplementation -rather than a translation — so the planner should scope it as such, and on a CDNA -host, not on gfx1100. - -## Intake re-screen (2026-08-14, linux-gfx90a) - -Independent second-host confirmation of the 2026-08-13 screen. **The decline -recommendation stands, unchanged: `cant-port`.** Verified against a fresh shallow -clone at upstream head `1ce47ea` ("optimize kda prepare cu_seqlens scan with -prefix-sum and binary search (#13)"), one commit ahead of the first screen. Nothing -material changed. - -Re-verified, each independently rather than taken on trust: - -- Licence MIT, tier 1 (`licenses.py check` → `license=MIT tier=1`). `scan-nvidia` - over the main tree: clean. -- The `cutlass` submodule is still pinned at `5c149f52`, and that pin's `LICENSE.txt` - still carries the per-part clause ("The files located in the `python/CuTeDSL` - directory are licensed under the NVIDIA End User License Agreement"), with - `EULA.txt` alongside it. FlashKDA still uses none of it: no `cutedsl`, - `cutlass.cute`, or `import cutlass` anywhere in `csrc/`, `flash_kda/`, `setup.py`, - `tests/`, or `benchmarks/`, and `setup.py` adds only the BSD-3-Clause C++ include - paths. **Still a person's ruling, still not cleared here.** -- No AMD-Ecosystem or ROCm FlashKDA repo (both 404). Repo search adds nothing new - toward AMD. Note for whoever re-runs this: a fork-list regex for `amd` matches - `Shamdon/FlashKDA` on the substring in the owner's name — a false positive, not an - AMD fork. -- `grep -rniE 'amd|rocm|hip|gfx[0-9]' README* docs/ BENCHMARK*.md` still yields - exactly one hit, still the substring in "on-chip". -- The AMD route to the capability is real and still live: `fla-org/flash-linear-attention` - has `fla/ops/kda/` (Triton, with a `backends/` subpackage) and - `.github/workflows/amd-mi300.yml`. [**CORRECTED by the gfx942 screen:** confirming - the workflow file exists is not confirming CI coverage. It is disabled and has - never run.] -- `SUPPORTED_CUDA_ARCHS = ["90a", "100a", "103a", "120a"]`; README requirements - "SM90 and above", CUDA 12.9+, PyTorch 2.4+. Source is 2,346 lines across six files. -- CuTe/TMA density confirmed: 24 `SM90_TMA` references, 24 `make_tma` sites, 8 - cluster/sm90-pipeline references. The first screen said 25 TMA sites; the correct - count is 24, so the queue summary is corrected. The difference changes nothing. - -**One refinement this host adds, and it is the reason the re-screen was worth -running.** The first screen was performed on gfx1100 and noted that RDNA3 is "the -worst fit of the family" — wave32, no MFMA. That invites the objection that the -decline is an artifact of screening on unsuitable hardware. It is not. This host is -gfx90a: CDNA2, wave64, MFMA present — the `wave64` gate architecture, and a far -better-suited target. The decline is unchanged from here, because it never rested on -the host's wavefront or matrix-core support. It rests on two things neither CDNA2 nor -CDNA3 changes: - -1. NVIDIA CUTLASS/CuTe has no ROCm backend, and every layout, tensor, and GEMM in - this codebase is a CuTe type. AMD's Composable Kernel is a different library with - a different API, not a drop-in. -2. TMA has no AMD equivalent on any current architecture, and TMA is the design here, - not an optimization layered on top — descriptors are threaded through every kernel - entry point as `CUTE_GRID_CONSTANT` parameters and the pipeline and barrier - structure is built around them. - -So a CDNA host would face the same from-scratch kernel reimplementation that gfx1100 -would. There is no platform in the fleet on which this becomes a translation rather -than a rewrite, which is precisely what `cant-port` means. - -**Dependencies re-checked:** `depends_on` stays empty. No MOAT project provides -CUTLASS or flash-linear-attention (checked against `moatlib.py projects` and -`DEPENDENCIES.md`), so there is no unknown hard dependency needing an intake request. -PyTorch, CUDA 12.9+, and the vendored CUTLASS submodule are external build -dependencies only. - -**Upstream health:** not archived, not disabled, 1211 stars, 115 forks, last push -2026-07-30. A PR would have a live destination if the recommendation were overridden. - -**Unrelated data defect noticed while screening (first noted 2026-08-14, gfx90a).** -`data/candidates.json` carries two entries named `FlashKDA`: -`MoonshotAI/FlashKDA` (the one screened here) and `0xwilliamortiz/FlashKDA` -(195 stars, MIT, "memory-efficient KDA kernels for training and decode"). The second -is a **404 on the GitHub API** — deleted or never public. Beyond being stale, it is a -name collision: adopting it would scaffold to the same `projects/FlashKDA` and the -same `port/FlashKDA` claim as this project. Not fixed here, since candidate curation -is not the intake role's to edit and the decline makes it non-urgent. - -## Intake re-screen (2026-08-14, linux-gfx942) - -Third-host screen, on CDNA3/MI300 — the architecture the previous two screens -pointed at as the best-case target and the one the claimed AMD fallback route names. -**The decline recommendation stands: `cant-port`.** But this screen corrects a -supporting fact that both earlier screens recorded as true, and that correction is -the reason this run was worth doing. - -Verified against a fresh shallow clone at upstream head `1ce47ea` — the same commit -the gfx90a screen saw, so the code has not moved. - -### Correction: the "CI-tested on MI300" claim is false - -Both prior screens, and the recorded `intake.duplicate_effort` and `summary` fields, -stated that `fla-org/flash-linear-attention` tests its Triton KDA kernels on AMD in -CI, citing `.github/workflows/amd-mi300.yml`. The file exists. It has never run. - -- The job is guarded `if: false`, and the file's own header comment says: "Disabled - by default: fla does not yet operate an amd-mi300 runner. Flip the `if:` guard - below (or replace with a real condition) once a runner is wired in. Workflow lives - here so the install / sanity-check pattern is documented." -- Its trigger is `workflow_dispatch` only — no push or PR trigger. -- Run count via the Actions API: **0**. For contrast, `nvidia-h100.yml` has 3028. - -The existence of a workflow file is not evidence of CI coverage. Checking the run -count is the cheap disambiguation, and it should be the habit whenever an -existing-support claim rests on a CI file. - -**What is actually true about the AMD route**, stated at the strength the evidence -supports: fla ships KDA as Triton kernels (`fla/ops/kda/`, the default path) and -offers a documented ROCm install extra (`pip install -e ".[rocm]"`, which -deliberately does not pin a Triton flavor so the ROCm wheel index supplies -`pytorch-triton-rocm`). So the AMD path is intended and plausible — Triton targets -ROCm — but it is **unvalidated by upstream on AMD hardware**. Nobody should record -that AMD users have a tested KDA path today. - -This weakens the "capability already reaches AMD" argument rather than the decline. -It also points at the tractable work, which is worth saying plainly in the queue: if -someone wants Kimi Delta Attention working on MI300, the cheap, high-value task is -validating fla's existing Triton `fla/ops/kda` on gfx942 — a different project, and -one whose upstream has already built the scaffolding and is visibly waiting for a -runner. Rewriting FlashKDA's CuTe kernels is the expensive way to the same capability. - -### Re-verified independently on this host - -- Licence MIT, tier 1 (`licenses.py check` → `license=MIT tier=1`); `LICENSE` read - directly, verbatim MIT, "Copyright (c) 2026 MoonshotAI". `scan-nvidia` over the - main tree: clean. -- `cutlass` submodule still pinned at `5c149f52a436782210263fb2f19b354443a61c6a` - (`git ls-tree HEAD` gitlink), url `https://github.com/NVIDIA/cutlass.git`. The - per-part CuTeDSL EULA finding from the first screen is unchanged and **still not - cleared here — it remains a person's ruling if this is ever adopted.** FlashKDA - uses none of it: a case-insensitive search for `cutedsl`, `cutlass.cute`, - `import cutlass`, `from cutlass` across the whole tree returns nothing, and - `setup.py` adds only the BSD-3-Clause C++ include paths (`cutlass/include`, - `cutlass/examples/common`, `cutlass/tools/util/include`). -- No `FlashKDA` repo in AMD-Ecosystem or ROCm (both 404), nor `AMD-Ecosystem/flash-kda` - or `ROCm/flash-linear-attention`. Repo search returns upstream, `vllm-project/FlashKDA` - (12 stars), two empty personal repos, and `popfido/FlashKDA-mlx`. Nothing toward AMD. -- No MOAT disposition for FlashKDA in `data/dispositions.json` (282 entries), no - opt-out (`optout.py list` → nobody), no other `port/` branch. -- `SUPPORTED_CUDA_ARCHS = ["90a", "100a", "103a", "120a"]`; 2346 lines across six - source files; three inline PTX uses. -- CuTe/TMA density: 24 `SM90_TMA` uses, 24 `make_tma` sites, 22 `CUTE_GRID_CONSTANT` - parameters, 5 `ClusterTransactionBarrier`, plus `cutlass/pipeline/sm90_pipeline.hpp` - and `cutlass/cluster_launch.hpp`. - -**The 24-vs-25 TMA discrepancy between the first two screens is settled**, since it -cost a correction once already: case-sensitive `SM90_TMA` in `csrc/` is 24; -case-insensitive is 25, the extra match being the include filename -`cute/arch/copy_sm90_tma.hpp` on line 14 of `utils.cuh`. Both counts were right about -different questions. 24 is the number of code uses. - -### Why CDNA3 does not rescue it - -gfx942 is the strongest case AMD can make here — CDNA3, wave64, MFMA, and the -`wave64` gate architecture. The decline is unchanged from this host for the reason -the gfx90a screen already gave, which this host confirms rather than repeats: the -blocker is not wavefront width or matrix-core availability. It is that every layout, -tensor, and GEMM is a CuTe type and NVIDIA CUTLASS has no ROCm backend, and that TMA -is the design rather than an optimization — descriptors threaded through every kernel -entry point as `CUTE_GRID_CONSTANT` parameters, with the pipeline and barrier -structure built around them. CDNA3 has no TMA equivalent. Three hosts spanning RDNA3, -CDNA2 and CDNA3 now agree, which is as much as screening can establish: there is no -platform in the fleet on which this becomes a translation rather than a from-scratch -kernel reimplementation. That is what `cant-port` means. - -**Dependencies:** `depends_on` stays empty; no MOAT project provides CUTLASS or -flash-linear-attention. PyTorch, CUDA 12.9+, and the vendored CUTLASS submodule are -external build dependencies only. - -**Upstream health:** not archived, not disabled, 1211 stars, 115 forks, last push -2026-07-30, default branch `master`. A PR would have a live destination if the -recommendation were overridden. - -**Minor data observation, not load-bearing (gfx942 screen).** `data/retired_stats.jsonl` line 83 -carries a single token record for project `FlashKDA` dated 2026-06-04 with source -`porter` — two months before this project's `adopted_at` (2026-08-07) and before -upstream's own recorded activity. There is no matching disposition and no other -trace of a prior lifecycle. Flagged for whoever maintains telemetry; it does not -affect the recommendation and I did not edit it. - -## Intake delta-check (2026-08-19, linux-gfx1100) — fourth dispatch, NOT a fourth screen - -This host was dispatched FlashKDA at `stage: unclaimed` for the fourth time. Three -screens (2026-08-13 gfx1100, 2026-08-14 gfx90a, 2026-08-14 gfx942) already agree on -`decline` / `cant-port`, and that recommendation is already recorded in -`status.json.intake`. So this run deliberately did **not** re-derive the screen. It -checked only what could have changed since 2026-08-14, and stopped. - -**Nothing changed. The recommendation stands, unchanged: `decline`, `cant-port`.** - -Delta checks, all read-only, no clone: - -- **Upstream head is still `1ce47ea`** ("optimize kda prepare cu_seqlens scan with - prefix-sum and binary search (#13)", authored 2026-07-29), the same commit the - gfx90a and gfx942 screens verified. `master` is the only branch. The tree is - therefore byte-identical to what was screened three times, so every code-level - finding — 24 `SM90_TMA` uses, 24 `make_tma` sites, 22 `CUTE_GRID_CONSTANT` - descriptors, `SUPPORTED_CUDA_ARCHS = ["90a","100a","103a","120a"]`, 2346 lines - across six files, the `cutlass` submodule pinned at `5c149f52` — holds by - construction. Re-cloning to recount them would produce the same numbers from the - same bytes, which is why it was skipped rather than repeated. -- **Licence unchanged: MIT, tier 1.** GitHub API `license.spdx_id = MIT`; the - `LICENSE` file was read directly by all three prior screens at this same SHA. - `license_spdx` was already recorded and stays `MIT`. -- **The CuTeDSL EULA finding is unchanged and still uncleared.** Same submodule pin, - so same per-part licensing. It remains a person's ruling if this is ever adopted, - and this host does not clear it either. -- **No AMD effort appeared.** `AMD-Ecosystem/FlashKDA`, `ROCm/FlashKDA`, and - `AMD-Ecosystem/flash-kda` are all still 404. A fresh repo search surfaces one name - the earlier screens did not see, `Unitflexmed1821/FlashKDA` (0 stars, not a fork, - pushed 2026-08-19, head "Update README.md", description advertising CUTLASS - kernels). Its README has zero matches for `amd|rocm|hip|gfx[0-9]` — an unrelated - re-upload, not an AMD port. Also new: `atomicmilkshake/godzilla-llama.cpp`, an - MSVC+CUDA Windows llama.cpp fork carrying KDA, likewise nothing toward AMD. -- **Upstream still healthy**, still a live PR destination if the recommendation is - overridden: not archived, not disabled, 1218 stars (was 1211), 117 forks (was 115), - last push 2026-07-30, default branch `master`. -- `depends_on` stays empty. - -### Why this project keeps coming back, which is the finding worth acting on - -The screen is not stuck — the *decision* is. FlashKDA is row 1 of the open intake -queue, issue AMD-Ecosystem/moat#8, opened 2026-08-07 and carrying **zero comments -after 12 days**. Its recommendation has been correct and complete since 2026-08-13. - -An agent may not record a decline, so `stage` correctly stays `unclaimed` — and the -selector offers `unclaimed` projects to whatever host asks next. A decline -recommendation awaiting a person is therefore indistinguishable, to the selector, -from a project nobody has looked at. That is the whole mechanism, and it is the same -one that had spconv screened six times on four platforms; spconv is row 2 of the same -unanswered issue. Four sessions have now been spent on FlashKDA to reach the answer -the first one reached. - -Two things would stop it, and both belong to a person, not to this role: - -1. **Answer issue #8.** One reply decides this batch, and `intake_queue.py apply` - records it. This is the real fix; everything else is a workaround. -2. **`moatlib.py set-hold FlashKDA`** if the answer will be a while. AGENTS.md - reserves `set-hold` to a person, so this host did not set it — but a held project - is skipped by the selector on every platform with no state touched, which is - exactly the behaviour wanted for a screened project queued behind a pending - decision. - -Neither was done here. The only write this run made was refreshing the queue summary -so the row itself reports that four hosts agree and the screen is not the bottleneck. - -## Intake re-screen (2026-08-20, linux-gfx90a) — fifth dispatch, and the reason changes - -Fifth dispatch at `stage: unclaimed`. The decline is unchanged, but **the recorded -reason moves from `cant-port` to `ported-elsewhere`**, because a check no earlier -screen performed found that AMD itself now ships this capability on ROCm. The -technical `cant-port` argument still holds independently — see below — so the decline -is robust under either reason; `ported-elsewhere` is recorded because it is the -stronger and less arguable of the two, and because it names the standing rule that -applies: we do not duplicate AMD's own work. - -### The new finding: ROCm/aiter ships Kimi Delta Attention - -`ROCm/aiter` — AMD's own kernel library, MIT, "Copyright (C) 2024-2026, Advanced Micro -Devices, Inc." — contains a Triton KDA implementation, including a path named after -this very project: - -- `aiter/ops/triton/kimi_delta_attn/` — public entry point `chunk_kimi_delta_attn`, - documented as "Kimi Delta Attention Operations (Forward Only) ... for the KDA - linear-attention mixer used by Kimi-Linear / Kimi-K3. The chunked prefill op - mirrors `fla.ops.kda.chunk_kda`." -- `aiter/ops/triton/_triton_kernels/chunk_delta_attn/flash_kda.py` — the fused - variant, exporting `flash_kda_fwd` and `flash_kda_supported`. -- `op_tests/triton_tests/chunk_delta_attn/test_flash_kda.py` — tests the fused - FlashKDA path against the default chunk_delta_attn pipeline as reference, on - relative error because the two differ at bf16 level by design. -- `op_tests/op_benchmarks/triton/bench_flash_kda.py` — a benchmark. -- Listed in `aiter/ops/triton/README.md`: "kimi_delta_attn/ # Kimi Delta Attention - (chunked delta attention)". - -It is live work, not a stub, and it runs on real AMD hardware: - - 01ddb01 2026-08-10 [chunk_delta_attn] Add Triton kernels for chunk delta attention forward pass (#4568) - 0f6766f 2026-08-12 [Triton] Fix LDS OOM issue on MI300 (#4671) - 3679a25 2026-08-13 [Triton] Optimize chunk_delta_attn performance. (#4683) - -An LDS-OOM fix specific to MI300 is only written by someone running the kernel on -MI300. That is the hardware evidence the `fla` route never had. - -**Why the earlier screens missed it, which is the transferable lesson.** All four -prior screens searched for the project NAME (`FlashKDA` in AMD-Ecosystem, in ROCm, in -repo search) and for forks of the upstream. AMD's support is neither: it is a -differently-named op inside a large general-purpose kernel library, added on -2026-08-10 — after the three real screens (08-13, 08-14, 08-14) had already run, and -the 08-19 delta-check looked only at the project name again. Searching the CAPABILITY -inside `ROCm/aiter` (`repo:ROCm/aiter kda in:path`, `repo:ROCm/aiter "kimi delta"`) -is two API calls and is what found it. For any attention/GEMM/quantization kernel -project, `ROCm/aiter` should be searched by capability, not just the org by name. - -### Scope parity: it is the same op, not a loose equivalent - -Upstream FlashKDA's entire public API is one function. `flash_kda/__init__.py` exports -`fwd(q, k, v, g, beta, scale, out, A_log, dt_bias, lower_bound, initial_state, -final_state, cu_seqlens)` — forward only, bf16, K = V = 128 — backed by four CUDA -files in `csrc/smxx/` (`fwd_kernel1.cuh`, `fwd_kernel2.cuh`, `fwd_launch.cu`, -`utils.cuh`). The README's stated purpose is to be auto-dispatched as -`flash-linear-attention`'s `chunk_kda` backend (fla-org/flash-linear-attention#852). - -aiter's `chunk_kimi_delta_attn` is forward-only and explicitly "mirrors -`fla.ops.kda.chunk_kda`" — the same op, in the same role, on ROCm. Upstream itself is -forward-only too: training kernels are not in the tree, they are an unmerged -third-party PR (#28, "Add CUDA training kernels (fwd+bwd) for KDA"). So there is no -"but AMD only has part of it" gap to point at. - -Honest limits on the claim: aiter's implementation is Triton rather than CUTLASS-class -tuned, it is 10 days old, and forward/prefill only. It is authoritative (AMD-official, -actively maintained, MI300-fixed) rather than mature-by-age. Per -`assess-existing-support.md`, authoritativeness is the deciding axis, not age — an -AMD-official effort means reuse-and-improve or skip, never re-port from scratch. - -**Where the work goes if someone wants more KDA on AMD:** contribute to -`ROCm/aiter/aiter/ops/triton/kimi_delta_attn/` — backward/training, decode, or -performance. That is a different project with a different intake, and it is a far -cheaper path than reimplementing 2,346 lines of CuTe/TMA kernels against Composable -Kernel. - -### Delta-checks, all re-verified this run - -- **Licence: MIT, tier 1** (`licenses.py check MoonshotAI/FlashKDA` → - `license=MIT tier=1`, cleared to contribute). GitHub API `license.spdx_id = MIT`. - `license_spdx` stays `MIT`. Read directly from the `LICENSE` file by all three full - screens at this same SHA. -- **Upstream head is still `1ce47ea`** ("optimize kda prepare cu_seqlens scan with - prefix-sum and binary search (#13)", 2026-07-29), and `master` is still the only - branch. The tree is byte-identical to what was screened three times, so the - code-level counts (24 `SM90_TMA`, 24 `make_tma`, 22 `CUTE_GRID_CONSTANT`, - `SUPPORTED_CUDA_ARCHS = ["90a","100a","103a","120a"]`, 2346 lines / six files) hold - by construction and were not recounted. No shallow clone was taken for the same - reason; re-running `scan-nvidia` over identical bytes would reproduce the identical - result recorded on 08-13, 08-14 and 08-14. -- **The CuTeDSL EULA finding is unchanged and still uncleared.** The `cutlass` - submodule is still pinned at `5c149f52a436782210263fb2f19b354443a61c6a`, url - `https://github.com/NVIDIA/cutlass.git` (verified through the contents API at - `ref=1ce47ea`, no clone needed). Same pin means the same per-part licensing: the - ~40 `python/CuTeDSL/` files under the NVIDIA EULA, everything FlashKDA actually - builds against under BSD-3-Clause. Still a person's ruling if this is ever adopted; - this host does not clear it either. Moot if the decline is upheld. - -### Duplicate effort, including the open-PR/branch check no earlier screen ran - -This screen adds the upstream open-PR and branch check that the role now requires -(added after FLAMEGPU2 was adopted while the maintainer's own HIP draft PR sat open). - -- **Upstream open PRs: 10, none AMD-related.** `gh pr list --state open` returns #16, - #17, #18, #19, #27, #28, #29, #30, #31, #32 — numerics tolerance, store - parallelization, FLA validation, CUTLASS error reporting, a CPU training backend, a - CUDA training-kernel PR, tile-prefix reuse, CP transition, reusable workspace, and - empty varlen batches. `--search "hip OR rocm OR amd"` over open PRs returns **empty**, - over ALL states returns **empty**, and the same search over issues returns **empty**. - No maintainer HIP draft, no competing AMD PR. -- **Upstream branches: `master` only.** No rocm/hip/amd branch. -- `AMD-Ecosystem/FlashKDA`, `ROCm/FlashKDA`, `AMD-Ecosystem/flash-kda`, - `ROCm/flash-kda`, `AMD-Ecosystem/flash-linear-attention`, - `ROCm/flash-linear-attention`: all still 404. Org-scoped repo search for `kda` in - AMD-Ecosystem and in ROCm: 0 each — which is exactly why the name search keeps - missing aiter, whose repo name contains no "kda". -- **Forks: 118, none AMD-directed.** Scanning all fork full names for - `amd|rocm|hip|gfx[0-9]` yields two substring false positives and nothing else: - `Shamdon/FlashKDA` (known, "amd" inside the owner name) and, new this run, - `woshipapa/FlashKDA` ("hip" inside the owner name). Neither is an AMD port. -- `vllm-project/FlashKDA` (12 stars, pushed 2026-08-19) re-checked directly since vLLM - does care about ROCm: branches are `master` (at upstream's `1ce47ea`), `dev`, - `agent/export-intermediate-recurrent-checkpoints`, - `codex/eagle-dual-checkpoint-pr7`, `thien-codex/state-store-refactor`; open PRs #7 - and #8 are smem layout and recurrent-state export. Nothing toward AMD. -- Repo search for `flashkda` adds nothing new: upstream, vllm-project's copy, - `Unitflexmed1821/FlashKDA` and `allanleewh/FlashKDA` (re-uploads), - `popfido/FlashKDA-mlx` (Apple MLX), `atomicmilkshake/godzilla-llama.cpp` (Windows - MSVC+CUDA). -- No MOAT disposition for FlashKDA (282 entries, no `kda` key), no opt-out - (`optout.py list` → nobody), no other `port/` branch. - -**The `fla` amd-mi300 CI correction still stands, and has not moved.** Workflow run -counts re-read this run: `amd-mi300.yml` = **0 runs**, `nvidia-h100.yml` = **3051** -(was 3028 on 08-14). Both workflows report `state=active`, which is the trap — the job -body is guarded `if: false`, so "active" describes the file, not any execution. Read -the run count, never the state field. The `fla` Triton path remains -installable-but-unvalidated on AMD; aiter, not fla, is now the credible AMD route. - -### The `cant-port` argument, unchanged and still independently sufficient - -Recorded here so that changing the reason does not lose it. Every layout, tensor and -GEMM is a CuTe type and NVIDIA CUTLASS has no ROCm backend; TMA is the design rather -than an optimization, with descriptors threaded through every kernel entry point as -`CUTE_GRID_CONSTANT` parameters and the pipeline and barrier structure built around -them; no AMD architecture has a TMA equivalent. Three hosts spanning RDNA3 (gfx1100), -CDNA2 (gfx90a) and CDNA3 (gfx942) agreed there is no fleet platform on which this -becomes a translation rather than a from-scratch kernel reimplementation. If a person -prefers to record the decline on technical grounds instead, `cant-port` is equally -defensible — the answer is the same either way. - -### Viability and dependencies - -Genuine CUDA, confirmed again: four CUDA sources in `csrc/smxx/`, a `flash_kda_C` -compiled torch extension, `SUPPORTED_CUDA_ARCHS = ["90a","100a","103a","120a"]`, -CUDA 12.9+ / PyTorch 2.4+ requirements. `viable = no` for MOAT's purposes. - -`depends_on` stays empty. No MOAT project provides CUTLASS, aiter, or -flash-linear-attention, so there is no unknown hard dependency needing an intake -request. PyTorch, CUDA 12.9+ and the vendored CUTLASS submodule are external build -dependencies only. - -**Upstream health:** not archived, not disabled, 1221 stars (was 1218), 118 forks (was -117), last push 2026-07-30, default branch `master`. A PR would still have a live -destination if the recommendation were overridden. - -### Queue status - -Still row 1 of intake queue issue AMD-Ecosystem/moat#8, opened 2026-08-07, unanswered -after 13 days. Five sessions have now been spent reaching the answer the first one -reached — though this one earned its keep by finding the aiter implementation and by -running the open-PR check for the first time. The two things that would stop the -re-dispatch loop remain a person's: answer issue #8, or `moatlib.py set-hold FlashKDA` -while it waits. Neither was done here. No fork was requested, no disposition was -written, no PR was opened. diff --git a/projects/FlashKDA/stats.jsonl b/projects/FlashKDA/stats.jsonl deleted file mode 100644 index 4836b202..00000000 --- a/projects/FlashKDA/stats.jsonl +++ /dev/null @@ -1,10 +0,0 @@ -{"kind":"session","ts":"2026-08-13T03:12:45Z","epoch":1786590765.138518808,"event":"start","platform":"linux-gfx1100"} -{"kind":"session","ts":"2026-08-13T03:19:12Z","epoch":1786591152.066405121,"event":"end","platform":"linux-gfx1100"} -{"kind":"session","ts":"2026-08-14T06:59:31Z","epoch":1786690771.859204086,"event":"start","platform":"linux-gfx90a"} -{"kind":"session","ts":"2026-08-14T07:01:55Z","epoch":1786690915.683138792,"event":"end","platform":"linux-gfx90a"} -{"kind":"session","ts":"2026-08-14T14:22:06Z","epoch":1786717326.587487960,"event":"start","platform":"linux-gfx942"} -{"kind":"session","ts":"2026-08-14T14:25:21Z","epoch":1786717521.427508328,"event":"end","platform":"linux-gfx942"} -{"kind":"session","ts":"2026-08-19T05:01:27Z","epoch":1787115687.749938491,"event":"start","platform":"linux-gfx1100"} -{"kind":"session","ts":"2026-08-19T05:04:05Z","epoch":1787115845.057749640,"event":"end","platform":"linux-gfx1100"} -{"kind":"session","ts":"2026-08-20T06:29:26Z","epoch":1787207366.578456516,"event":"start","platform":"linux-gfx90a"} -{"kind":"session","ts":"2026-08-20T06:33:56Z","epoch":1787207636.778818635,"event":"end","platform":"linux-gfx90a"} diff --git a/projects/FlashKDA/status.json b/projects/FlashKDA/status.json deleted file mode 100644 index 62058fbd..00000000 --- a/projects/FlashKDA/status.json +++ /dev/null @@ -1,26 +0,0 @@ -{ - "schema_version": 3, - "name": "FlashKDA", - "upstream_url": "https://github.com/MoonshotAI/FlashKDA", - "fork_url": null, - "fork_default_branch": "main", - "priority": 0.0, - "ext_type": "torch-extension", - "adopted_at": "2026-08-07T17:05:59Z", - "updated_at": "2026-08-20T06:42:24Z", - "head_sha": null, - "depends_on": [], - "porting": null, - "waivers": {}, - "stage": "unclaimed", - "platforms": {}, - "license_spdx": "MIT", - "intake": { - "verdict": "decline", - "reason": "ported-elsewhere", - "duplicate_effort": "ROCm/aiter ships Kimi Delta Attention on ROCm: aiter/ops/triton/kimi_delta_attn/ (chunk_kimi_delta_attn, forward-only, mirrors fla.ops.kda.chunk_kda) plus _triton_kernels/chunk_delta_attn/flash_kda.py (flash_kda_fwd), with tests and a benchmark; added 2026-08-10 (#4568), MI300 LDS OOM fix 2026-08-12 (#4671), perf work 2026-08-13 (#4683) -- AMD-official and running on MI300. No FlashKDA repo in AMD-Ecosystem or ROCm (404), no AMD-directed fork among 118, no ROCm/HIP/AMD branch or PR upstream (10 open PRs, HIP/ROCm/AMD search empty in all states), vllm-project/FlashKDA has 5 branches none AMD. fla's amd-mi300 CI is still 0 runs vs 3051 h100, so aiter not fla is the credible AMD route", - "viable": false, - "summary": "MIT (tier 1) Hopper/Blackwell-only KDA forward kernel, entirely CUTLASS/CuTe with TMA as the design -- decline as ported-elsewhere, NEW since the earlier screens: ROCm/aiter (AMD's own library) already ships the same op on ROCm as Triton kimi_delta_attn/chunk_kimi_delta_attn with a flash_kda fused path, tests, benchmark and an MI300 LDS fix, forward-only exactly like upstream's single fwd() export, so a port would duplicate AMD's own work; the older cant-port argument (no ROCm CuTe backend, no AMD TMA, from-scratch CK rewrite on any fleet platform) still holds independently, and if more KDA on AMD is wanted the cheap path is contributing backward/decode to aiter", - "at": "2026-08-20T06:42:24Z" - } -}