diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 5bc144bf6..069c10c47 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -718,7 +718,8 @@ extern "C" __global__ void experience_action_select( float eps_ord_mult, /* per-branch epsilon multiplier: order type */ float eps_urg_mult, /* per-branch epsilon multiplier: urgency */ 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__ 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 i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= N) return; @@ -773,7 +774,16 @@ extern "C" __global__ void experience_action_select( int ps_base = i * 23; /* PORTFOLIO_STRIDE = 23 */ float hold_time_val = portfolio_states[ps_base + 10]; cur_position = portfolio_states[ps_base + 0]; - in_hold = (hold_time_val > 0.0f && hold_time_val < (float)min_hold_bars); + /* Adaptive hold: base + ISV-driven extension (same formula as Layer 2) */ + int adaptive_hold = min_hold_bars; + if (isv_signals_ptr != NULL) { + float stability = 1.0f / (1.0f + isv_signals_ptr[1]); + float confidence = 1.0f - fminf(isv_signals_ptr[3], 1.0f); + int extension = (int)(stability * confidence * 8.0f); + if (isv_signals_ptr[5] < -0.001f) extension = 0; + adaptive_hold = min_hold_bars + extension; + } + in_hold = (hold_time_val > 0.0f && hold_time_val < (float)adaptive_hold); } if (in_hold) { diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index da04bd3e0..88bec88d5 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -2217,6 +2217,7 @@ impl GpuExperienceCollector { .arg(&config.eps_urg_mult) // per-branch epsilon multiplier: urgency .arg(&(t as i32)) // timestep for stateless Philox RNG .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) .launch(launch_cfg) .map_err(|e| MLError::ModelError(format!( "experience_action_select t={t}: {e}"