[rocm-libraries] ROCm/rocm-libraries#12173 (commit 4606f83)

feat(ck-tile): make dispatcher LDS capacity budget
 architecture-aware (#12173)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

JIRA ID : AICK-2249

## Summary

The dispatcher's codegen rejects any GEMM tile whose LDS staging
footprint exceeds a cap. The cap was keyed on the **pipeline only** and
carried no architecture term, so newer targets with substantially more
LDS than gfx942 were all held to gfx942's budget.

The effect is that the largest and deepest tiles — the ones most likely
to win on large GEMM shapes — were never generated, never benchmarked
and never selectable on those targets. On a gfx950 fp16 sweep, **160 of
288 offered configurations (56%) were rejected by this cap alone**.

This PR makes the budget architecture-aware, adds gfx1250 to the spec,
fixes both the Python and C++ validators together, and makes the check
account for ping-pong LDS staging.

**Scope:** the check is not GEMM-specific. `_validate_lds_capacity` sits
in the common `validate_kernel` chain, ungated by operator, and
`OperatorType` covers 11 operators (5 GEMM variants, 6 conv variants).
In Python it is reached by `unified_gemm_codegen.py:1595` (universal
GEMM) and `unified_grouped_conv_codegen.py:1941` (grouped convolution);
in C++ by `Registry::filter_by_arch()` (`registry.cpp:160`), which is
operator-agnostic and covers every registered kernel. The ctypes bridge
paths (quant variants, batched contraction, multi-ABD, FMHA) validate
via `validate_kernel_config` in `python/ctypes_utils.py`, which has no
LDS check at all, so they are unaffected either way.

Implements AICK-2249. Baseline evidence from AICK-2244.

## Motivation

`arch_specs.json` had `pipeline_lds_limits` as a **top-level** key — a
sibling of `architectures`, not nested inside it:

```json
"pipeline_lds_limits": {
  "mem": <bytes>, "compv3": <bytes>, "compv4": <bytes>, ... "default": <bytes>
}
```

Every value was a fixed byte count, derived from gfx942 and applied to
every target.

There was no architecture axis, and nowhere to put one.
`_validate_lds_capacity` did `LDS_CAPACITY_LIMITS.get(config.pipeline,
...)` and nothing else.

Meanwhile `arch.hpp` has declared a per-architecture LDS capacity all
along, via `get_lds_size()`. On gfx942 the hardcoded cap and that
capacity coincide; on the newer targets they do not, and the shortfall
is large. This PR consumes the existing declaration rather than
introducing any new hardware figure.

This is the usual shape of this defect: a value that was correct when
there was exactly one architecture, frozen into a schema with no slot
for a second one. It survived because on gfx942 the correct answer and
the hardcoded answer are the same number, so it only ever fails
*silently, by generating less*.

The evidence that it is binding rather than theoretical: on the gfx950
sweep the cutoff falls **exactly** on the hardcoded cap for compv3/mem,
and exactly on the tighter cap for compv4, with zero exceptions in
either direction. And the winning kernels on that target were
`128x256x64` and `128x128x128` under compv3 — precisely tiles that
compv4 was forbidden from using. When the measured optimum sits on the
constraint boundary, the constraint is probably binding.

## Design note

**I did not invent a new schema.** A per-architecture `lds_capacity_kb`
field already existed in `arch_specs.json` (gfx950 already correctly
said `160`), and `ADDING_NEW_GPU.md` already documented it as a required
onboarding field. Nothing read it. Every new-GPU onboarding has been
filling in a mandatory field that no code consumed — the contract was
documented and unhonoured. So the fix is to honour it.

Pipelines now declare a **basis** rather than a byte count:

```json
"compv3": { "basis": "fraction", "value": 1.0 },
"compv4": { "basis": "fraction", "value": 0.5 }
```

resolved against each architecture's capacity at generation time. A new
GPU needs **one** number, not ten. The alternative — nesting the byte
table under each architecture — would have reintroduced the same failure
mode one level down: 8 arches x 10 pipelines of hand-maintained
constants free to drift from `arch.hpp` independently.

Resolution happens in the generator, once, so the emitted Python and C++
get literal byte counts and cannot disagree about how to read the
schema.

### On the tighter-capped pipelines — I read the sources rather than
guessing

Four pipelines sat at the tighter cap and the ticket flagged it as an
open question: genuine double-buffering that should scale with capacity,
or an independent absolute limit? On gfx942 the two readings are
numerically identical, which is why it was never forced. **The answer is
not uniform across the four:**

| pipeline | doubles LDS? | evidence |
|---|---|---|
| `compv4` | **yes** | `GemmPipelineAgBgCrCompV4::GetSmemSize()` returns
`2 * Policy::GetSmemSize<Problem>()` |
| `preshufflev2` | **yes** |
`WeightPreshufflePipelineAGmemBGmemCRegV2::GetSmemSize()` returns
`DoubleSmemBuffer ? 2 * smem_size : smem_size` |
| `compv6` | **no** | returns the policy size unmultiplied |
| `preshufflev1` | n/a | no such pipeline exists;
`preshuffle_pipelines.supported` lists only `preshufflev2` |

This is corroborated by `DOUBLE_SMEM_PIPELINES = {"compv4",
"preshufflev2", "comp_async"}`, which already exists in
`unified_batched_contraction_codegen.py` and names exactly the same two.

So for compv4 and preshufflev2, half-of-capacity is the **exact** model,
not a safety margin: the validator checks `A+B` against `capacity/2`,
which is precisely equivalent to checking the real `2*(A+B)` allocation
against full capacity.

For compv6 and preshufflev1 the halved cap is **unexplained by
buffering**. I preserved their ratio rather than widening them, because
no evidence supports a larger budget and a blind raise that regresses is
worse than the status quo. Widening those two is a separate change that
needs a measurement behind it. The reasoning is recorded in
`arch_specs.json` next to each value.

### Ping-pong staging had to be threaded in

The check modelled a *single* staging buffer. But
`unified_grouped_conv_codegen.py:148` documents that `mem`, `compv3`,
`compv5` and `compv6` make double buffering a **configuration choice**
(`--double-smem-buffer`), not a property of the pipeline — and the conv
path passed `pipeline=` into the validator without that flag, so the
budget could not see it.

That was harmless while the cap matched gfx942's capacity, because twice
that still fits the larger parts. **Widening the budget removes the
accidental headroom**, so I had to close it in the same PR:

Writing `C` for a target's LDS capacity:

| | compv3 budget | actually allocated | vs capacity | |
|---|---|---|---|---|
| develop | gfx942's cap | 2x that | below `C` | fits, by accident |
| arch-aware alone | `C` | `2C` | **over `C`** | **overflows** |
| with this fix | `C/2` | `C` | exactly `C` | fits exactly |

So the flag is now threaded into both validators. Conv passes the value
it already tracks; C++ reads `algorithm.double_buffer`, which registered
kernels populate from `SelectedKernel::DoubleSmemBuffer`
(`kernel_registration.hpp:60`).

Pipelines that *always* double already carry the halving in their
per-pipeline budget, so the two signals are combined with a `min()` —
`compv4` and `preshufflev2` are never halved twice, which is what keeps
gfx942 byte-identical. Independent confirmation that those are exactly
the two: `unified_gemm_codegen.py:1531` sets `double_buffer = pipeline
in ("compv4", "preshufflev2")`.

One deliberate non-use: `KernelConfig::build_key()`
(`kernel_config.hpp:280`) and `utils.hpp:676` hardcode `double_buffer =
true` for *any* pipeline. Those build **query** keys, which never reach
`validate_lds` — only registered-instance keys do — so the flag is
trustworthy at the one site that reads it. Worth fixing separately.

### A landmine worth calling out

`arch_specs_generated.py` **already contained gfx1250** — family, warp
configs, warp tile combos — but `arch_specs.json`, the file it is
generated from, did not. The generated C++ header did not either.
Three-way drift, in a file stamped `AUTO-GENERATED - DO NOT EDIT
DIRECTLY`.

This means following the documented workflow (edit JSON, regenerate,
commit) would have **silently deleted gfx1250's warp tables**. I
transplanted them into the JSON verbatim rather than authoring anything,
and the regeneration check below confirms every table came back
byte-identical.

## Test plan

- [x] New regression suite passes (15/15). Asserts budgets differ across
gfx942/gfx950/gfx1250, that gfx942's budget is byte-identical to the
historical table, that no budget exceeds the declared hardware capacity,
that unknown targets get the *smallest* budget rather than the largest,
and end-to-end that a large staging tile is rejected on gfx942 and
accepted on gfx950
- [x] Double-buffer cases covered: configurable pipelines halve, `2 x
budget <= capacity` on every arch/pipeline pair, always-double pipelines
are not halved twice, single-buffered remains the default, and
end-to-end a large staging tile is accepted on gfx950 single-buffered
but rejected double-buffered
- [x] Existing suites pass: `test_arch_filter_constraints`,
`test_gemm_utils`, `test_codegen_common`, `test_dispatcher_common`,
`test_tile_math`, `test_grouped_conv_codegen`, `test_grouped_conv_utils`
- [x] Generator is idempotent; regeneration leaves all non-LDS tables
**byte-identical** (`ARCH_FAMILY_MAP`, `WARP_SUPPORTED_COMBINATIONS`,
`WARP_TILE_SUPPORTED_COMBINATIONS`,
`PRESHUFFLE_WARP_TILE_SUPPORTED_COMBINATIONS`,
`TRAIT_UNSUPPORTED_COMBINATIONS`, `ELEMENT_SIZE_MAP`,
`DTYPE_COMBINATIONS`, `PRESHUFFLE_PIPELINES`)
- [x] C++ compiles standalone and returns values identical to Python for
every architecture and pipeline, single- and double-buffered, including
the unknown-architecture fallback
- [x] `clang-format-18 -style=file` clean on both headers
- [ ] **GPU validation pending** — see below

### Equivalence proof

This edits a shared table, so I enumerated the validator's survivor set
per architecture, before and after. The axes now include **every
pipeline the validators can see** — `comp_async` and `wavelet` as well
as the nine that were in the old table — and both settings of the
ping-pong staging flag.

| arch | develop | this PR | gained | lost |
|---|---:|---:|---:|---:|
| gfx908 | 4,613 | 7,386 | +3,003 | **230** |
| gfx90a | 2,701 | 4,138 | +1,595 | **158** |
| gfx942 | 5,324 | 8,360 | +3,322 | **286** |
| gfx950 | 7,116 | 19,016 | +11,900 | 0 |
| gfx1100 | 1,155 | 1,870 | +770 | **55** |
| gfx1200 | 442 | 668 | +253 | **27** |
| gfx1201 | 442 | 668 | +253 | **27** |
| gfx1250 | 2,012 | 10,926 | +8,914 | 0 |
| **total** | | | **+30,010** | **783** |

**There are losses, and an earlier version of this section claimed there
were none.** That claim was measured over the old parameter space, which
could not contain `comp_async` because the pipeline had no entry in the
old table to enumerate. It was true of what it measured and wrong as a
general statement.

**All 783 losses are `comp_async`, on the six architectures whose budget
is unchanged**, and they are intended. Attribution by pipeline across
those six:

```
comp_async   783
(no other pipeline appears)
```

`comp_async` had no entry before, so it inherited the full-capacity
default. It allocates two LDS buffers unconditionally — `GetSmemSize()`
returns `num_lds_buffers * smem_size` with `num_lds_buffers = 2` — so it
was budgeted for twice the staging it can actually use. A tile such as
conv `128x64x128` at fp16 needs 48 KB of staging and was accepted, while
the kernel would then have asked the hardware for twice that. Rejecting
it is the fix, not a regression.

The two architectures with a raised budget lose nothing, and no default
configuration set in tree generates `comp_async`, so no shipped kernel
disappears.

## Explicitly out of scope

Two other filters in the same validation chain share this defect class
(gfx942-derived constraints applied to every architecture). Neither
rejected anything in the baseline campaign, so I left them alone:

- `TRAIT_UNSUPPORTED_COMBINATIONS` — 10 tuples, no arch key
- `_cshuffle_store_ok` — docstring says "GPU-verified on gfx942",
applied to all architectures

A wider audit of the validation chain found this same pattern in several
more places, including a hand-written warp-tile lookup on the native
side that disagrees with the generated Python table on a majority of
supported architectures. That work is tracked separately in AICK-2264
and AICK-2266, under epic AICK-2265, and is deliberately not part of
this PR: it is a different table on a different axis, and unlike this
change it removes configurations, so it needs its own before/after
enumeration and device validation.

Also left alone: `get_smem_capacity()` in `arch.hpp` special-cases a
single architecture and returns a fixed value otherwise, which disagrees
with `get_lds_size()` for at least one target. It has three live
consumers outside this ticket's scope (`cshuffle_epilogue.hpp`,
`moe_sorting_kernel.hpp`, `grouped_convolution_forward_kernel.hpp`) and
may legitimately encode a per-workgroup addressable limit rather than
per-CU capacity. Flagging rather than changing it.

Author: ozturkosu

dispatcher/codegen/arch_filter.py

@@ -159,7 +159,8 @@ class OperatorType(Enum):
         WARP_TILE_SUPPORTED_COMBINATIONS,
         PRESHUFFLE_WARP_TILE_SUPPORTED_COMBINATIONS,
         PRESHUFFLE_PIPELINES,
-        LDS_CAPACITY_LIMITS,
+        LDS_CAPACITY_LIMITS_BY_ARCH,
+        get_lds_limit,
         TRAIT_UNSUPPORTED_COMBINATIONS,
         DTYPE_COMBINATIONS,
     )
@@ -241,7 +242,43 @@ class OperatorType(Enum):
 
     PRESHUFFLE_PIPELINES = ["preshufflev2"]
 
-    LDS_CAPACITY_LIMITS = {"compv4": 32768, "preshufflev2": 32768, "default": 65536}
+    # Conservative fallback: the historical 64 KB / 32 KB budget, applied to
+    # every architecture. It deliberately understates gfx950 and gfx1250 rather
+    # than overstating anything, because a budget larger than the silicon
+    # produces kernels that cannot launch. The generated module carries the
+    # real per-architecture numbers; regenerate it rather than relying on this.
+    _FALLBACK_LDS_BUDGET = {
+        "mem": 65536,
+        "compv1": 65536,
+        "compv2": 65536,
+        "compv3": 65536,
+        "compv4": 32768,
+        "compv5": 65536,
+        "compv6": 32768,
+        "preshufflev1": 32768,
+        "preshufflev2": 32768,
+        # Mandatory double buffering (num_lds_buffers = 2), so half the budget.
+        "comp_async": 32768,
+        "wavelet": 65536,
+        "default": 65536,
+    }
+
+    LDS_CAPACITY_LIMITS_BY_ARCH = {
+        arch: dict(_FALLBACK_LDS_BUDGET) for arch in ARCH_FAMILY_MAP
+    }
+
+    def get_lds_limit(
+        gpu_arch: str, pipeline: str, double_smem_buffer: bool = False
+    ) -> int:
+        """Get the LDS staging budget in bytes for an architecture and pipeline."""
+        per_pipeline = LDS_CAPACITY_LIMITS_BY_ARCH.get(
+            gpu_arch.lower(), _FALLBACK_LDS_BUDGET
+        )
+        budget = per_pipeline.get(pipeline.lower(), per_pipeline["default"])
+        if double_smem_buffer:
+            # Conservative: the fallback assumes the smallest capacity we ship.
+            budget = min(budget, _FALLBACK_LDS_BUDGET["default"] // 2)
+        return budget
 
     TRAIT_UNSUPPORTED_COMBINATIONS = {
         ("compv3", "cshuffle", "interwave"),
@@ -353,6 +390,11 @@ class KernelConfig:
     epilogue: str = "cshuffle"
     scheduler: str = "intrawave"
 
+    # Ping-pong LDS staging. Only meaningful for the pipelines that make it a
+    # choice (mem, compv3, compv5, compv6); the ones that always double already
+    # carry it in their per-pipeline budget.
+    double_smem_buffer: bool = False
+
     # Layout (for whole-workgroup cover validation)
     layout: str = "rcr"
 
@@ -533,6 +575,7 @@ def is_kernel_valid(
         scheduler: str = "intrawave",
         layout: str = "rcr",
         operator: Optional[OperatorType] = None,
+        double_smem_buffer: bool = False,
     ) -> bool:
         """
         Quick validation check for a kernel configuration.
@@ -544,6 +587,8 @@ def is_kernel_valid(
             warp_tile_m, warp_tile_n, warp_tile_k: Warp tile dimensions
             pipeline, epilogue, scheduler: Kernel traits
             layout: Matrix layout (e.g., "rcr")
+            double_smem_buffer: Ping-pong LDS staging. Halves the staging
+                     budget for the pipelines that make it a choice.
             operator: Operator type (GEMM, CONV_FWD, CONV_BWD_DATA, etc.)
                      Affects validation rules for tile constraints.
                      Defaults to GEMM if not specified.
@@ -568,6 +613,7 @@ def is_kernel_valid(
             epilogue=epilogue.lower(),
             scheduler=scheduler.lower(),
             layout=layout.lower(),
+            double_smem_buffer=double_smem_buffer,
             operator=operator if operator is not None else OperatorType.GEMM,
         )
         return self.validate_kernel(config).valid
@@ -709,17 +755,29 @@ def _validate_lds_capacity(self, config: KernelConfig, result: ValidationResult)
         elem_size_a = ELEMENT_SIZE_MAP.get(config.datatype_a, 2)
         elem_size_b = ELEMENT_SIZE_MAP.get(config.datatype_b, 2)
 
+        # When the B cast policy runs before the LDS write, B is staged as
+        # ADataType rather than BDataType (GetSmemSizeB in
+        # gemm_universal_pipeline_ag_bg_cr_policy.hpp). Charging B at the wider
+        # of the two keeps a mixed-precision pair from being under-counted; for
+        # equal dtypes it is the same number.
+        elem_size_b_staged = max(elem_size_a, elem_size_b)
+
         matrix_a_size = config.tile_m * config.tile_k * elem_size_a
-        matrix_b_size = config.tile_n * config.tile_k * elem_size_b
+        matrix_b_size = config.tile_n * config.tile_k * elem_size_b_staged
         total_lds = matrix_a_size + matrix_b_size
 
-        max_lds = LDS_CAPACITY_LIMITS.get(
-            config.pipeline, LDS_CAPACITY_LIMITS["default"]
+        # The budget depends on the target, not just the pipeline: a tile that
+        # overflows one architecture's LDS may fit comfortably in another's.
+        max_lds = get_lds_limit(
+            self.gpu_arch, config.pipeline, config.double_smem_buffer
         )
 
         if total_lds > max_lds:
+            staging = " double-buffered" if config.double_smem_buffer else ""
             result.add_error(
-                f"LDS capacity exceeded: {total_lds} bytes > {max_lds} bytes limit. "
+                f"LDS capacity exceeded on {self.gpu_arch} "
+                f"(pipeline={config.pipeline}{staging}): "
+                f"{total_lds} bytes > {max_lds} bytes limit. "
                 f"Matrix A: {config.tile_m}x{config.tile_k}x{elem_size_a}={matrix_a_size}B, "
                 f"Matrix B: {config.tile_n}x{config.tile_k}x{elem_size_b}={matrix_b_size}B"
             )

dispatcher/codegen/arch_specs.json

@@ -172,9 +172,57 @@
         "bf8_fp8_fp32": [[16, 16, 16]],
         "int8_int8_int32": [[16, 16, 16]]
       }
+    },
+
+    "gfx1250": {
+      "family": "cdna5",
+      "target_family": "gfx125",
+      "architecture": "rdna",
+      "_architecture_note": [
+        "architecture is the ISA lineage, not the product generation. arch.hpp",
+        "states that gfx1250 shares the RDNA architecture with the GFX12 family",
+        "while being its own standalone target family, because its MMA builtins",
+        "and data-type ABI differ. family records the product generation."
+      ],
+      "description": "AMD gfx1250 (wave32, WMMA rather than MFMA)",
+      "warp_size": 32,
+      "_lds_capacity_kb_source": "get_lds_size(gfx125_t) in include/ck_tile/core/arch/arch.hpp",
+      "lds_capacity_kb": 320,
+      "_warp_tables_provenance": [
+        "warp_configs and warp_tile_combos below are transplanted verbatim from",
+        "arch_specs_generated.py, where gfx1250 was already present and in use.",
+        "They are NOT newly authored here. The generated module had drifted ahead",
+        "of this file; without this entry, regenerating would delete them.",
+        "Any change to these lists belongs to gfx1250 enablement, not here."
+      ],
+      "warp_configs": [
+        [2, 4, 1],
+        [1, 8, 1],
+        [8, 1, 1],
+        [4, 2, 1],
+        [2, 1, 1],
+        [1, 2, 2],
+        [4, 1, 1],
+        [1, 4, 1],
+        [2, 2, 1]
+      ],
+      "_warp_tile_8bit_source": [
+        "The 8-bit shapes are the dense WmmaTraits<gfx125_t, ...> specialisations in",
+        "include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl_8bit_traits.hpp:",
+        "fp8/fp8 and bf8/bf8 each have a 16x16x64 and a 16x16x128 form.",
+        "The mixed fp8/bf8 and bf8/fp8 pairs exist there too but are not listed",
+        "here; enabling them would newly generate kernels for combinations this",
+        "bridge has never exercised, which needs its own device validation."
+      ],
+      "warp_tile_combos": {
+        "fp16_fp16_fp32": [[16, 16, 32]],
+        "bf16_bf16_fp32": [[16, 16, 32]],
+        "fp8_fp8_fp32": [[16, 16, 64], [16, 16, 128]],
+        "bf8_bf8_fp32": [[16, 16, 64], [16, 16, 128]]
+      }
     }
   },
-  
+
   "element_sizes": {
     "fp16": 2,
     "bf16": 2,
@@ -221,18 +269,93 @@
     "c": "ck_tile::tensor_layout::gemm::ColumnMajor"
   },
 
-  "pipeline_lds_limits": {
-    "_comment": "LDS capacity limits in bytes for different pipeline types",
-    "mem": 65536,
-    "compv1": 65536,
-    "compv2": 65536,
-    "compv3": 65536,
-    "compv4": 32768,
-    "compv5": 65536,
-    "compv6": 32768,
-    "preshufflev1": 32768,
-    "preshufflev2": 32768,
-    "default": 65536
+  "pipeline_lds_budget": {
+    "_comment": [
+      "Per-pipeline budget for the A+B staging tiles, resolved at generation",
+      "time against each architecture's lds_capacity_kb.",
+      "",
+      "  basis 'fraction'    -> bytes = floor(lds_capacity_kb * 1024 * value)",
+      "  basis 'absolute_kb' -> bytes = value * 1024, identical on every arch",
+      "",
+      "Use 'absolute_kb' only for a limit that genuinely does not scale with",
+      "how much LDS the silicon has. Anything proportional to capacity must be",
+      "a 'fraction', otherwise a new architecture silently inherits the budget",
+      "of whichever one the constant was written for.",
+      "",
+      "The check models the A+B staging footprint only, matching the kernel's",
+      "own accounting; the epilogue reuses that space rather than adding to it."
+    ],
+
+    "mem": { "basis": "fraction", "value": 1.0 },
+    "compv1": { "basis": "fraction", "value": 1.0 },
+    "compv2": { "basis": "fraction", "value": 1.0 },
+    "compv3": { "basis": "fraction", "value": 1.0 },
+    "compv5": { "basis": "fraction", "value": 1.0 },
+
+    "compv4": {
+      "basis": "fraction",
+      "value": 0.5,
+      "_reason": [
+        "Double-buffered. GemmPipelineAgBgCrCompV4::GetSmemSize() returns",
+        "2 * Policy::GetSmemSize<Problem>(), so the kernel allocates two",
+        "staging buffers and only half the capacity is available to A+B.",
+        "Half-of-capacity is therefore the exact model, not a safety margin."
+      ]
+    },
+    "preshufflev2": {
+      "basis": "fraction",
+      "value": 0.5,
+      "_reason": [
+        "Double-buffered when DoubleSmemBuffer is set:",
+        "WeightPreshufflePipelineAGmemBGmemCRegV2::GetSmemSize() returns",
+        "DoubleSmemBuffer ? 2 * smem_size : smem_size. Budgeted for the",
+        "doubling case, which is the one that can overflow."
+      ]
+    },
+    "comp_async": {
+      "basis": "fraction",
+      "value": 0.5,
+      "_reason": [
+        "Mandatory double buffering, no configuration choice:",
+        "GemmPipelineAgBgCrCompAsync::GetSmemSize() returns",
+        "num_lds_buffers * smem_size with num_lds_buffers = 2.",
+        "Has no Pipeline enumerator on the C++ side, so this entry only",
+        "affects the Python validator; it is emitted to C++ only for",
+        "pipelines that exist in that enum."
+      ]
+    },
+    "wavelet": {
+      "basis": "fraction",
+      "value": 1.0,
+      "_reason": [
+        "Single-buffered. Listed explicitly so that every pipeline the",
+        "validators can see has a deliberate entry rather than silently",
+        "inheriting 'default'."
+      ]
+    },
+
+    "compv6": {
+      "basis": "fraction",
+      "value": 0.5,
+      "_reason": [
+        "NOT double-buffered: GemmPipelineAgBgCrCompV6::GetSmemSize() returns",
+        "the policy size unmultiplied. The historical halved cap is therefore",
+        "unexplained by buffering. Its ratio is preserved rather than widened,",
+        "because no evidence supports a larger budget; raising it needs a",
+        "measurement, not an assumption."
+      ]
+    },
+    "preshufflev1": {
+      "basis": "fraction",
+      "value": 0.5,
+      "_reason": [
+        "Vestigial. No WeightPreshufflePipelineAGmemBGmemCRegV1 exists and",
+        "preshuffle_pipelines.supported lists only preshufflev2, so nothing",
+        "generates this pipeline. Ratio preserved so the entry stays inert."
+      ]
+    },
+
+    "default": { "basis": "fraction", "value": 1.0 }
   },
 
   "unsupported_trait_combos": {