Skip to content

feat(torchwave): Per-column decomposition and chain gather folding (#18864) - #18864

Open
oerling wants to merge 2 commits into
facebookincubator:mainfrom
oerling:export-D116667979
Open

feat(torchwave): Per-column decomposition and chain gather folding (#18864)#18864
oerling wants to merge 2 commits into
facebookincubator:mainfrom
oerling:export-D116667979

Conversation

@oerling

@oerling oerling commented Sep 5, 2026

Copy link
Copy Markdown
Contributor

Summary:

Decomposes list-producing ops into per-tensor nodes, so each column is a value
the compiler can see: consumer counting, aliasing, CSE and cost-based block
shares all work per column rather than per bundle. On top of that, a column
folds its producer's gather into its own, so a run of subset ops
(batch_flip_and_truncate_sparse, grouped_masked_select_jagged_1d,
group_length_guard_sparse) becomes one gather over the original source with
nothing materialized in between.

Metadata::decompose plus decomposeListOps give one traversal that calls each
op's own rewrite rule, after CSE (cheaper on the bundled form) and before
partitioning, which the per-column nodes exist to inform.

A per-column gather is spelled as a chain of steps -- kStepRange for a flip,
kStepSelect for a select -- carried in step_descs / step_scalars with each
step's own tensors. A column absorbing its producer takes that chain whole: its
own step becomes outermost and the producer's steps append after it with their
tensor and scalar bases rebased. Which of the three getters runs then depends on
whether the composed chain contains a range ANYWHERE, not on what its outermost
step is. A select over a flip starts with a select and still needs the row walk,
because a range is defined per row.

Absorbing is refused unless the rows line up, which is a kernel constraint and
not a formality: prepareRowSteps prepares every step at one row number taken
from the outer walk, so row r of the outer has to be row r of the inner. Two
columns of one flip head hold the same lengths when they were built from the
same input lengths under the same cap -- and because a run of flips feeds each
column's new lengths to the next flip as that column's input lengths, that test
recurses down the run instead of bottoming out at tensor identity. Identity
alone decides only the last link and refuses everything deeper.

group_length_guard_sparse's final stage is a range step written twice. For
output element idx of row r it computes

colOffsets[r - 1] + (idx - resultOffsets[r - 1])

which is what rangeRowDescriptor computes for a range at mode 0 with no cap
over those same two offset arrays. Saying so -- one
tw.chain_gather_guard_column per column in place of the list-form final --
lets a consumer that already has a range absorb the guard as a SUB-LEVEL of its
own rather than as a second step, so the intermediate is never written and the
depth does not grow. The row proof is the head's: element 1 of its output list
is the offsets a chain over those rows walks, so element 0 is the lengths they
were scanned from, which is what chainRowsOf now returns. That arm goes ahead
of the flip head's, because N + 2 is a multiple of three for a seven-column
guard and only the flip arm's target check would otherwise keep it out.

The guard's getter is its own registration rather than a reuse of either
neighbour, for two reasons and one trap. Every column of one guard keeps the
same count -- the per-row minimum summed over rows -- which the head already
has as a scalar, where the flip reads a per-column slice of cumulative lengths.
And the head writes row ENDS, as an inclusive scan, so start_offsets is 0
where a select's offsets stage passes 1. The trap is sizeOrdinal: the select's
registration names the step tensors, which for a guard column are the row
offsets, so the grid would be sized by the ROW count -- three orders of
magnitude short of the output on a real batch. Naming nothing is right here,
because the default kMax over the inputs picks the source, which is at least as
long as the output.

WaveConfig::foldSharedChains folds a producer into every consumer that can
absorb one rather than only into a sole reader. Two foldable consumers of one
column can never satisfy the sole-reader rule: whichever rewrite runs first sees
the other as an outside reader and declines, and the second then finds a gather
already reading the buffer, so the column is materialized and NEITHER folds.
Deciding that needs readers counted through prim.ListPack / prim.ListUnpack
to the ops that really read the data, since a pack is a live user that reads
nothing and a pack immediately unpacked is the identity.

It defaults ON, which only measurement could settle: a reader that folds beside
one that then declines for its own reasons leaves the buffer AND duplicates the
work. On the ROO preproc it takes every column it is offered and leaves none
behind -- see the numbers below. sharedChain now forces it OFF rather than
taking the default, so the unfolded arm stays covered.

RangeRowDesc carries its within-row fields in 32 bits. Only the two origins
are offsets into a whole column and need 64; outLen and the valid window are
bounded by the step's maxLen, and dir is +1 or -1. At four steps per chain the
struct is copied per row, so at 48 bytes it was the largest single contributor
to the chain gathers' register and spill footprint, and it is now 24. The window
bounds narrow through narrowBound, which saturates rather than wraps: a bound
past the field's range can only make the window empty, and an empty window stays
empty at the clamp, whereas a wrap could turn it back into a live one and read
out of range.

Reviewed By: Yuhta

Differential Revision: D116667979

@netlify

netlify Bot commented Sep 5, 2026

Copy link
Copy Markdown

Deploy Preview for meta-velox canceled.

Name Link
🔨 Latest commit bd2907e
🔍 Latest deploy log https://app.netlify.com/projects/meta-velox/deploys/6a9cf2fa7a886b0008c820fd

@meta-cla meta-cla Bot added the CLA Signed This label is managed by the Facebook bot. Authors need to sign the CLA before a PR can be reviewed. label Sep 5, 2026
@meta-codesync

meta-codesync Bot commented Sep 5, 2026

Copy link
Copy Markdown
Contributor

@oerling has exported this pull request. If you are a Meta employee, you can view the originating Diff in D116667979.

@github-actions

github-actions Bot commented Sep 5, 2026

Copy link
Copy Markdown

Selective Build Plan

Linux release with adapters is running a full build (PR has a standing approval). See the CI workflows README for what this means.


Selective build plan

oerling added a commit to oerling/velox-1 that referenced this pull request Sep 5, 2026
…acebookincubator#18864)

Summary:
Pull Request resolved: facebookincubator#18864

Decomposes list-producing ops into per-tensor nodes, so each column is a value
the compiler can see: consumer counting, aliasing, CSE and cost-based block
shares all work per column rather than per bundle. On top of that, a column
folds its producer's gather into its own, so a run of subset ops
(`batch_flip_and_truncate_sparse`, `grouped_masked_select_jagged_1d`,
`group_length_guard_sparse`) becomes one gather over the original source with
nothing materialized in between.

`Metadata::decompose` plus `decomposeListOps` give one traversal that calls each
op's own rewrite rule, after CSE (cheaper on the bundled form) and before
partitioning, which the per-column nodes exist to inform.

A per-column gather is spelled as a chain of steps -- `kStepRange` for a flip,
`kStepSelect` for a select -- carried in `step_descs` / `step_scalars` with each
step's own tensors. A column absorbing its producer takes that chain whole: its
own step becomes outermost and the producer's steps append after it with their
tensor and scalar bases rebased. Which of the three getters runs then depends on
whether the composed chain contains a range ANYWHERE, not on what its outermost
step is. A select over a flip starts with a select and still needs the row walk,
because a range is defined per row.

Absorbing is refused unless the rows line up, which is a kernel constraint and
not a formality: `prepareRowSteps` prepares every step at one row number taken
from the outer walk, so row r of the outer has to be row r of the inner. Two
columns of one flip head hold the same lengths when they were built from the
same input lengths under the same cap -- and because a run of flips feeds each
column's new lengths to the next flip as that column's input lengths, that test
recurses down the run instead of bottoming out at tensor identity. Identity
alone decides only the last link and refuses everything deeper.

`group_length_guard_sparse`'s final stage is a range step written twice. For
output element idx of row r it computes

    colOffsets[r - 1] + (idx - resultOffsets[r - 1])

which is what `rangeRowDescriptor` computes for a range at mode 0 with no cap
over those same two offset arrays. Saying so -- one
`tw.chain_gather_guard_column` per column in place of the list-form final --
lets a consumer that already has a range absorb the guard as a SUB-LEVEL of its
own rather than as a second step, so the intermediate is never written and the
depth does not grow. The row proof is the head's: element 1 of its output list
is the offsets a chain over those rows walks, so element 0 is the lengths they
were scanned from, which is what `chainRowsOf` now returns. That arm goes ahead
of the flip head's, because N + 2 is a multiple of three for a seven-column
guard and only the flip arm's target check would otherwise keep it out.

The guard's getter is its own registration rather than a reuse of either
neighbour, for two reasons and one trap. Every column of one guard keeps the
same count -- the per-row minimum summed over rows -- which the head already
has as a scalar, where the flip reads a per-column slice of cumulative lengths.
And the head writes row ENDS, as an inclusive scan, so `start_offsets` is 0
where a select's offsets stage passes 1. The trap is `sizeOrdinal`: the select's
registration names the step tensors, which for a guard column are the row
offsets, so the grid would be sized by the ROW count -- three orders of
magnitude short of the output on a real batch. Naming nothing is right here,
because the default kMax over the inputs picks the source, which is at least as
long as the output.

`WaveConfig::foldSharedChains` folds a producer into every consumer that can
absorb one rather than only into a sole reader. Two foldable consumers of one
column can never satisfy the sole-reader rule: whichever rewrite runs first sees
the other as an outside reader and declines, and the second then finds a gather
already reading the buffer, so the column is materialized and NEITHER folds.
Deciding that needs readers counted through `prim.ListPack` / `prim.ListUnpack`
to the ops that really read the data, since a pack is a live user that reads
nothing and a pack immediately unpacked is the identity.

It defaults ON, which only measurement could settle: a reader that folds beside
one that then declines for its own reasons leaves the buffer AND duplicates the
work. On the ROO preproc it takes every column it is offered and leaves none
behind -- see the numbers below. `sharedChain` now forces it OFF rather than
taking the default, so the unfolded arm stays covered.

`RangeRowDesc` carries its within-row fields in 32 bits. Only the two origins
are offsets into a whole column and need 64; `outLen` and the valid window are
bounded by the step's maxLen, and `dir` is +1 or -1. At four steps per chain the
struct is copied per row, so at 48 bytes it was the largest single contributor
to the chain gathers' register and spill footprint, and it is now 24. The window
bounds narrow through `narrowBound`, which saturates rather than wraps: a bound
past the field's range can only make the window empty, and an empty window stays
empty at the clamp, whereas a wrap could turn it back into a live one and read
out of range.

Reviewed By: Yuhta

Differential Revision: D116667979
@meta-codesync meta-codesync Bot changed the title feat(torchwave): Per-column decomposition and chain gather folding feat(torchwave): Per-column decomposition and chain gather folding (#18864) Sep 5, 2026
@github-actions

github-actions Bot commented Sep 5, 2026

Copy link
Copy Markdown

CI Failure Analysis

Auto-generated by the CI Failure Analysis workflow. This comment is updated in place each time CI fails on a new commit, so it always reflects the latest run — re-pushing or re-running CI will refresh the analysis below. Last updated 2026-09-06 03:24:10 UTC from workflow run 34007798269.

🟡 Window Fuzzer with Presto as source of truth — FUZZER Failure View logs

Fuzzer: Window Fuzzer (velox_window_fuzzer_test)
Failed instance: 1 of 4 (seed=899708054)
Instances 2, 3, 4: Passed (seeds: 610554962, 416633362, 260062)

Error: approx_percentile custom result verification failure — the fuzzer found that a logically equivalent streaming window plan produced an approx_percentile result outside the acceptable accuracy range.

approx_percentile(pct: 0.75, accuracy: 0.0133) is more than 0.0133 away
from acceptable range of [0.767409, 0.856198]. Difference: 0.0174095

Expression: verifier->verify(actual.result)
Reason:     Result of a logically equivalent plan failed custom verification
File:       velox/exec/fuzzer/AggregationFuzzerBase.cpp:622
Function:   compare

Stack trace (key frames):

AggregationFuzzerBase::compare()       — AggregationFuzzerBase.cpp:625
AggregationFuzzerBase::testPlan()      — AggregationFuzzerBase.cpp:569
WindowFuzzer::testAlternativePlans()   — WindowFuzzer.cpp:676
WindowFuzzer::verifyWindow()           — WindowFuzzer.cpp:851
WindowFuzzer::go()                     — WindowFuzzer.cpp:557

Failed plan: approx_percentile(c0, c1, c2) with STREAMING window, partitioned by 15 columns, ROWS between UNBOUNDED PRECEDING and CURRENT ROW.


Correlation with PR changes:
This failure is not related to the PR changes. PR #18864 modifies files exclusively under velox/experimental/torchwave/ (AllocGroup, Cat, Compile, CompiledOp, Elementwise, KernelOperation, ParallelExpr, Registry, WaveConfig, WaveGraph, and their tests). The failing fuzzer exercises approx_percentile in the window operator (velox/exec/fuzzer/), which is completely unrelated to the torchwave subsystem.

Known issues:

  • 🔗 #8733 — ApproxPercentile fails with xxx away from acceptable range — This is a known, pre-existing bug tracked since July 2024. The ApproxPercentileResultVerifier occasionally rejects results that are marginally outside the declared accuracy bound.
  • Pre-existing/flaky on main: The identical failure (approx_percentile ... away from acceptable rangecustom verification failed) occurred on the main branch in Fuzzer Jobs run 33996507245 (same job, different seed). This confirms the failure is not introduced by this PR.

Reproduce locally:

./_build/debug/velox/functions/prestosql/fuzzer/velox_window_fuzzer_test \
  --seed 899708054 \
  --duration_sec 300 \
  --batch_size=50 \
  --minloglevel=0 \
  --enable_window_reference_verification \
  --presto_url=http://127.0.0.1:8080

Note: A running Presto server is required for this fuzzer since it uses Presto as the source of truth.

Recommended fix:
No action needed on this PR. The failure is tracked by #8733 and is a known flaky failure in the approx_percentile verification logic.

oerling added a commit to oerling/velox-1 that referenced this pull request Sep 5, 2026
…acebookincubator#18864)

Summary:
Pull Request resolved: facebookincubator#18864

Decomposes list-producing ops into per-tensor nodes, so each column is a value
the compiler can see: consumer counting, aliasing, CSE and cost-based block
shares all work per column rather than per bundle. On top of that, a column
folds its producer's gather into its own, so a run of subset ops
(`batch_flip_and_truncate_sparse`, `grouped_masked_select_jagged_1d`,
`group_length_guard_sparse`) becomes one gather over the original source with
nothing materialized in between.

`Metadata::decompose` plus `decomposeListOps` give one traversal that calls each
op's own rewrite rule, after CSE (cheaper on the bundled form) and before
partitioning, which the per-column nodes exist to inform.

A per-column gather is spelled as a chain of steps -- `kStepRange` for a flip,
`kStepSelect` for a select -- carried in `step_descs` / `step_scalars` with each
step's own tensors. A column absorbing its producer takes that chain whole: its
own step becomes outermost and the producer's steps append after it with their
tensor and scalar bases rebased. Which of the three getters runs then depends on
whether the composed chain contains a range ANYWHERE, not on what its outermost
step is. A select over a flip starts with a select and still needs the row walk,
because a range is defined per row.

Absorbing is refused unless the rows line up, which is a kernel constraint and
not a formality: `prepareRowSteps` prepares every step at one row number taken
from the outer walk, so row r of the outer has to be row r of the inner. Two
columns of one flip head hold the same lengths when they were built from the
same input lengths under the same cap -- and because a run of flips feeds each
column's new lengths to the next flip as that column's input lengths, that test
recurses down the run instead of bottoming out at tensor identity. Identity
alone decides only the last link and refuses everything deeper.

`group_length_guard_sparse`'s final stage is a range step written twice. For
output element idx of row r it computes

    colOffsets[r - 1] + (idx - resultOffsets[r - 1])

which is what `rangeRowDescriptor` computes for a range at mode 0 with no cap
over those same two offset arrays. Saying so -- one
`tw.chain_gather_guard_column` per column in place of the list-form final --
lets a consumer that already has a range absorb the guard as a SUB-LEVEL of its
own rather than as a second step, so the intermediate is never written and the
depth does not grow. The row proof is the head's: element 1 of its output list
is the offsets a chain over those rows walks, so element 0 is the lengths they
were scanned from, which is what `chainRowsOf` now returns. That arm goes ahead
of the flip head's, because N + 2 is a multiple of three for a seven-column
guard and only the flip arm's target check would otherwise keep it out.

The guard's getter is its own registration rather than a reuse of either
neighbour, for two reasons and one trap. Every column of one guard keeps the
same count -- the per-row minimum summed over rows -- which the head already
has as a scalar, where the flip reads a per-column slice of cumulative lengths.
And the head writes row ENDS, as an inclusive scan, so `start_offsets` is 0
where a select's offsets stage passes 1. The trap is `sizeOrdinal`: the select's
registration names the step tensors, which for a guard column are the row
offsets, so the grid would be sized by the ROW count -- three orders of
magnitude short of the output on a real batch. Naming nothing is right here,
because the default kMax over the inputs picks the source, which is at least as
long as the output.

`WaveConfig::foldSharedChains` folds a producer into every consumer that can
absorb one rather than only into a sole reader. Two foldable consumers of one
column can never satisfy the sole-reader rule: whichever rewrite runs first sees
the other as an outside reader and declines, and the second then finds a gather
already reading the buffer, so the column is materialized and NEITHER folds.
Deciding that needs readers counted through `prim.ListPack` / `prim.ListUnpack`
to the ops that really read the data, since a pack is a live user that reads
nothing and a pack immediately unpacked is the identity.

It defaults ON, which only measurement could settle: a reader that folds beside
one that then declines for its own reasons leaves the buffer AND duplicates the
work. On the ROO preproc it takes every column it is offered and leaves none
behind -- see the numbers below. `sharedChain` now forces it OFF rather than
taking the default, so the unfolded arm stays covered.

`RangeRowDesc` carries its within-row fields in 32 bits. Only the two origins
are offsets into a whole column and need 64; `outLen` and the valid window are
bounded by the step's maxLen, and `dir` is +1 or -1. At four steps per chain the
struct is copied per row, so at 48 bytes it was the largest single contributor
to the chain gathers' register and spill footprint, and it is now 24. The window
bounds narrow through `narrowBound`, which saturates rather than wraps: a bound
past the field's range can only make the window empty, and an empty window stays
empty at the clamp, whereas a wrap could turn it back into a live one and read
out of range.

Reviewed By: Yuhta

Differential Revision: D116667979
oerling added a commit to oerling/velox-1 that referenced this pull request Sep 6, 2026
…acebookincubator#18864)

Summary:
Pull Request resolved: facebookincubator#18864

Decomposes list-producing ops into per-tensor nodes, so each column is a value
the compiler can see: consumer counting, aliasing, CSE and cost-based block
shares all work per column rather than per bundle. On top of that, a column
folds its producer's gather into its own, so a run of subset ops
(`batch_flip_and_truncate_sparse`, `grouped_masked_select_jagged_1d`,
`group_length_guard_sparse`) becomes one gather over the original source with
nothing materialized in between.

`Metadata::decompose` plus `decomposeListOps` give one traversal that calls each
op's own rewrite rule, after CSE (cheaper on the bundled form) and before
partitioning, which the per-column nodes exist to inform.

A per-column gather is spelled as a chain of steps -- `kStepRange` for a flip,
`kStepSelect` for a select -- carried in `step_descs` / `step_scalars` with each
step's own tensors. A column absorbing its producer takes that chain whole: its
own step becomes outermost and the producer's steps append after it with their
tensor and scalar bases rebased. Which of the three getters runs then depends on
whether the composed chain contains a range ANYWHERE, not on what its outermost
step is. A select over a flip starts with a select and still needs the row walk,
because a range is defined per row.

Absorbing is refused unless the rows line up, which is a kernel constraint and
not a formality: `prepareRowSteps` prepares every step at one row number taken
from the outer walk, so row r of the outer has to be row r of the inner. Two
columns of one flip head hold the same lengths when they were built from the
same input lengths under the same cap -- and because a run of flips feeds each
column's new lengths to the next flip as that column's input lengths, that test
recurses down the run instead of bottoming out at tensor identity. Identity
alone decides only the last link and refuses everything deeper.

`group_length_guard_sparse`'s final stage is a range step written twice. For
output element idx of row r it computes

    colOffsets[r - 1] + (idx - resultOffsets[r - 1])

which is what `rangeRowDescriptor` computes for a range at mode 0 with no cap
over those same two offset arrays. Saying so -- one
`tw.chain_gather_guard_column` per column in place of the list-form final --
lets a consumer that already has a range absorb the guard as a SUB-LEVEL of its
own rather than as a second step, so the intermediate is never written and the
depth does not grow. The row proof is the head's: element 1 of its output list
is the offsets a chain over those rows walks, so element 0 is the lengths they
were scanned from, which is what `chainRowsOf` now returns. That arm goes ahead
of the flip head's, because N + 2 is a multiple of three for a seven-column
guard and only the flip arm's target check would otherwise keep it out.

The guard's getter is its own registration rather than a reuse of either
neighbour, for two reasons and one trap. Every column of one guard keeps the
same count -- the per-row minimum summed over rows -- which the head already
has as a scalar, where the flip reads a per-column slice of cumulative lengths.
And the head writes row ENDS, as an inclusive scan, so `start_offsets` is 0
where a select's offsets stage passes 1. The trap is `sizeOrdinal`: the select's
registration names the step tensors, which for a guard column are the row
offsets, so the grid would be sized by the ROW count -- three orders of
magnitude short of the output on a real batch. Naming nothing is right here,
because the default kMax over the inputs picks the source, which is at least as
long as the output.

`WaveConfig::foldSharedChains` folds a producer into every consumer that can
absorb one rather than only into a sole reader. Two foldable consumers of one
column can never satisfy the sole-reader rule: whichever rewrite runs first sees
the other as an outside reader and declines, and the second then finds a gather
already reading the buffer, so the column is materialized and NEITHER folds.
Deciding that needs readers counted through `prim.ListPack` / `prim.ListUnpack`
to the ops that really read the data, since a pack is a live user that reads
nothing and a pack immediately unpacked is the identity.

It defaults ON, which only measurement could settle: a reader that folds beside
one that then declines for its own reasons leaves the buffer AND duplicates the
work. On the ROO preproc it takes every column it is offered and leaves none
behind -- see the numbers below. `sharedChain` now forces it OFF rather than
taking the default, so the unfolded arm stays covered.

`RangeRowDesc` carries its within-row fields in 32 bits. Only the two origins
are offsets into a whole column and need 64; `outLen` and the valid window are
bounded by the step's maxLen, and `dir` is +1 or -1. At four steps per chain the
struct is copied per row, so at 48 bytes it was the largest single contributor
to the chain gathers' register and spill footprint, and it is now 24. The window
bounds narrow through `narrowBound`, which saturates rather than wraps: a bound
past the field's range can only make the window empty, and an empty window stays
empty at the clamp, whereas a wrap could turn it back into a live one and read
out of range.

Reviewed By: Yuhta

Differential Revision: D116667979
oerling added a commit to oerling/velox-1 that referenced this pull request Sep 6, 2026
…acebookincubator#18864)

Summary:
Pull Request resolved: facebookincubator#18864

Decomposes list-producing ops into per-tensor nodes, so each column is a value
the compiler can see: consumer counting, aliasing, CSE and cost-based block
shares all work per column rather than per bundle. On top of that, a column
folds its producer's gather into its own, so a run of subset ops
(`batch_flip_and_truncate_sparse`, `grouped_masked_select_jagged_1d`,
`group_length_guard_sparse`) becomes one gather over the original source with
nothing materialized in between.

`Metadata::decompose` plus `decomposeListOps` give one traversal that calls each
op's own rewrite rule, after CSE (cheaper on the bundled form) and before
partitioning, which the per-column nodes exist to inform.

A per-column gather is spelled as a chain of steps -- `kStepRange` for a flip,
`kStepSelect` for a select -- carried in `step_descs` / `step_scalars` with each
step's own tensors. A column absorbing its producer takes that chain whole: its
own step becomes outermost and the producer's steps append after it with their
tensor and scalar bases rebased. Which of the three getters runs then depends on
whether the composed chain contains a range ANYWHERE, not on what its outermost
step is. A select over a flip starts with a select and still needs the row walk,
because a range is defined per row.

Absorbing is refused unless the rows line up, which is a kernel constraint and
not a formality: `prepareRowSteps` prepares every step at one row number taken
from the outer walk, so row r of the outer has to be row r of the inner. Two
columns of one flip head hold the same lengths when they were built from the
same input lengths under the same cap -- and because a run of flips feeds each
column's new lengths to the next flip as that column's input lengths, that test
recurses down the run instead of bottoming out at tensor identity. Identity
alone decides only the last link and refuses everything deeper.

`group_length_guard_sparse`'s final stage is a range step written twice. For
output element idx of row r it computes

    colOffsets[r - 1] + (idx - resultOffsets[r - 1])

which is what `rangeRowDescriptor` computes for a range at mode 0 with no cap
over those same two offset arrays. Saying so -- one
`tw.chain_gather_guard_column` per column in place of the list-form final --
lets a consumer that already has a range absorb the guard as a SUB-LEVEL of its
own rather than as a second step, so the intermediate is never written and the
depth does not grow. The row proof is the head's: element 1 of its output list
is the offsets a chain over those rows walks, so element 0 is the lengths they
were scanned from, which is what `chainRowsOf` now returns. That arm goes ahead
of the flip head's, because N + 2 is a multiple of three for a seven-column
guard and only the flip arm's target check would otherwise keep it out.

The guard's getter is its own registration rather than a reuse of either
neighbour, for two reasons and one trap. Every column of one guard keeps the
same count -- the per-row minimum summed over rows -- which the head already
has as a scalar, where the flip reads a per-column slice of cumulative lengths.
And the head writes row ENDS, as an inclusive scan, so `start_offsets` is 0
where a select's offsets stage passes 1. The trap is `sizeOrdinal`: the select's
registration names the step tensors, which for a guard column are the row
offsets, so the grid would be sized by the ROW count -- three orders of
magnitude short of the output on a real batch. Naming nothing is right here,
because the default kMax over the inputs picks the source, which is at least as
long as the output.

`WaveConfig::foldSharedChains` folds a producer into every consumer that can
absorb one rather than only into a sole reader. Two foldable consumers of one
column can never satisfy the sole-reader rule: whichever rewrite runs first sees
the other as an outside reader and declines, and the second then finds a gather
already reading the buffer, so the column is materialized and NEITHER folds.
Deciding that needs readers counted through `prim.ListPack` / `prim.ListUnpack`
to the ops that really read the data, since a pack is a live user that reads
nothing and a pack immediately unpacked is the identity.

It defaults ON, which only measurement could settle: a reader that folds beside
one that then declines for its own reasons leaves the buffer AND duplicates the
work. On the ROO preproc it takes every column it is offered and leaves none
behind -- see the numbers below. `sharedChain` now forces it OFF rather than
taking the default, so the unfolded arm stays covered.

`RangeRowDesc` carries its within-row fields in 32 bits. Only the two origins
are offsets into a whole column and need 64; `outLen` and the valid window are
bounded by the step's maxLen, and `dir` is +1 or -1. At four steps per chain the
struct is copied per row, so at 48 bytes it was the largest single contributor
to the chain gathers' register and spill footprint, and it is now 24. The window
bounds narrow through `narrowBound`, which saturates rather than wraps: a bound
past the field's range can only make the window empty, and an empty window stays
empty at the clamp, whereas a wrap could turn it back into a live one and read
out of range.

Reviewed By: Yuhta

Differential Revision: D116667979
oerling added a commit to oerling/velox-1 that referenced this pull request Sep 6, 2026
…acebookincubator#18864)

Summary:
Pull Request resolved: facebookincubator#18864

Decomposes list-producing ops into per-tensor nodes, so each column is a value
the compiler can see: consumer counting, aliasing, CSE and cost-based block
shares all work per column rather than per bundle. On top of that, a column
folds its producer's gather into its own, so a run of subset ops
(`batch_flip_and_truncate_sparse`, `grouped_masked_select_jagged_1d`,
`group_length_guard_sparse`) becomes one gather over the original source with
nothing materialized in between.

`Metadata::decompose` plus `decomposeListOps` give one traversal that calls each
op's own rewrite rule, after CSE (cheaper on the bundled form) and before
partitioning, which the per-column nodes exist to inform.

A per-column gather is spelled as a chain of steps -- `kStepRange` for a flip,
`kStepSelect` for a select -- carried in `step_descs` / `step_scalars` with each
step's own tensors. A column absorbing its producer takes that chain whole: its
own step becomes outermost and the producer's steps append after it with their
tensor and scalar bases rebased. Which of the three getters runs then depends on
whether the composed chain contains a range ANYWHERE, not on what its outermost
step is. A select over a flip starts with a select and still needs the row walk,
because a range is defined per row.

Absorbing is refused unless the rows line up, which is a kernel constraint and
not a formality: `prepareRowSteps` prepares every step at one row number taken
from the outer walk, so row r of the outer has to be row r of the inner. Two
columns of one flip head hold the same lengths when they were built from the
same input lengths under the same cap -- and because a run of flips feeds each
column's new lengths to the next flip as that column's input lengths, that test
recurses down the run instead of bottoming out at tensor identity. Identity
alone decides only the last link and refuses everything deeper.

`group_length_guard_sparse`'s final stage is a range step written twice. For
output element idx of row r it computes

    colOffsets[r - 1] + (idx - resultOffsets[r - 1])

which is what `rangeRowDescriptor` computes for a range at mode 0 with no cap
over those same two offset arrays. Saying so -- one
`tw.chain_gather_guard_column` per column in place of the list-form final --
lets a consumer that already has a range absorb the guard as a SUB-LEVEL of its
own rather than as a second step, so the intermediate is never written and the
depth does not grow. The row proof is the head's: element 1 of its output list
is the offsets a chain over those rows walks, so element 0 is the lengths they
were scanned from, which is what `chainRowsOf` now returns. That arm goes ahead
of the flip head's, because N + 2 is a multiple of three for a seven-column
guard and only the flip arm's target check would otherwise keep it out.

The guard's getter is its own registration rather than a reuse of either
neighbour, for two reasons and one trap. Every column of one guard keeps the
same count -- the per-row minimum summed over rows -- which the head already
has as a scalar, where the flip reads a per-column slice of cumulative lengths.
And the head writes row ENDS, as an inclusive scan, so `start_offsets` is 0
where a select's offsets stage passes 1. The trap is `sizeOrdinal`: the select's
registration names the step tensors, which for a guard column are the row
offsets, so the grid would be sized by the ROW count -- three orders of
magnitude short of the output on a real batch. Naming nothing is right here,
because the default kMax over the inputs picks the source, which is at least as
long as the output.

`WaveConfig::foldSharedChains` folds a producer into every consumer that can
absorb one rather than only into a sole reader. Two foldable consumers of one
column can never satisfy the sole-reader rule: whichever rewrite runs first sees
the other as an outside reader and declines, and the second then finds a gather
already reading the buffer, so the column is materialized and NEITHER folds.
Deciding that needs readers counted through `prim.ListPack` / `prim.ListUnpack`
to the ops that really read the data, since a pack is a live user that reads
nothing and a pack immediately unpacked is the identity.

It defaults ON, which only measurement could settle: a reader that folds beside
one that then declines for its own reasons leaves the buffer AND duplicates the
work. On the ROO preproc it takes every column it is offered and leaves none
behind -- see the numbers below. `sharedChain` now forces it OFF rather than
taking the default, so the unfolded arm stays covered.

`RangeRowDesc` carries its within-row fields in 32 bits. Only the two origins
are offsets into a whole column and need 64; `outLen` and the valid window are
bounded by the step's maxLen, and `dir` is +1 or -1. At four steps per chain the
struct is copied per row, so at 48 bytes it was the largest single contributor
to the chain gathers' register and spill footprint, and it is now 24. The window
bounds narrow through `narrowBound`, which saturates rather than wraps: a bound
past the field's range can only make the window empty, and an empty window stays
empty at the clamp, whereas a wrap could turn it back into a live one and read
out of range.

Reviewed By: Yuhta

Differential Revision: D116667979
…rnel wrote (facebookincubator#18862)

Summary:
Pull Request resolved: facebookincubator#18862

Two independent changes, folded into one diff. The first is the barrier fix
this diff was opened for; the second is the sym_size border experiment that
was D118061301.

## Barrier before an elementwise op reads a leaf this kernel wrote

An elementwise op's operands are all marked `isRegister` by
`registerElementwise`, and `callNeedsBarrier` skips register operands on
the grounds that they flow inline rather than through memory. That holds
for a value fused into the expression tree; it does not hold for a leaf,
which is a load like any other. The test was that the two are the same
thing, and they are not, so no pure-elementwise op ever asked for a
barrier.

What that costs is visible whenever the leaf does not map index to index.
The ROO preproc graph fuses

    (%3274)  = where.self(...)
    (%11769) = transpose.int(%3274, dim0=0, dim1=1)
    (%11800) = mul(%11771, %11769)

into one kernel with nothing but `__syncthreads()` between the write and
the read. Through the transpose the element thread i of the `mul` wants
was written by a different block, so `%11800` came out part stale and
part uninitialized -- and only in the multi-block modes, which is what
made it look like a miscompile rather than a race.

The operand test `callNeedsBarrier` already had is now
`valueNeedsBarrier`, and `generateElementwise` runs it over each
subgraph's own memory leaves.

## Stop a returned sym_size bordering its operand (was D118061301)

A metadata getter (`sym_size` / `sym_numel`) whose only role is to be a graph
output still counts as a use of its operand when the partitioner builds its
levels. `makeLevelsInner` increments the operand's producer refCount, and
`makeCseBorder` turns anything above one into a border: the producer moves to
its own earlier layer and every other consumer of it follows a layer later than
it needed to. The getter's own work is one scalar field read on the host; the
whole cost is the layer split it forces.

`WaveConfig::deferSizeOutputs`, OFF BY DEFAULT, drops those getters from `top`
before `makeExprLevels` so they stop being a reference, and puts them back
before the last layer is built so they still run and still fill their output
slot.

The predicate counts users through the set reachable from the output node, not
through `users()`. torch.export leaves dead `_operator.ge` / `_operator.le`
shape guards behind once the asserts are stripped -- 428 of them on the ROO
preproc graph -- and they appear in a value's `users()` while contributing to no
level. Counting them refuses every candidate; with them filtered out, 255 of the
graph's 297 top-level exprs qualify and 248 of those have an operand whose only
other reachable user is a single consumer.

It is off by default because on that graph it is a net loss, and the measurement
is the point of the commit rather than the feature. See the test plan.

Reviewed By: Yuhta

Differential Revision: D118061298
oerling added a commit to oerling/velox-1 that referenced this pull request Sep 6, 2026
…acebookincubator#18864)

Summary:
Pull Request resolved: facebookincubator#18864

Decomposes list-producing ops into per-tensor nodes, so each column is a value
the compiler can see: consumer counting, aliasing, CSE and cost-based block
shares all work per column rather than per bundle. On top of that, a column
folds its producer's gather into its own, so a run of subset ops
(`batch_flip_and_truncate_sparse`, `grouped_masked_select_jagged_1d`,
`group_length_guard_sparse`) becomes one gather over the original source with
nothing materialized in between.

`Metadata::decompose` plus `decomposeListOps` give one traversal that calls each
op's own rewrite rule, after CSE (cheaper on the bundled form) and before
partitioning, which the per-column nodes exist to inform.

A per-column gather is spelled as a chain of steps -- `kStepRange` for a flip,
`kStepSelect` for a select -- carried in `step_descs` / `step_scalars` with each
step's own tensors. A column absorbing its producer takes that chain whole: its
own step becomes outermost and the producer's steps append after it with their
tensor and scalar bases rebased. Which of the three getters runs then depends on
whether the composed chain contains a range ANYWHERE, not on what its outermost
step is. A select over a flip starts with a select and still needs the row walk,
because a range is defined per row.

Absorbing is refused unless the rows line up, which is a kernel constraint and
not a formality: `prepareRowSteps` prepares every step at one row number taken
from the outer walk, so row r of the outer has to be row r of the inner. Two
columns of one flip head hold the same lengths when they were built from the
same input lengths under the same cap -- and because a run of flips feeds each
column's new lengths to the next flip as that column's input lengths, that test
recurses down the run instead of bottoming out at tensor identity. Identity
alone decides only the last link and refuses everything deeper.

`group_length_guard_sparse`'s final stage is a range step written twice. For
output element idx of row r it computes

    colOffsets[r - 1] + (idx - resultOffsets[r - 1])

which is what `rangeRowDescriptor` computes for a range at mode 0 with no cap
over those same two offset arrays. Saying so -- one
`tw.chain_gather_guard_column` per column in place of the list-form final --
lets a consumer that already has a range absorb the guard as a SUB-LEVEL of its
own rather than as a second step, so the intermediate is never written and the
depth does not grow. The row proof is the head's: element 1 of its output list
is the offsets a chain over those rows walks, so element 0 is the lengths they
were scanned from, which is what `chainRowsOf` now returns. That arm goes ahead
of the flip head's, because N + 2 is a multiple of three for a seven-column
guard and only the flip arm's target check would otherwise keep it out.

The guard's getter is its own registration rather than a reuse of either
neighbour, for two reasons and one trap. Every column of one guard keeps the
same count -- the per-row minimum summed over rows -- which the head already
has as a scalar, where the flip reads a per-column slice of cumulative lengths.
And the head writes row ENDS, as an inclusive scan, so `start_offsets` is 0
where a select's offsets stage passes 1. The trap is `sizeOrdinal`: the select's
registration names the step tensors, which for a guard column are the row
offsets, so the grid would be sized by the ROW count -- three orders of
magnitude short of the output on a real batch. Naming nothing is right here,
because the default kMax over the inputs picks the source, which is at least as
long as the output.

`WaveConfig::foldSharedChains` folds a producer into every consumer that can
absorb one rather than only into a sole reader. Two foldable consumers of one
column can never satisfy the sole-reader rule: whichever rewrite runs first sees
the other as an outside reader and declines, and the second then finds a gather
already reading the buffer, so the column is materialized and NEITHER folds.
Deciding that needs readers counted through `prim.ListPack` / `prim.ListUnpack`
to the ops that really read the data, since a pack is a live user that reads
nothing and a pack immediately unpacked is the identity.

It defaults ON, which only measurement could settle: a reader that folds beside
one that then declines for its own reasons leaves the buffer AND duplicates the
work. On the ROO preproc it takes every column it is offered and leaves none
behind -- see the numbers below. `sharedChain` now forces it OFF rather than
taking the default, so the unfolded arm stays covered.

`RangeRowDesc` carries its within-row fields in 32 bits. Only the two origins
are offsets into a whole column and need 64; `outLen` and the valid window are
bounded by the step's maxLen, and `dir` is +1 or -1. At four steps per chain the
struct is copied per row, so at 48 bytes it was the largest single contributor
to the chain gathers' register and spill footprint, and it is now 24. The window
bounds narrow through `narrowBound`, which saturates rather than wraps: a bound
past the field's range can only make the window empty, and an empty window stays
empty at the clamp, whereas a wrap could turn it back into a live one and read
out of range.

Reviewed By: Yuhta

Differential Revision: D116667979
…acebookincubator#18864)

Summary:
Pull Request resolved: facebookincubator#18864

Decomposes list-producing ops into per-tensor nodes, so each column is a value
the compiler can see: consumer counting, aliasing, CSE and cost-based block
shares all work per column rather than per bundle. On top of that, a column
folds its producer's gather into its own, so a run of subset ops
(`batch_flip_and_truncate_sparse`, `grouped_masked_select_jagged_1d`,
`group_length_guard_sparse`) becomes one gather over the original source with
nothing materialized in between.

`Metadata::decompose` plus `decomposeListOps` give one traversal that calls each
op's own rewrite rule, after CSE (cheaper on the bundled form) and before
partitioning, which the per-column nodes exist to inform.

A per-column gather is spelled as a chain of steps -- `kStepRange` for a flip,
`kStepSelect` for a select -- carried in `step_descs` / `step_scalars` with each
step's own tensors. A column absorbing its producer takes that chain whole: its
own step becomes outermost and the producer's steps append after it with their
tensor and scalar bases rebased. Which of the three getters runs then depends on
whether the composed chain contains a range ANYWHERE, not on what its outermost
step is. A select over a flip starts with a select and still needs the row walk,
because a range is defined per row.

Absorbing is refused unless the rows line up, which is a kernel constraint and
not a formality: `prepareRowSteps` prepares every step at one row number taken
from the outer walk, so row r of the outer has to be row r of the inner. Two
columns of one flip head hold the same lengths when they were built from the
same input lengths under the same cap -- and because a run of flips feeds each
column's new lengths to the next flip as that column's input lengths, that test
recurses down the run instead of bottoming out at tensor identity. Identity
alone decides only the last link and refuses everything deeper.

`group_length_guard_sparse`'s final stage is a range step written twice. For
output element idx of row r it computes

    colOffsets[r - 1] + (idx - resultOffsets[r - 1])

which is what `rangeRowDescriptor` computes for a range at mode 0 with no cap
over those same two offset arrays. Saying so -- one
`tw.chain_gather_guard_column` per column in place of the list-form final --
lets a consumer that already has a range absorb the guard as a SUB-LEVEL of its
own rather than as a second step, so the intermediate is never written and the
depth does not grow. The row proof is the head's: element 1 of its output list
is the offsets a chain over those rows walks, so element 0 is the lengths they
were scanned from, which is what `chainRowsOf` now returns. That arm goes ahead
of the flip head's, because N + 2 is a multiple of three for a seven-column
guard and only the flip arm's target check would otherwise keep it out.

The guard's getter is its own registration rather than a reuse of either
neighbour, for two reasons and one trap. Every column of one guard keeps the
same count -- the per-row minimum summed over rows -- which the head already
has as a scalar, where the flip reads a per-column slice of cumulative lengths.
And the head writes row ENDS, as an inclusive scan, so `start_offsets` is 0
where a select's offsets stage passes 1. The trap is `sizeOrdinal`: the select's
registration names the step tensors, which for a guard column are the row
offsets, so the grid would be sized by the ROW count -- three orders of
magnitude short of the output on a real batch. Naming nothing is right here,
because the default kMax over the inputs picks the source, which is at least as
long as the output.

`WaveConfig::foldSharedChains` folds a producer into every consumer that can
absorb one rather than only into a sole reader. Two foldable consumers of one
column can never satisfy the sole-reader rule: whichever rewrite runs first sees
the other as an outside reader and declines, and the second then finds a gather
already reading the buffer, so the column is materialized and NEITHER folds.
Deciding that needs readers counted through `prim.ListPack` / `prim.ListUnpack`
to the ops that really read the data, since a pack is a live user that reads
nothing and a pack immediately unpacked is the identity.

It defaults ON, which only measurement could settle: a reader that folds beside
one that then declines for its own reasons leaves the buffer AND duplicates the
work. On the ROO preproc it takes every column it is offered and leaves none
behind -- see the numbers below. `sharedChain` now forces it OFF rather than
taking the default, so the unfolded arm stays covered.

`RangeRowDesc` carries its within-row fields in 32 bits. Only the two origins
are offsets into a whole column and need 64; `outLen` and the valid window are
bounded by the step's maxLen, and `dir` is +1 or -1. At four steps per chain the
struct is copied per row, so at 48 bytes it was the largest single contributor
to the chain gathers' register and spill footprint, and it is now 24. The window
bounds narrow through `narrowBound`, which saturates rather than wraps: a bound
past the field's range can only make the window empty, and an empty window stays
empty at the clamp, whereas a wrap could turn it back into a live one and read
out of range.

Reviewed By: Yuhta

Differential Revision: D116667979
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

CLA Signed This label is managed by the Facebook bot. Authors need to sign the CLA before a PR can be reviewed. meta-exported

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants