From 4d4cd996db7b4ee35ca4fc75702227749acca335 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 10 May 2026 20:35:39 +0200 Subject: [PATCH] =?UTF-8?q?feat(sp21):=20T3.3=20ISV=20defrost=20=E2=80=94?= =?UTF-8?q?=20hold=5Freward=5Fema=20(atomic=20Phase=203.2=20wireup)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Closes the Phase 3.2 forward-reference loop in the SP20 aggregator. Previously `out_inputs->per_bar_hold_reward = 0.0f` was hardcoded; the per-bar Hold opportunity-cost producer existed (experience_kernels.cu:3823 — `per_bar_opp_cost = -aux_conf × cost_scale`) and wrote to `hold_baseline_buffer` and `r_micro` directly, but never reached HOLD_REWARD_EMA. Result: HOLD_REWARD_EMA frozen at sentinel 0.0 across all observed epochs; the SP20 reward centering loop (`r_micro += per_bar_opp_cost - HOLD_REWARD_EMA`) stayed uncentered, biasing the policy's reward signal away from zero-mean. T3.3 wireup mirrors the T2.2 alpha refactor: a per-env scratch buffer that the producer writes at every step (alongside the existing `hold_baseline_buffer` write), and the aggregator reads at the same step. Sums opp_cost over Hold-direction envs only; emits the mean as `per_bar_hold_reward`. The HOLD_REWARD_EMA's gate (`hold_fraction > 0.5f` from T3.2) preserves the strict-majority semantic from before. Atomic across producer site, kernel signature, aggregator, launcher, collector, and tests (per feedback_no_partial_refactor): - experience_kernels.cu — new `float* per_bar_opp_cost_per_env` kernel arg, NULL-tolerant write at line ~3823. - sp20_aggregate_inputs_kernel.cu — new arg, 6th shmem stripe (`sh_opp_cost_sum`), per-thread Hold-gated accumulation, tree reduction extended to 6 stripes, output `per_bar_hold_reward = opp_cost_sum / hold_count` when hold_count > 0 else 0. - sp20_aggregate_inputs.rs — launcher signature, dynamic_shmem_bytes 5→6 stripes, doc table, internal test. - gpu_experience_collector.rs — new `per_bar_opp_cost_per_env: CudaSlice` field, alloc, struct construction, kernel-arg pass at both experience_env_step and sp20_aggregate_inputs launches (raw_ptr). - sp20_aggregate_inputs_test.rs — helper renamed `run_kernel_with_is_win_and_opp_cost`, 4 existing sites pass NULL, 2 NEW oracle tests verifying real producer + NULL fallback. - sp20_phase1_4_wireup_test.rs — NULL fallback at the wireup site; HOLD_REWARD_EMA-stays-at-zero assertion remains valid. Verification: - cargo check -p ml --tests: passes (warnings only) - cargo test -p ml --test sp20_aggregate_inputs_test --features cuda -- --ignored: 12/12 GPU oracle tests pass on RTX 3050 Ti, including both new T3.3 tests (per_bar_hold_reward_means_over_hold_envs_only, null_per_bar_opp_cost_emits_zero). - cargo test -p ml --test sp20_phase1_4_wireup_test --features cuda -- --ignored: 2/2 pass under the new NULL-tolerant contract. Plan reference: docs/plans/2026-05-10-sp21-train-eval-coherence-isv-defrost.md Tier 3 status: T3.1 ✓, T3.2 ✓, T3.3 ✓ (this commit), T3.4 withdrawn, T3.5 cascade-pending, T3.6 withdrawn. Co-Authored-By: Claude Opus 4.7 (1M context) --- .../src/cuda_pipeline/experience_kernels.cu | 21 ++++ .../cuda_pipeline/gpu_experience_collector.rs | 43 +++++++ .../cuda_pipeline/sp20_aggregate_inputs.rs | 23 +++- .../sp20_aggregate_inputs_kernel.cu | 61 +++++++-- crates/ml/tests/sp20_aggregate_inputs_test.rs | 118 +++++++++++++++++- crates/ml/tests/sp20_phase1_4_wireup_test.rs | 4 + docs/dqn-wire-up-audit.md | 79 ++++++++++++ ...0-sp21-train-eval-coherence-isv-defrost.md | 4 +- 8 files changed, 334 insertions(+), 19 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 6c259b37f..14a4cc0c4 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -2212,6 +2212,17 @@ extern "C" __global__ void experience_env_step( * sp20_aggregate_inputs_kernel reads this in place of the broken * per-bar `step_ret > 0` predicate. NULL-tolerant for tests. */ int* __restrict__ is_win_per_env, + /* SP21 T3.3 (2026-05-10) — per-env per-bar Hold opportunity-cost. + * `per_bar_opp_cost_per_env` ← [N] f32: written at the per-bar + * Hold opp-cost computation site (line ~3823) alongside the + * existing `hold_baseline_buffer` write. Each slot is + * `-aux_conf × cost_scale ∈ [-0.25, 0]` for the env's current + * step. Consumer: sp20_aggregate_inputs_kernel sums over Hold- + * direction envs and emits the mean as + * `out_inputs->per_bar_hold_reward` → HOLD_REWARD_EMA. NULL- + * tolerant for tests (consumer falls back to 0.0f matching the + * pre-T3.3 Phase 3.2-forward-reference contract). */ + float* __restrict__ per_bar_opp_cost_per_env, /* SP20 Phase 3 Task 3.2 (2026-05-10) — Hold opportunity-cost dual * emission (Component 2) per spec §4.2. * @@ -3822,6 +3833,16 @@ extern "C" __global__ void experience_env_step( : 0.0f; const float per_bar_opp_cost = -aux_conf_i * cost_scale; + /* SP21 T3.3 (2026-05-10): write the per-env per-bar opp_cost + * to the new per_bar_opp_cost_per_env buffer (consumed by + * sp20_aggregate_inputs_kernel for the HOLD_REWARD_EMA + * input). NULL-tolerant — when the buffer isn't wired + * (oracle tests / pre-T3.3 callers), the slot stays + * untouched. Race-free per-i write (one thread per env). */ + if (per_bar_opp_cost_per_env != NULL) { + per_bar_opp_cost_per_env[i] = per_bar_opp_cost; + } + /* Path 2 — always written when buffer wired. */ if (has_hold_buf) { const long long row_off = diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 4ed32ea84..ec1279ba8 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -743,6 +743,17 @@ pub struct GpuExperienceCollector { /// `step_ret > 0` predicate when NULL for test scaffolds). pub(crate) is_win_per_env: CudaSlice, // [alloc_episodes] + /// SP21 T3.3 (2026-05-10): per-env per-bar Hold opportunity-cost + /// scratch. Written every step by `experience_env_step` as + /// `per_bar_opp_cost = -aux_conf × cost_scale ∈ [-0.25, 0]`. + /// Consumer: `sp20_aggregate_inputs_kernel` reads `[env]` slot only + /// for envs in `dir == DIR_HOLD` and emits the mean as + /// `out_inputs->per_bar_hold_reward`. NULL-tolerant for tests + /// (consumer falls back to `0.0f` matching pre-T3.3 Phase 3.2- + /// forward-reference contract). Defrosts HOLD_REWARD_EMA — was + /// stuck at 0 because the aggregator hardcoded 0.0f for this slot. + pub(crate) per_bar_opp_cost_per_env: CudaSlice, // [alloc_episodes] + /// SP20 Phase 3 Task 3.2 (2026-05-10): per-env circular Hold /// opportunity-cost buffer (Component 2) per spec §4.2. /// @@ -1893,6 +1904,19 @@ impl GpuExperienceCollector { "sp20: alloc is_win_per_env ({alloc_episodes} i32): {e}" )))?; + // SP21 T3.3 (2026-05-10): per-env per-bar Hold opportunity-cost + // scratch. Sentinel 0.0; written by experience_env_step every + // step as `per_bar_opp_cost = -aux_conf × cost_scale ∈ [-0.25, 0]`. + // Consumed by sp20_aggregate_inputs_kernel for envs in DIR_HOLD, + // emitted as mean → SP20EmaInputs.per_bar_hold_reward → drives + // HOLD_REWARD_EMA (which was previously hardcoded 0 by the Phase + // 3.2 forward-reference contract). Sized [alloc_episodes]. + let per_bar_opp_cost_per_env = stream + .alloc_zeros::(alloc_episodes) + .map_err(|e| MLError::ModelError(format!( + "sp21: alloc per_bar_opp_cost_per_env ({alloc_episodes} f32): {e}" + )))?; + // SP20 Phase 3 Task 3.2 (2026-05-10): per-env Hold opportunity- // cost circular buffer. Layout `[alloc_episodes × // HOLD_BASELINE_BUFFER_SIZE = 30]` f32 row-major. Sentinel 0.0 @@ -2954,6 +2978,7 @@ impl GpuExperienceCollector { label_at_open_per_env, alpha_per_env, is_win_per_env, + per_bar_opp_cost_per_env, hold_baseline_buffer, states_out, actions_out, @@ -6396,6 +6421,16 @@ impl GpuExperienceCollector { // tick at close bars is dominated by tx_cost deduction // → systematically negative even on winning trades). .arg(&mut self.is_win_per_env) + // SP21 T3.3 (2026-05-10): per-env per-bar Hold + // opportunity-cost scratch buffer. Producer: + // experience_env_step writes per_bar_opp_cost = -aux_conf + // × cost_scale at line ~3823 alongside the existing + // hold_baseline_buffer write. Consumer: + // sp20_aggregate_inputs_kernel sums over Hold-direction + // envs and emits mean → SP20EmaInputs.per_bar_hold_reward + // → HOLD_REWARD_EMA (defrosts the previously-hardcoded + // Phase 3.2 forward-reference 0.0f). + .arg(&mut self.per_bar_opp_cost_per_env) // SP20 Phase 3 Task 3.2 (2026-05-10): Hold // opportunity-cost dual emission (Component 2). Six // new args at the end of the kernel signature: @@ -6733,6 +6768,13 @@ impl GpuExperienceCollector { // `step_ret > 0` predicate that pinned WR_EMA at 0 in // production. let is_win_per_env_dev = self.is_win_per_env.raw_ptr(); + // SP21 T3.3 (2026-05-10): per-env per-bar Hold opp-cost + // pointer. Populated every step by `experience_env_step` + // (launched immediately above on the same stream — stream- + // implicit producer→consumer ordering). Replaces the + // hardcoded `out_inputs->per_bar_hold_reward = 0.0f` + // placeholder atomically per `feedback_no_partial_refactor`. + let per_bar_opp_cost_per_env_dev = self.per_bar_opp_cost_per_env.raw_ptr(); unsafe { launch_sp20_aggregate_inputs( &self.stream, @@ -6747,6 +6789,7 @@ impl GpuExperienceCollector { aux_dir_acc_dev, alpha_per_env_dev, is_win_per_env_dev, + per_bar_opp_cost_per_env_dev, n_envs_i32_sp20, dir_divisor, hold_action_id, diff --git a/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs b/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs index 70841ba0e..6efedecfc 100644 --- a/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs +++ b/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs @@ -18,7 +18,7 @@ //! | `trade_duration` | `round(mean(hold_at_exit) over closed envs)` if `is_close == 1`, else `0` | //! | `hold_fraction` | `hold_count / n_envs ∈ [0, 1]` (SP21 T3.2: was binary `action_is_hold` strict-majority vote) | //! | `alpha` | `mean(alpha_per_env) over closed envs` if `is_close == 1`, else `0` — **SP20 Phase 2 Task 2.2** (R_event - hold_baseline). | -//! | `per_bar_hold_reward` | `0.0` — **Phase 3.2 forward reference** (-aux_conf * cost_scale) | +//! | `per_bar_hold_reward` | `mean(per_bar_opp_cost over Hold-direction envs)` when `hold_count > 0`, else `0.0` (SP21 T3.3, was Phase 3.2 forward reference) | //! | `aux_logits_p50` | `aux_logits_p50_dev[0]` (sp20_stats output) | //! | `aux_logits_std` | `aux_logits_std_dev[0]` (sp20_stats output) | //! | `aux_dir_acc` | `aux_dir_acc_dev[0]` (existing aux_dir_acc_reduce_kernel output [0]) | @@ -67,13 +67,14 @@ pub const SP20_AGGREGATE_BLOCK: u32 = 64; /// Compute the dynamic shared-memory bytes the aggregation kernel -/// needs for the configured block size: 5 stripes × `bdim` × 4 bytes -/// (3 i32 stripes for closed/wins/hold counts + 2 f32 stripes for the -/// hold_at_exit sum and the SP20 Phase 2 Task 2.2 alpha sum). +/// needs for the configured block size: 6 stripes × `bdim` × 4 bytes +/// (3 i32 stripes for closed/wins/hold counts + 3 f32 stripes for the +/// hold_at_exit sum, the SP20 Phase 2 Task 2.2 alpha sum, and the +/// SP21 T3.3 per-bar Hold opp-cost sum). #[must_use] pub fn dynamic_shmem_bytes() -> u32 { SP20_AGGREGATE_BLOCK - .checked_mul(5) + .checked_mul(6) .and_then(|x| x.checked_mul(std::mem::size_of::() as u32)) .expect("SP20 aggregate: shmem bytes overflow") } @@ -144,6 +145,15 @@ pub unsafe fn launch_sp20_aggregate_inputs( // the per-bar `step_ret` at close bars is dominated by tx-cost so // the legacy predicate was essentially always false. is_win_per_env_dev: u64, + // SP21 T3.3 (2026-05-10): per-env per-bar Hold opportunity-cost + // device pointer. `[n_envs]` f32 buffer written every step by + // `experience_env_step` as `per_bar_opp_cost = -aux_conf × cost_scale`. + // The aggregator reads it only for Hold-direction envs and emits + // mean opp-cost when `hold_count > 0`. NULL-tolerant — when `0` + // the kernel emits `out_inputs->per_bar_hold_reward = 0.0f`, + // matching the pre-T3.3 Phase 3.2-forward-reference contract used + // by oracle tests that don't wire the producer. + per_bar_opp_cost_per_env_dev: u64, n_envs: i32, dir_divisor: i32, hold_action_id: i32, @@ -191,6 +201,7 @@ pub unsafe fn launch_sp20_aggregate_inputs( .arg(&aux_dir_acc_dev) .arg(&alpha_per_env_dev) .arg(&is_win_per_env_dev) + .arg(&per_bar_opp_cost_per_env_dev) .arg(&n_envs) .arg(&dir_divisor) .arg(&hold_action_id) @@ -218,7 +229,7 @@ mod tests { fn shmem_bytes_match_layout() { // SP20 Phase 2 Task 2.2: 5 stripes × bdim × 4 bytes (added the // alpha-sum stripe atomically with the per-env alpha producer). - let expected = SP20_AGGREGATE_BLOCK as usize * 5 * 4; + let expected = SP20_AGGREGATE_BLOCK as usize * 6 * 4; assert_eq!(dynamic_shmem_bytes() as usize, expected); } diff --git a/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu b/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu index af388a0da..e45970f37 100644 --- a/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu +++ b/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu @@ -180,6 +180,22 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel( * with the per-env buffer populated by `experience_env_step`'s * segment_complete branch as `(segment_return > 0.0f) ? 1 : 0`. */ const int* __restrict__ is_win_per_env, /* [n_envs] or NULL */ + /* SP21 T3.3 (2026-05-10) — per-env per-bar Hold opportunity-cost + * input. `[n_envs]` f32 device buffer; written every step by + * `experience_env_step` as `per_bar_opp_cost = -aux_conf × cost_scale ∈ + * [-0.25, 0]`. The aggregator reads `per_bar_opp_cost_per_env[env]` + * only for envs in `dir == hold_action_id` (Hold-direction envs) so + * non-Hold envs' opp_costs don't contaminate the average. The + * downstream HOLD_REWARD_EMA's gate (`hold_fraction > 0.5f` at the + * EMA kernel) ensures the EMA fires only when majority Hold — + * preserving the strict-majority semantic of the pre-T3.2 binary + * `action_is_hold` field. + * + * NULL-tolerant — when NULL the kernel emits + * `out_inputs->per_bar_hold_reward = 0.0f` (matches the pre-T3.3 + * Phase 3.2-forward-reference contract used by aggregate-kernel + * oracle tests that don't wire the per-bar producer). */ + const float* __restrict__ per_bar_opp_cost_per_env, /* [n_envs] or NULL */ int n_envs, /* Action decoding parameters (production packed factored action * encoding). `dir_divisor = NUM_MAGNITUDES * NUM_ORD * NUM_URG`, @@ -197,19 +213,20 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel( const int tid = threadIdx.x; const int bdim = blockDim.x; - /* Dynamic shared memory: 5 stripes of bdim × 4 bytes each (4 stripes - * pre-Task-2.2 + 1 stripe for the SP20 Phase 2 Task 2.2 alpha sum). + /* Dynamic shared memory: 6 stripes of bdim × 4 bytes each (5 stripes + * pre-T3.3 + 1 stripe for the SP21 T3.3 per-bar Hold opp-cost sum). * We use `int` storage for the 3 count reductions and `float` storage - * for the hold-at-exit and alpha sums (alias-compatible because both - * are 4 bytes wide; we cast between views as needed). */ + * for the hold-at-exit, alpha, and opp-cost sums (alias-compatible + * because both are 4 bytes wide; we cast between views as needed). */ extern __shared__ int sh_int[]; /* Layout: [closed_count | wins_count | hold_count | - * hold_at_exit_sum_f32 | alpha_sum_f32] */ + * hold_at_exit_sum_f32 | alpha_sum_f32 | opp_cost_sum_f32] */ int* sh_closed_count = &sh_int[0]; int* sh_wins_count = &sh_int[bdim]; int* sh_hold_count = &sh_int[2 * bdim]; float* sh_hold_at_exit_sum = (float*)&sh_int[3 * bdim]; float* sh_alpha_sum = (float*)&sh_int[4 * bdim]; + float* sh_opp_cost_sum = (float*)&sh_int[5 * bdim]; /* Per-thread strided accumulation over [0, n_envs). */ int local_closed_count = 0; @@ -217,6 +234,7 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel( int local_hold_count = 0; float local_hold_at_exit_sum = 0.0f; float local_alpha_sum = 0.0f; + float local_opp_cost_sum = 0.0f; for (int env = tid; env < n_envs; env += bdim) { const int tc = trade_close_per_env[env * env_stride]; @@ -262,7 +280,17 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel( local_alpha_sum += alpha_per_env[env]; } } - if (dir == hold_action_id) local_hold_count += 1; + if (dir == hold_action_id) { + local_hold_count += 1; + /* SP21 T3.3 (2026-05-10): accumulate per-bar Hold opp-cost + * for Hold-direction envs only. NULL-tolerant — when the + * producer isn't wired (oracle tests / pre-T3.3 callers), + * the accumulator stays at 0 → `out_inputs->per_bar_hold_reward + * = 0.0f`, matching the Phase 3.2 forward-reference contract. */ + if (per_bar_opp_cost_per_env != NULL) { + local_opp_cost_sum += per_bar_opp_cost_per_env[env]; + } + } } sh_closed_count[tid] = local_closed_count; @@ -270,10 +298,11 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel( sh_hold_count[tid] = local_hold_count; sh_hold_at_exit_sum[tid] = local_hold_at_exit_sum; sh_alpha_sum[tid] = local_alpha_sum; + sh_opp_cost_sum[tid] = local_opp_cost_sum; __syncthreads(); /* Tree reduction across `bdim`. Standard log2(bdim) pattern. We - * reduce all five stripes in one loop body — single barrier per + * reduce all six stripes in one loop body — single barrier per * pass keeps the cost equal to a single reduction. */ for (int s = bdim / 2; s > 0; s >>= 1) { if (tid < s) { @@ -282,6 +311,7 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel( sh_hold_count[tid] += sh_hold_count[tid + s]; sh_hold_at_exit_sum[tid] += sh_hold_at_exit_sum[tid + s]; sh_alpha_sum[tid] += sh_alpha_sum[tid + s]; + sh_opp_cost_sum[tid] += sh_opp_cost_sum[tid + s]; } __syncthreads(); } @@ -293,6 +323,7 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel( const int hold_count = sh_hold_count[0]; const float hold_at_exit_sum_f = sh_hold_at_exit_sum[0]; const float alpha_sum_f = sh_alpha_sum[0]; + const float opp_cost_sum_f = sh_opp_cost_sum[0]; const int is_close_out = (closed_count > 0) ? 1 : 0; @@ -366,7 +397,21 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel( out_inputs->trade_duration = trade_duration_out; out_inputs->hold_fraction = hold_fraction_out; /* SP21 T3.2 */ out_inputs->alpha = alpha_out; /* Task 2.2 — was 0.0f */ - out_inputs->per_bar_hold_reward = 0.0f; /* Phase 3.2 fwd ref */ + /* SP21 T3.3 (2026-05-10): per_bar_hold_reward = mean per-bar Hold + * opp-cost over Hold-direction envs at this step. Replaces the + * Phase 3.2 forward-reference `0.0f` placeholder. The accumulator + * `opp_cost_sum_f` only summed over envs with `dir == hold_action_id` + * (per-thread gate above), so dividing by `hold_count` gives the + * Hold-conditional mean. When no envs are in Hold OR the producer + * is NULL (oracle tests), the sum is 0 and we emit 0.0f matching + * pre-T3.3 contract. The downstream HOLD_REWARD_EMA's gate + * (`hold_fraction > 0.5f` at the EMA kernel) ensures the EMA fires + * only when majority Hold — preserving strict-majority semantic. */ + float per_bar_hold_reward_out = 0.0f; + if (hold_count > 0) { + per_bar_hold_reward_out = opp_cost_sum_f / (float)hold_count; + } + out_inputs->per_bar_hold_reward = per_bar_hold_reward_out; out_inputs->aux_logits_p50 = p50_in; out_inputs->aux_logits_std = std_in; out_inputs->aux_dir_acc = dir_acc_in; diff --git a/crates/ml/tests/sp20_aggregate_inputs_test.rs b/crates/ml/tests/sp20_aggregate_inputs_test.rs index cb50ad24c..acf263508 100644 --- a/crates/ml/tests/sp20_aggregate_inputs_test.rs +++ b/crates/ml/tests/sp20_aggregate_inputs_test.rs @@ -149,6 +149,7 @@ mod gpu { dir_buf.dev_ptr, /* alpha_per_env_dev = */ 0, // NULL — test scaffold without SP20 reward producer /* is_win_per_env_dev = */ 0, // NULL — falls back to `sr > 0` legacy predicate + /* per_bar_opp_cost_per_env_dev = */ 0, // SP21 T3.3 NULL — emits per_bar_hold_reward=0 /* n_envs = */ n as i32, DIR_DIVISOR, HOLD_ACTION_ID, @@ -218,6 +219,7 @@ mod gpu { dir_buf.dev_ptr, alpha_buf.dev_ptr, // SP20 Phase 2 Task 2.2 alpha producer /* is_win_per_env_dev = */ 0, // NULL — alpha-only variant; legacy `sr > 0` predicate + /* per_bar_opp_cost_per_env_dev = */ 0, // SP21 T3.3 NULL — emits per_bar_hold_reward=0 /* n_envs = */ n as i32, DIR_DIVISOR, HOLD_ACTION_ID, @@ -237,12 +239,17 @@ mod gpu { /// per-bar `step_ret > 0` predicate at close (which is dominated by /// tx-cost so almost always false at close bars, pinning WR_EMA at /// 0 in production). - fn run_kernel_with_is_win( + /// + /// SP21 T3.3 (2026-05-10): also accepts `per_bar_opp_cost_per_env` + /// to exercise the new Hold opportunity-cost producer path. Pass + /// an empty slice to use the NULL fallback (matches pre-T3.3 contract). + fn run_kernel_with_is_win_and_opp_cost( trade_close: &[i32], step_ret: &[f32], hold_at_exit: &[f32], actions: &[i32], is_win_per_env: &[i32], + per_bar_opp_cost_per_env: &[f32], aux_logits_p50: f32, aux_logits_std: f32, aux_dir_acc: f32, @@ -267,6 +274,22 @@ mod gpu { let win_buf = unsafe { MappedI32Buffer::new(n) }.expect("alloc is_win"); win_buf.write_from_slice(is_win_per_env); + // SP21 T3.3 (2026-05-10): per-env per-bar Hold opp-cost producer + // buffer. When the caller passes an empty slice, we pass NULL + // (preserving pre-T3.3 contract). When non-empty, we wire the + // producer so the aggregator emits a non-zero per_bar_hold_reward + // for tests that exercise the new T3.3 path. + let opp_buf_opt = if per_bar_opp_cost_per_env.is_empty() { + None + } else { + assert_eq!(per_bar_opp_cost_per_env.len(), n, + "per_bar_opp_cost_per_env len must match n_envs (or be empty for NULL)"); + let buf = unsafe { MappedF32Buffer::new(n) }.expect("alloc opp_cost"); + buf.write_from_slice(per_bar_opp_cost_per_env); + Some(buf) + }; + let opp_buf_dev = opp_buf_opt.as_ref().map_or(0_u64, |b| b.dev_ptr); + let p50_buf = unsafe { MappedF32Buffer::new(1) }.expect("alloc p50"); p50_buf.write_from_slice(&[aux_logits_p50]); let std_buf = unsafe { MappedF32Buffer::new(1) }.expect("alloc std"); @@ -292,6 +315,7 @@ mod gpu { dir_buf.dev_ptr, /* alpha_per_env_dev = */ 0, win_buf.dev_ptr, // SP20 Phase 2 Task 2.2-fix is_win producer + opp_buf_dev, // SP21 T3.3 — wired iff per_bar_opp_cost_per_env non-empty /* n_envs = */ n as i32, DIR_DIVISOR, HOLD_ACTION_ID, @@ -566,9 +590,10 @@ mod gpu { // and accumulates `wins_count = 3`. let is_win_per_env = vec![1_i32, 1, 1, 0]; - let (ic, iw, _, _, _, _, _, _, _) = run_kernel_with_is_win( + let (ic, iw, _, _, _, _, _, _, _) = run_kernel_with_is_win_and_opp_cost( &trade_close, &step_ret, &hold_at_exit, &actions, - &is_win_per_env, 0.0, 0.0, 0.0, + &is_win_per_env, &[], // SP21 T3.3 NULL: opp-cost producer not exercised here + 0.0, 0.0, 0.0, ); assert_eq!(ic, 1, "is_close = 1 (4 envs closed)"); @@ -615,4 +640,91 @@ mod gpu { fraction not binary majority. got iw={iw}", ); } + + /// SP21 T3.3 (2026-05-10) — per_bar_hold_reward producer wired. + /// + /// Constructs 4 envs, 3 in DIR_HOLD with synthetic per-bar opp-costs + /// `[-0.10, -0.20, -0.05, _]` (1 env in Long has its slot ignored — + /// only Hold-direction envs contribute to the mean). Expected: + /// `per_bar_hold_reward = mean(-0.10, -0.20, -0.05) = -0.1167`. + /// Pins the contract that the aggregator (a) sums opp_cost only + /// over Hold-dir envs, (b) divides by hold_count (NOT n_envs), and + /// (c) emits the mean as a real signal (not 0.0 placeholder). + #[test] + #[ignore = "requires GPU"] + fn per_bar_hold_reward_means_over_hold_envs_only() { + let trade_close = vec![0_i32; 4]; // no closes this step + let step_ret = vec![0.0_f32; 4]; + let hold_at_exit = vec![0.0_f32; 4]; + // 3 Hold envs (DIR_HOLD=1), 1 Long env (DIR_LONG=2). + let actions = vec![ + encode_dir(HOLD_ACTION_ID), + encode_dir(HOLD_ACTION_ID), + encode_dir(HOLD_ACTION_ID), + encode_dir(2), + ]; + let is_win_per_env = vec![0_i32; 4]; // no wins (no closes) + // Hold envs have opp_costs; Long env's slot is read by the + // kernel but the per-thread accumulator is gated on + // `dir == hold_action_id`, so its value (here -1.0 sentinel) + // is correctly excluded from the sum. + let per_bar_opp_cost = vec![-0.10_f32, -0.20, -0.05, -1.0]; + + let (_, _, _, ah, _, hr, _, _, _) = run_kernel_with_is_win_and_opp_cost( + &trade_close, &step_ret, &hold_at_exit, &actions, + &is_win_per_env, &per_bar_opp_cost, + 0.0, 0.0, 0.0, + ); + + // hold_fraction = 3/4 = 0.75 (T3.2 contract). + assert!( + (ah - 0.75).abs() < 1e-6, + "hold_fraction = 3/4 = 0.75 exact; got {ah}", + ); + // per_bar_hold_reward = mean over Hold-dir envs only. + // = (-0.10 + -0.20 + -0.05) / 3 = -0.35 / 3 ≈ -0.11667. + // The Long env's -1.0 slot MUST NOT contribute (gated out). + let expected = (-0.10_f32 - 0.20 - 0.05) / 3.0; + assert!( + (hr - expected).abs() < 1e-6, + "per_bar_hold_reward = mean over Hold-dir envs = {expected:.6}; \ + got {hr}. The Long env's -1.0 sentinel MUST be excluded from \ + the sum (per-thread gate on `dir == hold_action_id`).", + ); + } + + /// SP21 T3.3 (2026-05-10) — NULL producer fallback contract. + /// + /// When `per_bar_opp_cost_per_env` is NULL (oracle scaffolds without + /// the new producer), the aggregator MUST emit + /// `per_bar_hold_reward = 0.0f` regardless of hold_count or any + /// other state. Preserves the pre-T3.3 Phase 3.2 forward-reference + /// placeholder behavior bit-identically for backward compat. + #[test] + #[ignore = "requires GPU"] + fn null_per_bar_opp_cost_emits_zero() { + let trade_close = vec![0_i32; 4]; + let step_ret = vec![0.0_f32; 4]; + let hold_at_exit = vec![0.0_f32; 4]; + // All 4 in Hold — hold_count = 4, hold_fraction = 1.0. + let actions = vec![encode_dir(HOLD_ACTION_ID); 4]; + let is_win_per_env = vec![0_i32; 4]; + + let (_, _, _, ah, _, hr, _, _, _) = run_kernel_with_is_win_and_opp_cost( + &trade_close, &step_ret, &hold_at_exit, &actions, + &is_win_per_env, &[], // NULL — no opp-cost producer wired + 0.0, 0.0, 0.0, + ); + + assert!( + (ah - 1.0).abs() < 1e-6, + "hold_fraction = 4/4 = 1.0; got {ah}", + ); + assert_eq!( + hr, 0.0, + "NULL per_bar_opp_cost_per_env MUST emit per_bar_hold_reward=0.0 \ + (pre-T3.3 Phase 3.2 forward-reference contract preserved); \ + got {hr}", + ); + } } diff --git a/crates/ml/tests/sp20_phase1_4_wireup_test.rs b/crates/ml/tests/sp20_phase1_4_wireup_test.rs index 9554b16ef..938c89e70 100644 --- a/crates/ml/tests/sp20_phase1_4_wireup_test.rs +++ b/crates/ml/tests/sp20_phase1_4_wireup_test.rs @@ -178,6 +178,10 @@ mod gpu { // (e.g. WR_EMA = 1.0 with one // winning trade at step_ret > 0) // remain valid. + /* per_bar_opp_cost_per_env_dev = */ 0, // SP21 T3.3 — wireup test stays + // NULL; consumer falls back to + // per_bar_hold_reward = 0.0f + // matching pre-T3.3 contract. n_envs as i32, DIR_DIVISOR, HOLD_ACTION_ID, diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 6503cd69f..0113d2bed 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,85 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +## 2026-05-10 — SP21 T3.3: hold_reward_ema defrost (atomic Phase 3.2 producer wireup) + +**Branch:** `sp20-aux-h-fixed` (off `1790a31b6`). +**Commits:** 1 atomic (this commit). + +Closes the Phase 3.2 forward-reference loop in the SP20 aggregator. +Previously `out_inputs->per_bar_hold_reward = 0.0f` was hardcoded as a +deferred Phase 3.2 placeholder; the per-bar Hold opportunity-cost +producer existed (`experience_kernels.cu:3823` — `per_bar_opp_cost = +-aux_conf × cost_scale`) and wrote to `hold_baseline_buffer` and +`r_micro` directly, but never reached HOLD_REWARD_EMA. Result: +HOLD_REWARD_EMA frozen at sentinel 0.0 across all observed epochs, +the SP20 reward centering loop (`r_micro += per_bar_opp_cost - +HOLD_REWARD_EMA`) stayed uncentered, biasing the policy's reward +signal away from zero-mean. + +**T3.3 wireup mirrors the T2.2 alpha refactor**: a per-env scratch +buffer that the producer writes at every step (alongside the existing +`hold_baseline_buffer` write), and the aggregator reads at the same +step. Sums opp_cost over Hold-direction envs only; emits the mean as +`per_bar_hold_reward`. The HOLD_REWARD_EMA's gate (`hold_fraction > +0.5f` from T3.2) preserves the strict-majority semantic from before +the binary→fractional refactor. + +**Affected files (atomic per `feedback_no_partial_refactor`):** + +- `crates/ml/src/cuda_pipeline/experience_kernels.cu` — new + `float* per_bar_opp_cost_per_env` kernel arg; per-env producer + write at the existing per-bar opp-cost computation site (line + ~3823) alongside the `hold_baseline_buffer` write. +- `crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu` — + new kernel arg, 6th shmem stripe (`sh_opp_cost_sum`), per-thread + Hold-gated accumulation (`if (dir == hold_action_id) + local_opp_cost_sum += per_bar_opp_cost_per_env[env]`), tree + reduction extended to 6 stripes, output computation + (`per_bar_hold_reward = opp_cost_sum / hold_count`). +- `crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs` — launcher + signature accepts `per_bar_opp_cost_per_env_dev: u64` (NULL = `0`), + `dynamic_shmem_bytes()` 5→6 stripes, doc table, internal test. +- `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` — new + `pub(crate) per_bar_opp_cost_per_env: CudaSlice` field, alloc + in constructor, struct-construction line, kernel-arg pass at both + the experience_env_step launch and the sp20_aggregate_inputs + launch (raw_ptr). +- `crates/ml/tests/sp20_aggregate_inputs_test.rs` — `run_kernel_with_ + is_win` renamed `run_kernel_with_is_win_and_opp_cost`; 4 existing + call sites pass NULL; **2 new GPU oracle tests** verifying + (a) `per_bar_hold_reward = mean over Hold-dir envs` with synthetic + opp_costs and a Long env whose slot must be excluded; (b) NULL + producer fallback emits 0.0 preserving the pre-T3.3 contract. +- `crates/ml/tests/sp20_phase1_4_wireup_test.rs` — NULL fallback at + the wireup test's `launch_sp20_aggregate_inputs` site (the + `alpha_and_hold_reward_emas_stay_at_zero_with_null_producers` + assertion remains valid). + +**Verification:** + +- `cargo check -p ml --tests` — passes (warnings only). +- `cargo test -p ml --test sp20_aggregate_inputs_test --features cuda + -- --ignored` — **12/12 GPU oracle tests pass on RTX 3050 Ti**, + including both new T3.3 tests (`per_bar_hold_reward_means_over_ + hold_envs_only`, `null_per_bar_opp_cost_emits_zero`). +- `cargo test -p ml --test sp20_phase1_4_wireup_test --features cuda + -- --ignored` — 2/2 pass; HOLD_REWARD_EMA assertion holds under + the new NULL-tolerant contract. + +**Pearls applied:** + +- `pearl_first_observation_bootstrap` — HOLD_REWARD_EMA's existing + bootstrap path (sentinel 0.0 → first observation) still applies; + T3.3 only changes WHAT the observation is (real opp_cost mean + instead of hardcoded 0). +- `feedback_no_partial_refactor` — kernel signature, launcher, + collector, tests all migrated atomically in this single commit. +- `pearl_blend_formulas_must_have_permanent_floor` — the consumer + side at `experience_kernels.cu:3838` (`r_micro += per_bar_opp_cost + - HOLD_REWARD_EMA`) is now correctly centered; previously the + EMA=0 made the centering a no-op. + ## 2026-05-10 — SP21 Tier-3 foundation: T3.1+T3.2 ISV defrost (wr_ema, hold_pct_ema) **Branch:** `sp20-aux-h-fixed` (off `6df44e4c6`). diff --git a/docs/plans/2026-05-10-sp21-train-eval-coherence-isv-defrost.md b/docs/plans/2026-05-10-sp21-train-eval-coherence-isv-defrost.md index 5a4039463..56c4ceebd 100644 --- a/docs/plans/2026-05-10-sp21-train-eval-coherence-isv-defrost.md +++ b/docs/plans/2026-05-10-sp21-train-eval-coherence-isv-defrost.md @@ -132,9 +132,9 @@ Each ISV slot below stayed constant across all 3 observed epochs of `d7bj7`. The | T3.1 | `wr_ema` | 0.0000 (frozen) | EMA of segment-close win rate ∈ [0, 1] | `pearl_per_bar_vs_segment_pnl_signal_mismatch` patch added `is_win_per_env` buffer; verify producer is reading it AND writing to `wr_ema_idx` | | T3.2 | `hold_pct_ema` | 0.0000 (frozen) | EMA of fraction of bars where policy holds an active position | Likely an EMA producer reading from action_counts or position_active_mask | | T3.3 | `hold_reward_ema` | 0.0000 (frozen) | EMA of opportunity-cost reward magnitude per bar | Likely populated alongside hold_pct | -| T3.4 | `loss_cap` | -1.0000 (sentinel) | Should be replaced with adaptive value after first observation per Pearl-A | First-observation bootstrap kernel — sentinel-detection guard may have wrong condition | +| ~~T3.4~~ | ~~`loss_cap`~~ | ~~-1.0000 (sentinel)~~ | **WITHDRAWN 2026-05-10 (not a bug)**: investigation shows the value `-1.0` is NOT a sentinel — it's the LOSS_CAP controller's correct output for `WR_EMA < 0.50` (`-1.0 - clamp((wr_ema-0.50)/0.05, 0, 1)`). The actual sentinel per `state_reset_registry` is `0.0` (FoldReset bootstrap). After T3.1 fix, WR_EMA moves to actual val WR ≈ 0.46, but LOSS_CAP stays at -1.0 because WR is still below the 0.50 ramp threshold. **Will move to -2.0 only when WR_EMA crosses 0.55** (linear ramp in [0.50, 0.55]). GPU test `loss_cap_ramp_boundaries` verified the controller logic is correct (passed in our run). | **No fix needed at this slot.** Same shape as T3.6 — downstream of upstream WR signal being below the controller's design threshold. Will defrost automatically when WR climbs above 0.50. | | T3.5 | `hold_cost_scale` | 0.0100 (at min floor) | Controller-driven within [min, max] bounds | Controller reads `hold_pct_ema` (which is also frozen at 0) — chain failure: T3.2 must fix first | -| T3.6 | `aux_conf_threshold` | 0.0100 (at min floor) | Controller drives based on aux confidence distribution | Producer kernel reads aux confidence histogram; check if histogram producer is firing | +| ~~T3.6~~ | ~~`aux_conf_threshold`~~ | ~~0.0100 (at min floor)~~ | **WITHDRAWN 2026-05-10 (not a bug)**: investigation shows the controller at `sp20_controllers_compute_kernel.cu:152-157` (`clamp(aux_dir_acc_ema - 0.50, 0.01, 0.20)`) is *working correctly*. With observed `aux_dir_acc ≈ 0.46` (below random in d7bj7 logs), raw = -0.04, controller clamps at floor 0.01. Chain `aux_dir_acc_reduce_kernel → aggregator → EMA blend → controller` is intact and the EMA tracks the actual ~0.46 value. The slot will defrost on its own when aux improves above 0.51. **Downstream symptom of the upstream aux signal problem, NOT an independent producer-chain bug.** | **No fix needed at this slot.** Upstream problem (aux_dir_acc < 0.51) addressed by upstream feature/architecture work — separate from SP21 scope. | **Strategy**: T3.1 fixed first as canonical example (it has a recent pearl + recent commit `64bbbe418` claiming the fix). Verify the patch actually propagated to the ISV slot, not just to `is_win_per_env` buffer. Once T3.1 is verified working, use the same pattern (find producer → check input availability → check write-to-isv → check sentinel guard) for T3.2-T3.6 in parallel.