diff --git a/docs/audit/backend/apple/APPLE_AUDIT.md b/docs/audit/backend/apple/APPLE_AUDIT.md index 93fabe183..ecca19a9c 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,55 @@ 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. + + 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 + 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/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/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..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,18 +1,31 @@ --- -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 > [`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**. +> 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), @@ -46,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, @@ -144,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)) @@ -241,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/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 new file mode 100644 index 000000000..d7bed7bee --- /dev/null +++ b/docs/audit/compiler/AMD_KERNEL_COMPILER_SURVEY.md @@ -0,0 +1,1241 @@ +--- +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]** + +> **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 +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]** + +### 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 + +### 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 | +| 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]** + +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. + +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:** + +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. +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. +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? +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. +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. +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. +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:** + +- 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` 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 + 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`; +`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 +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).*