foxhunt/docs/superpowers/notes/2026-06-02-determinism-phase2.6-backward-kernel-fix.md

Determinism Phase 2.6 — backward-pipeline investigation + DETERMINISM ACHIEVED

Date: 2026-06-02 Dispatch: Phase 2.6 follow-up to Phase 2.5's RL-loop falsification. Outcome: Two distinct non-determinism sources identified and fixed. Same-seed runs now produce bit-equal training checksums + bit-equal eval summaries. Determinism foundation goal achieved. Verdict: Phase 2 → 2.6 falsification chain converged. No further phase needed.

---

Recap (Phase 2.5 → Phase 2.6)

Phase 2.5 falsified all 6 RL-loop candidate groups (CfC carry, PER state, action RNG, LOB sim, ISV controllers, reward shaping). At step 2 dump every forward activation + every RL-loop scalar was EQUAL, yet every trunk weight + every trunk gradient + every mamba2 backward output DIVERGED. The bug therefore lived in step 2's BACKWARD KERNEL CHAIN, in a layer not yet instrumented. Phase 2.5's dispatch ask listed 4 not-yet-dumped buffer classes (G/H/I/J).

Task 1 — Extended instrumentation

Added IntegratedTrainer::dump_backward_state_for_debug(step) invoked from build_diag_value for step ∈ {0, 1, 2, 3}. Captures the 4 new buffer groups via DtoH copies on self.stream:

• Group G — PER replay sample outputs: sampled_h_t, sampled_h_tp1, sampled_actions, sampled_rewards, sampled_dones, sampled_log_pi_old (end-of-step snapshot — these are overwritten each k_iter, so last-k_iter-wins).
• Group H — Head forward outputs: q_logits, pi_logits, v_pred, iqn_q_values, dueling_q_composed (post-fwd, pre-bwd).
• Group I — Pre-reduce per-batch / per-row param-grad scratches: cfc_grad_w_in/_w_rec/_b/_tau_scratch, vsn_grad_w/_b_scratch, attn_grad_q_scratch, grad_ln_a/_b_gain/_bias_per_row (zeroed at step start, accumulated by K-loop +=, reduced before Adam, live until next step's memset).
• Group J — Per-head grad_h_t + combined: ss_q/_pi/_v/_frd/ _outcome_grad_h_t_d, grad_h_t_combined_d (the encoder-backward upstream gradient).

scripts/compare-backward-state.py reads run_a/ vs run_b/ dumps and emits a per-group verdict + sub-case interpretation per the Phase 2.6 decision tree. scripts/determinism-check.sh gained a new --debug-dump-backward flag composing with the existing --debug-dump, --debug-dump-mamba2, --debug-dump-rl flags.

Added Group K mid-investigation when the initial Group I/J verdict at step 2 pointed at grad_h_t_outcome (Δ ~ 0.003, far above fp32 noise) — Group K dumps outcome_head.labels_d, w_d, b_d to localise WHICH input drove the outcome-head divergence.

Task 2 — Verdict 1 (after first run)

FOXHUNT_DETERMINISTIC=1 ./scripts/determinism-check.sh --quick \
  --debug-dump-backward

Per-step verdict:

Group	Step 0	Step 1	Step 2
G — PER sample buffers	EQUAL	EQUAL	EQUAL
H — Head fwd outputs	iqn_q_values DIVERGE	iqn_q_values DIVERGE	iqn_q_values DIVERGE
I — Pre-reduce grad scratches	EQUAL	EQUAL	ALL DIVERGE
J — Per-head grad_h_t	EQUAL	EQUAL	grad_h_t_outcome + grad_h_t_combined DIVERGE

Decision-tree sub-case (2): Group H diverges with EQUAL Group G → head forward kernel is non-deterministic. Specifically: iqn_q_values first diverges at step 0 with max|Δ| = 4.5e-5. No upstream RL-loop perturbation needed; the divergence enters during IQN's own forward.

Bug 1 — IQN tau-sampling read-write race

crates/ml-alpha/cuda/rl_iqn_forward.cu's rl_iqn_tau_cos_features was launched as Grid=(B, N_TAU, 1) Block=(EMBED_DIM, 1, 1). For each batch=b, all N_TAU=32 blocks for that batch ran concurrently across multiple SMs.

The kernel had a read-write race on prng_state[batch]:

• Every block read seed = prng_state[batch] at line 88.
• Block (b, 0) computed adv = seed; xorshift32 × 8; prng_state[batch] = adv at lines 102-107.
• No inter-block barrier. With 32 blocks per batch and SMs scheduling blocks in run-dependent order, blocks with tau_idx > 0 sometimes read the original seed (if scheduled before block 0's write) and sometimes the advanced seed (if scheduled after). Different runs picked different orderings → different local_state for blocks tau_idx > 0 → different τ values → different iqn_q_values.

The race produced run-dependent values at step 0 already — no upstream training divergence required to enter the picture. The divergence didn't cascade into trunk gradients (Groups I/J) until step 2 because IQN's backward only writes IQN-internal weight grads (no grad_h_t flow back into the encoder), so the encoder ran from identical inputs at step 0/1. The cascade started at step 2 once the second bug fired (see below).

Fix 1 — Split sampling and state-advance into two kernels

Made rl_iqn_tau_cos_features READ-ONLY on prng_state (added const qualifier to the pointer; removed the tau_idx==0 advance block at lines 102-107).

Added a sibling kernel rl_iqn_advance_prng_state launched on the same stream immediately AFTER rl_iqn_tau_cos_features. The companion advances prng_state[batch] by 8 xorshift32 steps with Grid=(ceil(B/256), 1, 1) Block=(256, 1, 1) — one thread per batch element, sole writer per batch. Kernel-launch ordering on a single stream is a grid-wide barrier (replaces the racy "tau_idx==0 block writes prng_state at end of monolithic kernel" pattern).

Per feedback_no_atomicadd: no atomic added. Per feedback_no_nvrtc: the new kernel goes through build.rs (same cubin as the existing forward kernels). Per feedback_no_legacy_aliases: the old kernel signature changed (uint32_t* → const uint32_t* on prng_state); no alias kept.

Task 2 — Verdict 2 (after Fix 1)

After Bug 1 fix, iqn_q_values matches at step 0/1/2. But step 2 still diverged in Group I (all pre-reduce scratches) + Group J (grad_h_t_outcome with Δ = 0.003 — orders of magnitude above fp32 noise).

The fused outcome kernel is sole-writer-per-(batch, c) so it cannot introduce divergence from equal inputs. Added Group K dump (outcome_head.labels_d, w_d, b_d) to identify the upstream input drift.

=== step 0 ===
  -- group K (outcome head inputs) --
    outcome_labels: EQUAL  (128 i32 elements — all sentinel -1)
    outcome_w:      DIVERGE  max|Δ|=0.00309763  (384 f32 elements)
    outcome_b:      EQUAL    (3 f32 elements — zero-init)

Bug 2 — OutcomeHead init missing scoped_init_seed: OutcomeHead::new called ml_core::cuda_autograd::init:: near_zero_xavier(...) without first installing a scoped_init_seed(seed) guard. near_zero_xavier defers to generate_uniform which falls back to a time+thread-id-seeded xoshiro256++ chain when no thread-local INIT_RNG is set (crates/ml-core/src/cuda_autograd/init.rs:144-164). Two same-seed processes therefore drew DIFFERENT initial outcome weights.

The divergence stayed dormant for the first 1-2 steps because outcome_head.labels_d starts at sentinel -1 and only gets overwritten by assign_labels on done events (line 7876). With no done events yet, the fused outcome kernel writes s_grad_logits[k] = 0 (label < 0 branch) → grad_h_t = 0 regardless of w. The first done event in mid-smoke fired at step 2, which is when grad_h_t_outcome started flowing the divergent weights into the encoder gradient and the cascade began.

Every other head (dqn, iqn, policy, value, dueling_q) already installed the scoped_init_seed guard around its weight draws (per pearl_scoped_init_seed_for_reproducibility). The outcome head was the only omission.

Fix 2 — Add seed parameter to OutcomeHead::new + install guard

Changed OutcomeHead::new(b_size, stream) → OutcomeHead::new(b_size, seed, stream). Wrapped the near_zero_xavier call in a { let _seed_guard = ml_core::cuda_autograd::init::scoped_init_seed(seed); ... } block so the thread-local RNG is seeded for the duration of the W draw.

Trainer call site in IntegratedTrainer::new now passes cfg.dqn_seed.wrapping_add(0x0CE0) (= "OCE0", Outcome Classifier Encoder 0) for an init draw stream independent of the other heads.

Per feedback_no_legacy_aliases: signature changed directly. Per feedback_no_partial_refactor: all 3 callers (trainer + 2 tests) migrated atomically.

Task 3 — Validation

Post-fix ./scripts/determinism-check.sh --quick (200 train + 50 eval, b=128, RTX 3050):

DETERMINISTIC: all `checksums.*` leaves match across all 200 rows
(rel-tol=1e-05, abs-tol=1e-07).

Tested at seeds 42, 43, 44 — all PASS.

diff <(jq -c 'del(.elapsed_s)' run_a/diag.jsonl) <(jq -c 'del(.elapsed_s)' run_b/diag.jsonl) returns empty — same-seed diag.jsonl is byte-equal modulo wall-clock timing.

diff /tmp/foxhunt-determinism-a/eval_summary.json /tmp/foxhunt-determinism-b/eval_summary.json returns empty:

{
  "b_size": 128,
  "max_drawdown_usd": 0.0,
  "n_eval_trades_dropped": 0,
  "n_eval_trades_seen": 1,
  "n_pre_eval_trades_wrapped": 0,
  "n_trades": 1,
  "profit_factor": 0.0,
  "sharpe_ann": -359035207680.0,
  "total_pnl_usd": -12.5,
  "win_rate": 0.0
}

Eval pnl: -$12.50 in both runs — bit-equal.

diff /tmp/foxhunt-determinism-a/alpha_rl_train_summary.json /tmp/foxhunt-determinism-b/alpha_rl_train_summary.json returns empty — train summary also bit-equal.

Backward-pipeline verdict (post fix)

=== step 0 ===
  -- group G (PER replay sample) --       ALL EQUAL
  -- group H (head fwd outputs) --        ALL EQUAL  (q/pi/v/iqn/dueling)
  -- group I (pre-reduce grad scratch) -- ALL EQUAL
  -- group J (per-head grad_h_t) --       ALL EQUAL
  -- group K (outcome head inputs) --     ALL EQUAL  (labels/w/b)
=== step 1 ===  ALL EQUAL across all groups
=== step 2 ===  ALL EQUAL across all groups
=== step 3 ===  ALL EQUAL across all groups

Speed cost

200-step quick mid-smoke on RTX 3050:

Configuration	Wall-clock
FOXHUNT_DETERMINISTIC=0 (TF32, no fix)	not measured — pre-fix divergent
FOXHUNT_DETERMINISTIC=1 (PEDANTIC, pre-Phase-2.6)	73.5s
FOXHUNT_DETERMINISTIC=1 (PEDANTIC, post-Phase-2.6 fixes)	73.8s

The IQN sampling split (one additional kernel launch per IQN forward call, 2 calls per K-iter) adds ~0.3s over 200 steps = ~1.5ms/step amortised. No measurable cost from the OutcomeHead seed guard (one-shot at construction).

Operational rule

A multi-block kernel that reads-then-writes a shared scalar with no inter-block barrier is a race. Canonical fix: split the kernel so the shared scalar is either pure-read or written by a single narrow kernel on the same stream. Kernel-launch ordering on a single stream is the cleanest grid-wide barrier available without cooperative_groups::grid().sync().

Every weight-init call site must install scoped_init_seed. Without it, init::* functions fall back to time+thread-id RNG and two same-seed processes draw different initial weights, producing silent training-trajectory divergence that doesn't appear in checksums until the divergent weights first affect the loss (potentially many steps later — outcome_head had a 2-step lag because labels stayed at sentinel until the first done event).

Per feedback_extending_existing_code_audits_for_existing_violations: when adding a new head, audit existing heads for the scoped_init_seed pattern. The omission in OutcomeHead was a 2026-pre-Phase-2.6 oversight; this dispatch retroactively fixed it.

Phase summary: PER → cuBLAS → mamba2 → RL state → backward chain

Five-phase falsification chain (Phase 2 → 2.6) achieves bit-equal same-seed determinism:

Phase	Bug	Fix LOC
2 (PER tree rebuild)	__threadfence between levels with parallel blocks	~30 LOC: single-block + __syncthreads
2.3 (cuBLAS GEMM)	CUBLAS_GEMM_DFALT with TF32 split-K non-determinism	~80 LOC: CUBLAS_PEDANTIC_MATH at all 3 cuBLAS sites
2.4 (mamba2 falsified)	— (Phase 2.3 already deterministic)	—
2.5 (RL-loop falsified)	— (all 6 candidate groups EQUAL through step 2)	—
2.6 — Bug 1 (IQN tau sampling race)	Multi-block RW race on prng_state[batch]	~50 LOC: split kernel + new sibling kernel
2.6 — Bug 2 (OutcomeHead unseeded init)	Missing scoped_init_seed guard around near_zero_xavier	~10 LOC: add seed param + install guard

Total LOC across all 5 phases: ~170 LOC (5 .cu / .rs touches), zero behavioral changes (training math unchanged), zero performance regression (≤ 1.5ms/step overhead measured).

Files touched (Phase 2.6, uncommitted)

Path	Change	LOC
crates/ml-alpha/cuda/rl_iqn_forward.cu	rl_iqn_tau_cos_features made READ-ONLY on prng_state; new rl_iqn_advance_prng_state sibling kernel	+50
crates/ml-alpha/src/rl/iqn.rs	Load rl_iqn_advance_prng_state symbol; call after rl_iqn_tau_cos_features in forward_inner	+20
crates/ml-alpha/src/rl/outcome_head.rs	OutcomeHead::new signature +seed parameter; install scoped_init_seed guard around near_zero_xavier	+15
crates/ml-alpha/src/trainer/integrated.rs	Pass cfg.dqn_seed.wrapping_add(0x0CE0) to OutcomeHead::new; new dump_backward_state_for_debug method (env-gated by FOXHUNT_DETERMINISM_DEBUG_BACKWARD=1) + invocation in build_diag_value	+290
crates/ml-alpha/src/trainer/perception.rs	Added 11 debug_*_view accessor methods for the Phase 2.6 pre-reduce grad scratches	+30
scripts/determinism-check.sh	New --debug-dump-backward flag composing with existing PER/mamba2/RL dump flags	+30
scripts/compare-backward-state.py	New: bytewise compare for 4 backward-pipeline buffer classes + Group K outcome inputs with verdict + decision-tree interpretation	+290 (new file)

No commit. Discipline checklist:

• No new files in repo root
• Scripts in scripts/, docs in docs/superpowers/notes/
• One .cu file touched: kernel split + new sibling (no functionality removed)
• Per feedback_no_atomicadd.md: no atomic added
• Per feedback_no_nvrtc.md: cubin route via existing build.rs entry
• Per feedback_no_quickfixes.md: actual root-cause fixes (race + seed), not tolerance masks or f64 accumulators
• Per feedback_local_smoke_before_cluster.md: validated on RTX 3050, seeds 42/43/44, no cluster submit
• Per feedback_systematic_debugging + dispatch STOP rule: dump first, identify, then fix. Two distinct bugs surfaced in sequence; both fixes anchored on raw-byte evidence
• No commit — left uncommitted for human review

Reproduction

# Build with Phase 2.6 fixes
SQLX_OFFLINE=true cargo build --release --example alpha_rl_train -p ml-alpha

# Run determinism check (no extra flags needed)
FOXHUNT_DETERMINISTIC=1 ./scripts/determinism-check.sh --quick

# Expected output:
#   DETERMINISTIC: all `checksums.*` leaves match across all 200 rows
#   (rel-tol=1e-05, abs-tol=1e-07).

# Run for seeds 43, 44 too
FOXHUNT_DETERMINISTIC=1 ./scripts/determinism-check.sh --quick --seed 43
FOXHUNT_DETERMINISTIC=1 ./scripts/determinism-check.sh --quick --seed 44

# Verify eval-summary bit-equality
diff /tmp/foxhunt-determinism-a/eval_summary.json \
     /tmp/foxhunt-determinism-b/eval_summary.json

# Verify train-summary bit-equality
diff /tmp/foxhunt-determinism-a/alpha_rl_train_summary.json \
     /tmp/foxhunt-determinism-b/alpha_rl_train_summary.json

# Run the backward-state dump diagnostic if you want to see the
# now-deterministic state at step 0/1/2/3 explicitly
FOXHUNT_DETERMINISTIC=1 ./scripts/determinism-check.sh --quick \
  --debug-dump-backward
# Expected: VERDICT: all dumped backward-pipeline buffers match
# across all dumped steps.