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

feat(dqn-v2): B.1 Flat opp-cost scales with ISV[Q_DIR_ABS_REF] — self-scaling no-tuning

Browse Source 
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-24 22:18:14 +02:00
parent 44539d8f4b
commit 12bba98ecd
4 changed files with 25 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files

23
crates/ml/src/cuda_pipeline/experience_kernels.cu
View File

@@ -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;
}
}

6
crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
View File

@@ -603,8 +603,8 @@ pub struct GpuExperienceCollector {
total_reward_per_sample: CudaSlice<f32>, // [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<f32>, // [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::<f32>(total_output * 6)

6
crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu
View File

@@ -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).

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

@@ -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