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

fix(kelly): Task 2.Z — adaptive warmup_floor from policy conviction

Browse Source 
Replaces static warmup_floor=0.5f in kelly_position_cap (trade_physics.cuh
line ~293) with an adaptive signal derived 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 policy at cold start). Low conviction
-> low floor (safety dominates). Clamped to [0, 1] - structural bound;
conviction only matters until maturity->1 (10+ trades) when the blend
flows to pure kelly_f and the floor contribution vanishes.

Wiring:
  1. kelly_position_cap / apply_kelly_cap / unified_env_step_core all gain
     a float `conviction` parameter (threaded through, no default).
  2. experience_action_select gains a new out_conviction[N] output buffer,
     computed as (max(q_dir) - min(q_dir)) / fmaxf(isv[21], 1e-6f) clamped
     [0,1]. Fallback when ISV[21]<=1e-6f: use q_range itself as denom,
     conviction=1 (trust policy face-value - same outcome as old static
     0.5 at health=1, but from a real signal shape).
  3. experience_env_step & backtest_env_step{,_batch} gain a
     conviction_ptr[N] (or [chunk_len*N]) input buffer, NULL-tolerant
     with fallback 1.0.
  4. Rust launch side: GpuExperienceCollector allocates conviction_buf[N]
     alongside q_gaps_buf; GpuBacktestEvaluator allocates chunked
     conviction buffer cn=n_windows*CHUNK_SIZE. Both wired into the 4
     kernel launches (experience_action_select + experience_env_step;
     experience_action_select + backtest_env_step_batch).

The previous static 0.5 pinned cold-start cap to <=0.375*max_pos at
health=0.5 (safety_multiplier=0.75), which combined with the
`abs_pos < 0.375f -> actual_mag = 0` threshold in the unified-env-core
magnitude decoder pinned realised magnitude to Quarter for the first
~10 trades regardless of what the policy's mag_idx requested. Smoke
test EVAL_DIST=[1.0, 0.0, 0.0] pre-fix was a downstream symptom of
this physics gate, not a magnitude Q-head failure.

Per feedback_adaptive_not_tuned.md: no hard-coded numeric knobs.
Conviction flows from the network's own Q-spread signal, evolving
temporally. Per feedback_no_functionality_removal.md: Kelly cap is
modified, not removed; warmup_floor is made adaptive, not deleted.

Test plan:
  SQLX_OFFLINE=true CARGO_INCREMENTAL=0 cargo check -p ml
    --example train_baseline_rl --tests  -> passes.
  Smoke (magnitude_distribution, 20 epochs) shows:
    [MAG_DIST] Quarter~0.62-0.70 Half~0.15-0.19 Full~0.13-0.21
    Training-mode magnitude distribution is now healthy (>5% floor
    for Half and Full each). Eval-mode smoke is non-deterministic in
    this horizon (EVAL_DIST Quarter collapse observed 2/3 runs; one
    run EVAL_DIST=[0.573, 0.325, 0.102]). Direction regression NOT
    triggered - Hold stays ~0 in most runs, Flat occasionally high
    (this is known H10 eval tie-break variance, unrelated to the
    Kelly change). q_dir_abs_ref observed in ISV_DIR_MEANS:
    ~0.14-0.60 across runs - conviction signal is flowing.

The Kelly fix removes a structural pin; downstream EVAL_DIST variance
now reflects Q-head conviction honestly rather than being clamped.
This commit is contained in:
jgrusewski
2026-04-23 00:32:30 +02:00
parent b8cd4e1d94
commit ff683470e7
5 changed files with 173 additions and 22 deletions
Download Patch File Download Diff File Expand all files Collapse all files

25
crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
View File

@@ -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) */

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

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

18
crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs
View File

@@ -328,6 +328,11 @@ pub struct GpuBacktestEvaluator {
/// Chunked Q-gap output buffer (chunked path: [n_windows * CHUNK_SIZE]).
chunked_q_gaps_buf: Option<CudaSlice<f32>>,
/// 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<CudaSlice<f32>>,
/// 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::<f32>(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::<f32>(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}"))

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

@@ -528,6 +528,11 @@ pub struct GpuExperienceCollector {
batch_actions: CudaSlice<i32>,
/// Q-gap conviction per episode [N]: output by action_select, input to env_step.
q_gaps_buf: CudaSlice<f32>,
/// 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<f32>,
/// Kelly optimal fraction (per-epoch scalar from CPU).
/// Current timestep counter per episode: [N].
current_timesteps: CudaSlice<i32>,
@@ -1087,6 +1092,11 @@ impl GpuExperienceCollector {
let q_gaps_buf = stream
.alloc_zeros::<f32>(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::<f32>(n)
.map_err(|e| MLError::ModelError(format!("alloc conviction_buf: {e}")))?;
let current_timesteps = stream
.alloc_zeros::<i32>(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}"

76
crates/ml/src/cuda_pipeline/trade_physics.cuh
View File

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