fix(sp15-p1.3.b-followup-OOB): wire ISV bus + SP15 control pointers BEFORE first collect — fixes "load sp15_dd_state cubin" CUDA_ILLEGAL_ADDRESS
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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) => {
|
||||
|
||||
@@ -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`.
|
||||
|
||||
Reference in New Issue
Block a user