restructure gather fixtures: parallel gather_kernel becomes default - #116
restructure gather fixtures: parallel gather_kernel becomes default#116lasch wants to merge 1 commit into
Conversation
Adds new gather_kernel, old default to gather_kernel_1core Adds gather_kernel, a row-tiled/parallel variant, and makes it the fixture's default so DistributeWork has real work to do in the common case. gather_kernel is renamed to gather_kernel_1core (base: None, unchanged in meaning); the six column-slice edge variants rebase onto the new default and become parallel; large_k stays single-program to avoid changing its fan-out semantics. Updates fixtures/README.md and fixtures/gather/README.md to match, and regenerates the CI-gated docs/patterns/memory.md. Signed-off-by: Lars Schneidenbach <schneidenbach@us.ibm.com> Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
fc48b68 to
2197dc6
Compare
| "1core": { | ||
| # gather_kernel_1core: the one variant in this fixture that pins | ||
| # the no-scf.for shape — a single descriptor_gather call consumes | ||
| # the whole index array in one shot, no tl.program_id, no row |
There was a problem hiding this comment.
i think we can reduce the amount of comments here, keep it short like 1-2 lines
| }, | ||
| "tags": ["descriptor-gather"], | ||
| "tags": ["descriptor-gather", "1core"], | ||
| "summary": ( |
There was a problem hiding this comment.
just say what it is no need to sound apologetic.
| # alone wouldn't catch a bug that triggers only at the | ||
| # smallest legal block. Allows duplicates so an aliasing bug | ||
| # at the minimum size shows up. | ||
| # Smallest legal sizes per the verifier: BLOCK_ROWS=8 (verifier |
There was a problem hiding this comment.
BLOCK_ROWS=8 is locked in because of a check that no longer applies to Spyre.**
meta.py:1077 justifies BLOCK_ROWS=8 with "Triton frontend requires
x_offsets.shape[0] >= 8 for tt.descriptor_gather". That was true for NVIDIA
TMA, but semantic.py:1163-1164 shows the check is now guarded by
if not target_info.is_spyre() — the minimum was released for Spyre when the
guard was added. All seven new gather_kernel variants are pinned to 8 for no
reason. The fixture could use smaller values to cover more of the space (e.g.
BLOCK_ROWS=4 for a two-row tile) and the comment should be corrected.
(meta.py:1075-1078, semantic.py:1163)
2 participants
Summary
Adds new gather_kernel, old default to gather_kernel_1core
Adds gather_kernel, a row-tiled/parallel variant, and makes it the
fixture's default so DistributeWork has real work to do in the common
case. gather_kernel is renamed to gather_kernel_1core (base: None,
unchanged in meaning); the six column-slice edge variants rebase onto
the new default and become parallel; large_k stays single-program to
avoid changing its fan-out semantics.
Updates fixtures/README.md and fixtures/gather/README.md to match,
and regenerates the CI-gated docs/patterns/memory.md.
Note: if the scope of #11 includes 3D+ kernels, then this is not yet fully closing issue #11
This is only doing the first few gather kernels and their
parallelissue.Scope
Testing
Commands and results:
Checklist
third_party/spyre, I explained why in the summary.