restructure gather fixtures: parallel gather_kernel becomes default

third_party/spyre/docs/patterns/memory.md

@@ -129,13 +129,13 @@ Expected diagnostics:
 
 **Round-trip evidence**
 
-- `gather::default` — M=1024, N=64, K_INDICES=32, BLOCK_COLS=32, y_offset=16
+- `gather::default` — M=1024, N=64, K_INDICES=256, BLOCK_ROWS=8, BLOCK_COLS=32, y_offset=16
+- `gather::1core` — M=1024, N=64, K_INDICES=32, BLOCK_COLS=32, y_offset=16 (also demonstrates: 1core)
 - `gather::2d` — M=1024, N=128, K_INDICES=64, BLOCK_ROWS=8, BLOCK_COLS=16 (also demonstrates: program-id-2d, num-programs-fold)
 - `gather::2d_serial`
 - `gather::1d` — K=1024, K_INDICES=256, BLOCK_ROWS=8 (also demonstrates: 1d-source)
-- `gather::2d_large_table_serial`
 
-_+ 4 more variants_
+_+ 5 more variants_
 
 ## descriptor-gather-2d-indices
 

third_party/spyre/test/fixtures/README.md

@@ -32,9 +32,20 @@ reference oracle and input generator. Different functions
   partial-override rule for nested fields; a variant that wants to
   change `constexpr` (or `params`) replaces the whole list / dict.
 - The base is `"default"` unless the variant declares `"base": "<name>"`,
-  which names another variant in the same `VARIANTS` dict. The `"base"`
-  key is consumed at load time and does not appear in the registry entry.
-  Circular chains are caught at collection time.
+  which names another variant in the same `VARIANTS` dict, or
+  `"base": None`, which opts out of the fallback entirely (the variant's
+  dict is used as-is, with no merge). `"base"` is consumed at load time
+  and does not appear in the registry entry (`"base": None` is the one
+  exception — see the field reference below). Circular chains are caught
+  at collection time.
+- The `"default"`-fallback is implicit and applies to *every* variant that
+  omits `"base"`, not just ones that look like they want to inherit
+  something. If `"default"` sets a field like `parallel: False` with a
+  `summary` justifying it, every sibling variant silently inherits that
+  justification too unless it sets its own or opts out with
+  `"base": None`. Prefer keeping `"default"` a representative, unexotic
+  case for exactly this reason — a `"default"` that is itself an edge
+  case makes every accidental non-opt-out sibling misrepresent itself.
 - Registry keys: `<folder>` for the default variant and for
   single-variant kernels; `<folder>__<variant>` for every other entry.
   e.g. `vector_add`, `vector_add__dynamic`.
@@ -46,15 +57,17 @@ reference oracle and input generator. Different functions
 | `kernel_fn` | `@triton.jit` function | Compiled on demand via `compile_to_ttir` → `make_ktir_mod`. |
 | module-level `SIGNATURE` | `dict[str, str]` | Dtype per `@triton.jit` arg. Pure types — no values. Declared at module scope in `meta.py`, not inside `VARIANTS`. Used by every variant that doesn't redeclare it. |
 | variant `SIGNATURE` | `dict[str, str]` | Optional per-variant override. Replaces the module-level map wholesale — use when the variant's kernel has a different arg list (e.g. softmax's `multi_tile` has `BLOCK_N` where `single_tile` has `BLOCK_SIZE`). |
-| `base` | `str` | Optional. Name of another variant in the same `VARIANTS` dict to use as the merge base instead of `"default"`. Consumed at load time; not stored in the registry entry. |
+| `base` | `str \| None` | Optional. Name of another variant in the same `VARIANTS` dict to use as the merge base instead of `"default"`. Consumed at load time; not stored in the registry entry. `None` opts out of the implicit `"default"`-fallback entirely — use this when the variant must not inherit something `"default"` sets (e.g. a `SIGNATURE` it doesn't share, or a `parallel`/`summary` pair that wouldn't apply to it). Unlike a named base, an explicit `"base": None` currently *does* survive into the resolved registry entry as a stray key (a minor asymmetry in the loader — harmless since nothing reads it back). |
 | `constexpr` | `list[str]` | Which arg names are Triton constexprs for this variant. Each variant declares the full list explicitly (no partial override over default's list). Values for constexprs come from `params`. |
 | `params` | `dict[str, list[Any]]` | Single source of truth for argument values. Lists today carry one element each; future Cartesian expansion (one registry entry per product) is deferred — when it lands, the `constexpr` vs runtime partition stays the same per expansion. |
 | `grid` | `list[int]` | Per-axis partition of the 32-core Spyre grid. One entry per `tl.program_id` axis the kernel reads; `prod(grid)` equals the hardware core count. Defaults to the backend's `(32,)` (1D on all cores) when omitted. |
 | `reference` | `(inputs) -> np.ndarray` | NumPy oracle for the numerical test. Omit for structure-only variants. Defined alongside `VARIANTS` in the same `meta.py`. |
 | `inputs` | `(**param_values) -> {"arg_name": np.array, ...}` | Pointer/tensor input generator. Called with kwargs matching `params` keys; returns pointer/tensor args only. Runtime scalars (params that aren't in `constexpr`) are merged in by the framework. |
 | `output_key` | `str` | Which `inputs` key holds the output buffer compared against `reference(inputs)`. |
 | `func_name` | `str` | KTIR function name for `ktir_cpu`. Defaults to `kernel_fn.__name__`. |
-| `parallel` | `bool`, default `True` | Set `False` for single-program kernels that do not call `tl.program_id` — skips the DistributeWork-presence check in `TestExample`. |
+| `parallel` | `bool`, default `True` | Set `False` for single-program kernels that do not call `tl.program_id` — skips the DistributeWork-presence check in `TestExample`. When `False`, also set `summary` (below) to justify it — and remember `"default"`-fallback means every sibling without its own `"base"` inherits both fields together unless it opts out. |
+| `tags` | `list[str]` | Free-form categorization strings (e.g. `"descriptor-gather"`, `"1core"`). Consumed by `scripts/gen_patterns_docs.py` to group variants into the generated `docs/patterns/*.md` — see that script's `--check` mode, which CI runs to catch stale generated docs. |
+| `summary` | `str` | Human-readable justification, surfaced in the `pytest.skip` reason when `parallel: False` (see `test_work_distribution` in `test_ktir_examples.py`). Explains *why* the variant is legitimately single-program rather than leaving that to be inferred. Because of the `"default"`-fallback rule above, a `summary` set on `"default"` is inherited by every sibling that doesn't set its own — which is exactly the failure mode that motivated documenting `"base": None` as an opt-out. |
 | `extra_checks` | `(tester) -> None` | Optional. Runs alongside the shared structural suite for variant-specific assertions (e.g. `memref<?x` only in the dynamic variant). |
 | `xfail_numerical` | `str \| dict` | Optional. `str` is shorthand for `{"reason": str, "strict": True}`; `dict` is forwarded to `pytest.mark.xfail(**d)` (so `raises=ValueError` etc. work). Attached at collection time so failures show as `XFAIL`, not `SKIP`. Use this when the kernel compiles but the numerical comparison fails (e.g. `ktir_cpu` can't parse a dynamic memref shape). |
 | `disabled` | `dict` | Optional. `{"reason": str, "tracking_test": "file.py::ClassName"}`. Marks a variant as unable to compile through the TTIR→KTIR pipeline today. Every structural and numerical test skips with `reason`. `tracking_test` points at the single-pass test that pins the underlying gap (e.g. a `test_lower_desc_memory.py` class asserting the expected verification failure). The meta-test `test_disabled_variants_tracking_tests_exist` fails if `tracking_test` no longer resolves, so a closed gap can't leave a stale `disabled` block behind. Use this instead of `xfail` when the kernel does not yet compile — it keeps the failure documented in one place (the tracking test) rather than duplicated across every structural test. |

third_party/spyre/test/fixtures/gather/README.md

@@ -1,44 +1,69 @@
 # gather
 
 End-to-end test fixture for **`tl.descriptor_gather`** — Triton's
-indirect row-indexed load. The fixture carries two `@triton.jit`
+indirect row-indexed load. The fixture carries three `@triton.jit`
 functions sharing one source file:
 
-- **`gather_kernel`** — single-program. One kernel invocation pulls
-  `K_INDICES` rows from a 2D source matrix into a contiguous
-  `[K_INDICES, BLOCK_COLS]` output tile. Implements
+- **`gather_kernel`** — the fixture's `default`. Fixed
+  `[y_offset, y_offset + BLOCK_COLS)` column slice, same row-level
+  contract as `gather_kernel_1core` below, but the `K_INDICES` rows are
+  tiled `BLOCK_ROWS` at a time and distributed across a 1D Spyre core
+  grid. Implements
   `out[i, :] = in[idx[i], y_offset : y_offset + BLOCK_COLS]`.
+- **`gather_kernel_1core`** — single-program. One kernel invocation pulls
+  `K_INDICES` rows from a 2D source matrix into a contiguous
+  `[K_INDICES, BLOCK_COLS]` output tile in one `descriptor_gather` call,
+  no `tl.program_id`. Same row-level contract as `gather_kernel`, minus
+  the distribution — the simplest kernel in the file.
 - **`gather_2d_kernel`** — tiled across a 2D Spyre core grid. Each
   program instance `(pid_m, pid_n)` produces one
   `BLOCK_ROWS × BLOCK_COLS` tile of the output; together the cores
   materialize the full `[K_INDICES, N]` result. Implements
-  `out[i, j] = in[idx[i], j]`.
+  `out[i, j] = in[idx[i], j]` — no `y_offset`; gathers the full row
+  width by column-tiling instead.
 
-Both back the same downstream pattern: **embedding lookups** and
+All three back the same downstream pattern: **embedding lookups** and
 **indirect row-gather access into a 2D source**.
 
 ### Pythonic semantics
 
-The two kernels differ in *what* they gather, not in the underlying
+The three kernels differ in *what* they gather, not in the underlying
 `tt.descriptor_gather` mechanism. The mechanism takes two coords —
 `x_offsets` (row indices, fanned out across `BLOCK_ROWS`) and
 `y_offset` (a scalar column start) — and returns a
-`[BLOCK_ROWS, BLOCK_COLS]` tile from the source. The two kernels
+`[BLOCK_ROWS, BLOCK_COLS]` tile from the source. The three kernels
 expose that primitive in different ways.
 
 ```python
-# gather_kernel  — single-program, fixed column slice
+# gather_kernel_1core — single-program, fixed column slice
 for i in range(K_INDICES):
     out[i, :] = in[idx[i], y_offset : y_offset + BLOCK_COLS]
 # out shape: [K_INDICES, BLOCK_COLS]
 # y_offset is a runtime kernel argument; chosen once per launch.
 ```
 
-`gather_kernel` takes a fixed `[y_offset, y_offset + BLOCK_COLS)`
+`gather_kernel_1core` takes a fixed `[y_offset, y_offset + BLOCK_COLS)`
 column window of every gathered row — so `BLOCK_COLS` is part of the
 output shape and `y_offset` flows directly through to the underlying
 gather op as a kernel arg.
 
+`gather_kernel` implements the *exact same* row-level loop body
+— same fixed column window, same `y_offset` kernel argument — but
+splits the `range(K_INDICES)` loop into `BLOCK_ROWS`-sized chunks and
+assigns chunks to cores via `tl.program_id(0)`, so each core's
+`descriptor_gather` call only ever touches its own row tile:
+
+```python
+# gather_kernel — same row-level contract as gather_kernel_1core,
+# distributed across cores in BLOCK_ROWS-row chunks
+for m_block in range(m_start, m_end):           # ← this core's chunk
+    offset_m = m_block * BLOCK_ROWS
+    out[offset_m : offset_m + BLOCK_ROWS, :] = \
+        in[idx[offset_m : offset_m + BLOCK_ROWS],
+           y_offset : y_offset + BLOCK_COLS]
+# out shape: [K_INDICES, BLOCK_COLS] — same as gather_kernel_1core
+```
+
 `gather_2d_kernel` writes the **full row width** of every gathered
 row, but it builds that row by walking `BLOCK_COLS`-wide column tiles
 in an inner loop — i.e. it calls the same gather primitive multiple
@@ -67,22 +92,33 @@ for m_block in range(K_INDICES // BLOCK_ROWS):
 ```
 
 Notice that the body of the inner loop is *exactly* what
-`gather_kernel` does in one shot — same `x_offsets`/`y_offset` call
+`gather_kernel_1core` does in one shot — same `x_offsets`/`y_offset` call
 into the same underlying `tt.descriptor_gather`. The 2D kernel
 schedules many such calls (one per `(m_block, n_block)` tile, sharded
 across cores) so that the union of their outputs covers the full
 `[K_INDICES, N]` matrix.
 
 ## Why this fixture exists
 
-`gather_kernel` pins the `tt.descriptor_gather` →
-`ktdp.construct_indirect_access_tile` lowering at the simplest possible
-shape — no `tl.program_id`, one descriptor_gather over the whole
-output, `DistributeWork` is a no-op. Seven variants cover edge cases
-of the column-slice machinery (`y_offset = 0`, full-row, minimum legal
-block sizes, slice ending at `N`, wider slice, larger fan-out).
-
-`gather_2d_kernel` adds coverage the single-program kernel cannot:
+`gather_kernel_1core` (variant `1core`) pins the `tt.descriptor_gather`
+→ `ktdp.construct_indirect_access_tile` lowering at the simplest possible
+shape — no `tl.program_id`, one `descriptor_gather` over the whole
+output, `DistributeWork` is a no-op. Its only sibling, `large_k`, pins a
+larger single-shot fan-out; it is kept single-program rather than
+rebased onto the parallel kernel below, since row-tiling would turn its
+per-gather fan-out into `BLOCK_ROWS` and silently change what it tests.
+
+`gather_kernel` (the `default`) pins the same lowering
+*distributed*: `tl.program_id(0)` tiles `K_INDICES` into `BLOCK_ROWS`-row
+chunks across the core grid, so `DistributeWork` has real work to do —
+unlike `gather_kernel_1core`'s no-op case — while every gathered row still
+takes the same fixed `y_offset` column slice. Six variants cover edge
+cases of the column-slice machinery (`y_offset = 0`, full-row, a wider
+full-row, minimum legal block sizes, slice ending at `N`, wider slice);
+all are rebased onto this kernel, so they are parallel too.
+
+`gather_2d_kernel` adds coverage neither single-program kernel nor
+`gather_kernel` can:
 
 1. **Two-axis `tl.program_id`.** Most other fixtures use a 1D grid;
    this is the first to exercise a 2D grid (`[4, 8]`) with
@@ -106,22 +142,39 @@ the degenerate path where `rows_per_core = m_blocks` and
 
 ## Variants
 
-### Single-program (`gather_kernel`)
-
-| Variant            | M    | N   | K_INDICES | BLOCK_COLS | y_offset | y_off+BLOCK | dups | Pinned bug class                                |
-|--------------------|------|-----|-----------|------------|----------|-------------|------|-------------------------------------------------|
-| `default`          | 1024 | 64  | 32        | 32         | 16       | 48          | no   | sanity, non-zero offset, slice strictly inside  |
-| `y_offset_zero`    |  256 | 32  | 16        | 16         |  0       | 16          | no   | `y_offset = 0` path                             |
-| `full_row`         |  128 | 16  | 16        | 16         |  0       | 16          | no   | `BLOCK_COLS == N` (full-row gather)             |
-| `min_block_cols`   |   64 | 64  |  8        |  8         | 32       | 40          | yes  | verifier minimums (`K=8`, `BLOCK_COLS=8`)       |
-| `slice_at_end`     |  256 | 64  | 16        | 16         | 48       | 64          | no   | slice ends exactly at column `N` (off-by-one)   |
-| `wide_slice`       |  128 | 256 | 16        | 128        | 64       | 192         | no   | larger `BLOCK_COLS` (size-dependent bugs)       |
-| `large_k`          |  512 | 64  | 128       | 32         | 16       | 48          | yes  | larger fan-out + duplicate indices              |
-
-All seven set `parallel: False` — `DistributeWork` is a no-op, so
+### Parallel (`gather_kernel`)
+
+| Variant           | M    | N   | K_INDICES | BLOCK_ROWS | BLOCK_COLS | y_offset | y_off+BLOCK | dups | Pinned bug class                                     |
+|-------------------|------|-----|-----------|------------|------------|----------|-------------|------|-------------------------------------------------------|
+| `default`         | 1024 | 64  | 256       | 8          | 32         | 16       | 48          | no   | sanity, non-zero offset, slice strictly inside         |
+| `y_offset_zero`   |  256 | 32  |  16       | 8          | 16         |  0       | 16          | no   | `y_offset = 0` path                                    |
+| `full_row`        |  128 | 16  |  16       | 8          | 16         |  0       | 16          | no   | `BLOCK_COLS == N` (full-row gather)                    |
+| `slice_large_row` |  128 | 256 |  16       | 8          | 256        |  0       | 256         | no   | `full_row` at a 16× wider embedding dim                |
+| `min_block_cols`  |   64 | 64  |   8       | 8          | 8          | 32       | 40          | yes  | verifier minimums (`BLOCK_ROWS=8`, `BLOCK_COLS=8`)     |
+| `slice_at_end`    |  256 | 64  |  16       | 8          | 16         | 48       | 64          | no   | slice ends exactly at column `N` (off-by-one)          |
+| `wide_slice`      |  128 | 256 |  16       | 8          | 128        | 64       | 192         | no   | larger `BLOCK_COLS` (size-dependent bugs)              |
+
+All seven run on `grid=[32]` and set `parallel: True` — each core
+gathers its own `BLOCK_ROWS`-row chunk via `tl.program_id(0)`.
+`min_block_cols` reduces to one busy core (`K_INDICES == BLOCK_ROWS`,
+so `m_blocks=1`); still legitimately `parallel: True`, since
+`DistributeWork` emits `ktdp.get_compute_tile_id` and the distribution
+`scf.for` regardless of how many cores end up with work to do.
+
+### Single-program (`gather_kernel_1core`)
+
+| Variant   | M    | N  | K_INDICES | BLOCK_COLS | y_offset | y_off+BLOCK | dups | Pinned bug class                        |
+|-----------|------|----|-----------|------------|----------|-------------|------|--------------------------------------------|
+| `1core`   | 1024 | 64 | 32        | 32         | 16       | 48          | no   | sanity — no `scf.for`, one-shot gather      |
+| `large_k` |  512 | 64 | 128       | 32         | 16       | 48          | yes  | larger fan-out + duplicate indices          |
+
+Both set `parallel: False` — `DistributeWork` is a no-op, so
 `test_work_distribution` is skipped. The whole index array is consumed
 in one `descriptor_gather`, so `K_INDICES` is the *total* number of
-rows gathered (not a tile size).
+rows gathered (not a tile size). `large_k` stays on this kernel rather
+than being rebased onto `gather_kernel`: row-tiling would turn
+its per-gather fan-out into `BLOCK_ROWS`, silently changing what it
+pins.
 
 ### 2D-tiled (`gather_2d_kernel`)
 
@@ -152,25 +205,40 @@ through Spyre today:
 idx = tl.load(idx_ptr + tl.arange(0, K_INDICES))
 ```
 
-Both kernels therefore load the index tensor via a 1D
+All three kernels therefore load the index tensor via a 1D
 `tl.make_tensor_descriptor`, which lowers cleanly through
 `LowerDescriptorMemory`.
 
 ## Preconditions
 
-### Single-program kernel (unchecked → enforce in variant params)
+### Single-program kernel — `gather_kernel_1core` (unchecked → enforce in variant params)
 
 Verifier rules from `tt.descriptor_gather` (the Triton frontend at
 `python/triton/language/semantic.py:descriptor_gather`):
 
 - The source descriptor's `block_shape` leading dim is exactly **1**.
-- `K_INDICES ≥ 8`.
+- `K_INDICES ≥ 8` — the whole index array is loaded as one `x_offsets`
+  tile, so the verifier's minimum binds directly on `K_INDICES`.
 - `BLOCK_COLS ≥ 32 / bitwidth * 8` (i.e. ≥ **8** for f32, ≥ **16** for f16).
 - `BLOCK_COLS` is a power of two.
 - `y_offset + BLOCK_COLS ≤ N` (slice fits in the source row — the
   kernel does not zero-pad).
 
-### 2D kernel (unchecked → enforce in variant params)
+### Parallel kernel — `gather_kernel` (unchecked → enforce in variant params)
+
+Same verifier rules apply, but row-tiling shifts where the `x_offsets`
+minimum binds:
+
+- `BLOCK_ROWS ≥ 8` — each core loads a `[BLOCK_ROWS]` index tile per
+  gather call, so the verifier's minimum binds on `BLOCK_ROWS`, not
+  `K_INDICES`.
+- `K_INDICES % BLOCK_ROWS == 0` (tiles exactly cover the index array —
+  no masking).
+- `BLOCK_COLS ≥ 32 / bitwidth * 8` and a power of two (same as the
+  single-program kernel).
+- `y_offset + BLOCK_COLS ≤ N` (same as the single-program kernel).
+
+### 2D kernel — `gather_2d_kernel` (unchecked → enforce in variant params)
 
 In addition to the verifier rules above (with `BLOCK_COLS`'s size and
 power-of-two constraints), the 2D kernel assumes: