diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index cd93df9da..abd9b2dc9 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -1119,22 +1119,48 @@ extern "C" __global__ void experience_env_step( float book_aggression_weight, /* micro-reward: book aggression weight (default 0.3) */ float hold_quality_weight, /* micro-reward: retrospective hold quality weight (default 0.2) */ float micro_reward_temp, /* micro-reward: tanh temperature (default 3.0) */ - float cf_ratio /* D4/N4: adaptive counterfactual flip ratio [0,1]; 0.5=healthy, 0.8=collapsed */ + float cf_ratio, /* D4/N4: adaptive counterfactual flip ratio [0,1]; 0.5=healthy, 0.8=collapsed */ + /* ── Phase 3 unified-env control scalars (pinned device-mapped, host-writable) ── + * exploration_scale ∈ [0, 1]: gates training-only stochastic perturbations + * (saboteur cost noise, plan_params position scaling, counterfactual flips). + * 1.0 = full training mode; 0.0 = deterministic / validation-equivalent. + * Saboteur multipliers lerp toward 1.0: sab_eff = 1 + exp_s * (sab - 1). + * shaping_scale ∈ [0, 1]: gates additive reward-shaping bundles + * (drawdown_penalty, churn_penalty, micro_reward, asymmetric clamp on segment return). + * 1.0 = full shaping; 0.0 = pure per-bar P&L (raw_returns_out is always pure). + * Both pointers may be NULL — defaults to 1.0 for backward compatibility. */ + const float* __restrict__ exploration_scale_ptr, + const float* __restrict__ shaping_scale_ptr ) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= N) return; - /* #33 Per-episode saboteur overrides */ + /* ── Phase 3 unified-env control scalars ── + * Read once per thread. NULL-tolerant for backward compatibility (defaults + * to 1.0 = full training mode). Validation flips both to 0.0 to disable + * saboteur, plan-params nudge, counterfactuals, and reward-shaping bundles + * — leaving pure trade physics + raw P&L. */ + float exploration_scale = (exploration_scale_ptr != NULL) ? exploration_scale_ptr[0] : 1.0f; + float shaping_scale = (shaping_scale_ptr != NULL) ? shaping_scale_ptr[0] : 1.0f; + /* Defensive clamp — host writers may not enforce [0,1]. */ + exploration_scale = fminf(1.0f, fmaxf(0.0f, exploration_scale)); + shaping_scale = fminf(1.0f, fmaxf(0.0f, shaping_scale)); + + /* #33 Per-episode saboteur overrides — lerped by exploration_scale. + * sab_eff = 1 + exploration_scale * (sab - 1) → identity at scale=0, + * full saboteur at scale=1. Same lerp form for both spread and slippage. */ if (saboteur_params != NULL) { float sab_spread = saboteur_params[i * 3 + 0]; float sab_fill = saboteur_params[i * 3 + 1]; float sab_slip = saboteur_params[i * 3 + 2]; - spread_cost *= sab_spread; - tx_cost_multiplier *= sab_slip; + float eff_spread = 1.0f + exploration_scale * (sab_spread - 1.0f); + float eff_slip = 1.0f + exploration_scale * (sab_slip - 1.0f); + spread_cost *= eff_spread; + tx_cost_multiplier *= eff_slip; /* fill_ioc_fill_prob is not a kernel param — it's baked into the - * portfolio sim via trade_physics.cuh. We can't override it per-episode - * without adding it as a kernel parameter. For now, the saboteur affects - * spread and slippage (the two most impactful microstructure levers). */ + * portfolio sim via trade_physics.cuh. The exploration_scale gating + * here covers spread + slippage; fill probability stays at base. */ + (void)sab_fill; /* still read for register layout consistency */ } /* G2: Transaction cost curriculum — scale spread_cost and tx_cost_multiplier */ @@ -1162,11 +1188,16 @@ extern "C" __global__ void experience_env_step( /* G7: Counterfactual return augmentation — 50% flip direction. * Doubles effective dataset, forces symmetric strategies. * Negate directional features (indices 0-15) in the REPLAY BUFFER copy - * so the original batch_states remain untouched for portfolio sim. */ + * so the original batch_states remain untouched for portfolio sim. + * + * Gated by exploration_scale: at 0.0 the effective flip probability is 0 + * (validation/deterministic mode), at 1.0 the configured cf_ratio applies. + * Linear blend: effective_cf = exploration_scale * cf_ratio. */ int do_flip = 0; { float flip_u = philox_uniform(i, bar_idx + 8888, 7777); - do_flip = (flip_u < cf_ratio) ? 1 : 0; + float effective_cf = exploration_scale * cf_ratio; + do_flip = (flip_u < effective_cf) ? 1 : 0; } if (do_flip) { for (int f = 0; f < 16 && f < state_dim; f++) { @@ -1446,9 +1477,15 @@ extern "C" __global__ void experience_env_step( /* ── Trade Plan: readiness-gated activation + enforcement ── * Plan head is randomly initialized → garbage plans early in training. * Gate by model readiness: readiness < 0.5 → plan disabled (model trades freely). - * readiness ≥ 0.5 → plan activates and enforces. Smooth transition. */ + * readiness ≥ 0.5 → plan activates and enforces. Smooth transition. + * + * Phase 3 unified-env: also gated by exploration_scale > 0.5. Validation + * runs with exploration_scale=0 → plan disabled → position is the raw + * Kelly/margin-capped target, matching the backtest kernel. */ float readiness = (readiness_ptr != NULL) ? readiness_ptr[0] : 1.0f; - int plan_enabled = (readiness >= 0.5f && plan_params_ptr != NULL); + int plan_enabled = (readiness >= 0.5f + && plan_params_ptr != NULL + && exploration_scale >= 0.5f); /* Plan activation on Flat → Positioned (only when model is ready) */ int was_flat_plan = (fabsf(ps0_f) < 0.001f); @@ -1796,18 +1833,22 @@ extern "C" __global__ void experience_env_step( + w_book * sign_pos * book_aggression + w_hold * hold_quality - spread_penalty; - reward = micro_reward_scale * tanhf(quality / micro_reward_temp); + reward = shaping_scale * micro_reward_scale * tanhf(quality / micro_reward_temp); } else if (!segment_complete && fabsf(position) > 0.001f) { /* Fallback: original holding cost when micro_reward_scale=0 or no OFI */ - reward = -holding_cost_rate * fabsf(position); + reward = -shaping_scale * holding_cost_rate * fabsf(position); } /* Flat (no position, no trade): reward stays 0.0f from initialization. * No signal is correct — the model should not be rewarded or penalized * for correctly staying flat when there's no edge. */ - /* ---- Drawdown penalty (from trade_physics.cuh) ---- */ + /* ---- Drawdown penalty (from trade_physics.cuh) ────────────────────── + * Phase 3: gated by shaping_scale. The drawdown penalty is *behavioral* + * (steers the policy away from large equity drops); the *physics* of + * drawdown impact is already in the per-bar P&L. At shaping_scale=0 + * (validation mode) we measure pure equity-change reward only. */ float equity_now = cash + position * raw_close; - reward += compute_drawdown_penalty(equity_now, peak_equity, dd_threshold, w_dd); + reward += shaping_scale * compute_drawdown_penalty(equity_now, peak_equity, dd_threshold, w_dd); /* ---- Portfolio value for floor check ---- */ float new_portfolio_value_pre_floor = cash + position * raw_close; @@ -1909,16 +1950,18 @@ extern "C" __global__ void experience_env_step( * No hardcoded gates — the model sees costs and learns to minimize them. */ /* Inventory penalty: continuous cost proportional to position size. - * Teaches: small positions cheap to hold, large positions expensive. */ + * Teaches: small positions cheap to hold, large positions expensive. + * Phase 3: gated by shaping_scale (behavioral; not in validation reward). */ float inventory_penalty = holding_cost_rate * fabsf(position); - reward -= inventory_penalty; + reward -= shaping_scale * inventory_penalty; /* Churn penalty: graduated cost for rapid position flips. * Not a hard gate — the model CAN exit early, but it costs extra. - * Zero penalty after churn_threshold bars. */ + * Zero penalty after churn_threshold bars. + * Phase 3: gated by shaping_scale (behavioral; not in validation reward). */ if (entering_trade && hold_time < churn_threshold && hold_time > 0.0f) { float churn_penalty = churn_penalty_scale * (churn_threshold - hold_time) / churn_threshold; - reward -= churn_penalty; + reward -= shaping_scale * churn_penalty; } /* Plan conviction as reward scaling — ONLY when plan head is mature. diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 8bcd8c3c5..2f99c3b04 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -681,6 +681,20 @@ pub struct GpuExperienceCollector { /// readiness < 0.5 → plan disabled. readiness ≥ 0.5 → plan enforces. readiness_dev_ptr: u64, + /// Phase 3 unified-env: pinned device-mapped exploration_scale [1] f32. + /// 1.0 = full training mode (saboteur, plan_params, counterfactuals on). + /// 0.0 = deterministic / validation-equivalent (those perturbations off). + /// Lerped through saboteur multipliers and counterfactual flip probability. + exploration_scale_pinned: *mut f32, + exploration_scale_dev_ptr: u64, + + /// Phase 3 unified-env: pinned device-mapped shaping_scale [1] f32. + /// 1.0 = full reward shaping (drawdown/churn/inventory/micro-reward applied). + /// 0.0 = pure per-bar P&L only — what the validation kernel measures. + /// Multiplies the additive shaping bundle in experience_env_step. + shaping_scale_pinned: *mut f32, + shaping_scale_dev_ptr: u64, + /// v8: Per-bar curriculum difficulty scoring kernel. difficulty_scores_kernel: CudaFunction, /// v8: Hindsight experience relabeling kernel (optimal exit). @@ -1158,6 +1172,26 @@ impl GpuExperienceCollector { (host_ptr as *mut f32, dev_ptr_out) }; + // Phase 3 unified-env: exploration_scale + shaping_scale pinned scalars. + // Same alloc pattern as cost_anneal. Default to 1.0 (full training mode); + // validation-mode callers flip both to 0.0 by writing through the *mut. + let alloc_pinned_scalar = |label: &str, init: f32| -> (*mut f32, u64) { + let mut host_ptr: *mut std::ffi::c_void = std::ptr::null_mut(); + let mut dev_ptr_out: u64 = 0; + unsafe { + let rc = cudarc::driver::sys::cuMemAllocHost_v2(&mut host_ptr, std::mem::size_of::()); + assert_eq!(rc, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemAllocHost for {label}"); + let rc2 = cudarc::driver::sys::cuMemHostGetDevicePointer_v2(&mut dev_ptr_out, host_ptr, 0); + assert_eq!(rc2, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemHostGetDevicePointer for {label}"); + *(host_ptr as *mut f32) = init; + } + (host_ptr as *mut f32, dev_ptr_out) + }; + let (exploration_scale_pinned, exploration_scale_dev_ptr) = + alloc_pinned_scalar("exploration_scale", 1.0); + let (shaping_scale_pinned, shaping_scale_dev_ptr) = + alloc_pinned_scalar("shaping_scale", 1.0); + Ok(Self { stream, market_dim, @@ -1257,6 +1291,10 @@ impl GpuExperienceCollector { isv_signals_dev_ptr: 0, // NULL until trainer sets it plan_params_dev_ptr: 0, // NULL until trainer sets it readiness_dev_ptr: 0, // NULL until trainer sets it + exploration_scale_pinned, + exploration_scale_dev_ptr, + shaping_scale_pinned, + shaping_scale_dev_ptr, difficulty_scores_kernel, hindsight_relabel_kernel, td_lambda_kernel, @@ -1286,6 +1324,22 @@ impl GpuExperienceCollector { unsafe { *self.cost_anneal_pinned = factor; } } + /// Phase 3 unified-env: set exploration_scale ∈ [0, 1]. + /// 1.0 = full training mode (saboteur, plan_params, counterfactuals on). + /// 0.0 = deterministic / validation-equivalent (those perturbations off). + /// Pinned device-mapped — zero memcpy, kernel reads next launch. + pub fn set_exploration_scale(&mut self, scale: f32) { + unsafe { *self.exploration_scale_pinned = scale.clamp(0.0, 1.0); } + } + + /// Phase 3 unified-env: set shaping_scale ∈ [0, 1]. + /// 1.0 = full reward shaping bundle (drawdown/churn/inventory/micro-reward). + /// 0.0 = pure per-bar P&L only (validation reward). + /// Pinned device-mapped — zero memcpy, kernel reads next launch. + pub fn set_shaping_scale(&mut self, scale: f32) { + unsafe { *self.shaping_scale_pinned = scale.clamp(0.0, 1.0); } + } + /// Set ISV signals device pointer for adaptive hold enforcement. /// The experience collector's env_step kernel reads ISV signals to compute /// dynamic hold time: base + f(stability, confidence). Pass 0 for static hold. @@ -2446,6 +2500,8 @@ impl GpuExperienceCollector { .arg(&config.hold_quality_weight) // micro-reward: hold quality weight .arg(&config.micro_reward_temp) // micro-reward: tanh temperature .arg(&cf_ratio) // D4/N4: adaptive counterfactual flip ratio + .arg(&self.exploration_scale_dev_ptr) // Phase 3 unified-env: gates saboteur/plan/CF flips + .arg(&self.shaping_scale_dev_ptr) // Phase 3 unified-env: gates additive reward shaping .launch(launch_cfg) .map_err(|e| MLError::ModelError(format!( "experience_env_step t={t}: {e}"