diff --git a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu index 445578902..0e1ed5296 100644 --- a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu +++ b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu @@ -112,7 +112,11 @@ extern "C" __global__ void backtest_env_step( * 1ffdf38dd). Kept for ABI alignment. * Both pointers may be NULL — kernel ignores them (NO-OP today). */ const float* __restrict__ exploration_scale_ptr, - const float* __restrict__ shaping_scale_ptr + const float* __restrict__ shaping_scale_ptr, + /* Per-sample conviction [n_windows] for adaptive Kelly warmup floor. + * Written by experience_action_select (out_conviction). NULL → fallback + * 1.0 (trust policy face-value). Threaded into unified_env_step_core. */ + const float* __restrict__ conviction_ptr ) { /* Both scales are no-ops in backtest today — ignored. The signature * alignment is so the future unified_env_kernel can be a drop-in @@ -214,6 +218,11 @@ extern "C" __global__ void backtest_env_step( &trail_vol_scale, &trail_trend_scale ); + /* Per-sample conviction for Kelly warmup floor — matches training. */ + float conviction_core = (conviction_ptr != NULL) + ? fminf(1.0f, fmaxf(0.0f, conviction_ptr[w])) + : 1.0f; + unified_env_step_core( action_val, close, b1_size, b2_size, b3_size, @@ -221,6 +230,7 @@ extern "C" __global__ void backtest_env_step( contract_multiplier, margin_pct, health_k, trail_vol_scale, trail_trend_scale, + conviction_core, /* adaptive Kelly warmup floor */ &position, &cash, &entry_price, &hold_time, &max_equity, value, /* prev_equity (snapshot) */ @@ -350,7 +360,11 @@ extern "C" __global__ void backtest_env_step_batch( /* Phase 3 unified-env scalars — see single-step variant for rationale. * NO-OP in backtest; kept for ABI alignment with future unified kernel. */ const float* __restrict__ exploration_scale_ptr, - const float* __restrict__ shaping_scale_ptr + const float* __restrict__ shaping_scale_ptr, + /* Per-sample conviction [chunk_len * n_windows] — layout matches + * chunked_actions ([step_offset * n_windows + window]). NULL → fallback + * 1.0. Threaded into unified_env_step_core for adaptive Kelly warmup. */ + const float* __restrict__ conviction_ptr ) { /* Phase 3 unified-env: shaping params retained for ABI but no longer * applied to step_returns. See single-step variant header for rationale. */ @@ -447,6 +461,12 @@ extern "C" __global__ void backtest_env_step_batch( &trail_vol_scale_b, &trail_trend_scale_b ); + /* Per-sample conviction for adaptive Kelly warmup floor. Layout + * matches chunked_actions: [s * n_windows + w]. */ + float conviction_core = (conviction_ptr != NULL) + ? fminf(1.0f, fmaxf(0.0f, conviction_ptr[s * n_windows + w])) + : 1.0f; + unified_env_step_core( action_val, close, b1_size, b2_size, b3_size, @@ -454,6 +474,7 @@ extern "C" __global__ void backtest_env_step_batch( contract_multiplier, margin_pct, health_k_batch, trail_vol_scale_b, trail_trend_scale_b, + conviction_core, /* adaptive Kelly warmup floor */ &position, &cash, &entry_price, &hold_time, &max_equity, value, /* prev_equity (snapshot for step_return) */ diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 9ef950e38..ba3b72178 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -726,6 +726,10 @@ extern "C" __global__ void experience_state_gather( * @param q_values [N, q_stride] Q-values from cuBLAS (bf16) * @param out_actions [N] selected factored action index (output) * @param out_q_gaps [N] output: Q-gap per episode for conviction sizing (bf16) + * @param out_conviction [N] output: per-sample direction-branch conviction ∈ [0, 1] + * = (max(q_dir) − min(q_dir)) / fmaxf(q_dir_abs_ref, 1e-6f). + * Consumed by env_step → unified_env_step_core → apply_kelly_cap + * as the adaptive warmup floor. NULL-tolerant. * @param epsilon exploration probability in [0, 1] * @param N number of episodes * @param b0_size direction branch size (4) @@ -760,7 +764,8 @@ extern "C" __global__ void experience_action_select( int timestep, /* current timestep for stateless RNG */ const float* __restrict__ per_sample_epsilon, /* [N] IQL expectile gap epsilon, NULL=use cosine schedule */ const float* __restrict__ isv_signals_ptr, /* [8] pinned ISV signals for adaptive hold. NULL = static. */ - int contrarian_active /* D7/N7: when non-zero, negate Q values before Boltzmann → argmin sampling */ + int contrarian_active, /* D7/N7: when non-zero, negate Q values before Boltzmann → argmin sampling */ + float* out_conviction /* [N] output: direction-branch conviction ∈ [0,1] for Kelly warmup. NULL=skip. */ ) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= N) return; @@ -1125,6 +1130,45 @@ extern "C" __global__ void experience_action_select( float gap = best_q - (q_sign_local * q_b0[flat_idx]); out_q_gaps[i] = (gap > 0.0f) ? gap : 0.0f; } + + /* Per-sample direction-branch conviction for adaptive Kelly warmup floor. + * conviction = (max(q_dir) − min(q_dir)) / fmaxf(q_dir_abs_ref, 1e-6f) + * clamped to [0, 1]. 0 = uniform Q (no preference → cold caution dominates). + * 1 = full spread at the observed Q scale (network trusts its pick). + * + * q_dir_abs_ref comes from ISV slot 21 — an EMA of max(|Q_mean|) across + * direction bins, updated every epoch. It's the natural scale for + * normalising Q-spreads into a temporal policy-confidence signal. + * + * Fallback when isv_signals_ptr is NULL or ISV[21] ≤ 1e-6f (pre-first- + * update / cold-start): use q_range itself as the denominator, yielding + * conviction = 1.0. Trust the policy at face value when we cannot measure + * its scale — same outcome as the previous static warmup_floor=0.5 would + * have produced at health=1, but driven by a real signal shape rather + * than a tuned constant. Once ISV[21] graduates above the epsilon floor + * the adaptive normalisation takes over. + * + * Uses raw Q (no contrarian sign flip) — conviction is about the policy's + * confidence in the decision, regardless of which side of the distribution + * argmax lands on. */ + if (out_conviction != NULL) { + float q_max_raw = q_b0[0]; + float q_min_raw = q_b0[0]; + for (int a = 1; a < b0_size; a++) { + float qv = q_b0[a]; + q_max_raw = fmaxf(q_max_raw, qv); + q_min_raw = fminf(q_min_raw, qv); + } + float q_range = q_max_raw - q_min_raw; + float q_dir_abs_ref = (isv_signals_ptr != NULL) + ? isv_signals_ptr[21] + : 0.0f; + float denom = (q_dir_abs_ref > 1e-6f) + ? q_dir_abs_ref + : fmaxf(q_range, 1e-6f); + float conviction = q_range / denom; + out_conviction[i] = fminf(1.0f, fmaxf(0.0f, conviction)); + } } /* ================================================================== */ @@ -1271,7 +1315,12 @@ extern "C" __global__ void experience_env_step( * 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 + const float* __restrict__ shaping_scale_ptr, + /* Per-sample policy conviction ∈ [0, 1] for adaptive Kelly warmup floor. + * Written by experience_action_select (out_conviction). NULL → fallback + * 1.0 (trust the policy face-value; same outcome as the old static + * warmup_floor=0.5 at health=1). Threaded into unified_env_step_core. */ + const float* __restrict__ conviction_ptr ) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= N) return; @@ -1604,6 +1653,15 @@ extern "C" __global__ void experience_env_step( &trail_vol_scale, &trail_trend_scale ); + /* Per-sample conviction for Kelly warmup floor. NULL fallback → 1.0 + * (trust policy face-value when action_select didn't write a value — + * e.g. future kernels that skip the conviction buffer). The action_select + * kernel clamps to [0, 1] already; defensive clamp here guards against + * direct device writes bypassing that path. */ + float conviction_core = (conviction_ptr != NULL) + ? fminf(1.0f, fmaxf(0.0f, conviction_ptr[i])) + : 1.0f; + unified_env_step_core( action_idx, raw_close, b1_size, b2_size, b3_size, @@ -1613,6 +1671,7 @@ extern "C" __global__ void experience_env_step( contract_multiplier, margin_pct, health_k_train, trail_vol_scale, trail_trend_scale, + conviction_core, /* adaptive Kelly warmup floor */ &position, &cash, &entry_price, &hold_time, &peak_equity, /* helper updates peak_equity via fmaxf */ prev_equity, /* snapshot for step_return denominator */ diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs index 8548590f8..35dc4def5 100644 --- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs @@ -328,6 +328,11 @@ pub struct GpuBacktestEvaluator { /// Chunked Q-gap output buffer (chunked path: [n_windows * CHUNK_SIZE]). chunked_q_gaps_buf: Option>, + /// Per-sample direction-branch conviction ∈ [0, 1], chunked layout + /// `[chunk_len * n_windows]` matching `chunked_actions_buf`. Written by + /// `experience_action_select` each chunk; consumed by `backtest_env_step_batch` + /// as the adaptive Kelly warmup-floor signal (via `unified_env_step_core`). + chunked_conviction_buf: Option>, /// Cached CUDA graph for `evaluate_dqn_graphed()` step loop. /// @@ -584,6 +589,7 @@ impl GpuBacktestEvaluator { b3_size: 0, q_gap_threshold: 0.0, chunked_q_gaps_buf: None, + chunked_conviction_buf: None, dqn_graph: None, chunked_q_values: None, chunked_states_buf: None, @@ -1002,6 +1008,9 @@ impl GpuBacktestEvaluator { let ch_q_gaps = self.chunked_q_gaps_buf.as_ref().ok_or_else(|| { MLError::ModelError("chunked_q_gaps_buf unexpectedly None".to_owned()) })?; + let ch_conviction = self.chunked_conviction_buf.as_ref().ok_or_else(|| { + MLError::ModelError("chunked_conviction_buf unexpectedly None".to_owned()) + })?; let n = self.n_windows; let state_dim = self.state_dim; @@ -1107,6 +1116,7 @@ impl GpuBacktestEvaluator { .arg(&0u64) // per_sample_epsilon: NULL → use cosine schedule .arg(&self.isv_signals_ptr) // ISV signals for adaptive hold .arg(&0i32) // F6: contrarian_active — always off during backtest + .arg(ch_conviction) // out_conviction: per-sample Kelly warmup signal .launch(launch_cfg) .map_err(|e| MLError::ModelError(format!( "chunked action_select chunk_start={chunk_start}: {e}" @@ -1159,6 +1169,7 @@ impl GpuBacktestEvaluator { // Phase 3 unified-env scalars (NULL = no-op in val mode). .arg(&0u64) // exploration_scale_ptr (NULL) .arg(&0u64) // shaping_scale_ptr (NULL) + .arg(ch_conviction) // conviction_ptr: adaptive Kelly warmup signal .launch(LaunchConfig { grid_dim: (env_blocks.max(1), 1, 1), block_dim: (256, 1, 1), @@ -1376,6 +1387,11 @@ impl GpuBacktestEvaluator { let ch_q_gaps = self.stream.alloc_zeros::(cn) .map_err(|e| MLError::ModelError(format!("alloc chunked q_gaps: {e}")))?; + /* Per-sample conviction for adaptive Kelly warmup floor. Layout matches + * chunked_actions/chunked_q_gaps — [chunk_len * n_windows]. */ + let ch_conviction = self.stream.alloc_zeros::(cn) + .map_err(|e| MLError::ModelError(format!("alloc chunked conviction: {e}")))?; + info!( "GpuBacktestEvaluator: action-select initialised — {} windows, \ branches=[{},{},{},{}], chunk_size={chunk}, chunked_batch={cn}", @@ -1392,6 +1408,7 @@ impl GpuBacktestEvaluator { self.chunked_states_buf = Some(ch_states_buf); self.chunked_actions_buf = Some(ch_actions_buf); self.chunked_q_gaps_buf = Some(ch_q_gaps); + self.chunked_conviction_buf = Some(ch_conviction); Ok(()) } @@ -1493,6 +1510,7 @@ impl GpuBacktestEvaluator { // be 0.0 anyway. .arg(&0u64) // exploration_scale_ptr (NULL) .arg(&0u64) // shaping_scale_ptr (NULL) + .arg(&0u64) // conviction_ptr (NULL) — fallback 1.0 in-kernel .launch(env_cfg) .map_err(|e| { MLError::ModelError(format!("backtest_env_step launch step {step}: {e}")) diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 70d46eb54..d459b686e 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -528,6 +528,11 @@ pub struct GpuExperienceCollector { batch_actions: CudaSlice, /// Q-gap conviction per episode [N]: output by action_select, input to env_step. q_gaps_buf: CudaSlice, + /// Per-sample direction-branch conviction ∈ [0, 1] per episode [N]: output by + /// action_select, input to env_step. Derived from direction Q-spread normalised + /// by q_dir_abs_ref (ISV[21]). Drives adaptive Kelly warmup floor in + /// `unified_env_step_core` (replaces the old static 0.5 floor). + conviction_buf: CudaSlice, /// Kelly optimal fraction (per-epoch scalar from CPU). /// Current timestep counter per episode: [N]. current_timesteps: CudaSlice, @@ -1087,6 +1092,11 @@ impl GpuExperienceCollector { let q_gaps_buf = stream .alloc_zeros::(n) .map_err(|e| MLError::ModelError(format!("alloc q_gaps_buf: {e}")))?; + /* Per-sample conviction [N] for adaptive Kelly warmup floor — written by + * experience_action_select, consumed by experience_env_step. */ + let conviction_buf = stream + .alloc_zeros::(n) + .map_err(|e| MLError::ModelError(format!("alloc conviction_buf: {e}")))?; let current_timesteps = stream .alloc_zeros::(n) .map_err(|e| MLError::ModelError(format!("alloc current_timesteps: {e}")))?; @@ -1343,6 +1353,7 @@ impl GpuExperienceCollector { batch_states, batch_actions, q_gaps_buf, + conviction_buf, current_timesteps, q_values, q_select_buf, @@ -2867,6 +2878,7 @@ impl GpuExperienceCollector { .arg(&self.per_sample_epsilon_ptr) // IQL expectile gap epsilon (0=cosine schedule) .arg(&self.isv_signals_dev_ptr) // ISV signals for adaptive hold (0=NULL=static) .arg(&(self.contrarian_active_cache as i32)) // D7/N7: Q-negation flag + .arg(&mut self.conviction_buf) // out_conviction: per-sample Kelly warmup signal .launch(launch_cfg) .map_err(|e| MLError::ModelError(format!( "experience_action_select t={t}: {e}" @@ -2998,6 +3010,7 @@ impl GpuExperienceCollector { .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 + .arg(&self.conviction_buf) // conviction_ptr: adaptive Kelly warmup signal .launch(launch_cfg) .map_err(|e| MLError::ModelError(format!( "experience_env_step t={t}: {e}" diff --git a/crates/ml/src/cuda_pipeline/trade_physics.cuh b/crates/ml/src/cuda_pipeline/trade_physics.cuh index 20d93dc8c..5fe68bb90 100644 --- a/crates/ml/src/cuda_pipeline/trade_physics.cuh +++ b/crates/ml/src/cuda_pipeline/trade_physics.cuh @@ -244,6 +244,18 @@ __device__ __forceinline__ float compute_drawdown( * less reliable. Same temporal-coupling pattern as the distillation SAXPY * alpha and the Q-target label smoothing eps. Applied uniformly in both * training and backtest envs so there's no train/val mismatch. + * + * The `conviction` parameter (∈ [0, 1], caller-normalised) drives the + * cold-start warmup floor — no longer a static 0.5f. Derived by the + * caller from the policy's per-sample direction Q-spread normalised by + * the `q_dir_abs_ref` ISV EMA (isv_signals[21]). High conviction → high + * floor (trust the policy to size up even before Kelly stats accumulate). + * Low conviction → low floor (safety dominates). `warmup_floor` only + * matters during the cold-start window; once `maturity → 1` (trade count + * ≥ 10) the blend flows to pure kelly_f and the floor contribution + * vanishes. Conviction is [0, 1] clamped here — a structural bound + * (fraction of max_position), not a tuned numeric knob; the raw signal + * IS the floor. */ __device__ __forceinline__ float kelly_position_cap( float win_count, @@ -251,7 +263,8 @@ __device__ __forceinline__ float kelly_position_cap( float sum_wins, float sum_losses, float max_position, - float safety_multiplier + float safety_multiplier, + float conviction ) { const float prior_wins = 2.0f; const float prior_losses = 2.0f; @@ -273,31 +286,41 @@ __device__ __forceinline__ float kelly_position_cap( * accumulate. That would clamp the agent to zero position forever * and starve Q-learning before any signal can form. * - * Blend towards a warmup floor (50% of max_position) weighted by - * stats maturity: maturity = min(1, total_trades / 10). Early on - * the cap is dominated by the floor; as real trades accumulate it - * transitions to true Kelly. At 10+ completed trades the floor - * contribution is zero — pure data-driven Kelly. + * Blend towards a warmup floor weighted by stats maturity: + * maturity = min(1, total_trades / 10). Early on the cap is dominated + * by the floor; as real trades accumulate it transitions to true + * Kelly. At 10+ completed trades the floor contribution is zero — + * pure data-driven Kelly. * - * Floor at 0.5 (not full 1.0): combined with the health-coupled - * safety multiplier (0.5 + 0.5×health ∈ [0.5, 1.0]), cold-start - * cap ranges from 25% (health=0, safety=0.5) to 50% (health=1, - * safety=1.0) of max_position — enough room for Q-learning to - * discover non-zero positions without reckless over-levering. + * Adaptive warmup floor: derived from policy-observed conviction + * (normalised direction Q-spread, computed at action_select and + * passed through by the caller). High conviction → trust the policy + * to size up even at cold start. Low conviction → safety dominates + * (floor collapses toward zero). Combined with the health-coupled + * safety multiplier (0.5 + 0.5×health ∈ [0.5, 1.0]), cold-start cap + * adapts to both network health AND per-sample Q-spread. + * + * Clamped to [0, 1] — structural bound (a fraction of max_position). + * No hard-coded numeric knob. Conviction flows adaptively from the + * network's own Q-spread signal, evolving temporally with training. + * Once maturity → 1 the floor term drops out entirely, so conviction + * only matters during the cold-start window. * * 10-trade maturity is short enough that learning isn't suppressed * for long; the priors (4 virtual trades) already ensure the formula * doesn't swing wildly on small samples. */ float total_trades = win_count + loss_count; float maturity = fminf(1.0f, total_trades * 0.1f); - const float warmup_floor = 0.5f; + float warmup_floor = fminf(1.0f, fmaxf(0.0f, conviction)); float effective_kelly = maturity * kelly_f + (1.0f - maturity) * warmup_floor; return effective_kelly * max_position * safety_multiplier; } /* Apply Kelly cap: clamp a signed target_position to [-cap, +cap]. - * Convenience for call sites that want a single clamp call. */ + * Convenience for call sites that want a single clamp call. + * `conviction` ∈ [0, 1] threads through to kelly_position_cap's adaptive + * warmup_floor — see docstring there. */ __device__ __forceinline__ float apply_kelly_cap( float target_position, float win_count, @@ -305,10 +328,11 @@ __device__ __forceinline__ float apply_kelly_cap( float sum_wins, float sum_losses, float max_position, - float safety_multiplier + float safety_multiplier, + float conviction ) { float cap = kelly_position_cap(win_count, loss_count, sum_wins, sum_losses, - max_position, safety_multiplier); + max_position, safety_multiplier, conviction); return fmaxf(-cap, fminf(cap, target_position)); } @@ -482,7 +506,11 @@ __device__ __forceinline__ void compute_regime_trail_scales( * 2. Hold action passthrough * 3. Target position via compute_target_position_4branch * 4. Margin cap (apply_margin_cap) - * 5. Kelly cap with health-coupled safety (apply_kelly_cap) + * 5. Kelly cap with health-coupled safety AND policy conviction + * (apply_kelly_cap). Conviction ∈ [0, 1] — per-sample, derived from + * direction Q-spread normalised by q_dir_abs_ref. Drives the cold- + * start warmup floor so Kelly cap adapts to policy conviction instead + * of a static 0.5 floor; once maturity → 1 it's a no-op. * 6. Trailing stop check — base 0.5%, regime-adaptive via vol_scale and * trend_scale (caller resolves from ADX/CUSUM features; identical * formulas in both training and val to preserve train/eval parity) @@ -530,6 +558,15 @@ __device__ __forceinline__ void unified_env_step_core( * with ADX, CUSUM read from features[bar * market_dim + 40/41]. */ float vol_scale, float trend_scale, + /* ── Policy conviction ∈ [0, 1] for adaptive Kelly warmup floor ── + * Per-sample direction Q-spread normalised by q_dir_abs_ref (ISV[21]). + * Drives `kelly_position_cap` warmup_floor — high conviction relaxes + * the cold-start cap so the policy can size up; low conviction keeps + * safety dominant. Only affects the cap when maturity < 1 (<10 trades); + * once Kelly stats graduate the floor term drops out. Clamped [0,1] + * inside kelly_position_cap; caller may pass out-of-range and the + * helper bounds it. */ + float conviction, /* ── Per-thread mutable state (in/out) ── */ float* position, float* cash, @@ -571,7 +608,9 @@ __device__ __forceinline__ void unified_env_step_core( ); } - /* ── Step 5: Kelly cap (health-coupled safety; uses physics scale) ── */ + /* ── Step 5: Kelly cap (health-coupled safety + conviction-adaptive + * warmup floor; uses physics scale). Conviction flows per-sample from + * the action_select kernel's direction Q-spread. */ if (!is_hold_action) { float h = fminf(1.0f, fmaxf(0.0f, health_for_kelly)); float safety = 0.5f + 0.5f * h; @@ -579,7 +618,8 @@ __device__ __forceinline__ void unified_env_step_core( target_position, *win_count, *loss_count, *sum_wins, *sum_losses, - max_position_physics, safety + max_position_physics, safety, + conviction ); }