Concurrent query execution: full integration (concurrency stack + late-mat + hardening tiers 1-5) - #1583
Draft
felipeblazing wants to merge 141 commits into
Draft
Concurrent query execution: full integration (concurrency stack + late-mat + hardening tiers 1-5)#1583felipeblazing wants to merge 141 commits into
felipeblazing wants to merge 141 commits into
Conversation
…ressing Squashed from the 16-commit fused scan-filter campaign, so the branch carries the mechanism as one change rather than as a series stacked on the SF1000 repro branch, whose work is already on dev. The campaign's captured run artifacts are left out, since later work on this branch deletes them again. Range, pair, dictionary and dynamic-filter conjuncts produce a selection mask during decode, and the output columns are decoded survivor-compacted instead of full width and gathered. Two waves: wave 1 ballots every row-selecting source into packed mask words and counts survivors, wave 2 decodes each column through the route its plan shape supports (bitpack/delta mask walk, index-list walk, dictionary key gather, str_split offsets + char gather) or falls back to a full decode plus a survivor gather. The whole path is behind SIRIUS_EXP_FUSED_SCAN_FILTER and is inert when the gate is off. Measured on GB300, TPC-H SF1000: the fused scan-filter work took the suite from 7.866 s to 6.918 s — the dictionary and delta routes, the index-list decode, the str_split masked gather (q12 -35.8%), and dynamic-filter masks with dual delivery. A decode-orchestration bundle measured +0.46% and was dropped rather than kept. The internal shorthand this code was written with (W1-W4 work items, K-numbered kernels, RULE 1/2, iteration numbering) is replaced by names that say what the code does in a later commit on this branch. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…gen dedup Design plan for the boundary and codegen work around PR sirius-db#1391 (fused scan-filter) and the closed PR sirius-db#1380 (dict predicate pushdown), whose content is folded into sirius-db#1391. Part A replaces the ~40-name public surface with an intent-level API (scan_read / scan_read_result / compressed_scan). Part B turns the BOOL8 type substitution into declared partial evaluation. Part D records codegen findings F1-F6 and the enumerator x consumer factoring. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…ce kernel args F1: emit_bitpack_mask_out/_mask_consume were byte-identical to their _generic counterparts (value_source already routes a Bitpack leaf to bitpack_value_source). Deleted; verified by rendering 4 dtypes x 6 variants x 2 shapes before/after -- zero diff once comment lines are stripped. The deleted RenderError throws were unreachable. F5: the renderer emitted trailing parameter DECLARATIONS from one switch while the launcher pushed ARGUMENTS from a second switch in another file, with cuLaunchKernel's untyped void** between them -- a mismatch was silent argument misalignment, not a compile error. DecodeKernelSpec now carries the list as TrailingParam tags emitted from the same table as the declaration text; both launch paths bind by tag. F4: the pair launcher open-coded the whole bind/launch sequence and OMITTED the per-chunk metadata bounds guard, risking an out-of-bounds read that faults the CUDA context. Both paths now share launch_rendered_spec, so the pair path inherits the guard. BEHAVIOUR CHANGE: a pair launch with under-sized per-chunk channels is now refused rather than proceeding. F6: five bespoke precondition blocks replaced by a VariantContract table and one checker; diagnostics now name the specific missing field. New test_render_signature_contract asserts declared params == buffers + 2 + trailing for every shape/dtype/variant and the pair path, parsing the signature rather than diffing a golden string. Mutation-tested: injecting one undeclared parameter fails it across every variant. Rendered source is byte-identical after F4/F5/F6; F1 changes only comment text, which does change the JIT cache key -- expect a one-time NVRTC recompile per (shape, dtype, variant). Tests: test_masked_decode_variants 28/28 PASS, test_render_signature_contract PASS, sirius_unittest [compression] 36 cases / 1222 assertions PASS. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…riants The K3 store, K5 dict gather and K6 offsets-meta emitters each open-coded the same survivor loop -- selection stage, strided row walk, mask-bit test, rank computation -- differing only in the one to four lines run per survivor. Extract emit_mask_survivor_loop(sink); each variant now supplies just its sink. Rendered source is byte-identical (dump-and-diff across 4 dtypes x 6 variants x 2 shapes), so no JIT cache churn. This is the emitter half of the enumerator x consumer factoring in DECODE_PUSHDOWN_PLAN.md section 5: the loop is the mask_bits enumerator, the sink is the consumer. Remaining: F3 (the delta emitter still duplicates ~100 lines of emit_delta_producer) and the index_list enumerator. Tests: test_masked_decode_variants 28/28, test_render_signature_contract PASS, sirius_unittest [compression] 36 cases / 1222 assertions PASS. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
sirius-db#1371 (84c49be) merged to dev and carried sirius-db#1380's dict predicate pushdown with it: decode_predicate, decode_equality_pushdown and extract_string_equality_pushdown are now in mainline, while sirius-db#1391's fused scan-filter machinery is not. The plan therefore covers refactoring APIs that are already shipped, not only unmerged ones -- notably P5, the BOOL8 type sniff, which becomes a fix to dev rather than a cleanup of a PR. Also strips progress markers, a status work-log and a struck-through correction from the document so it reads as a design doc rather than a running log. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The constructor resolved the dynamic-filter channel via &_ingestible->table_info() without checking the pointer, so constructing the operator with a null ingestible segfaulted. test_gpu_pipeline_task_history.cpp does exactly that -- it exercises the scheduling/reservation surface with no source -- so sirius_unittest [scan] crashed at test 62 of 265. Predates the dev merge: both the unguarded deref and the nullptr construction were present at 0e61668, and dev's constructor has no such deref, so the two only met on this branch. sirius_unittest [scan]: 265 test cases / 274587 assertions PASS (was SIGSEGV). Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…ucer emit_delta_mask_consume duplicated ~100 lines of emit_delta_producer's root-level striped path -- the same items[IPT] load, cub::BlockScan / BlockExchange union, three transposes and unsigned-counterpart reconstruction -- differing only in the final store. Its header comment said as much: kept standalone "so the plain rendered source stays byte-identical". Thread a DeltaStore policy through emit_delta_producer instead. plain writes every row to dst[row]; mask_compact writes survivors to dst[rank]. The masked emitter is now validation plus one call, and future delta work is done once rather than twice. Byte-identical rendered source, verified by dump-and-diff across 4 dtypes x 6 variants x 2 shapes (bitpack leaf, delta->bitpack) -- so no JIT cache churn. renderer.cpp -42 lines net. This completes section 5's emitter dedup (F1-F3): the mask_bits enumerator and its consumers now share one loop, and delta is no longer a special case. Tests: test_masked_decode_variants 27/27, test_render_signature_contract PASS, sirius_unittest [compression] 1222 assertions, [scan] 274587 assertions. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…mer shape
The renderer and the launcher described the same taxonomy twice: a flat
DecodeVariant enum plus a separate pair entry point on one side, seven flat
launcher functions on the other. Neither represented the underlying product,
so the two enumerations had to be kept in step by hand.
Every kernel is a point in a two-axis product:
Enumerator: all_rows | mask_bits | index_list (how rows are walked)
Consumer: write_column | ballot_range | ballot_pair | dict_gather |
offsets_meta (what happens per row)
plain = all_rows x write_column K3 = mask_bits x write_column
K1 = all_rows x ballot_range K4 = index_list x write_column
K1m2 = all_rows x ballot_pair K5 = mask_bits x dict_gather
K6.1 = mask_bits x offsets_meta
Everything now derives from the pair rather than from a variant tag:
- trailing parameters = enumerator_params ++ consumer_params, which
reproduces the previous per-variant order exactly;
- the `out` slot's type follows from the consumer alone (the ballot
consumers repurpose it for mask words);
- the launcher's precondition contract is two small switches over the axes
instead of one per-variant table;
- Walker::build dispatches on the axes.
shape_is_supported() rejects unsupported points at render, and names the
meaningful-but-unbuilt ones (index_list x dict_gather / offsets_meta, which
would give K4-speed dictionary and string decode below the K4 crossover) as
the combinations that light up when their emitters land.
Entry-symbol suffixes are unchanged and centralised in shape_symbol_suffix:
they key the JIT cache, so moving one forces a recompile.
Rendered source byte-identical across 4 dtypes x 6 shapes x 2 tree shapes.
Tests: test_masked_decode_variants 27/27, test_render_signature_contract PASS,
sirius_unittest [compression] 1222 assertions, [scan] 274587 assertions.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…ape axes Section 5 becomes a description of the shipped state (the enumerator x consumer product, the 3x5 table, what F1-F6 replaced) rather than a work plan, and the sequencing section marks Part D done while promoting Phase 1: since sirius-db#1371 put sirius-db#1380's pushdown on dev, the BOOL8 type-sniff hazard is in shipped mainline code rather than an unmerged PR. Adds "Can the axes grow?" to section 7. Growth is additive -- a new consumer works under every enumerator it composes with -- with run_list and range_slice as the plausible enumerators, and ballot_membership, aggregate, count, hash and null_ballot as consumers. ballot_membership is called out as the highest-value one: membership is the most expensive conjunct today precisely because it cannot work on packed bits, so wave 1 decodes the key column full width and the source cap is 1; a consumer sees the value in-register. Also records the constraints that stop the product being fully orthogonal: ballot consumers need a non-compacting enumerator (the mask layout depends on ballot lanes matching consecutive rows), arity currently hides inside the consumer and should become explicit rather than spawning ballot_triple, and some consumers constrain plan shape rather than the axis. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Replaces the ad-hoc list with a derivation from what each axis can express. Consumer, by what it emits per row: store the value (transform_store -- store f(v) not v, i.e. projection pushdown into decode; q1/q6 materialise two full columns to compute one), store a derived payload, reduce to a bit (ballot_membership, ballot_expression, ballot_null), reduce across rows (count, aggregate, zone_map), derive a value for elsewhere (hash). ballot_membership ranks highest -- membership is the worst conjunct today precisely because it cannot work on packed bits. Enumerator, by how the row set is described: adds run_list (beats index_list on clustered survivors), range_slice (LIMIT/OFFSET, mid-chunk splits), stride, and gather_list (arbitrary, non-ascending ids -- fusing a join's or sort's gather into decode). Records why gather_list is the one candidate that does not drop into the frame unchanged: K4 maps one block per chunk and relies on the mask->indices wave for two invariants -- ids partitioned by chunk (the kernel derives in-chunk position as idxs[k] - chunk_start) and per-chunk bitpack scalars loaded once per block. The output slot is already gather semantics; the partitioning is what breaks, so an arbitrary permutation needs either a stable pre-pass by chunk plus a scatter, or a launch over the list with per-element metadata lookup, losing both the scalar amortisation and coalescing. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The converters reported two facts about a decode by constructing two different
representation subclasses, and the scan recovered them with two dynamic_casts
testing type identity. Encoding a boolean in an object's type meant clone()
could not preserve it -- it "intentionally degraded" -- so the information was
only valid between the conversion and the scan's capture of it, which happened
to be adjacent on one thread.
Replace both marker classes with one decoded_batch_representation carrying a
decode_outcome{row_filtered, rule2_bailed}. The scan does one dynamic_cast to
fetch a struct instead of two to test identity, and the outcome survives a copy
because it is a property of the decode, which a copy shares.
The plain gpu_table_representation is still used when there is nothing to
report, so the ordinary path and the gate-off path are unchanged.
This is the P6 half of the plan's Phase 1; P5 (the BOOL8 type sniff in
parquet_gpu_ingestible) follows, and adds the substituted column list to the
same outcome.
Tests: sirius_unittest [compression] 1222 assertions, [scan] 274587
assertions, test_render_signature_contract PASS.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…r it build_filter_expression_for recovered which columns the decode had substituted by testing each candidate for cudf::type_id::BOOL8. That inference is only unambiguous while candidates are VARCHAR-only: extend the equality pushdown to numeric or boolean equality -- which the range work is adjacent to -- and a genuine BOOL8 column becomes indistinguishable from a substituted one, silently rewriting a filter conjunct into a bare boolean reference to the wrong column. The converter knows the answer exactly, so report it. decode_outcome gains predicate_columns, stamped onto the split by prepare_for_processing and forwarded to post_filter_and_project through filtered_table. The scan maps the reported positions onto the primary indices it nominated and rewrites only those. The BOOL8 test is gone from the scan; type substitution is no longer something the consumer has to detect. Covered end to end by the three "predicate pushdown" cases in sirius_unittest [compression] (9-11 of 36), which run real SQL and check the aggregate over rows the substituted predicate selects. Tests: [compression] 1222 assertions in 36 cases, [scan] 274587 assertions in 265 cases. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…hase 2 Closes P5 and P6 in the problem list with their commits, records what Phase 1 did and did not deliver (position and unapplied need the facade, so they land in Phase 2), and notes the caveat: predicate_columns is populated only on the compression converter path, so it is a converter fact rather than a universal one until the facade generalises it. Adds section 6b so Phase 2 can be picked up without this session's context -- which files to read in order, the shape of the change, the five pushdown setters that make P7 the load-bearing part, and the test loop that actually exercises the paths rather than just compiling them. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…ssion
The scan used to hand the decompressor five separate carriers -- an equality
pushdown, a range pushdown plus its coverage flag, a membership pushdown plus
its generation -- through five setter/getter pairs on each of the two compressed
representation classes, with clone() obliged to copy every one. It also digested
its own filter three times, through three extraction functions whose result
types leaked codegen's range_predicate and pair_compare_op into the scan's public
header, and the converter then re-assembled the decoder's request inline in ~190
lines that both converters duplicated.
Replaced by two objects and one entry point:
sirius::scan_decode_request what one scan asks of a decode: per column, an
equality set to answer in place, bounds to drop
rows against, and join filters to test -- one
struct per column instead of four parallel
vectors.
sirius::compressed_scan the request as the decoder holds it. Immutable
and shared by every batch; the per-batch
adjustments (a fresher join-filter snapshot,
giving up compaction, narrowing to what a chunk
can honour) hand back a new one.
decode_compressed_chunk() decodes a chunk under a scan and reports what it
managed to apply. Never returns null: every way
the filtering declines ends in the plain decode.
Both representations now carry one shared_ptr<const compressed_scan>, so clone()
copies a pointer and cannot forget a field. Both converters shrank to a single
call. The three extraction functions merged into one analysis pass
(op::analyze_scan_filters) run once per scan by the ingestible and only mapped
onto slots by the scan manager, which drops the duplicate extraction the manager
was doing; scan_utils.hpp no longer includes the codegen selection header.
The env gate and the selectivity ceiling had been copied per translation unit
and had drifted -- two readers accepted only "1" where the decoder accepts
anything but "0", so a value like "true" turned the feature on in one layer and
off in another. They now live in one place (decode_filter_policy.hpp).
Behaviour is otherwise unchanged, with two deliberate exceptions:
- coverage of the whole filter now also requires every range to reach a decoded
slot. A range that maps to no served column was silently dropped while the
batch could still be tagged as needing no further filtering.
- the parquet ingestible's candidate extraction ran twice (a merge duplicated
the block), appending each candidate to the position list twice.
Names lost their branch-local shorthand: rule2_bailed is now
selection_unprofitable, "the fused scan-filter pipeline" is decode-time
filtering, the diagnostic prefix is [decode-filter], and
row_filtered_table_representation.hpp is named after the class it actually
holds.
pixi run make test: 2403 passed, 1 skipped. [compression] and [scan] also pass
with SIRIUS_EXP_FUSED_SCAN_FILTER=1.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The decode-time filtering code was written in the shorthand of the branch that
built it: kernels called K1/K1m2/K3/K4/K5/K6, the enable policy called RULE 1
and RULE 2, worklog owners W1-W4, "iteration 3/4/5/7", "STATUS-W2", "track A/B",
and "bail" for a decision that is really "stop compacting". None of it is
derivable from the code, and Phase D had already given the kernels real names --
an enumerator x consumer product -- so the K-numbers were a second, redundant
taxonomy to keep in step by hand.
The shorthand is now spelled out at every use:
K1 / K1m2 the range ballot / the pair ballot
K3 / K4 the mask walk / the index walk
K5 / K6 the dictionary gather / the masked str_split route
RULE 1 the static output-shape check
RULE 2 / bail the selectivity ceiling / giving compaction up
classic the ordinary (unfiltered) decode
W1-W4, iteration N, STATUS-W2, track A/B deleted; they name people and
worklogs, not behaviour
Identifiers followed: tier_dict_k5 -> tier_dict_gather, tier_str_k6 ->
tier_str_split, scan_filter_status::bailed_high_selectivity ->
declined_unselective, tier_is_fused_capable -> tier_decodes_compacted, and the
orchestrator's k4_pick -> index_walk_pick, rule2_bailed -> unselective.
Also deletes column_decode_directive and make_scan_filter_request, an adapter
that collapsed the output-shape tags to a boolean and had no callers left once
the scan-side facade started building requests directly.
Comments only, plus those renames -- no emitted kernel source or control flow
changed. pixi run make test: 2403 passed, 1 skipped; the two codegen host tests
(masked_decode_variants, render_signature_contract) also pass.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Marks the facade done with what it did and did not deliver: position/unapplied
still need the analysis to speak sirius::ast, so they move to Phase 3, and the
residual is still a rebuilt DuckDB expression. Closes P7 and part of P8, and
notes that the gate readers had genuinely drifted before consolidation
("true" enabled the feature in one layer and not another).
Answers open question 1 by construction rather than by decision: the carrier
types are the facade's own and the conversion to codegen:: happens at one point
inside compressed_scan.cpp, so no shared header is needed.
Replaces the Phase 2 cold-start brief with a Phase 3 one — what is left of
P1-P4 and P8 and where it now lives — and records the test loop that actually
covers this area, including the ninja line simpatico's own ctest targets need
(without it ctest reports "Not Run", which reads like a pass).
Adds a fifth instance to the merge-duplication risk: the parquet ingestible's
candidate extraction existed twice and the compiler could not catch it.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Replaces the four encodings of "how does this column decode" with one enum and
the six plan probes with one call.
output_tier (5 values) + a parallel compact_capable vector
+ decode_selection's compact_capable/dict_compact/str_compact bools
-> codegen::decode_route { full, bitpack_mask, delta_mask, dict_codes,
str_split }
plan_supports_selection_decode / _dict_ / _str_ / _predicate_decode,
plan_selection_tier, column_supports_predicate_decode
-> simpatico::probe_column(tree) -> column_decode_caps
`full` IS "not compactable" and can_produce_mask() IS `route == bitpack_mask`,
so the invariants that were comment-enforced ("umbrella true iff classifier !=
tier_b", "the compacted modes are mutually exclusive") are now unstateable — the
runtime check for the latter is deleted along with the enum that allowed it.
decompress_column and the orchestrator both verify a requested route against
probe_column instead of re-deriving capability from a different probe than the
one that picked the route.
[compression] passes with the gate on and off; all 14 simpatico tests pass.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The four conjunct kinds each had their own counting loop, their own emission loop and, for the join filters only, their own ordering sort. The count fed a shared cap, so the arithmetic had to be kept in step by hand across three places. They are now one `selection_source` list built once: the count and the cap read its size, one comparator states the ordering rule (exact conjuncts, which cost no probe launch, ahead of join filters ranked by kind then build-side key count), and one switch emits it. Same order out as before, so no behaviour change — but the rule is now stated instead of implied by which loop ran first, and the diagnostic prints every source in the order the decode will keep them. [compression] passes with the gate on and off. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The decode handed back a ragged column vector — compacted-route columns survivor-sized, full-width ones not — and a separate compact_scan_filter_output call reconciled them. Nothing in the types said the second call was required, and no caller could have known which column was which. decompress_scan_filter now returns a cudf::table that is already uniformly survivor-sized, doing the reconciliation itself, and compact_scan_filter_output drops off the public API. If the assembly refuses (a null-masked column, a mis-sized output) the call falls back to the unfiltered decode and reports status = failed with a message, so a caller cannot end up holding a half-filtered batch either. [compression] passes with the gate on and off; all 14 simpatico tests pass. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Six env knobs were read in three translation units, in three different parse styles, with the defaults kept in sync by comment. Phase 2 consolidated the two the scan side needed; the rest still lived in the orchestrator, so the gate and the 0.35 ceiling each had two definitions. They now live once, with the decode that acts on them (codegen/selection/decode_policy.hpp): the gate, the diagnostic switch, the two selectivity ceilings, the index-walk crossover and the join-filter cap. The sirius-side header re-exports the two the scan uses instead of re-reading the environment. One parse helper each for flags and fractions, so "set it tiny is a kill switch" is a property of the parser rather than of each copy. [compression] passes with the gate on and off; all 14 simpatico tests pass. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Marks P1-P4 and P8 closed with their commits, and records what did NOT get done: unapplied/position still need the analysis to speak sirius::ast, and no phase since D has been measured — this box has no GPU budget, so the "no measurable delta" clause on Phases 2 and 3 is unverified and blocks merging. Replaces the Phase 3 brief with what is actually left: measure, then either the AST residual or §7's capability work. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
post_filter_and_project rebuilt its filter per batch: when the decode had answered a column in place, it re-ran convert_table_filters_to_expression with a substitution set to get a fresh DuckDB expression tree, then lowered that to Sirius AST — two full tree allocations on the batch path, and the substitution expressed by rebuilding rather than by choosing. The filter is now decomposed at bind (decompose_table_filters, which convert_table_filters_to_expression is reimplemented on top of, so the conjunct sets cannot diverge) and each conjunct lowered to AST once. Per batch, residual_filter::against picks a form per conjunct: a bare reference for the columns the decoder answered, the prebuilt comparison for the rest. No DuckDB expression is built and nothing is converted on the batch path. A conjunct that cannot be lowered now fails at bind. Before, the lowering ran per batch and a failure handed the evaluator a null AST to dereference; the straightforward port of that would have been an empty residual, which reads as "this scan has no filter" and silently returns unfiltered rows. make test: 2403 passed, 1 skipped ([scan]'s 265 cases cover this path); [compression] passes with the gate on and off. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Every equality the decode answers is ANDed into the row selection before wave 2 runs, so on an applied decode the surviving rows already satisfy it. The scan nevertheless re-evaluated it: the answer was delivered as a BOOL8 column purely so the residual could AND back a condition that could no longer be false. decode_outcome now reports predicates_enforced — the decode both answered those conjuncts and applied them — and the residual drops them instead of referencing the answer. False whenever the filtering declined and the answers came from the plain predicated rerun, which drops no rows; those still have to be evaluated. When that leaves nothing, post_filter_and_project marks the batch ROW_FILTERED rather than passing it through: a null residual means "already filtered", not "no filter". A scan whose whole filter is one dictionary equality now skips the post-decode pass entirely. The enforced path is unreachable under default thresholds at the sizes the end-to-end tests use (they decline at 0.25 selectivity against a 0.10 ceiling), so test_residual_filter pins the decision itself: keep / reference / drop, plus the empty case and a column the scan never nominated. Mutation-checked — removing the drop fails 3 of the 5. Forcing the applied path with a raised ceiling also leaves both dictionary pushdown tests' answers unchanged. make test: 2408 passed, 1 skipped; all 14 simpatico tests pass. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
unapplied is done: the residual is a Sirius AST predicate assembled per batch from conjuncts lowered at bind, with three outcomes per conjunct (keep / reference / drop). Notes that the drop case needed predicates_enforced and is unreachable at default thresholds, so it is pinned by a unit test rather than end to end. position is not done: the conjunct side landed, but never delivering the column needs a variable output arity. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…ange
The decode-range extraction had its own constant-lowering stack — decoded_bound
{int128 floor, ceil}, range_accumulator{int128 lo, hi} and
physical_integer_payload's switch over DuckDB physical types — alongside
helper/numeric_narrowing.hpp's sirius::numeric_range, which is the same int128
min/max with a domain tag and a decimal scale, and whose header calls itself
"the single fitting authority" for exactly this.
Now: numeric_range is the accumulator, the lowered bound and the return type,
and ast::constant_numeric_range does the payload extraction — it already knows
which payload alternative a declared type may carry and rejects a constant whose
payload disagrees. What stays local is the part narrowing has no notion of:
restating a constant at the COLUMN's scale, where floor and ceil straddle a
constant the column cannot represent.
Two behaviour changes, both widening:
- a UBIGINT constant is now accepted where physical_integer_payload refused it.
Sound: bounds outside int64 clamp to the full domain or to a provably empty
range, which is what those conjuncts mean on an int64-decoded column.
- DATE keeps a local arm; constant_numeric_range covers integer and decimal
literal domains only.
INT128-physical constants stay refused: the decoded domain is int64 and
rescaling one by a power of ten could overflow the accumulator before the clamp.
This logic had no direct test — it is reached only through a compressed GPU-tier
pin with the gate on — so test_scan_filter_ranges pins it first: strict vs
non-strict bounds, the empty range, DATE's day count, decimal rescale in both
directions and the equality no representable value satisfies, and what clears
whole-filter coverage. Mutation-checked: dropping the ceil bump fails 2
assertions, dropping a strict inequality's ±1 fails 2.
[compression], [scan] and all 14 simpatico tests pass.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
They bypass Sirius's ~62 typed SET options: invisible to duckdb_settings(), not settable per session, and cached on first read — which is why some decode paths are only reachable by a unit test rather than end to end. The obstacle is layering, so noting it rather than papering over it: simpatico has no DuckDB dependency, so the values cannot be pulled from the setting registry where they are read. They would have to arrive from Sirius, pushed down per query or passed into the decode call as a policy value. Deferred while the feature is experimental. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
None of combine_masks_and, run_selection_cnt, mask_to_row_indices or mask_from_bool8 had a direct test. Their only exercise was through the orchestrator, and only when the filtered decode applies — which the end-to-end tests mostly decline on selectivity. combine_masks_and was almost certainly never executed at all: the decode calls it only with two or more mask sources, and no test builds a request with more than one. Each is now checked against a host reference computed independently, plus the two invariants everything downstream assumes and nothing else would notice losing: tail-zero past num_rows (the destination is pre-filled with 1s so a kernel that skips the tail cannot read as zero by luck), and strictly ascending row ids partitioned into their own chunk's offset window. Empty and all-survive selections get their own cases — the first because zero survivors is a legitimate outcome rather than an error, the second because an identity permutation is where an off-by-one still looks plausible. Mutation-checked: dropping the num_rows guard fails the tail case, shifting the per-word output base by one popcount fails the ordering cases, and reducing the AND to a copy fails the combine case. 15 simpatico tests pass. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Evaluated all four against their cudf equivalents. None fits, for one structural reason worth stating rather than rediscovering: these are in-place, chunk-segmented operations over a caller-owned arena with a padding invariant, and cudf's bitmask API is allocating and whole-column — bitmask_and and bools_to_mask return fresh buffers sized by bit count, segmented_count_set_bits returns a host vector where we copy back 4 bytes, and a copy_if would redo the scan chunk_offsets already holds. The closest call is combine_masks_and, where cudf::detail::inplace_bitmask_and does have the right ownership; noted as the one to revisit if cudf promotes it out of detail. Also adds a reuse-review section to the plan recording what was folded, what is deferred, what is still open (the stream pools, the two TableFilterSet walkers) and one missed connection rather than duplication: a zone-map filter's bounds could feed the range ballot instead of a full-width probe. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…lkers collect_null_prune_predicates and decompose_table_filters are NOT duplicates and should not be merged — the first exists precisely because the second drops IS_NOT_NULL, so one keeps what the other skips. What they did duplicate is the bookkeeping underneath: column_index -> primary index -> is it ours -> which batch column. That was stated inconsistently. An out-of-range column_index skipped in one and threw std::out_of_range from .at() in the other; the bounds check on the batch-position map existed in one and not the other. resolve_filtered_column now states it once, with three outcomes rather than two, because the third is where the callers legitimately differ: a conjunct that must be EVALUATED cannot reference a column that was never materialized (a wiring bug worth failing on, with the primary index reported), while a filter used only to PRUNE drops silently. Each caller decides that for itself; neither re-derives how to find the column. Tests cover all three outcomes plus both skip reasons, and that decompose throws rather than silently dropping — dropping there would return rows the filter rejects. [scan] and [compression] pass. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
pin_table's compress_pool and compressed_scan's decode_pool were byte-identical apart from the error string, and both were a plain `thread_local stream_pool` initialised on first use. stream_pool::init creates its streams on whatever device happens to be current at that moment and then never re-creates them. CUDA streams are device-bound — simpatico's own stream_cache says so and keys its recycling by device, and test_multi_gpu_stream_affinity exists to hold that property — so a thread that decoded on device 0 and later worked on device 1 would submit device-1 work onto device-0 streams. The two pools were also independent, so a thread that both pinned and decoded held 8 streams from two thread_locals that did not know about each other. Both now call simpatico::thread_device_stream_pool(n), which keys one pool per (thread, device) and keeps the never-destroyed lifetime the buffers depend on: a buffer records the stream it was built on for its eventual async free, so the handle has to outlive it. NOT VERIFIED on hardware: this box has one GPU, so the multi-device path still self-skips. What is covered is what one GPU can check — that repeated calls return the same pool (the reuse the lifetime argument rests on) and that a second thread gets its own handles. make test: 2414 passed, 1 skipped; 15 simpatico tests pass. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
There is exactly one manager thread per GPU executor, and it performed the blocking make_reservation (and the downgrade .get()) itself while holding a reserved pool slot: one query's memory-hungry task stalled every query's dispatch to that device. The F1 fair pops made the QUEUE order fair, but a fair pop still stalled behind the manager's blocked reserve. process_task now only resolves the completion handler, attributes the slot to the task's query, and dispatches; the reservation, the downgrade-on-shortfall dance and the execution run on the pool worker (prepare_and_execute), where the slot attribution already covers them for the per-query drains. The manager's only blocking points are reserve() and pop(), both interruptible — which also makes the error-path quiesce bracket's manager join prompt instead of hostage to a memory wait. At most ONE task per executor parks in a memory wait (_memory_waiter_parked, preserving the historical one-blocking-reservation-per-device arbitration and keeping num_threads-1 slots dispatchable); overflow waiters re-queue through the executor's own queue after a 10 ms worker-held backoff, so a hungry query with many tasks cannot re-create the head-of-line blocking one level down, and the bounded backoff plus the F1 rotation rule out a requeue spin. The requeue rides the same lifecycle gate as the OOM reschedule and never touches gpu_reservation_max_retries accounting. Telemetry stays FSM-legal: a requeue emits routing (the queued->routing->queued loop of a scheduler->executor hop); reserving is entered only once the task commits to acquiring, and the manager-thread uuid is captured by value because a quiesce/resume bracket destroys the manager_loop stack frame the wrapper lives on. Evidence (test_gpu_pipeline_executor_memory_wait.cpp, real GPU space with a test-held reservation): pre-fix, 4 short tasks scheduled behind one parked hungry task never dispatch (0/4 in 30 s, manager blocked); post-fix all 4 complete (first at ~0.1 ms) while the hungry task is still waiting, and a second hungry task re-queues (tasks_requeued_on_memory_wait metric) instead of consuming the last worker slot. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…wn-safe schedule_lookahead hard-coded _query_task_global_states.begin(), so only the OLDEST live query ever received lookahead warm-up — every newer concurrent query started cold (the code carried its own TODO). It now rotates round-robin across the ACCEPTING queries, reusing the F1 rotation contract: a cursor remembers the last query served, the scan starts after it (upper_bound, wrapping, stale cursor harmless), quiescing/closed queries are skipped via the lifecycle registry, and within one call each live query is tried once — so a query with nothing warmable cannot pin the rotation and starve the rest. Single-query behavior is unchanged: one request per call, in queue order. The accounting half: schedule_lookahead dereferences the query's operators and pushes a creation request under only lookahead_mutex, from the task scheduler's management thread — a producer the per-query drains did not order themselves against. drain_pending_tasks now clears the lookahead queue under that mutex FIRST, before its request drain, which mirrors what slot-attach gives the creation path: a walk in progress blocks the clear (the plan outlives drain_pending_tasks, so the operators are alive) and its push lands before the drain, which drops it; a walk starting after the clear finds the queue empty and is a no-op. Previously (clear LAST) a racing push could land after the drain and survive reset(query_id) — a stale request holding a raw operator pointer into a plan the caller destroys next, with nothing waiting on the worker that would eventually dereference it. Evidence (test_task_creator_lookahead.cpp): the rotation and unwarmable-query cases fail pre-fix (all requests land on the oldest query); a looped schedule_lookahead-vs-reset race asserts the creation queue holds nothing for the dead query under every interleaving. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
batch_telemetry_registry::on_query_end() took no query id: it consumed every live placement across all 16 shards as reason=query_end and cleared every consumer port, so query A's end silently truncated query B's telemetry for the rest of B's life. Placements and ports now carry the owning query id — stamped from the registering port (emit_plan_telemetry knows the window id) or, for lazily-registered reschedule claims, from the claiming pipeline (sirius_pipeline::get_query_id()). on_query_end(query_id) consumes only that query's placements and erases only its ports, and returns the drained count for observability/tests. The ALL-queries drain survives as on_all_end(), used by terminate() and uninstall() — never by a single query's cleanup. Unit test: two queries' placements live; one ends; the other's placement AND port survive, keep transitioning, and close with their own reason (processed), verified down to the ndjson consumption records. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…n slot The plan-time SlotGuard in install_transparent_execution made plan GENERATION consume a full execution-window slot: with the counted gate saturated by executing queries, a peer could not even be PLANNED (OnFinalizePrepare parked in acquire_query_lifecycle_slot until a window freed). Planning is CPU-only; it needs the runtime-health/cancellation fast-fail, the E1 config snapshot, and a consistent pin-table view — which the scan manager's _pinned_entries_mutex + owning shared_ptr<pinned_entry> probes provide per call. New SiriusContext::PlanViewGuard: same fast-fail split as the slot acquire (stable runtime-unavailable error, InterruptException) and the thread-local config snapshot, but NO slot — plan generation no longer blocks on, counts toward, or is counted by query_lifecycle_peak or admission. The execution window (StandaloneQueryScope) still acquires its slot when execution begins, so lifecycle open/close bracketing is untouched (SlotGuard never opened lifecycle entries). To make the slot-free plan view airtight, find_pinned_entry_for_parquet_files now returns an OWNING shared_ptr (it returned a raw map pointer readable only "inside one slot-scoped window" — already unsound at slots > 1); the planner holds it across the residency gate like the duckdb probe. The unpin SlotGuard comment no longer claims plan generation is slot-serialized. Validation (test_concurrent_adversarial.cpp): with ONE slot, a heavy query holds the window while a peer's Prepare (full transparent rebind) completes strictly inside the hold — zero such prepares before the fix — and single-slot peak stays 1. Also adds the F5 isolation scenario: merge re-pins of table X beside continuous queries on table Y only, asserting zero spurious "currently reading" rejections. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…dows Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…an serve try_match_cached_entry snapshotted the WHOLE pin table (a shared_ptr to every pinned entry) before matching outside the lock, so any query anywhere inside a match held a use_count on every entry — and a concurrent re-pin MERGE (or MVCC attach) on a completely unrelated table spuriously failed its serving-refcount guard with "a query is currently reading the pinned entry". The identity gate (can_serve_with_columns: format / file-set / table + column-superset compares — cheap string/vector work) now runs UNDER _pinned_entries_mutex, and the snapshot takes references only to entries that pass it. Queries on table Y therefore hold no reference to table X's entry: zero spurious rejections is structural. The heavy match work (validation, zone-map plan building, the MVCC branch) still runs outside the lock on the owning snapshot, exactly as before. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…ed identity _pinned_entries and attach_mvcc_metadata are keyed by the bare user-supplied pin name, but a name says nothing about the source: two live same-named pins over DIFFERENT sources — bare-name duckdb pins from two ATTACHed databases (name == table ref, resolved via the search path), or two parquet file sets pinned under one name — fell into the merge/replace path. Same row count meant the second database's columns were silently spliced into the first database's entry and its MVCC metadata (v_base, checkpoint iteration from the OTHER database) was attached over the first's, so a query on the first table could serve the second table's data. The name stays the user-facing handle (unpin by name, per-name uniqueness); the merge/replace path now compares the RESOLVED identity — cache_entry_info::compression_plan_key(), the same catalog.schema.table or canonicalized-file-set key the serving matchers and the E5 plan registry use — and rejects a same-named pin of a different source with a clear error naming both identities. attach_mvcc_metadata re-checks the caller's identity key (defense in depth; the pin registry lock plus the insert guard already make a foreign attach unreachable from the pin flow). Test: two ATTACHed DBs with same-shaped, differently-valued tables; the bare-name collision is rejected, the first entry is provably untouched, and both pins live simultaneously under distinct names with each table serving its own values. Parquet same-name collision covered too. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
… the triage Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…per query, not per device Two operators sized against whole-device/whole-host free memory as if they were alone, so concurrent queries overshot together: - SORT_SAMPLE derived its partition count from get_available_memory() — two concurrent sorts each budgeted max_sort_partition_memory_fraction of the SAME free bytes. It now divides the free-memory read by the number of live queries (query_lifecycle_registry::size(), plumbed from SiriusContext at plan time). Each sort sizes against its 1/N share: more, smaller partitions under concurrency — slower, but every query completes. The registry pointer is nullptr in unit tests (no context), preserving single-query behavior bit-for-bit; SiriusContext owns the registry and outlives every plan, so the pointer cannot dangle. - The materialized result collector picked its HOST space by max_element over free bytes and reserved AFTER: two collectors reading the same snapshot picked the same space and the loser proceeded unreserved even when another space had room. It now reserves FIRST — probing every host space with make_reservation_or_null in most-free-first order and using whichever space granted it; the pre-existing unreserved WARN fallback remains only when every space refuses. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…ream work from operators The build does not enable per-thread default streams, so legacy default-stream work implicitly synchronizes with every blocking stream on the device, and a cudaDeviceSynchronize stalls every co-resident query. Four operator-layer sites, each with its ordering argument: - hash join, dynamic-filter publish fallback: a build batch with no writer event used to trigger cudaDeviceSynchronize(). The producing stream is unknown by construction there, so no narrower sync exists; publication is best-effort by contract (the non-GPU-residency path already skips), and publishing UNORDERED could read half-written keys — so skip publication with a WARN instead of stalling the device. - hash join, orphan-pairing empty batch: make_data_batch recorded the writer event via cudf::get_default_stream() (value 0 — no event at all, which is what armed the fallback above). Record it on a pool stream from the batch's own space instead; the empty table owns no device data, so any live stream on the right device correctly publishes it. - parquet footer probe (build_file_scan_info): the AST-literal device scalars and the row-group stats filter ran on the legacy default stream. They now run on a function-local non-blocking rmm::cuda_stream: every consumer is on that same stream, filter_row_groups_with_stats returns host data (internally synced), and the stream is declared before the translated expression so it outlives the stream-ordered scalar frees. - string decode exact-total readback: blocking cudaMemcpy (legacy-stream semantics) replaced by cudaMemcpyAsync on the task's stream + that stream's sync — the CUB scan producing the value ran on the same stream, so same-stream ordering covers the producer/consumer pair. Deliberately KEPT device-wide: SiriusContext::terminate() teardown syncs, the defragmenter's post-cudaMemPoolTrimTo sync, and the reservation-manager destructor drain — process-teardown/OOM-emergency paths where quiescing the whole device is the point. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…sult collectors
Two [sizing] scenarios on the adversarial harness (concurrent_test_utils):
- concurrent full-table ORDER BYs on a small absolute pool. Results are
verified client-side — exact row count, exact column checksums, and
lexicographic (v, id) sortedness — instead of comparing multi-million-row
reference strings. Asserts overlap (query_lifecycle_peak > 1) and zero
runtime fallbacks: a concurrent-overshoot OOM completing via the CPU
fallback would hide exactly the regression the scenario pins.
- concurrent wide-result queries, so every result collector wants host
space for its GPU->HOST clones in the same window. Asserts all workers
complete with the right row counts, zero runtime fallbacks, and that the
collector's unreserved-fallback WARN ("proceeding without reservation")
never fires — with reserve-first ordering, every clone at this sizing
must obtain a host reservation.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…ery lifecycle - pipeline-execution.md: task_scheduler::prepare_for_query is gone (per-query state registers via task_creator); per-query drain_after_error(query_id) / wait_for_completion(query_id); C4 manager loop (reservation moved to the worker, single memory-wait slot, re-queue metric); per-query OOM retry cap (gpu_reservation_max_retries, 50 ms backoff); queue type is the multi_index_priority_queue with fair pops. - architecture-overview.md: same deleted API; per-query ownership hierarchy (data_repository_manager_registry, query_lifecycle_registry, no context-owned query); C4 thread model; admission/cleanup lifecycle steps. - task-creator.md: per-query query_task_global_state (prepare_for_query / reset(query_id)); _kiosk is gone (bounded_thread_pool with per-query slot attribution); drain_pending_tasks(query_id) ordered steps; lookahead warm-up section (D3); per-query active_gpu_ids. - scan.md: per-query query_scan_manager_state, scoped_dispatcher, reset(query_id), max_concurrent_queries, prefetch-cache epochs (F3). - optimizations.md: _kiosk reference; PR sirius-db#507 counters that no longer exist; monitor->request->two-tier sweep replaces run_downgrade_pass; retry cap. - execution-flow.md: engine-owned completion handler + create_query registration; C4 executor steps; per-query error drain; sequence diagram. - memory-management.md: attributed request_downgrade(query_id, ...), TIER-1 cross-query sweep + sweep token, gated TIER-2 re-push (B1). - data-management.md: one shared_data_repository_manager per query; for_each_repository (deleted in cucascade hygiene) replaced by the real API. - multi-gpu-architecture.md: _active_gpu_ids member became per-query state. Verified against integration/concurrency-full @ 9061821. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
New docs/super-sirius/concurrency-model.md covering the execution model the concurrency effort produced, verified against integration/concurrency-full @ 9061821: - query identity (query_id_t, priority bands, the 2^31 band wrap) - admission: the counted slot pool sized by scan_manager.max_concurrent_queries, StandaloneQueryScope/SlotGuard windows, window-id wrap guard (H8), query_lifecycle_registry open -> quiescing -> closed - per-query state: query_task_global_state, query_scan_manager_state, data_repository_manager_registry, parked plans, prefetch epochs - config snapshots (E1/E2/E3) and SET-affects-queries-admitted-after semantics - scheduling: multi_index_priority_queue indexes, F1 fair pops, the C4 non-blocking manager + single memory-wait slot, D3 lookahead rotation - memory pressure: attributed downgrade requests, per-query drain(query_id), TIER-1/TIER-2 sweeps, the sweep gate + wait_inflight_request fences (explicitly marked interim pending steps 6+7), the retry-cap knob - query end: run_mandatory_cleanup's ordered steps, the drain_after_error path with the quiesce/resume bracket (marked volatile), D5 classification - transparent execution: per-connection capture, OnFinalizePrepare, PREPARE/EXECUTE interception, CPU fallback, E4 immutability, E7 RAII - testing: bring-up harness, adversarial suite, co-tenancy rules README gains the TOC row and reading-order slot. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
One clearly-marked addendum block at the top of 99-execution-summary.md pointing readers at 01-bringup-triage.md (current per-item status) and the new docs/super-sirius/concurrency-model.md; the original content is untouched. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
The docs refresh in the previous commits was verified against integration/concurrency-full @ 9061821. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…rent code The enforcement chain is unchanged, but the quoted layers drifted: prepare_for_processing returns void and stores read_only_data_batch accessors (throwing -> oom_reschedule instead of returning nullopt), lock_or_prepare_batch returns optional<read_only_data_batch>, the enclosing function is gpu_pipeline_task::execute, and the Phase-15 INVARIANT comments were dropped from the existing read sites during later refactoring (noted; the convention stands for new code). Stale line-number citations replaced with function-level ones. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…cking primitives Bump the cucascade pointer to branch f9/fair-waits (based on the previously recorded cc7badcb), two commits: - 20c8cf2 fix(memory): FIFO ticket handoff for exclusive_stream_pool BLOCK checkout. A released stream goes to the longest-waiting caller; a release-and-reacquire caller queues at the tail; GROW callers mint a fresh stream instead of taking one a parked waiter is owed. - a9b7aba fix(memory): per-space FIFO wait lists for blocking reservations. notification_channel gains a ticketed scoped_waiter (NOTIFIED only ever goes to the head of the list); make_reservation keeps its fast path when nobody is parked and otherwise joins the FIFO instead of barging, so a heavy query's release-and-re-request loop can no longer perpetually beat a light query's single wait. make_reservation_or_null/_upto keep their try-semantics. Also removes memory_reservation_manager's never-notified cross-space _wait_cv (a latent permanent hang when no candidate space exists); request_reservation returns nullptr for that case, which every caller already handles. API shapes are unchanged on both primitives; this is an internal wait-discipline change only. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…e pool itask_executor::wait_and_drain_query ended in _bounded_pool->wait_all(), so one query's error-path cleanup waited for EVERY query's in-flight tasks — including a co-tenant's parked memory wait, which lasts until memory frees somewhere (flagged during the C4 fix). The bracket's manager join closes the pop-to-attach window: after it, no task is in-hand, so every task of the failing query is either running under a slot attached to it, queued (swept by the per-query drain), or was dispatched untagged (no pipeline — its query is unknowable). The wait is now bounded_thread_pool::wait_for_query_and_untagged: this query's attributed slots plus the untagged count, which can only shrink while the manager (the sole reserve() caller) is down. Co-tenants' attributed slots are ignored, so a peer's parked wait no longer extends the erroring query's cleanup. The invariant is preserved verbatim: when wait_and_drain_query returns, no thread is still executing a task that references the failing query's plan. The pool tracks attached_active_ alongside active_ (untagged = difference), attach/release notify the per-query CV on untagged transitions, and active_untagged() is exposed as a test/diagnostic aid. quiesce_manager() no longer waits for pool work itself — the caller picks the wait — and stale "kiosk" wording in the executor/pool docs now describes the bounded_thread_pool reality. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
- test_reservation_fairness.cpp: a heavy caller's release-and-re-request loop cannot starve a parked waiter (served within a bounded number of grants); parked waiters are served in arrival order (step-released holds make the expected order fully deterministic); shutdown wakes EVERY parked waiter, not just one. - test_stream_pool_fifo.cpp: BLOCK checkout serves waiters in arrival order; a release-and-reacquire caller queues behind a parked waiter; GROW mints a fresh stream instead of taking one a parked waiter is owed. - test_task_executor_error_bracket.cpp: wait_and_drain_query returns while a co-tenant's memory wait is still parked (times out under a whole-pool wait), with the plan-safety invariant asserted — the failing query's running task completes before the bracket returns and its queued task never runs; an untagged (pipeline-less) in-flight task is still waited for conservatively. - test_bounded_thread_pool.cpp: wait_for_query_and_untagged waits out untagged slots but not co-tenants; active_untagged tracks the reserve-to-attach window. Every case unblocks and joins its threads before asserting, so a failing discipline reports cleanly instead of hanging or terminating on a joinable thread; the GPU-backed fixtures use small absolute pools (256 MiB) and skip with a WARN when no device is usable. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…-end
Every path that used to BORROW a data_repository* across blocking work now
takes shared ownership, so a query's teardown never has to fence against a
borrower and no borrower can dangle:
- cucascade (submodule -> ed8d315, branch steps67/shared-ownership, based
on a9b7aba): get_repositories() returns shared_ptr snapshots; repositories
gain a destructor-side leak callback; the manager attributes those reports
per {operator_id, port_id} via set_leak_handler().
- TIER-1 downgrade sweeps own what they borrow: the manager snapshot AND the
per-manager repository snapshot are shared_ptr for the sweep's duration
(batches were already shared_ptr inside each candidate). The interim sweep
gate is still in place but no longer protects anything on this path; it is
retired in the next commit.
- convertible_data_batch_provider holds its repository by shared_ptr.
- Pipeline wiring: operator ports co-own their repository (port::repo_owner;
the raw port::repo stays as a cached alias so existing operator readers are
untouched). repository_wiring_materializer wires both.
- gpu_pipeline_task::_data_repos is vector<shared_ptr> (B4): a task crosses
queue hops, OOM-reschedule sleeps and TIER-2 extractions; it now keeps its
destination repositories alive instead of carrying raw pointers.
- data_repository_manager_registry::erase(query_id) drops ONLY the map entry:
in-flight holders keep their repositories alive until they naturally finish.
Leaked-batch accounting moved to the repository destructors, attributed by
query via the handler installed in create_for_query(); run_mandatory_cleanup
and drop_query_runtime_state_best_effort no longer consume a leak report
from erase().
Tests: registry tests updated for the void erase (plus a new
erase-destroys-unborrowed-manager case); provider/task construction sites
migrated to the shared types.
Closes B3 and B4 (issue register group B).
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…time fence survives
With sweeps co-owning everything they borrow (previous commit), the
registry's interim sweep gate no longer protects anything:
- data_repository_manager_registry loses begin_sweep()/sweep_guard and the
teardown-intent bracket; erase()/clear() drop the map entry (or map) and
release outside the lock, never waiting for a sweep. The downgrade
executor's TIER-1 sweep no longer takes a token.
- RE-DERIVED FENCES, documented in downgrade_executor.hpp and
run_mandatory_cleanup: wait_inflight_request() SURVIVES — it is a PLAN
fence, not a repository fence. A peer's in-flight TIER-2 pass can hold the
ending query's task in a convertible wrapper across a blocking conversion;
the wrapper's gate-refused drop DESTROYS the task, and ~gpu_pipeline_task
walks the plan (mark_task_completed -> notify_downstream_pipelines).
Plan parking (B5) defers the plan's death to cleanup, so cleanup must not
destroy it while such a wrapper is alive. The extraction-time-keys gate
check stays with it: together they bound the wait (requests starting after
quiesce() cannot extract the ending query's tasks).
- docs/super-sirius/concurrency-model.md: the teardown-fences section now
describes the steps 6+7 end-state (retired gate, surviving plan fence,
destructor-side leak attribution); the volatility note narrows to the
executor blocking primitives.
Tests: erase-while-a-sweep-holds (shared borrows survive a concurrent erase;
un-consumed batches accounted in the repository destructor with {operator,
port} attribution) and erase-during-an-in-flight-sweep on a live executor
(pre-step-6 the erase blocked on the gate; now it returns at once and the
sweep completes unharmed).
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
downgrade_executor::drain() (cancel EVERY queued request — failing healthy peers' waiters — and stop-join-restart the processing thread) loses its last legitimate callers and is deleted: - terminate() already stops executors outright via stop(); - per-query cleanup uses drain(query_id); - tests migrate to stop()/start() cycles, or to drain(make_query_id(0)) for the unattributed requests (the monitor's and external byte targets carry query id 0), which routes cancellation through fail_request() — the D6 re-arm path — without touching any thread. _lifecycle_mutex goes with it: it existed to serialize drain()'s stop-join-restart against stop(); stop() itself self-serializes through the _running CAS (one caller wins the true->false transition, a concurrent second returns immediately, exactly as before). The deterministic per-query-drain and monitor re-arm tests now wedge the processing LOOP (two candidates against a 1-thread pool: the worker parks in the predicate, the loop parks in reserve() for the second candidate), so queued requests provably stay queued under any request-dispatch model — this keeps them deterministic when request processing becomes concurrent (F8). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…t victim selection
The downgrade executor's serial one-request-at-a-time loop (candidate
collection, dispatch, then a trailing pool wait_all before the promise) is
replaced by per-request dispatch:
- Each popped request gets a shared request_context (its stats, target
spaces, and an outstanding-token count: one for the dispatch phase plus
one per conversion worker). Whoever drops the count to zero completes the
request — monitor re-arm, stats log, promise fulfilment, THEN clearing the
in-flight entry — so the processing thread moves on to the next request
while conversions are still running. Concurrency is bounded by the pool.
- In-flight publishing becomes a seq-keyed {request -> query} map:
drain(query_id) waits for exactly that query's entries (a per-query
in-flight COUNT, not a single slot), and wait_inflight_request() becomes a
BARRIER — it waits for the requests published before entry, not for
idleness, so a steady stream of monitor/peer requests cannot starve a
query's cleanup. Every worker destroys its convertible wrapper before
releasing its token, preserving the plan-lifetime fence's contract that a
completed request has no live wrappers.
- TIER-2 victim selection prefers tasks NOT belonging to the requesting
query (extracting the requester's own queued work to satisfy its request
is self-defeating): a peers-first pass with the requester excluded, then —
after waiting for that pass's conversions to land, since satisfied lags
the dispatches — an unfiltered last-resort pass. Budget per pass is the
queue size at pass start. Unattributed requests (query 0) own nothing and
take one unfiltered pass. convertible_gpu_pipeline_task_provider grows the
exclusion overload of get_next_convertible.
Preserved semantics, each pinned by the deterministic tests: D6 monitor
re-arm (complete_request clears the flag for consumed requests,
fail_request for destroyed ones), A7 per-query promise cancellation, and
the OOM retry-cap interplay (promises always resolve).
New tests: requests process concurrently (a second query's request completes
while another's is provably in flight; per-query drain and the in-flight
barrier wait for exactly the right requests); provider exclusion filter; a
request satisfied by a peer's task leaves the requester's own task resident,
and takes it only as the last resort.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…cleanup run_mandatory_cleanup now classifies failures per STEP instead of leaving every throw to the catch-site default: - SHARED-verdict steps: the downgrade executors' per-query drain + in-flight barrier, and the repository registry's erase. A failure there is not a this-query-only event (every co-tenant's cleanup drains the same executors and registry next), so these steps rethrow as the typed SiriusSharedCleanupStepFailure and the catch sites (finish() and the destructor backstop) latch the process-wide runtime_health via the new mark_shared_cleanup_step_failure() - loud, like the CUDA-corruption path. - Per-query steps (the query's task_creator reset, queue sweeps, scan reset, plan destruction) keep the existing default: classify_query_failure() probes for sticky CUDA corruption and otherwise contains the failure to the query (counted by per_query_cleanup_failures()). Testability: inject_cleanup_step_fault_for_testing() invokes a step-labeled hook at the head of each classified step group, and run_mandatory_cleanup_backstop_for_testing() drives the private backstop, so the dispatch is unit-tested on a bare context with no subsystem corruption (test_cleanup_step_classification.cpp): a downgrade_drain or repository_erase fault latches the runtime with no per-query count; a task_creator_reset fault leaves the runtime healthy and counts one contained failure; a clean cleanup classifies nothing. docs: the concurrency model's D5 section records the step-level tier. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…g-up triage Four new rows in the fixed table (B3/B4 shared ownership, the step-7 fence retirement with the surviving plan-lifetime fence called out, F8 concurrent downgrade requests, D5 step-level classification), and the stale steps-6+7-pending references updated in the MUST FIX items. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Enables and hardens concurrent GPU query execution end-to-end. This branch integrates @wmalpica's concurrency stack (#1364–#1366, #1369, plus the steps 3–12 enablement from the draft #1372 work), the late-materialization work (#1409, based on #1474), and five tiers of concurrency hardening driven by the issue register in
docs/concurrency/00-issue-register.md.Final validation: full suite 2562 cases / 2561 passed / 1 env-dependent skip / 32.9M assertions, on a clean build, plus an adversarial concurrency grid (spill storms, teardown races, SET storms, pin churn, prepared-statement races, fairness scenarios) run 5x each.
What this enables
scan_manager.max_concurrent_queries; concurrent transparent SQL queries genuinely overlap (asserted via aquery_lifecycle_peak()watermark, not just correctness).PREPARE/EXECUTEnow runs transparently (parameterless SELECTs).Hardening highlights (full record:
docs/concurrency/01-bringup-triage.md; architecture:docs/super-sirius/concurrency-model.md, new)runtime_unavailable_latch; per-query telemetry query-end; pin-name identity collisions rejected (was silent cross-database column splicing).gpu_reservation_max_retries(was hard-coded 100);gpu_execution(enable_optimizer=...)honored.Dependencies / merge order
🤖 Generated with Claude Code