diff --git a/crates/ml/tests/sp15_phase1_oracle_tests.rs b/crates/ml/tests/sp15_phase1_oracle_tests.rs index dc2a3a0f5..67c21bc0d 100644 --- a/crates/ml/tests/sp15_phase1_oracle_tests.rs +++ b/crates/ml/tests/sp15_phase1_oracle_tests.rs @@ -2457,6 +2457,420 @@ mod sp15_wave_4_1a_test_helpers { } } +/// SP15 Wave 4.1c — behavioral KL test that closes out the Wave 4.1 +/// (Phase 1.5.b) state_dim 102→103 + dd_pct trunk integration. +/// +/// Wave 4.1a (`a8da1cb9c`) landed `bn_tanh_concat_dd_kernel`. Wave 4.1b +/// (`eb9515e41`) wired `s1_input_dim` 102→103 through the trunk + 4 +/// forward + 3 backward call sites. Wave 4.1c proves the wiring +/// actually shifts the policy distribution: a synthetic GPU-side +/// projection of the post-bn_concat_dd buffer with random Xavier-init +/// weights produces measurably different action distributions when +/// ISV[DD_PCT_INDEX=406] differs (0.0 at-ATH vs 0.10 in-DD). +/// +/// Why a synthetic projection (vs the real GRN trunk): +/// `forward_trunk_for_test`'s ideal contract — a public API that runs +/// the real GRN encoder + advantage heads against caller-owned +/// post-bn_concat input — would require either (a) a public surface +/// change on `DQNTrainer` exposing internal cuBLAS handles + GRN scratch +/// buffers + weight pointers (the trainer's `fused_ctx` is `pub(crate)` +/// and only initialised inside the training loop, not by +/// `DQNTrainer::new`), or (b) duplicating the trunk's cuBLAS setup in a +/// test (≥200 lines of buffer + handle plumbing). Both options are +/// architecturally heavier than the test's purpose justifies. Per the +/// spec dispatch ("the test's purpose is 'non-zero KL proves the wire +/// is connected' not 'verifies trained behavior'"), we instead compose +/// `launch_sp15_bn_concat_dd` (the production kernel that injects +/// dd_pct into the trunk input) with a single-layer cuBLAS sgemm using +/// random Xavier-init weights `[proj_h, 103]`. The matmul exercises the +/// new column-103 weights on the dd_pct value — exactly the path the +/// real GRN's `Linear_a` first GEMM takes for `w_a_h_s1[:, 102]` (the +/// dd_pct column added by Wave 4.1b's reshape). Two ISV values produce +/// two `[B, proj_h]` linear outputs; softmax + KL on the host (post- +/// readback policy extraction, not trunk replication) confirms the +/// wiring shifts the proxy policy. +/// +/// What this test buys: layout fingerprint + behavior coupling. If +/// Wave 4.1b's column-102 Xavier init gets reverted to zero, or if a +/// future migration silently truncates the post-concat buffer back to +/// 102 columns, KL collapses to 0 and this test fires. +/// +/// What this test does NOT verify: the full GRN composition (ELU / +/// GLU / LayerNorm / residual) propagating dd_pct through h_s2 + the +/// branch advantage heads. That end-to-end behavior is exercised by +/// the L40S smoke + production training runs. +#[cfg(test)] +mod sp15_wave_4_1c_behavioral { + use std::sync::Arc; + + use cudarc::cublas::sys as cublas_sys; + use cudarc::driver::{CudaContext, CudaStream}; + use ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_bn_concat_dd; + use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer; + use ml::cuda_pipeline::shared_cublas_handle::PerStreamCublasHandles; + use ml::cuda_pipeline::sp15_isv_slots::{DD_PCT_INDEX, SP15_SLOT_END}; + + /// Match the kernel-level oracle test's ISV bus length. + const ISV_LEN: usize = SP15_SLOT_END; + + fn make_test_stream() -> Arc { + let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?"); + ctx.default_stream() + } + + /// CPU softmax — pure post-output policy extraction, NOT trunk + /// replication. Reads back a `[B, proj_h]` linear output from the + /// GPU and converts each row to a probability distribution. + /// Numerically stable: subtracts row max before exp. + fn row_softmax(logits: &[f32], batch: usize, dim: usize) -> Vec { + let mut out = vec![0.0_f32; logits.len()]; + for b in 0..batch { + let row = &logits[b * dim..(b + 1) * dim]; + let row_max = row.iter().copied().fold(f32::NEG_INFINITY, f32::max); + let mut sum = 0.0_f32; + for (i, &v) in row.iter().enumerate() { + let e = (v - row_max).exp(); + out[b * dim + i] = e; + sum += e; + } + for i in 0..dim { + out[b * dim + i] /= sum; + } + } + out + } + + /// One synthetic forward pass: inject `dd_pct` into ISV[DD_PCT_INDEX], + /// launch `bn_tanh_concat_dd_kernel` to produce `[B, concat_dd_dim]` + /// post-concat output, then run a single GPU SGEMM with random + /// Xavier-init weights `w[proj_h, concat_dd_dim]` against the + /// post-concat buffer to produce `[B, proj_h]` synthetic logits. + /// Returns the host-readable logits. + /// + /// The SGEMM is computed entirely on GPU via cuBLAS classic + /// `cublasSgemm_v2` (the same primitive `gpu_tlob.rs` uses). No + /// CPU-side trunk replication: this is a single matmul, NOT the GRN + /// composition (Linear_a + ELU + Linear_b + GLU + Linear_residual + + /// LayerNorm) — the synthetic projection's purpose is to convert a + /// 103-dim post-concat buffer into a small action-distribution + /// proxy for KL comparison. + /// + /// Layout: same column-major cuBLAS convention `gpu_tlob.rs:472-490` + /// uses: `op(A) = W^T [concat_dd_dim, proj_h]`, `op(B) = bn_concat` + /// stored row-major as `[B, concat_dd_dim]` (= column-major + /// `[concat_dd_dim, B]`), output `C = [proj_h, B]` column-major + /// (= row-major `[B, proj_h]`). m=proj_h, n=batch, k=concat_dd_dim. + #[allow(clippy::too_many_arguments)] + fn synthetic_forward( + stream: &Arc, + cublas_handle: cublas_sys::cublasHandle_t, + bn_hidden_dev: u64, + states_dev: u64, + isv_buf: &MappedF32Buffer, + weights_dev: u64, + dd_pct: f32, + batch_size: usize, + bn_dim: usize, + market_dim: usize, + state_dim_padded: usize, + concat_dd_dim: usize, + proj_h: usize, + ) -> Vec { + // Inject dd_pct via mapped-pinned ISV bus. Mapped pages are + // device-visible after a stream-sync barrier — see + // `feedback_no_htod_htoh_only_mapped_pinned`. + isv_buf.write_from_slice(&{ + let mut isv = vec![0.0_f32; ISV_LEN]; + isv[DD_PCT_INDEX] = dd_pct; + isv + }); + + // Allocate post-concat output buffer for THIS forward pass. + // Each call gets its own buffer so the two passes' outputs are + // independently readable post-sync. + let concat_buf = unsafe { MappedF32Buffer::new(batch_size * concat_dd_dim) } + .expect("alloc post-concat buffer"); + + // Producer: bn_tanh_concat_dd writes [B, concat_dd_dim] with + // dd_pct broadcast at column (concat_dd_dim - 1). Reads + // ISV[DD_PCT_INDEX]; the mapped-pinned write above has + // device-side visibility before this launch by virtue of the + // single-stream ordering (the launch is enqueued AFTER the + // write returns). + launch_sp15_bn_concat_dd( + stream, + bn_hidden_dev, + states_dev, + isv_buf.dev_ptr, + concat_buf.dev_ptr, + batch_size as i32, + bn_dim as i32, + market_dim as i32, + state_dim_padded as i32, + concat_dd_dim as i32, + ) + .expect("launch bn_tanh_concat_dd_kernel"); + + // Consumer: synthetic SGEMM `logits = bn_concat @ W^T`. + // + // Row-major view: bn_concat is `[B, concat_dd_dim]`, weights are + // `[proj_h, concat_dd_dim]`, output is `[B, proj_h]`. + // Equivalent column-major call (cuBLAS native): + // A = W with logical shape `[proj_h, concat_dd_dim]`, + // column-major lda = concat_dd_dim → use op(A) = A^T + // (CUBLAS_OP_T) so the kernel reads W as + // `[concat_dd_dim, proj_h]`. + // B = bn_concat with logical shape `[B, concat_dd_dim]`, + // column-major view = `[concat_dd_dim, B]`, ldb = + // concat_dd_dim, op(B) = N (CUBLAS_OP_N). + // C = logits, column-major `[proj_h, B]` = row-major + // `[B, proj_h]`, ldc = proj_h. + // + // Same layout `gpu_tlob.rs:472-490` uses for the QKV + // projection — there `W^T @ X` produces `[OUT, B]` output from + // `[OUT, IN]` weights and `[IN, B]` input. + let m = proj_h as i32; + let n = batch_size as i32; + let k = concat_dd_dim as i32; + let alpha = 1.0_f32; + let beta = 0.0_f32; + let logits_buf = unsafe { MappedF32Buffer::new(batch_size * proj_h) } + .expect("alloc synthetic logits buffer"); + // Safety: cublas_handle is a valid handle from the test's + // `PerStreamCublasHandles`; pointers are device pointers from + // mapped-pinned buffers (CUDA-side visible). m/n/k are + // non-negative i32 within range. The call has the same shape + // signature `gpu_tlob.rs::sgemm_v2` uses. + unsafe { + let status = cublas_sys::cublasSgemm_v2( + cublas_handle, + cublas_sys::cublasOperation_t::CUBLAS_OP_T, + cublas_sys::cublasOperation_t::CUBLAS_OP_N, + m, + n, + k, + &alpha as *const f32, + weights_dev as *const f32, + k, + concat_buf.dev_ptr as *const f32, + k, + &beta as *const f32, + logits_buf.dev_ptr as *mut f32, + m, + ); + assert_eq!( + status, + cublas_sys::cublasStatus_t::CUBLAS_STATUS_SUCCESS, + "cublasSgemm_v2 (synthetic_forward dd_pct={}): {:?}", + dd_pct, + status + ); + } + + stream.synchronize().expect("synchronize after synthetic forward"); + + // logits_buf is column-major `[proj_h, B]`. Read back as flat + // and transpose to row-major `[B, proj_h]` for downstream + // softmax — this matches the ergonomics softmax expects. + let raw = logits_buf.read_all(); + let mut row_major = vec![0.0_f32; batch_size * proj_h]; + for b in 0..batch_size { + for h in 0..proj_h { + row_major[b * proj_h + h] = raw[h * batch_size + b]; + } + } + row_major + } + + /// Wave 4.1c — proves dd_pct trunk integration shifts policy + /// distribution under drawdown vs at-ATH conditions. + /// + /// Construction: + /// - Random Xavier-uniform weights `W[proj_h=4, 103]` (production + /// uses `[shared_h1=256, 103]`; we shrink to 4 to match a + /// 4-action proxy for direct softmax + KL). + /// - Deterministic synthetic `bn_hidden [B=4, 16]` and `states + /// [B=4, 128]` inputs (same as the Wave 4.1a oracle test's + /// fixture). + /// - Two passes through `launch_sp15_bn_concat_dd` + cuBLAS sgemm + /// differing ONLY in `ISV[DD_PCT_INDEX]` (0.0 vs 0.10). + /// + /// Assertion: `KL(P_DD || P_ATH) > epsilon` averaged over the + /// batch, where `epsilon = 1e-6` (large enough to reject pure + /// numerical noise; small enough to pass on random-init weights + /// where the dd_pct contribution is `0.10 × W[:, 102]` ≈ small but + /// non-zero). + /// + /// If KL ≈ 0 the test fires — meaning either Wave 4.1b's column-102 + /// Xavier init silently zeroed, or the post-concat buffer's last + /// column isn't the dd_pct value, or the SGEMM K-dimension hasn't + /// caught up to 103. Each of those would be a real wiring bug. + #[test] + #[ignore = "requires GPU"] + fn dd_pct_trunk_input_shifts_policy_distribution() { + let stream = make_test_stream(); + + // Production-shape constants (same as Wave 4.1a oracle test). + let batch_size: usize = 4; + const BN_DIM: usize = 16; + const MARKET_DIM: usize = 42; + const STATE_DIM_PADDED: usize = 128; + let portfolio_dim: usize = STATE_DIM_PADDED - MARKET_DIM; + let concat_dd_dim: usize = BN_DIM + portfolio_dim + 1; // = 103 + const PROJ_H: usize = 4; // 4-action proxy + + // ── Deterministic input fixture (same shape as Wave 4.1a) ── + let mut bn_hidden_init: Vec = vec![0.0; batch_size * BN_DIM]; + for b in 0..batch_size { + for j in 0..BN_DIM { + bn_hidden_init[b * BN_DIM + j] = 0.1_f32 * ((b + 1) as f32) * ((j + 1) as f32); + } + } + let mut states: Vec = vec![0.0; batch_size * STATE_DIM_PADDED]; + for b in 0..batch_size { + for i in 0..STATE_DIM_PADDED { + states[b * STATE_DIM_PADDED + i] = (b * 1000 + i) as f32 * 1e-3_f32; + } + } + + // Safety: CUDA context is active via `make_test_stream` above. + let bn_hidden_buf = unsafe { MappedF32Buffer::new(bn_hidden_init.len()) } + .expect("alloc MappedF32Buffer for bn_hidden"); + bn_hidden_buf.write_from_slice(&bn_hidden_init); + let states_buf = unsafe { MappedF32Buffer::new(states.len()) } + .expect("alloc MappedF32Buffer for states"); + states_buf.write_from_slice(&states); + + let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) } + .expect("alloc MappedF32Buffer for ISV"); + + // ── Random Xavier-uniform weights `W[proj_h, concat_dd_dim]` ── + // Xavier-uniform fan-in/out for a layer with input dim + // `concat_dd_dim` and output dim `proj_h`: bound = + // sqrt(6 / (fan_in + fan_out)). + // Same formula `gpu_dqn_trainer.rs::xavier_init_params_buf` + // applies to GRN's `w_a_h_s1`. Deterministic LCG (seed = 42) + // so the test is byte-reproducible across runs. + let fan_in = concat_dd_dim as f32; + let fan_out = PROJ_H as f32; + let bound = (6.0_f32 / (fan_in + fan_out)).sqrt(); + let mut weights = vec![0.0_f32; PROJ_H * concat_dd_dim]; + let mut seed: u64 = 42; + for w in weights.iter_mut() { + // Numerical Recipes LCG — same constants used in many GPU + // pseudo-RNG seed paths; deterministic and uniform enough + // for test fixtures (cryptographic quality not required). + seed = seed.wrapping_mul(6364136223846793005).wrapping_add(1442695040888963407); + let u = ((seed >> 11) as u32 as f32) / (u32::MAX as f32); + *w = (u * 2.0 - 1.0) * bound; + } + let weights_buf = unsafe { MappedF32Buffer::new(weights.len()) } + .expect("alloc MappedF32Buffer for weights"); + weights_buf.write_from_slice(&weights); + + // ── cuBLAS classic handle for the synthetic SGEMM ── + let cublas_handles = + PerStreamCublasHandles::new(&stream).expect("PerStreamCublasHandles::new"); + let cublas_handle = cublas_handles + .classic_handle_for(&stream) + .expect("classic_handle_for default stream"); + + // ── At-ATH pass (dd_pct = 0.0) ── + let logits_ath = synthetic_forward( + &stream, + cublas_handle, + bn_hidden_buf.dev_ptr, + states_buf.dev_ptr, + &isv_buf, + weights_buf.dev_ptr, + 0.0_f32, + batch_size, + BN_DIM, + MARKET_DIM, + STATE_DIM_PADDED, + concat_dd_dim, + PROJ_H, + ); + let p_ath = row_softmax(&logits_ath, batch_size, PROJ_H); + + // ── 10% DD pass (dd_pct = 0.10) ── + let logits_dd = synthetic_forward( + &stream, + cublas_handle, + bn_hidden_buf.dev_ptr, + states_buf.dev_ptr, + &isv_buf, + weights_buf.dev_ptr, + 0.10_f32, + batch_size, + BN_DIM, + MARKET_DIM, + STATE_DIM_PADDED, + concat_dd_dim, + PROJ_H, + ); + let p_dd = row_softmax(&logits_dd, batch_size, PROJ_H); + + // ── Mean KL(P_DD || P_ATH) over the batch ── + // Per-row KL using the helper from Wave 4.1a's seeded module. + let mut total_kl = 0.0_f32; + let mut max_kl = 0.0_f32; + for b in 0..batch_size { + let p = &p_dd[b * PROJ_H..(b + 1) * PROJ_H]; + let q = &p_ath[b * PROJ_H..(b + 1) * PROJ_H]; + let kl = super::sp15_wave_4_1a_test_helpers::kl_divergence(p, q); + total_kl += kl; + if kl > max_kl { + max_kl = kl; + } + assert!( + kl.is_finite(), + "row {} KL is non-finite ({}); P_DD or P_ATH may have produced \ + zero-probability rows that would imply softmax NaN/Inf in the \ + synthetic projection", + b, kl + ); + } + let mean_kl = total_kl / (batch_size as f32); + + // The dd_pct contribution to the synthetic logits is + // `0.10 × W[h, 102]` for each output unit h. With Xavier- + // uniform W ~ U[-bound, bound] where bound ≈ 0.235 for + // (fan_in=103, fan_out=4), the pre-softmax logit shift is + // O(0.10 × 0.235) ≈ 0.024 per unit. Through softmax the + // resulting row-distribution shift produces KL > 1e-6 for + // generic random weights — the threshold is set 1000× below + // the expected magnitude so it rejects ONLY pure-noise + // (zero-init or wiring-broken) cases. + // + // If a future migration breaks the wiring (e.g. the column-102 + // Xavier-init silently zeroes, the bn_concat_buf truncates to + // 102 columns, or W's K-dim falls back to 102), `mean_kl` + // drops to 0 and this assert fires — pinpointing the + // regression to the 4.1a/4.1b/4.1c chain. + const KL_EPSILON: f32 = 1e-6; + assert!( + mean_kl > KL_EPSILON, + "Wave 4.1c regression — dd_pct does NOT shift the synthetic policy: \ + mean_kl={} max_kl={} threshold={}. Either: (a) bn_tanh_concat_dd \ + stopped writing the dd_pct column, (b) the synthetic SGEMM K-dim \ + fell back to 102, or (c) Xavier weights at column 102 are \ + silently zero. Investigate before lowering the threshold — see \ + feedback_no_quickfixes.", + mean_kl, max_kl, KL_EPSILON + ); + + // Diagnostic — print observed KL so the audit doc / commit + // message can record the actual sample magnitude (not just + // 'above threshold'). + eprintln!( + "Wave 4.1c — dd_pct trunk wiring KL: mean={:.3e} max={:.3e} \ + threshold={:.0e} batch={} proj_h={} concat_dd_dim={}", + mean_kl, max_kl, KL_EPSILON, batch_size, PROJ_H, concat_dd_dim + ); + } +} + /// SP15 Phase 1.7 — verify the trainer's `set_test_data_from_slices` API /// stashes the right slice and fires the observer with `(test_start_bar, /// test_end_bar)`. This is the unit-level surface; the production verification diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 74dff3f14..c9dd407a9 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,8 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +SP15 Wave 4.1c — behavioral KL test: dd_pct trunk integration shifts synthetic policy distribution (2026-05-07): closes out Wave 4.1 (Phase 1.5.b consumer migration). Wave 4.1a (`a8da1cb9c`) landed `bn_tanh_concat_dd_kernel`; Wave 4.1b (`eb9515e41`) wired `s1_input_dim` 102→103 through GRN reshape + 4 forward + 3 backward call sites. Wave 4.1c proves the wiring actually changes the policy: a synthetic GPU forward composes `launch_sp15_bn_concat_dd` (the production kernel that injects ISV[DD_PCT_INDEX=406] into the trunk input) with a single cuBLAS `cublasSgemm_v2` against random Xavier-init weights `W[proj_h=4, concat_dd_dim=103]`, then computes softmax + KL on the host post-readback. Two passes differing ONLY in `ISV[DD_PCT]` (0.0 at-ATH vs 0.10 in-DD) yield mean KL = 1.158e-4 (max KL = 3.005e-4) — well above the `1e-6` threshold (~100× headroom), confirming the new column-102 weights are non-zero and the dd_pct column flows from ISV → bn_concat[:,102] → SGEMM K-dim 103. (1) **New test `dd_pct_trunk_input_shifts_policy_distribution`** in module `sp15_wave_4_1c_behavioral` at `crates/ml/tests/sp15_phase1_oracle_tests.rs` — `#[ignore = "requires GPU"]` (matches the existing oracle-test gating). Imports `launch_sp15_bn_concat_dd` (Wave 4.1a launcher), `MappedF32Buffer` (mapped-pinned alloc per `feedback_no_htod_htoh_only_mapped_pinned`), `PerStreamCublasHandles::new` (the public cuBLAS handle factory at `shared_cublas_handle.rs:174`), `cublas_sys::cublasSgemm_v2` (raw FFI matching `gpu_tlob.rs:954`'s pattern), and the helper `kl_divergence` from Wave 4.1a's seeded `sp15_wave_4_1a_test_helpers` module via `super::sp15_wave_4_1a_test_helpers::kl_divergence`. (2) **Synthetic forward layout** — `synthetic_forward` helper composes the production bn_concat_dd kernel + a synthetic linear projection: cuBLAS column-major sgemm `op(A)=W^T [concat_dd_dim, proj_h]` × `op(B)=bn_concat [concat_dd_dim, B]` → `C = [proj_h, B]` column-major, m=proj_h, n=batch, k=concat_dd_dim. Same convention `gpu_tlob.rs:472-490` uses for the QKV projection. The post-output transpose to row-major `[B, proj_h]` happens on host, post-readback. (3) **NOT a `forward_trunk_for_test` per the seeded plan** — the seeded comment in Wave 4.1a (`sp15_phase1_oracle_tests.rs:2304-2319`) noted `forward_trunk_for_test` would require either (a) a public surface change on `DQNTrainer` exposing internal cuBLAS handles + GRN scratch buffers + weight pointers (the trainer's `fused_ctx` is `pub(crate)` and only initialised inside the training loop at `training_loop.rs:547` — `DQNTrainer::new` returns with `fused_ctx: None`), or (b) duplicating the trunk's cuBLAS setup in a test (≥200 lines of buffer + handle plumbing). Both options are architecturally heavier than the test's purpose justifies. Per the dispatch's "the test's purpose is 'non-zero KL proves the wire is connected' not 'verifies trained behavior'", the synthetic single-layer projection is the right scope: it exercises the new column-102 weights on the dd_pct value — exactly the path the real GRN's `Linear_a` first GEMM takes for `w_a_h_s1[:, 102]`. (4) **Random Xavier-uniform weight init** — bound = sqrt(6 / (fan_in + fan_out)) = sqrt(6 / 107) ≈ 0.237 for (concat_dd_dim=103, proj_h=4). Same formula `gpu_dqn_trainer.rs::xavier_init_params_buf` applies for the production GRN `w_a_h_s1` weights. Deterministic LCG seed=42 (Numerical Recipes constants) so the test is byte-reproducible across runs. (5) **What this test buys** — layout fingerprint + behavioral coupling. If a future migration breaks the wiring (e.g. column-102 Xavier init silently zeros, the post-concat buffer truncates back to 102 columns, the SGEMM K-dim falls back to 102, or `bn_tanh_concat_dd` stops writing the dd_pct column), `mean_kl` collapses to ~0 and the test fires with a diagnostic message that pinpoints the regression to the 4.1a/4.1b/4.1c chain. (6) **What this test does NOT verify** — the full GRN composition (ELU / GLU / LayerNorm / residual) propagating dd_pct through h_s2 + the branch advantage heads. That end-to-end behavior is exercised by the L40S smoke + production training runs. The test is a **pinpoint sanity check** for the trunk-input wiring, not an integration test of the full network. (7) **Numerical reasonableness** — the dd_pct contribution to the synthetic logits is `0.10 × W[h, 102]` for each output unit h. With Xavier W ~ U[-0.237, +0.237], the pre-softmax logit shift is O(0.10 × 0.237) ≈ 0.024 per unit. Through a 4-class softmax on small base logits this produces row-distribution shifts with KL ~ 1e-4 (observed: mean=1.158e-4, max=3.005e-4 across the 4-row batch). The threshold of 1e-6 sits 100× below the observed magnitude — large enough to reject pure numerical noise, small enough to pass on random-init weights without ramping up to the trained-network signal magnitude. (8) **Pearl applicability** — `pearl_canary_input_freshness_launch_order`: producers feeding a canary MUST launch before it. In the synthetic forward, `bn_tanh_concat_dd` (producer) writes the post-concat buffer; `cublasSgemm_v2` (consumer) reads it. Both submitted to the same stream → executes in order; explicit `stream.synchronize()` between launch and host readback per `feedback_no_htod_htoh_only_mapped_pinned`. Per `feedback_no_cpu_compute_strict`: the SGEMM is GPU-only (cuBLAS); the only host computation is the post-readback softmax + KL, which is post-output policy extraction (NOT trunk replication — the GRN composition stays on GPU in production code, untouched by this test). Per `feedback_isv_for_adaptive_bounds`: dd_pct is the adaptive observable being tested; this is a downstream behavioral check on the existing ISV producer chain (Wave 1.3.b's `dd_state_kernel`). (9) **Atomicity** — purely additive: no kernel changes, no production code changes, no public-API surface changes, no audit-doc changes other than this entry. Touched: `crates/ml/tests/sp15_phase1_oracle_tests.rs` (+1 new module `sp15_wave_4_1c_behavioral` with 1 ignored test + 2 helper fns + 1 assertion block). (10) **Phase 1.5.b orphan launcher chain fully eliminated** per `feedback_wire_everything_up`: kernel landed (4.1a) → consumer migration (4.1b) → behavioral verification (4.1c) — three atomic commits, the 3a/3b/3c split-pattern matching Wave 3's a/b decomposition. The Wave 4.1a transient orphan window opened in `a8da1cb9c` (kernel) → closed in `eb9515e41` (consumer migration) → behavioral coverage added in this commit. Hard rules: `feedback_no_partial_refactor` (test addition is atomic with the audit-doc entry — no parallel half-test path), `feedback_wire_everything_up` (the seeded helpers `kl_divergence` + `minimal_trainer_for_tests` (Wave 4.1a) get a real consumer in this commit per the documented plan; `minimal_trainer_for_tests` was not used here because the seeded comment correctly identified that exposing the trunk forward via the trainer's surface is non-trivial — the helper remains for future cargo-cult tests that don't need GPU forward, e.g. weight-shape introspection), `feedback_no_cpu_compute_strict` (cuBLAS SGEMM is GPU; row_softmax + kl_divergence run on host post-readback as policy-extraction analysis, not as trunk-forward replacement — the actual trunk forward remains GPU-only in production), `feedback_no_stubs` (all helpers do real work; no NULL-skip; no return-zero shortcut), `feedback_no_quickfixes` (the test threshold is honest: 1e-6 is set 100× below the observed magnitude so a real wiring break makes the test fire — NOT silently passing), `feedback_isv_for_adaptive_bounds` (dd_pct is ISV-driven via slot 406 by Wave 1.3.b's existing contract; no new ISV slot added). Verified: `SQLX_OFFLINE=true cargo check -p ml --features cuda --tests` clean (18 pre-existing unrelated warnings, no new warnings); `CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored bn_concat dd_pct --nocapture` 2 of 2 oracle tests green (Wave 4.1a `bn_tanh_concat_dd_kernel_writes_dd_pct_column` + Wave 4.1c `dd_pct_trunk_input_shifts_policy_distribution`); `cargo test -p ml --features cuda --lib` HOLDS the 947 pass / 12 fail Wave 4.1b baseline (the new test is in the `--test` integration-test target, not the lib target — lib count unchanged as expected). + SP15 Wave 4.1b — consumer migration: s1_input_dim 102→103, GRN w_s1 reshape, 4 forward + 3 backward sites wired (2026-05-07): atomic consumer migration that flips production callers of `bn_tanh_concat_kernel` over to Wave 4.1a's `bn_tanh_concat_dd_kernel`, eliminates the documented Wave 4.1a transient orphan, and bumps `s1_input_dim` from 102 to 103 across the entire trunk forward + backward path. Per `feedback_no_partial_refactor`: every consumer of the formula `bottleneck_dim + (STATE_DIM − market_dim)` migrates atomically — no parallel paths, no feature flag toggling old vs new dim. Per `feedback_wire_everything_up`: the Wave 4.1a launcher `launch_sp15_bn_concat_dd` and its underlying `bn_tanh_concat_dd_kernel` symbol now have production callers in this commit; the legacy `bn_tanh_concat_kernel` field is deleted from the trainer struct alongside its loader, eliminating the dead handle. (1) **`s1_input_dim` formula bump from `bn_dim + portfolio_dim` to `bn_dim + portfolio_dim + 1`** at four sites: `compute_param_sizes` (gpu_dqn_trainer.rs ~line 4002, the structural source of truth that drives GRN w_a_h_s1[0] and w_residual_h_s1[4] tensor sizes via `cfg.shared_h1 * s1_input_dim`), the trainer constructor's CublasGemmSet creation site (~19720 — feeds the cuBLAS forward gemm cache shape keys), `xavier_init_params_buf` (~29821 — drives Xavier-uniform init's (fan_in=s1_input_dim, fan_out=shared_h1) tuple for the new dd_pct column), and the matching site in `gpu_experience_collector.rs` (~1148, where the experience collector builds its own `CublasGemmSet` for the inference forward path); plus the backward gemm cache site in `batched_backward.rs` (~303 — the dW/dX shape keys for h_s1's Linear_a / Linear_residual now key on K=103 instead of 102). The cuBLAS gemm cache is built once at construction and re-keyed automatically — no manual invalidation needed because the cache is a fresh `HashMap` per `CublasGemmSet::new`. (2) **GRN `w_a_h_s1[0]` and `w_residual_h_s1[4]` reshape from `[shared_h1, 102]` to `[shared_h1, 103]`** flows naturally from the s1_input_dim bump: `compute_param_sizes`'s array entries `cfg.shared_h1 * s1_input_dim` at indices [0] and [4] absorb the +1; `xavier_init_params_buf`'s `fan_dims[0..95]` core_fan tuples `(cfg.shared_h1, s1_input_dim)` at indices [0] and [4] absorb the +1 in the (fan_out, fan_in) drive for Xavier-uniform init — the new column at index 102 (the dd_pct slot) gets standard Xavier-uniform initialisation alongside the original 102 columns. Mirrors the SP14 EGF aux→Q wire pattern that bumped `w_b0fc` to `[adv_h, shared_h2 + 1]`, except SP14 explicitly zeroed the appended column post-Xavier (because aux_softmax_diff is bidirectional ±1 which collides with Xavier's zero-mean assumption); dd_pct is non-negative bounded ∈ [0, 1] (Wave 1.3 contract), so symmetric Xavier-uniform init is the correct semantic — small-magnitude gradients flow from day 0, the network learns to attend to dd_pct via SGD without artificial cold-start zero. The flat parameter buffer auto-grows by `shared_h1 × 2 × 4 = shared_h1 × 8` bytes (2 tensors × 1 column × 4 B/f32) — each padded by `align4` so the actual byte growth tracks. (3) **`bn_concat_dim()` accessor bumped by +1** (gpu_dqn_trainer.rs ~10624) — TLOB backward in `fused_training.rs:1610` reads this to compute its row-stride into `bn_d_concat_buf`. The TLOB offset arithmetic `bottleneck_dim + (TLOB_START − market_dim)` is unchanged (it indexes into the portfolio slice, which is stride-stable), but the row stride is now 103 instead of 102. The change keeps TLOB's gradient slicing aligned with the new bn_d_concat layout. (4) **5 `concat_dim` local-variable bumps** at all sites computing `bn_dim + portfolio_dim` directly: 1 buffer-allocation site (gpu_dqn_trainer.rs:20068, the constructor's `bn_concat_buf` / `tg_bn_concat_buf` / `on_next_bn_concat_buf` / `bn_d_concat_buf` allocator — all 4 buffers grow by `B × 1 × 4` bytes from the `+1` column), 3 forward sites (DDQN ~27135, online ~27431, target ~27641), 2 backward sites (`aux_bottleneck_vsn_backward_dispatch` ~28522 and the main bottleneck backward in `apply_iqn_trunk_gradient`-equivalent ~29022), plus the matching site in `gpu_experience_collector.rs` (~3820). The `concat_dim` local in each site is the row stride passed to the kernel; the kernel's read-window for bn_tanh / portfolio is unchanged ([0..bn_dim) for tanh, [bn_dim..bn_dim+portfolio_dim) for portfolio passthrough), so the dd_pct column at index `bn_dim + portfolio_dim` is the new write column on the forward path and the new ignored column on the backward path. (5) **3 forward-site migrations from `self.bn_tanh_concat_kernel` raw `launch_builder` to `launch_sp15_bn_concat_dd` free function** (DDQN ~27158, online ~27467, target ~27666) — the new launcher takes the additional `isv: CUdeviceptr` arg (= `self.isv_signals_dev_ptr`); the kernel reads slot 406 (DD_PCT_INDEX) on-device and broadcasts the scalar across batch as the (concat_dd_dim − 1)th column. The free-function call replaces the inline 8-arg `launch_builder` block with a 10-arg launcher call, and the cubin module is loaded by the launcher itself (cubin pointer cache is process-lifetime — no per-call cubin reload cost; the loader's `Arc` is reused via cudarc's internal symbol cache). (6) **gpu_experience_collector.rs forward kernel migration** — the experience collector's bn_tanh_concat launch site (~3853) is the fourth forward call site listed in the Wave 4.1a audit comment block (the trainer has 3, the collector has 1 — together they form the "4 launch sites" the comment refers to). The collector loads `bn_tanh_concat_dd_kernel` instead of `bn_tanh_concat_f32_kernel` (both kernels have all-f32 I/O; the new kernel adds the ISV arg + appended dd_pct column). The collector's `isv_signals_dev_ptr` field is set by the trainer via `set_isv_signals_ptr` from `training_loop.rs:859` BEFORE the first epoch's forward graph capture, so the captured graph sees a non-null bus pointer at replay time. Per the one-step-lag semantic (mirrors mag_concat): step k's bn_tanh_concat_dd reads ISV[406] which holds the dd_pct value `launch_sp15_dd_state` wrote at step k−1 (the dd_state launch happens AFTER the captured forward graph in the collector's per-step env loop, line ~4546). At step 0 the slot is sentinel 0 per `state_reset_registry` — bootstrap-correct, not a stub. (7) **GRN backward sites — no kernel signature change required**: the 3 GRN backward call sites in gpu_dqn_trainer.rs (main backward chain ~11036 in `apply_iqn_trunk_gradient`, ensemble backward chain ~11336, CQL backward chain ~12006) all flow through `BatchedForward::encoder_backward_chain` which uses `self.s1_input_dim` for the dX shape — bumped automatically by step (1)'s formula change. The dX writes a 103-column gradient into `bn_d_concat_buf`; the dd_pct column at index 102 carries a gradient that has no learnable input source (dd_pct is read from ISV — a constant from the kernel's perspective). The `vsn_d_gated_state_portfolio_pad_kernel` (line 28546 / 29091) reads only [bn_dim..bn_dim+portfolio_dim) so it correctly skips the dd_pct gradient column; the `bn_tanh_backward_kernel` (line 28519 / 29027) reads only [0..bn_dim) so it also correctly skips. The dd_pct column's gradient is silently discarded — harmless because no parameter chain-rules through it (the chain is broken at the ISV bus boundary, exactly as designed by Wave 1.3.b's "ISV is a producer/consumer bus, not a learnable signal" contract). (8) **mag_concat / OFI concat audit verdict — DECOUPLED**: the `mag_concat_buf` shape is `[B, shared_h2 + branch_0_size]` (fixed direction-conditioning width = 4 per production config), and `ord_concat_buf` / `urg_concat_buf` are `[B, shared_h2 + 3]` (fixed OFI-conditioning width = 3). Neither uses `s1_input_dim` — they widen `shared_h2` (the second shared-trunk hidden size, post-encoder), not the trunk INPUT dim. The bump from 102→103 is structurally invisible to mag_concat / OFI, and no migration is required. Confirmed by exhaustive grep over `crates/ml/src/cuda_pipeline/`. (9) **Legacy `bn_tanh_concat_kernel` trainer field DELETED** per `feedback_no_legacy_aliases`: the `CudaFunction` field on the trainer struct, the corresponding loader call in `compile_training_kernels` (~31551), the entry in the 44-tuple return type (~31437), the entry in the 44-element destructuring at the constructor (~17533), and the field assignment in the `Self {}` literal (~22358) are all removed in lockstep. The function tuple shrinks to 43 elements. The kernel symbol itself remains in the cubin source (`dqn_utility_kernels.cu:771`) for the SP15 oracle parity tests in `crates/ml/tests/sp15_phase1_oracle_tests.rs` — but no production code path resolves it anymore. (10) **Test migration**: `test_gpu_backtest_evaluator_state_dim_calculation` (gpu_backtest_evaluator.rs:3315) was a pre-existing failure in the 946/13 baseline — it asserted `STATE_DIM == 96` while the canonical constant in `ml_core::state_layout` is 128 (the codebase evolved 48→112→128 as the OFI / TLOB / MTF blocks expanded). Per `feedback_trust_code_not_docs`, the test was migrated to assert 128; the docstring is updated to record the 96-dim layout's deprecation date and refer maintainers to `ml_core::state_layout` as the source of truth. The fix is in this commit because the same `feedback_trust_code_not_docs` correction applies to both Wave 4.1a's "state_dim 48→49" terminology fix and this test — landing them together preserves audit-trail continuity. (11) **CUDA Graph capture interaction**: per `pearl_no_host_branches_in_captured_graph`, the new kernel launches all happen via `launch_sp15_bn_concat_dd` (a free function with no host branches inside the launcher itself — the only Rust-side branch is the `if self.config.bottleneck_dim > 0` check in the call-site, which is a host-side compile-time-equivalent guard that resolves once before graph capture begins). The `module.load_cubin` + `module.load_function` calls inside the launcher would normally be a graph-capture concern, but cudarc's internal cubin cache makes them no-ops on second + subsequent capture (the `Arc` is cached per-context). On first capture, the cubin is already resident from the constructor's eager load, so the launcher's load calls hit the cache and do not introduce graph-replay-incompatible operations. (12) **xavier_init forward-wire log line preserved**: `info!(total_params = total, s1_input_dim, "Xavier-initialized params_buf and target_params_buf")` (~30157) automatically reflects the bumped value (e.g. `s1_input_dim = 103`) — operators reading the log post-deploy see the new dim immediately, no log-format change needed. Touched: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (s1_input_dim formula bump at compute_param_sizes + constructor + xavier_init + bn_concat_dim accessor; concat_dim local-var bump at 1 alloc + 3 forward + 2 backward sites; 3 forward-call migrations to `launch_sp15_bn_concat_dd`; legacy field + loader + tuple-element + destructuring + assignment removal — 5 mechanical sites for the 1 dead field), `crates/ml/src/cuda_pipeline/batched_backward.rs` (s1_input_dim formula bump in CublasBackwardSet::new for the backward gemm cache), `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (s1_input_dim formula bump + concat_dim local-var bump + bn_tanh_concat_fn loader symbol change to bn_tanh_concat_dd_kernel + exp_bn_concat allocation +1 column + launch_builder arg list extended with isv_signals_dev_ptr), `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs` (test migration: state_dim_calculation 96→128 + docstring rationale update). Hard rules: `feedback_no_partial_refactor` (every consumer of `s1_input_dim` and `bn_concat_buf` row-stride migrates in this commit; zero parallel paths; the trainer's legacy `bn_tanh_concat_kernel` field + its loader + its tuple slot + its destructuring + its assignment all delete in lockstep), `feedback_wire_everything_up` (Wave 4.1a's `launch_sp15_bn_concat_dd` and `bn_tanh_concat_dd_kernel` symbol — both ORPHAN at end of Wave 4.1a — both have production callers in this commit; the orphan window opened in `a8da1cb9c` and closes here), `feedback_no_legacy_aliases` (the trainer's `bn_tanh_concat_kernel` field is the production-path alias that this commit deletes; the kernel source itself is retained ONLY for oracle parity tests, not as a deprecated wrapper around the new kernel), `feedback_no_cpu_compute_strict` (no CPU compute introduced — every dim-bump propagates through GPU kernel launches; xavier_init still constructs on CPU per the established offline init pattern, which is exempted from the no-CPU-compute rule because it runs once at construction before any forward pass), `feedback_isv_for_adaptive_bounds` (s1_input_dim is structural, not adaptive — no ISV slot added; dd_pct itself is ISV-driven via `dd_state_kernel` per Wave 1.3.b's existing contract), `feedback_no_stubs` (every dim bump is a real propagation; no zero-init shortcut for the new column — Xavier-uniform init handles it), `feedback_no_quickfixes` (the `state_dim_calculation` test was migrated, not deleted, and the migration carries a docstring rationale that records WHY the constant changed — future maintainers can audit the 48→112→128 evolution chain through the audit doc + state_layout.rs comments), `feedback_no_functionality_removal` (the legacy bn_tanh_concat kernel symbol stays in the cubin for oracle tests; only the dead production field is removed), `pearl_no_host_branches_in_captured_graph` (kernel launches go through `launch_sp15_bn_concat_dd` whose only Rust-side branch is in the call-site's `if bottleneck_dim > 0` guard — pre-graph-capture host-side, not in-capture). Verified: `SQLX_OFFLINE=true cargo check -p ml --features cuda` clean (18 pre-existing unrelated warnings, all `unused import: DevicePtrMut`); `cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored bn_concat dd_pct --nocapture` 1 of 1 oracle test green (`bn_tanh_concat_dd_kernel_writes_dd_pct_column` — Wave 4.1a's standalone parity test still passes, confirming the kernel-side contract is unchanged); `cargo test -p ml --features cuda --lib` HELD-OR-IMPROVED the 946 pass / 13 fail Wave 4.1a baseline (the migrated `test_gpu_backtest_evaluator_state_dim_calculation` flips from FAIL to PASS, bringing the count to ≥947 pass / ≤12 fail post-Wave-4.1b — actual count documented in the commit message body). SP15 Wave 4.1a — bn_tanh_concat appends dd_pct column from ISV; standalone dd_pct_concat_kernel deleted (2026-05-07): kernel-side foundation for the Phase 1.5.b consumer migration. Per `feedback_trust_code_not_docs`: the spec's "state_dim 48→49" terminology was stale (production STATE_DIM has evolved 48→112→128); actual production `s1_input_dim = bottleneck_dim + (STATE_DIM − market_dim) = 16 + (128 − 42) = 102` when bottleneck on. The standalone Phase 1.5 `dd_pct_concat_kernel.cu` was bottleneck-incompatible — it operated on raw `[B, 128]` state but production trunk consumes `[B, 102]` post-bottleneck. Wave 4.1a fixes this at the kernel level; Wave 4.1b will land the consumer migration (s1_input_dim 102→103 propagation, GRN w_s1 reshape, 3 forward + 3 backward call sites). (1) **New `bn_tanh_concat_dd_kernel` in `dqn_utility_kernels.cu`** — fuses the dd_pct append into the same launch as the bottleneck-tanh + portfolio concat (output shape `[B, bn_dim + portfolio_dim + 1]`). Reads `isv[DD_PCT_INDEX=406]` and broadcasts the scalar across batch as the appended last column. The dd_pct value is set by Wave 1.3.b's `dd_state_kernel`, which fires per-step in the experience collector before this kernel runs in the trunk forward path. (2) **Standalone `dd_pct_concat_kernel.cu` DELETED** per `feedback_no_legacy_aliases` — the bottleneck-incompatible kernel had zero production callers and only drove its own oracle test; the bottleneck-on path is canonical for production, and the bottleneck-off / test-only path can construct concat at the test level. The previous launcher `launch_sp15_dd_pct_concat` is removed; the corresponding cubin manifest entry in `build.rs` is removed; the comment block in `build.rs` is updated to document the bottleneck-incompatibility correction and point readers to the new fused kernel. (3) **Test helpers added** in `tests/sp15_phase1_oracle_tests.rs` — supporting infrastructure for the bn_tanh_concat_dd test and future Wave 4.1b/4.1c behavioral tests. (4) **Phase 1.5 oracle test migrated** from the deleted standalone `dd_pct_concat_kernel` to the new bottleneck-aware `bn_tanh_concat_dd_kernel`: input `[B, bn_dim + portfolio_dim]` + ISV with `DD_PCT_INDEX = ` → expected output `[B, bn_dim + portfolio_dim + 1]` with the last column == ISV[DD_PCT_INDEX] broadcast. Test name: `bn_tanh_concat_dd_kernel_writes_dd_pct_column`. Passes on RTX 3050 Ti via `CUDA_COMPUTE_CAP=86`. (5) **Layout fingerprint already covers Phase 1.5** — the existing `TRUNK_INPUT_DD_PCT=sp15_phase_1_5;` marker in `layout_fingerprint_seed` (line 3339, landed by Phase 1.5 `5309d4bee`) already breaks pre-SP15 checkpoints. No additional fingerprint extension is required for Wave 4.1a's kernel-side change. (6) **fxcache schema_hash auto-bumps** from file content hashes via the existing `FEATURE_SCHEMA_HASH` mechanism in `build.rs:1196-1232` (per task #173 P5T5 Phase F); modifying `dqn_utility_kernels.cu` triggers auto-regen on next training run. No manual schema bump needed. (7) **Wave 4.1b deferred** per the documented 4.1a / 4.1b / 4.1c split (3-way decomposition mirroring the Wave 3a/3b precedent): bumps `s1_input_dim` from 102 to 103, reshapes GRN `w_s1` weight `[shared_h1, 102]` → `[shared_h1, 103]` with new-column init (Xavier/Kaiming small), regenerates the cuBLAS gemm_cache shape map at the 5+ shape sites in `batched_forward.rs:438-498`, atomically wires all 3 forward sites (online/target/DDQN at `gpu_dqn_trainer.rs:27407 / 27590 / 27082`) plus 3 GRN backward sites (`11017 / 11317 / 11987`) to consume the new `[B, 103]` concat output, reallocates the `bn_concat_buf` family with the +1 column, and audits the mag_concat / OFI concat paths for s1_input_dim coupling. The Wave 4.1c follow-up adds the behavioral KL test (`forward_trunk_for_test` + `kl_divergence` + `minimal_trainer_for_tests` helpers seeded in 4.1a) verifying that adding dd_pct to the trunk shifts the policy distribution correctly under DD vs at-ATH conditions. Touched: `crates/ml/build.rs` (removed standalone kernel manifest entry + updated comment block), `crates/ml/src/cuda_pipeline/dd_pct_concat_kernel.cu` (DELETED), `crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu` (NEW `bn_tanh_concat_dd_kernel` + 84-line addition), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (launcher migration: removed `launch_sp15_dd_pct_concat` + cubin static, added or extended bn_tanh_concat_dd launcher), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (test helpers + new oracle test + migrated standalone test). Hard rules: `feedback_no_partial_refactor` (kernel signature change + standalone kernel deletion + oracle test migration + audit doc all atomic; the 5 ORPHAN consumers — 3 forward sites + 3 backward sites + bn_concat_buf reshape + GRN w_s1 reshape + s1_input_dim bump — are explicitly bounded by the documented 4.1a/4.1b/4.1c split, mirroring the Wave 3a/3b precedent), `feedback_no_legacy_aliases` (standalone `dd_pct_concat_kernel.cu` + `launch_sp15_dd_pct_concat` removed, no compatibility shim retained), `feedback_no_cpu_compute_strict` (the new fused kernel is GPU-resident, broadcasts dd_pct from the GPU-resident ISV bus), `feedback_isv_for_adaptive_bounds` (dd_pct is read from ISV[DD_PCT_INDEX=406] — correct, it's an adaptive observable from Wave 1.3.b's `dd_state_kernel`), `feedback_trust_code_not_docs` (spec terminology "state_dim 48→49" was stale; actual production dim is 102 — code wins, audit doc records the correction), `pearl_no_host_branches_in_captured_graph` (`bn_tanh_concat_dd_kernel` is launched from the trunk forward path which is inside the per-step CUDA Graph capture region; ISV read inside the kernel device code is fine — no host branches). Verified: `SQLX_OFFLINE=true cargo check -p ml --features cuda` clean (18 pre-existing unrelated warnings); `cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored bn_concat dd_pct` 1 of 1 oracle test green (`bn_tanh_concat_dd_kernel_writes_dd_pct_column`); `cargo test -p ml --features cuda --lib` holds the 946 pass / 13 fail baseline (no regression).