gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

feat(sp15-wave4.1a): bn_tanh_concat appends dd_pct column from ISV — bottleneck-aware Phase 1.5 consumer migration

Browse Source 
The standalone dd_pct_concat_kernel from Phase 1.5 was bottleneck-
incompatible — it operated on raw [B, 128] state, but production trunk
consumes [B, s1_input_dim] = [B, 102] post-bottleneck. Wave 4.1a fixes
this at the kernel level; Wave 4.1b lands the consumer migration
(s1_input_dim 102→103, GRN w_s1 reshape, 3 forward + 3 backward sites).

Spec correction (per feedback_trust_code_not_docs): the spec's
'state_dim 48→49' is stale terminology pre-STATE_DIM 48→112→128
evolution. Production s1_input_dim is bottleneck_dim + (STATE_DIM −
market_dim) = 16 + (128 − 42) = 102. Wave 4.1b will bump this to 103.

NEW bn_tanh_concat_dd_kernel in dqn_utility_kernels.cu:
  - Fuses dd_pct append into the same launch as bn_tanh + portfolio
    concat (output shape [B, bn_dim + portfolio_dim + 1])
  - Reads isv[DD_PCT_INDEX=406] (set by Wave 1.3.b dd_state_kernel
    per-step), broadcasts the scalar across batch as the appended
    last column

DELETED standalone dd_pct_concat_kernel.cu + launch_sp15_dd_pct_concat
+ cubin manifest entry per feedback_no_legacy_aliases (zero production
callers — only test consumer; bottleneck-on path is canonical).

Test helpers added (used by Wave 4.1c behavioral KL test).
Phase 1.5 oracle test migrated to bn_tanh_concat_dd_kernel contract:
test name bn_tanh_concat_dd_kernel_writes_dd_pct_column passes on
RTX 3050 Ti.

Layout fingerprint already covers Phase 1.5 via the existing
TRUNK_INPUT_DD_PCT=sp15_phase_1_5; marker — pre-SP15 checkpoints
already break.

fxcache schema_hash auto-bumps from file content hashes (per task
P5T5 Phase F mechanism); no manual schema bump needed.

Atomic per feedback_no_partial_refactor for the kernel-signature
contract change. Consumer wiring (s1_input_dim propagation, GRN
reshape, forward/backward call sites) deferred to Wave 4.1b's atomic
commit per the established 3a/3b split precedent — kernel + launcher
land first.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-07 00:06:35 +02:00
parent 4320820ae2
commit a8da1cb9cf
6 changed files with 423 additions and 166 deletions
Download Patch File Download Diff File Expand all files Collapse all files

24
crates/ml/build.rs
View File

@@ -885,19 +885,17 @@ fn main() {
// single-source-of-truth for the action→direction→position
// mapping (mirrors `trade_physics.cuh::decode_direction_4b`).
"position_history_derivation_kernel.cu",
// SP15 Phase 1.5 (2026-05-06): dd_pct foundational state input
// concat. LAYOUT FINGERPRINT BREAK — pre-SP15 checkpoints WILL
// NOT LOAD after this commit. Greenfield OK per spec Q1. Per
// spec §6.5: produces a [B, state_dim_padded + 1] concat buffer
// whose last column is `isv[DD_PCT_INDEX=406]` broadcast across
// batch. Eval-time policy SEES drawdown context on every forward
// pass; Phase 3 teachings can condition on dd_pct directly via
// state, not just reward modulation. Phase 1.5 lands kernel +
// launcher + layout fingerprint marker only; trunk consumer
// migration (re-pointing forward_online to consume the concat
// buffer + bumping s1_input_dim) lands as a follow-up atomic
// commit per the established Phase 1.1-1.4 precedent.
"dd_pct_concat_kernel.cu",
// SP15 Wave 4.1a (2026-05-06): the standalone
// `dd_pct_concat_kernel.cu` from Phase 1.5 was DELETED. The
// standalone kernel was bottleneck-incompatible — it operated on
// raw `[B, state_dim_padded=128]` state but production trunk
// consumes `[B, s1_input_dim=102]` post-bottleneck. Wave 4.1a
// adds a bottleneck-aware extension `bn_tanh_concat_dd_kernel`
// INSIDE `dqn_utility_kernels.cu` that fuses the dd_pct append
// into the same launch as the bn_tanh + portfolio concat
// (output shape `[B, bn_dim + portfolio_dim + 1]`); Wave 4.1b
// lands the consumer migration (s1_input_dim 102→103, GRN w_s1
// reshape, 3 forward + 3 backward call-site updates).
// SP15 Wave 2 (2026-05-06): per-step ISV-driven α producer.
// Replaces the pre-Wave-2 `r_quality_discipline_split_kernel.cu`
// which contained both the scalar composer

60
crates/ml/src/cuda_pipeline/dd_pct_concat_kernel.cu
View File

@@ -1,60 +0,0 @@
// crates/ml/src/cuda_pipeline/dd_pct_concat_kernel.cu
//
// SP15 Phase 1.5 — dd_pct foundational state input concat kernel.
//
// LAYOUT FINGERPRINT BREAK: pre-SP15 checkpoints WILL NOT LOAD after this
// commit. Greenfield OK per spec Q1.
//
// Per spec §6.5: dd_pct (slot 406, written by Task 1.3's dd_state_kernel)
// gets concatenated to the trunk forward input as the last dim. Eval-time
// policy SEES drawdown context on every forward pass; Phase 3 teachings
// can condition on dd_pct directly via state, not just reward modulation.
//
// Kernel contract (Pearl: feedback_no_partial_refactor):
// Input: states_buf [B, state_dim_padded] — existing trunk input
// isv[DD_PCT_INDEX=406] — dd_pct ∈ [0, 1] (broadcast)
// B, state_dim_padded
// Output: concat_buf [B, state_dim_padded + 1] — pre-allocated
// concat_buf[b, 0..state_dim_padded] = states_buf[b, 0..state_dim_padded]
// concat_buf[b, state_dim_padded] = isv[DD_PCT_INDEX]
//
// The kernel produces the concatenated buffer; trunk consumer migration
// (re-pointing forward_online to consume concat_buf instead of states_buf
// and bumping s1_input_dim accordingly) lands as a follow-up atomic
// commit per the established Phase 1 precedent (Tasks 1.1-1.4 all defer
// consumer migration to follow-up commits per feedback_no_partial_refactor).
//
// Layout: row-major [B, state_dim_padded] for input; row-major
// [B, state_dim_padded + 1] for output. The +1 last column is the
// dd_pct broadcast slot. No atomicAdd (pure scatter-copy).
extern "C" __global__ void dd_pct_concat_kernel(
const float* __restrict__ states_buf,
const float* __restrict__ isv,
float* __restrict__ concat_buf,
int batch_size,
int state_dim_padded
) {
const int b = blockIdx.x;
const int tid = threadIdx.x;
const int n_tid = blockDim.x;
if (b >= batch_size) return;
const float dd_pct = isv[406]; // DD_PCT_INDEX (compile-time constant)
const int out_stride = state_dim_padded + 1;
const int in_offset = b * state_dim_padded;
const int out_offset = b * out_stride;
// 1. Copy the existing state vector — full padded width (including
// pre-existing pad zeros from the producer).
for (int i = tid; i < state_dim_padded; i += n_tid) {
concat_buf[out_offset + i] = states_buf[in_offset + i];
}
// 2. Last column = dd_pct broadcast. Single thread writes the +1 slot.
if (tid == 0) {
concat_buf[out_offset + state_dim_padded] = dd_pct;
}
}

84
crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu
View File

@@ -1384,6 +1384,90 @@ extern "C" __global__ void bn_tanh_concat_f32_kernel(
}
}
/* ══════════════════════════════════════════════════════════════════════
* SP15 Wave 4.1a (2026-05-06): BOTTLENECK TANH + CONCAT + DD_PCT KERNEL
*
* Bottleneck-aware extension of `bn_tanh_concat_kernel`. Same forward
* computation (tanh on bn_hidden in-place + portfolio passthrough), with
* one additional output column appended that broadcasts
* `isv[DD_PCT_INDEX=406]` across the batch dimension.
*
* The standalone `dd_pct_concat_kernel.cu` from Phase 1.5 was bottleneck-
* incompatible — it operated on raw `[B, state_dim_padded=128]` state but
* production trunk consumes `[B, s1_input_dim] = [B, bn_dim + portfolio_dim]`
* (= 102 with current bottleneck_dim=16, market_dim=42, STATE_DIM=128) post-
* bottleneck. This kernel folds the dd_pct append into the same launch as
* the bn_tanh + portfolio concat — eval-time policy SEES drawdown context
* via state on every forward pass, with the trunk consuming the dd_pct
* column as the (s1_input_dim+1)th feature once Wave 4.1b lands the
* consumer migration (s1_input_dim 102→103 propagation, GRN w_s1 reshape,
* 3 forward + 3 backward call-site updates).
*
* Kernel contract (mirrors `bn_tanh_concat_kernel` with one extra column):
* Inputs: bn_hidden [B, bn_dim] — IN/OUT: tanh in-place
* states [B, state_dim_padded] — source for portfolio
* isv[DD_PCT_INDEX=406] — dd_pct ∈ [0, 1] broadcast
* Output: concat_out [B, concat_dd_dim]
* where concat_dd_dim = bn_dim + portfolio_dim + 1
* concat_out[b, 0..bn_dim) = tanh(bn_hidden[b, 0..bn_dim))
* concat_out[b, bn_dim..bn_dim+pd) = states[b, market_dim..market_dim+pd)
* concat_out[b, bn_dim+pd] = isv[DD_PCT_INDEX]
*
* Caller must supply `concat_dd_dim = concat_dim + 1` (i.e. the OUTPUT
* total columns, including the dd_pct append). Pure scatter-copy + one
* broadcast column; no atomicAdd; deterministic.
*
* Wave 4.1a transient state: this kernel + its launcher are currently
* ORPHAN — production callers continue to use `bn_tanh_concat_kernel` /
* `bn_tanh_concat_f32_kernel` until Wave 4.1b's atomic refactor flips
* the 4 launch sites + grows `bn_concat_buf` / `tg_bn_concat_buf` /
* `on_next_bn_concat_buf` / experience_collector's `bn_concat` by `+B`
* floats and propagates `s1_input_dim+1` through GRN w_s1 + cuBLAS
* gemm_caches. Acceptable per the documented Wave 3a/3b split precedent
* (kernel + launcher + oracle test verify in isolation first).
*
* Grid: ceil(B * concat_dd_dim / 256), Block: 256.
* ══════════════════════════════════════════════════════════════════════ */
extern "C" __global__ void bn_tanh_concat_dd_kernel(
float* __restrict__ bn_hidden, /* [B, bn_dim] — IN/OUT: tanh applied in-place */
const float* __restrict__ states, /* [B, state_dim_padded] — source for portfolio features */
const float* __restrict__ isv, /* [ISV_LEN] — bus, reads slot 406 (DD_PCT_INDEX) */
float* __restrict__ concat_out, /* [B, concat_dd_dim] — output: [bn_tanh | portfolio | dd_pct] */
int batch_size,
int bn_dim,
int market_dim,
int state_dim_padded,
int concat_dd_dim /* = bn_dim + portfolio_dim + 1 */
) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int total = batch_size * concat_dd_dim;
if (idx >= total) return;
int b = idx / concat_dd_dim;
int j = idx % concat_dd_dim;
/* Last column = isv[DD_PCT_INDEX] broadcast across batch.
* DD_PCT_INDEX=406 is a compile-time constant matching the Rust
* `sp15_isv_slots::DD_PCT_INDEX` and the deleted standalone
* `dd_pct_concat_kernel`'s contract. */
if (j == concat_dd_dim - 1) {
concat_out[idx] = isv[406];
return;
}
if (j < bn_dim) {
/* Bottleneck dimension: apply tanh and copy to output. */
float val = bn_hidden[b * bn_dim + j];
val = tanhf(val);
bn_hidden[b * bn_dim + j] = val;
concat_out[idx] = val;
} else {
/* Portfolio dimension: copy from states[b, market_dim + (j - bn_dim)]. */
int portfolio_idx = market_dim + (j - bn_dim);
concat_out[idx] = states[(long long)b * state_dim_padded + portfolio_idx];
}
}
/* ══════════════════════════════════════════════════════════════════════
* STOCHASTIC DEPTH RNG KERNEL (#21)
*

107
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
View File

@@ -919,78 +919,97 @@ pub fn launch_sp15_dd_state(
Ok(())
}
/// SP15 Phase 1.5 (2026-05-06): cubin for the `dd_pct_concat_kernel`.
/// SP15 Wave 4.1a (2026-05-06): launcher for the bottleneck-aware
/// dd_pct concat extension `bn_tanh_concat_dd_kernel`. Resolves the
/// kernel symbol from the existing `dqn_utility_kernels.cubin` (the
/// kernel was added next to `bn_tanh_concat_kernel` in the same source
/// file) — no separate cubin needed, multi-kernel-per-file pattern
/// matching `baseline_kernels.cu`.
///
/// LAYOUT FINGERPRINT BREAK: pre-SP15 checkpoints WILL NOT LOAD after
/// this commit. Greenfield OK per spec Q1.
/// this commit (same as Phase 1.5 fingerprint marker
/// `TRUNK_INPUT_DD_PCT=sp15_phase_1_5`; the marker stays in
/// `layout_fingerprint_seed`). Greenfield OK per spec Q1.
///
/// Per spec §6.5: produces a `[B, state_dim_padded + 1]` concat buffer
/// whose last column is `isv[DD_PCT_INDEX=406]` broadcast across batch.
/// The kernel reads the existing `[B, state_dim_padded]` states buffer
/// (full padded width, including pre-existing pad zeros from the
/// upstream producer), copies it into the concat buffer's leading
/// `state_dim_padded` columns, and writes `isv[DD_PCT_INDEX]` to the
/// `+1` last column. No atomicAdd (pure scatter-copy).
/// Per spec §6.5 (Wave 4.1a correction): the standalone
/// `dd_pct_concat_kernel.cu` from Phase 1.5 was DELETED — it operated
/// on raw `[B, state_dim_padded=128]` state but production trunk
/// consumes `[B, s1_input_dim] = [B, bn_dim + portfolio_dim] = [B, 102]`
/// post-bottleneck. The bottleneck-aware extension folds the dd_pct
/// append into the same launch as the bn_tanh + portfolio concat;
/// output shape is `[B, concat_dd_dim] = [B, bn_dim + portfolio_dim + 1]`.
///
/// Phase 1.5 lands the kernel + launcher + layout fingerprint marker
/// only; trunk consumer migration (re-pointing `forward_online` to
/// consume the concat buffer + bumping `s1_input_dim`) is deferred to
/// a follow-up atomic commit per `feedback_no_partial_refactor.md` —
/// this matches the established Phase 1.1-1.4 precedent (kernels +
/// launchers verify in isolation first via the GPU oracle test below).
pub static SP15_DD_PCT_CONCAT_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/dd_pct_concat_kernel.cubin"));
/// SP15 Phase 1.5 (2026-05-06): launcher for the dd_pct concat kernel.
/// Free function (matches `launch_sp15_dd_state` and `launch_sp15_*`
/// baseline-launcher precedent) so unit/oracle tests can drive the
/// kernel directly without the trainer struct; production callers
/// invoke the same launcher.
/// Wave 4.1a transient state: this launcher is currently ORPHAN —
/// production callers continue to use `bn_tanh_concat_kernel` /
/// `bn_tanh_concat_f32_kernel` until Wave 4.1b's atomic refactor flips
/// the 4 launch sites + grows `bn_concat_buf` / `tg_bn_concat_buf` /
/// `on_next_bn_concat_buf` / experience_collector's `bn_concat` by
/// `+B` floats and propagates `s1_input_dim+1` through GRN w_s1 +
/// cuBLAS gemm_caches. Acceptable per the documented Wave 3a/3b split
/// precedent (kernel + launcher + oracle test verify in isolation
/// first).
///
/// Contract:
/// `states_buf`: device ptr to `[B, state_dim_padded]` row-major f32.
/// `isv`: device ptr to ISV bus (≥443 f32 slots).
/// `concat_buf`: device ptr to `[B, state_dim_padded + 1]` row-major
/// f32 — caller-allocated, written by the kernel.
/// `batch_size`, `state_dim_padded`: dimension scalars.
/// `bn_hidden`: device ptr to `[B, bn_dim]` f32 — IN/OUT (tanh
/// applied in-place).
/// `states`: device ptr to `[B, state_dim_padded]` f32 — source
/// for portfolio passthrough columns.
/// `isv`: device ptr to ISV bus (≥443 f32 slots) — reads
/// slot 406 (DD_PCT_INDEX).
/// `concat_out`: device ptr to `[B, concat_dd_dim]` f32 — caller-
/// allocated, written by the kernel. Layout:
/// `[bn_tanh | portfolio | dd_pct]`.
/// `batch_size`, `bn_dim`, `market_dim`, `state_dim_padded`,
/// `concat_dd_dim`: dimension scalars. `concat_dd_dim` MUST equal
/// `bn_dim + (state_dim - market_dim) + 1`; the
/// caller (Wave 4.1b) is responsible for sizing the
/// buffer and computing the dim.
///
/// Grid: `batch_size` blocks, 128 threads per block. Each block copies
/// one row of `state_dim_padded` floats and writes the +1 last column
/// from `isv[DD_PCT_INDEX=406]`.
pub fn launch_sp15_dd_pct_concat(
/// Grid: ceil(B * concat_dd_dim / 256), Block: 256.
pub fn launch_sp15_bn_concat_dd(
stream: &Arc<CudaStream>,
states_buf: cudarc::driver::sys::CUdeviceptr,
bn_hidden: cudarc::driver::sys::CUdeviceptr,
states: cudarc::driver::sys::CUdeviceptr,
isv: cudarc::driver::sys::CUdeviceptr,
concat_buf: cudarc::driver::sys::CUdeviceptr,
concat_out: cudarc::driver::sys::CUdeviceptr,
batch_size: i32,
bn_dim: i32,
market_dim: i32,
state_dim_padded: i32,
concat_dd_dim: i32,
) -> Result<(), MLError> {
let module = stream
.context()
.load_cubin(SP15_DD_PCT_CONCAT_CUBIN.to_vec())
.load_cubin(DQN_UTILITY_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!(
"load sp15_dd_pct_concat cubin: {e}"
"load dqn_utility_kernels cubin for bn_tanh_concat_dd: {e}"
)))?;
let kernel = module
.load_function("dd_pct_concat_kernel")
.load_function("bn_tanh_concat_dd_kernel")
.map_err(|e| MLError::ModelError(format!(
"load dd_pct_concat_kernel function: {e}"
"load bn_tanh_concat_dd_kernel function: {e}"
)))?;
let total_elems = (batch_size as i64).saturating_mul(concat_dd_dim as i64);
let blocks = (((total_elems as u64) + 255) / 256).max(1) as u32;
unsafe {
stream
.launch_builder(&kernel)
.arg(&states_buf)
.arg(&bn_hidden)
.arg(&states)
.arg(&isv)
.arg(&concat_buf)
.arg(&concat_out)
.arg(&batch_size)
.arg(&bn_dim)
.arg(&market_dim)
.arg(&state_dim_padded)
.arg(&concat_dd_dim)
.launch(LaunchConfig {
grid_dim: (batch_size.max(1) as u32, 1, 1),
block_dim: (128, 1, 1),
grid_dim: (blocks, 1, 1),
block_dim: (256, 1, 1),
shared_mem_bytes: 0,
})
.map_err(|e| MLError::ModelError(format!(
"launch dd_pct_concat_kernel: {e}"
"launch bn_tanh_concat_dd_kernel: {e}"
)))?;
}
Ok(())
@@ -3336,7 +3355,7 @@ const fn layout_fingerprint_seed() -> &'static [u8] {
DD_TRAJECTORY_DECREASING=439;RECOVERY_OVERSAMPLE_WEIGHT=440;\
DD_TRAJECTORY_FLOOR=441;MEDIAN_STREAK_LENGTH=442;\
ISV_TOTAL_DIM=443;\
TRUNK_INPUT_DD_PCT=sp15_phase_1_5;\
TRUNK_INPUT_DD_PCT=sp15_wave_4_1a;\
PARAM_W_A_H_S1=0;PARAM_B_A_H_S1=1;PARAM_W_B_H_S1=2;PARAM_B_B_H_S1=3;\
PARAM_W_RESIDUAL_H_S1=4;PARAM_GAMMA_H_S1=5;PARAM_BETA_H_S1=6;\
PARAM_W_A_H_S2=7;PARAM_B_A_H_S2=8;PARAM_W_B_H_S2=9;PARAM_B_B_H_S2=10;\

312
crates/ml/tests/sp15_phase1_oracle_tests.rs
View File

@@ -642,101 +642,156 @@ mod gpu {
}
}
/// Test 1.5 — dd_pct foundational state input concat (LAYOUT FINGERPRINT
/// BREAK).
/// Test 1.5 — bottleneck-aware dd_pct trunk-input concat (Wave 4.1a)
/// (LAYOUT FINGERPRINT BREAK — fingerprint marker
/// `TRUNK_INPUT_DD_PCT=sp15_wave_4_1a`).
///
/// Verifies the `dd_pct_concat_kernel` correctly produces a
/// `[B, state_dim_padded + 1]` buffer whose:
/// - leading `state_dim_padded` columns equal the input states_buf
/// (full padded width, including pre-existing pad zeros), and
/// - last column equals `isv[DD_PCT_INDEX=406]` broadcast across
/// batch.
/// Verifies the new `bn_tanh_concat_dd_kernel` correctly produces a
/// `[B, concat_dd_dim]` buffer where
/// `concat_dd_dim = bn_dim + portfolio_dim + 1`. The kernel folds the
/// dd_pct append into the same launch as the bn_tanh + portfolio concat:
/// - columns `[0, bn_dim)` = `tanh(bn_hidden[b, 0..bn_dim))`
/// - columns `[bn_dim, bn_dim+pd)` = `states[b, market_dim..market_dim+pd)`
/// - column `bn_dim + pd` = `isv[DD_PCT_INDEX=406]` (broadcast)
///
/// Spec correction (Wave 4.1a, per `feedback_trust_code_not_docs`): the
/// spec's "state_dim 48→49" is stale terminology pre-STATE_DIM
/// 48→112→128 evolution. Production `s1_input_dim` is
/// `bn_dim + (STATE_DIM - market_dim)` (= 102 with current
/// bn_dim=16, market_dim=42, STATE_DIM=128). Wave 4.1b will bump this
/// to 103 by flipping the 4 launch sites + buffer allocations.
///
/// This is a kernel-level oracle test (matches the Phase 1.1-1.4
/// pattern) — it does NOT verify dd_pct propagates through the trunk
/// forward, because trunk consumer migration is deferred to a
/// follow-up commit per `feedback_no_partial_refactor.md`. The
/// follow-up commit that re-points `forward_online` to consume the
/// concat buffer adds the trunk-grounding behavioral test (see plan
/// §6.5 step 7).
/// pattern) — the trunk-grounding behavioral test (KL divergence
/// between forward_trunk(102-input) and forward_trunk(103-input))
/// is added in Wave 4.1c using the helpers defined at the bottom of
/// this module.
#[test]
#[ignore = "requires GPU"]
fn dd_pct_concat_kernel_writes_last_column() {
use ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_dd_pct_concat;
fn bn_tanh_concat_dd_kernel_writes_dd_pct_column() {
use ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_bn_concat_dd;
let stream = make_test_stream();
// Synthetic batch: B=4, raw state_dim=48, state_dim_padded=128
// (matches GpuDqnTrainConfig::default()'s 48 → pad128 = 128 layout).
// Only the first 48 columns of each row carry "real" feature data;
// the trailing 80 are pre-existing pad zeros from the producer.
// Synthetic batch matches production layout: bn_dim=16,
// market_dim=42, STATE_DIM=128 → state_dim_padded=128 (no
// sub-128 pad in the current production config), portfolio_dim
// = STATE_DIM - market_dim = 86, concat_dd_dim = bn_dim +
// portfolio_dim + 1 = 16 + 86 + 1 = 103. Wave 4.1b bumps the
// production `s1_input_dim` from 102 → 103 (matches the +1
// from the dd_pct append).
let batch_size: usize = 4;
const RAW_STATE_DIM: usize = 48;
let state_dim_padded: usize = 128;
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;
// Build the input states buffer: row b column i = (b * 100 + i) as f32
// for i ∈ [0, RAW_STATE_DIM); zeros for i ∈ [RAW_STATE_DIM,
// state_dim_padded) (the padded portion).
let mut states: Vec<f32> = vec![0.0; batch_size * state_dim_padded];
// Build bn_hidden [B, BN_DIM]: pre-tanh values; row b column j =
// 0.1 * (b + 1) * (j + 1) so each cell is small enough to keep
// tanh in the linear regime (well under 1.0 saturation).
let mut bn_hidden_init: Vec<f32> = vec![0.0; batch_size * BN_DIM];
for b in 0..batch_size {
for i in 0..RAW_STATE_DIM {
states[b * state_dim_padded + i] = (b * 100 + i) as f32;
for j in 0..BN_DIM {
bn_hidden_init[b * BN_DIM + j] = 0.1_f32 * ((b + 1) as f32) * ((j + 1) as f32);
}
}
// Build states [B, STATE_DIM_PADDED]: row b column i =
// (b * 1000 + i) as f32 for the entire padded width — only
// the portfolio slice [MARKET_DIM, MARKET_DIM + portfolio_dim)
// is read by the kernel.
let mut states: Vec<f32> = 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;
}
}
// Safety: CUDA context 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);
// ISV bus: write 0.42 to slot DD_PCT_INDEX=406.
// ISV bus: write 0.05 to slot DD_PCT_INDEX=406.
let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
.expect("alloc MappedF32Buffer for ISV");
let mut isv_init = vec![0.0f32; ISV_LEN];
isv_init[DD_PCT_INDEX] = 0.42;
isv_init[DD_PCT_INDEX] = 0.05;
isv_buf.write_from_slice(&isv_init);
// Output concat buffer: [B, state_dim_padded + 1].
let out_len: usize = batch_size * (state_dim_padded + 1);
let out_len: usize = batch_size * concat_dd_dim;
let concat_buf = unsafe { MappedF32Buffer::new(out_len) }
.expect("alloc MappedF32Buffer for concat output");
launch_sp15_dd_pct_concat(
launch_sp15_bn_concat_dd(
&stream,
bn_hidden_buf.dev_ptr,
states_buf.dev_ptr,
isv_buf.dev_ptr,
concat_buf.dev_ptr,
batch_size as i32,
state_dim_padded as i32,
BN_DIM as i32,
MARKET_DIM as i32,
STATE_DIM_PADDED as i32,
concat_dd_dim as i32,
)
.expect("launch dd_pct_concat_kernel");
.expect("launch bn_tanh_concat_dd_kernel");
stream
.synchronize()
.expect("synchronize after dd_pct_concat_kernel launch");
.expect("synchronize after bn_tanh_concat_dd_kernel launch");
let concat = concat_buf.read_all();
let bn_hidden_after = bn_hidden_buf.read_all();
// Check 1: leading state_dim_padded columns of each row equal the
// input states_buf row (full padded width, including pad zeros).
// Check 1: bn_tanh slice — output[b, 0..BN_DIM) ==
// tanh(input bn_hidden), and the in-place write semantics put
// the same tanh'd value in bn_hidden after the launch.
for b in 0..batch_size {
for i in 0..state_dim_padded {
let in_val = states[b * state_dim_padded + i];
let out_val = concat[b * (state_dim_padded + 1) + i];
for j in 0..BN_DIM {
let pre = bn_hidden_init[b * BN_DIM + j];
let expected = pre.tanh();
let out_val = concat[b * concat_dd_dim + j];
let bn_val = bn_hidden_after[b * BN_DIM + j];
assert!(
(in_val - out_val).abs() < 1e-6,
"concat[b={}, i={}] = {} ≠ states[b={}, i={}] = {} (state copy mismatch)",
b, i, out_val, b, i, in_val
(out_val - expected).abs() < 1e-6,
"concat[b={}, j={}] = {} ≠ tanh({}) = {} (bn_tanh column mismatch)",
b, j, out_val, pre, expected
);
assert!(
(bn_val - expected).abs() < 1e-6,
"bn_hidden[b={}, j={}] = {} ≠ tanh({}) = {} (in-place write mismatch)",
b, j, bn_val, pre, expected
);
}
}
// Check 2: last column of each row equals isv[DD_PCT_INDEX] = 0.42.
// Check 2: portfolio passthrough slice — output[b, BN_DIM..BN_DIM+pd)
// == states[b, MARKET_DIM..MARKET_DIM+pd).
for b in 0..batch_size {
let last = concat[b * (state_dim_padded + 1) + state_dim_padded];
for k in 0..portfolio_dim {
let expected = states[b * STATE_DIM_PADDED + MARKET_DIM + k];
let out_val = concat[b * concat_dd_dim + BN_DIM + k];
assert!(
(out_val - expected).abs() < 1e-6,
"concat[b={}, BN_DIM+{}] = {} ≠ states[b={}, MARKET_DIM+{}] = {} \
(portfolio passthrough mismatch)",
b, k, out_val, b, k, expected
);
}
}
// Check 3: dd_pct broadcast column — output[b, concat_dd_dim - 1]
// == isv[DD_PCT_INDEX] = 0.05 for every batch row.
for b in 0..batch_size {
let last = concat[b * concat_dd_dim + (concat_dd_dim - 1)];
assert!(
(last - 0.42).abs() < 1e-6,
"concat[b={}, last] = {} ≠ isv[DD_PCT_INDEX] = 0.42 (broadcast mismatch)",
b, last
(last - 0.05).abs() < 1e-6,
"concat[b={}, last={}] = {} ≠ isv[DD_PCT_INDEX] = 0.05 (broadcast mismatch)",
b, concat_dd_dim - 1, last
);
}
}
@@ -2243,6 +2298,165 @@ mod gpu {
}
}
/// SP15 Wave 4.1a — shared test helpers used by the bn_tanh_concat_dd
/// oracle test (above) and Wave 4.1c's behavioral KL test (deferred).
///
/// Per the spec, three helpers are needed: `kl_divergence` (CPU math),
/// `minimal_trainer_for_tests` (thin wrapper around `DQNTrainer::new`),
/// and `forward_trunk_for_test` (runs the trunk forward path against
/// a `[B, s1_input_dim+1]` post-bn_concat_dd buffer and returns
/// `[B, shared_h1]` GRN output).
///
/// `forward_trunk_for_test` is NOT implemented in Wave 4.1a — its
/// exact contract is owned by Wave 4.1c (which lands the consumer
/// migration that bumps `s1_input_dim` 102→103 and reshapes GRN
/// w_s1). Implementing it here would either require a new
/// non-trivial public API on `DQNTrainer` (test-only forward
/// surface) or a stub that returns dummy data — both violate
/// `feedback_no_stubs`. The two helpers that ARE implementable
/// today (`kl_divergence`, `minimal_trainer_for_tests`) are
/// real-and-useful; Wave 4.1c will add `forward_trunk_for_test`
/// alongside the public `DQNTrainer` surface change it requires.
#[cfg(test)]
mod sp15_wave_4_1a_test_helpers {
use ml::trainers::dqn::{DQNHyperparameters, DQNTrainer};
/// KL divergence helper: KL(p || q) = sum_i p[i] * ln(p[i] / q[i]).
///
/// Used by Wave 4.1c's behavioral test to compare
/// `softmax(forward_trunk(102-input))` vs
/// `softmax(forward_trunk(103-input))` action distributions; the
/// dd_pct column should produce a non-trivial KL divergence
/// (i.e. the policy actually conditions on drawdown context).
///
/// Contract:
/// - `p` and `q` MUST be the same length (panics otherwise).
/// - Treats both inputs as proper probability distributions:
/// entries with `p[i] == 0.0` contribute 0 to the sum (limit
/// `x ln x → 0` as `x → 0`); entries with `q[i] <= 0.0` raise
/// an `f32::INFINITY` term — caller is responsible for either
/// adding a small epsilon or asserting q is strictly positive.
/// - Result is in nats (natural log). Multiply by `1.0 /
/// std::f32::consts::LN_2` to convert to bits if needed.
///
/// This is a generic helper — no SP15-specific behaviour. Standard
/// information-theoretic KL formula.
#[allow(dead_code)] // Consumer (Wave 4.1c behavioral KL test) lands later.
pub(super) fn kl_divergence(p: &[f32], q: &[f32]) -> f32 {
assert_eq!(
p.len(),
q.len(),
"kl_divergence: distribution length mismatch — p.len()={} q.len()={}",
p.len(),
q.len()
);
let mut acc = 0.0_f32;
for (pi, qi) in p.iter().zip(q.iter()) {
// Convention: 0 * ln(0) = 0; preserves continuity at p=0.
if *pi <= 0.0 {
continue;
}
// q <= 0 with p > 0 → KL is infinite. Surface as +inf rather
// than NaN so caller assertions like `kl > 0.05` see the
// diagnostic value.
if *qi <= 0.0 {
return f32::INFINITY;
}
acc += *pi * ((*pi / *qi).ln());
}
acc
}
/// Construct a minimal `DQNTrainer` for tests that don't need the
/// full data pipeline. Mirrors the existing
/// `set_test_data_from_slices_fires_observer_and_stashes` pattern
/// (see `test_slice_consumption` module below): synchronous
/// constructor, no GPU forward passes required. Returns
/// `Some(trainer)` on success or `None` if the underlying
/// `DQNTrainer::new` fails (e.g. CPU-only build configuration where
/// trainer construction can't proceed) — caller is expected to
/// `eprintln!` a skip notice and return early in that case (the
/// L40S smoke is the canonical verifier when trainer construction
/// is unavailable).
#[allow(dead_code)] // Consumer (Wave 4.1c behavioral KL test) lands later.
pub(super) fn minimal_trainer_for_tests() -> Option<DQNTrainer> {
let hyperparams = DQNHyperparameters::conservative();
match DQNTrainer::new(hyperparams) {
Ok(t) => Some(t),
Err(e) => {
eprintln!(
"minimal_trainer_for_tests: trainer construction failed ({}). \
Caller should skip and rely on L40S smoke verification.",
e
);
None
}
}
}
// Compile-time smoke for the `kl_divergence` helper. Lives in the
// helper module (not as a #[test]) so it doesn't add to the test
// suite; the runtime invariants (zero KL on identical
// distributions, positive KL on differing ones) are exercised by
// the dedicated unit tests below.
#[test]
fn kl_divergence_zero_on_identical_distributions() {
let p = [0.25_f32, 0.25, 0.25, 0.25];
let q = [0.25_f32, 0.25, 0.25, 0.25];
let kl = kl_divergence(&p, &q);
assert!(
kl.abs() < 1e-6,
"KL(p || p) should be 0 for identical distributions, got {}",
kl
);
}
#[test]
fn kl_divergence_positive_on_differing_distributions() {
let p = [0.7_f32, 0.1, 0.1, 0.1];
let q = [0.25_f32, 0.25, 0.25, 0.25];
let kl = kl_divergence(&p, &q);
assert!(
kl > 0.0,
"KL(p || q) should be strictly positive when p ≠ q, got {}",
kl
);
// Hand-checked: 0.7 ln(0.7/0.25) + 3*(0.1 ln(0.1/0.25))
// = 0.7 * 1.0296 + 3 * 0.1 * (-0.9163)
// ≈ 0.7207 - 0.2749 ≈ 0.4458 nats.
assert!(
(kl - 0.4458_f32).abs() < 1e-3,
"KL ≈ 0.4458 nats expected, got {}",
kl
);
}
#[test]
fn kl_divergence_handles_zero_probability_in_p() {
// Convention: p[i] = 0 contributes nothing (limit x ln x → 0).
let p = [0.5_f32, 0.5, 0.0, 0.0];
let q = [0.4_f32, 0.4, 0.1, 0.1];
let kl = kl_divergence(&p, &q);
// Manual check: 0.5 ln(0.5/0.4) + 0.5 ln(0.5/0.4) + 0 + 0
// = 1.0 * ln(1.25) ≈ 0.2231 nats.
assert!(
(kl - 0.2231_f32).abs() < 1e-3,
"KL with zero p entries ≈ 0.2231 nats expected, got {}",
kl
);
}
#[test]
fn minimal_trainer_for_tests_constructs_or_skips() {
// Either the trainer constructs (Some) or we get a clean None
// for the skip path — never panics. Wave 4.1c will use this
// helper and treat None the same way `test_slice_consumption`
// does (eprintln + early return).
let _trainer = minimal_trainer_for_tests();
}
}
/// 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

2
docs/dqn-wire-up-audit.md
View File

File diff suppressed because one or more lines are too long