diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index e03dc93..f95751c 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -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** `` — 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** | @@ -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** ``. 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 `` | +| 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 ` 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= -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 diff --git a/README.md b/README.md index b4dc960..a5325fe 100644 --- a/README.md +++ b/README.md @@ -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 | diff --git a/docs/ARCHITECTURE.md b/docs/ARCHITECTURE.md index 5168777..51d59be 100644 --- a/docs/ARCHITECTURE.md +++ b/docs/ARCHITECTURE.md @@ -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) diff --git a/docs/BACKPORT.md b/docs/BACKPORT.md index 467e4e8..aef79fa 100644 --- a/docs/BACKPORT.md +++ b/docs/BACKPORT.md @@ -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** ``. +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 `` | BlivionIaG commits; later `agent@opencode.local` FA zeros stay extras (serve) | +| `fa_rdna2.cu` | `b1b3fa938` BlivionIaG `` | 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** `` (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** `` | 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 | diff --git a/tests/test_no_b12x_cuda.py b/tests/test_no_b12x_cuda.py index 900807c..4ce9589 100644 --- a/tests/test_no_b12x_cuda.py +++ b/tests/test_no_b12x_cuda.py @@ -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 == [] diff --git a/tiles/gemm/exl3_3inst/README.md b/tiles/gemm/exl3_3inst/README.md index 6dacd63..9fa71a2 100644 --- a/tiles/gemm/exl3_3inst/README.md +++ b/tiles/gemm/exl3_3inst/README.md @@ -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` | @@ -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** -`` 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** `` as Author **and** Committer. Do not +squash them into a Cursor rewrite. diff --git a/tiles/sequence/causal_conv/README.md b/tiles/sequence/causal_conv/README.md index aefdee7..434b389 100644 --- a/tiles/sequence/causal_conv/README.md +++ b/tiles/sequence/causal_conv/README.md @@ -8,7 +8,7 @@ do not retarget one onto the other. This stub is the conv contract only. | Lock | Status | |---|---| -| LDS bytes | extras `causal_conv1d_rdna2.cu` add `5eb84b4fa` **agent@opencode.local**: **0** — register-only, no global scratch. Keep that Author/Committer if picked; do not re-author as Cursor Agent | +| LDS bytes | extras `causal_conv1d_rdna2.cu` add `5eb84b4fa` **BlivionIaG**: **0** — register-only, no global scratch | | `__launch_bounds__` | unset in extras (one warp per dim-block) | | Wave | 32 on gfx1030 / gfx1100; extras: 32 threads / dim-block, dim multiple of 32 | | Math | scalar FMA, fp16 in/out, fp32 acc — not `fdot2`, not `fdot2.bf16`, not WMMA |