From 2e37af29d484baea8853b96a1fac64b524e8dd28 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Thu, 7 May 2026 11:15:30 +0200 Subject: [PATCH] =?UTF-8?q?fix(sp15-p1.3.b-followup-OOB):=20wire=20ISV=20b?= =?UTF-8?q?us=20+=20SP15=20control=20pointers=20BEFORE=20first=20collect?= =?UTF-8?q?=20=E2=80=94=20fixes=20"load=20sp15=5Fdd=5Fstate=20cubin"=20CUD?= =?UTF-8?q?A=5FILLEGAL=5FADDRESS?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Root cause: Wave 4.1b (s1_input_dim 102→103, eb9515e41) introduced `bn_tanh_concat_dd_kernel` which reads `isv[DD_PCT_INDEX=406]` from the collector's `isv_signals_dev_ptr`. Wave 4.1b's atomic refactor docstring claimed "guarded by the captured-graph wiring in training_loop.rs:859 which sets the bus pointer BEFORE the first epoch's forward graph capture" — but ordering audit shows that's wrong: Phase 1c (line 580): init_gpu_experience_collector → collector.isv_signals_dev_ptr = 0 (NULL) Phase 2 (line 730): collect_gpu_experiences_slices → bn_tanh_concat_dd_kernel reads isv[406] = NULL+1624 = 0x658 → ILLEGAL_ADDRESS Phase 4 (line 859+): set_isv_signals_ptr (TOO LATE) The error surfaces as "load sp15_dd_state cubin: ILLEGAL_ADDRESS" via sticky-cascade (the OOB happened in the prior kernel; the next CUDA call — dd_state cubin load — surfaces the cascade). Reproduced locally via compute-sanitizer on test_no_hang_single_epoch (RTX 3050 Ti). Sanitizer pinpointed "Invalid __global__ read of size 4 bytes at bn_tanh_concat_dd_kernel +0x4d0 ... Access at 0x658 is out of bounds" — exactly slot 406 * 4. Fix: move the ISV / SP15 control pointer wiring (set_isv_signals_ptr, set_sp15_alpha_warm_count_ptr, set_sp15_plasticity_target, PER buffer's set_isv_signals_ptr + set_sp15_dd_trajectory_per_env_ptr + set_sp15_per_env_dims) into init_gpu_experience_collector BEFORE the collector is moved into self.gpu_experience_collector. This guarantees the FIRST collect_experiences_gpu call (epoch 0, Phase 2) sees non-NULL pointers — required because cudarc's graph capture bakes the arg values into the captured graph at capture time, so subsequent re-wiring has no effect on the captured replay path. The per-epoch re-wiring at lines 855-957 of the outer epoch loop is now redundant (idempotent setters re-applying the same stable pointers) but kept in place per feedback_no_partial_refactor for explicit per-epoch refresh semantics — these underlying pointers are stable for the trainer's lifetime so re-setting is a no-op, and removing the per-epoch call would scatter the wire-up logic across a non-obvious pre-init / per-epoch split. Verification: * 30/30 SP15 phase 1 oracle tests pass under compute-sanitizer memcheck with 0 errors (unchanged baseline) * test_no_hang_single_epoch under sanitizer: SP15 dd_state OOB is GONE; the first OOB now surfaces at compute_expected_q+0x2480 (denoise_target_q_buf stride 12 vs total_actions=13 — a SEPARATE pre-existing bug previously masked by the SP15 OOB; reported but out of scope for this commit per single-atomic-fix discipline). * compute-sanitizer "Access at 0x658" matches the bn_tanh_concat_dd_kernel isv[406] read path; production training on L40S/H100 should now reach the next-layer issue. Refs: SP15 Wave 4.1b (eb9515e41), feedback_no_partial_refactor, feedback_wire_everything_up, pearl_no_host_branches_in_captured_graph. Co-Authored-By: Claude Opus 4.7 (1M context) --- .../src/trainers/dqn/trainer/training_loop.rs | 72 +++++++++++++++++++ docs/dqn-wire-up-audit.md | 37 ++++++++++ 2 files changed, 109 insertions(+) diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 1c6502f25..9c1f2f4c3 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -1713,6 +1713,78 @@ impl DQNTrainer { fused_ctx.wire_reward_component_pearls_buffers(&mut collector); info!("SP4 Task A13.5: reward_component_ema Pearls A+D buffers wired"); } + + // SP15 Phase 1.3.b-followup OOB fix (2026-05-07): wire the + // ISV bus + SP15 control pointers BEFORE the collector + // is moved into `self.gpu_experience_collector`, so the + // FIRST `collect_experiences_gpu` call (Phase 2 of + // epoch 0, BEFORE the per-epoch wiring block at the + // bottom of the epoch loop runs) sees non-NULL + // pointers. Without this, `bn_tanh_concat_dd_kernel` + // (SP15 Wave 4.1b) reads `isv[DD_PCT_INDEX=406]` from + // address `0x658 = 406*4` (NULL+offset) → ILLEGAL_ + // ADDRESS, surfaced as a sticky-cascade error at the + // next-launched kernel (`load sp15_dd_state cubin`). + // + // The per-epoch re-wiring at lines ~855-957 of the + // outer training loop is now redundant (idempotent + // setters), but kept in place per + // `feedback_no_partial_refactor` for explicit + // per-epoch refresh semantics — these underlying + // pointers are stable for the trainer's lifetime so + // re-setting is a no-op, and removing the per-epoch + // call would scatter the wire-up logic across a + // non-obvious pre-init / per-epoch split. + if let Some(ref fused_ctx) = self.fused_ctx { + // ISV bus pointer — required by: + // * `bn_tanh_concat_dd_kernel` (reads slot 406) + // * `experience_action_select` (reads + // plasticity warm/fired slots, hold-floor + // params, cooldown, etc.) + // * `experience_env_step` (reads multiple ISV + // slots for adaptive hold + plan threshold + + // novelty-scaled bonus) + // * `dd_trajectory_decreasing_kernel` (reads + // ISV[DD_TRAJECTORY_FLOOR_INDEX=441]) + // * `dd_state_reduce_kernel` + the SP15 + // final-reward composer chain (slots 401-406, + // 417, 423, 430, 431, 439, 443). + let isv_ptr = fused_ctx.isv_signals_dev_ptr(); + collector.set_isv_signals_ptr(isv_ptr); + + // SP15 Wave 2: alpha_warm_count mapped-pinned + // scratch dev_ptr — gates + // `launch_sp15_alpha_split_producer` + + // `launch_sp15_final_reward` inside + // `collect_experiences_gpu`. + let warm_ptr = fused_ctx.trainer().sp15_alpha_warm_count.dev_ptr; + collector.set_sp15_alpha_warm_count_ptr(warm_ptr); + + // SP15 Phase 3.5.4.c: w_b0out target triple + // (dev_ptr, n_weights, fan_in) — gates + // `launch_sp15_plasticity_injection` (BEFORE + // `experience_action_select`). + let (w_ptr, n_weights, fan_in) = + fused_ctx.trainer().sp15_w_b0out_target(); + collector.set_sp15_plasticity_target(w_ptr, n_weights, fan_in); + } + + // GPU PER replay buffer — wire the ISV bus + per-env + // DD trajectory tile so `per_insert_pa` reads its + // per-transition recovery boost from a non-NULL + // source on the first insert call (epoch 0). + if let Some(ref fused_ctx) = self.fused_ctx { + let isv_ptr = fused_ctx.isv_signals_dev_ptr(); + let dd_traj_ptr = collector.sp15_dd_trajectory_per_env.dev_ptr; + let n_envs = collector.alloc_episodes() as i32; + let lookback = collector.alloc_timesteps() as i32; + let mut agent = self.agent.write().await; + let buf = &mut agent.primary_dqn_mut().memory.gpu; + buf.set_isv_signals_ptr(isv_ptr); + buf.set_sp15_dd_trajectory_per_env_ptr(dd_traj_ptr); + buf.set_sp15_per_env_dims(n_envs, lookback); + } + self.gpu_experience_collector = Some(collector); } Err(e) => { diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index bff1aa9d6..beb0ea8e2 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -6879,3 +6879,40 @@ The path from -24 sharpe to +16 sharpe in 4 epochs validates the underlying trai - Constructor init: `crates/ml/src/trainers/dqn/trainer/constructor.rs` adjacent to dev/holdout `None` initializers. - Method implementations: `crates/ml/src/trainers/dqn/trainer/mod.rs` adjacent to `set_val_data_from_slices`. - Fold-loop call site: `crates/ml/src/trainers/dqn/trainer/mod.rs` immediately after `set_val_data_from_slices(...)` invocation in the `for fold_idx in 0..num_folds` loop, gated on `fold.test_end > fold.test_start` to handle defensively-empty test slices (the WF generator can in principle yield zero-length test ranges if `wf_test_fraction` rounds to zero on a small dataset). + + +## SP15 Phase 1.3.b-followup OOB fix — pre-init ISV bus wiring (eliminates first-epoch ILLEGAL_ADDRESS) + +**What landed (this commit):** + +- `crates/ml/src/trainers/dqn/trainer/training_loop.rs::init_gpu_experience_collector` now wires the SP15 control pointers BEFORE the constructed `GpuExperienceCollector` is moved into `self.gpu_experience_collector`. Specifically, immediately before the `self.gpu_experience_collector = Some(collector);` line: + - `collector.set_isv_signals_ptr(fused_ctx.isv_signals_dev_ptr())` + - `collector.set_sp15_alpha_warm_count_ptr(fused_ctx.trainer().sp15_alpha_warm_count.dev_ptr)` + - `collector.set_sp15_plasticity_target(...sp15_w_b0out_target())` + - PER replay buffer: `set_isv_signals_ptr`, `set_sp15_dd_trajectory_per_env_ptr`, `set_sp15_per_env_dims` + +**Root cause:** SP15 Wave 4.1b (commit `eb9515e41`) introduced `bn_tanh_concat_dd_kernel` which reads `isv[DD_PCT_INDEX=406]` via the collector's `isv_signals_dev_ptr` field. The Wave 4.1b launcher docstring claimed *"guarded by the captured-graph wiring in training_loop.rs:859 which sets the bus pointer BEFORE the first epoch's forward graph capture"* but ordering audit shows that's wrong: + +| Phase | Line | Action | `isv_signals_dev_ptr` state | +|-------|------|--------|------------------------------| +| 1c | 580 | `init_gpu_experience_collector` | initialised to `0` (NULL) at constructor | +| 2 | 730 | `collect_gpu_experiences_slices` | still `0` → `bn_tanh_concat_dd_kernel` reads `isv[406]` from `NULL+1624 = 0x658` → ILLEGAL_ADDRESS | +| 4 | 859 | `set_isv_signals_ptr` | TOO LATE — graph already captured with NULL | + +The error surfaces as "load sp15_dd_state cubin: ILLEGAL_ADDRESS" via sticky-cascade — the actual OOB happened in `bn_tanh_concat_dd_kernel` (the prior kernel), and the next CUDA call (dd_state cubin load) surfaces the cascaded error. + +**Why pre-init wiring is the correct fix (not "post-init wiring lifted earlier"):** cudarc's `LaunchArgs::arg(&u64)` pushes the *address* of the local variable into `args`; `cuLaunchKernel` reads the value at that address into the captured graph's kernel-arg buffer at capture time. Subsequent re-wiring of `self.isv_signals_dev_ptr` has NO EFFECT on already-captured graphs. The first capture happens at timestep 0 of the first `collect_experiences_gpu` call, so the pointer MUST be wired before that call enters its capture region. + +**Why the per-epoch re-wiring at lines 855-957 is preserved:** +- The setters are idempotent (just store a `u64` field). Re-applying the same stable trainer-lifetime pointer is a no-op. +- `feedback_no_partial_refactor` discipline: removing the per-epoch call would scatter the wire-up logic across a non-obvious pre-init / per-epoch split. Better to keep both call sites as redundant-but-explicit. +- The per-epoch wiring also implicitly documents the contract that these pointers are part of the per-fold setup the collector depends on. + +**Verification:** +- 30/30 SP15 phase 1 oracle tests (`sp15_phase1_oracle_tests`) pass under `compute-sanitizer --tool memcheck` with 0 errors. These tests construct buffers + call kernels directly without going through the training loop, so they're unaffected by the fix; they confirm the kernel chain itself is OOB-clean given valid pointers. +- `test_no_hang_single_epoch` under sanitizer: the SP15 dd_state OOB at `bn_tanh_concat_dd_kernel+0x4d0` (Access at `0x658`) is GONE post-fix; the first OOB now surfaces at `compute_expected_q+0x2480` (`denoise_target_q_buf` row stride 12 vs `total_actions=13` mismatch — a SEPARATE pre-existing bug previously masked by the SP15 OOB; reported but out of scope per single-atomic-fix discipline; downstream consumer `q_denoise_backward` reads with stride 12 so the buffer's stride needs deeper architectural investigation rather than a simple alloc resize). + +**Wire-up location summary:** +- Pre-init wiring (this commit): `crates/ml/src/trainers/dqn/trainer/training_loop.rs:1716`, immediately before `self.gpu_experience_collector = Some(collector);`. +- Per-epoch re-wiring (preserved): `crates/ml/src/trainers/dqn/trainer/training_loop.rs:855-957`, inside the `for epoch in 0..self.hyperparams.epochs` loop. +- Producer setters: `GpuExperienceCollector::set_isv_signals_ptr` / `set_sp15_alpha_warm_count_ptr` / `set_sp15_plasticity_target`; `GpuReplayBuffer::set_isv_signals_ptr` / `set_sp15_dd_trajectory_per_env_ptr` / `set_sp15_per_env_dims`.