Investigate making the gather examples parallel

[fabianlim]

Summary

The gather fixture has 22 parallel: False variants, all skipping
test_work_distribution. Every other fixture (matmul, softmax,
vector_add) always uses tl.program_id and never needs this flag.
The parallel flag only exists in gather because gather_kernel was
written as a single-program kernel — an inconsistency that should be
resolved.

Goal

Make the gather fixture consistent with every other fixture: all
kernels use tl.program_id, all variants run test_work_distribution,
and the parallel flag goes away entirely.

This requires deciding for each kernel whether:

1. It should be rewritten to use tl.program_id (preferred — matches
   all other fixtures), or
2. It genuinely tests a single-program lowering that has no parallel
   equivalent, in which case the skip is legitimate and the variant
   should be removed or kept with a clear justification.

Questions to answer

• Does gather_kernel (the original single-program kernel) still serve
  a purpose distinct from gather_2d_kernel, or can its variants be
  replaced by parallel equivalents?
• Do the *_serial variants (2d_serial, 2d_large_table_serial)
  test anything not already covered by their parallel twins?

Acceptance

• No parallel: False entries remain in fixtures/gather/meta.py, or
  each remaining one has an explicit justification and a decision was
  made not to remove it.
• test_work_distribution skip count drops to zero for gather variants.
• test_numerical still green for all variants.

Pointers

• third_party/spyre/test/fixtures/gather/kernel.py
• third_party/spyre/test/fixtures/gather/meta.py
• third_party/spyre/test/test_ktir_examples.py — skip site (line 307)

[fabianlim]

@lasch these tests are fixture tests, so they are not candidate for #55, as those are limited to transform level tests.

[lasch]

Does gather_kernel (the original single-program kernel) still serve
a purpose distinct from gather_2d_kernel, or can its variants be
replaced by parallel equivalents?

I'd say, we keep a distinct gather_kernel for the following reason(s):

• keeps the lowering at its simplest shape
• other edge-case variants use a dedicated y_offset too
• tiling in gather_2d_kernel can hardly express it (or at least it's harder to debug)

Plan:

• add one-line justification comment to default's "parallel": False

@fabianlim let me know if that aligns with your thoughts, too

[fabianlim]

@lasch Initially we wanted to not do this, because we wanted our fixtures to be representative of spyre kernels. Since spyre has a fixed number of cores, we cannot have kernels with symbolic shapes, where if we did not have the scf.for iterations we will not be able to cover the extent of input data, so that was the reason,

but having said that, we have now found a use case in this PR #105, where I have now introduced such a kernel, because the current state of the dataflow-scheduler does not yet handle scf.for.

So im thinking yes maybe we could, and if we have a better way to document this parallel: False is one way to (or maybe we can rename parallel to something else to make it more clear why we choose to have these fixtures

[fabianlim]

@lasch Agreed on keeping a distinct gather_kernel. On the rename — not important, and I'd now say don't rename it at all. The reason no good word turned up is that the flag is being asked to do two jobs: gate a test, and explain that the example is deliberately simple. So keep parallel for the first job only, and add a tag for the second:

1. Leave "parallel" exactly as it is. It's an explicit, greppable predicate for test_work_distribution and that's all it needs to mean.
2. Tag the semantic separately: "tags": ["descriptor-gather", "simplified"], or "simplified:single-program" if the reason is worth carrying. Nothing validates tags, so there's no registry to update — and it scales when the next example is simplified for a different reason.
3. Put your one-line justification in summary rather than a comment, so it sits with the variant and can be surfaced in the skip message.

That way parallel stops carrying an implied "…and here's why", which is the part that reads as a TODO today.

Leave direct_access_tile alone — it looks similar but it's a fact about what survives the lowering, not about how the example was written. Separate issue if we want it.

[lasch]

The current hierarchy of fixtures poses a new question how to handle parallel and summary with respect to the default-fixture.

• the default fixture is gather_kernel and thus is parallel=False and needs a summary=<justification>
• However, this causes every other fixture that doesn't have an explicit base to use these defaults
  • e.g. 3d_group has no base entry and sets parallel=False. Therefore, without being careful, it would just inherit the summary=<justification> from the default (implicit parent) unless we also set summary=<3d_group_justification>.
• dropping the default summary justification is also not a good idea because it has parallel=False and thus needs a justification

TL;DR: if fixtures have no explicitly set summary, then those fixtures will happily come with the default justification in the parent summary. Is that something we can live with or shall we restructure the hierarchy of the fixtures such that

[fabianlim]

@lasch you are having this problem because the faulty fixture is your default fixture. I see you have 2 options

1. you fix the gather_kernel and remove the parallel = False
2. you add a proper gather fixture that as no simplifications and make that the default fixture, then we keep the parallel = False fixtures as derivations of that default

[lasch]

@fabianlim calling it 'faulty fixture' is a contradiction of what we discussed above (the way I understand it). Therefore your option 1 would need clarification or revision of earlier agreements. Option 2 is what I asked for ('... shall we restructure the hierarchy...')

[fabianlim]

@lasch

herefore your option 1 would need clarification or revision of earlier agreements.

Yes sorry my bad wrong choice of words.

Option 2 is what I asked for ('... shall we restructure the hierarchy...')

I think this is the better way

[lasch]

How about this consideration:

• add: gather_basic_kernel as new 'default'
  • same row-level contract as gather_kernel (fixed [y_offset, y_offset+BLOCK_COLS) slice per row)
  • but row-tiled via tl.program_id(0) across the core grid (parallel=True)
• current gather_kernel becomes a derived fixture renamed to single_program and stays parallel=False
• currently derived kernels y_offset_zero, full_row, slice_large_row, etc. still derive from default and therefore become parallel=True variants now
• only large_k remains child of single_program

[fabianlim]

@lasch what you propose seems reasonable, but it will be easier for me to check once you have the meta.py drafted out in a PR.

• But i will prefer to call item 1 gather_kernel_1core - you can look at the existing fixtures (e.g., matmul) and we do not have this 'basic' naming convention. Basic is not very descriptive. Something that 1core or something is better