From 12bba98ecd0f032f8ceb362aa4d7718a5ef99c25 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 24 Apr 2026 22:18:14 +0200 Subject: [PATCH] =?UTF-8?q?feat(dqn-v2):=20B.1=20Flat=20opp-cost=20scales?= =?UTF-8?q?=20with=20ISV[Q=5FDIR=5FABS=5FREF]=20=E2=80=94=20self-scaling?= =?UTF-8?q?=20no-tuning?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Current Flat opp-cost was conviction-driven but |Q|-scale was implicit. When training drifts Q magnitudes into ±50, opp-cost becomes invisible relative to action Q-values → Flat wins argmax. Fix: multiply by isv_signals_ptr[21] (Q_DIR_ABS_REF_INDEX, EMA of max(|Q_mean|) across direction bins, populated by update_eval_v_range / q_stats_kernel). Self-scaling: opp-cost tracks |Q| proportionally throughout training. No tuned multiplier; relies on existing ISV slot. Floor 1e-3 for cold- start protection before EMA is warm. rc[4] is now written in the Flat branch so reward_component_ema kernel picks it up into ISV[67]. No parallel paths — the old formula IS modified in place. Plan 3 Task 2. Spec §4.B.1. Co-Authored-By: Claude Opus 4.7 (1M context) --- .../src/cuda_pipeline/experience_kernels.cu | 23 +++++++++++++++---- .../cuda_pipeline/gpu_experience_collector.rs | 6 ++--- .../reward_component_ema_kernel.cu | 6 ++--- docs/dqn-wire-up-audit.md | 1 + 4 files changed, 25 insertions(+), 11 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 8db45946c..25b1339e8 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -1412,9 +1412,9 @@ extern "C" __global__ void experience_env_step( * Layout [N*L, 6] row-major; component indices: * [0] popart — final on-policy reward (denominator for PopArt drift) * [1] cf — counterfactual reward written to CF slot - * [2] trail — 0.0 (placeholder; explicit term added in Plan 3 B.1) + * [2] trail — 0.0 (placeholder; future term) * [3] micro — dense OFI micro-reward (R5) - * [4] opp_cost — 0.0 (placeholder; wired in Plan 3 B.1) + * [4] opp_cost — Flat opportunity-cost reward (Plan 3 Task 2 B.1) * [5] bonus — 0.0 (placeholder; wired in Plan 3 B.2/C.4) * Producer: this kernel (experience_env_step). * Consumer: reward_component_ema GPU kernel → ISV[63..69). */ @@ -1499,9 +1499,9 @@ extern "C" __global__ void experience_env_step( float* rc = reward_components_per_sample + out_off * 6; rc[0] = 0.0f; /* popart (final reward) — overwritten below */ rc[1] = 0.0f; /* cf reward — overwritten below at cf_off */ - rc[2] = 0.0f; /* trail penalty (Plan 3 B.1 placeholder) */ + rc[2] = 0.0f; /* trail penalty (future term) */ rc[3] = 0.0f; /* micro reward — overwritten below */ - rc[4] = 0.0f; /* opp_cost (Plan 3 B.1 placeholder) */ + rc[4] = 0.0f; /* opp_cost — overwritten below at Flat branch */ rc[5] = 0.0f; /* bonus (Plan 3 B.2/C.4 placeholder) */ } /* Task 2.X prerequisite defaults — written BEFORE any early-return so every @@ -2306,10 +2306,23 @@ extern "C" __global__ void experience_env_step( float atr_pct_flat = expf(log_atr_flat) / fmaxf(raw_close, 1.0f); float vol_proxy_flat = fmaxf(atr_pct_flat, 0.0001f); if (vol_proxy_flat > 0.01f) vol_proxy_flat = 0.01f; + /* B.1 (Plan 3 Task 2): scale Flat opp-cost by ISV[21] (q_dir_abs_ref + * EMA, max(|Q_mean|) across direction bins). Self-scaling — as Q + * magnitudes drift during training, opp-cost tracks proportionally. + * When Q ≈ ±50, conviction_core keeps the penalty visible relative + * to action Q-values. No tuned multiplier. Populated by + * update_eval_v_range / q_stats_kernel. + * Floor 1e-3 ensures a non-zero penalty during the first few + * epochs before the EMA is warm (cold-start protection). */ + const float q_abs_ref = (isv_signals_ptr != NULL) + ? fmaxf(1e-3f, isv_signals_ptr[21]) + : 1e-3f; /* conviction_core ∈ [0, 1] — ISV-driven per-sample adaptive * multiplier. Replaces the prior 0.5 hardcoded constant. */ reward = -shaping_scale * holding_cost_rate - * conviction_core * vol_proxy_flat; + * conviction_core * vol_proxy_flat * q_abs_ref; + /* C.2 attribution: record opp_cost component. */ + reward_components_per_sample[out_off * 6 + 4] = reward; } } diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 8203b6369..22010c97e 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -603,8 +603,8 @@ pub struct GpuExperienceCollector { total_reward_per_sample: CudaSlice, // [alloc_episodes * alloc_timesteps] /// C.2 Reward-component attribution (Plan 3 Task 1, spec §4.C.2). /// [alloc_episodes * alloc_timesteps * 6] row-major. - /// Components: [0]=popart(total), [1]=cf, [2]=trail(0), [3]=micro, [4]=opp_cost(0), [5]=bonus(0). - /// Producer: experience_env_step (rc[0,1,3] live; rc[2,4,5]=0 placeholder until B.1/B.2/C.4). + /// Components: [0]=popart(total), [1]=cf, [2]=trail(0), [3]=micro, [4]=opp_cost, [5]=bonus(0). + /// Producer: experience_env_step (rc[0,1,3,4] live; rc[2,5]=0 placeholder until B.2/C.4). /// Consumer: reward_component_ema GPU kernel writes ISV[63..69) after each epoch. reward_components_per_sample: CudaSlice, // [alloc_episodes * alloc_timesteps * 6] @@ -1056,7 +1056,7 @@ impl GpuExperienceCollector { .map_err(|e| MLError::ModelError(format!("alloc total_reward_per_sample: {e}")))?; // C.2 Reward-component attribution (Plan 3 Task 1, spec §4.C.2). // [total_output * 6] row-major: [popart, cf, trail, micro, opp_cost, bonus] per (i,t). - // Producer: experience_env_step (writes rc[0], rc[1], rc[3]; rc[2,4,5] = 0.0 placeholder). + // Producer: experience_env_step (writes rc[0], rc[1], rc[3], rc[4]; rc[2,5] = 0.0 placeholder). // Consumer: reward_component_ema GPU kernel → ISV[63..69). let reward_components_per_sample = stream .alloc_zeros::(total_output * 6) diff --git a/crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu b/crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu index 7be7a6c06..8c9c63597 100644 --- a/crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu +++ b/crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu @@ -7,10 +7,10 @@ * * [c=0] popart — final on-policy reward (pre-PopArt; denominator signal) * [c=1] cf — counterfactual replay reward written to the CF slot - * [c=2] trail — 0.0 until Plan 3 B.1 adds an explicit trailing-stop term + * [c=2] trail — 0.0 placeholder (future trailing-stop term) * [c=3] micro — dense OFI-alignment micro-reward (R5) - * [c=4] opp_cost — 0.0 until Plan 3 B.1 wires Flat opportunity cost - * [c=5] bonus — 0.0 until Plan 3 B.2/C.4 wire trade-attempt/timing bonus + * [c=4] opp_cost — Flat opportunity cost (Plan 3 Task 2 B.1; ISV[21]-scaled) + * [c=5] bonus — 0.0 placeholder (future trade-attempt/timing bonus) * * Reduces mean |r_c| over the batch per component, then applies an * adaptive-rate EMA into ISV slots [REWARD_POPART_EMA_INDEX..+6). diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index e8450ab83..2500b8513 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -58,6 +58,7 @@ | `cuda_pipeline/q_quantile_kernel.cu` | `GpuDqnTrainer::launch_q_quantile_reduce` → `FusedTrainingCtx::launch_q_quantile_reduce` → `training_loop.rs` epoch boundary; reads q_out_buf [N, total_actions], writes ISV[50..58) P5/P95 EMAs per branch | Wired | Plan 2 Task 1 C.1 — cold-path (per-epoch); consumer: `update_eval_v_range` reads Q_P05/Q_P95 to set quantile-based half-width | — | | `trainers/dqn/monitors/reward_component_monitor.rs` | Read-only observer for `reward_component_ema` kernel output (ISV slots 63..69); consumers: HEALTH_DIAG `reward_split` line + `controller_activity` smoke fire-rate tracking | Wired | Plan 3 Task 1 C.2 | — | | `cuda_pipeline/reward_component_ema_kernel.cu` | `GpuExperienceCollector::launch_reward_component_ema_inplace` → `training_loop.rs` (called after `reward_contrib_fractions`, before HEALTH_DIAG); reads `reward_components_per_sample [N*L, 6]`, writes ISV[63..69) adaptive EMA (α=0.05) | Wired | Plan 3 Task 1 C.2 — hot-path (per-step); single-block 6-thread kernel | — | +| B.1 Flat opp-cost ISV scaling (`experience_kernels.cu` Flat branch) | `experience_env_step` kernel (Flat position=0, not segment_complete branch); writes `rc[4]` to `reward_components_per_sample`; consumed by `reward_component_ema` → ISV[67] | Wired | Plan 3 Task 2 B.1 — opp-cost multiplied by `isv_signals_ptr[21]` (Q_DIR_ABS_REF_INDEX, EMA of max(|Q_mean|) across direction bins). Self-scaling: as Q magnitudes drift during training, opp-cost tracks proportionally. Floor 1e-3 for cold-start. No tuned multiplier. | — | ## CUDA Pipeline — Rust Wrappers