Retrospective: M8a (the ESCHER CBST store)#
Status: done on the branch m8a-escher-store (from main at 8b9067e, which contains M7
and M6b; origin/main has not moved since, so there was nothing to merge), waiting for the
author’s review and merge. M8a is the first half of PLAN 11.3’s M8 row: the internal ESCHER store,
proven against both originals by CBST trace replay. M8b (the hypergraph container, its I/O and
Python, the header freeze) follows on the accepted store; both are in 0.3.0 with M9. Each step
adds its section; the milestone’s acceptance table and the open items close the record.
How M8a reached main (2026-10-14). The author merged pull request #14 with Squash and
merge: main’s M8a is the squash commit 2de6ebe (one parent, 30c3eef; its tree equals
b45e685, the head of m8a-escher-store). The M8a commits this record cites are not ancestors of
main through 2de6ebe; they live in refs/pull/14/head (= b45e685; git fetch origin pull/14/head brings them into a clone) and in the ancestry of m8b-hypergraph, whose merge commit
brings them back into main’s history. Squash merging is off since then (the author’s decision,
GOVERNANCE.md; ADR 0019’s amendment).
Planning change (not an author decision; reported to the author on 2026-10-07). M8 runs as
M8a then M8b, both in 0.3.0 (0.3.0 = M8a + M8b + M9; PLAN Appendix F unchanged in content), as
M1a/M1b, M2a/M2b and M6a/M6b did: the store port and the staged merge are internal and fully
checkable by trace replay, and the public hypergraph headers are frozen only after the
container, the line graph, the update path, the I/O and Python exist. docs/developer/plan.md
and docs/roadmap.md show the split; together M8a and M8b carry every item of PLAN 11.3’s M8
row (the old single row omitted the row views, the sorted_merge line graph, the hypergraph
generators, the CBST host model and the header freeze).
The author’s decisions for M8 (2026-10-07; GOVERNANCE.md approvals log). (1) The memory gate
is read as store accounting: dynG’s store device bytes against ESCHER-GPU’s own
DynamicHypergraph::deviceBytes() on coauth-DBLP after the build and after each of 3 x 50K
batches (seed 1); the Nsight Systems peak of live allocations of both processes is recorded,
not gated, in M8 (it is gated in M9). (2) O16: the hypergraph id_reuse default is
erase_first; the H-SOSP parity runs set insert_first.
Step 1: references, recorders, traces (refs-traces)#
What was built#
Commit |
What |
|---|---|
|
|
|
the CBST trace recorders as link-time twins ( |
|
the readers of trace format 1: |
|
|
|
smaller traces: integrity reports only on failure, |
|
the cross-check drivers of the two cores ( |
|
|
|
the fixes of the resume (below): output-option files are not inputs; the summary skips hash-only inputs; the comparator sees compactions |
|
the golden sets ( |
|
ADR 0036, |
(the certify commit) |
|
Deviations from the plan and the step’s text (each with the same intent)#
MOSP_ESCHER’s
bin/mainis not incbst_trace/mosp_escher_test. It draws its graph fromstd::random_device(mosp/src/generateGraph.cu:89-90), so two runs of the unpatched copy print different numbers and record different traces; no set that contains it can be exported twice with identical manifests. It is recorded in every export and listed in the set’sinfo.jsonunderexcludedby its invocation and the reason (not its op counts, which vary from run to run: see “The acceptance fix”); its CBST calls areDynamicGraph’s, which the seededstressTest,parallelStressTestandtest_dynamicgraph_roundtripmake too (all in the set). The twin output check compares it on its exit status and verdict lines.Section names. The sets are named
cbst_trace/<set>in the milestone text and live in$DYNG_SCRATCH/goldens/cbst_trace/<set>; theirgoldens.tomlsections are[sets.cbst_trace_<set>], because the goldens.toml writers split sections by a name without/. The writers now accept digits in set names (coauth50k).Digest modes. “Full mode for small runs” is applied per tree: in the default
automode a tree with more than 4,096 keys or row values gets delta digests (the changed and named keys in full, the 64-bit hash of every other key’s record). Delta digests are lossless (both readers rebuild every state and check every hash); the test suites’ 100K-hyperedge runs would otherwise need several GB per trace.Targeted sets.
cbst_trace/escher_gpu_targetedandcbst_trace/mosp_escher_targeted(parity/exporters/escher/cbst_targeted.cpp) drive each original’s unchangedCBSTOperationsAPI through the paths the suites miss or that are too large to commit (the coverage table ofparity/results/M8a.md).ADR number. The trace format and recorder ADR is 0036: the 0.2.0 release branch took 0035.
The cross-check of A’s line graph. A rebuilds its H2H (and V2V) from the device every batch (
rebuildFromDevice), which B’s store has no call for: A’s H2H and V2V traces are compared through B’s store up to their first rebuild (rebuild_only_in_a); the step’s minimum (H2V and V2H of A, and B’s h2h through A) is fully replayed. A’s coauth H2H inputs are hash-only (the 10 GB budget), so they could not be replayed anyway.The trace sets in the release certificate.
certify.pylists thecbst_tracesets (kindcbst_trace) but keeps a set without a replay record out of the golden verdict: the store that replays them does not exist before step 2. A replay record of one is checked as usual; the store’s replay writes them later (parity/results/M8a.md, section 8).Subset size. The committed subset is 307,170 bytes (“at most about 300 KB”); the repository’s test data stays at 470 KB, under 1 MB.
The interruption and the resume (2026-10-07)#
The step was interrupted by a usage limit after its first full chain of checks. Two exports of that
chain had failed and were fixed on resume, each by a recorder- or reader-side change, never by
relaxing a comparison: cbst_trace_mosp_escher_test differed between its two exports because the
recorder took the value of hsospBench --out (a CSV an earlier run had appended to) as an input
file; cbst_trace_escher_gpu_syn1m failed in the summary, which read hash-only inputs. Both sets
were exported twice again with identical manifests; the sets that passed were not re-exported
(no other invocation names an existing output file, so their traces are unchanged by the
recorder fix). The cross-checks that read the changed set, and A through A, were re-run with the
final comparator: 0 unexplained divergences.
Measurements#
The goldens take 6.75 GB of zstd data (limit 10 GB; largest: B’s DBLP hyperedge run, 2.75 GB);
the full chain (fresh builds and suites, census, twin checks, datasets, 8 double exports,
cross-checks) runs in about 80 minutes on GPU 1. See parity/results/M8a.md (the originals’ suites, the census, the twin check, the golden sets
and their sizes, the dataset pins, the cross-check, the coverage table).
Notes for step 2 (the host model)#
Read the traces through
dyng::test::cbst_trace::readerandstate_tracker(cpp/tests/parity/escher/cbst_trace.hpp): the tracker already rebuilds every state and checks the recorded hashes; the comparator compares a store’s state with the tracker’s.Large sets are zstd files: decompress to a temporary file (
zstd -dc --long=31) or pipe into the reader’sstd::istreamconstructor; the side file (trace.blob.zst) is only needed for construct and rebuild inputs.Hash-only arrays (
DYNG_CBST_TRACE_HASH_ONLY; listed per set ingoldens.toml) cannot be replayed: a tree whose rebuild or construct input is hash-only is checked by its hashes only.
Step 2: the CBST host model and the replay comparator (host-model)#
What was built#
Commit |
What |
|---|---|
|
|
|
|
(this commit) |
ADR 0037, |
Deviations from the step’s text (each with the same intent)#
The model keeps the flat payload, not only the per-key segments and rows the step lists: the originals’ observable rows depend on payload bytes a row-level model loses (dead rows keep their values, a reuse clears only the first segment, the unfill kernels differ by degree bin). The per-key segments and rows are derived from it (
records()); ADR 0037, section 1.Three field sets instead of “criterion 8 plus offsets when asked”.
semanticis exactly criterion 8,offsetsadds the offsets (for step 3), andexactcompares every recorded field; the model is held toexact, the strictest. Nothing is compared less.Unreplayable operations are reported, not compared. 3 ops of
escher_gpu_coauth50k(the H2H rebuilds, hash-only inputs) and the 54 ops of the batches-mode setescher_gpu_syn1mcannot be replayed by any store; the replay reports them with their reason and fails on any other. The acceptance criterion’s “100% agreement on every set” holds for every replayable op.The public-header rule.
dyng/testing/cbst_host_model.hppexposes CBST keys, offsets and the layout, the documented exception to criterion 9’s rule (which is about the hypergraph API).Model records are not release-certificate records. The replay records of this step are
parity/results/M8a-host-model-<set>-<fields>.json;certify.pykeeps waiting for the store’sparity-cbst_trace_<set>-*.jsonrecords (step 3 on), so the host model cannot stand in for the store in the 0.3.0 certificate.dyng_default()is provisional. It composes B’s compaction triggers with A’s growth where B throws; no trace exists for it. Step 4 aligns it with the store’s dyng preset.
Measurements#
The Release-built replay of all 8 sets through the model takes 4 min 45 s with 8 jobs, the manifest checks included (the largest, B’s DBLP hyperedge run: 2 min, about 63 million key states compared). The layout check of the 17 large sizes takes 20 s in the Debug build.
Notes for step 3 (the verbatim port)#
Drive the store through
store_adapter(cpp/tests/parity/escher/cbst_compare.hpp): implementcreate_tree,destroy_tree,run(the seven ops;InsertMappinginop_outcome::mapping, degree-bin launches inbinswithhas_bins),scalarsanddigest(read back the device state: keys in BST order,subtree_availby heap position, per key the record and the row walked as the recorder walks it), thenreplay(reader, adapter, {field_set::offsets or exact}). The model’s adapter (host_model_adapter) is the reference implementation of the interface.parity/compare.py cbst_trace --store deviceis reserved (it exits 2 now): point it at a device replay tool, and write the certificate recordsparity-cbst_trace_<set>-*.jsonthere.Under
escher_gpu, A’s sets must passexact(offsets, occupancy,deviceBytes(), bin launches are all in it); the scenarios of step 4 can mirror every op oncbst_host_modeland comparerecords()with the store’s digest.The 3 coauth H2H rebuilds stay unreplayable until M8b computes the line-graph merge: their inputs’ recorded hashes can then check dynG’s merge.
Step 3: the verbatim port of ESCHER-GPU’s store (verbatim-port)#
What was built#
Commit |
Class |
What |
|---|---|---|
|
c1 transplant |
ESCHER-GPU@abcf9b2’s store in |
|
c2 errors |
|
|
c3 memory |
|
|
c4 streams |
the resources’ stream, |
|
c5 registration |
|
|
harness |
|
|
dynG additions |
the empty tree; clone; download; |
|
tests |
the cbst_test layout sizes, kernel-path coverage, errors, two streams, host syncs per op, the source rules |
|
policy |
|
|
test |
the early memory reading on coauth-DBLP, |
|
harness, CI |
|
(the docs commit) |
docs |
ADR 0038, the store’s README, |
Deviations from the step’s text (each with the same intent)#
c1 carries the replay harness. The step asks for the replay after every commit, so the device adapter, its tests and
--store deviceof the replay tool are in c1 next to the transplant; the store files of c1 are the pure transplant (transplant.py --check).The exception of a construct below its payload is
invalid_argument_error(the caller’s capacity contradicts its own payload;capacity_erroris reserved for the offset limit and full payloads). A kept a tree without nodes there, so the replay reports it as the documented divergenceconstruct_below_payload(one op of the golden sets, A’s targeted scenariosmall_capacity), by a precise rule and not as a tolerated mismatch.The device error word has no user yet. No kernel of A detects a contract violation; B’s bounded chain walker (step 4) is its first user.
The comparator learned three things (test code, used by every later step): the store’s own integrity verdict is compared with the trace’s
check.rows(after the state, so that a first divergence names a state field), documented divergences are reported apart from mismatches, andcompare.pyjudges MOSP_ESCHER’s sets underescher_gpu()by their first divergences per tree against the host model under the same preset (the step’s “up to their first divergence, listed with its cause”: the causes are the behaviour classes of ADR 0037).store_policyis its own commit after the additions, not part of the transplant: c1-c5 keep A’s growth inline as A had it, and the policy struct then takes it over with the same numbers (the replays are unchanged).Thrust’s temporary storage comes from a Thrust allocator over
memory_resource_ref(util/thrust_allocator.cuh): Thrust takes an allocator, not a CCCL memory resource, so the adapters ofcccl_memory_resource.cuhare not needed for it.The layout test calls two host launchers of the library (
launch_build_layout,launch_extract_keys_from_positions): nvcc keeps the host stubs of a library’s kernels hidden, so a test cannot launch them directly as A’scbst_testdid.find.cu. A’s debug lookups printed from the device;find()returns the heap position of every key, andfindContents(which printed rows) is gone (rows are read withdownload()).Untested at full scale:
flatten_rows’ 2^31 check and the growth beyondINT_MAXneed more than 8 GB of host or device memory; the policy’s numbers (grown_capacityat and beyondINT_MAX) are unit-tested on the host,flatten_rows’ check is code-reviewed only.The memory reading replays the coauth H2H rebuilds by their sizes. Their inputs are hash-only (A’s line-graph merge); the accounting depends on the row count, the used payload and the capacity only, which are recorded in full.
The merge-step script builds a reduced library (
DYNG_ALGORITHMS=sssp, minimal index types, no tools or tutorials, Release): only the replay targets are needed, and every commit builds in a few minutes.ADR number. 0038 (0035 is the release branch’s, 0036-0037 this milestone’s).
Names. Every camelCase identifier of the ported files became its snake_case spelling (locals included), the kernels
<verb>_<object>_kernel; all listed inrename_map.toml. The store’s stages arehypergraph.cbst.<op>, nested under the container’shypergraph.apply.*of the timed-region map.
Measurements#
The replay of the small sets (escher_gpu_test, escher_gpu_targeted, mosp_escher_test, mosp_escher_targeted) on the device takes about 4 minutes with a Release tool and 6 jobs on one GPU; a commit of the merge-step table takes about 8 minutes (the reduced build included). The large sets take about 25 minutes (coauth-DBLP’s 33 million key states, B’s DBLP runs with 63 and 48 million).
The store accounting on coauth-DBLP equals A’s
deviceBytes()at every op (ratio 1.0000); the nsys peaks: A’s whole program 3,018.9 MiB, the store replay 2,772.6 MiB (parity/results/M8a.md, section 11).Sanitizers: racecheck over the store tests takes 85 s (the 2^24-key layout), the other tools a few seconds each.
ci/gpu_local.sh(dev-cuda, GPU 1, at7e397a2): all 14 steps passed, the new stepcbstincluded (about 2 hours; memcheck over every GPU executable is most of it).
Notes for step 4 (MOSP_ESCHER’s behaviours)#
The store takes a
store_policy; addmosp_escher()and grow the struct (B’s fixed capacityflat * 1.3 + extra, compaction triggers, construct validation, the fill clamps, scratch growth, caller counts) one behaviour per commit;parity/cbst_merge_steps.shreplays any list of commits.The device adapter accepts
escher_gpuonly (device_store_adapter);compare.py cbst_trace --store device --policyrejects the others. Map B’sescher::EscherErrorto the dynG exception of each throw site in the comparator when the B preset arrives (the exception type is compared).MOSP_ESCHER’s traces now diverge under
escher_gpu()at three classes only (caller counts, scratch growth, compaction / capacity / validation): each commit of step 4 should move the first divergence of exactly the trees of its class, whichfirst_divergencesper tree shows.cbst_store.hppstill includes<cuda_runtime.h>(as A’sstructure.hpp); the container of M8b, compiled by the host compiler in CPU builds, needs a host-only facade or a forward-declared context.
Step 4: MOSP_ESCHER’s behaviours, part 1 (merge-invariant)#
What was built#
Ten behaviours of MOSP_ESCHER@4b86159 merged into the store, one per commit (ADR 0039): (a) construct validation; (b) the fill clamps; (c) caller row counts as a fast path of construct; (d) one bounded chain walk,
chain.cuh, for every reader and the unfill writer; (e) B’s integrity checks and dynG’s index checks; (f)rows()(B’sfind.cu) and B’sdelete_avail.cucomment; (g) the scratch growth rule; (h) an explicitcompact()(compact.cu); (i) the storage counters (stats(), B’s accounting scope) without B’s process-global mempool setting; (j) a rule for appends and removals on dead or absent keys. (k), the disposition of every other differing file, is ADR 0039 section 6 (atomicAdd_sllverified unused and not ported).The store’s control block: PLAN 4.7.3’s device error word, its counters and a pinned mirror read back with every read point;
check_errors().The harness: the staging preset
mosp_escher_staged()and its host-model twin (staged_host_policy()), so that B’s traces are judged after every commit; the merge-step rows list the first divergences by field; every construct of every replay built both ways (twin constructs); the replay compares three integrity checks;decode_opshared by the replay and the new compact-after-every-op lockstep (dyng_cbst_compact_tests,dyng_cbst_compact_large_tests).23 new tests (17 in
cbst_merge_test.cpp; the twin constructs, the lockstep (2), the validation, the staging preset and the scratch rules), the merge-step table (parity/results/M8a.md, section 12), ADR 0039, the store README and ADR 0038 updated.
Deviations from the step’s text (each with the same intent)#
A staging preset. The step asks for B’s sets “under the most complete B preset of that commit”; there was no B preset before step 5, so
store_policy::mosp_escher_staged()(B’s behaviours merged so far, B’s fixed capacity, no automatic compaction) is that preset, and the comparator’sstaged_host_policy()gives the host model the same behaviours. Each commit then shows B’s divergences moving: 50 + 7 trees before, 18 + 4 after (only automatic compaction and the constructs without row counts remain). Step 5 replaces the preset bymosp_escher().Harness commits in the table.
2bc0e24(the staging preset),1e01ad8(the rows by field) and4402fcf(the provenance identifiers) change no store behaviour; they have rows of their own.The chain bound is the used payload, not B’s 1,024 hops (as the step asks; ADR 0039 section 3): the longest chain of any trace is 98 hops, so no trace changes, and the stop condition (B’s bound observable in its own traces) does not apply.
The clamps go further than B’s literal code: dynG clamps the occupancy (the write position and an overflow’s chain pointer stay in the segment), B only the room; the warp and block unfill kernels clamp it too. Only inconsistent metadata is affected; no trace state has it.
The device error word is a control block of 4 words outside ESCHER-GPU’s accounting scope:
device_bytes()must stay A’sdeviceBytes()(criterion 10(a)) andallocated_bytes()keeps its scope; the accounting test counts the block separately.Dead or absent keys: the rule chosen of the two the step offers is “count them and surface them as internal_error at the next existing sync point”; the items are also skipped (nothing is written). An empty tree’s append and remove_values now throw (ADR 0038 had them change nothing); the test was changed accordingly.
Caller counts are also range-checked on the host (dynG; B trusted them), and the Debug check against the scan costs one synchronization in Debug builds only. A’s traces have no counts: their twin constructs use counts computed on the host.
compact()grows or throwscapacity_errorwhere B threwEscherError, before anything changed (not poisoning); a tracedcompact()is unsupported underescher_gpu()(A had none), except in the lockstep, which enables it.The host model’s
compact()already existed (step 2 wrote it from B’scompactCBST, independently of the store); the step uses it as the oracle of the lockstep instead of adding a second one.Validation: a capacity below the payload is now rejected for a tree of 0 records too (B’s rule; step 3 allowed it), and a record count above
INT_MAX / 3is acapacity_error(dynG: the scratch’s int32 indexing).The integrity checks in the replay:
check.subtree_availis compared with B’s recorded verdict and with 0 for A’s states (A had no such check);check.index(dynG’s) with 0.Per-commit checks: before each commit the CUDA build,
ctest -L cbst -LE slow, memcheck on the touched tests (and racecheck after (d) and (h)) and the fastci/check.shsteps (format, reuse, provenance, pre-commit); the replays right after each commit (a temporary worktree needs the commit); the fullci/check.sh, the portability checks and the large sets at the end of the step.The 11
test_cbst_smokeconstructs without row counts stay listed divergences: B’snullptrpath (occupancy 0) is never merged; step 5 must document them for criterion 10(b).ADR number 0039. The scratch rule of dynG’s default preset (step 5) is chosen here: geometric.
Measurements#
A commit of the merge-step table takes about 8-10 minutes (the reduced Release build and the small sets under both presets; GPU 1, shared with the 0.2.0 release’s test runs).
The compact-after-every-op lockstep over the four small sets takes 211 s (Release, GPU 1): 15,385 ops, 15,383 compactions.
Racecheck over
dyng_cbst_store_teststakes about 170 s (the 2^24-key layout and the new 2,000-segment row).
Notes for step 5 (automatic compaction and the presets)#
Replace
mosp_escher_staged()bymosp_escher()(store_policy::compaction_rule: B’s trigger points infillImplandinsertCBST, ADR 0039 section 2 (h) gives the kernels) and its host-model twin by the model’smosp_escher();compare.py’sagainst_modelthen stops judging B’s sets against the model.Map
escher::EscherErrorto the dynG exception of each throw site in the comparator (construct validation:invalid_argument_error; fill, insert, compact capacity:capacity_error).Document the 11
rowOccupancy = nullptrconstructs oftest_cbst_smoke(B’s legacy occupancy 0, never merged) as a divergence class, or criterion 10(b) cannot hold for them.dynG’s default: compact first, then grow (A’s rule), the geometric scratch,
capacity_erroronly at the int32 offsets or on OOM;cbst_host_policy::dyng_default()must be aligned with it.The mutations of criterion 17 that touch this part (compaction padding, surplus padding, the tie order, requested keys) can use
dyng_cbst_compact_testsand the replays: a changed compaction padding fails the lockstep.
Step 5: automatic compaction, the presets, scenarios, mutations (merge-policies)#
What was built#
Commit |
What |
|---|---|
|
B’s automatic compaction at its two trigger points ( |
|
|
|
an injected allocation failure at every allocation of an insert and a fill under every preset (partial failure, poisoning) |
|
B’s 9 |
|
the graph-on-CBST test (PLAN 6.1 D8) and the replay of B’s |
|
the 7 CBST mutation hooks, the CTests |
|
|
|
|
|
the compact-after-every-op lockstep stops a tree at a verified dead-key item (a step-4 failure found here) |
(the docs commit) |
ADR 0040, the store’s README, |
Deviations from the step’s text (each with the same intent)#
construct_without_row_countsis a field-level exemption, not a tree-level one. The resumed work had made the whole tree stop at such a construct; that would have left the 22 later ops of the 11test_cbst_smoketrees uncompared. The rule now exempts only the occupancy of the keys B built withrowOccupancy = nullptr(while the trace still has 0 and the store its scanned count) and B’scheck.rowswhen it is exactly one metadata violation per live such key (B’s own checker,integrity.cu:160-186), and keeps comparing everything else (ADR 0040, section 3). The difference is mandated by criterion 11 (“never B’snullptrpath”) and ADR 0039 (c). With it, B’smosp_escher_testreplays all 14,104 ops with 0 mismatches; 31 ops carry the documented exemption.B’s sets under
escher_gpu()diverge in 39 trees, not 50: the same field-level rule applies to the host model underescher_gpu()(which also scans), so the smoke trees no longer count as diverging at their constructs on either side; the judgement (store = model, tree by tree) is unchanged.The device store’s default preset is the trace’s original’s (
escher_gpufor A,mosp_escherfor B): “the original’s preset” in the records now means that.Criterion 10(c) needed a definition, the field set
policy(ADR 0040, section 3): after a policy difference the rows of deleted keys are the policy’s (a compaction empties them, A keeps their stale values), and an InsertMapping that differs after a length changed is documented and ends the tree, because its later rows sit under other keys. That happens in 5 trees, all H2V trees of A’s medium graph (043-main,051-type1,059-type2,067-type3,075-coarseTriangle: the insert at op 4, after the default compacted at op 2 where A grew). A mapping difference before a policy difference is still a mismatch; there is none.max_payloadis a field ofstore_policyand of the public test policycbst_host_policy(the API snapshot is updated); it is INT_MAX for every preset.Poisoning is conservative: the store does not tell a failure before its first write from a later one, so every failure of an updating operation poisons it; only the precondition checks, the construct validation and
compact()’scapacity_errorare strong (ADR 0040, section 4).The mutations are hooks in the sources (
#if defined(DYNG_MUTATION_CBST_...)), as cycle_count’s are, plus one in the graph test (the graph-on-CBST is a test, PLAN 6.1 D8); the CTests rebuild a copy of the hypergraph module per mutation (8 copies in every CUDA build withDYNG_CBST_MUTATION_TESTS, default on);mutate.py’s suitecbstruns GoogleTest executables instead ofcompare.py.A mutation found a blind spot: with the absent value of every removal list placed last, the random lockstep did not see “unfill ignores the last value” (only the replays and the graph test did); the absent value is now placed at random.
Reduced sizes under the sanitizers (
DYNG_CBST_REDUCED=1, ADR 0040 section 5): the full scenarios (scale up to 2^20, random with 70,000 rows) would take hours under racecheck; the full sizes run in everyctest.The early-warning table splits the timelines at
cudaDeviceSynchronize(A’s ScopedTimer stages bracket every store operation, and the replay synchronizes after every op), classifies each segment by its kernels, and counts blocking and allocating runtime calls per segment; it is reported, not gated.A step-4 test was failing:
dyng_cbst_compact_large_tests(slow) had failed sincebe449b6(the dead-key rule), which step 4 did not re-run it after (its retrospective says the lockstep passed: true atc05d7cc). The lockstep now ends a tree at a fill or unfill that names a key the model’s state holds deleted or absent (f03b3eb); the four small sets: 14,007 ops, 13,964 compactions, 0 problems, 41 trees stopped at such an item (15,385 ops before the rule existed).The ADR is new (0040) rather than an extension of 0039’s text, which keeps ADR 0039 as the record of part 1.
No pull request is opened by this step: the orchestrator opens it and the author merges (the session rules forbid pushing); item 9 of the step is left to the orchestrator.
Per-commit checks as in step 4 (deviation 39): the CUDA build,
ctest -L cbst -LE slow, the fastci/check.shsteps before each commit, the replays (cbst_merge_steps.sh) right after it; the fullci/check.sh, the portability checks, the sanitizers, the large sets,ci/gpu_local.shand the fresh-clone builds at the end. After the documents were committed, the whole merge-step table (34 commits,477d592toe7157af) was replayed once more with the final comparator, and the large sets at the head’s code (f03b3eb): every rowok(parity/results/M8a.md, section 13, “The checks of step 5”).
Measurements#
Replays (
parity/results/M8a.md, section 13): every commit of the stepok; A’s sets 1,17652 ops, 0 mismatches; B’s sets under
mosp_escher()14,104 + 53 ops, 0 mismatches, 51 compactions and 3 exceptions at B’s ops; the large sets at00319daand at the final commit: coauth50k 15 ops, 0 mismatches; DBLP hyperedge 27 and vertex 21 ops, 0 mismatches undermosp_escher().
Memory (criterion 15; GPU 0, exclusive lock,
043a6dd): every tree’sdevice_bytes()after its construct and after every op of the 3 batches equals ESCHER-GPU’sdeviceBytes()(18 readings, worst ratio 1.0000) underescher_gpu()and underdyng_default()(the same values at every point: no compaction happens on coauth50k, and the scratch never grows past its size after the construct, so the geometric rule and A’s exact one hold the same bytes); the nsys peak of live device allocations is 3,018.9 MiB for A’s whole program and 2,772.6 MiB for the store replay under either preset (recorded, not gated).Early warning (same run): device time per op h2v insert 3.18 / 3.07 ms (A / dynG, 0.96x), h2v erase 0.10 / 0.10 ms, v2h unfill 0.71 / 0.79 ms (1.11x), v2h fill 0.105 / 0.100 ms; blocking runtime calls per op insert 17 / 7, fill 2 / 1, erase and unfill 0 / 0; allocation calls per op insert 25 / 0, erase 2 / 0, unfill 1 / 0, fill 3 / 0 (dynG allocates from the stream-ordered pool); registers per thread equal for every hot kernel but
unfill_thread(23 / 24); the warp and block bins do not run on coauth50k. (Corrected by the review fixes: one run at default clocks, and the dynG allocations were not counted; see “The review fixes”.)Mutations: all 7 killed (4-13 failed tests each); a mutant suite runs in 2-7 s.
Sanitizers (reduced sizes, GPU 1): memcheck, racecheck, synccheck and initcheck clean on
dyng_cbst_scenario_tests,_policy_tests,_graph_tests,_store_testsand_device_tests(racecheck 32-236 s per executable).Effort: step 5 took about 4 hours of wall-clock time after the first resume and about 2 more for the final checks after the second (the whole merge-step table takes 77 minutes, the large sets 24; the lock and GPU 1 shared with the release branch’s runs); M8a as a whole about 2 days for its five steps against the plan’s 1-1.5 weeks for M8 + M9 (
docs/developer/plan.mdsection 4 re-baselines 0.3.0: 1.5-2 weeks of focused effort, 3-4 weeks calendar).
The M8b handoff#
The store’s contract (
cpp/src/hypergraph/escher/README.md): 1-based keys and+1values;insert_mapping::item_to_keyis the only source of a new row’s key; requested keys are ignored; a failed update poisons the store (ADR 0040, section 4);check_errors()is the read point of the end of an update.Sizing: the container calls
store_policy::initial_capacity()of its preset (in-place trees with slack, merged trees exact;mosp_escher()flat * 1.3 + extra).cbst_store.hppincludes<cuda_runtime.h>: a host-compiled container needs a host-only facade (open since step 3).The line graph: A’s H2H rebuild inputs are hash-only in the golden set coauth50k; M8b’s
sorted_mergecan be checked against the recorded hashes and the store accounting.Watch: v2h unfill at 1.14x of A’s device time per op at locked clocks (1.11x in the first single run; the bounded walk and the dead-key lookup, visible in the unfill kernels’ SASS); M9 gates the
updateregion and the nsys peak.The store’s contracts after the review: a poisoned store refuses
rows(),find()andcheck_errors(); batches are validated (validate_batch);dyng_default()grows after a compaction that leaves less than 1/8 free (size the container’s trees knowing that a long-churn tree under the default ends up with at least 1/8 of slack).
The review fixes (2026-10-08)#
Independent reviewers confirmed 26 problems in the milestone as submitted at 2ee8bcc; all are
fixed on the branch, each in its own small commit (aa26772 .. the commit of this record), with
the affected tests, replays and measurements re-run after it (parity/results/M8a.md, section 14).
What was fixed#
Store correctness (
cpp/src/hypergraph/escher, ADR 0038 “Later changes”, ADR 0040):a construct of 0 records after an earlier tree kept the old arrays (stale
find()positions, appends into stale rows);rebuild_from_devicewith 0 records launched 0-block grids and poisoned the store;row_locator::build_sorted_copy(max_key 0)launched zero grids (aa26772);no operation checked its batch’s shape (an out-of-bounds device read that lost the CUDA context); the construct accepted zero-length, overlapping or unordered rows and unsorted keys (a fill then overwrote the next row); the fill’s overflow plan was summed in int32 and could wrap (
9175297);rows(),find()andcheck_errors()succeeded on a poisoned store;clone()dropped a pending device error and failed onto another device; the Debug check of the caller’s row counts synchronized uncounted (e4ceaff);dyng_default()had no hysteresis and compacted the whole payload more and more often before it grew (3961dc0);test gaps:
row_locator::build,flatten_rows’ limit, two streams with compaction and growth, a launch on a literal default stream (0feda27).
Parity honesty (
parity/,cpp/tests/parity/escher):the criterion-10(c) replay labelled any layout difference a policy difference (a planted compaction-padding defect passed): policy points now need differing events, the store’s layout under every preset is checked in lockstep with the host model, and the documented divergences are pinned per set and rule (
8214aab);B’s sets judged against the model compared only the first field name of each tree’s first divergence: now the whole set of (field, key, expected, actual), with the unreplayable check (
301dffc);the release certificate left the CBST sets out of its verdict and no step wrote their records: now required under the original’s preset and
dyng_default(16a7251,6bb7782);B’s
bin/mainwas never replayed: recorded and replayed once byci/gpu_local.sh(cbst_trace_goldens.py once main;16a7251).
Provenance and records: the provenance check’s CBST symbols derived from the rename map, and
cbst_merge_test.cpp’s header (9e2a3db); two wrong citations (5440447); the “every recorded field” claims qualified (d8a2a25); section 12’s lockstep record dated (7668ed5); personal paths out of the replay records (3439022and the regenerated records);cbst_merge_steps.shno longer replaces the records it is cited by (6865a42).Measurements: the early-warning table re-measured with PLAN 8.6’s protocol (20 alternating runs per side at locked clocks, medians and spread; registers, occupancy and SASS of every hot kernel from cuobjdump; every allocation call counted,
cudaMallocFromPoolAsyncincluded) (f12e58dand the record).
The findings and their commits#
# |
finding (severity) |
commit |
|---|---|---|
1 |
the criterion-10(c) replay labelled layout differences as policy differences (high) |
|
2 |
the release certificate never gated the CBST sets (medium) |
|
3 |
B’s sets judged against the model on one field name per tree (medium) |
|
4 |
B’s |
|
5 |
construct of 0 records kept the old node arrays (high) |
|
6 |
|
|
7 |
|
|
8 |
no validation of an operation’s batch shape (medium) |
|
9 |
construct accepted zero-length or unordered rows (medium) |
|
10 |
|
|
11 |
reads of a poisoned store succeeded (medium) |
|
12 |
|
|
13 |
the Debug check of the row counts synchronized uncounted (low) |
|
14 |
the fill’s overflow prefix sum was int32 (low) |
|
15 |
test gaps: |
|
16 |
provenance SYMBOLS missed about 20 port identifiers (medium) |
|
17 |
|
|
18 |
the “every recorded field” claims named no exemptions (low) |
|
19 |
section 12 dated no commit for its lockstep pass (low) |
|
20 |
ADR 0040 cited B’s fill throw at the wrong lines (low) |
|
21 |
the README credited A’s 64-bit rank to the port (low) |
|
22 |
personal paths in 32 replay records (low) |
|
23 |
|
|
24 |
the early-warning table did not follow PLAN 8.6’s protocol (medium) |
|
25 |
the early-warning table counted no dynG allocation (medium) |
|
26 |
|
|
Deviations (each with the same intent)#
The hysteresis is a new rule of the default preset, not of PLAN 6.4.5’s literal text (“compact first, then grow”): after an automatic compaction the default also grows when fewer than 1/8 of the slots are free (ADR 0040 section 1). It keeps the plan’s order and its failure rule, changes no original’s preset, and is mirrored by the host model, so the criterion-10(c) replays document the extra growths (B’s
mosp_escher_test: 5 documented divergences instead of 1, pinned).The policy lockstep’s model of
mosp_escher()scans where B’s construct had no row counts (the store never takes B’snullptrpath, ADR 0039 (c)); the replay against B’s trace still documents that divergence (construct_without_row_counts). Trees whose next op the model cannot predict (an op the preset’s original lacks, e.g. A’s rebuilds undermosp_escher()) stop as counted agreed stops.The comparison with the model after a trace divergence (the review’s “better” fix) is the policy lockstep at the head; the merge-step replays of earlier commits compare the whole first divergence only, because the lockstep test does not exist in those commits’ trees.
The fill’s int32 overflow path is not exercised by a test: it needs batches of several GB. It is reached only when the bump pointer plus the item sizes plus 3 per item passes INT_MAX (code inspection; the traces never come near), and then it reads the plan back exactly before deciding.
flatten_rows’ limit is tested through its new parameter instead.download()stays readable on a poisoned store (the forensic exception, documented): the tests use it to show what a failed operation wrote.One intermediate commit failed two Python tests:
16a7251made the certificate require the CBST records, and the full-certificate fixture got them in6bb7782.Commits were staged from the tested final tree, hunk by hunk; every store commit was then built and its CBST tests run in a temporary worktree (section 14).
The whole-first-divergence comparison reads a thrown op’s status by its exception type (
32cf67d). Its first full re-run of the merge-step table failed every commit from2bc0e24on, in B’s sets under the staging preset; in every failing tree the store and the host model threwcapacity_errorat the same op with every compared value equal, and the sets differed only in the two stores’ wording of the message and in the InsertMapping of a thrown insert, which neither returns. The table’s divergences were described correctly. The comparison now takes the status of a thrown op as “exception” (criterion 3’s recorded field) and a thrown op’s mapping as none; a different type or any other value still fails ( a7fd8e8). The re-run then failed one more row,c981416, at the targetedcompact()that the store merges one commit later: the store reports it unsupported, the model’s staging policy “has no compact()”;4b1c519reads exactly these two wordings as one status (the model’s other “undefined” statuses stay distinct;9acb3b5). Nothing else of the comparison changed.3961dc0changed the testing header without the API baseline (compaction_headroom):ci/check.sh’s docs step failed on the branch untilf0eb0caregeneratedpublic_api.txt(ADR 0023; the CHANGELOG already listed the field).
Measurements#
The planted defect (the review’s compaction padding): the default-preset replay now fails (
mosp_escher_test10,519 mismatches, first at op 197;escher_gpu_targetedoff its pinned count; the policy lockstep fails), where it passed with 2,401 “documented” divergences before.Replays at the head: every set under its original’s preset and
dyng_default()with 0 mismatches and exactly the pinned documented divergences; B’s sets underescher_gpu()agree with the host model in the whole first divergence of every tree; the policy lockstep over the small sets 14,319 / 15,237 / 15,385 ops (escher_gpu()/mosp_escher()/dyng_default()), 0 problems; B’sbin/mainrecorded once, 111 ops replayed, 0 mismatches.The merge-step table replayed again with the whole-first-divergence comparison: every row
okafter32cf67dand4b1c519(the first run’s failures were the two stores’ wording of one exception or missing operation, deviation 63); at the head every set, the large ones included, with 0 mismatches and the pinned documented divergences (results section 14).Memory (criterion 15): unchanged, 1.0000 of ESCHER-GPU’s accounting at all 18 readings under both presets.
Early warning (20 alternating runs per side, locked clocks): v2h unfill 1.14x of A (0.847 / 0.744 ms, the spread under 2%), h2v insert 0.99x, erase 1.00x, fill 0.94x; the unfill kernels’ SASS grew (thread 192 -> 376, warp 352 -> 552, block 568 -> 752 instructions; registers 23 -> 24, 23 -> 25, 36 -> 36), all at full theoretical occupancy; dynG makes as many allocation calls per op as A, all stream-ordered (step 5’s “0” was a counting error).
Effort: about 9 hours over two sessions (a usage limit interrupted the first), most of it waiting for the perf lock and replaying the merge steps between the release branch’s certificate runs.
Open items#
The hosted CI has not run on this branch (nothing is pushed); the orchestrator opens the PR.
The acceptance fix (2026-10-13)#
The independent acceptance verifier passed 17 of the 18 criteria at 953fdbf and failed
criterion 4: the manifest of cbst_trace_mosp_escher_test was not reproducible. The set’s
info.json listed the op counts of the excluded bin/main (deviation 1), and those counts follow
the graph that std::random_device draws (108 or 111 ops), so a re-export matched the pin only
when bin/main happened to draw a graph of the same size.
Fix (
c210ed3): theexcludedrecord holds the invocation and the reason only, tested with two runs of different op counts. The set was exported again twice and re-pinned (5e8555b, manifested6400698e3d...; all 24 case digests unchanged). Its replay records were re-run against the new manifest with unchanged results (96193bf).Evidence: three more exports from a third fresh scratch gave the same manifest while
bin/mainrecorded 108, 111 and 111 ops (parity/results/M8a.md, section 15).Lesson: anything measured from an invocation that is excluded because its inputs vary must stay out of a hashed manifest. The
--twicecheck passed by chance (both exports drew graphs of the same size), so it could not catch this; a test of the record now does.Deviations: none; the plan’s “exported twice, byte-identical manifests” now holds for every run.
The acceptance criteria#
# |
criterion |
where it is met |
|---|---|---|
1 |
base and references |
step 1; |
2 |
additive recorders |
step 1; ADR 0036; the census and twin checks |
3 |
trace format |
ADR 0036; |
4 |
golden sets |
step 1; |
5 |
dataset pins |
step 1; |
6 |
cross-check of the originals |
step 1; section 6 |
7 |
timed regions |
|
8 |
host model, the go/no-go |
step 2; ADR 0037; section 10 |
9 |
verbatim port |
step 3; ADR 0038; section 11 |
10 |
replay at every merge-step commit |
sections 11-13 ( |
11 |
one B behaviour per commit |
steps 4-5; ADR 0039 section 2, ADR 0040 |
12 |
presets and error policy |
ADR 0040 sections 1, 2, 4; |
13 |
scenarios and unit tests |
|
14 |
kernel-path coverage |
|
15 |
early store memory |
|
16 |
GPU robustness |
the sanitizers (sections 11-13), two streams, host syncs, double replay |
17 |
mutations |
|
18 |
CI and records |
|