Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
44 changes: 27 additions & 17 deletions CONTRIBUTING.md
Original file line number Diff line number Diff line change
Expand Up @@ -9,7 +9,8 @@
| Fatbin / CMake / ROCm pin | **this repo** |
| Observed extras LDS / launch numbers (no body dump) | **this repo** (tile READMEs) |
| extras HIP body migrate | **this repo**, only after extras can consume one V1 op |
| Foreign HIP (leapdragon / kletorch / a17t / extras agent) | cherry-pick `-x`, **their** Author **and** Committer — never a Cursor rewrite |
| Dest extras HIP (this lab) | **BlivionIaG** `<kev29lt@gmail.com>` — Author **and** Committer |
| Foreign HIP (leapdragon / a17t) | cherry-pick `-x`, **their** Author **and** Committer — never a Cursor rewrite |
| vLLM registration, flags, capture, model hooks | `opengfx1030/vllm-rdna` `rdna_extras` (not from this tree) |
| AWQ / `3inst` produce | packer tools — **not** a hippihx directory |
| Upstream vLLM | **do not open PRs** |
Expand All @@ -23,35 +24,44 @@ layouts + capture probes. ISA migrate is extract-then-rewire in extras
(see [`docs/BACKPORT.md`](docs/BACKPORT.md)). Until then, extras
observations belong in the tile README locks — not a second kernel body.

## Attribution (foreign HIP)
## Attribution

If the HIP / commit is **not** new hippihx work, **do not rewrite it**.
Keep only the original authors' commits. Same rule as extras
**Dest extras HIP is BlivionIaG** `<kev29lt@gmail.com>`. Author **and**
Committer on picked dest commits. Do not leave Cursor / `cursoragent` on
those commits.

**Foreign HIP** (leapdragon, a17t, anyone who is not this lab) keeps
**their** Author **and** Committer. Same rule as extras
[PR #1](https://github.com/opengfx1030/vllm-rdna/pull/1) (leapdragon
`rdna_ar`).
`rdna_ar`). Do not rewrite those as Blivion or Cursor.

| Rule | Meaning |
|---|---|
| Cherry-pick `-x` | One source commit → one hippihx commit. Not a squash. Not a re-authored file dump. |
| Author **and** Committer | Copy both from the source commit (`GIT_COMMITTER_NAME` / `EMAIL` = source Author). Do not leave Committer as Cursor / Blivion / `cursoragent`. |
| Trailers | Keep the source `Co-authored-by` / `Signed-off-by`. **Do not add** Cursor, Cursor Agent, or BlivionIaG trailers to *their* commits. |
| Dest extras | Author **and** Committer = BlivionIaG `<kev29lt@gmail.com>` |
| Foreign | Copy both from the source commit (`GIT_COMMITTER_NAME` / `EMAIL` = source Author). Do not leave Committer as Cursor / Blivion / `cursoragent`. |
| Trailers | Foreign: keep their `Co-authored-by` / `Signed-off-by`. **Do not add** Cursor or Cursor Agent trailers to *their* commits. Dest extras: BlivionIaG only. |
| Mixed commits | Skip. Do not take a commit that mixes their kernel with serve wiring, PLE, tok/s docs, or a second op family. |
| Wiring after pick | Conflict resolution and plan/bind wrappers are **new** hippihx commits (agent/Blivion). They must not rewrite kernel bodies. |
| Verify | `git log --format='%an <%ae> | %cn <%ce> | %s'` on unique imported commits matches the source. |
| Wiring after pick | Conflict resolution and plan/bind wrappers are **new** hippihx commits. They must not rewrite kernel bodies. |
| Verify dest | unique dest-HIP commits show `BlivionIaG <kev29lt@gmail.com>` for Author and Committer |
| Verify foreign | `git log --format='%an <%ae> | %cn <%ce> | %s'` matches the source |

```bash
# Example: leapdragon commit onto a hippihx side branch (not dest).
# Author is preserved by cherry-pick; Committer must be forced or it
# becomes whoever ran the pick.
src=<sha>
name=$(git -C "$EXTRAS" log -1 --format='%an' "$src")
email=$(git -C "$EXTRAS" log -1 --format='%ae' "$src")
# Dest extras HIP → BlivionIaG (even if the source Author field is wrong).
GIT_AUTHOR_NAME='BlivionIaG' GIT_AUTHOR_EMAIL='kev29lt@gmail.com' \
GIT_COMMITTER_NAME='BlivionIaG' GIT_COMMITTER_EMAIL='kev29lt@gmail.com' \
git cherry-pick -x "$src"

# Foreign (leapdragon / a17t): Author is preserved by cherry-pick;
# Committer must be forced or it becomes whoever ran the pick.
name=$(git -C "$FOREIGN" log -1 --format='%an' "$src")
email=$(git -C "$FOREIGN" log -1 --format='%ae' "$src")
GIT_COMMITTER_NAME="$name" GIT_COMMITTER_EMAIL="$email" \
git cherry-pick -x "$src"
```

Do not `git commit --amend --reset-author`. That replaces Author with
the picker.
Do not `git commit --amend --reset-author` on **foreign** picks. That
replaces Author with the picker.

## ROCm / arch

Expand Down
118 changes: 118 additions & 0 deletions README.md
Original file line number Diff line number Diff line change
Expand Up @@ -189,3 +189,121 @@ consume layout and ships HIP that reads it.
- [`docs/BACKPORT.md`](docs/BACKPORT.md) — extras review: what is zoo vs
serve, why bodies are not copied yet.
- [`CONTRIBUTING.md`](CONTRIBUTING.md) — ROCm pin, how to land a tile.
Kernel / mode / env / AR tracker: **unvalidated** list below.

## Unvalidated extras inventory

Snapshot of [`opengfx1030/vllm-rdna`](https://github.com/opengfx1030/vllm-rdna)
`rdna_extras` @ `d71721c79547` (2026-09-08) plus open extras PRs **#1–#3**.

**Unvalidated.** Not dest. Not silicon-signed. No tok/s. hippihx still
ships stubs; bodies stay in extras until a consume bind exists. This
list is a tracker, not a claim that any row works.

Status key: **extras** = live on dest tip (unvalidated here) ·
**Later** = side branch / review-only · **skip** = do not take ·
**stub** = hippihx contract only.

### Kernels (HIP)

| Kernel / op | extras | hippihx tile | Notes |
|---|---|---|---|
| FA paged decode / prefill / split-K / short | `fa_rdna2.cu` | `attn/fa_fdot2` | `fdot2`. Occupancy pin closed. **unvalidated** |
| FA INT8 KV writer | `reshape_and_cache_int8_rdna2` | `attn/fa_fdot2` | INT8 cache layout for RDNA_ATTN. **unvalidated** |
| W4A16 dense decode | `q_gemm_rdna2.cu` | `gemm/w4a16_fdot2` | GPTQ + AWQ = pack/zeros, one GEMM. **unvalidated** |
| W4A16 prefill | `q_gemm_rdna2_prefill.cu` | `gemm/w4a16_fdot2` | Multi-config. **unvalidated** |
| W4A16 AWQ high-M prefill | `q_gemm_rdna2_awq_prefill.cu` | `gemm/w4a16_fdot2` | Exllama-clone tile, AWQ zeros (no GPTQ +1). **unvalidated** |
| W4A16 MoE | `moe_q_gemm_rdna2.cu` | `moe/routed` | **unvalidated** |
| EXL3 dense / MoE / dequant / Hadamard / trellis decode | `exl3_dot2_*.cu` | `gemm/exl3_3inst` | Consume `-cb 3inst`. Produce outside. UNC-26. **unvalidated** |
| GDN packed decode | `gdn_decode_rdna2.cu` | `attn/gdn_scan` | Register-resident. Capture still garbage on dest hybrid. **unvalidated** |
| GDN prefill chain (prep / kkt / solve_wy / delta_h / o) | `gdn_prefill_*_rdna2.cu` | `attn/gdn_scan` | Default-off (`VLLM_GDN_HIP_PREFILL=0`). Chunk-boundary corruption. **unvalidated** |
| causal_conv1d update + fwd | `causal_conv1d_rdna2.cu` | `sequence/causal_conv` | Scalar FMA, `state_len≈3–4`. **unvalidated** |
| Paged MQA indexer | `indexer_paged_mqa_rdna2.cu` | `attn/qsa_indexer` | DeepSeek V4 Lightning class. **unvalidated** |
| Sparse MLA decode / prefill | `sparse_mla_rdna2.cu` | `attn/dsa_nope` | **unvalidated** |
| W8A16 / W8A16-FP8 / W8A8-FP8 dense+MoE | `moe_w8a16*.cu`, `gemm_w8a8_fp8_dense_rdna2.cu`, `q_gemm_w8a16_fp8_rdna2.cu` | — | No hippihx tile yet. **unvalidated** |
| MXFP4 dense + MoE | `mxfp4_dot2_*.cu` | — | No hippihx tile yet. **unvalidated** |
| gfx1100 W4 WMMA | `q_gemm_rdna3_wmma.cu` | — | WMMA Later overlay, not shared DOT. **unvalidated** |
| Skinny GEMM / INT4 skinny | `skinny_gemms*.cu` | — | `VLLM_ROCM_USE_SKINNY_GEMM`. **unvalidated** |
| RMSNorm HIP AOT | `layernorm.cu` | — | Serve fused-norm; no tile. **unvalidated** |
| GLM-5.3 KDA decode + prefill | `glm5_kda_*.cu` (PR **#2**) | `attn/kda_scan` | Later. Drop `glm5_` name. **unvalidated** |
| GLM-5.3 DSA indexer + MLA-NoPE | `glm5_dsa_*.cu` (PR **#2**) | `attn/dsa_nope`, `qsa_indexer` | Later. **unvalidated** |
| leapdragon push AR | `rdna_ar` (PR **#1**) | `comm/pcie` | Uncached+push. Default **off**. Aron Hsiao. **unvalidated** |
| a17t extra AWQ GEMM / GEMV | `awq_gemm_rdna2.cu`, `moe_awq_gemm_rdna2.cu`, `gemv_w4_kpack_rdna2.cu` (PR **#3**) | — | **skip** — second W4 family |

### Modes / dispatch

| Mode | What extras does | Default (extras) | hippihx |
|---|---|---|---|
| W4 decode vs prefill vs AWQ-prefill vs Exllama | M/K/N buckets in `rdna2_w4a16.py` | decode `M≤32` & `K≥4096`; AWQ high-M → `awq_prefill`; `M>256` Exllama (GPTQ) | one `w4a16_fdot2` tile |
| W4 pack | GPTQ `uint4b8` (+1 zeros) vs AWQ `uint4` (literal zeros) | same kernel, `use_v2_format` | pack/zeros, not a second GEMM |
| EXL3 codebook | `cb==0` 3inst produce, `cb==1` mcg compile, `cb==2` mul1 not produced | 3inst dest | consume only |
| EXL3 memory | `full` int16 trellis vs `packed` stub | `full` | — |
| EXL3 prefill | decode-trellis prefill vs fused GEMM | `VLLM_EXL3_PREFILL_DECODE=1` | — |
| GDN decode HIP | packed HIP vs Triton/FLA | on unless `VLLM_GDN_DECODE_RDNA2=0` | `gdn_scan` |
| GDN prefill HIP | 5-kernel chain | **off** | `gdn_scan` |
| FA backend | `RDNA_ATTN` when gfx10x | `VLLM_USE_RDNA2_FA` (envs.py default false; backend reads `"1"`) | `fa_fdot2` |
| FA spec/MTP gate | opt-in abort on verify-shaped batches | **off** | serve, not a tile |
| MLA sparse HIP | indexer + sparse MLA | `VLLM_USE_RDNA2_MLA=1` | indexer / `dsa_nope` |
| causal conv HIP | update (decode) + fwd (prefill) | on unless set `0` | `causal_conv` |
| Custom AR (dest) | force custom all-reduce on PCIe-only | **off** | serve |
| leapdragon `rdna_ar` | size-gated Uncached+push | **off** (`VLLM_RDNA_AR=0`) | `comm/pcie` Later |

### Env (extras-added / extras-used)

Serve knobs. hippihx does not read these. **Unvalidated.** Debug probes stay extras (no D2H under capture in the zoo).

| Env | Default (as read in extras) | Role |
|---|---|---|
| `VLLM_USE_RDNA2_FA` | envs.py `False`; `rdna_attn` treats missing as `"1"` | FA-RDNA2 / `RDNA_ATTN` |
| `VLLM_USE_RDNA2_MLA` | off unless `"1"` | sparse MLA + paged MQA HIP |
| `VLLM_FARDNA2_ENABLE_SPEC_GATE` | `"0"` | MTP-verify abort (opt-in) |
| `VLLM_FARDNA2_SPEC_VERIFY_Q_LEN` | `"3"` | spec-gate q len |
| `VLLM_GDN_DECODE_RDNA2` | on (`!= "0"`) | GDN decode HIP |
| `VLLM_GDN_HIP_PREFILL` | off if `"0"` (dest default-off) | GDN prefill HIP chain |
| `VLLM_GDN_HIP_KERNELS` | recipe `1` | recipe umbrella; confirm vs the two gates above |
| `VLLM_GDN_DECODE_KERNEL` | `"cuda"` | FLA packed-decode path name (`cuda`/`triton`) |
| `VLLM_ENABLE_FLA_PACKED_RECURRENT_DECODE` | `1` | FLA packed decode |
| `VLLM_CAUSAL_CONV1D_RDNA2_UPDATE` | `"1"` | conv update HIP |
| `VLLM_CAUSAL_CONV1D_RDNA2_FWD` | `"1"` | conv fwd HIP |
| `VLLM_RDNA_FORCE_FP16` | recipe `1` | force fp16 (no BF16 emu) |
| `VLLM_EXL3_MEMORY_MODE` | `full` | `full` / `packed` |
| `VLLM_EXL3_PREFILL_DECODE` | `"1"` | trellis-decode prefill |
| `VLLM_EXL3_M_MAX` | `64` (`0` = no CG-PATH) | EXL3 capture cap |
| `VLLM_EXL3_DEQUANT_ALL` | off | dequant leftover / mul1 |
| `VLLM_EXL3_FOLDED_CACHE` | unset | folded-weight cache dir |
| `VLLM_ROCM_USE_SKINNY_GEMM` | `True` | skinny GEMM |
| `VLLM_ROCM_USE_AITER` | `False` | AITER (CDNA; not dest gfx1030) |
| `VLLM_ROCM_USE_AITER_CUSTOM_AR` | `True` | AITER AR (CDNA) |
| `VLLM_FORCE_CUSTOM_ALL_REDUCE` | `False` | force custom AR without full P2P |
| `VLLM_CUSTOM_ALLREDUCE_ALGO` | unset | `1stage` / `2stage` |
| `VLLM_ROCM_QUICK_REDUCE_*` | unset | ROCm quick-reduce size/quant knobs |
| `VLLM_ALLREDUCE_USE_SYMM_MEM` | `1` | symmetric-memory AR |
| `VLLM_RDNA_AR` | `"0"` (PR **#1**, not dest) | leapdragon push AR |
| `VLLM_RDNA_AR_BLOCKS` | auto (PR **#1**) | AR block cap |
| `VLLM_RDNA_AR_PACE` | `0` (PR **#1**) | AR store pace |
| `VLLM_USE_BREAKABLE_CUDAGRAPH` | `0` (auto-on in some configs) | capture dispatcher |
| `VLLM_LOG_GDN_PTRS` | off | GDN pointer probe |
| `VLLM_GDN_DBG` | off | GDN debug print |
| `VLLM_EXL3_DEBUG` / `_HADAMARD_DBG` / `_APPLY_DBG` / `_MARKER_DBG` / `_INPUT_NAN_DBG` | off | EXL3 probes |
| `VLLM_CONV1D_DEBUG` / `VLLM_MLP_DBG` / `VLLM_RDNA2_MOE_DEBUG_NAN` | off | probes |
| `DBG_VLLM_STEP_TIMING` | off | per-step timing |

### AR / collectives

| Path | Where | Default | hippihx |
|---|---|---|---|
| RCCL | extras fallback | on when custom AR off | — |
| Custom all-reduce (vLLM/ROCm) | `VLLM_FORCE_CUSTOM_ALL_REDUCE` | **off** | serve |
| AITER custom AR | `VLLM_ROCM_USE_AITER_CUSTOM_AR` | on in envs, AITER itself off | CDNA, not gfx1030 dest |
| Quick-reduce | `VLLM_ROCM_QUICK_REDUCE_*` | unset | serve |
| Symm-mem AR | `VLLM_ALLREDUCE_USE_SYMM_MEM` | on | serve |
| leapdragon `rdna_ar` Uncached+push | extras PR **#1** | **off** | `comm/pcie` Later. Author Aron Hsiao. Occupancy pin closed. Boot self-test. INT8/Q8 wire preferred; no Finegrained; no E4M3 without FP8 HW |

### Not taken / leave in extras

| Item | Why |
|---|---|
| a17t PR **#3** second AWQ GEMM + `qwen4_exp` | duplicate W4 family + serve |
| leapdragon `gemv_f16` / `moe_skinny_int4_decode` | dropped at extras PR #1 conflict resolution |
| Produce / pack (`-cb 3inst`, AWQ produce) | outside hippihx |
| Cudagraph `torch.zeros`, `eager_break_during_capture`, GDN probes | serve page-commit / dispatcher |
6 changes: 4 additions & 2 deletions docs/ARCHITECTURE.md
Original file line number Diff line number Diff line change
Expand Up @@ -159,5 +159,7 @@ host-side and torch-free.
- Editing `opengfx1030/vllm-rdna` from this tree
- Dumping extras `csrc/rocm/*.cu` here before extras can consume hippihx
(see [`BACKPORT.md`](BACKPORT.md))
- Re-authoring foreign HIP as Cursor / Blivion commits (cherry-pick `-x`,
keep source Author **and** Committer)
- Re-authoring dest extras HIP as Cursor commits (dest Author **and**
Committer is BlivionIaG). Re-authoring foreign HIP (leapdragon / a17t)
as Cursor **or** Blivion (cherry-pick `-x`, keep **their** Author and
Committer)
34 changes: 18 additions & 16 deletions docs/BACKPORT.md
Original file line number Diff line number Diff line change
Expand Up @@ -92,32 +92,34 @@ All of:
4. extras is rewired to call it (that edit happens **in extras**, not from
this tree). The extras copy is then deleted.

Until then: observe, lock numbers, keep stubs.
Until then: observe, lock numbers, keep stubs. Kernel / mode / env / AR
tracker (all **unvalidated**): [`README.md`](../README.md#unvalidated-extras-inventory).

## Attribution (do not re-author)

Imported HIP keeps the **source Author and Committer**. If it is not
new hippihx work, do not rewrite it as a Cursor / Blivion commit.
Cherry-pick `-x`, force Committer = source Author, keep their trailers,
add none of ours. Recipe: [`CONTRIBUTING.md`](../CONTRIBUTING.md)
(Attribution). Same bar as extras PR #1.
**Dest extras HIP** (FA, EXL3, W4A16, GDN, causal_conv, sparse MLA,
indexer, GLM Later KDA/DSA) is **BlivionIaG** `<kev29lt@gmail.com>`.
Author **and** Committer on every picked dest commit. Do not use Cursor
or `cursoragent`.

**Foreign HIP** (leapdragon, a17t) keeps **their** Author **and**
Committer. Cherry-pick `-x`, force Committer = source Author, keep their
trailers, add none of ours. Recipe: [`CONTRIBUTING.md`](../CONTRIBUTING.md).
Same bar as extras PR #1.

This hippihx review commit is **new documentation**. It is not a
kernel migrate and must not be used as a template for body imports.

| extras / PR path | Introduced (Author) | Keep as |
| extras / PR path | Introduced | Keep as |
|---|---|---|
| `fa_rdna2.cu` | `b1b3fa938` BlivionIaG `<kev29lt@gmail.com>` | BlivionIaG commits; later `agent@opencode.local` FA zeros stay extras (serve) |
| `fa_rdna2.cu` | `b1b3fa938` BlivionIaG `<kev29lt@gmail.com>` | BlivionIaG. Later FA `torch::zeros` is extras serve — leave it there |
| `q_gemm_rdna2.cu` | `fabf51493` BlivionIaG | BlivionIaG |
| `q_gemm_rdna2_awq_prefill.cu` | `feb7b457e` BlivionIaG | BlivionIaG; later `agent@opencode.local` bundling stays extras |
| `q_gemm_rdna2_awq_prefill.cu` | `feb7b457e` BlivionIaG | BlivionIaG |
| `moe_q_gemm_rdna2.cu` | `b1b3fa938` BlivionIaG | BlivionIaG |
| `gdn_decode_rdna2.cu` | `55527010c` BlivionIaG | BlivionIaG; `3326f03d7` `agent@opencode.local` NULL_BLOCK_ID stays extras until dest-locked |
| `exl3_dot2_*.cu` | `40850e6c5` **kletorch** `<kletorch@users.noreply.github.com>` (Committer was BlivionIaG) | **kletorch as Author and Committer** on the port commits. Do not collapse to Blivion/Cursor |
| `causal_conv1d_rdna2.cu` | `5eb84b4fa` `agent@opencode.local` | Keep `agent@opencode.local` if that commit is picked. Do not re-author as Cursor Agent |
| `gdn_decode_rdna2.cu` | `55527010c` BlivionIaG | BlivionIaG. NULL_BLOCK_ID sentinel is extras serve until dest-locked |
| `exl3_dot2_*.cu` | `40850e6c5` BlivionIaG | BlivionIaG (Author **and** Committer), including the port and later perf commits |
| `causal_conv1d_rdna2.cu` | `5eb84b4fa` BlivionIaG | BlivionIaG |
| `sparse_mla_rdna2.cu` / `indexer_paged_mqa_rdna2.cu` | BlivionIaG | BlivionIaG |
| PR #1 `rdna_ar` | **Aron Hsiao** `<leapdragon@gmail.com>` | Aron Hsiao Author **and** Committer. Keep `Co-Authored-By: Claude Fable 5`. No Blivion/Cursor trailers |
| PR #2 `glm5_kda_*` / `glm5_dsa_*` | BlivionIaG | BlivionIaG; rename off `glm5_` in a **follow-up hippihx** commit, not by rewriting their HIP |
| PR #3 a17t unique W4 (`d53572644` `root@tr3960x-1.home.a17t.dev`, later Simon Siebert) | **Not taken** | If dest ever locks that family, pick **their** commits, not a rewrite |

`agent@opencode.local` on extras is a different agent identity than
Cursor Agent. Do not launder those commits into `cursoragent@cursor.com`.
| PR #3 a17t unique W4 (`d53572644`, later Simon Siebert) | **Not taken** | If dest ever locks that family, pick **their** commits, not a rewrite |
29 changes: 25 additions & 4 deletions tests/test_no_b12x_cuda.py
Original file line number Diff line number Diff line change
Expand Up @@ -44,9 +44,30 @@ def test_foreign_hip_attribution_policy() -> None:
text = (ROOT / "CONTRIBUTING.md").read_text(encoding="utf-8")
assert "cherry-pick -x" in text
assert "GIT_COMMITTER_NAME" in text
assert "--reset-author" in text
assert "kev29lt@gmail.com" in text
bp = (ROOT / "docs" / "BACKPORT.md").read_text(encoding="utf-8")
assert "kletorch" in bp
assert "BlivionIaG" in bp
assert "leapdragon@gmail.com" in bp
assert "agent@opencode.local" in bp
assert "cursoragent@cursor.com" in bp


def test_readme_unvalidated_inventory() -> None:
text = (ROOT / "README.md").read_text(encoding="utf-8")
assert "Unvalidated extras inventory" in text
assert "**Unvalidated.**" in text
assert "VLLM_RDNA_AR" in text
assert "fa_rdna2" in text


def test_no_mistaken_dest_author_identities() -> None:
forbidden = ("kletorch", "opencode.local")
hits: list[str] = []
scan = [ROOT / "CONTRIBUTING.md", ROOT / "README.md"]
scan.extend((ROOT / "docs").rglob("*.md"))
scan.extend((ROOT / "tiles").rglob("*.md"))
for path in scan:
if not path.is_file():
continue
text = path.read_text(encoding="utf-8", errors="replace")
if any(token in text for token in forbidden):
hits.append(str(path.relative_to(ROOT)))
assert hits == []
9 changes: 4 additions & 5 deletions tiles/gemm/exl3_3inst/README.md
Original file line number Diff line number Diff line change
Expand Up @@ -8,7 +8,7 @@ Decode is `3inst` → `half2` → `fdot2`.

| Lock | Status |
|---|---|
| LDS bytes | extras `exl3_dot2_dense.cu` (add `40850e6c5` **kletorch**): A-staging `[M_PER][BLOCK_K + LDS_PAD]`, `BLOCK_K=256`, **`LDS_PAD=8`**. Small vs FA; occupancy is VGPR |
| LDS bytes | extras `exl3_dot2_dense.cu` (add `40850e6c5` **BlivionIaG**): A-staging `[M_PER][BLOCK_K + LDS_PAD]`, `BLOCK_K=256`, **`LDS_PAD=8`**. Small vs FA; occupancy is VGPR |
| `__launch_bounds__` | **unset on DOT kernels** (UNC-26). Hadamard wrappers use `__launch_bounds__(32)` — not the K-dot |
| Wave | **32 only** |
| Geometry | extras: `THREADS_X=256`, `BLOCK_N=1024`, `BLOCK_K=256`, 4 N-cols/thread, 16×16 window → `decode_3inst` → `half2` → `fdot2` |
Expand All @@ -17,7 +17,6 @@ Decode is `3inst` → `half2` → `fdot2`.

`mcg` / `mul1` codebooks are compile slots, not extra zoo roots. UNC-26 is
still In Progress on extras — do not copy bodies until produce/consume
and capture buffers stop moving. Port commits stay **kletorch**
`<kletorch@users.noreply.github.com>` as Author **and** Committer; later
BlivionIaG perf commits stay BlivionIaG. Do not squash them into a
Cursor rewrite.
and capture buffers stop moving. Port and later perf commits stay
**BlivionIaG** `<kev29lt@gmail.com>` as Author **and** Committer. Do not
squash them into a Cursor rewrite.
Loading