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

phase3(env-unification): add exploration_scale + shaping_scale control scalars

Browse Source 
Foundation for the unified train/val env kernel (Phase 3 of the env-unification
design). Adds two pinned device-mapped scalars to the training kernel that gate
the asymmetries currently separating training from validation:

  exploration_scale ∈ [0, 1] gates training-only stochastic perturbations:
    - Saboteur cost noise: sab_eff = 1 + exploration_scale × (sab - 1)
      → identity at scale=0, full saboteur at scale=1
    - Counterfactual flip rate: effective_cf = exploration_scale × cf_ratio
    - Plan-params position scaling: only active when scale ≥ 0.5

  shaping_scale ∈ [0, 1] gates additive reward-shaping bundles:
    - Drawdown penalty (× shaping_scale)
    - Inventory penalty (× shaping_scale)
    - Churn penalty (× shaping_scale)
    - Micro-reward composite (× shaping_scale)
    - Holding-cost fallback (× shaping_scale)
    Capital-floor reward and segment-completion P&L are NOT gated — those
    are physics/safety, not behavioral shaping.

Both scalars live in pinned host memory mapped to the device, following the
existing pattern used for cost_anneal_pinned, isv_signals_dev_ptr, etc.
Default value is 1.0 (full training mode); zero memcpy on update — the kernel
reads the current value on its next launch via cuMemHostGetDevicePointer.

Setters exposed:
  GpuExperienceCollector::set_exploration_scale(f32)
  GpuExperienceCollector::set_shaping_scale(f32)

Backward compatibility: scalars default to 1.0, kernel pointers are
NULL-tolerant (falls back to 1.0 inside the kernel if pointer is NULL).
The existing TD-propagation smoke test passes unchanged with default scales
(Best Sharpe 19.34 at epoch 17 in this run; was 15.19 baseline — within
run-to-run variance, no regression).

What this unblocks (deferred to next session, task #18):
  - Wire validation backtest paths to call experience_env_step with both
    scalars at 0.0 instead of using the separate backtest_env_kernel.
  - Verify step_returns are byte-equivalent between scales=0 path and the
    legacy backtest kernel (one source of truth for env physics).
  - Delete backtest_env_kernel.cu (~678 LOC) and its launcher.

Files touched:
  crates/ml/src/cuda_pipeline/experience_kernels.cu          (+62 / -19)
  crates/ml/src/cuda_pipeline/gpu_experience_collector.rs    (+56)

Verified: SQLX_OFFLINE=true cargo check -p ml --lib --tests passes.
TD-propagation smoke test runs cleanly end-to-end (32.89s).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-21 08:02:12 +02:00
parent 41858c31df
commit 3f6eb006ca
2 changed files with 118 additions and 19 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

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

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

@@ -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::<f32>());
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}"