Skip to content

feat(ilp): GA hardening gates and documentation alignment - #2

Merged
levi770 merged 14 commits into
mainfrom
feat/dilp-ga-hardening-docs-and-gates
Mar 4, 2026
Merged

feat(ilp): GA hardening gates and documentation alignment#2
levi770 merged 14 commits into
mainfrom
feat/dilp-ga-hardening-docs-and-gates

Conversation

@levi770

@levi770 levi770 commented Mar 4, 2026

Copy link
Copy Markdown
Contributor

Summary

This PR packages the dILP GA-hardening workstream into three logical commits and aligns top-level docs with current behavior:

  1. feat(ilp): harden GA trainer gates and reproducibility controls
  2. test(ilp): add GA reliability/performance gates and hardening coverage
  3. docs(ilp): align README, changelog, roadmap, and architecture with GA hardening

Core changes

  • Deterministic trainer path (TrainConfig.deterministic) with reproducible attempt seeding
  • selected_hard persistence and deterministic ordering in artifact flow
  • Holdout strategy execution (LOO for <=20 positives, k-fold for >20)
  • Holdout threshold enforcement in promotion (holdout_threshold)
  • Typed schema promotion gate + waiver/manual-review fallback (typed_schema_required, waiver_untyped)
  • Host transfer telemetry exposure via PyO3 (host_transfer_stats, reset_host_transfer_stats)
  • Forward timing telemetry (forward_p95_us) propagated through trainer artifacts
  • Added GA reliability/performance suites:
    • python/tests/test_ilp_ga_reliability.py
    • python/tests/test_ilp_performance.py

Documentation updates

Updated to match current code behavior and gate semantics:

  • README.md
  • CHANGELOG.md
  • docs/ROADMAP.md
  • docs/architecture/dilp-training.md

Verification

Fresh CUDA-enabled verification run:

LD_LIBRARY_PATH=/usr/local/cuda/lib64:/usr/lib/wsl/lib:$LD_LIBRARY_PATH \
PYTHONPATH=crates/pyxlog/python \
.venv/bin/python -m pytest python/tests/test_ilp_types.py python/tests/test_ilp_trainer.py python/tests/test_ilp_promoter.py python/tests/test_ilp_holdout.py python/tests/test_ilp_robustness.py -q
# 44 passed in 276.53s

LD_LIBRARY_PATH=/usr/local/cuda/lib64:/usr/lib/wsl/lib:$LD_LIBRARY_PATH \
PYTHONPATH=crates/pyxlog/python \
.venv/bin/python -m pytest python/tests/test_ilp_d2h_gate.py python/tests/test_ilp_sparse.py -q
# 19 passed in 20.24s

LD_LIBRARY_PATH=/usr/local/cuda/lib64:/usr/lib/wsl/lib:$LD_LIBRARY_PATH \
PYTHONPATH=crates/pyxlog/python \
.venv/bin/python -m pytest python/tests/test_ilp_beta_gate.py -q --timeout=1200
# 20 passed in 754.62s

LD_LIBRARY_PATH=/usr/local/cuda/lib64:/usr/lib/wsl/lib:$LD_LIBRARY_PATH \
PYTHONPATH=crates/pyxlog/python \
.venv/bin/python -m pytest python/tests/test_ilp_reliability.py -q --timeout=1800
# 20 passed in 784.62s

LD_LIBRARY_PATH=/usr/local/cuda/lib64:/usr/lib/wsl/lib:$LD_LIBRARY_PATH \
PYTHONPATH=crates/pyxlog/python \
.venv/bin/python -m pytest python/tests/test_ilp_performance.py -q --timeout=1800
# 3 passed in 14.68s

LD_LIBRARY_PATH=/usr/local/cuda/lib64:/usr/lib/wsl/lib:$LD_LIBRARY_PATH \
PYTHONPATH=crates/pyxlog/python \
.venv/bin/python -m pytest python/tests/test_ilp_ga_reliability.py -q --timeout=1800
# 1 passed in 1497.67s

Additional quick smoke after commit split:

  • python/tests/test_ilp_performance.py -q --timeout=1800 -> 3 passed in 14.75s

levi770 added 12 commits March 1, 2026 22:59
Design doc addressing two implementation gaps found in dILP beta audit:
1. SparseMaskBackend N³ elimination via IlpMask enum + DLPack API
2. Row-count cache wiring across provider/executor/registry
…asks)

Detailed task-by-task plan with exact file paths, code, and test commands
for: row-count cache wiring, IlpMask enum, sparse executor path,
DLPack-native set_rule_mask_sparse, SparseMaskBackend rewrite, guard
tests, and reliability re-validation.
Copilot AI review requested due to automatic review settings March 4, 2026 11:16

Copilot AI left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pull request overview

This PR advances the dILP “GA hardening” workstream by tightening trainer/promotion gates, improving reproducibility and artifact auditability, adding telemetry/transfer accounting, and aligning documentation and tests with the updated semantics.

Changes:

  • Add deterministic training controls, persist selected_hard, and propagate forward-pass timing + memory telemetry summaries.
  • Move sparse mask application to a DLPack-based API (set_rule_mask_sparse) with Rust-owned top‑k ranking and a sparse executor path.
  • Add GA reliability/performance test suites and update docs (README/roadmap/architecture/changelog) to reflect current gate behavior.

Reviewed changes

Copilot reviewed 25 out of 25 changed files in this pull request and generated 8 comments.

Show a summary per file
File Description
python/tests/test_ilp_trainer.py Adds telemetry assertions (forward p95, step timings) and deterministic reproducibility checks.
python/tests/test_ilp_sparse_guard.py Adds guard coverage to prevent sparse path regressions (dense API usage / D2H counter behavior).
python/tests/test_ilp_sparse.py Updates sparse mask tests to pass CUDA tensors and adds sparse-backend behavioral assertions.
python/tests/test_ilp_robustness.py Extends robustness coverage (candidate-space degeneracy, CUDA OOM wrapping, NaN injection path).
python/tests/test_ilp_promoter.py Adds tests for typed-schema gate behavior and explicit holdout-f1 gate failure reporting.
python/tests/test_ilp_performance.py Adds telemetry + host-transfer accounting smoke tests and an optional slow SLO check.
python/tests/test_ilp_holdout.py Adds k-fold holdout scoring/variance coverage and propagates variance into TrainResult.
python/tests/test_ilp_ga_reliability.py Adds a statistical GA reliability gate (Clopper–Pearson lower bound).
python/tests/test_ilp_d2h_gate.py Verifies host transfer stats API exposure/reset semantics.
docs/plans/2026-03-01-sparse-executor-transfer-fix.md Adds an implementation plan for sparse executor + transfer elimination work.
docs/plans/2026-03-01-sparse-executor-transfer-fix-design.md Adds design doc detailing sparse-native mask flow and row-count cache wiring.
docs/architecture/dilp-training.md Updates architecture doc to reflect new holdout strategy, typed schema gate, and telemetry semantics.
docs/ROADMAP.md Updates milestone status and enumerates GA-hardening deliverables.
crates/xlog-runtime/src/ilp_registry.rs Refactors ILP masks to Dense/Sparse enum and caches row-count reads.
crates/xlog-runtime/src/executor.rs Adds sparse mask execution path and caches buffer row-count reads.
crates/xlog-cuda/src/provider.rs Wires row-count caching and adds untracked f64 downloads with host-transfer tracking.
crates/pyxlog/src/lib.rs Adds DLPack-native set_rule_mask_sparse, relation type annotations, and host transfer stats methods.
crates/pyxlog/python/pyxlog/ilp/types.py Extends TrainConfig and result types for GA gates/holdout/typed-schema controls.
crates/pyxlog/python/pyxlog/ilp/trainer.py Adds determinism controls, forward timing telemetry, selected_hard persistence, and holdout computation plumbing.
crates/pyxlog/python/pyxlog/ilp/promoter.py Adds holdout-f1 threshold and typed-schema gates to promotion flow.
crates/pyxlog/python/pyxlog/ilp/holdout.py Adds k-fold scoring + variance and a unified holdout_f1_and_variance() entry point.
crates/pyxlog/python/pyxlog/ilp/backend.py Updates backend protocol to return (cand_probs, selected_hard) and moves sparse masking to set_rule_mask_sparse.
crates/pyxlog/pyproject.toml Adds SciPy as an optional test dependency.
README.md Updates dILP feature summary and documents GA-hardening semantics/telemetry.
CHANGELOG.md Updates unreleased notes for GA hardening, telemetry, and reliability/performance gates.
Comments suppressed due to low confidence (1)

crates/pyxlog/python/pyxlog/ilp/backend.py:110

  • DenseMaskBackend.decode_argmax falls back to returning candidate index 0 if the argmax (i,j,k) is not found in candidates. If argmax lands on a non-candidate cell (which is possible given W is N^3), this will report the wrong candidate and can incorrectly drive stability/convergence logic. Consider computing argmax only over candidate positions, or raising/handling the non-candidate argmax explicitly instead of silently returning 0.
    def decode_argmax(self, W, candidates, n):
        with torch.no_grad():
            flat = W.view(-1)
            idx = flat.argmax().item()
            i = idx // (n * n)
            j = (idx % (n * n)) // n
            k = idx % n
        for ci, c in enumerate(candidates):
            if c["i"] == i and c["j"] == j and c["k"] == k:
                return ci
        return 0

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

Comment thread crates/pyxlog/python/pyxlog/ilp/backend.py Outdated
Comment thread crates/pyxlog/python/pyxlog/ilp/types.py Outdated
Comment thread crates/pyxlog/python/pyxlog/ilp/trainer.py
Comment thread docs/plans/2026-03-01-sparse-executor-transfer-fix.md Outdated
Comment thread crates/xlog-runtime/src/executor.rs
Comment thread crates/pyxlog/python/pyxlog/ilp/promoter.py Outdated
Comment thread python/tests/test_ilp_performance.py Outdated
Comment thread crates/pyxlog/python/pyxlog/ilp/holdout.py Outdated
@levi770
levi770 merged commit 50b9aa8 into main Mar 4, 2026
@levi770
levi770 deleted the feat/dilp-ga-hardening-docs-and-gates branch March 4, 2026 12:45
levi770 added a commit that referenced this pull request Mar 5, 2026
Add a CUDA kernel that builds a histogram of fact indices from sorted
COO data, producing the count array needed for CSR row_offsets via
prefix-sum. This eliminates Transfer #2 (sorted COO D2H for host CSR
construction) in the sparse executor pipeline.

- CUDA kernel: ilp_csr_histogram in kernels/ilp.cu
- Rust launcher: ilp_csr_histogram_launch on CudaKernelProvider
- Kernel registration in ILP module loading
- 3 tests: basic, empty, single_element (all pass)
- Certification suite: 206/206 pass (no regressions)
levi770 added a commit that referenced this pull request Apr 28, 2026
…done

ROADMAP.md still bundled hash-join + GroupBy + sort + dedup as
one unchecked item even though slices #5 and #6 already landed
the latter three. Splits that line into per-operator entries
and reflects the actual state of the v0.6.0 stream-safety
section:

* `[x]` Migrate the fused compare+scan+compact filter path
  (slice #3, `filter_fused_scan_recorded`) — was already
  unchecked.
* `[x]` Migrate sort (slice #5, `sort_recorded`).
* `[x]` Migrate dedup full-row (slice #5,
  `dedup_full_row_recorded`).
* `[x]` Migrate GroupBy (slice #6,
  `groupby_multi_agg_recorded` / `groupby_agg_recorded`).
* `[ ]` Migrate hash-join — the only remaining operator-
  surface migration in the v0.6.0 milestone, and the
  prerequisite for retaking the deferred GPU-resident
  binary-join materialization prototype.
* `[x]` Wire `filter_recorded` / `filter_columns_recorded`
  into a runtime / provider opt-in selector (slice #2,
  `XLOG_USE_RECORDED_FILTERS`).

Each entry annotates the slice it landed in and the
explicit narrowness of the migration so future work can
pick up the deferred bullets (multi-type sort, key-based
dedup, diff_full_row, union, LogSumExp, >4 key-column
GroupBy, ILP / ILP-exact, host-mask compact entry).

No code changes.
levi770 added a commit that referenced this pull request Apr 28, 2026
…in + cert mode

Slice 8 of the v0.6 stream-safety milestone: extends the env
opt-in pattern from slice #2 (filter) to the four remaining
operator classes. Per-operator env vars give real callers
opt-in routing through the recorded launch paths; an umbrella
`XLOG_USE_RECORDED_OPS=1` flag activates all five at once for
the integration cert mode. Defaults unchanged — legacy paths
remain the production default until the runtime stack is
certified end-to-end.

# Env vars

* `XLOG_USE_RECORDED_FILTERS` (existing)
* `XLOG_USE_RECORDED_SORT` (new) — gates `sort()`
* `XLOG_USE_RECORDED_DEDUP` (new) — gates `dedup_full_row()`
* `XLOG_USE_RECORDED_GROUPBY` (new) — gates
  `groupby_multi_agg()` (and `groupby_agg()` which forwards)
* `XLOG_USE_RECORDED_HASH_JOIN` (new) — gates
  `hash_join_v2()`, `hash_join_v2_with_limit()`, and
  `hash_join_v2_with_index()`
* `XLOG_USE_RECORDED_OPS` (new umbrella) — when set, all
  five gates activate via OR

Each `use_recorded_*_env()` helper checks its specific flag
OR the umbrella.

# Wiring

`recorded_filter_stream` field renamed to
`recorded_op_stream` (operator-agnostic) — one cached
non-default launch stream is shared across all five
recorded paths. The recorder serializes work on it; multiple
operations chain through commit-order events.

Each dispatcher checks env + runtime-backed manager + an
eligibility predicate matching the recorded variant's narrow
constraints (slice numbers in parens):

* sort: U32 / Symbol key columns only (#5)
* dedup_full_row: every column U32 / Symbol (#5)
* groupby_multi_agg: U32 / Symbol keys + Count / Sum / Min /
  Max aggs + ≤4 keys (#6)
* hash_join_v2*: ≤4 key columns (pack_keys constraint)
  (#7A / #7B / #7C / #7D)

Mismatches fall through to legacy. No public API changes.

# Bug fixes (surfaced by the cert run, not by drop+reuse)

Two pre-existing correctness bugs in earlier slices were
caught when env-gated dispatch fed real workloads:

1. **`sort_recorded` output `d_num_rows` was wrong** when
   input row_cap > logical row count. Slice #5 used
   `upload_device_row_count(input.num_rows())` —
   `input.num_rows()` returns row_cap, not the device-resident
   logical count. Downstream operators that read d_num_rows
   for the actual row count saw phantom rows with garbage data.
   Fix: alloc fresh `output_d_num_rows` and dtod-async copy
   from `input.num_rows_device()` on launch_stream
   (matching legacy `sort`'s `clone_device_row_count`).
   Output recorded via post-preflight fresh write.

2. **`compact_buffer_by_device_mask_counted_recorded`
   over-checked `d_mask.len() >= input.num_rows()`** (row_cap).
   The `mask_clamp_rows` kernel only reads `d_mask[i]` for
   `i < num_rows_device` (logical count); it tolerates
   `d_mask.len() < row_cap` as long as `d_mask.len() >=
   logical count`. Recorded hash_join semi/anti's mask
   sized to logical count was triggering false rejections.
   Fix: relax the check to `d_mask.is_empty() → error`;
   document that the kernel's own bounds against the
   device-resident logical count is authoritative. Same
   constraint as legacy `filter_by_device_mask`.

# Acceptance gates (all met)

* `cargo fmt --check` clean.
* `cargo test -p xlog-cuda --tests --release -- --test-threads=1`
  (env unset, default behavior): 403/403.
* Per-operator cert: each individual env var enables its
  target dispatcher and `XLOG_USE_DEVICE_RUNTIME=1
  cargo test -p xlog-integration --test real_world_tests
  --release -- --test-threads=8` passes 13/13.
* Umbrella cert mode: `XLOG_USE_RECORDED_OPS=1
  XLOG_USE_DEVICE_RUNTIME=1 cargo test -p xlog-integration
  --test real_world_tests --release -- --test-threads=8`:
  20/20 stress runs.

# Out of scope (deferred)

* GPU-resident binary-join materialization prototype
  (still archived per `archive/gpu-resident-binary-join-prototype-*`
  refs). Hash-join migration is now the prerequisite that's
  in place; the prototype retake is the next slice.
* Multi-type sort_recorded (i32, i64, u64, f32, f64, bool)
  for the recorded ILP / non-U32 GroupBy paths.
* LogSumExp + >4 key-column GroupBy in the recorded path.
* DLPack / Arrow runtime identity beyond what slice #4
  delivered.
levi770 added a commit that referenced this pull request Apr 29, 2026
Implements v0.6.0 release blocker #2: a public stress harness for the
recorded launch discipline that exercises both in-process parallel
scheduling (A3) and fresh subprocess fork (A4) under
`XLOG_USE_DEVICE_RUNTIME=1 XLOG_USE_RECORDED_OPS=1`.

# Harness shape

  * Two distinct workloads (per user direction):
    - `friends`: sort + hash-join sensitive
      (`fof(X,Z) :- friend(X,Y), friend(Y,Z), X != Z`)
    - `reach`: recursive fixed-point + joins
      (transitive closure over a randomized DAG)
  * Both workloads deterministic given (workload, nodes,
    edges_per_node) — graphs are seeded by params alone.
  * Stable FNV-1a checksum over sorted result rows; intentionally
    not `DefaultHasher` so cross-process A4 forks agree across
    mildly different builds.
  * Fixed schedule + seeded-random tail per worker. Each
    worker/child reports `base_seed`, `worker_id`, `iter`,
    `workload`, graph params on failure for repro.
  * Gate command (header-documented):
    `XLOG_USE_DEVICE_RUNTIME=1 XLOG_USE_RECORDED_OPS=1 \
       cargo test -p xlog-integration --test test_a3_a4_stress \
       --release -- --test-threads=1 --nocapture`
  * Env-unset entry surfaces a SKIPPED message rather than passing
    silently on the legacy path.

# A3 (in-process parallel)

8 threads × 32 iters. Each iter builds a fresh runtime/provider/
executor stack against the shared CUDA primary context, runs one
workload, and compares its checksum against a per-param reference
the parent computed serially. Compile is serialized via a
process-global mutex (Compiler symbol-interning is not concurrent).
CUDA kernel modules are pre-warmed in the parent thread before
worker spawn (cudarc / CUDA driver first-launch is not concurrent).

# A4 (subprocess fork)

16 fresh subprocesses × 4 iters per child. Parent re-invokes
`current_exe()` with `XLOG_A3A4_CHILD=<id>` and serializes the
reference table to env. Each child runs from a cold CUDA primary
context, exits 0/non-zero with a structured stderr line on
failure. Aggregates pass/fail in the parent.

# Current gate result on `83ece8bd` + this commit

  * A4: **16/16 PASS** — fresh-context cross-process is clean.
  * A3: **8 drift failures out of 256** measurements
    (~3% under contention).
  * Symptom tally: stream-misuse=0, uaf=0, drift=7-8, leak=0.

The harness is doing exactly what blocker #2 was designed for: it
surfaces real concurrent-execution drift in the engine that the
existing umbrella ×50 (which uses smaller workloads and lower
contention) does not catch. The drift is NOT in the recorded-launch
contract (`StreamMisuse`/`UseAfterFree` symptom counts are zero) —
it sits above the GPU layer, likely in the Datalog runtime's
interaction with the shared CUDA primary context under thread
contention.

Per the user direction "If A3/A4 reveals a stream or allocator bug,
docs written before it will be wrong" — the harness has revealed a
real correctness issue that needs a scoping decision (in-scope for
v0.6.0 vs. relocated to a follow-up release blocker) before
declaring blocker #2 closed.

# Files

  * crates/xlog-integration/Cargo.toml: dev-dependency on
    `rand` + `rand_chacha` (matches `xlog-cuda-tests`); new
    `[[test]]` entry for `test_a3_a4_stress`.
  * crates/xlog-integration/tests/test_a3_a4_stress.rs: harness.

No code changes outside the test target. The umbrella ×50 and
xlog-cuda --tests gates remain green.
levi770 added a commit that referenced this pull request Apr 29, 2026
The A3/A4 stress harness (commit 27ec3bd) was run through a 5-mode
diagnostic matrix to classify the A3 in-process thread-of-N drift:

| Mode | Fixture    | Runtime | Recorded | A3 drift |
|------|------------|---------|----------|----------|
| 1    | per_iter   | on      | on       | 7-8      |
| 2    | per_thread | on      | on       | 8        |
| 3    | shared     | on      | on       | 2 + 6 other |
| 4    | per_iter   | on      | off      | 7        |
| 4b   | per_thread | on      | off      | 7-8      |
| 5    | per_iter   | off     | off      | 0 + 8 other |
| 5b   | per_thread | off     | off      | 6-7      |

A4 fork-isolated stress passes 16/16 in every mode. Drift fires at
comparable rates with both env vars unset (legacy cudarc + legacy
operator dispatch + per-thread runtime), proving the bug class is
NOT introduced by the v0.6.0 stream runtime or recorded launches.
It is pre-existing same-process multi-executor / multi-provider
concurrency against a shared CUDA primary context.

Roadmap changes:

  v0.6.0 Tests and Certification:
    * A3/A4 reproducer suite item: marked DONE with a structured
      summary of the matrix evidence and the explicit re-scope.

  v0.6.0 Known Non-Blocking Residuals:
    * New entry for the A3 thread-of-N drift documenting the
      matrix finding and pointing to the v0.7.0 follow-up.

  v0.6.0 Release Gate:
    * "cert suite passes against runtime-backed manager" → checked
      (3361785).
    * "A3/A4 stress observes zero UAF / StreamMisuse" → checked.
      The symptom tally is `stream-misuse=0 uaf=0` in every matrix
      mode; the ROADMAP wording specifically covered UAF /
      StreamMisuse, and that bar is met. The note clarifies that
      A3 thread-of-N drift is a documented residual.

  v0.6.0 Release Blockers Remaining:
    * Blocker #2 closed (A3/A4 stress harness). Remaining: docs
      (#3) and host-mask/ILP recorded migration decision (#4).

  v0.7.0 Concurrency Hardening (NEW subsection):
    * "Certify same-process multi-executor concurrency against
      one CUDA primary context" with the matrix-derived
      re-target candidates and a pass criterion ("A3 thread-of-N
      drift drops to zero on per_thread / shared modes").

The v0.6.0 stream-safety release gate is now explicitly
**A4 + cert suite + umbrella ×50**, not "A3 must be zero drift".
levi770 added a commit that referenced this pull request Apr 29, 2026
The A3 portion of `test_a3_a4_stress.rs` was extended with three
selectors during the blocker #2 re-scope investigation. They need
to be in tree because the v0.6.0 ROADMAP now cites them as part
of the matrix evidence trail.

Added:

  * `run_workload_in(fx, p)` — runs one workload iter against a
    caller-provided `RuntimeFixture` so the harness can isolate
    cross-runtime churn from shared-runtime concurrency.
  * `run_workload_once(p)` retained as the convenience entry that
    builds a fresh fixture per call (PerIter mode equivalent).
  * `A3FixtureMode` enum + `XLOG_A3_FIXTURE_MODE` env var
    (`per_iter` default / `per_thread` / `shared`).
  * `XLOG_A3_DIAGNOSTIC` env var: bypasses the
    `XLOG_USE_DEVICE_RUNTIME=1 XLOG_USE_RECORDED_OPS=1` gate so
    the same binary can re-run on legacy / recorded-off paths
    for direct comparison against the runtime + recorded path.
  * `run_a3` rewritten to honor the fixture mode: PerIter
    (current behavior), PerThread (build one fixture per
    thread, reuse for all 32 iters), Shared (one
    `Arc<RuntimeFixture>` cloned across all threads).

The default test invocation
(`XLOG_USE_DEVICE_RUNTIME=1 XLOG_USE_RECORDED_OPS=1
XLOG_A3_FIXTURE_MODE` unset) still runs PerIter under the
recorded runtime, matching the v0.6.0 release-gate command in
the file header. The diagnostic capability is OFF by default;
nothing about the existing harness contract changes.

# Validation

  * `cargo fmt --check` — clean.
  * Default invocation still surfaces the same A3 thread-of-N
    drift documented in ROADMAP.md as a non-blocking residual;
    A4 still passes 16/16.

# Why this is committed but the harness is NOT a release gate

The harness committed at `27ec3bd9` was already documented in
ROADMAP.md as informational on the A3 dimension (matrix evidence
showed A3 drift is pre-existing legacy concurrency, not v0.6.0
stream safety). The diagnostic selectors are the apparatus that
produced that classification. Future re-runs of the matrix —
e.g., when the v0.7.0 "Certify same-process multi-executor
concurrency" backlog item is picked up — re-use these selectors.
levi770 added a commit that referenced this pull request Apr 29, 2026
Closes v0.6.0 release blocker #3.

# New files

  * docs/architecture/device-runtime.md — runtime stack
    (AsyncCudaResource → LoggingResource → GlobalDeviceBudget
    → XlogDeviceRuntime + StreamPool), the access-aware
    prepare/finish API, alloc-ready event semantics,
    GlobalDeviceBudget reservation + retry-after-reap,
    singleton (try_get) vs composed (with_resource)
    construction modes, env-gated dispatch matrix
    (XLOG_USE_RECORDED_FILTERS / SORT / DEDUP / GROUPBY /
    HASH_JOIN / OPS, plus XLOG_USE_DEVICE_RUNTIME on the
    integration / cert fixtures), and the supported
    certification modes (legacy 206/206 + runtime+recorded
    206/206). Closes with a pointer to the documented A3
    pre-existing concurrency residual.
  * docs/architecture/recorded-launch-migration.md —
    operator-author checklist for LaunchRecorder. API
    one-pager, the Access matrix (Read/Write/ReadWrite × what
    preflight waits on / what finish records to), helper-
    internal-scratch pattern via prepare_first_use /
    finish_first_use, host scalar reads + cu_stream.synchronize
    fence, external DLPack/Arrow handling (strict reject
    /permissive skip), and an explicit anti-pattern section
    covering: write_post_preflight_fresh removal, recording
    after preflight, using the singleton in production,
    long-term DeviceBlock retention, allocating helper scratch
    on launch_stream. Ends with the four-gate validation
    command sequence for migrating an operator.

# ARCHITECTURE.md update

Added a "v0.6 Device Runtime + Recorded Launch Discipline"
subsection inside Memory Management, linking both new docs.
No broader rewrite — surgical scope as directed.

# ROADMAP.md update

Marked the two v0.6.0 Documentation boxes complete, with
links to the new docs. Updated the v0.6.0 Release Blockers
Remaining summary: blocker #3 closed; remaining blocker #4 is
the host-mask compact / ILP / ILP-exact recorded-migration
scoping decision.

# v0.6.0 status snapshot after this commit

  * Blocker #1 (formal cert harness) — DONE (3361785).
  * Blocker #2 (A3/A4 stress) — DONE (27ec3bd + a55fb11
    re-scope: A4 fork-isolated PASS; A3 thread-of-N drift
    confirmed pre-existing against legacy default → v0.7.0).
  * Blocker #3 (docs) — DONE (this commit).
  * Blocker #4 (host-mask + ILP scoping) — pending.

No code changes.
levi770 added a commit that referenced this pull request May 14, 2026
…extension

Two blocking findings from user iteration-4 review of the
W3.2 implementation:

**Finding #1 — Tier-1 wrapper contract violated.** The plan
locked at iteration 4 §345 requires each k=6 ABI wrapper body
to contain **exactly one statement** that calls the shared
template, with NO conditionals (no `if`/`switch`/ternary), NO
loops. The implementation had 3-statement count wrappers and
5-statement materialize wrappers (thread-idx + bound check +
template call); the audit cert had been weakened from 1-stmt
to 3/5-stmt to match.

Restructured `wcoj.cu` to introduce two NEW grid-level
templates:

  template <int K_VAL, typename T>
  __device__ __forceinline__ void wcoj_clique_template_count_grid_t(...);

  template <int K_VAL, typename T>
  __device__ __forceinline__ void wcoj_clique_template_materialize_grid_t(...);

Both absorb the thread-idx + bounds checks. Each ABI wrapper
body is now EXACTLY ONE statement: a single template call into
the grid-level template. No `if`, no `for`, no `?:` anywhere
in the wrappers.

Tier-1 cert updated: shared `assert_wrapper_is_single_template_call`
helper enforces stmts == 1, contains `<grid_template>`<K, ...`>,
and rejects every conditional + loop token. All 4 Tier-1 cells
pass on the new tightened contract. Tier-2's
`no_six_literal_in_template_body` audit list extends to cover
the new grid-level templates.

**Finding #2 — Recursive WCOJ helper extension reverted.** The
plan iteration 4 §177 explicitly says the recursive WCOJ helper
is NOT extended for clique-keyed dispatch in W3.2. The earlier
implementation added `try_dispatch_wcoj_clique5_on_body` /
`_clique6_on_body` calls into `execute_wcoj_or_fallback_node`
in `recursive.rs:46-52`. Removed those entries.

Recursive clique bodies are still correctly rejected by the
promoter's `recursive_scan_count == 0` gate in `promote_multiway`,
so they fall through to binary-join. The non-recursive SCC
dispatch chain (in `execute_stratum_impl`) retains the clique
dispatch entries — that's the only path W3.2 wires up.

Verification:
* cargo fmt --check --all: clean.
* cargo build -p xlog-cuda --release: clean (compile-budget
  gate still passing with the grid-level templates).
* W3.2 acceptance: 33/33 still PASS
  (8 source-audit + 6 provider + 15 promoter + 4 dispatch).
* Workspace tests: 1990 PASS / 0 FAIL / 17 ignored.
* CUDA cert suite: 1/1.

Evidence README updated to reflect the corrected wrapper
structure + non-recursive-only dispatcher integration.
levi770 added a commit that referenced this pull request May 14, 2026
…t fix

Three findings on the recon doc:

* **Major**: spike scope expanded with duplicate-key fixture
  amendment. Original 1-col-sorted-unique design only measured
  set intersection (every match is 1:1). W4.3 is a binary-
  relation join operator; the spike must also exercise pair
  emission with run-length matching to justify production
  direction. Amended Lock #2 (kernel scope) to specify TWO
  arity regimes — (a) sorted-unique 1-col for upper-bound
  set-intersection measurement, (b) 2-col `(key, payload)`
  with duplicate keys for run-length pair-emission
  measurement; both regimes use the same kernel. Amended
  Lock #4 (parity) to specify
  `BTreeSet<(u32, u32, u32)>` for the duplicate-key cell.
  Amended Lock #5 (matrix) with a duplicate-rate symmetric
  cell (e.g. L=R=1000 with 4× key duplication → 16000
  output rows, output ≫ input).

* **Minor**: "Three tractable options" → "Four tractable
  options" (the table lists A/B/C/D).

* **Minor**: sortedness-check kernel precedent miscited as
  W3.2. Recited as "WCOJ layout fast-path / sorted-unique
  checker" with file refs:
  `crates/xlog-cuda/kernels/wcoj.cu` (kernels) and
  `crates/xlog-cuda/src/provider/wcoj.rs:3137` (u32),
  `:3265` (u64) (provider entry points). The originating
  closure-board attribution wasn't verified at recon time
  and isn't load-bearing for the spike's reference
  precedent.

Recon doc now correctly reflects the user-authorized spike
scope. Awaiting authorization satisfied; proceeding to
spike worktree creation + Phase 1 (kernel + manifest).

Refs: docs/plans/2026-05-08-w43-sort-merge-join-recon.md
levi770 added a commit that referenced this pull request May 14, 2026
User finding on iteration-5 commit 2db729a: the live plan still
carried soft `>= 1` counter criteria for Cert A/B/E/F and the
"either path" / "no kernel-launch crash; empty output via either
path" loophole text for Cert G across multiple canonical sites
(D7 row, Acceptance Grid, Step 6 + Step 7 + Step 10 prose). The
executed certs landed via the Step 6 patch (c665bd0 — Cert A
`>= 1` → `== 1`) and Step 10 patch (6f25377 — Cert G adds D7
route assertions per fresh-executor subcase) with stricter
exact-equality discipline. Plan was stale relative to certs.

Logged as **F-W43-13 (Major)** inside iteration 5. This is the
same class as F-W43-7/8/9/10 (file-wide-concern drift) but
applied to a contract tightening rather than a label bump.

Patches (in-place, all live sites):

D7 row (line 35):
* #1 (Cert A): `>= 1` → `== 1`, adds `nested_loop_dispatch_count == 0`.
* #2 (Cert B): `>= 1 or hash` → `nested_loop_dispatch_count == 1`
  (W4.2 fallback fired exactly once).
* #5 (Cert D'): adds explicit `nested_loop_dispatch_count == 0`
  (the certs already assert this; just makes it canonical).
* #6 (Cert E): `>= 1 + parity` → `== 1` + `nested_loop == 0` +
  parity.
* #7 (Cert F): `>= 1` → `== 1` + `nested_loop == 0` + output
  count == 4000 + all 4000 tuples distinct.
* #7' (Cert G): replaces "reflects the chosen short-circuit
  (either dispatched OR not-dispatched)" with explicit per-
  fresh-executor `sort_merge == 1 AND nested_loop == 0`
  proving the F-W43-4 contract end-to-end (sortedness probe
  short-circuits n<2 → Ok(true), dispatch admits, kernel empty
  fast path emits empty output).
* Row label provenance: now `per F-W43-3 + F-W43-4 + F-W43-12 +
  F-W43-13`.

Acceptance Grid (lines 225-232):
* Cert A: `>= 1` → `== 1`.
* Cert B: `>= 1` → `== 1`.
* Cert E: `>= 1 + parity` → `== 1 + nested_loop == 0 + parity`.
* Cert F: `>= 1 + output count == 4000` → `== 1 + nested_loop ==
  0 + output count == 4000 + all 4000 tuples distinct + parity`.
* Cert G: "no kernel-launch crash; empty output via either path;
  row-set parity" → "both subcases (empty L + empty R) per fresh
  executor: sort_merge == 1 + nested_loop == 0 + empty output +
  parity vs hash + no kernel-launch crash". Row label updated to
  `per F-W43-4 + F-W43-13`.

Step prose (lines 146, 159, 181-183):
* Step 6 (Cert A): `>= 1` → `== 1` (with "per F-W43-13 exact-
  equality discipline" rationale).
* Step 7 (Cert B): `>= 1` → `== 1` (with same rationale).
* Step 10 (Certs E + F + G): all three certs' counter assertions
  rewritten to match the executed exact-equality contract; Cert
  G prose specifically expanded to describe the F-W43-13 D7
  route assertion that closes the parity-only loophole.

Iteration-5 Amendment Log extended with F-W43-13 row (line 328)
+ updated process observation (line 332) noting that all three
findings (F-W43-11, F-W43-12, F-W43-13) are execution-discovered.
F-W43-13 specifically extends the F-W43-7/8/9/10 file-wide-
concern lesson: when a contract is tightened (not just labeled),
the same grep-everywhere discipline applies — D7 grid +
Acceptance Grid + prose Steps must all be inspected before the
iteration is declared closed.

Line 3 (Plan iteration metadata) and Plan-Approval Gate
paragraph (line 259) updated to reference all three findings
(was: two; now: three).

No code changes. No D-table design changes. No Step plan
structural changes. The +8 Acceptance Grid pass-count delta is
unchanged (Certs A, B, C, D, D', E, F, G still 8 cells).

Verification:
- Live plan grep for `>= 1` / `>=1` / "either path" returns
  hits ONLY inside the Iteration-5 Amendment Log (legitimate
  historical references describing the drift being patched).
- All cert-contract claims in the live plan now match the
  committed test bodies at
  `crates/xlog-integration/tests/test_w43_sort_merge_dispatch.rs`
  (Cert A line 348, Cert B line 482, Cert E line 911, Cert F
  line 1063, Cert G G1 line 1277, Cert G G2 line 1344).

Refs: docs/plans/2026-05-10-w43-sort-merge-join-plan.md (live
iteration 5, F-W43-13); commits c665bd0 (Step 6 patch — Cert A
exact-equality), 6f25377 (Step 10 patch — Cert G route
assertions).
levi770 added a commit that referenced this pull request Jun 12, 2026
…sics

XLOG_DEBUG_VERIFY_CLONES=1 byte-compares every cloned column against its
source inside clone_buffer, discriminating transport faults (clone wrong
at birth) from source faults (clone faithful, source already corrupt) in
the defect-#2 export-stage investigation. Off unless explicitly enabled,
same family as the XLOG_DEBUG_* probes from 4735ae9.
niveousdragon added a commit that referenced this pull request Aug 11, 2026
grep -rn -i maritime .github/workflows/ returned nothing: neither
test_maritime_convert.py nor test_maritime_cv.py was run by any of the 8
workflows, because every pytest invocation in CI is an explicit file
list (deep-review finding #2, HIGH — the same gap the #199 remediation
closed for the CAVIAR suites, repeated one PR later).

Extend the existing caviar-examples file list rather than adding a
separate maritime-examples job: the maritime suites need exactly the
environment that job already builds (pytest, CPU torch so the
torch-gated run_fold test runs for real, PYTHONPATH=crates/pyxlog/python
for the pure-Python pyxlog used by relational_search), and every other
job in ci.yml is one job per toolchain concern, not per example corpus
— duplicating checkout/setup/install for a second pytest list would be
the inconsistent choice. Both suites are synthetic-fixture-only, CPU
clean (local run: 24 passed).
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants