From ff23e38026c1f4f8e7a2471423e7358098ca3f06 Mon Sep 17 00:00:00 2001 From: Greg Stoner Date: Tue, 28 Jul 2026 23:05:52 -0600 Subject: [PATCH 1/3] Add gfx1250/MI450 target reference and AMD kernel-compiler survey MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Two primary-source reference documents from a source-level read of AMD's ROCm repositories, plus the Apple follow-on they imply. Documentation only; no code changes. backend/rocm/GFX1250_MI450_COMPILER_REFERENCE.md — target facts for gfx1250 (MI450), derived from LLVM tablegen, ROCm library source, and the on-machine assembler rather than from vendor prose. Every claim carries a provenance marker and section 10 gives commands to re-derive it without AMD hardware. Findings that correct or extend what the tree currently assumes: - Our cluster_mode flag is inverted. FeatureClusters is on gfx1250/1251 and not on gfx950; rocm_target.py asserts the opposite. Latent today — supports_cluster_mode() has no codegen consumer — but wrong. - gfx1250 LDS is 327680 bytes and it has 1024 addressable VGPRs; our entries carry 65536 and 256, marked PROVISIONAL. The VGPR figure matters because rocm_tiling.py treats the register budget as the dominant lever. - gfx1250 and gfx1251 are complementary SKUs, not revisions: the same WMMA instruction runs 4x apart, and FP64 runs 6x apart the other way. _GFX1250_CLASS_ARCHES is right for ABI and wrong for any cost model. - Async and tensor waits are never inserted automatically by LLVM. They must come from llvm.amdgcn.{asyncmark,wait.asyncmark}, and a too-large index yields no wait at all rather than a conservative one. Our ROCM_WaitTokenOp has no immediate operand, so it can only express a full drain. - AMD clusters are multicast-into-own-LDS plus a barrier, scoped to a shader engine. There is no distributed shared address space, so an NVIDIA CGA kernel that reads a peer CTA's shared memory will not lower. compiler/AMD_KERNEL_COMPILER_SURVEY.md — StinkyTofu, rocRoller, Composable Kernel and hipBLASLt read for transferable architecture. Section 6 ranks 21 items take/skip. The two most useful are structural: rocRoller's observer scheduling gives a cost query before commitment, and CK derives vector width, access count and traversal order from a distribution encoding rather than having a kernel author write them. It also records two static, device-free quality metrics found in production AMD code — a step-distance locality histogram and a bank-conflict analyzer that computes N-way conflict from a descriptor alone. Both bear directly on the mock-cost-model finding in TILESIGHT_ASSESSMENT.md. APPLE_AUDIT.md gains backlog item 8 as the follow-on: Apple is the only backend that executes broadly enough to say whether such a metric predicts anything, so the action is to calibrate one against recorded latency rather than to add one. A metric that cannot rank Apple kernels should not be trusted to rank kernels we cannot measure. Co-Authored-By: Claude Opus 5 --- docs/audit/backend/apple/APPLE_AUDIT.md | 25 +- .../rocm/GFX1250_MI450_COMPILER_REFERENCE.md | 1151 ++++++++++++++++ docs/audit/backend/rocm/ROCM_AUDIT.md | 3 + .../rocm/ROCM_PATTERNS_FROM_AMD_ECOSYSTEM.md | 12 + .../compiler/AMD_KERNEL_COMPILER_SURVEY.md | 1178 +++++++++++++++++ 5 files changed, 2368 insertions(+), 1 deletion(-) create mode 100644 docs/audit/backend/rocm/GFX1250_MI450_COMPILER_REFERENCE.md create mode 100644 docs/audit/compiler/AMD_KERNEL_COMPILER_SURVEY.md diff --git a/docs/audit/backend/apple/APPLE_AUDIT.md b/docs/audit/backend/apple/APPLE_AUDIT.md index 93fabe183..d193894de 100644 --- a/docs/audit/backend/apple/APPLE_AUDIT.md +++ b/docs/audit/backend/apple/APPLE_AUDIT.md @@ -1,5 +1,5 @@ --- -last_updated: 2026-07-13 +last_updated: 2026-07-28 audit_role: sub_audit --- @@ -291,6 +291,29 @@ not current support counts. hardware-free bridge remains `python/tessera/compiler/microscaling.py`. 7. **Cross-backend real-hardware proof** for NVIDIA / ROCm remains tracked in the per-platform audits, not here. +8. **Be the validation site for the hardware-free cost model.** Added 2026-07-28 + from the AMD kernel-compiler survey + ([`../../compiler/AMD_KERNEL_COMPILER_SURVEY.md`](../../compiler/AMD_KERNEL_COMPILER_SURVEY.md)), + which found two *static, device-free* quality metrics in production AMD code: + a step-distance locality histogram over a materialized access order (§3.7) and + a bank-conflict analyzer that computes N-way conflict from a descriptor alone + (§3.8). Both are computable on any target and neither needs silicon. + + Apple is the only backend that executes broadly enough to say whether such a + metric *predicts anything*. The action is therefore not "add a metric" but + **calibrate one**: compute the locality/conflict score for kernel families the + Apple lane already measures, and check the score against recorded latency. A + metric that does not rank Apple kernels correctly should not be trusted to + rank ROCm or NVIDIA kernels we cannot measure. This is the concrete follow-on + to the mock-cost-model finding in + [`../../compiler/TILESIGHT_ASSESSMENT.md`](../../compiler/TILESIGHT_ASSESSMENT.md) + §2, and it gates how much weight the arbiter's hardware-free tier can carry. + + A second, smaller item from the same survey: the MSL synthesizer currently + *authors* access patterns, whereas CK *derives* vector width, access count and + traversal order from a distribution encoding (§3.9). That is a design question + for the synthesizer, not a task — record it when the codegen path is next + revisited, and do not treat it as blocking. ## Hardware capability reference (grounded 2026-06-17) diff --git a/docs/audit/backend/rocm/GFX1250_MI450_COMPILER_REFERENCE.md b/docs/audit/backend/rocm/GFX1250_MI450_COMPILER_REFERENCE.md new file mode 100644 index 000000000..08dee28ab --- /dev/null +++ b/docs/audit/backend/rocm/GFX1250_MI450_COMPILER_REFERENCE.md @@ -0,0 +1,1151 @@ +--- +last_updated: 2026-07-28 +audit_role: reference +--- + +# gfx1250 / MI450 — Source-Grounded Compiler Reference + +> **Purpose.** A primary-source reference for the gfx1250 (AMD Instinct MI450 +> series) target, assembled for compiler and backend work: hardware constants, +> the WMMA pipeline and its hazards, the three data-movement mechanisms, the +> split completion model, workgroup clusters, device-initiated SDMA, and the +> scale-up fabric. +> +> **This is not a status surface.** Nothing here claims Tessera support for +> anything. For ROCm status see [`ROCM_AUDIT.md`](ROCM_AUDIT.md); for counts see +> `docs/audit/generated/`. For general AMD-ecosystem design patterns see +> [`ROCM_PATTERNS_FROM_AMD_ECOSYSTEM.md`](ROCM_PATTERNS_FROM_AMD_ECOSYSTEM.md), +> which this document extends with gfx1250-specific depth. +> +> **Provenance per claim:** **[V]** verified from primary source (LLVM +> tablegen/source, ROCm repo source, or the on-machine assembler); **[A]** +> verified by assembling the instruction locally; **[S]** stated by an AMD +> engineering document in-repo (first-party but prose); **[I]** inference or +> recommendation — explicitly not established. +> +> Researched 2026-07-28 against `ROCm/llvm-project@amd-staging`, +> `ROCm/rocm-systems@develop`, `ROCm/rocm-libraries@develop`, and both local +> LLVMs (Homebrew `llvm/22.1.8` and LLVM 23.1.0-rc1) — see the toolchain note in §10. + +--- + +## 0. Evidence ladder + +When these sources disagree — and they do — resolve in this order: + +1. **The assembler** (`llvm-mc -arch=amdgcn -mcpu=gfx1250`). It either encodes or + it does not. Cheapest and most decisive. +2. **LLVM tablegen** (`AMDGPU.td`, `IntrinsicsAMDGPU.td`). Machine-consumed, so + errors surface as miscompiles rather than stale prose. +3. **ROCm library source** (`rocshmem`, `rccl`, `stinkytofu`, `rocke`). + Production code, but the comments can lag the code. +4. **AMD engineering prose** (in-repo design docs, blogs). Useful for intent and + measurements; demonstrably self-contradictory in places — see §8. + +A worked example of why the order matters: `rocke`'s MHA case study §6 states +gfx1250 has "no async global→LDS DMA" and "no `ds_read_tr`", while §5 of the +*same file* describes building a `ds_load_tr16_b128` path, and the capability +table in the sibling plan asserts both exist. The assembler settles it — both +exist. **[V]** + +--- + +## 1. Target identity and hardware constants + +gfx1250 is the AMD Instinct **MI450 series**. gfx1251 is a sibling ISA in the +same family — *not* an alias (see §2.4). Both derive from +`FeatureISAVersion12_50_Common` in `AMDGPU.td`. **[V]** + +Despite the `gfx12xx` numbering it is **not RDNA 4**. It is wave32 with WMMA and +no MFMA — a matrix pipeline lineage distinct from both RDNA 4 (gfx1200/1201, +16x16x16 WMMA) and CDNA (gfx942/gfx950, MFMA). **[V]** + +| Constant | Value | Source | +|---|---|---| +| Addressable LDS | **327680 B (320 KiB)** | `FeatureAddressableLocalMemorySize327680`, `AMDGPU.td:2292` **[V]** | +| Addressable VGPRs | **1024** | `Feature1024AddressableVGPRs`, `AMDGPU.td:2208` **[V]** | +| Wavefront | 32 | `FeatureWavefrontSize32` **[V]** | +| Virtual address bits | 57 | ROCKE capability table **[S]** | +| Data cache line | 128 B | `FeatureDataCacheLineSize128` **[V]** | +| Matrix path | WMMA; `has_mfma = false` | ROCKE **[S]**, corroborated by `AMDGPU.td:418` **[V]** | + +For scale: gfx950 is 163840 B LDS, RDNA 3/3.5/4 are 65536 B. gfx1250 doubles +CDNA 4 and is 5× RDNA. The VGPR count doubles CDNA's 512 combined +(256 VGPR + 256 AGPR); gfx1250 has no AGPR file, so 1024 is all architectural +VGPR. **[V]** + +AMD's own ROCKE capability table for the gfx1250 `ArchTarget` **[S]**: + +```text +wave_size = 32 matrix_path = wmma has_mfma = false +has_wmma = true has_swmma = true has_tdm = true +has_async_global_lds = true has_ds_load_tr = true max_lds_bytes = 320*1024 +wgp_cache_lds_shared = true virtual_address_bits = 57 +waitcnt_model = split_gfx1250 barrier_model = split_named_cluster +requires_shader_end_padding = true +``` + +`max_lds_bytes = 320*1024` = 327680 matches the LLVM feature exactly — an +independent confirmation from AMD's kernel team against AMD's compiler team. + +### 1.1 Notable subtarget features + +Present on gfx1250 **[V]**: + +`FeatureClusters`, `FeatureMcastLoadInsts`, `FeatureAsyncLoadToLDSInsts`, +`FeatureAsyncStoreFromLDSInsts`, `FeatureAsynccnt`, `FeatureWaitXcnt`, +`FeatureTransposeLoadF4F6Insts`, `FeatureSWMMACGfx1250Insts`, +`FeatureWMMACoexecutionHazards`, `FeatureTransCoexecutionHazard`, +`FeatureLdsBarrierArriveAtomic`, `FeatureSWakeupBarrier`, +`FeatureSetPrioIncWgInst`, `FeatureGloballyAddressableScratch`, +`FeatureKernargPreload`, `FeatureVmemPrefInsts`, `FeatureSmemPrefetchInsts`, +`FeatureXNACK`, `FeatureSupportsSRAMECC`, `FeatureRealTrue16Insts`, +`Feature45BitNumRecordsBufferResource`. + +Absent on gfx1250, present on gfx1251: `FeatureGFX125xLowestRateWMMA`, +`FeatureFullRate64Ops`, `FeaturePackedFP64Ops`, `FeatureGFX1251GEMMInsts`. **[V]** + +### 1.2 Cache hierarchy and the SCOPE ladder + +From `AMDGPUUsage.rst`, "Memory Model GFX125x". **[V]** + +``` +agent ── shader engines (SE) ── shader arrays (SA) ── work-group processors (WGP) + └── 4× SIMD32 (2 SIMD32-pairs) +``` + +- Each WGP has a **single write-through WGP$ shared between LDS and vector L0**. + Vector L0 holds clean data only. This is the hardware fact behind ROCKE's + `wgp_cache_lds_shared = true`, and behind the TDM guidance about bypassing vs + routing through the per-WGP cache. +- **Each WGP$ has two request queues, one per SIMD32-pair.** Each handles both + LDS and vector L0 requests; requests within a queue are serial and in-order, + but are *not* ordered against the other queue. (This is the "2 memory units per + WGP" figure that appears in AMD's marketing material.) +- Scalar memory uses a separate scalar L0, **not** kept coherent with vector L0 + by hardware — safe only because scalar ops are restricted to memory proven + not to change during the dispatch. +- All WGPs on an SE share an **L1 buffer**, with a separate request queue per + WGP$ (again in-order within a queue, unordered across queues). +- An agent may have **multiple L2 caches**. Virtual address ranges can be + configured non-hardware-coherent, read-write coherent with other L2s on the + same or other agents, or L2-bypassing for system coherence. + +The `SCOPE` field on vector memory operations names a cache level **[V]**: + +| `SCOPE` | Level | +|---|---| +| `SCOPE_CU` | WGP — the compiler's default, omitted in textual asm | +| `SCOPE_SE` | Shader Engine | +| `SCOPE_DEV` | Device / agent | +| `SCOPE_SYS` | System | + +An operation reaching a cache with a *smaller* scope is forwarded onward; at a +cache whose scope is ≥ its own it can complete locally (read hits, write +completes and reports, RMW done locally). Hardware assigns each cache a scope +per agent configuration, which is what lets `SCOPE_DEV` implement agent +coherence even with multiple non-coherent L2s. + +Also present: an `nv` ("non-volatile") bit marking memory not expected to change +during the kernel, propagated to cache lines as `$nv`; and `global_inv` / +`global_wb` / `global_wbinv` cache-control instructions whose affected levels are +selected by `SCOPE`, completing via `s_wait_storecnt`. **[V]** + +--- + +## 2. Matrix pipeline — WMMA + +### 2.1 Shapes + +The gfx1250 WMMA family is far wider than RDNA 4's single 16x16x16 **[V]**: + +| Class | Intrinsics | +|---|---| +| f16 / bf16, K=32 | `wmma_f32_16x16x32_{f16,bf16}`, `wmma_f16_16x16x32_f16`, `wmma_bf16_16x16x32_bf16`, `wmma_bf16f32_16x16x32_bf16` | +| fp8 / bf8, K=64 | `wmma_{f32,f16}_16x16x64_{fp8,bf8}_{fp8,bf8}` (4 combos each) | +| fp8 / bf8, K=128 | `wmma_{f32,f16}_16x16x128_{fp8,bf8}_{fp8,bf8}` | +| int, K=64 | `wmma_i32_16x16x64_iu8` | +| block-scaled MX | `wmma_f32_16x16x128_f8f6f4`, `wmma_scale_f32_16x16x128_f8f6f4`, `wmma_scale16_...` (i32 vs i64 scale) | +| f4 | `wmma_scale_f32_32x16x128_f4`, `wmma_scale16_f32_32x16x128_f4` | +| sparse | `SWMMAC` at 16x16x64 / 16x16x128 (`FeatureSWMMACGfx1250Insts`) | + +`wmma_f32_16x16x128_f8f6f4` uses `AMDGPUWmmaIntrinsicModsC_MatrixFMT` — a +**matrix-format operand**, so one instruction covers fp8/fp6/fp4 selected by a +format field rather than distinct opcodes. Any dtype contract modelling this +must treat element format as an *operand*, not part of the opcode identity. **[V]** + +### 2.2 Fragment ABI + +Per AMD's own probe **[S]**, corroborated by the intrinsic signatures **[V]**: + +- A/B operands: `<16 x half>` per lane — the 32 K-elements split across the two + lane-halves, 16 each. +- Accumulator: gfx12 column-distributed `<8 x float>`. +- bf16 uses native `<16 x bfloat>`, **not** `<_ x i16>`. +- The "v2" ModsC ABI carries an `i16` C-modifier immediate plus trailing + `i1, i1` reuse flags — RDNA 4 uses the plain 3-arg form. + +This matches what `python/tessera/compiler/rocdl_emit.py` already asserts. + +### 2.3 The WMMA co-execution hazard — mandatory + +`FeatureWMMACoexecutionHazards` and `FeatureTransCoexecutionHazard` are on +gfx1250. LLVM implements the fix in `GCNHazardRecognizer.cpp` +(`checkWMMACoexecutionHazards` / `fixWMMACoexecutionHazards`), including V_NOP +hoisting out of loops, citing hardware spec **SPG 4.6.12.1 "Requirements for +WMMA data hazards"**. **[V]** + +```c +const int WMMAWaitStates[] = {5, 9, 3, 5, 9, 17, 2}; // next op is a WMMA +const int VALUWaitStates[] = {4, 8, 2, 4, 8, 16, 1}; // next op is a VALU +``` + +Categories are derived from instruction latency **[V]**: + +| Cat | Latency | Instructions | +|---|---|---| +| 0 | 8 | WMMA f16/bf16; 16x16x128 fp8/bf8; f8f6f4 (not both-F4) | +| 1 | 16 | WMMA IU8 | +| 2 | 8 | SWMMAC f16/bf16/fp8/bf8 | +| 3 | 16 | SWMMAC IU8 | +| 4 | 16 | **gfx1251** 16-pass forms | +| 5 | 32 | **gfx1251** 32-pass forms | +| 6 | 4 | **gfx1250** 16x16x64 fp8/bf8; f8f6f4 (both-F4) | + +The hazard fires only on register overlap — WMMA0's `vdst` against WMMA1's +`src0`/`src1`/`Idx`, or against a co-executable VALU's operands. Independent +work already scheduled between them counts toward the requirement. + +**Consequences.** + +- Lowering through LLVM (Target IR → LLVM IR → AMDGPU) gets this handling for + free. **[V]** +- Hand-emitting assembly does not. AMD's own ROCKE hits the bug precisely + because it emits asm directly: at high occupancy it reports intermittent + garbage (>1e20) in the P·V accumulator in roughly 1/3 of runs, seed- and + timing-dependent, and names "proper backend WMMA scheduler" as the correct + fix. **[S]** +- This is a strong argument for Decision #19's discipline on this target, and a + reason not to add an asm fast path for gfx1250. **[I]** + +A second-order trap from the same source **[S]**: switching their softmax from +`ds_swizzle` to DPP *removed* an LDS serialization that had been +**incidentally** supplying the hazard gap, and the kernel began producing NaNs. +The fix was to auto-bump WMMA spacing whenever DPP softmax is enabled. A pure +performance change silently broke correctness through an unmodelled coupling. + +### 2.4 gfx1250 ≠ gfx1251 + +`FeatureGFX125xLowestRateWMMA` is on **gfx1251** and `gfx12-5-generic`, not on +gfx1250. It selects hazard categories 4/5 (latency 16/32) instead of 0/6 +(latency 8/4) — the **same instruction at 4× different throughput**. **[V]** + +Treating the two as one class is correct for **ABI** (identical intrinsics and +operand forms) and wrong for any **cost model, scheduler, or autotuner**. + +### 2.5 The machine model behind the hazard categories + +`SISchedule.td` carries three distinct models — `GFX1250SpeedModel`, +`GFX1251SpeedModel`, `GFX125xGenericSpeedModel` — and they are what +`computeInstrLatency` reads to pick a hazard category (§2.3). **[V]** + +**WMMA runs on a dedicated XDL pipe, not the VALU.** That is the structural +reason co-execution hazards exist at all: + +```tablegen +def HWXDL : ProcResource<1>; // matrix pipe +let ReleaseAtCycles = [4] in def : HWWriteRes; +let ReleaseAtCycles = [8] in def : HWWriteRes; +let ReleaseAtCycles = [16] in def : HWWriteRes; +let ReleaseAtCycles = [32] in def : HWWriteRes; + +def : HWWriteRes; // RDNA-style WMMA: VALU pipe +def : HWWriteRes; +def : HWWriteRes; +``` + +`ReleaseAtCycles == latency` on the XDL entries means the matrix pipe is held for +the full duration — back-to-back WMMA of the same class cannot overlap, but VALU +work *can* run alongside. That overlap is the co-execution the hazard table +governs. + +Note the contrast with RDNA: `Write*PassWMMA` targets `HWVALU`, so **gfx1151 WMMA +occupies the vector ALU** and competes with all vector work, whereas **gfx1250 +WMMA occupies a separate pipe**. A cost model ported from gfx1151 to gfx1250 +without this distinction will misprice every mixed WMMA/VALU loop. **[I]** + +Per-instruction assignment **[V]**: + +| Instruction class | gfx1250 | gfx1251 | +|---|---|---| +| 16x16x64 FP8/BF8 | XDL 1-pass — **4 cyc** | XDL 4-pass — 16 cyc | +| F16 / BF16 | XDL 2-pass — **8 cyc** | XDL 4-pass — 16 cyc | +| 16x16x128 FP8/BF8 | XDL 2-pass — **8 cyc** | XDL 8-pass — **32 cyc** | +| SWMMAC 16x16x128 FP8/BF8 | (2-pass class) | XDL 4-pass — 16 cyc | +| IU8 / IU4 | XDL 4-pass — 16 cyc | XDL 8-pass — 32 cyc | +| 32x16x128 F4 | XDL 2-pass — 8 cyc | XDL 8-pass — 32 cyc | +| F8F6F4 scaled | `WriteWMMAScale_16X16X128_F8F6F4` | `WriteWMMAScaleFP4_16X16X128_F8F6F4` | +| `V_WMMA_F64_16X16X4_F64` | *absent* | `Write4PassWMMA` (VALU, 16 cyc) | + +And the scalar side inverts: + +| | gfx1250 | gfx1251 | +|---|---|---| +| `WriteDouble` | 37 cyc | **6 cyc** | +| `WriteDoubleAdd` | 37 | **5** | +| `WriteTrans64` | 38 | **7** | + +So the two are **complementary SKUs, not revisions**: gfx1250 is the +low-precision matrix part (4–16 cycle WMMA, 37-cycle FP64); gfx1251 is the FP64 +part (16–32 cycle WMMA, 6-cycle FP64, plus an FP64 WMMA gfx1250 lacks). **[V]** + +Shared latencies from `GFX125xCommonWriteRes`, useful as a first-order cost +model **[V]**: + +| Class | Cycles | +|---|---| +| `WriteVMEM` | **320** | +| `WriteLDS` / `WriteSMEM` | 20 | +| `WriteBranch` | 32 | +| `WriteExport` | 16 | +| `WriteSFPU` | 4 | +| `WriteVALUDummy` | 5 | +| `WriteSALU` | 2 | +| `WriteBarrier` | **2000** | + +The 320-cycle VMEM and 2000-cycle barrier are the numbers that make the +pipelining arithmetic in AMD's decode work (§8) come out the way it does. + +--- + +## 3. Data movement — three distinct mechanisms + +gfx1250 has three ways to move data that a compiler must model separately. They +use different instructions, different completion counters, and have different +arch availability. + +### 3.1 Async global→LDS + +``` +global_load_async_to_lds_b{8,32,64,128} → ASYNCcnt +global_store_async_from_lds_b128 → ASYNCcnt +``` + +Intrinsics: `int_amdgcn_global_load_async_to_lds_b*`, signature +`(global_ptr, lds_ptr, offset_imm, cachepolicy_imm)`. **[V]** All assemble on +gfx1250. **[A]** + +**gfx1250 does not have `global_load_lds`** — that is the gfx950 mechanism. The +instruction *and* the counter both differ. Code keyed on "async global→LDS" as a +single concept across gfx950 and gfx1250 is wrong on both axes. **[A]** + +### 3.2 TDM — `tensor_load_to_lds` + +Descriptor-driven bulk tensor transfer; the AMD analog of NVIDIA's TMA. +**gfx1250-exclusive** — rejected by the assembler on gfx950, gfx1151, and +gfx1201. **[A]** + +```c +int_amdgcn_tensor_load_to_lds(v4i32 D#g0, v8i32 D#g1, v4i32 D#g2, + v4i32 D#g3, v8i32 D#g4, i32 cachepolicy) +``` + +- **28 i32 (112 B) of descriptor state** across five register groups. Groups 2 + and 3 are zero for ≤2D tensors; group 4 is reserved and silently ignored. +- **No pointer operands.** Both the global address and the LDS destination live + inside the descriptor. This is a genuinely different op shape from a + `(dst, src, bytes)` copy — not something a byte-count copy op grows into. +- `cachepolicy`: bits[0-2] `th`, bits[3-4] `scope`. +- `IntrConvergent` — a wave-collective operation. +- Completion is **TENSORcnt**, not ASYNCcnt. **[V]** + +Assembly forms **[A]**: + +```asm +tensor_load_to_lds s[0:3], s[4:11] ; ≤2D short form +tensor_load_to_lds s[0:3], s[4:11], s[12:15], s[16:19] ; full 4-group +tensor_load_to_lds s[0:3], s[4:11] th:TH_LOAD_NT scope:SCOPE_SYS +tensor_store_from_lds s[0:3], s[4:11] +``` + +12-byte encoding; unused descriptor groups encode as the null-SGPR `0x7c`. + +Production usage: AMD's Tensile gfx1250 pipeline anchors its entire +cluster-barrier insertion pass on `tensor_load_to_lds` sites, and +`stinkytofu`'s hardware model registers it as a first-class instruction. **[V]** + +There are also `TENSOR_SAVE` / `TENSOR_STOP` VFLAT instructions (opcodes +0x06e/0x06f) for context save/preemption of the tensor unit — not part of the +data path. **[V]** + +**Open:** the D# *field* layout — what the 28 dwords actually contain — is not +in LLVM (the intrinsic passes register groups through opaquely). It needs the +gfx1250 ISA guide or a Tensile descriptor builder. See §11. + +### 3.3 Transposed LDS reads + +`ds_load_tr8_b64` and family (`FeatureTransposeLoadF4F6Insts`) — read LDS +directly into the WMMA B-operand layout. Assembles on gfx1250. **[A]** + +Measured caveat **[S]**: AMD built this path (`ds_load_tr16_b128`) for decode +and measured it **neutral** — the `ds_bpermute` lane-pair stitch needed to +assemble the K=32 B operand ate the LDS-read savings. See §8. + +--- + +## 4. Completion model — split wait counters + +gfx1250 replaces the single combined counter with a split model +(`waitcnt_model = split_gfx1250`). All of these assemble **[A]**: + +``` +s_wait_loadcnt s_wait_storecnt s_wait_dscnt s_wait_kmcnt +s_wait_asynccnt s_wait_tensorcnt s_wait_xcnt +``` + +Mapping, per `AMDGPUUsage.rst` "Memory Model GFX125x" **[V]**: + +| Operation | Counter | +|---|---| +| Load (global, scratch, flat, buffer) | `s_wait_loadcnt` | +| Store (global, scratch, flat, buffer) | `s_wait_storecnt` | +| non-ASYNC LDS | `s_wait_dscnt` | +| **ASYNC LDS** (`global_load_async_to_lds_*`, `cluster_load_async_to_lds_*`) | `s_wait_asynccnt` | +| **Tensor** (`tensor_load_to_lds` / `tensor_store_from_lds`) | `s_wait_tensorcnt` | +| scalar memory (`s_load_*`) | `s_wait_kmcnt` | + +**`s_wait_xcnt` is different in kind** — it increments when a memory operation is +*issued* and decrements when that instruction's **address translation** completes. +Waiting on any memory counter `s_wait_*cnt N` also waits on `s_wait_xcnt N`. It +carries one hard correctness requirement **[V]**: + +> `s_wait_xcnt 0x0` is required before flat and global atomic stores / +> read-modify-write operations to guarantee atomicity during an xnack replay. + +Two ordering caveats that constrain any scheduler **[V]**: + +- Completion (counter decrement) is reported **in issue order within a type**, but + in **no particular order between types**. +- The order in which *data reaches registers* can differ from issue order even + though completion is reported in order — so a `s_wait_*cnt` is required to stop + two in-flight loads targeting the same register from racing. + +### 4.1 Waits are positional, not identity-based + +This is the single most important semantic for anyone lowering async ops: + +> `s_wait_*cnt N` blocks until **at most `N`** ops of that kind remain +> outstanding (it keeps the `N` most-recently-issued in flight and drains +> everything older). **[S]** + +``` +w(D) = n - i - 1 # producer D at index i (0 = oldest) in a FIFO of size n + # emitted wait = countFrom(D) - 1 + # min across all deps constraining the same counter +``` + +**You cannot say "wait for this specific copy."** You can only bound how many +are outstanding, and computing that bound requires the producer's FIFO position +along every CFG path reaching the consumer. `N = 0` is always correct and always +a full drain — it forfeits exactly the overlap async copies exist to provide. + +### 4.2 How AMD lowers tokens into positions + +AMD's production solution is instructive because it starts from the same +abstraction we do. `StinkyBuildImplicitDependencyPass` attaches `MemTokenData` +token IDs to tensor loads, DS ops, and barriers, materializing LDS ordering as +pseudo-register defs/uses — because "a `tensor_load_to_lds` writes an LDS region +and a later `ds_read` of that region depends on it, but there is no vreg linking +them." Then SSA def-use with PHIs at dominance frontiers feeds a forward +dataflow solver. **[S]** + +The parts that make it work: + +- Per-counter FIFO models **tagged per CFG predecessor edge**, so a join + consumer sees each path's depth rather than a collapsed union. +- PHI summaries at merges, taking `min` of `countFrom(src) - 1` over constrained + incoming paths. +- A documented escape hatch: on iteration-cap hit, force every counter to 0 — + "a fully-drained, always-safe plan." +- Anti-dependencies (WAR-on-LDS, barrier ordering) still come from token overlap, + not the SSA RAW chain. + +**Wave-count-dependent policy** **[S]**: tensor-counter RAW deps drain only at +barriers in multi-wave kernels (cross-wave LDS visibility is the barrier's job), +but at *every consumer* when `NumWaves == 1`. Wave count changes wait placement. + +### 4.3 The LLVM lowering seam — `asyncmark` + +**This is the most important section in this document for backend work.** + +`SIInsertWaitcnts.cpp` models `ASYNC_CNT` and `TENSOR_CNT`, but **not the same +way it models the other counters**, and the difference defines the seam between +what LLVM does and what a frontend must do. **[V]** + +The official documentation states it outright: + +> ASYNC LDS and tensor vector memory operations are **not covered by the memory +> model** implemented by the AMDGPU backend. Neither `s_wait_asynccnt` nor +> `s_wait_tensorcnt` are inserted automatically. **They must be emitted using +> compiler built-in calls.** +> — `AMDGPUUsage.rst`, "Memory Model GFX125x" + +and the implementation agrees: + +> `AsyncCnt` and `TensorCnt` always default to `~0u` (don't wait for it). They +> are only updated when a call to `@llvm.amdgcn.wait.asyncmark()` is processed. +> — `SIInsertWaitcnts.cpp:231` + +LLVM's normal waitcnt machinery is **register-dependency driven** (score +brackets over vreg defs/uses). An async LDS-DMA or TDM load writes an *LDS +region*; there is no register linking it to a later `ds_read`. LLVM therefore +cannot infer the dependency — the same reason AMD's asm-level solver needed +`MemTokenData` pseudo-registers (§4.2). So for these two counters LLVM does not +attempt inference, and instead exposes an explicit marker protocol. + +**Two meta intrinsics** — both emit *no hardware instruction*; they are consumed +by the pass, which emits the real `s_wait_asynccnt` / `s_wait_tensorcnt` +immediates **[V]**: + +```llvm +; "Sets a marker in the stream of async requests" +declare void @llvm.amdgcn.asyncmark() ; __builtin_amdgcn_asyncmark +; "Waits until the Nth previous marker is completed, if it exists" +declare void @llvm.amdgcn.wait.asyncmark(i16 immarg) ; __builtin_amdgcn_wait_asyncmark +``` + +Gated by `hasAsyncMark()` = `HasVMemToLDSLoad && GFX1250Plus` (`AMDGPU.td:1458`). + +**Mechanism** **[V]**: + +1. `AsyncScore[T]` accumulates a per-counter snapshot as async/tensor ops are + seen (`shouldUpdateAsyncMark` routes TDM → `TENSOR_CNT`, async LDS-DMA → + `ASYNC_CNT`, non-async LDS-DMA → `LOAD_CNT`). +2. `ASYNCMARK` → `recordAsyncMark()`: pushes `AsyncScore` onto the `AsyncMarks` + vector and **resets it**. Each mark therefore captures the *batch* of async + ops issued since the previous mark. +3. `WAIT_ASYNCMARK N` → `determineAsyncWait(N)`: indexes + `AsyncMarks[size - N - 1]`, derives the real per-counter immediate via + `determineWaitForScore`, then **erases that mark and all older ones**. + +**Division of labour:** + +| LLVM owns | The frontend owns | +|---|---| +| Counter arithmetic and score brackets | **Where the markers go** | +| CFG join merging (`mergeAsyncMarks`) | The value of `N` | +| Loop handling and mark truncation | | +| Overflow guards (`min(UB - Score, getLimit(T) - 1)`) | | +| Emitting the final `s_wait_*cnt N` | | + +So the answer to "must we build a FIFO dataflow solver?" is **no** — LLVM does +the dataflow, including joins. What a frontend must supply is *batch marking*, +which is a substantially smaller and more local problem, and one that maps +naturally onto a token-typed async op. + +**Semantics that bite** **[V]**: + +- **A too-large `N` produces no wait at all**, silently: `if (AsyncMarks.size() + <= N) return {};`. This matches the intrinsic's "if it exists" wording. The + failure mode of a wrong `N` is therefore a **silent race, not a conservative + stall** — correctness sits entirely with the producer of the marks, so any + lowering needs its own verification rather than trusting LLVM to catch it. +- **`MaxAsyncMarks = 16`.** At the cap, `N = min(N, 15)`. The comment is explicit + that this exists to *ensure a non-trivial wait is still generated* after a + merge truncation — so the clamp errs toward waiting more, not less. +- **Marks are consumed.** After servicing, the waited mark and all older ones are + erased, so indices are relative to the live set, not to absolute program order. +- **At joins**, `mergeAsyncMarks` pads the shorter list with zero-marks at the + *front* and merges pairwise from the end — marks align **by recency**, not by + absolute index. +- **Calls do not drain these counters.** `Inst.isCall()` applies a blanket wait + "but `AsyncCnt` and `TensorCnt` are never included in such blanket waits" + (`:2733`). Async state survives a call boundary. +- `ASYNCMARK` blocks waitcnt merging across it, so a mark is also a scheduling + barrier for wait placement. + +--- + +## 5. Workgroup clusters + +`FeatureClusters` appears in exactly two feature sets in all of `AMDGPU.td`: +`FeatureISAVersion12_50_Common` (gfx1250/gfx1251) and `FeatureISAVersion13`. +**gfx950 does not have it.** **[V]** + +### 5.1 Barrier encoding + +There is no distinct cluster-barrier opcode in emitted assembly. Scope is +selected by **barrier ID** **[S]**, all four forms assembling on gfx1250 **[A]**: + +```asm +s_barrier_signal -1 / s_barrier_wait -1 ; workgroup scope +s_barrier_signal -3 / s_barrier_wait -3 ; cluster scope +``` + +LLVM also exposes `int_amdgcn_s_cluster_barrier` / +`__builtin_amdgcn_s_cluster_barrier` as a convenience wrapper **[V]**, but +Tensile emits the raw `-3` form. + +### 5.2 Clusters are Shader-Engine scoped — the key architectural fact + +`cluster` is a **first-class LLVM IR syncscope**, and on gfx125x it lowers to +`scope:SCOPE_SE` **[V]**: + +| LLVM syncscope | ISA | +|---|---| +| *none*, `one-as` | `scope:SCOPE_SYS` | +| `system`, `system-one-as` | `scope:SCOPE_SYS` | +| `agent`, `agent-one-as` | `scope:SCOPE_DEV` | +| **`cluster`, `cluster-one-as`** | **`scope:SCOPE_SE`** | +| `workgroup` / `wavefront` / `singlethread` (+ `-one-as`) | `scope:SCOPE_CU` (default, omitted in asm) | + +That single row explains the whole cluster design. A cluster lives **within one +Shader Engine**, and the SE-shared **L1 buffer** (§1.2) is its coherence point. +Multicast-into-peer-LDS works because the participating WGPs sit behind a common +L1 — which is also why there is no distributed shared address space (§5.7): the +sharing happens in the cache hierarchy, not the address space. + +Semantics of the scope **[V]**: `cluster` synchronizes with `system`, `agent`, or +`cluster` operations executed by a thread **on the same cluster**, plus +`workgroup`/`wavefront` operations in the same work-group/wavefront, for all +address spaces except private. Critically: + +> On targets that do not support workgroup cluster launch mode, this behaves like +> `agent` scope instead. + +So `cluster` syncscope is **portable by construction** — it degrades to `agent` +rather than failing to compile. That makes it safe to emit unconditionally from a +target-independent layer. + +### 5.3 Compile-time cluster declaration + +| Mechanism | Meaning | +|---|---| +| `"amdgpu-cluster-dims"="x,y,z"` fn attr | `"0,0,0"` = cluster disabled; `"1024,1024,1024"` = enabled but dimensions unknown at compile time; anything else = explicit dims. Only meaningful on targets with cluster support. **[V]** | +| `.cluster_dims` | Cluster dimensions recorded in the code-object metadata. **[V]** | +| `"amdgpu-no-cluster-id-{x,y,z}"` | Asserts the kernel never reads the corresponding `llvm.amdgcn.cluster.id.*`, enabling preload/ABI trimming. **[V]** | + +The `"1024,1024,1024"` sentinel is worth noting: it distinguishes "clusters are +on but the shape is dynamic" from "clusters are off", which a lowering must not +conflate. + +### 5.4 Documented gaps — read before planning cluster work + +`AMDGPUUsage.rst`'s GFX125x memory model carries an explicit incompleteness note +**[V]**: + +> This section is currently incomplete as work on the compiler is still ongoing. +> The following is a non-exhaustive list of unimplemented/undocumented features: +> non-volatile bit code sequences, globally accessing scratch atomics, +> **multicast loads**, **barriers (including split barriers) and cooperative +> atomics**. Scalar operations memory model needs more elaboration as well. + +So the two mechanisms clusters are *built on* — multicast loads and cluster +barriers — are not yet covered by the backend's documented memory model, even +though the instructions encode and the intrinsics exist. Anything built on them +today is ahead of the formal model. **[I]** + +### 5.5 Cooperative atomics + +A separate gfx125x mechanism worth recording: wide cooperative load/store across +naturally-aligned, contiguous lane groups within one wave32 **[V]**: + +| Intrinsic | Lane groups | +|---|---| +| `llvm.amdgcn.cooperative.atomic.{load,store}.32x4B` | `0-31` | +| `llvm.amdgcn.cooperative.atomic.{load,store}.16x8B` | `0-15`, `16-31` | +| `llvm.amdgcn.cooperative.atomic.{load,store}.8x16B` | `0-7`, `8-15`, `16-23`, `24-31` | + +Undefined behaviour if used outside the global address space, across a bus that +cannot carry 128B/256B requests (e.g. host memory over PCIe), with an unsupported +lane group, or with more lane groups per wave than the maximum. + +### 5.6 Identity + +``` +__builtin_amdgcn_cluster_id{_x,_y,_z} +__builtin_amdgcn_cluster_workgroup_id{_x,_y,_z} +__builtin_amdgcn_cluster_workgroup_flat_id +__builtin_amdgcn_cluster_workgroup_max_id{_x,_y,_z} +__builtin_amdgcn_cluster_workgroup_max_flat_id +``` +`IntrinsicsAMDGPU.td:168-178`. **[V]** + +### 5.7 Multicast — and what clusters are *not* + +```c +AMDGPUAsyncClusterLoadLDS: (global_ptr, lds_ptr, offset, cachepolicy, + workgroup_broadcast_mask → M0) +int_amdgcn_cluster_load_async_to_lds_b{8,32,64,128} +int_amdgcn_cluster_load_b{32,64,128} +``` + +**AMD clusters are a replication/broadcast primitive plus a barrier — not a +shared address space.** A grep of the full intrinsics file and the HIP API +surfaces no peer-LDS addressing, no `dsmem` analog, no cluster address space. +NVIDIA's `cluster.map_shared_rank` has no counterpart. Data is *multicast into +each workgroup's own LDS*, selected by a broadcast bitmask. **[V]** + +Porting an NVIDIA CGA kernel that reads a peer CTA's shared memory will not +lower. The correct mental model is TMA multicast, not distributed shared memory. + +### 5.8 Host launch + +HIP ≥ 7.0. **[V]** + +```c +hipLaunchAttribute attr[1]; +attr[0].id = hipLaunchAttributeClusterDimension; +attr[0].val.clusterDim = {2, 1, 1}; // must evenly divide gridDim +config.attrs = attr; config.numAttrs = 1; +hipLaunchKernelExC(&config, kernel, params); +``` + +Supporting surface: `hipClusterSchedulingPolicy{Default,Spread,LoadBalancing}`, +`hipOccupancyMaxActiveClusters`, `hipOccupancyMaxPotentialClusterSize`, +`hipFuncAttributeRequiredCluster{Width,Height,Depth}`, +`hipFuncAttributeClusterDimMustBeSet`, +`hipFuncAttributeNonPortableClusterSizeAllowed`, `hipErrorInvalidClusterSize`. + +Device-side directive: `__attribute__((cluster_dims(x,y,z)))` (`Attr.td:1626`), +mutually exclusive with `no_cluster`; HIP wraps it as `CLUSTER_DIMS(X,Y,Z)`. **[V]** + +**Capability detection is runtime, not arch-keyed** (`hip_device.cpp:739`) **[V]**: + +```c +// A cluster of size 1 is a regular single-block launch (legal on all GPUs); +// clusterLaunch advertises multi-block cluster support... +deviceProps.clusterLaunch = info.clusterMaxSize_ > 1; +``` + +The canonical guard in AMD's own tests is `devProp.clusterLaunch != 0`. + +### 5.9 The correctness discipline + +From AMD's Tensile cluster-barrier insertion pass **[S]** — five rules whose +entire purpose is keeping `signal -3` / `wait -3` **paired on every +control-flow path**: + +- Every cluster signal is preceded by a workgroup-scope `signal -1` / `wait -1` + pair, so all waves reach the join before any wave issues the cluster signal. +- Only `WaveIdx == 0` executes the cluster signal. +- The hard cases are loop-entry guards and drain iterations where the paired + `tensor_load_to_lds` is disabled — the handshake must be suppressed on exactly + the same paths, or the pairing breaks. +- Each rule carries its own idempotency check so re-running is a no-op. + +Unbalanced signal/wait is the failure mode, and it is a control-flow problem, +not a local one. + +--- + +## 6. Device-initiated SDMA (inter-GPU) + +Distinct from TDM. TDM is shader-issued, intra-GPU, global→LDS. SDMA is a +separate copy engine, and this code drives it **from inside a kernel** for +inter-GPU transfer. Implemented in `rocshmem/src/sdma/`, consumed by RCCL's +`anvil_sdma` GIN backend. **[V]** + +**Arch support is a closed set**: gfx90a, gfx942, gfx950, gfx1250. Every other +target hits `LOGD_ERROR_ABORT("SDMA is not supported on this architecture")`. +**gfx1151 is not supported.** **[V]** + +### 6.1 Packet ISA — OSS7.0 + +`sdma_pkt_struct_mi4.h`: "OSS7.0 SDMA packet structures (CDNA4 / MI350X and +later)", auto-generated from `OSS_70-sDMA_MAS.md`. **[V]** + +Base opcodes (all generations): `NOP=0`, `COPY=1`, `WRITE=2`, `FENCE=5`, +`TRAP=6`, `POLL_REGMEM=8`, `ATOMIC=10`, `CONST_FILL=11`, `TIMESTAMP=13`. + +MI4-specific sub-opcodes **[V]**: + +| Sub-op | Value | Capability | +|---|---|---| +| `COPY_LINEAR_WAIT_SIGNAL` | 0x0 | **fused wait → copy → signal** | +| `COPY_LINEAR_PHY` | 0x8 | physical-address copy | +| `COPY_SWAP_WAIT_SIGNAL` | 0x9 | swap + fused wait/signal | +| `COPY_LINEAR_MULTICAST` | 0xa | **one copy, many destinations** | +| `COPY_MULTICAST_WAIT_SIGNAL` | 0xa | multicast + fused wait/signal | +| `COPY_PAGE_TRANSFER` | 0xc | page-granular transfer | +| `FENCE_64B` / `POLL_MEM_64B` | 0x2 / 0x5 | 64-bit fence / poll | +| `CONSTANT_FILL_PAGE` | 0x4 | page-granular fill | + +Plus `SDMA_SIGNAL_OP_ADD64_MI4 = 111`, `SDMA_WAIT_FUNC_GEQ_MI4 = 5`. + +### 6.2 The fused packet + +`SDMA_PKT_COPY_LINEAR_WAIT_SIGNAL_MI4` — **19 DWORDs**, `static_assert`-checked **[V]**: + +``` +header : op, subop, tmz, npd, wait, signal ← wait/signal are enable bits +wait : wait_function, wait_scope, wait_temporal_hint, + wait_addr_31_3 / _63_32, ← bit 3: 8-byte aligned + wait_reference_31_0 / _63_32, ← 64-bit reference + wait_mask_31_0 / _63_32 ← 64-bit mask +copy : copy_count, + src_scope, src_temporal_hint, + dst_scope, dst_temporal_hint, + src_addr_31_0 / _63_32, dst_addr_31_0 / _63_32 +``` + +One descriptor expresses: **block until `(mem[wait_addr] & mask) GEQ reference`, +then copy `src → dst`, then signal** — a complete producer-consumer handoff with +independent cache scope and temporal hint on each of the wait, source, and +destination sides. + +Sizes for comparison: `COPY_LINEAR_PHY` 8 DW, `FENCE` 4 DW, `FENCE_64B` 5 DW. +The fused packet costs ~2.4× a plain copy in ring space, which is why callers +gate it behind `useSdmaFusedSignal(...)`. + +### 6.3 Device-side ring protocol + +1 MB ring (`SDMA_QUEUE_SIZE`, "matches rocm-xio sdma-ep"), queue created through +`hsakmt`. **[V]** + +- **Reservation**: lock-free multi-producer CAS on `cachedWptr`. `CanWriteUpto` + caches `rptr` and only re-reads the hardware register when the cached view says + full. +- **Wraparound**: pads with NOPs, count encoded in the first DWORD as + `((numOffsetDwords - 1) & 0xFFFF) << 16`. +- **Packet write**: `static_assert(sizeof(PacketType)/sizeof(uint32_t) <= 64)` — + "Ensure that one warp can write the whole packet." +- **Commit is strictly in reservation order**: spin until `committedWptr == base`. +- **Publish**: three stores, each separated by a full drain, + `__builtin_amdgcn_wave_barrier()`, and a signal fence — + `wptr` (AGENT) → `doorbell` (SYSTEM) → `committedWptr` (AGENT). +- **Completion**: `quietAll()` spin-polls `rptr` until it reaches the target. +- A `SdmaQueueSingleProducerDeviceHandle` subclass drops the CAS for the + one-thread-per-queue case, with identical binary layout. + +### 6.4 gfx1250-specific divergences + +Two places where the chip differs from CDNA, both confirming §4 **[V]**: + +```c +// memory drains +#if defined(__gfx1250__) + asm volatile("s_wait_loadcnt 0x0\n s_wait_storecnt 0x0" ::: "memory"); +#elif defined(__gfx90a__) || defined(__gfx942__) || defined(__gfx950__) + __builtin_amdgcn_s_waitcnt(0); +#endif + +// atomics — different mnemonic AND different cache-modifier syntax +// gfx1250: flat_atomic_cmpswap_b64 %0, %1, %2 scope:SCOPE_SYS nt +// CDNA: flat_atomic_cmpswap_x2 %0, %1, %2 sc0 nt +``` + +gfx1250 replaced CDNA's `sc0`/`sc1` bit flags with **named scopes** +(`scope:SCOPE_SYS`) across both the DMA path and the TDM instructions. + +### 6.5 SIMD reconvergence deadlock + +From `submitPacket` **[S]**: + +> All stores inside the loop to avoid SIMD reconvergence deadlock: the +> `committedWptr` update must complete before this lane becomes inactive, so +> that other lanes in the same wavefront can proceed. + +A lock-free in-order commit protocol executed by lanes of the same wavefront can +deadlock on reconvergence: if a lane wins its turn, exits the loop, and goes +inactive before publishing, the remaining lanes spin forever on a turn that is +never released. The fix is structural — keep every store inside the loop body. +**Any device-side lock-free protocol has this failure mode.** + +--- + +## 7. Scale-up fabric — UALink / IFoE + +Not compiler-facing today, recorded for the distributed track. + +**Stack** **[V]**: + +| Layer | Artifact | +|---|---| +| Kernel | `ifoe.ko`, `ifoe-cfg.ko`, `ifoe-cmd.ko`; `/dev/cbl-cfg-ifoe.cfg.0` | +| sysfs | `/sys/class/drm/renderD/device/ualink/` — presence gates everything | +| amdsmi | `projects/amdsmi/{src,include}/ualoe_lib/` (~350 KB of headers) | +| ROCr | `props->FabricHandleSupported` in HSA node properties | +| RCCL | `ARSMI_get_fabric_info()` → `struct ARSMI_fabricInfo` | +| HIP | `hipMemFabricHandle_t`, `hipMemHandleTypeFabric = 0x8`, `hipDeviceAttributeHandleTypeFabricSupported` | +| rocSHMEM | `src/memory/hip_allocator_vmm_fabric.cpp` | + +**Topology model** **[V]**: + +```c +typedef enum { ARSMI_FABRIC_TYPE_UALOE = 0, // over Ethernet + ARSMI_FABRIC_TYPE_UALLINK = 1, // native + ARSMI_FABRIC_TYPE_UNKNOWN = 2 } ARSMI_fabric_type_t; + +struct ARSMI_fabricInfo { + int supported; ARSMI_fabric_type_t fabric_type; + ARSMI_fabric_accelerator_vpod_state_t accel_state; // UNCONFIGURED→CONFIGURED→READY→ACTIVE→ERROR + ARSMI_fabric_npa_address_mode_t addr_mode; + uint32_t accel_id; + uint8_t ppod_id[16]; uint32_t ppod_size; // physical pod + uint32_t bandwidth; // Mb/s + uint32_t latency; // ns + uint32_t vpod_id; uint32_t vpod_size; // virtual pod (the partition you run in) +}; +``` + +Three points worth carrying forward **[I]**: + +1. Ethernet-transported and native UALink are two values of one enum behind one + query — software does not branch on transport, it reads bandwidth/latency. +2. **The fabric reports its own bandwidth and latency.** A planner can query the + interconnect cost model instead of hardcoding it. +3. Physical vs virtual pod are separate; collective planning must key on the + **virtual** pod, since a rack may be partitioned. + +**Device-initiated comms** (RCCL GIN — upstream NCCL device API with pluggable +backends: `gdaki`, `proxy`, `rocshmem_gda`, `anvil_sdma`) **[V]**: + +- Peer addressing is **arithmetic**: `base + (peer - rank) * vmmStride + off`. + VMM fabric handles map every peer's buffer at uniform stride in one VA space. + Fallbacks: `remote_vas[peer]`, then an IPC table scan. +- "LSA memory is VMM-mapped fine-grain (cache-coherent via Infinity Fabric). + Plain stores are visible to all peers through the HW coherence domain." +- **Threshold-driven transport arbiter**: `bytes <= rsCtx->sdmaThreshold` → + `ipcPut` (coherent stores); else the SDMA queue; degrades to `ipcPut` if no + queue handle. This is a measured arbiter over two implementations of one + logical op — the same shape as Decision #28's kernel arbiter, applied to + transport. +- **Fencing is path-dependent**: `sdma_anvil::quiet(handle)` vs + `__builtin_amdgcn_fence(RELEASE, "agent")` vs `__threadfence_system()`, + selected by which data path the put took, plus a `cuda::thread_scope` + negotiation that skips the fence when the caller already gave a stronger scope. + +An **[I]** worth flagging: MCDI, and block names like the **EX** switch, +**XRSEC** crypto, and **XRPFC** flow control, are distinctly Solarflare/Xilinx +vocabulary. This reads like Xilinx/Solarflare NIC IP repurposed as the fabric +NIC, but the source never says so — treat as hypothesis. + +--- + +## 8. Measured lessons from AMD's own gfx1250 decode kernels + +From `rocke`'s MHA/decode optimization case study **[S]**. The decode kernel is +memory-latency/bandwidth bound (GEMV-like, arithmetic intensity ≈ 1), so +compute- and LDS-oriented levers do not touch the bottleneck. + +**Worked:** + +| Lever | Result | +|---|---| +| **DPP `row_xmask` softmax butterfly** | replaced a 4-stage `ds_swizzle_b32` chain (LDS-port, `lgkmcnt`-serialized) with a VALU DPP butterfly: **128 `ds_swizzle_b32` → 0**; **1.4–1.56×** at low wave count; numerically exact (`max_abs ≈ 4e-5`) | +| Cooperative multi-wave CTA | up to ~1.5× on **small batch only** | +| `num_segments` (split-KV) | AITER-style occupancy knob; helps small batch, clamp for large | + +**Did not work** — the recurring cause is the **wave32 cross-lane tax**: +optimizations that win on gfx950 (wave64, MFMA) lose on gfx1250 because +assembling operands costs more shuffles. + +| Lever | Result | +|---|---| +| Register-resident P (skip P→LDS) | **~2× slower** (259→671 µs) | +| HW transpose LDS reads (`ds_load_tr16_b128`) | **neutral** — `ds_bpermute` stitch eats the savings | +| SW pipeline + `iglp_opt` / `sched_group_barrier` | **neutral** — waitcnt already tuned; softmax chain was the real critical path | +| Double-buffered V | **slower** (259→335 µs) — halves occupancy to hide already-hidden latency | +| Multi-wave at large batch | **4–10× slower** (259→1061 µs) — device already saturated | +| Native fp8 PV GEMM | not built — ceiling probe showed PV+V-staging is **0.6%** of the kernel | + +**Methodology worth adopting** **[S]**: AMD's WMMA layout probe uses **random +asymmetric A/B** deliberately — "a row/col swap in the lane map transposes the +result and fails verify, so a PASS at multiple tile counts uniquely confirms the +mapping." A symmetric or structured test matrix passes with a transposed lane +map. They also treat their own lane maps as "a hypothesis until proven on +silicon." + +**Also worth adopting** **[S]**: the `ablate_pv` ceiling probe — before building +a from-scratch optimization, measure its *ceiling* by ablating the region it +would improve. That is what killed the native-fp8 PV work at 0.6% before any +implementation cost was paid. + +--- + +## 9. Observations against the current Tessera tree + +**These are observations from this survey, not an accepted backlog.** The +authoritative ROCm queue is the open-actions table in +[`ROCM_AUDIT.md`](ROCM_AUDIT.md); anything below needs its own verification and +triage before it earns an ID. + +| # | Observation | Where | +|---|---|---| +| 1 | `cluster_mode` is inverted: asserted `"ready"` for gfx950 (which lacks `FeatureClusters`) and `"tba"` for gfx1250/1251 (which have it). `supports_cluster_mode()` is a live predicate over that dict, but has no codegen consumer today — latent, not miscompiling. | `python/tessera/compiler/rocm_target.py` | +| 2 | gfx1250 LDS is `65536` marked PROVISIONAL; grounded value is `327680`. | `rocm_target.py` | +| 3 | gfx1250 VGPR budget is `256` (RDNA-derived); grounded value is `1024`. Matters because `rocm_tiling.py`'s thesis is that the register budget is the dominant tiling lever. | `rocm_target.py` | +| 4 | `ROCM_WaitTokenOp` has **no wait immediate** — only a counter name, and the enum admits just `vmcnt`/`lgkmcnt`. It can express `N = 0` (full drain) and nothing else, which is correct but forfeits async overlap. The token-typed design matches AMD's own `MemTokenData`. **Scoped by §4.3:** LLVM owns the dataflow, so the work is *batch marking* — lower `async_copy` groups to `llvm.amdgcn.asyncmark()` and `wait(token)` to `llvm.amdgcn.wait.asyncmark(N)`, where `N` is the number of marks between the token's batch and the wait. No FIFO solver needed. Note the failure mode: a too-large `N` yields **no wait**, so this needs its own verification. | `src/compiler/codegen/Tessera_ROCM_Backend/include/TesseraROCM/IR/TesseraROCMOps.td` | +| 5 | `_GFX1250_CLASS_ARCHES = {gfx1250, gfx1251}` is correct for ABI, wrong for cost models (§2.4). Worth a boundary comment. | `python/tessera/compiler/rocdl_emit.py` | +| 6 | `capabilities.py` lists gfx1250 dtypes as bf16/fp16/fp32/int8 — no fp8 — while gfx950 above it carries fp8 and a `wmma_f8` flag. gfx1250 is the arch with the scaled-fp8 matrix path. | `python/tessera/compiler/capabilities.py` | +| 7 | Module docstring calls gfx950 "MI325X"; MI325X is CDNA 3 / gfx942. `gpu_target_map.py` has it right (MI350X/MI355X/MI350P). | `rocm_target.py` | +| 8 | Generated softmax/reduce kernels use `gpu.shuffle xor` butterflies, which lower to `ds_bpermute`/`ds_swizzle` — the exact LDS-port pattern AMD replaced for 1.4–1.56×. The DPP `row_xmask` lever is available on **gfx1151** with an arch-keyed mnemonic (`v_max_f32_dpp` on gfx1151; `v_max_num_f32_dpp` on gfx12+; `v_add_f32_dpp` is common). Needs the ROCM-6 A/B ratchet — DPP is a latency-hiding lever and can lose when VALU-throughput-bound. | `GenerateROCMSoftmaxKernel.cpp:89`, `GenerateROCMReduceKernel.cpp:125`, `GenerateROCMArgReduceKernel.cpp:113` | +| 9 | `attn_split_kv.py` (`plan_split_kv`) has no consumer outside its own unit test. AMD's `num_segments` is the same knob and is occupancy-gated in their production dispatcher. | `python/tessera/compiler/attn_split_kv.py` | +| 10 | WMMA fragment fixtures should use **random asymmetric** operands (§8) or they can pass with a transposed lane map. | `tests/unit/test_rocm_wmma_gemm_generated.py` | +| 11 | `distributed_planner.py` has no link bandwidth/latency terms. If it grows them, `(bandwidth, latency, ppod, vpod)` is the shape the hardware offers (§7). | `python/tessera/compiler/distributed_planner.py` | + +### 9.1 The cross-cutting theme + +Three independent mechanisms, one pattern: + +| Mechanism | How the dependency is carried | +|---|---| +| TDM (`tensor_load_to_lds`) | descriptor register groups, no pointer operands | +| waitcnt (`s_wait_*cnt N`) | FIFO **position**, not identity | +| SDMA (`COPY_LINEAR_WAIT_SIGNAL`) | wait predicate fused **into** the copy descriptor | + +AMD hardware consistently wants the dependency expressed **inside the +descriptor**, not as a separate ordering instruction the compiler infers. +Token-typed async ops are the right *source* abstraction — AMD's own design +confirms it — but every lowering target on this vendor converts tokens into a +positional count or a descriptor field. That argues for a single "async +dependency → target completion model" lowering interface rather than three +ad-hoc ones. **[I]** + +--- + +## 10. Reproducing every claim + +The assembler probes need no AMD hardware and no ROCm install — any LLVM with +the AMDGPU target works (verified on Homebrew LLVM 22.1.8, arm64 macOS). + +```bash +LLC=$(brew --prefix llvm)/bin/llc +MC=$(brew --prefix llvm)/bin/llvm-mc + +# Subtarget features visible to this LLVM +$LLC -march=amdgcn -mcpu=gfx1250 -mattr=help 2>&1 | grep -iE "cluster|async|1024|tensor" + +# Does an instruction exist on this arch? +echo "tensor_load_to_lds s[0:3], s[4:11]" | $MC -arch=amdgcn -mcpu=gfx1250 -show-encoding +echo "s_barrier_signal -3" | $MC -arch=amdgcn -mcpu=gfx1250 -show-encoding +echo "s_wait_tensorcnt 0" | $MC -arch=amdgcn -mcpu=gfx1250 -show-encoding +``` + +Fetching primary sources (needs an authenticated `gh`): + +```bash +gh api -H "Accept: application/vnd.github.raw" \ + "repos/ROCm/llvm-project/contents/llvm/lib/Target/AMDGPU/AMDGPU.td?ref=amd-staging" +``` + +Useful paths: `llvm/lib/Target/AMDGPU/{AMDGPU.td,GCNHazardRecognizer.cpp}`, +`llvm/include/llvm/IR/IntrinsicsAMDGPU.td`, `clang/include/clang/Basic/Attr.td` +(all `ROCm/llvm-project@amd-staging`); +`projects/{hip,rccl,rocshmem,amdsmi,clr,hip-tests}` (`ROCm/rocm-systems@develop`); +`shared/stinkytofu`, `dnn-providers/hip-kernel-provider/rocke` +(`ROCm/rocm-libraries@develop`). + +**Toolchain note.** Two LLVMs are present on the current Mac: the Homebrew keg +`llvm/22.1.8` (on `PATH`), and **LLVM 23.1.0-rc1 at +`/opt/homebrew/llvm-23.1.0-rc1/`** — the latter is what `build/CMakeCache.txt` +pins as `LLVM_DIR`. Note it is *not* at the `/opt/homebrew/opt/llvm@23` path +`CLAUDE.md` cites, so that reference is stale even though the toolchain itself is +present. + +Both versions were probed and agree on everything in this document: each encodes +the gfx1250 WMMA family, TDM (`tensor_load_to_lds`), the cluster-scope barrier +IDs, and the split wait counters (`s_wait_tensorcnt`, `s_wait_asynccnt`), and +both reject the `s_cluster_barrier` convenience mnemonic. Either is sufficient +for the probes above. + +--- + +## 11. Unexplored sources, ranked + +Everything below is a real gap in this survey, not a formality. + +> **Two former top items have been read.** +> **`SIInsertWaitcnts.cpp`** → §4.3: LLVM owns the counter dataflow including CFG +> joins; the frontend owns marker placement via +> `llvm.amdgcn.{asyncmark,wait.asyncmark}`. Resolves the question behind §9 #4. +> **`AMDGPUUsage.rst`** → §1.2 (cache hierarchy + SCOPE ladder), §4 (official +> counter table, `s_wait_xcnt`), §4.3 (official confirmation), §5.2–5.5 (cluster +> syncscope, `amdgpu-cluster-dims`, documented gaps, cooperative atomics). +> +> **A dead end worth recording:** `AMDGPUUsage.rst` refers three times to a +> section `amdgpu-dma-operations` for "full documentation" of the async-LDS and +> tensor intrinsics — and that anchor **is never defined anywhere in the file**. +> The TDM descriptor layout is not merely hard to find in LLVM's docs; the +> section that would hold it is an unwritten placeholder. **[V]** + +**High value for compiler work:** + +1. **`SIMemoryLegalizer.cpp`** — still worth reading for *how* the fence code + sequences in the GFX125x tables are actually emitted, though + `AMDGPUUsage.rst` now gives us the normative sequences themselves. +2. **`SISchedule.td` / the gfx1250 scheduler model** — the WMMA latencies + (4/8/16/32) that drive §2.3 come from `computeInstrLatency`; the model behind + them is unread. +3. **The gfx1250 ISA guide** — now confirmed as the *only* remaining path to the + TDM D# field layout (§3.2), LDS/VGPR allocation granularity, and the + SPG 4.6.12.1 hazard section LLVM cites. Neither LLVM's docs nor its tablegen + carry the descriptor field definitions. `rocm_target.py` grounds gfx1151 in + the published RDNA3.5 guide, so an equivalent document may exist. + +**High value for kernel/codegen design:** + +4. **`shared/stinkytofu` in full** — an entire asm-level optimizing compiler from + AMD: pass pipeline, DAG scheduler (`ReadyQueue.hpp`, and a `CDNA5.hpp` we did + not open), `AsmVerifierPass`, `stinkytofu-opt` tool, `waitcnt-check`. We read + two docs out of many. +5. **`shared/rocroller`** — AMD's GEMM kernel generator with its own IR. It uses + `tensor_load_to_lds` in `MemoryInstructions.cpp`, so it likely contains the + descriptor construction we could not find. +6. **Triton's AMD backend** — `third_party/amd` in `triton-lang/triton`, + including `mxfp_fa_gfx1250.py` (the blog's kernel) and the Gluon layout + system. The concrete fragment layouts and the `cga_layout` field live here. +7. **`rocke`'s attention kernels themselves** — `wmma_attention_fwd.py`, + `_wmma_attention_common.py`, `attention_tiled_2d.py`. We read the design docs, + not the code. +8. **`projects/composablekernel`** — AMD's tile-programming library. Relevant to + a tile-centric compiler on general principle; unexamined here. +9. **rocWMMA's gfx1250 support** — fragment layouts in a portable, documented form. + +**Lower priority / situational:** + +10. **MES (Micro Engine Scheduler)** for cluster scheduling — we already keep an + RDNA MES write-up at `docs/reference/isa/rdna/mes/`; the gfx1250 cluster + dispatch path likely involves it. +11. **`clr`/`hipamd` cluster implementation** — how `hipLaunchKernelExC` actually + programs cluster dims into the AQL packet. +12. **`rocprofiler-sdk`** counters for TDM/cluster/SDMA — required before any + measured gfx1250 work, irrelevant before it. +13. **KFD / kernel driver ABI** for TDM and cluster enablement. + +--- + +## 12. Sources + +**LLVM** (`ROCm/llvm-project@amd-staging`) — `llvm/lib/Target/AMDGPU/AMDGPU.td`, +`GCNHazardRecognizer.cpp`, `FLATInstructions.td`; +`llvm/include/llvm/IR/IntrinsicsAMDGPU.td`; `clang/include/clang/Basic/Attr.td`. + +**ROCm systems** (`ROCm/rocm-systems@develop`) — +`projects/hip/include/hip/hip_runtime_api.h`; +`projects/clr/hipamd/src/hip_device.cpp`; +`projects/hip-tests/catch/unit/cluster/{hipClusterLaunch.cc,hipClusterCompilerDirective.cc,ClusterHelper.hpp}`; +`projects/rocshmem/src/sdma/{anvil_device.hpp,sdma_opcodes.h,sdma_pkt_struct_mi4.h}`; +`projects/rccl/src/include/nccl_device/gin/{gin_device_api.h,anvil_sdma/*}`; +`projects/rccl/src/include/alt_rsmi.h`; +`projects/rocr-runtime/libhsakmt/src/topology.c`; +`projects/amdsmi/{src,include}/ualoe_lib/*`. + +**ROCm libraries** (`ROCm/rocm-libraries@develop`) — +`shared/stinkytofu/docs/developer/cluster-barrier.md`, +`shared/stinkytofu/docs/user/stinky-waitcnt-insertion-pass.md`, +`shared/stinkytofu/hardware/src/gfx/Gfx1250/Gfx1250.cpp`; +`dnn-providers/hip-kernel-provider/rocke/library/builders/gfx1250/attention/{gfx1250_mha_optimization_case_study.md,gfx1250_universal_attention_plan.md}`; +`dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/examples/gfx1250/wmma_probe.py`. + +**AMD blog** — "Attention Decode on MI450 with Gluon", +https://rocm.blogs.amd.com/software-tools-optimization/gluon-attention-decode-mi450/README.html +(secondary; every claim taken from it was re-derived above or is marked **[S]**). + +--- + +*See also: [`ROCM_AUDIT.md`](ROCM_AUDIT.md) (status) · +[`ROCM_PATTERNS_FROM_AMD_ECOSYSTEM.md`](ROCM_PATTERNS_FROM_AMD_ECOSYSTEM.md) +(ecosystem patterns) · [`STRIX_HALO_EXECUTION_PLAN.md`](STRIX_HALO_EXECUTION_PLAN.md) +(gfx1151 bring-up) · [`../BACKEND_AUDIT.md`](../BACKEND_AUDIT.md) (cross-backend).* diff --git a/docs/audit/backend/rocm/ROCM_AUDIT.md b/docs/audit/backend/rocm/ROCM_AUDIT.md index 4061f679a..79d2b707b 100644 --- a/docs/audit/backend/rocm/ROCM_AUDIT.md +++ b/docs/audit/backend/rocm/ROCM_AUDIT.md @@ -16,6 +16,9 @@ log. Detailed Strix Halo bring-up history lives in [`STRIX_HALO_EXECUTION_PLAN.md`](STRIX_HALO_EXECUTION_PLAN.md), and reusable AMD design guidance lives in [`ROCM_PATTERNS_FROM_AMD_ECOSYSTEM.md`](ROCM_PATTERNS_FROM_AMD_ECOSYSTEM.md). +Primary-source gfx1250 / MI450 target facts — WMMA shapes and hazards, TDM, +workgroup clusters, the split wait-counter model, device-initiated SDMA — live in +[`GFX1250_MI450_COMPILER_REFERENCE.md`](GFX1250_MI450_COMPILER_REFERENCE.md). > **Status authority.** ROCm proof is recorded at exact-target granularity. The > generic `rocm` name is a family rollup and never inherits compile, execution, diff --git a/docs/audit/backend/rocm/ROCM_PATTERNS_FROM_AMD_ECOSYSTEM.md b/docs/audit/backend/rocm/ROCM_PATTERNS_FROM_AMD_ECOSYSTEM.md index dad093b09..f002538c8 100644 --- a/docs/audit/backend/rocm/ROCM_PATTERNS_FROM_AMD_ECOSYSTEM.md +++ b/docs/audit/backend/rocm/ROCM_PATTERNS_FROM_AMD_ECOSYSTEM.md @@ -11,6 +11,18 @@ audit_role: reference > [`ROCM_AUDIT.md`](ROCM_AUDIT.md); for the hardware bring-up ladder see > [`STRIX_HALO_EXECUTION_PLAN.md`](STRIX_HALO_EXECUTION_PLAN.md). > +> **gfx1250 / MI450 depth lives in a sibling document.** §1.8 below covers the Gluon +> GEMM tutorial at survey depth; the primary-source target reference — hardware +> constants, WMMA shapes/ABI/hazards, the three data-movement mechanisms, split wait +> counters, workgroup clusters, device-initiated SDMA, and the UALink/IFoE fabric — is +> [`GFX1250_MI450_COMPILER_REFERENCE.md`](GFX1250_MI450_COMPILER_REFERENCE.md). +> +> **AMD's two kernel *compilers* are surveyed separately.** This document covers +> kernel *libraries*; StinkyTofu (asm-level pass optimizer, gfx1250+) and rocRoller +> (dual-graph kernel generator) are read for compiler architecture, algorithms, and +> codegen technique in +> [`../../compiler/AMD_KERNEL_COMPILER_SURVEY.md`](../../compiler/AMD_KERNEL_COMPILER_SURVEY.md). +> > Surveyed: **AITER**, **ATOM**, **hipBLASLt**, **rocWMMA**, **Mori**, **Iris**, **XIO**, > and the **AMD Gluon GEMM tutorial**. > diff --git a/docs/audit/compiler/AMD_KERNEL_COMPILER_SURVEY.md b/docs/audit/compiler/AMD_KERNEL_COMPILER_SURVEY.md new file mode 100644 index 000000000..587b28c6a --- /dev/null +++ b/docs/audit/compiler/AMD_KERNEL_COMPILER_SURVEY.md @@ -0,0 +1,1178 @@ +--- +last_updated: 2026-07-28 +audit_role: reference +--- + +# AMD Kernel-Compiler Survey — StinkyTofu, rocRoller, Composable Kernel & hipBLASLt + +> **Purpose.** Four production AMD kernel *compilers, tile frameworks and +> selection engines* read at source level for transferable architecture, +> algorithms, and codegen technique. +> This is design input for Tessera's compiler, not a target-facts reference and +> not a status surface. +> +> For gfx1250/MI450 **target facts** (ISA, hazards, counters, clusters, SDMA) see +> [`../backend/rocm/GFX1250_MI450_COMPILER_REFERENCE.md`](../backend/rocm/GFX1250_MI450_COMPILER_REFERENCE.md). +> For the AMD **ecosystem survey** (AITER, hipBLASLt, rocWMMA, Mori, Iris, XIO, +> Gluon) see +> [`../backend/rocm/ROCM_PATTERNS_FROM_AMD_ECOSYSTEM.md`](../backend/rocm/ROCM_PATTERNS_FROM_AMD_ECOSYSTEM.md). +> For our own direction see +> [`COMPILER_THEORY_OF_OPERATION.md`](COMPILER_THEORY_OF_OPERATION.md). +> +> **Provenance:** **[V]** verified from source; **[S]** stated by an in-repo AMD +> design document; **[I]** inference or recommendation. +> +> Surveyed 2026-07-28 against `ROCm/rocm-libraries@develop`. + +--- + +## 0. Why these four + +Most of the AMD ecosystem is *kernel libraries*. These four are compilers, +compile-time frameworks, or selection engines, and together they cover the whole +pipeline from problem to chosen kernel: + +| | StinkyTofu | rocRoller | Composable Kernel (`ck_tile`) | hipBLASLt / TensileLite | +|---|---|---|---|---| +| Input | existing AMD **assembly** | a high-level `Command` | C++ template kernel source | a runtime GEMM call | +| Output | optimized assembly | assembly → Comgr → HIP execution | compile-time-specialized kernels | a *selected* kernel + launch | +| Core artifact | Logical IR → Asm IR | `KernelGraph` (dual graph), lowered in place | coordinate algebra in templates | `SolutionLibrary` selection tree | +| Role | post-pass optimizer for hipBLASLt/TensileLite | standalone kernel generator | tile programming model + kernel library | the shipping GEMM library + arbiter | +| Arch scope | **gfx1250 and beyond** (older use `rocisa`) | gfx9 → gfx12 | gfx9 → gfx12 | all supported | + +StinkyTofu is what AMD built *specifically because gfx1250 needed a new asm-level +optimizer*. rocRoller is the closest existing thing to what Tessera is trying to +be. **Composable Kernel is the most mature tile-programming model in the +ecosystem, and its coordinate framework (§3) is the single most transferable idea +in this survey.** **hipBLASLt/TensileLite (§4)** is the production arbiter — how a +shipping library actually chooses among generated, hand-written and predicted +kernels. Reading all four shows which of our design choices are convergent and +which are idiosyncratic. + +--- + +## 1. StinkyTofu — an LLVM-inspired asm-level pass optimizer + +> "StinkyTofu is an LLVM-inspired pass-based IR optimizer for AMD GPU assembly +> kernels, used by hipBLASLt/TensileLite via Python bindings." **[S]** + +### 1.1 Two IR levels — independent convergence with Decision #19 + +**Logical IR** (architecture-agnostic, high-level) → **Asm IR** (concrete, +architecture-specific). Passes operate on Asm IR; core types are +`StinkyInstruction`, `Function`, `BasicBlock`. **[S]** + +That is exactly Decision #19 ("backends expose a hardware-free Target IR before +hardware-specific lowering"), arrived at independently by a team whose only goal +was making hipBLASLt faster. Worth citing the next time the two-level cost is +questioned. + +### 1.2 Cost is *in* the IR + +The textual form carries scheduling cost as instruction attributes **[S]**: + +``` +st.func @name() { +^entry: + v0 = "st.v_mul_f32"(v1, v2) { issueCycles = 1, latencyCycles = 5 } +} +``` + +`issueCycles` / `latencyCycles` are materialized on the instruction rather than +looked up in a side table at scheduling time. Consequences worth stealing **[I]**: +the IR is self-describing for cost, a dumped kernel is a complete scheduling +problem, and cost-model changes are visible as IR diffs rather than invisible +behavioural drift. Our `#tile.layout`-style attributes already prove we can carry +this kind of data; cost is a natural next tenant. + +### 1.3 The hardware description is data, not code + +``` +hardware/src/gfx/GfxXXX/ + GfxXXXInstructions.def (DEF_T / DEF_BATCH) + GfxXXXFormats.def + arch.cmake (ARCH_MAJOR, ARCH_WAVEFRONT, costs, register limits) + │ tablegen + ▼ +hardware/generated/ GfxXXX_{init,costs,operands,block}.inc + ▼ + gfxisa lib → stinkytofu lib → tools / Python bindings +``` + +> "New architectures require only adding a `hardware/src/gfx/GfxXXX/` directory +> with `.def` files — **no C++ edits** for instruction definitions." **[S]** + +This is Decision #28's `TargetPlugin` seam realized as **data plus codegen** +rather than as a C++ interface. The test of the seam is falsifiable and cheap: +*can a new arch land without touching the optimizer's source?* Worth adopting as +an explicit acceptance criterion for our own backend-plugin work. **[I]** + +### 1.4 Region-scoped scheduling via `ScopeAdaptor` + +> "`ScopeAdaptor` extracts instruction regions (identified by named groups like +> `loopWithPrefetch`) into temporary `Function` objects for isolated scheduling, +> then splices results back." **[S]** + +A cheap way to get region-local scheduling without a real outlining pass or a +region dialect: lift, schedule in isolation, splice. The named-group tagging +(`loopWithPrefetch`) is applied by the *kernel generator* upstream — the +optimizer never has to rediscover loop structure heuristically. That +producer-tags/consumer-trusts split recurs (see §1.6) and is the single most +repeated idea in this codebase. **[I]** + +### 1.5 Declarative peephole patterns + +Optimizations are compiled from `.pattern` files with a three-block grammar +**[V]**: + +``` +pattern PatternName { + match { $mul = v_mul_f32 $tmp, $a, $b + $add = v_add_f32 $dst, $tmp, $c } // temporal order, dest first + constraints { ... } + rewrite { ... } +} +``` + +Variables are untyped (`$x`) and the system infers instruction-vs-register-vs- +constant from context. Instructions in `match` are listed in execution order, and +data flow is expressed by reusing a variable (`$tmp`). + +The transferable part is not the syntax but the **separation of the pattern +corpus from the pass**: peepholes become a reviewable, testable data file rather +than C++ that only its author can audit. **[I]** + +### 1.6 `StinkyWmmaVgprReorderPass` — the best-structured pass in either project + +A **read-only analysis** that finds VGPR savings in N-buffered GEMM loops. It +never mutates anything; downstream passes act on its result. **[S]** + +**The optimization.** In an N-buffered GEMM the WMMA instructions split into N +pools, each with its own registers for the pool-varying operand. When that +operand is the *outer* loop dimension its registers stay live across every inner +iteration, blocking cross-pool aliasing. Making it the *inner* dimension — +grouping its instructions contiguously within each pool — tightens the liveness +interval so pool N's registers die before pool 0 next needs them, and the two +pools can alias the same physical registers. + +**The structure — three swappable layers**, each replaceable without touching the +others: + +| Layer | Interface | Built-in | +|---|---|---| +| 1 | `IRegLivenessAnalysis::computeLiveness(bb, wmmaSeq)` | `WmmaIntervalLiveness` — interval = `[first WMMA reading the group, last]`. A full-instruction liveness backend drops in unchanged. | +| 2 | `IWmmaReorderAlgorithm::solve(pools, liveness)` → `{desiredOrder, aliases}` | `PoolVaryingReorderAlgorithm` — fires when `interval width > number of distinct B groups` | +| 3 | `WmmaReorderAnalysisResult` | `{applicable, desiredWmmaOrder, replacements, totalVgprSaved}` — the stable contract | + +Both layers are injected at construction; `nullptr` selects the default. + +Two details worth copying wholesale **[I]**: + +- **Pool tagging over heuristics.** TensileLite stamps each WMMA with a + `WmmaPoolData{poolIndex}` modifier at generation time. *"If any wmma + instruction is missing this modifier the pass bails out for the entire basic + block. Partial tagging indicates a misconfigured pipeline and the pass must not + proceed on incomplete information."* Fail loudly on partial metadata rather + than silently degrading — that is Decision #21's spirit applied to an analysis. +- **Symmetry via detection, not special-casing.** `detectABIndices` decides which + operand is pool-varying by checking register-group intersection across pools, + so the pass handles hardware-A-varying and hardware-B-varying kernels with one + code path. + +### 1.7 Optimization remarks as a first-class output + +`LoopRegionRemarkPass` emits remarks on **loop health**: region count, boundary +causes, **`s_nop` waste**, branch count. **[S]** + +The compiler reporting *why* a schedule is bad — and quantifying wasted issue +slots — is a diagnostics idea we do not currently have. It pairs naturally with +the Evaluator: a remark stream is a cheap, structured signal that does not need +hardware to be useful. **[I]** + +### 1.8 Other notes + +- **Pseudo-PHI nodes.** `buildUseDefChain()` inserts PHIs at CFG joins for + cross-block def-use; they are never emitted (`AsmEmitter` skips them), and any + code that counts instructions must skip `GFX::PHI`. Same device as the + `MemTokenData` pseudo-registers used for wait insertion. **[S]** +- **Intrinsic system.** High-level ops (ReLU, Clamp) live in + `src/ir/logical/Intrinsics.intrinsic`, compile to a binary `intrinsics.st.bc` + at build time, and load at runtime through `IntrinsicRegistry` — a two-stage + build that avoids a circular dependency with TableGen. **[S]** +- **`stinkytofu-opt` + `stinkytofu-check`.** An opt-style driver with a pass + registry plus FileCheck-style tests. Structurally identical to + `tessera-opt` + lit. **[S]** + +--- + +## 2. rocRoller — a dual-graph kernel generator + +> "RocRoller transforms high-level kernel specifications (`Command`) through a +> dual-graph IR (`KernelGraph` with `ControlGraph` + `CoordinateGraph`) into +> optimized GPU assembly, then assembles to binary via AMD Comgr and executes via +> HIP." **[S]** + +### 2.1 The KernelGraph — the headline idea + +A single `KernelGraph` encodes **three** things at once, and is *iteratively +rewritten in place* until it is low-level enough to emit from **[S]**: + +1. **Coordinate transforms** — how vector/matrix/tensor indices are computed from + each other. +2. **Control flow and operations** — operations and their dependencies. +3. **Data flow and distribution** — how data moves through the GPU and how it is + distributed. + +**The coordinate graph is a hypergraph.** Nodes are `Dimension` variants (a size +and stride; or a `for`-loop index). Edges are *index transforms* — e.g. `Flatten` +takes several source indices and produces a row-major contiguous index — and an +edge connects a **tail set** of sources to a **head set** of destinations. **[S]** + +**The control graph** has `Operation` nodes (`LoadVGPR`, `Multiply`, `ForLoop`, +`If`), where control constructs **contain nested control graphs as their bodies**. + +Why this matters to us **[I]**: this is a first-class *layout/index algebra* held +as a graph, separate from control flow, in the same IR. It is the concrete +realization of the "layout algebra with named hardware axes" idea we noted from +the MLSys GPU book. Our stack expresses layout through attributes +(`#tile.layout`) hung on ops in a dialect tower; rocRoller makes the index +transform itself the graph and lowers by rewriting it. That is a genuinely +different point in the design space, and the hypergraph edge (many→many) is what +makes fusion/splitting of dimensions natural rather than encoded. + +It is worth being clear about the trade: one mutable graph lowered in place gives +up MLIR's verifiable stage boundaries and round-trippable textual dialects, which +are load-bearing for our lit-test discipline. **The idea to take is the +coordinate hypergraph as a representation, not the single-IR architecture.** **[I]** + +### 2.2 `Expression` and `EvaluationTime` — staging as a type + +Expressions are `std::variant` trees visited with `std::visit`, with +transformations in `ExpressionTransformations.hpp`. The key annotation **[S]**: + +> "they can only contain certain types of values at certain points. The +> `EvaluationTime` enum can be used to describe when a certain expression can be +> used." + +An explicit *when-is-this-knowable* tag on every expression. That is precisely +what Decision #28's symbolic-dim policy (`static | bucket | dynamic`) needs to be +enforceable rather than conventional — the same concept, and evidence it belongs +on the expression rather than on the op. **[I]** + +### 2.3 Observer-based scheduling — `peek` / `modify` / `observe` + +The single best structural idea in either codebase. `Context::schedule()` runs +every instruction through a `MetaObserver` composing many `IObserver`s **[S]**: + +| Method | Role | +|---|---| +| `peek` | *what would happen* if this were scheduled — stall cycles, errors — **without committing** | +| `modify` | mutate the instruction (e.g. attach a `WaitCount`) | +| `observe` | update machine state after scheduling (queue occupancy, hazard windows) | + +`WaitcntObserver` is just one observer, reading `GPUArchitecture` to decide waits. +Hazard handling, wait insertion, and cost accounting become **composable observers +over one instruction stream** instead of separate passes that must agree with each +other. + +Why this is the right shape for us **[I]**: `peek` is a **cost query before +commitment**, which is exactly the primitive a measured arbiter or a scheduling +search needs and which a pass pipeline cannot express. Our three AMD completion +mechanisms (waitcnt positions, TDM descriptors, SDMA fused packets — see the +gfx1250 reference §9.1) are three observers over one stream, not three passes. +This is the most directly actionable idea in this document. + +### 2.4 `Component` — a runtime plugin registry + +A `ComponentBase` declares an `Argument` type and a `Basename`. Concrete +components supply `Match(args)` (a predicate), `Build(args)` (a factory), and a +`Name`. `Component::Get(args)` searches registrations, caches instances for +non-single-use components, and is thread-safe via reader-writer locks. Used +throughout for architecture-specific codegen. **[S]** + +The distinguishing feature versus a plain virtual interface is that **selection is +a predicate over arguments**, not a type switch — so a new backend registers a +`Match` and never edits a dispatch site. Compare Decision #28's +`KernelEmitter`/`TargetPlugin` seam. **[I]** + +### 2.5 `Generator` — instruction streams as C++20 coroutines + +Lazy sequences via `co_yield`, modelling `std::ranges::input_range`, with +`map`/`filter`/`take`/`only`/`empty` and `.to()`. Movable, not +copyable; nothing executes until the first value is pulled. **[S]** + +Instruction generation is a *stream* rather than a materialized vector, which is +what lets `peek` interleave with generation. Mostly a C++ implementation choice, +but it is the mechanism that makes §2.3 ergonomic. **[I]** + +### 2.6 The IR serializes to YAML + +`Command`, `KernelGraph`, `Expression`, and `Operation` all implement +serialization traits (`toYAML` / `fromYAML` / `writeYAML`), *"particularly useful +for debugging, caching compiled kernels, and inspecting internal structures."* +**[S]** + +A serializable IR is a **cache key, a bug report, and a regression fixture in one +artifact**. Our AOT/compilation-cache and Evaluator work both want this; worth +checking what our Graph/Schedule/Tile IR can currently round-trip. **[I]** + +--- + +## 3. Composable Kernel — the P/Y/X/D coordinate framework + +CK ships roughly **400 KB of conceptual documentation** across ~28 files under +`docs/conceptual/ck_tile/`. It is the most thoroughly explained tile-programming +model in the AMD ecosystem, and the framework below is why. + +### 3.1 Four coordinate spaces, two transformations + +CK's thesis is that thread→data assignment should be a **composition of +well-defined mappings between named coordinate spaces**, resolved at compile +time. **[S]** + +| Space | Meaning | Coordinates | +|---|---|---| +| **P** | Thread Position — the hardware execution hierarchy | `thread_x`, `thread_y`, `warp_id`, `block_id` | +| **Y** | Local Data — the *algorithm's* view of its own work | `y0, y1, y2, y3` (algorithm-specific) | +| **X** | Global Position — coordinates in the problem domain | matrix row/col, image spatial coords | +| **D** | Memory Address — linearized, after padding/interleaving | a single offset | + +``` +P ─┐ + ├─► (P + Y → X) ─► X ─► (X → D) ─► D +Y ─┘ distribution layout + strategy optimization +``` + +The separation of responsibility is the payload **[S]**: + +- **`P + Y → X`** encodes the **distribution strategy** — how work is partitioned + across threads. Structuring this transform correctly is what *guarantees memory + coalescing*. +- **`X → D`** encodes **layout optimization** — padding, interleaving, address + space. Designing this is what *minimizes bank conflicts*. +- Architecture portability comes from **changing transform parameters while the + algorithm text stays fixed**. + +Why this matters to us **[I]**: Y-space is the one that does not appear in our +IR. We have global/tile coordinates and we have layouts, but the *algorithm's +local view of its own work*, as a coordinate space distinct from both thread +position and global position, is not named. It is what lets CK express an +algorithm "in its natural form" while the distribution stays a separate, +swappable object. Naming it would give the tile dialect a place to put the +per-thread iteration space that today is implicit in lowering. + +This framework is strictly more refined than rocRoller's coordinate graph (§2.1): +same underlying idea, but with the spaces *named* and each transformation +assigned a distinct optimization duty. + +### 3.2 The transform algebra — a closed operator set + +Transforms map between a **lower** (source) and **upper** (target) dimension +space, and are **bidirectional** **[V]**: + +- `calculate_lower_index()` — upper → lower (where to actually find the element) +- `calculate_upper_index()` — lower → upper (recover the original coordinate) +- `update_lower_index()` — **incremental** movement by a delta, without + recomputing from scratch + +The operator set is small and closed **[V]**: + +| Transform | Role | +|---|---| +| `MergeTransform` | multi-D → linear | +| `UnmergeTransform` | linear → multi-D | +| `EmbedTransform` | linear → multi-D strided | +| `ReplicateTransform` | 0-D → multi-D broadcast | +| `OffsetTransform` | translation | +| `PassThroughTransform` | identity | +| `PadTransform` | boundaries | +| `XorTransform` | **swizzle** | +| `SliceTransform`, `ModuloTransform` | sub-ranges, wrapping | + +Three observations **[I]**: + +1. **`XorTransform` is a first-class member of the algebra.** The bank-conflict + swizzle is not a special case bolted onto the LDS path — it is an index + transform like any other, so it composes and inverts. Compare §4.1, where the + same optimization is a hand-written parameter block. +2. **`ReplicateTransform` makes broadcast a coordinate operation**, which is how + a value shared across threads stays inside one framework rather than becoming + a separate mechanism. +3. **Everything is zero-copy and logical.** *"The actual tensor data remains + stored in memory in linear fashion, exactly as specified by the original tensor + shape and strides at creation time."* Transforms create views; they never move + data. That invariant is what makes composition safe. + +`update_lower_index()` deserves separate mention: incremental coordinate movement +is what makes sliding a tile window cheap, and CK gives it a dedicated 18 KB +document (`coordinate_movement.rst`). A tile IR that can only recompute absolute +coordinates will pay for it in every loop. + +### 3.3 `tile_distribution_encoding` — the distribution as compile-time data + +The whole thread↔data mapping is one declarative template parameter pack **[V]**: + +```cpp +using Encoding = tile_distribution_encoding< + sequence<>, // R — replication dims (none here) + tuple, // H — hierarchical lengths, X dim 0 (M) + sequence<4, 2, 8, 4>>, // hierarchical lengths, X dim 1 (N) + tuple, sequence<1, 2>>, // P → RH major + tuple, sequence<2, 2>>, // P → RH minor + sequence<1, 1, 2, 2>, // Y → RH major + sequence<0, 3, 0, 3>>; // Y → RH minor +``` + +`sequence<4, 2, 8, 4>` is a four-level decomposition reading directly onto +hardware: *four repetitions per thread, two warps per block, eight threads per +warp, four elements per vector op.* The `P → RH` and `Y → RH` maps are a **wiring +diagram** stating which hierarchy level of which X dimension each thread-position +coordinate and each local-data coordinate indexes. **[S]** + +The transferable property **[I]**: the distribution is a **value**, separable from +the kernel body, comparable, and enumerable. That is exactly the shape an +autotuner or a measured arbiter wants — a distribution becomes a search +coordinate rather than a code variant. + +**How the wiring actually works.** P and Y both index a single unified +**RH-space** — the concatenation of the R (replication) dims and the H +(hierarchy) groups — addressed by a `(major, minor)` pair **[V]**: + +- **major** — which RH group: `0` = R, `1..N` = the H group for X dimension *n* +- **minor** — which component within that group + +So `Ps2RHssMajor`/`Minor` and `Ys2RHsMajor`/`Minor` are literally a permutation +table: *"P coordinate i indexes component `minor[i]` of group `major[i]`."* The +whole distribution strategy is a wiring diagram over a hierarchical index space. + +The H hierarchy has a canonical four-level reading for GEMM **[S]**: + +```cpp +using HsLengthss = Tuple< + Sequence, // M + Sequence>; // N +// ^iterations ^warps ^threads ^vector width +// per thread in M per warp per access +``` + +R-dimensions exist for three stated purposes **[S]**: data reuse (the same input +feeding multiple output computations), reduction (several threads collaborating +on one result), and bandwidth reduction. + +**The encoding → transform chain is mechanical.** `make_ps_ys_to_xs_adaptor()` +builds a fixed three-stage chain straight from the encoding, with no per-kernel +hand-authoring **[V]**: + +``` +combine(P, Y) → ReplicateTransform (if R dims) → UnmergeTransform (into H dims) + → MergeTransform (into X dims) → X +``` + +That is the implementability result: **given the encoding, the coordinate +machinery is generated, not written.** A distribution search therefore explores +encodings, and the transform chain follows for free. **[I]** + +Per-thread storage is handled separately by a `ys_to_d_descriptor` — a plain +lengths/strides pair giving `offset = Σ y[i] * stride[i]`, chosen so that vector +loads land contiguously (e.g. layout `[M/VectorSize][N][VectorSize]`). Y→D is +where register-level layout is decided, distinct from the X→D global layout. **[V]** + +### 3.4 Pipeline = Problem + Policy + +A CK kernel is composed from **a pipeline, a tile partitioner, and an epilogue**, +and the pipeline itself splits in two **[S]**: + +| Component | Question it answers | +|---|---| +| **Problem** | *What* to compute — shapes, dtypes, the math (GEMM, conv) | +| **Policy** | *How* to move data — access patterns, hardware-specific choices | +| **Tile Partitioner** | problem dims (M, N, K) → workgroup tiles (kM, kN, kK) → grid | +| **Epilogue** | activation, bias, post-processing | + +Holding *what* and *how* apart at the type level is the same instinct as our +Graph IR / Schedule IR split, arrived at inside a template library. The useful +detail is that **Policy is a substitutable object**, so the same Problem can be +retargeted or retuned without touching the algorithm. + +Supporting vocabulary worth adopting for its precision **[S]**: **Tile Window** +(a viewport into a larger tensor defining the current tile's position and +bounds), **Block Tile** / **Wave Tile** (workgroup- and wave-granularity +sub-tiles), **Load Tile** / **Store Tile** (global↔LDS↔register transfers as +named operations). + +### 3.5 Intrawave vs interwave scheduling + +A clean statement of a real scheduling dichotomy for K-loop accumulation **[S]**: + +| | Mechanism | Best for | +|---|---|---| +| **Interwave** | K split into chunks; the *same* chunk loaded into every wave; all waves sync between chunks | **memory-bound** — coordinated accesses, optimized cache hit rate | +| **Intrawave** | full K loaded per wave; each wave runs independently, no sync; the CU interleaves | **compute-bound** — CU has scheduling freedom | + +Both ship, selected per workload. This is a *policy* choice in the §3.4 sense, +and it is the kind of discrete, nameable schedule axis our Schedule IR should be +able to carry rather than rediscover. **[I]** + +### 3.6 Coordinate movement — why tile iteration is cheap + +A coordinate in CK is not a position. It is the **materialized state of an entire +transform chain** **[V]**: + +```cpp +class TensorAdaptorCoordinate { + MultiIndex top_index_; // input position + MultiIndex bottom_index_; // output after all transforms + MultiIndex hidden_index_; // cached intermediate results +}; +``` + +`hidden_index_` — the cached intermediates of each stage — is what makes partial +recomputation possible. Movement then has a fast path **[V]**: + +```cpp +coord.top_index_ += step; +if (transformation_affects_movement(desc, step)) { + coord.hidden_index_ = desc.calculate_bottom_index(coord.top_index_); + coord.offset_ = desc.calculate_offset(coord.top_index_); +} else { + coord.offset_ += calculate_step_offset(desc, step); // single add +} +``` + +**If a step does not cross a transform boundary — no carry into a merged +dimension — the address update is one addition.** That is the entire reason +sliding a tile window is cheap, and it is a property of holding the chain state +rather than the coordinate alone. + +The consequence for us **[I]**: an IR whose only operation is "compute the +address at coordinate C" cannot express this. It needs a *move-by-delta* +operation over a coordinate that carries chain state, plus the ability to decide +statically whether a given delta is boundary-crossing. This is the concrete +mechanism behind the `update_lower_index()` note in §3.2, and it is the part most +likely to be missed when porting the coordinate algebra alone. + +### 3.7 Space-filling curves and a compile-time locality metric + +`SpaceFillingCurve` maps a 1-D access index to multi-dimensional coordinates, so +"traverse this tile well" becomes a linear loop **[S]**. Parameters: dimension +`Order` (row- vs column-major traversal), `scalars_per_access` (vector width per +dimension), and a **snake** flag that reverses direction on alternate rows/planes +to keep consecutive accesses spatially close. + +Stated best practices **[S]**: match traversal order to storage order; size the +vector as `min(fast_dim_length, cache_line_size / sizeof(T))`; enable snake for +large tensors; and `static_assert` that vector access aligns to cache lines. + +**The part worth taking is the analyzer.** CK ships a compile-time traversal- +quality metric that walks the curve, takes the Manhattan distance between +consecutive accesses, and bins them **[V]**: + +```cpp +const auto step = sfc.get_step_between(i, i + 1); +index_t distance = Σ |step[d]|; +if (distance <= 1) sequential_steps++; +else if (distance <= 16) cache_line_jumps++; // within a cache line +else large_jumps++; +``` + +This is a **hardware-free, static locality score for an access pattern**. That is +precisely the gap [`TILESIGHT_ASSESSMENT.md`](TILESIGHT_ASSESSMENT.md) recorded — +that the analytical cost model our arbiter falls back on was a mock. A +step-distance histogram over a materialized access order is cheap, needs no +device, and is a real signal. **[I]** + +### 3.8 Two worked swizzles — and what they prove about the algebra + +CK documents two independent bank-conflict swizzles. Both are built **entirely +from the §3.2 operator set** — no new primitives — which is the strongest +available evidence that the algebra is actually closed. **[V]** + +#### XOR preshuffle (`lds_index_swapping.rst`) + +Operates on a 3-D LDS coordinate `[K0, M, K1]`, where `K1 = KPack` (the thread's +vector width along K) and `K0 = KPerBlock/KPack`. Three stages: + +``` +1. XOR K0' = K0 XOR (M % (KPerBlock/KPack * MLdsLayer)) +2. Unmerge L = K0' / (KPerBlock/KPack) // MLdsLayer == 1 ⟹ L = 0 + K0'' = K0' % (KPerBlock/KPack) +3. Merge (L, M) → M' (K0'', K1) → K' +``` + +Stage 1 mixes M-dimension bits into the K0 index, redistributing accesses across +banks. `MLdsLayer` is the knob that lets **several rows share one bank set**, which +is what keeps small tiles from wasting bank capacity — stage 2 exists only to +carve that layer index back out. + +In code it is ordinary descriptor composition — `make_xor_transform` alongside a +`make_pass_through_transform`, with explicit lower/upper dimension index lists: + +```cpp +transform_tensor_descriptor( + BaseDescriptor{}, + make_tuple(make_xor_transform(Sequence{}), + make_pass_through_transform(Number{})), + Sequence<1, 0>{}, // XOR consumes dims [1,0] + Sequence<2>{}); // pass through dim 2 +``` + +Stated configuration heuristics **[S]**: `MLdsLayer` = 1 / 2 / 4 for tile sizes +≤32 / ≤64 / larger; `KPack` = 8 for fp16/bf16, 4 for fp32, 2 otherwise; with +`static_assert(TileSize % (MLdsLayer * KPack) == 0)`. + +#### Morton / Z-order (`swizzling_example.rst`) + +The result worth internalizing: + +> "MergeTransform enables Morton ordering by reordering and merging coordinate +> bits." + +```cpp +using SplitTransform = UnmergeTransform>; // coord → bits +using MortonMergeTransform = MergeTransform>; // bits → index +// merge computes: morton_idx = y₁·8 + x₁·4 + y₀·2 + x₀ +``` + +**Morton ordering is not a special function — it is a `Merge` over bit-split +dimensions taken in a permuted order.** Unmerge each coordinate down to +individual bits, then Merge them back in interleaved order, and bit-interleaving +falls out of the ordinary linearization arithmetic. The 4×4 tile layout it +produces: + +``` + 0 1 4 5 + 2 3 6 7 + 8 9 12 13 +10 11 14 15 +``` + +Why this matters for us **[I]**: it upgrades take-list item 2 from "add an XOR +transform" to something stronger — **a `Merge`/`Unmerge` pair that can address +individual bits covers an entire family of swizzles** (Z-order, and by +construction other bit-interleavings) with no new operators at all. If our layout +algebra gets bit-granular split/join, the swizzle family comes free; if it only +handles whole dimensions, every swizzle stays a special case. + +#### A second static metric — and this one is assertable + +Both documents ship the same analyzer shape **[V]**: + +```cpp +for (tid = 0; tid < WarpSize; ++tid) { + offset = desc.calculate_offset(coords_for(tid)); + bank = (offset * sizeof(T) / BankWidth) % NumBanks; + bank_access[bank]++; +} +max_conflict = max over banks; // "N-way bank conflict" +``` + +The Morton document runs it **comparatively** — row-major versus Morton — turning +"is this layout better" into a number computed from the descriptor alone. + +This is a companion to §3.7's locality histogram, and it is the more immediately +useful of the two for us because it is **a property a unit test can assert**: +*this descriptor is conflict-free for a warp-wide access on N banks.* No device, +no fixture, no measurement — just the descriptor and the bank count. That is +exactly the kind of gate our LDS-layout work currently lacks. **[I]** + +### 3.9 The user-facing surface — `TileWindow` and `sweep_tile` + +These two APIs are where the framework pays off, and reading them settles what +the encoding is actually *for*. + +**`TileWindow` separates "which data is mine" from "how to fetch it."** **[S]** + +> "While TileDistribution solves the problem of work assignment by mapping +> threads to tensor coordinates, it does not address *how* threads access the +> data at those coordinates. TileWindow serves as the critical bridge." + +```cpp +template +struct tile_window_with_static_distribution { + TensorView tensor_view_; + Distribution distribution_; + array origin_; // runtime + static constexpr auto window_lengths = ...; // compile-time +}; +``` + +Note the static/dynamic split: **window lengths are compile-time, the origin is +runtime.** That is exactly the `static | bucket | dynamic` distinction Decision +#28 requires, landed on the natural seam — shape is static, position is not. **[I]** + +**`LoadStoreTraits` derives the access pattern from the distribution.** It is a +compile-time engine performing three analyses **[S]**: + +1. **Vector dimension identification** — which Y dimension has stride 1 +2. **Access pattern calculation** — how many memory operations, in what order +3. **Space-filling curve construction** — the traversal order itself + +So the SFC of §3.7 is **not hand-picked — it is computed from the encoding.** +`scalar_per_vector` likewise. This is the payoff of distribution-as-a-value: once +the encoding exists, vector width, access count, and traversal order all follow. + +The load loop makes the entire chain concrete **[V]**: + +```cpp +static_for<0, Traits::num_access, 1>{}([&](auto i_access) { + const auto y_indices = Traits::get_y_indices(i_access); // ← SFC + const auto x_indices = distribution_.calculate_x_from_y(y_indices); // ← P+Y→X + const auto global_indices = add_arrays(origin_, x_indices); + if constexpr (Traits::scalar_per_vector > 1) { /* vector load */ } + else { /* scalar load */ } +}); +``` + +`SFC → Y → (P+Y→X) → +origin → global`, with vectorization as a `constexpr` +branch rather than a runtime one. Window movement is `set_window_origin(...)`, +O(1) on the precomputed coordinates from §3.6. + +**`sweep_tile` is the iteration surface** — "load once, use many times": load the +X data into registers once, then sweep Y positions while X stays resident. +Implemented as compile-time recursive `static_for` over Y-space lengths, so it +unrolls with zero runtime overhead. Named use cases: matmul (reuse A columns), +convolution (reuse filter weights), reduction (accumulate over Y), broadcast +(apply X across all Y). **[S]** + +CK states the layering as a four-line contract **[S]**: + +> 1. **TileDistribution**: "Here's how to divide work" +> 2. **TileWindow**: "Here's the data, loaded efficiently" +> 3. **Sweep operations**: "Here's how to process every element" +> 4. **User code**: "Thanks! *does computation*" + +**Why this is the important section for us** **[I]**: it is a clean separation of +four concerns our stack partly conflates — + +| Concern | CK owns it in | Derived or written? | +|---|---|---| +| which data is mine | `TileDistribution` encoding | **written** (the tunable) | +| how to fetch it | `TileWindow` + `LoadStoreTraits` | **derived** | +| what order to traverse | space-filling curve | **derived** | +| what to compute | the `sweep_tile` lambda | **written** (the algorithm) | + +Only two of the four are authored. Vectorization, coalescing, access count, and +traversal order are *consequences* of the encoding, not independent knobs a +kernel author sets. That is the strongest argument in this survey for +distribution-as-a-value (take-list item 11): it is not merely an autotuning +convenience, it is what makes everything else derivable. + +--- + +## 4. hipBLASLt / TensileLite — the arbiter, in production + +This is the closest existing thing to Decision #28's measured arbiter, and it is +richer than the "Tensile generates kernels" summary suggests. + +### 4.1 Five pathways, not one + +The runtime flow is `hipblaslt.cpp` → `rocblaslt_mat.cpp` → `tensile_host.cpp` → +TensileLite host → a lazily-loaded `.hsaco`/`.co` from a *device library* +directory. But several kernel sources coexist **[S]**: + +| Pathway | Where | Status | +|---|---|---| +| **TensileLite** — generated assembly kernels | `tensilelite/` | primary; "what new work uses" | +| **rocRoller** custom kernels | `library/src/amd_detail/rocblaslt/src/rocroller/` | ships alongside, gated by `HIPBLASLT_ENABLE_ROCROLLER`, **ON by default** | +| **ExtOps** (softmax, layernorm, AMax) | `device-library/extops/`, generated by `SoftmaxGenerator.py` etc. | separate device library | +| **Matrix transform** | `device-library/matrix-transform/`, `rocblaslt_transform.cpp` | separate op family | +| **User-driven tuning** | `UserDrivenTuningParser.cpp` | runtime override of selection | + +So a production GEMM library runs **a generated-kernel tier, an alternate-generator +tier, and a user-override tier simultaneously**, with a build flag to drop one. That +is Decision #28's three-tier model observed in the wild, plus a fourth tier we do +not currently model: explicit user override at runtime. **[I]** + +### 4.2 Selection is a composable tree of single-concern nodes + +The design statement, from `SolutionLibrary.hpp` **[V]**: + +> A complete SolutionLibrary is a **tree of objects which each handles a single +> aspect of selecting a solution** for a given problem. Each node in the tree will +> handle an aspect such as: +> - Compatibility with a particular model of GPU +> - Selecting kernels that solve a particular type of problem (transpose, data +> type, etc.) +> - Selecting the fastest kernel based on benchmark results or other logic +> - Ensuring that a problem is compatible with any assumptions made by a +> particular kernel (e.g. size or stride requirements) + +with the documented example composition: + +``` +MasterSolutionLibrary (serialization) + └─ GPU selection + └─ Problem type selection + └─ Predicated logic for specific sizes + └─ Matching library based on benchmarks + └─ Individual kernels +``` + +The node types available are effectively a **taxonomy of selection strategies**, +each a `SolutionLibrary` subclass **[V]**: + +| Node | Strategy | +|---|---| +| `MasterSolutionLibrary` | root; owns serialization | +| `MapLibrary` | dispatch by key | +| `ExactLogicLibrary` | exact match on a predicate set | +| `ProblemMatchingLibrary` | *"Uses a distance function to select solutions based on benchmarks… At runtime, we find the benchmarked size that is closest to the size asked for."* | +| `GranularitySelectionLibrary` | *"Compares the tile sizes of each kernel, the dimensions of the problem, and the number of compute units… to select a kernel that fits the best on the GPU with the lowest amount of waste"* | +| `FreeSizeLibrary` | free-size problems | +| `PredictionLibrary` / `MLPClassificationLibrary` | learned models (§4.5) | +| `PlaceholderLibrary` | lazy code-object loading | +| `CachingLibrary` | memoized selection | +| `EmbeddedLibrary` | compiled-in library data | + +This is the structural answer to "how do you combine hand-tuned, generated, and +predicted kernels without a tangle of special cases": **you don't build one +arbiter, you build a tree whose nodes each decide one thing.** Analytical +selection, benchmark lookup, and a learned model are sibling node types, not +competing designs. **[I]** + +### 4.3 The interface has the shape an arbiter needs + +```cpp +virtual std::shared_ptr + findBestSolution(MyProblem const&, Hardware const&, double* fitness = nullptr) const = 0; + +virtual SolutionSet + findAllSolutions(MyProblem const&, Hardware const&, SolutionLibrarySearchType) const = 0; + +virtual SolutionVector + findTopSolutions(MyProblem const&, Hardware const&, int numSolutions) const; +``` + +Three details worth copying **[V]**: + +- **`fitness` is an out-parameter of "find best."** Selection returns not only a + choice but a *score for that choice* — which is what lets a caller decide + whether to trust it, and what an accuracy-budgeted arbiter needs. +- **`findTopSolutions(N)` is first-class**, alongside `findAllSolutions` and + `getSolutionByIndex`. An arbiter that can only ask "what's best" cannot + benchmark candidates; top-N is the primitive that makes measurement possible. +- **Search strictness is an enum**, not a boolean: `DEFAULT` (full predicates), + `GEMM_TYPE_ONLY` (dtype/transpose/grouped/mx-block match only), `HARDWARE_ONLY` + (accept everything). Progressive relaxation is built into the query. + +Grouped-GEMM gets parallel overloads throughout (`findAllSolutionsGroupedGemm`, +`findTopSolutionsGroupedGemm`), i.e. "a batch of problems selecting one solution" +is a modelled case, not an afterthought. + +### 4.4 A pluggable distance-metric set + +`Distance.hpp` provides, as interchangeable policies **[V]**: `Equality`, +`Range`, `RatioDistance`, `ManhattanDistance`, `EuclideanDistance`, +**`JSDivergence`**, `RandomDistance`, `GridBasedDistance`. + +`RatioDistance` is the interesting default-case metric for GEMM — problem sizes +matter multiplicatively, not additively, so nearest-neighbour in log space beats +Euclidean. `RandomDistance` exists for exploration/testing. That a +Jensen–Shannon divergence is on the list at all indicates how much of this is +treated as a genuine statistical matching problem. **[I]** + +### 4.5 The learned path — a residual MLP over *derived* features + +`MLPClassification.hpp` **[V]**: + +> Neural net used to **estimate efficiency values for solutions in the library**. + +Structure: `StandardScaler` (mean/scale normalisation) → `DenseLayer`s → +`ResBlock`s, in `float`. Note it predicts a **per-solution efficiency**, not a +kernel-id class — so its output composes with the rest of the tree as a score +rather than replacing selection. + +The feature set (`MLFeatures.hpp`) is the transferable part **[V]**: + +| Feature | What it is | +|---|---| +| `FreeSizeA`, `FreeSizeB` | M, N | +| `BatchSize`, `BoundSize` | batch, K | +| `Tile0Granularity`, `Tile1Granularity` | `1/mt0`, `1/mt1` | +| `CUGranularity` | how well the tile grid fills the CUs | +| `WavesPerSIMD` | occupancy | + +Only the first four are problem dimensions. The rest are **derived +occupancy/granularity ratios** — quantities that plausibly *cause* performance +rather than merely correlate with it. Any model we build, learned or analytical, +should be fed granularity and occupancy terms rather than raw shapes. **[I]** + +### 4.6 Reject-and-continue, with a one-time diagnostic + +A worked example of a hardware quirk handled inside the selection tree **[V]**: + +> `streamKDynamicQueueSupported()` excludes StreamK dynamic-queue / work-stealing +> solutions (SK4 and the dynamic sub-path of SK5) on devices whose **XCD count is +> not a power of two**, warning the user once. This is **reject-and-continue**: +> selection falls through to another (SK3-static / non-StreamK) solution for the +> GEMM. + +Three things at once: a hardware-shape predicate (XCD count parity), graceful +degradation to a different algorithm rather than failure, and a **warn-once** +diagnostic so the fallback is visible without being noisy. That combination is a +good template for our capability gates, which today tend to be binary +supported/unsupported. **[I]** + +--- + +## 5. Concrete algorithms worth taking + +### 5.1 LDS swizzle for bank-conflict elimination + +The most immediately usable algorithm in the survey, and it lands directly on the +ROCm audit's open "bounds-aware swizzle" item. **[V]** + +**Hardware model (GFX950).** LDS has 64 banks × 4 B = 256 B per bank row = exactly +16 `dwordx4` columns. Other architectures have 32 banks → 128 B rows → 8 columns. +When a tile row's K spans fewer than the columns-per-bank-row, multiple tile rows +pack into one bank row, and reads from different rows at the same column offset +collide. + +**Parameters** (`LDSSwizzleParams` in `LowerTile.cpp`): + +``` +numColumns = tileK / (128 / elementBits) // dwordx4 chunks per tile row +columnsPerBankRow = numBanks / 4 // 16 on GFX950, 8 on 32-bank parts +rowsPerBankRow = columnsPerBankRow / numColumns +bankRowIdx = row / rowsPerBankRow +``` + +| dtype / macK | cols/row × rows/bank-row | +|---|---| +| FP4 macK=128 | 4 × 4 | +| FP4 macK=256 | 8 × 2 | +| FP8 macK=128 | 8 × 2 | +| FP16 macK=128 | 16 × 1 — **no conflict, swizzle skipped** | +| FP16 macK=64 | 8 × 2 | +| FP16 macK=32 | 4 × 4 | + +**The transform:** column-level pair-swap plus circular rotation. *Only column +indices are permuted; row assignments are unchanged.* Skipped entirely when +`numColumns >= columnsPerBankRow` (`LDSSwizzleParams::noConflicts()`). + +**The non-obvious hardware detail** — and the reason a naive swizzle is wrong: + +> "The LDS unit processes a `ds_read_b128` from a 64-lane wave in **4 phases, +> each executing 16 threads** simultaneously. The 16 threads in a phase access LDS +> in parallel, so **bank conflicts only occur between threads within the same +> phase**." + +And the phases are **not contiguous lane ranges**: + +| Phase | Threads | +|---|---| +| 0 | T0-3, T12-15, T20-23, T24-27 | +| 1 | T32-35, T44-47, T52-55, T56-59 | +| 2 | T4-7, T8-11, T16-19, T28-31 | + +Any conflict analysis that assumes lanes 0–15 form a phase will compute the wrong +answer. **[V]** + +Note this table is GFX950/wave64; gfx1151 is wave32 with a different bank count, +so the *method* transfers but every constant must be re-derived per target. **[I]** + +### 5.2 Workgroup mapping for cache locality + +> "Workgroup mapping parameters specify how rocRoller maps the GPU workgroup +> number (hardware) to tile numbers (software). Workgroup mapping is done to +> increase cache efficiency." **[S]** + +Mechanically: the hardware workgroup number is a pre-populated SGPR, exposed in +the coordinate graph as a `Workgroup` node, attached to dangling +`MacroTileNumber` leaves during the `ConnectWorkgroups` transformation — and the +mapping policy is applied **in that same pass**. + +The idea worth taking is the *placement* **[I]**: grid→tile remapping is a +**coordinate-graph rewrite**, not a special case buried in the launcher. Our +existing `head_first_xcd` swizzle is the same class of optimization implemented +ad hoc; making it a layout/index transform would let the cost model see it. + +### 5.3 Register placeholders — a codegen anti-pattern + +`Register::Value::Placeholder` forces a specific physical register. It is +*required* when a value must persist in the same register across loop iterations. +Used unnecessarily, it defeats an optimization inside `Expression::generate()` +that would otherwise assign the destination directly — costing an extra register +and an extra `v_mov` per use. The optimization only fires when `nullptr` is passed +to `generate()`. **[S]** + +The general lesson **[I]**: *pinning is a constraint, and constraints must be +justified per use.* A codegen API that makes pinning the ergonomic default will +silently inflate register pressure — the exact quantity `rocm_tiling.py` treats as +the dominant tiling lever. + +--- + +## 6. Ranked take / skip for Tessera + +**Take — high value, low coupling:** + +1. **Observer scheduling (`peek`/`modify`/`observe`)** — §2.3. Gives a + before-commitment cost query and unifies our three AMD completion mechanisms + as observers rather than passes. Biggest single *mechanism* here. +2. **Bit-granular `Merge`/`Unmerge`, plus `XorTransform`** — §3.2, §3.8. Swizzle + stops being a special case: XOR preshuffle is one transform in a chain, and + Morton/Z-order is *just* a `Merge` over bit-split dimensions in permuted + order. Bit-granular split/join buys the whole swizzle family with no new + operators. +3. **The static bank-conflict analyzer** — §3.8. Computes N-way conflict from a + descriptor alone, so "this layout is conflict-free for a warp-wide access" is + a **unit-testable assertion** with no device. The cheapest correctness gate we + could add to LDS-layout work. +4. **The LDS swizzle algorithm** — §5.1. Directly serves the open bounds-aware + swizzle item; method transfers, constants must be re-derived per target. +5. **Fail-loud on partial metadata** — §1.6. Cheap, and it converts a class of + silent mis-optimization into a diagnostic. +6. **`EvaluationTime` on expressions** — §2.2. Makes the symbolic-dim policy + enforceable instead of conventional. +7. **Optimization remarks / loop health** — §1.7. Structured "why is this + schedule bad" output that needs no hardware. +8. **Intrawave/interwave as a named schedule axis** — §3.5. A discrete, + documented policy choice our Schedule IR should carry explicitly. +9. **The compile-time locality metric** — §3.7. A step-distance histogram over a + materialized access order is a hardware-free static cost signal, and it lands + squarely on the mock-analytical-cost-model finding in + [`TILESIGHT_ASSESSMENT.md`](TILESIGHT_ASSESSMENT.md). Cheapest real answer to + that gap. + +10. **Selection as a composable tree of single-concern nodes** — §4.2. The + structural answer to combining hand-tuned, generated and predicted kernels: + not one arbiter, but a tree where analytical selection, benchmark lookup and + a learned model are sibling node types. +11. **`fitness` out-param and `findTopSolutions(N)`** — §4.3. An arbiter that + only answers "what's best" cannot measure; top-N plus a confidence score are + the primitives that make an accuracy budget enforceable. +12. **Derived granularity features, not raw shapes** — §4.5. Tile/CU granularity + and waves-per-SIMD over M/N/K, for any cost model we build, learned or + analytical. +13. **Reject-and-continue with warn-once** — §4.6. Capability gates that degrade + to another algorithm and say so once, instead of binary supported/unsupported. + +**Take — larger, worth designing toward:** + +14. **Name Y-space** — §3.1. The algorithm's local view of its own work is the one + coordinate space our IR does not have, and it is what lets the distribution + become a separate swappable object. Highest-leverage *conceptual* item in this + document. +15. **Distribution-as-a-value** — §3.3, §3.9. `tile_distribution_encoding` makes + the thread↔data mapping comparable and enumerable — a search coordinate for + the autotuner rather than a code variant. §3.9 raises the stakes: vector + width, access count, and traversal order are all **derived** from the + encoding, so this is what makes the rest of the machinery generated rather + than authored. +16. **Derive the access pattern, don't author it** — §3.9. Only *which data is + mine* and *what to compute* are written; how to fetch it and in what order + are consequences. Worth testing our tile lowering against: how many of those + four are currently hand-specified? +17. **Static shape / dynamic origin as the symbolic-dim seam** — §3.9. CK puts + window *lengths* in the type and the *origin* in a runtime field. That is a + clean, load-bearing instance of Decision #28's `static | bucket | dynamic` + policy. +18. **A closed transform operator set with bidirectional + incremental ops** — + §3.2. Especially `update_lower_index()`: a tile IR that can only recompute + absolute coordinates pays for it in every loop. +19. **The coordinate hypergraph as a representation** — §2.1. Take the index + algebra; do *not* take the single-mutable-IR architecture, which trades away + the verifiable stage boundaries our lit discipline depends on. +20. **Serializable IR as cache key + fixture** — §2.6. +21. **"New arch = data only, no source edits"** as an explicit acceptance test for + the backend-plugin seam — §1.3. + +**Skip / already have:** + +- Two IR levels (§1.1) and an opt-style driver with FileCheck tests (§1.8) — we + have both; note them as convergence evidence, not work. +- Problem/Policy split (§3.4) — our Graph IR / Schedule IR split is the same + instinct; the only delta is making Policy substitutable as a value. +- `Component` registry (§2.4) — our capability/pipeline registries already cover + this; the only delta is predicate-based selection. +- `Generator` coroutines (§2.5) — a C++ ergonomics choice, not an architecture. +- Declarative peephole DSL (§1.5) — attractive, but MLIR PDL/rewrite patterns + already occupy this slot for us. +- CK's template-metaprogramming implementation strategy — the *framework* is the + idea; C++ compile-time specialization is not a model we should copy into an + MLIR-based compiler. + +--- + +## 7. Not read + +**Within these three projects.** StinkyTofu: DAG scheduler internals +(`src/transforms/asm/dag/` — `ReadyQueue.hpp`, and a `CDNA5.hpp` implying a +generation beyond gfx1250), `AsmVerifierPass`, `adding-architecture.md`, the +`python_module` bindings. rocRoller: `CoordinateGraph/` docs, `lib/` source, +`GPUArchitectureGenerator`. + +Composable Kernel is the largest partial read — roughly **17 of its ~28 +conceptual documents remain unread**. Still on-topic: + +| Doc | Why it matters | +|---|---| +| `static_distributed_tensor.rst` (16 KB), `load_store_traits.rst` (16 KB) | the zero-overhead implementation strategy | +| `buffer_views.rst` (23 KB), `tensor_views.rst` (16 KB), `descriptors.rst`, `adaptors.rst` | the layered view stack under the algebra | +| `thread_mapping.rst` (18 KB), `hardware/` | logical→physical thread mapping | +| `coordinate_systems.rst` (20 KB), `tensor_coordinates.rst` (14 KB) | the P/Y/X/D treatment in full | +| `convolution_example.rst` (20 KB) | the framework applied to a non-GEMM op | + +Read in this survey: `introduction_motivation`, `transforms`, `tile_distribution`, +`encoding_internals`, `coordinate_movement`, `space_filling_curve`, +`swizzling_example`, `lds_index_swapping`, `tile_window`, `sweep_tile`, +`CK-Tile-intra-inter-wave`, `TERMINOLOGY` — 12 of ~28. + +Also unread in CK: `tile_engine/`, `codegen/`, `experimental/`, and the actual +`include/ck_tile/` headers — this survey read the *documentation*, not the +implementation. + +**Beyond these three**, still unread: + +1. **hipBLASLt kernel *generation*** — §4 read the selection side only. + `KernelWriter.py` / `KernelWriterAssembly.py`, the `rocisa` Nanobind assembly + module, `Components/` (modular MAC / global-read / scheduling blocks), and the + three-phase `BenchmarkProblems → LibraryLogic → ClientWriter` tuning pipeline + are all unread. `ContractionProblemPredicates.hpp` alone is 119 KB. +2. **`ExactLogicLibrary` / `MapLibrary` / `CachingLibrary` internals** — §4.2 + names them from their headers; only `MatchingLibrary` and + `GranularitySelectionLibrary` were read in any detail. +3. **rocWMMA** — cooperative-matrix fragments in portable form. +4. **AITER** — surveyed at brief depth in the ecosystem doc, never at source + level. + +--- + +## 8. Sources + +All `ROCm/rocm-libraries@develop`. + +**StinkyTofu** (`shared/stinkytofu/`) — `docs/developer/architecture.md`, +`wmma-vgpr-reorder-pass.md`, `pattern-grammar.md`, `cluster-barrier.md`; +`docs/user/stinky-waitcnt-insertion-pass.md`; `hardware/src/gfx/Gfx1250/`; +`include/stinkytofu/core/Types.hpp`. + +**rocRoller** (`shared/rocroller/`) — `CLAUDE.md`; `docs/src/DesignOverview.md`, +`LDSSwizzling.md`, `RegisterPlaceholders.md`, `WorkgroupMapping.rst`. + +**Composable Kernel** (`projects/composablekernel/`) — `TERMINOLOGY.md`; +`docs/conceptual/CK-Tile-intra-inter-wave.rst`; +`docs/conceptual/ck_tile/{introduction_motivation,transforms,tile_distribution, +encoding_internals,coordinate_movement,space_filling_curve,swizzling_example, +lds_index_swapping,tile_window,sweep_tile}.rst`. + +**hipBLASLt / TensileLite** (`projects/hipblaslt/`) — `CLAUDE.md`, +`tensilelite/CLAUDE.md`; `tensilelite/include/Tensile/{SolutionLibrary,Distance, +MatchingLibrary,GranularitySelectionLibrary,MLPClassification,MLFeatures}.hpp`. + +**LLVM** (`ROCm/llvm-project@amd-staging`) — `llvm/lib/Target/AMDGPU/SISchedule.td` +(the gfx1250/gfx1251 machine model, recorded in the gfx1250 target reference §2.5 +rather than here). + +--- + +*See also: +[`../backend/rocm/GFX1250_MI450_COMPILER_REFERENCE.md`](../backend/rocm/GFX1250_MI450_COMPILER_REFERENCE.md) +(target facts) · +[`../backend/rocm/ROCM_PATTERNS_FROM_AMD_ECOSYSTEM.md`](../backend/rocm/ROCM_PATTERNS_FROM_AMD_ECOSYSTEM.md) +(ecosystem survey) · +[`COMPILER_THEORY_OF_OPERATION.md`](COMPILER_THEORY_OF_OPERATION.md) (our +direction) · [`TILESIGHT_ASSESSMENT.md`](TILESIGHT_ASSESSMENT.md) (external +assessment, same reference role).* From e8b60d54f7ea891bc72fbdf6726c405f080ca109 Mon Sep 17 00:00:00 2001 From: Greg Stoner Date: Wed, 29 Jul 2026 07:53:01 -0600 Subject: [PATCH 2/3] Survey rocisa, rocFFT, rocPRIM, rocRAND, rocALUTION; re-read rocWMMA MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Extends the two survey documents with six more ROCm projects. Documentation only. Compiler survey gains section 4.7 on rocisa, TensileLite's nanobind assembly generator — the same Python-driving-C++ shape we have. Three findings worth copying: IR nodes carry a mandatory clone() deep-copy contract; exporting a vector to Python is a copy, so elements are mutable through their shared_ptr but cannot be assigned or replaced; and import raises if any C++ source is newer than the built extension. That last one is added to the take list — we lost time this session to a tessera-opt binary that silently did not match its sources. Patterns doc gains four project briefs and a rocWMMA re-read: rocFFT has the best cache design in the ecosystem. The kernel name is the cache key, with every differentiating parameter encoded into it, so profiler output and cache identity are the same string and the cache needs no schema update when a new parameter appears. Three further key fields guard staleness — architecture, HIP version, and generator version. A read-only system cache ships with the library alongside a read-write user cache, the shipped one populated at build time by a helper that shares the generator but is not installed. AOT and JIT are one path with a policy knob rather than two lanes. Also records that hipRTC holds process-wide locks, so parallel compilation needs a helper process. rocPRIM turns tuning output into generated headers, and its fallback_config is a typed fallback ladder: an untuned type inherits the config of a representative matched on size range and floating-pointness. That is dtype bucketing, the same move Decision #28 makes for shapes. rocRAND is the one with a direct bearing on us. Under dynamic ordering it picks launch geometry per device, and AMD states plainly that the number of generators and the sequence of generated numbers can vary as a result. So reproducibility versus performance is a named opt-in mode, not an emergent property. Worth confirming the same holds for Decision #18: if a tuned launch configuration ever fed an RNG offset scheme, autotuning would silently change numerical output. rocALUTION is included as a contrast, not a pattern. It selects execution location at run time via RTTI and silently migrates an object back to the host when a routine is unavailable on the accelerator. That is the opposite of Decision #21, which requires a diagnostic naming the op and target. Both are defensible for their audience; the contrast is worth recording because silent host migration is how a performance cliff hides. rocWMMA re-read adds that collaborative fragments are a movement concept and are explicitly unsupported in MMA functions, that partial and oversized tiles became the library's problem in 2.0.0, and that the wavefront-centric contract is undefined behaviour rather than a hint. Co-Authored-By: Claude Opus 5 --- .../rocm/ROCM_PATTERNS_FROM_AMD_ECOSYSTEM.md | 120 +++++++++++++++++- .../compiler/AMD_KERNEL_COMPILER_SURVEY.md | 62 +++++++-- 2 files changed, 167 insertions(+), 15 deletions(-) diff --git a/docs/audit/backend/rocm/ROCM_PATTERNS_FROM_AMD_ECOSYSTEM.md b/docs/audit/backend/rocm/ROCM_PATTERNS_FROM_AMD_ECOSYSTEM.md index f002538c8..f782cfd74 100644 --- a/docs/audit/backend/rocm/ROCM_PATTERNS_FROM_AMD_ECOSYSTEM.md +++ b/docs/audit/backend/rocm/ROCM_PATTERNS_FROM_AMD_ECOSYSTEM.md @@ -1,11 +1,11 @@ --- -last_updated: 2026-07-13 +last_updated: 2026-07-28 audit_role: reference --- # ROCm Backend — Patterns from the AMD ROCm Ecosystem -> **Purpose.** A source-grounded survey of eight production AMD ROCm projects, read for +> **Purpose.** A source-grounded survey of twelve production AMD ROCm projects, read for > *transferable patterns* that improve Tessera's ROCm backend and the compiler at large. > This is an **ideas + design-vocabulary** document, not a status surface — for status see > [`ROCM_AUDIT.md`](ROCM_AUDIT.md); for the hardware bring-up ladder see @@ -24,7 +24,8 @@ audit_role: reference > [`../../compiler/AMD_KERNEL_COMPILER_SURVEY.md`](../../compiler/AMD_KERNEL_COMPILER_SURVEY.md). > > Surveyed: **AITER**, **ATOM**, **hipBLASLt**, **rocWMMA**, **Mori**, **Iris**, **XIO**, -> and the **AMD Gluon GEMM tutorial**. +> and the **AMD Gluon GEMM tutorial** (2026-06-18); plus **rocFFT**, **rocPRIM**, +> **rocRAND**, and **rocALUTION**, with rocWMMA re-read (2026-07-28). > > _Researched 2026-06-18. Provenance is flagged per claim: **[V]** verified from > repo/docs/paper source, **[D]** DeepWiki summary (high-confidence, not source-exact), @@ -58,7 +59,7 @@ original “artifact-only” frontier is obsolete. The patterns below are split into **(A) compiler objects that must reach a consumer**, **(B) measured kernel redesigns**, and **(C) the distributed/GPU-initiated-comm track**. -The throughline across all eight projects: **AMD's stack has converged on "make the +The throughline across all twelve projects: **AMD's stack has converged on "make the hardware concept a first-class object"** — layouts, fragments, symmetric heaps, epilogues, tuned-config rows. That is Tessera's founding thesis. Most of what follows is not new architecture; it is validation of the IR design plus a concrete vocabulary (exact enums, @@ -156,6 +157,18 @@ Header-only C++ template lib modeled on `nvcuda::wmma`; lowers to `amdgcn_mfma_* - **One code path for MFMA vs WMMA.** Same `fragment` + `mma_sync(d,a,b,c)` transparently lowers per arch; wave width auto (64 CDNA / 32 RDNA). API: `fill_fragment`, `load_matrix_sync`, `store_matrix_sync`, `load_matrix_coop_sync`. **[V]** +- **Collaborative fragments are a *movement* concept, not a compute one** (re-read + 2026-07-28). Their data is distributed across participating waves in the same thread + block, to balance shared responsibility for collective transfers such as global→LDS — + and they are **explicitly not supported in MMA functions**. Wave collaboration is + expressed as fragment-scheduler metadata supplied by the developer, so the collaboration + policy travels with the *type* rather than the call site. **[V]** +- **Partial and oversized tiles became the library's problem, not the caller's** — + from rocWMMA 2.0.0 (ROCm 7.0), `mma_sync` handles them piece-wise automatically, which is + why the performance samples now use larger tiles with simpler code. **[V]** +- **Wavefront-centric contract:** all threads in a wavefront must be active or behaviour is + undefined; small edge-case blocks are auto-padded rather than thread-masked. That is a + precondition our Tile IR would have to guarantee, not merely hope for. **[I]** ### 1.5 Mori — modular RDMA framework ([repo](https://github.com/ROCm/mori)) @@ -253,6 +266,105 @@ choices*, not micro-opts. The "obvious" double-buffer regressed badly; slicing t tile to fit the VGPR budget was the real lever. Tiling and register allocation are coupled and must be co-designed. +### 1.9 rocFFT — runtime compilation, and the best cache design in the ecosystem ([repo](https://github.com/ROCm/rocm-libraries/tree/develop/projects/rocfft)) + +Read 2026-07-28 for its `designdocs/`, which are unusually explicit about *why*. The RTC +design is the most directly reusable thing in this section. **[V]** + +- **The motivation is variant explosion, and it is quantified.** Stockham kernels need a + variant per `{arch} × {6 interleaved/planar × in-place/out-of-place} × {fwd,inv} × + {precisions} × {unit,non-unit stride} × {callbacks on/off}`. Beyond build time, this hits + a hard wall: the default `-fPIC` memory model caps shared objects at **2 GiB**, and they + were hitting build breaks. **[V]** +- **The kernel *name* is the cache key.** Every differentiating parameter is encoded into + the function name — scheme, length, placement, direction, formats, precision, stride type, + twiddle base, callback type. Two consequences AMD calls out explicitly: profilers and logs + then name exactly which kernel ran even when it was runtime-compiled, and *"the caching + code needn't be aware of all the possible parameters that kernels could differ by. New + parameters can be added at any time, and as long as the kernel names are updated + accordingly, the cache will just work."* **[V]** +- **Three more key fields guard staleness:** GPU architecture, HIP runtime version, and + **kernel generator version**. The last is the one that matters for a compiler — it + invalidates the cache when codegen changes. **[V]** +- **Two-tier cache.** A read-only system cache shipped with the library + (`ROCFFT_RTC_SYS_CACHE_PATH`, located relative to the shared object) plus a read-write + user cache (`ROCFFT_RTC_CACHE_PATH`). Static linking breaks the relative lookup, so the + env var becomes required — but the user cache still works, so behaviour degrades rather + than fails. **[V]** +- **The shipped cache is populated at build time by a helper executable** that shares the + generator and RTC code but is *not* installed, so no symbols are exported merely to serve + the build. *"The criteria for which kernels to pre-build can be arbitrary. Less common + choices will be runtime-compiled, and runtime compilation is still a fallback."* This is + AOT and JIT as one path with a policy knob, not two lanes. **[V]** +- **Serialize/deserialize APIs exist for diskless and distributed runs** — rank 0 populates + the cache, broadcasts the buffer, other ranks deserialize in-memory. **[V]** +- **A concrete gotcha worth knowing before we parallelize compilation:** hipRTC holds + process-wide locks, so multithreaded compilation does not help. rocFFT spawns a *helper + process* when a compile is already in flight, falling back to in-process if the helper + cannot be found. **[V]** + +**[I]** This maps almost one-for-one onto our compilation cache and the Apple AOT lane. The +name-as-key idea is the cheapest and most valuable: it makes the cache schema-free and makes +profiler output and cache identity the same string. + +### 1.10 rocPRIM — tuned configs as generated headers, with a typed fallback ladder ([repo](https://github.com/ROCm/rocm-libraries/tree/develop/projects/rocprim)) + +- **Tuning output is generated source.** `tuning/configs_header_generation.py` (~46 KB) + turns benchmark results into config headers compiled into the library; `run_tuning.py` + drives the sweep. Tuned configuration is therefore a build artifact under version control, + not a runtime database. **[V]** +- **The idea worth taking is `fallback_config.json`** — what happens for a type nobody + tuned. Rather than a single global default, it defines a **typed fallback ladder**: each + entry names a representative `based_on_type` plus a predicate + (`sizeof_min_exclusive`, `sizeof_max_inclusive`, `is_floating_point`). An untuned 2-byte + float falls back to `rocprim::half`'s config; a 12-byte non-float falls back to + `int128_t`'s; and so on down through `float`, `int64_t`, `int`, `short`. **[V]** + +**[I]** That is *dtype bucketing* — the same move Decision #28 makes for shapes, applied to +the type axis. It is how you tune a handful of representatives and still answer for the long +tail of dtypes, which is precisely the problem our 15-name canonical dtype set plus planned +low-precision formats will create. + +### 1.11 rocRAND — tuning that changes the answer, made an explicit contract ([repo](https://github.com/ROCm/rocm-libraries/tree/develop/projects/rocrand)) + +The important finding in this section, and a direct check on Decision #18. + +- Under `ROCRAND_ORDERING_PSEUDO_DYNAMIC`, rocRAND picks block and grid sizes to suit the + specific GPU model. AMD states the consequence plainly: *"the number of allocated + generators and the sequence of the generated numbers can also vary."* **[V]** +- Tuning is a benchmark sweep over a block-size × grid-size matrix, fastest wins per device + (`benchmark_rocrand_tuning`, with `BENCHMARK_TUNING_{THREAD,BLOCK}_OPTIONS` and a + minimum-grid filter). Grid candidates are generated as **multiples of the device's CU + count** by a helper that shells out to `rocminfo`. **[V]** + +**[I]** So rocRAND treats *reproducibility versus performance* as a named, opt-in ordering +mode rather than an emergent property. We should confirm the same is true for us: Decision +#18 fixes `stream_id = global_seed * num_ranks + rank` with non-overlapping Philox counter +offsets, and `@jit(deterministic=True)` is a declared knob — but if any tuned launch +configuration ever feeds an RNG offset scheme, autotuning would silently change numerical +output. That is a cheap thing to assert and an expensive thing to discover late. + +### 1.12 rocALUTION — runtime hardware abstraction with automatic host fallback ([repo](https://github.com/ROCm/rocm-libraries/tree/develop/projects/rocalution)) + +Sparse iterative solvers and preconditioners. Included as a deliberate **contrast** with our +own design rules rather than as a pattern to copy. **[V]** + +- Objects (matrix, vector, solver) are allocated on the host and *moved* to an accelerator by + a call. Execution location is chosen at **run time via RTTI**, which AMD explicitly + contrasts with template-based libraries that decide at compile time. The stated goal is one + source that runs unchanged with or without an accelerator. +- **Automatic fallback:** *"The library checks at run time whether a specific routine is + implemented on the selected accelerator. If the routine is not available, the associated + object is moved back to the host. The computation is then performed on the CPU + automatically."* + +**[I]** That last behaviour is the opposite of Decision #21, which requires a stable +diagnostic naming the op and target rather than a silent fall-through. Both choices are +defensible for their audience — rocALUTION optimizes for "it always runs", we optimize for +"you always know where it ran" — but the contrast is worth having on record, because silent +host migration is exactly how a performance cliff hides. Our `reference_cpu` labelling in the +Apple audit is the same instinct as rocALUTION's fallback, made visible. + --- ## 2. Patterns to adopt — ranked, mapped to Tessera diff --git a/docs/audit/compiler/AMD_KERNEL_COMPILER_SURVEY.md b/docs/audit/compiler/AMD_KERNEL_COMPILER_SURVEY.md index 587b28c6a..98a7ada2f 100644 --- a/docs/audit/compiler/AMD_KERNEL_COMPILER_SURVEY.md +++ b/docs/audit/compiler/AMD_KERNEL_COMPILER_SURVEY.md @@ -914,6 +914,41 @@ diagnostic so the fallback is visible without being noisy. That combination is a good template for our capability gates, which today tend to be binary supported/unsupported. **[I]** +### 4.7 `rocisa` — the Python/C++ boundary, and what it costs + +TensileLite's assembly generator is not Python. `rocisa/` is a C++ module bound via +**nanobind**, and `KernelWriter.py` calls into it to emit instructions. Reading its +developer docs is worthwhile because it is the same hybrid we are: a Python compiler +driving a C++ core. **[V]** + +- **IR nodes carry a mandatory deep-copy contract.** Anything inheriting `Item` or + `Instruction` must supply a copy constructor *and* override + `clone() -> std::shared_ptr`, with the instruction to "make sure you deepcopy any + pointer". Python-side `__deepcopy__` is wired explicitly per class. Deep copy is treated + as a first-class requirement of the node type, not an afterthought. +- **The boundary is sharper than it looks.** *"Vector memory management between Python and + C++ is different, so exporting vectors to Python is copy instead of reference."* So + `module.items()` returns a **copy** — elements are `shared_ptr`, so you can mutate what + they point at, but **you cannot assign or replace an element** through that handle. That + is exactly the class of bug that is silent in Python and impossible to reason about from + the C++ side. +- **Convenience costs throughput at the boundary.** `countType(module, Instruction)` is kept + because it is handy, with the explicit note that it *"runs slower than directly using + templates"* and that a templated `countInstruction(module)` should be added and exported + when it matters. Convenience wrappers are marked as prototyping tools rather than + quietly becoming the default. +- **Staleness is a hard error, not a warning.** `rocisa/__init__.py` compares source + timestamps against a generated `_build_info.py`; if any `.cpp/.hpp/.h/.def/.inc` is newer + than the loaded `.so`, import raises with a rebuild message. Pre-built wheels omit + `_build_info.py` and skip the check. + +**[I]** Three of these are directly worth stealing for our own Python↔C++ seam. The +import-time staleness check is the cheapest and highest-value — we have already lost time +this session to a `tessera-opt` binary that silently did not match its sources, and an +equivalent guard turns that from a confusing test failure into one clear message. The +copy-not-reference boundary rule and the "convenience wrapper is for prototyping" note are +both things better written down than rediscovered. + --- ## 5. Concrete algorithms worth taking @@ -1049,34 +1084,38 @@ the dominant tiling lever. 13. **Reject-and-continue with warn-once** — §4.6. Capability gates that degrade to another algorithm and say so once, instead of binary supported/unsupported. +14. **Import-time staleness check on the C++ extension** — §4.7. rocisa raises on + import if any source is newer than the built `.so`. We lost time this session to a + stale `tessera-opt`; this converts that into one clear message. + **Take — larger, worth designing toward:** -14. **Name Y-space** — §3.1. The algorithm's local view of its own work is the one +15. **Name Y-space** — §3.1. The algorithm's local view of its own work is the one coordinate space our IR does not have, and it is what lets the distribution become a separate swappable object. Highest-leverage *conceptual* item in this document. -15. **Distribution-as-a-value** — §3.3, §3.9. `tile_distribution_encoding` makes +16. **Distribution-as-a-value** — §3.3, §3.9. `tile_distribution_encoding` makes the thread↔data mapping comparable and enumerable — a search coordinate for the autotuner rather than a code variant. §3.9 raises the stakes: vector width, access count, and traversal order are all **derived** from the encoding, so this is what makes the rest of the machinery generated rather than authored. -16. **Derive the access pattern, don't author it** — §3.9. Only *which data is +17. **Derive the access pattern, don't author it** — §3.9. Only *which data is mine* and *what to compute* are written; how to fetch it and in what order are consequences. Worth testing our tile lowering against: how many of those four are currently hand-specified? -17. **Static shape / dynamic origin as the symbolic-dim seam** — §3.9. CK puts +18. **Static shape / dynamic origin as the symbolic-dim seam** — §3.9. CK puts window *lengths* in the type and the *origin* in a runtime field. That is a clean, load-bearing instance of Decision #28's `static | bucket | dynamic` policy. -18. **A closed transform operator set with bidirectional + incremental ops** — +19. **A closed transform operator set with bidirectional + incremental ops** — §3.2. Especially `update_lower_index()`: a tile IR that can only recompute absolute coordinates pays for it in every loop. -19. **The coordinate hypergraph as a representation** — §2.1. Take the index +20. **The coordinate hypergraph as a representation** — §2.1. Take the index algebra; do *not* take the single-mutable-IR architecture, which trades away the verifiable stage boundaries our lit discipline depends on. -20. **Serializable IR as cache key + fixture** — §2.6. -21. **"New arch = data only, no source edits"** as an explicit acceptance test for +21. **Serializable IR as cache key + fixture** — §2.6. +22. **"New arch = data only, no source edits"** as an explicit acceptance test for the backend-plugin seam — §1.3. **Skip / already have:** @@ -1127,8 +1166,8 @@ implementation. **Beyond these three**, still unread: 1. **hipBLASLt kernel *generation*** — §4 read the selection side only. - `KernelWriter.py` / `KernelWriterAssembly.py`, the `rocisa` Nanobind assembly - module, `Components/` (modular MAC / global-read / scheduling blocks), and the + `KernelWriter.py` / `KernelWriterAssembly.py`, the `rocisa` C++ *sources* (its + developer docs are read in §4.7; `include/` and `src/` are not), `Components/` (modular MAC / global-read / scheduling blocks), and the three-phase `BenchmarkProblems → LibraryLogic → ClientWriter` tuning pipeline are all unread. `ContractionProblemPredicates.hpp` alone is 119 KB. 2. **`ExactLogicLibrary` / `MapLibrary` / `CachingLibrary` internals** — §4.2 @@ -1160,7 +1199,8 @@ lds_index_swapping,tile_window,sweep_tile}.rst`. **hipBLASLt / TensileLite** (`projects/hipblaslt/`) — `CLAUDE.md`, `tensilelite/CLAUDE.md`; `tensilelite/include/Tensile/{SolutionLibrary,Distance, -MatchingLibrary,GranularitySelectionLibrary,MLPClassification,MLFeatures}.hpp`. +MatchingLibrary,GranularitySelectionLibrary,MLPClassification,MLFeatures}.hpp`; +`tensilelite/rocisa/{README.md,docs/}`. **LLVM** (`ROCm/llvm-project@amd-staging`) — `llvm/lib/Target/AMDGPU/SISchedule.td` (the gfx1250/gfx1251 machine model, recorded in the gfx1250 target reference §2.5 From bec0c0d2b88fff614a4dc2bde00d3b141f705178 Mon Sep 17 00:00:00 2001 From: Greg Stoner Date: Wed, 29 Jul 2026 08:25:35 -0600 Subject: [PATCH 3/3] Address PR 476 review: complete the LDS phase table, scope calibration correctly MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit P2 — the ds_read_b128 phase table listed only 3 of 4 phases, covering 48 of 64 lanes. An analyzer built from it would leave lanes 36-43, 48-51 and 60-63 unmodelled and could accept a layout that conflicts on exactly those lanes. Phase 3 is T36-39, T40-43, T48-51, T60-63; the structure is phase 1 = phase 0 + 32 and phase 3 = phase 2 + 32, which doubles as the completeness check (four disjoint sets covering 0-63 exactly once). P1 — APPLE_AUDIT.md scoped the hardware-free cost-model calibration to Apple alone, asserting that ROCm and NVIDIA kernels "cannot be measured". Both claims were false, written from the pre-bring-up framing CLAUDE.md itself retired: NVIDIA has a committed, consumed, device-keyed sm_120 autotune corpus and ROCm has measured gfx1151 retune/ratchet evidence. Excluding them would also have weakened the calibration - a score fitted on one architecture reproduces the single-arch overfit the TileSight assessment records for NeuSight. All four architecture queues now record a state under sync key COSTMODEL-CALIB-2026-07-29: NVIDIA-CALIB-1 (shape depth, cheapest to land - analysis over already-recorded data), ROCM-CALIB-1 (the metric's home ground, extracted from production AMD code; wave32 constants must be re-derived), APPLE-CALIB-1 (op breadth, with the conflict metric possibly not applicable to Metal threadgroup memory), X86-CALIB-1 (split verdict: locality applies, the bank-conflict analyzer does not - no software-managed scratchpad or wave phases). Co-Authored-By: Claude Opus 5 --- docs/audit/backend/apple/APPLE_AUDIT.md | 44 +++++++++++++++---- docs/audit/backend/apple/todo.md | 37 +++++++++++++++- docs/audit/backend/nvidia/todo.md | 40 ++++++++++++++++- docs/audit/backend/rocm/todo.md | 36 ++++++++++++++- docs/audit/backend/x86/todo.md | 36 ++++++++++++++- .../compiler/AMD_KERNEL_COMPILER_SURVEY.md | 23 ++++++++++ 6 files changed, 203 insertions(+), 13 deletions(-) diff --git a/docs/audit/backend/apple/APPLE_AUDIT.md b/docs/audit/backend/apple/APPLE_AUDIT.md index d193894de..ecca19a9c 100644 --- a/docs/audit/backend/apple/APPLE_AUDIT.md +++ b/docs/audit/backend/apple/APPLE_AUDIT.md @@ -299,15 +299,41 @@ not current support counts. a bank-conflict analyzer that computes N-way conflict from a descriptor alone (§3.8). Both are computable on any target and neither needs silicon. - Apple is the only backend that executes broadly enough to say whether such a - metric *predicts anything*. The action is therefore not "add a metric" but - **calibrate one**: compute the locality/conflict score for kernel families the - Apple lane already measures, and check the score against recorded latency. A - metric that does not rank Apple kernels correctly should not be trusted to - rank ROCm or NVIDIA kernels we cannot measure. This is the concrete follow-on - to the mock-cost-model finding in - [`../../compiler/TILESIGHT_ASSESSMENT.md`](../../compiler/TILESIGHT_ASSESSMENT.md) - §2, and it gates how much weight the arbiter's hardware-free tier can carry. + The action is not "add a metric" but **calibrate one**: compute the + locality/conflict score for kernel families that are already measured, and + check the score against recorded latency. A metric that does not rank measured + kernels correctly should not be trusted to rank unmeasured ones. This is the + concrete follow-on to the mock-cost-model finding in + [`TILESIGHT_ASSESSMENT.md`](../../compiler/TILESIGHT_ASSESSMENT.md) §2, and it + gates how much weight the arbiter's hardware-free tier can carry. + + **This is a cross-backend calibration, not an Apple-only one** (corrected + 2026-07-29 — an earlier draft of this item said Apple was "the only backend + that executes broadly enough" and referred to "ROCm or NVIDIA kernels we + cannot measure"; both claims were false, written from the pre-bring-up framing + CLAUDE.md itself retired). NVIDIA has a committed, consumed, device-keyed + `nvidia:sm_120` autotune corpus over 64/256/512/1024/2048 square buckets plus + fused GEMM and causal attention + ([`../nvidia/NVIDIA_AUDIT.md`](../nvidia/NVIDIA_AUDIT.md); generator at + `benchmarks/nvidia/record_autotune_corpus.py`). ROCm has measured gfx1151 + evidence including the size-adaptive grouped-GEMM tile selector, the hot-path + ratchet, and a measured resident-GPU crossover for sparse attention + ([`../rocm/ROCM_AUDIT.md`](../rocm/ROCM_AUDIT.md); + `rocm_gfx1151_compiler_retune_2026_07_15.json`). + + Excluding those corpora would not just be unfair to the sibling backends — it + would make the calibration *weaker*. A locality metric fitted against one + architecture is the exact failure mode the same assessment records for learned + predictors: TileSight §5.2 shows NeuSight leading on the A100 in its training + distribution and losing that lead on every newer part. A score validated on + Apple alone would carry the same defect by construction. + + Apple's actual distinctive contribution is **op breadth**, not exclusivity: + the widest F4-verified family envelope in the fleet, so it can test whether a + score generalizes *across op kinds*. NVIDIA and ROCm contribute **shape depth** + within GEMM/attention. The calibration needs both axes, so it is owned across + the queues under sync key `COSTMODEL-CALIB-2026-07-29` — see the per-backend + items named there. A second, smaller item from the same survey: the MSL synthesizer currently *authors* access patterns, whereas CK *derives* vector width, access count and diff --git a/docs/audit/backend/apple/todo.md b/docs/audit/backend/apple/todo.md index 4c915ac20..d1bb621c0 100644 --- a/docs/audit/backend/apple/todo.md +++ b/docs/audit/backend/apple/todo.md @@ -3,11 +3,46 @@ audit_role: plan plan_state: landing owner: Apple backend target: apple_gpu -last_updated: 2026-07-28 +last_updated: 2026-07-29 --- # Apple compiler, exact-device, and performance plan +## APPLE-CALIB-1: contribute op breadth to the hardware-free score calibration + +Cross-backend sync `COSTMODEL-CALIB-2026-07-29` — **follow-up required, owning +host M1 Max (apple7).** Apple is the *breadth* axis, not the sole site. + +**What is being calibrated.** Two static, device-free quality metrics found in +production AMD code and recorded in +[`../../compiler/AMD_KERNEL_COMPILER_SURVEY.md`](../../compiler/AMD_KERNEL_COMPILER_SURVEY.md) +§3.7–3.8: a step-distance locality histogram over a materialized access order, +and an N-way bank-conflict analyzer computed from a descriptor alone. Both are +computable on any target with no silicon. The question is whether either +*predicts measured latency* — which decides how much weight the arbiter's +hardware-free tier can carry, per +[`TILESIGHT_ASSESSMENT.md`](../../compiler/TILESIGHT_ASSESSMENT.md) §2. + +**Apple's role.** The widest F4-verified op-family envelope in the fleet, so it +answers *does the score generalize across op kinds* — norm chains, attention with +online softmax, pointwise-reduce, gated matmul, coopmat `simdgroup_matrix`, not +just GEMM. NVIDIA and ROCm supply shape depth within GEMM/attention +(`NVIDIA-CALIB-1`, `ROCM-CALIB-1`). Both axes are required: a score fitted on one +architecture reproduces the overfit that assessment §5.2 records for NeuSight, +which led on the A100 inside its training distribution and lost that lead on +every newer part. + +**Apple-specific caveat.** The bank-conflict half was derived for LDS with a +known bank count and a 4-phase wave64 access pattern. Metal threadgroup memory is +not LDS and its banking is not documented to the same level (Decision #27 — do +not assert a Metal hardware detail without a real source), so the conflict metric +may be **not applicable** on Apple even where the locality metric is not. Report +that split rather than one blended verdict. + +**Missing exact-device evidence.** Rank correlation between each score and +recorded M1 Max latency, per op family, over the families the Apple lane already +measures. A score that does not rank measured Apple kernels correctly is not +trustworthy for unmeasured kernels anywhere. ## APPLE-RASTER-1: reconcile the MLX-inherited swizzle with the shared contract Cross-backend sync `RASTER-CONTRACT-2026-07-28` — **follow-up required, owning diff --git a/docs/audit/backend/nvidia/todo.md b/docs/audit/backend/nvidia/todo.md index f274c5373..705823cd0 100644 --- a/docs/audit/backend/nvidia/todo.md +++ b/docs/audit/backend/nvidia/todo.md @@ -3,11 +3,49 @@ audit_role: plan plan_state: landing owner: NVIDIA backend target: nvidia_sm120 -last_updated: 2026-07-28 +last_updated: 2026-07-29 --- # NVIDIA compiler test-suite evaluation and rearchitecture +## NVIDIA-CALIB-1: supply the sm_120 corpus to the hardware-free score calibration + +Cross-backend sync `COSTMODEL-CALIB-2026-07-29` — **follow-up required, owning +host NR2 Pro (RTX 5070 Ti, sm_120).** + +**Correction that created this item.** `APPLE_AUDIT.md` originally scoped this +calibration to Apple alone, on the stated grounds that ROCm and NVIDIA kernels +"cannot be measured". That was false for NVIDIA: this backend already has a +committed, **consumed**, device-keyed `nvidia:sm_120` autotune corpus covering +64/256/512/1024/2048 square buckets plus fused GEMM and causal attention, +generated by `benchmarks/nvidia/record_autotune_corpus.py`. Excluding it would +have discarded the deepest per-shape latency evidence in the fleet. + +**What is being calibrated.** Two static, device-free scores from +[`../../compiler/AMD_KERNEL_COMPILER_SURVEY.md`](../../compiler/AMD_KERNEL_COMPILER_SURVEY.md) +§3.7–3.8 — a step-distance locality histogram and an N-way bank-conflict +analyzer — against recorded latency, to decide how much weight the arbiter's +hardware-free tier can carry +([`TILESIGHT_ASSESSMENT.md`](../../compiler/TILESIGHT_ASSESSMENT.md) §2). + +**NVIDIA's role: shape depth.** The committed corpus already varies the shape +axis within GEMM and attention at fixed op kind, which is exactly the axis Apple +cannot supply and the one a locality score most needs to be tested against — +locality changes with shape at constant op. Apple supplies op breadth +(`APPLE-CALIB-1`); ROCm supplies a second, independent architecture +(`ROCM-CALIB-1`). Fitting on any one of the three reproduces the single-arch +overfit the assessment records for NeuSight. + +**Note on translating the metrics.** Both were derived for AMD LDS with a known +bank count and a wave64 4-phase access pattern (survey §5.1). CUDA shared memory +is 32-bank and warp-synchronous; the *method* transfers but every constant must +be re-derived for sm_120 before a conflict number here means anything. Do not +port AMD constants. + +**Missing exact-device evidence.** Rank correlation between each score and the +recorded sm_120 latencies already in the corpus — this needs no new hardware run, +only an analysis pass over committed data, which makes it the cheapest of the +three contributions to land. ## NVIDIA-RASTER-1: consume the shared block-rasterization contract Cross-backend sync `RASTER-CONTRACT-2026-07-28` — **follow-up required, owning diff --git a/docs/audit/backend/rocm/todo.md b/docs/audit/backend/rocm/todo.md index 3d7f02703..53867bd89 100644 --- a/docs/audit/backend/rocm/todo.md +++ b/docs/audit/backend/rocm/todo.md @@ -1,5 +1,5 @@ --- -last_updated: 2026-07-28 +last_updated: 2026-07-29 audit_role: plan plan_state: open scope: ROCm backend implementation and exact-device proof @@ -7,6 +7,40 @@ scope: ROCm backend implementation and exact-device proof # ROCm backend TODO +## ROCM-CALIB-1: supply gfx1151 evidence to the hardware-free score calibration + +Cross-backend sync `COSTMODEL-CALIB-2026-07-29` — **follow-up required, owning +host Strix Halo (Radeon 8060S, gfx1151).** + +**Correction that created this item.** `APPLE_AUDIT.md` originally scoped this +calibration to Apple alone, stating that ROCm and NVIDIA kernels "cannot be +measured". False here: this backend holds the measured size-adaptive grouped-GEMM +tile selector, the hot-path perf ratchet, a measured resident-GPU crossover for +large-block sparse attention, and the committed +`rocm_gfx1151_compiler_retune_2026_07_15.json` retune corpus. + +**ROCm's special standing on this item.** Both metrics *originate here*. The +step-distance locality histogram and the N-way bank-conflict analyzer are +extracted from production AMD code +([`../../compiler/AMD_KERNEL_COMPILER_SURVEY.md`](../../compiler/AMD_KERNEL_COMPILER_SURVEY.md) +§3.7–3.8), so ROCm is the one backend where the metric can be checked against the +hardware model it was actually written for. If it fails to rank gfx1151 kernels, +that is a much stronger negative result than failing on a target it was never +designed for — and it should end the line of work rather than prompt retuning. + +**Constant re-derivation is mandatory, not optional.** The published phase table +is GFX950/wave64 with 64 banks (survey §5.1, now including all four phases — +`phase 1 = phase 0 + 32`, `phase 3 = phase 2 + 32`, the four disjoint and +covering lanes 0–63). gfx1151 is **wave32** with a different bank count, so every +constant must be re-derived before a conflict count means anything on this part. +ISA truth stays [`docs/reference/isa/rdna/`](../../../reference/isa/rdna/) +(gfx1151 = RDNA 3.5). + +**Missing exact-device evidence.** Rank correlation between each score and the +recorded gfx1151 latencies, over the retune corpus and hot-path ratchet rows. The +grouped-GEMM selector is the most informative subject: its tile choice is already +known to be size-adaptive under measurement, so a locality score that cannot +reproduce that ordering has failed on its home ground. ## ROCM-RASTER-1: consume the shared block-rasterization contract Cross-backend sync `RASTER-CONTRACT-2026-07-28` — **follow-up required, owning diff --git a/docs/audit/backend/x86/todo.md b/docs/audit/backend/x86/todo.md index fdf8f9b51..51491f040 100644 --- a/docs/audit/backend/x86/todo.md +++ b/docs/audit/backend/x86/todo.md @@ -1,5 +1,5 @@ --- -last_updated: 2026-07-28 +last_updated: 2026-07-29 audit_role: plan plan_state: open owner: x86 backend @@ -9,6 +9,40 @@ scope: x86 AMX/AVX-512 backend implementation and exact-device proof # x86 backend TODO +## X86-CALIB-1: split verdict on the hardware-free score calibration + +Cross-backend sync `COSTMODEL-CALIB-2026-07-29` — **split: bank-conflict metric +not applicable; locality metric follow-up required.** Owning host Zen 5 (Ryzen AI +Max+ 395 CPU complex, AVX-512, no AMX). + +Two static device-free scores are being calibrated against measured latency +([`../../compiler/AMD_KERNEL_COMPILER_SURVEY.md`](../../compiler/AMD_KERNEL_COMPILER_SURVEY.md) +§3.7–3.8; motivation in +[`TILESIGHT_ASSESSMENT.md`](../../compiler/TILESIGHT_ASSESSMENT.md) §2). They do +not get the same verdict here, and reporting one blended state would hide that. + +**Bank-conflict analyzer — not applicable, architecture-specific reason.** It +counts N-way conflicts across a fixed number of software-managed scratchpad banks +under a wave's phase-grouped access. The x86 lane has no software-managed +scratchpad and no wave phases: AVX-512 loads go through a hardware-managed +L1/L2/L3 hierarchy where the analogous hazards are 4 KiB aliasing, cache-set +associativity conflicts, and store-forwarding stalls. Those are real, but they +are a different model with different inputs — not this analyzer with different +constants. + +**Locality histogram — follow-up required.** The step-distance histogram over a +materialized access order is genuinely target-independent: it scores an access +*order*, not a memory technology. This is also the metric with the strongest +prior for CPUs, since blocked-algorithm cache analysis is a CPU literature +(Lam/Rothberg/Wolf 1991, cited in the same assessment). x86 executes natively and +has committed benchmarks (`benchmarks/benchmark_x86_e2e*.py`), so it can supply a +non-GPU architecture to the calibration — valuable precisely because a score that +holds across CPU *and* GPU is far less likely to be fitting an accelerator +artifact. + +**Missing exact-device evidence.** Rank correlation between the locality score +and recorded Zen 5 AVX-512 latencies over the e2e benchmark rows. No evidence is +owed for the conflict metric. Cross-backend sync `RASTER-CONTRACT-2026-07-28` — **not applicable, with an architecture-specific reason.** Schedule IR gained `raster_order` / `raster_group` on `schedule.tile` / `schedule.knob` (arch-neutral definition in diff --git a/docs/audit/compiler/AMD_KERNEL_COMPILER_SURVEY.md b/docs/audit/compiler/AMD_KERNEL_COMPILER_SURVEY.md index 98a7ada2f..d7bed7bee 100644 --- a/docs/audit/compiler/AMD_KERNEL_COMPILER_SURVEY.md +++ b/docs/audit/compiler/AMD_KERNEL_COMPILER_SURVEY.md @@ -582,6 +582,21 @@ that the analytical cost model our arbiter falls back on was a mock. A step-distance histogram over a materialized access order is cheap, needs no device, and is a real signal. **[I]** +> **Calibration is owned across all four queues** under sync key +> `COSTMODEL-CALIB-2026-07-29`, not by whichever backend happens to be nearest. +> A score is only worth what it predicts, and one fitted against a single +> architecture reproduces the overfit `TILESIGHT_ASSESSMENT.md` §5.2 records for +> NeuSight — top of the table on the A100 it trained on, beaten on every newer +> part. So: NVIDIA `NVIDIA-CALIB-1` (committed device-keyed sm_120 corpus — +> shape depth, and the cheapest to land since it is an analysis pass over data +> already recorded), ROCm `ROCM-CALIB-1` (gfx1151 retune corpus and hot-path +> ratchet — and the metric's *home ground*, since it was extracted from +> production AMD code), Apple `APPLE-CALIB-1` (widest F4-verified op-family +> envelope — op breadth), x86 `X86-CALIB-1` (**split verdict**: the locality +> histogram applies and adds a non-GPU architecture; the §3.8 bank-conflict +> analyzer does not, as the AVX-512 lane has no software-managed scratchpad or +> wave phases). **[I]** + ### 3.8 Two worked swizzles — and what they prove about the algebra CK documents two independent bank-conflict swizzles. Both are built **entirely @@ -1000,6 +1015,14 @@ And the phases are **not contiguous lane ranges**: | 0 | T0-3, T12-15, T20-23, T24-27 | | 1 | T32-35, T44-47, T52-55, T56-59 | | 2 | T4-7, T8-11, T16-19, T28-31 | +| 3 | T36-39, T40-43, T48-51, T60-63 | + +All four phases are listed on purpose: the four are a *partition* of the 64 lanes, +so an analyzer built from a three-row table would leave lanes 36–43, 48–51 and +60–63 unmodelled and could accept a layout that conflicts on exactly those lanes. +The structure is `phase 1 = phase 0 + 32` and `phase 3 = phase 2 + 32`, which is +also the check that the rows are complete — the four sets are disjoint and cover +0–63 exactly once. **[V]** Any conflict analysis that assumes lanes 0–15 form a phase will compute the wrong answer. **[V]**