feat(torchwave): Per-column decomposition and chain gather folding (#18864) - #18864
feat(torchwave): Per-column decomposition and chain gather folding (#18864)#18864oerling wants to merge 2 commits into
Conversation
✅ Deploy Preview for meta-velox canceled.
|
|
@oerling has exported this pull request. If you are a Meta employee, you can view the originating Diff in D116667979. |
Selective Build Plan
Selective build plan |
…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
400a986 to
079479d
Compare
CI Failure Analysis
🟡 Window Fuzzer with Presto as source of truth — FUZZER Failure View logsFuzzer: Window Fuzzer ( Error: Stack trace (key frames): Failed plan: Correlation with PR changes: Known issues:
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:8080Note: A running Presto server is required for this fuzzer since it uses Presto as the source of truth. Recommended fix: |
…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
079479d to
86c05d5
Compare
…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
86c05d5 to
f4d15eb
Compare
…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
f4d15eb to
9cc60cd
Compare
…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
9cc60cd to
fb4dca8
Compare
…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
…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
fb4dca8 to
825c7ac
Compare
…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
825c7ac to
bd2907e
Compare
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 withnothing materialized in between.
Metadata::decomposeplusdecomposeListOpsgive one traversal that calls eachop'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 --
kStepRangefor a flip,kStepSelectfor a select -- carried instep_descs/step_scalarswith eachstep'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:
prepareRowStepsprepares every step at one row number takenfrom 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. Foroutput element idx of row r it computes
which is what
rangeRowDescriptorcomputes for a range at mode 0 with no capover those same two offset arrays. Saying so -- one
tw.chain_gather_guard_columnper 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
chainRowsOfnow returns. That arm goes aheadof 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_offsetsis 0where a select's offsets stage passes 1. The trap is
sizeOrdinal: the select'sregistration 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::foldSharedChainsfolds a producer into every consumer that canabsorb 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.ListUnpackto 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.
sharedChainnow forces it OFF rather thantaking the default, so the unfolded arm stays covered.
RangeRowDesccarries its within-row fields in 32 bits. Only the two originsare offsets into a whole column and need 64;
outLenand the valid window arebounded by the step's maxLen, and
diris +1 or -1. At four steps per chain thestruct 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 boundpast 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