From 39efacf77dfa07ff08b3c13752606b7595c32afa Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sat, 30 May 2026 22:06:22 +0200 Subject: [PATCH] fix(rl): CMDP gates per-batch (one independent session per b) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Each batch element is an independent backtest session with its own $35k starting capital. The CMDP kernel had been summing per-batch rewards into a single ISV slot — at b=1024 the accumulator grew b_size× faster than the -$3500 single-account DD limit, tripping session_dd_triggered globally in ~250 steps (cluster alpha-rl-2j9k9 step 320: session_pnl_usd=-$5775, all 1024 sessions locked out). Same flaw on consec_loss_count: any 10 losses across the batch in a single step triggered cooldown for the entire fleet. Refactor — every Layer-1 state shard is per-batch: session_pnl_per_batch_d [b_size] f32 session_dd_triggered_per_batch_d [b_size] f32 consec_loss_per_batch_d [b_size] f32 cooldown_remaining_per_batch_d [b_size] f32 `actions_to_market_targets` now reads per-batch DD-triggered + cooldown, so one account hitting its limit does not lock out the other 1023. Summary ISV slots (RL_SESSION_PNL_USD, RL_CONSEC_LOSS_COUNT, RL_SESSION_DD_TRIGGERED, RL_COOLDOWN_REMAINING_STEPS) are now kernel-OUT only — they expose worst-account / any-triggered / max aggregates for diag + the single IQN-τ consumer ("be pessimistic when ANY session is in trouble"). 3 raw `actions_to_market_targets_fn.cu_function()` launch sites in integrated.rs were updated atomically per `feedback_no_partial_refactor` — public wrapper + 2 internal step_with_lobsim sites. The first attempt hit CUDA_ERROR_INVALID_VALUE in integrated_trainer_smoke because only the public wrapper had been updated. Validation: - 12/12 risk_stack_invariants pass on per-batch semantics (G1-G4 rewritten to seed per-batch buffers via write_slice_f32_d_pub). - 20/20 trade_management_kernels still pass (actions gating intact). - integrated_trainer_smoke passes end-to-end. - Local b=128 1k smoke: 1000/1000 steps clean, no NaN. CMDP behavior matches design: worst-of-128 hit DD (-$34k > $3500 limit), other 127 unaffected. IQN τ floored at 0.1 (worst-account drawdown >>10%), Kelly = 0.0 (observed edge negative: wr=0.21, R=1.57). --- .../cuda/actions_to_market_targets.cu | 17 ++- .../cuda/rl_cmdp_constraints_check.cu | 104 +++++++++++------- crates/ml-alpha/src/trainer/integrated.rs | 42 +++++++ .../ml-alpha/tests/risk_stack_invariants.rs | 96 ++++++++++------ 4 files changed, 180 insertions(+), 79 deletions(-) diff --git a/crates/ml-alpha/cuda/actions_to_market_targets.cu b/crates/ml-alpha/cuda/actions_to_market_targets.cu index dc1327049..e82b174b4 100644 --- a/crates/ml-alpha/cuda/actions_to_market_targets.cu +++ b/crates/ml-alpha/cuda/actions_to_market_targets.cu @@ -41,9 +41,11 @@ #define RL_ANTIMARTINGALE_MAX_INDEX 511 // ── Layer 1 (CMDP hard constraints — spec 2026-05-30-adaptive-risk-management) ── -#define RL_SESSION_DD_TRIGGERED_INDEX 664 +// Session-level DD-triggered + cooldown are PER-BATCH (one buffer slot per +// independent backtest session). The two ISV slots below carry only +// canonical-summary aggregates (worst per-batch DD, max per-batch +// cooldown) for diag visibility — gating reads the per-batch arrays. #define RL_MAX_OPEN_UNITS_INDEX 665 -#define RL_COOLDOWN_REMAINING_STEPS_INDEX 668 #define RL_NET_INVENTORY_LIMIT_USD_INDEX 670 // ── Layer 4 (Kelly fraction sizing — spec 2026-05-30-adaptive-risk-management) ── @@ -104,6 +106,11 @@ extern "C" __global__ void actions_to_market_targets( const int* __restrict__ pyramid_units_count, // [b_size] const int* __restrict__ close_unit_index, // [b_size] (-1 = auto oldest) const float* __restrict__ outcome_ema, // [b_size] per-batch + // CMDP per-batch state (Layer 1, spec 2026-05-30-adaptive-risk-management). + // Each batch element is an independent session; one account hitting + // its DD limit or cooldown does not gate the other 1023 sessions. + const float* __restrict__ session_dd_triggered_per_batch, // [b_size] + const float* __restrict__ cooldown_remaining_per_batch, // [b_size] int b_size, int pos_bytes ) { @@ -120,9 +127,11 @@ extern "C" __global__ void actions_to_market_targets( // ──────────────────────────────────────────────────────────────────── // Layer 1 (CMDP hard constraints — spec 2026-05-30-adaptive-risk-management). // Session-level overrides force Hold regardless of agent's choice. + // Per-batch (one independent backtest session per b); a tripped DD + // or active cooldown on session `b` does NOT lock out the others. // ──────────────────────────────────────────────────────────────────── - const bool dd_triggered = isv[RL_SESSION_DD_TRIGGERED_INDEX] >= 0.5f; - const bool in_cooldown = isv[RL_COOLDOWN_REMAINING_STEPS_INDEX] > 0.0f; + const bool dd_triggered = session_dd_triggered_per_batch[b] >= 0.5f; + const bool in_cooldown = cooldown_remaining_per_batch[b] > 0.0f; if (dd_triggered || in_cooldown) { market_targets[b * 2 + 0] = 2; // no-op side market_targets[b * 2 + 1] = 0; // size = 0 diff --git a/crates/ml-alpha/cuda/rl_cmdp_constraints_check.cu b/crates/ml-alpha/cuda/rl_cmdp_constraints_check.cu index 4f5c6eb1a..9d608d8c9 100644 --- a/crates/ml-alpha/cuda/rl_cmdp_constraints_check.cu +++ b/crates/ml-alpha/cuda/rl_cmdp_constraints_check.cu @@ -36,57 +36,81 @@ #define RL_COOLDOWN_REMAINING_STEPS_INDEX 668 #define RL_COOLDOWN_DURATION_INDEX 669 -// Outcome is derived inline: loss = (done & reward < 0), win = (done & reward > 0). +// Per-batch session-pnl + consec-loss tracking. Each batch element is an +// independent backtest "session" with its own $35k starting capital; +// `session_pnl_per_batch[b]` is the running PnL of that session and +// `consec_loss_per_batch[b]` is its losing-trade streak. The kernel +// writes per-batch state out for `actions_to_market_targets` to read +// (per-batch DD-triggered / cooldown gating) and also writes the +// canonical summary slots so diag + IQN-τ keep a single-account view. extern "C" __global__ void rl_cmdp_constraints_check( float* __restrict__ isv, - const float* __restrict__ rewards, // [b_size] shaped pnl - const float* __restrict__ dones, // [b_size] 1.0 = close + const float* __restrict__ rewards, // [b_size] shaped pnl + const float* __restrict__ dones, // [b_size] 1.0 = close + float* __restrict__ session_pnl_per_batch, // [b_size] IN/OUT + float* __restrict__ consec_loss_per_batch, // [b_size] IN/OUT + float* __restrict__ session_dd_triggered_per_batch,// [b_size] IN/OUT (0/1) + float* __restrict__ cooldown_remaining_per_batch, // [b_size] IN/OUT int b_size ) { if (threadIdx.x != 0 || threadIdx.y != 0 || threadIdx.z != 0) return; if (blockIdx.x != 0 || blockIdx.y != 0 || blockIdx.z != 0) return; - // ──────────────────────────────────────────────────────────── - // 1. Session pnl accumulation + DD check. - // Sum per-batch rewards across batch (the per-step pnl delta). - // ──────────────────────────────────────────────────────────── - float step_pnl = 0.0f; - for (int b = 0; b < b_size; ++b) { - step_pnl += rewards[b]; - } - const float session_pnl_new = isv[RL_SESSION_PNL_USD_INDEX] + step_pnl; - isv[RL_SESSION_PNL_USD_INDEX] = session_pnl_new; - if (session_pnl_new < isv[RL_SESSION_DD_LIMIT_USD_INDEX]) { - isv[RL_SESSION_DD_TRIGGERED_INDEX] = 1.0f; - } + const float dd_limit = isv[RL_SESSION_DD_LIMIT_USD_INDEX]; + const float consec_lim = isv[RL_CONSEC_LOSS_LIMIT_INDEX]; + const float cool_dur = isv[RL_COOLDOWN_DURATION_INDEX]; - // ──────────────────────────────────────────────────────────── - // 2. Cooldown decrement. - // ──────────────────────────────────────────────────────────── - const float cooldown_prev = isv[RL_COOLDOWN_REMAINING_STEPS_INDEX]; - if (cooldown_prev > 0.0f) { - isv[RL_COOLDOWN_REMAINING_STEPS_INDEX] = cooldown_prev - 1.0f; - } + // Aggregates for the canonical summary slots (single-account view). + float worst_pnl = 0.0f; // most-negative per-batch session_pnl + float worst_consec = 0.0f; // max per-batch streak + float any_dd_trig = 0.0f; + float max_cool_remain = 0.0f; - // ──────────────────────────────────────────────────────────── - // 3. Consecutive loss tracking. Outcome derived from done & reward sign. - // ──────────────────────────────────────────────────────────── - float consec = isv[RL_CONSEC_LOSS_COUNT_INDEX]; for (int b = 0; b < b_size; ++b) { - if (dones[b] < 0.5f) continue; - const float r = rewards[b]; - if (r < 0.0f) { - consec += 1.0f; - } else if (r > 0.0f) { + // ── 1. Per-batch session pnl + DD check ───────────────── + const float pnl_new = session_pnl_per_batch[b] + rewards[b]; + session_pnl_per_batch[b] = pnl_new; + if (pnl_new < dd_limit) { + session_dd_triggered_per_batch[b] = 1.0f; + } + + // ── 2. Per-batch cooldown decrement ───────────────────── + const float cool_prev = cooldown_remaining_per_batch[b]; + if (cool_prev > 0.0f) { + cooldown_remaining_per_batch[b] = cool_prev - 1.0f; + } + + // ── 3. Per-batch consec-loss tracking on closes ───────── + // r == 0 with done: ambiguous break-even, leave streak as-is. + float consec = consec_loss_per_batch[b]; + if (dones[b] >= 0.5f) { + const float r = rewards[b]; + if (r < 0.0f) consec += 1.0f; + else if (r > 0.0f) consec = 0.0f; + } + if (consec >= consec_lim) { + // Streak limit on THIS account hit — open its cooldown, + // reset streak. Other accounts unaffected. + cooldown_remaining_per_batch[b] = cool_dur; consec = 0.0f; } - // r == 0 with done: ambiguous — treat as no-change (rare break-even close) - } - if (consec >= isv[RL_CONSEC_LOSS_LIMIT_INDEX]) { - // Limit hit — start cooldown, reset counter. - isv[RL_COOLDOWN_REMAINING_STEPS_INDEX] = isv[RL_COOLDOWN_DURATION_INDEX]; - isv[RL_CONSEC_LOSS_COUNT_INDEX] = 0.0f; - } else { - isv[RL_CONSEC_LOSS_COUNT_INDEX] = consec; + consec_loss_per_batch[b] = consec; + + // Track aggregates for summary slots. + if (pnl_new < worst_pnl) worst_pnl = pnl_new; + if (consec > worst_consec) worst_consec = consec; + if (session_dd_triggered_per_batch[b] + >= 0.5f) any_dd_trig = 1.0f; + if (cooldown_remaining_per_batch[b] + > max_cool_remain) max_cool_remain = cooldown_remaining_per_batch[b]; } + + // ── Canonical summary slots ──────────────────────────────────── + // `RL_SESSION_PNL_USD_INDEX` exposes the worst-account PnL — this + // is the signal IQN-τ uses for `drawdown_frac` and what diag shows. + // Single-account semantics are preserved everywhere downstream. + isv[RL_SESSION_PNL_USD_INDEX] = worst_pnl; + isv[RL_SESSION_DD_TRIGGERED_INDEX] = any_dd_trig; + isv[RL_CONSEC_LOSS_COUNT_INDEX] = worst_consec; + isv[RL_COOLDOWN_REMAINING_STEPS_INDEX] = max_cool_remain; } diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index 85ff2570e..12b8a5869 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -743,6 +743,16 @@ pub struct IntegratedTrainer { rl_write_u64_fn: CudaFunction, ts_ns_d: CudaSlice, pub outcome_ema_d: CudaSlice, + /// CMDP Layer 1 per-batch state (spec 2026-05-30-adaptive-risk-management). + /// Each batch element is an independent backtest "session"; one session + /// hitting its DD limit or losing-streak cooldown must not gate the + /// other 1023 sessions. Canonical-summary aggregates (worst-account + /// PnL, max cooldown) are mirrored into ISV[662/664/667/668] for diag + /// + IQN-τ consumption. + pub session_pnl_per_batch_d: CudaSlice, + pub session_dd_triggered_per_batch_d: CudaSlice, + pub consec_loss_per_batch_d: CudaSlice, + pub cooldown_remaining_per_batch_d: CudaSlice, pub trade_context_d: CudaSlice, pub multires_output_d: CudaSlice, multires_state_d: CudaSlice, @@ -1983,6 +1993,21 @@ impl IntegratedTrainer { let outcome_ema_d = stream .alloc_zeros::(b_size) .context("alloc outcome_ema_d")?; + // CMDP per-batch state (spec 2026-05-30-adaptive-risk-management). + // Each backtest session in the batch tracks its own running PnL, + // DD-triggered flag, losing-streak counter, and cooldown remaining. + let session_pnl_per_batch_d = stream + .alloc_zeros::(b_size) + .context("alloc session_pnl_per_batch_d")?; + let session_dd_triggered_per_batch_d = stream + .alloc_zeros::(b_size) + .context("alloc session_dd_triggered_per_batch_d")?; + let consec_loss_per_batch_d = stream + .alloc_zeros::(b_size) + .context("alloc consec_loss_per_batch_d")?; + let cooldown_remaining_per_batch_d = stream + .alloc_zeros::(b_size) + .context("alloc cooldown_remaining_per_batch_d")?; let trade_context_d = stream .alloc_zeros::(b_size * 4) .context("alloc trade_context_d")?; @@ -2693,6 +2718,10 @@ impl IntegratedTrainer { rl_write_u64_fn, ts_ns_d, outcome_ema_d, + session_pnl_per_batch_d, + session_dd_triggered_per_batch_d, + consec_loss_per_batch_d, + cooldown_remaining_per_batch_d, trade_context_d, multires_output_d, multires_state_d, @@ -3583,6 +3612,10 @@ impl IntegratedTrainer { args.push_ptr(self.isv_dev_ptr); args.push_ptr(rewards_d.raw_ptr()); args.push_ptr(dones_d.raw_ptr()); + args.push_ptr(self.session_pnl_per_batch_d.raw_ptr()); + args.push_ptr(self.consec_loss_per_batch_d.raw_ptr()); + args.push_ptr(self.session_dd_triggered_per_batch_d.raw_ptr()); + args.push_ptr(self.cooldown_remaining_per_batch_d.raw_ptr()); args.push_i32(b_size as i32); let mut ptrs = args.build_arg_ptrs(); unsafe { @@ -3946,6 +3979,9 @@ impl IntegratedTrainer { args.push_ptr(pyramid_units_count_d.raw_ptr()); args.push_ptr(close_unit_index_d.raw_ptr()); args.push_ptr(outcome_ema_d.raw_ptr()); + // CMDP per-batch state — one DD-trigger / cooldown per session. + args.push_ptr(self.session_dd_triggered_per_batch_d.raw_ptr()); + args.push_ptr(self.cooldown_remaining_per_batch_d.raw_ptr()); args.push_i32(b_size_i); args.push_i32(pos_bytes_i); let mut ptrs = args.build_arg_ptrs(); @@ -6484,6 +6520,9 @@ impl IntegratedTrainer { args.push_ptr(self.pyramid_units_count_d.raw_ptr()); args.push_ptr(self.close_unit_index_d.raw_ptr()); args.push_ptr(self.outcome_ema_d.raw_ptr()); + // CMDP per-batch state — one DD-trigger / cooldown per session. + args.push_ptr(self.session_dd_triggered_per_batch_d.raw_ptr()); + args.push_ptr(self.cooldown_remaining_per_batch_d.raw_ptr()); args.push_i32(b_size_i); args.push_i32(pos_bytes_i); let mut ptrs = args.build_arg_ptrs(); @@ -8518,6 +8557,9 @@ impl IntegratedTrainer { args.push_ptr(self.pyramid_units_count_d.raw_ptr()); args.push_ptr(self.close_unit_index_d.raw_ptr()); args.push_ptr(self.outcome_ema_d.raw_ptr()); + // CMDP per-batch state — one DD-trigger / cooldown per session. + args.push_ptr(self.session_dd_triggered_per_batch_d.raw_ptr()); + args.push_ptr(self.cooldown_remaining_per_batch_d.raw_ptr()); args.push_i32(b_size_i); args.push_i32(pos_bytes_i); let mut ptrs = args.build_arg_ptrs(); diff --git a/crates/ml-alpha/tests/risk_stack_invariants.rs b/crates/ml-alpha/tests/risk_stack_invariants.rs index 616dfddc2..3d1b84e90 100644 --- a/crates/ml-alpha/tests/risk_stack_invariants.rs +++ b/crates/ml-alpha/tests/risk_stack_invariants.rs @@ -70,88 +70,113 @@ fn upload_f32(stream: &Arc, host: &[f32]) -> Result> // LAYER 1 — CMDP hard constraints // ═════════════════════════════════════════════════════════════════════ -// G1 — DD breaker sets sticky triggered flag when session_pnl crosses -// dd_limit. Single batch with one large loss drives session_pnl below -// the limit; one launch flips RL_SESSION_DD_TRIGGERED_INDEX to 1.0. +// CMDP gates Layer-1 state is per-batch (one independent backtest session +// per b). The four tests below exercise the per-batch buffers via b_size=1 +// (one account) to match the single-account semantics of the limits. + +/// Write a single value into a per-batch f32 device buffer through the +/// canonical mapped-pinned path (no raw HtoD on un-pinned slices per +/// `feedback_no_htod_htoh_only_mapped_pinned`). +fn write_pb( + stream: &Arc, + dst: &mut CudaSlice, + values: &[f32], +) -> Result<()> { + write_slice_f32_d_pub(stream, values, dst)?; + Ok(()) +} + +// G1 — DD breaker sets sticky triggered flag when an account's running +// session_pnl crosses dd_limit. b=1 single-account semantics; one launch +// with a -$4000 PnL delta drives session_pnl_per_batch[0] below the limit. #[test] #[ignore = "requires CUDA (MlDevice::cuda(0))"] fn g1_cmdp_dd_breaker_triggers_at_limit() -> Result<()> { let Some((dev, trainer)) = build_trainer() else { return Ok(()) }; let stream = dev.cuda_stream()?.clone(); - - // Confirm bootstrap seeding. sync(&trainer); - assert_eq!(trainer.read_isv_host(RL_SESSION_PNL_USD_INDEX), 0.0); - assert_eq!(trainer.read_isv_host(RL_SESSION_DD_TRIGGERED_INDEX), 0.0); + let dd_limit = trainer.read_isv_host(RL_SESSION_DD_LIMIT_USD_INDEX); assert!((dd_limit - SESSION_DD_LIMIT_USD).abs() < 1e-3); - // Single batch step delivering pnl = -4000 (below -3500 limit). + // Per-batch buffers start at 0 (alloc_zeros); confirm summary slots match. + assert_eq!(trainer.read_isv_host(RL_SESSION_PNL_USD_INDEX), 0.0); + assert_eq!(trainer.read_isv_host(RL_SESSION_DD_TRIGGERED_INDEX), 0.0); + let rewards_d = upload_f32(&stream, &[-4000.0])?; let dones_d = upload_f32(&stream, &[1.0])?; trainer.launch_rl_cmdp_constraints_check(&rewards_d, &dones_d, 1)?; sync(&trainer); + let pnl_summary = trainer.read_isv_host(RL_SESSION_PNL_USD_INDEX); assert!( - (trainer.read_isv_host(RL_SESSION_PNL_USD_INDEX) - (-4000.0)).abs() < 1e-3, - "session_pnl accumulated: expected -4000, got {}", - trainer.read_isv_host(RL_SESSION_PNL_USD_INDEX) + (pnl_summary - (-4000.0)).abs() < 1e-3, + "worst per-account PnL summary should be -4000; got {pnl_summary}" ); assert_eq!( trainer.read_isv_host(RL_SESSION_DD_TRIGGERED_INDEX), 1.0, - "dd_triggered must flip to 1.0 when session_pnl < dd_limit" + "any-DD-triggered summary must flip to 1.0 when the account breaches" ); - eprintln!("G1 OK — DD breaker tripped at session_pnl = -4000 < dd_limit = -3500"); + eprintln!("G1 OK — per-batch DD breaker tripped; summary mirrors worst account"); Ok(()) } -// G2 — Cooldown starts when consec_loss_count crosses limit. Feed 10 -// done-losses in a single launch (b_size=10) and verify cooldown_remaining -// snaps to cooldown_duration while consec resets to 0. +// G2 — Cooldown starts when an account's consec_loss streak crosses limit. +// b=1, 10 successive done-loss launches drive the per-batch streak to 10 +// → cooldown_remaining_per_batch[0] snaps to cooldown_duration; the +// kernel also resets the per-batch streak to 0 on the same launch. #[test] #[ignore = "requires CUDA (MlDevice::cuda(0))"] fn g2_cmdp_cooldown_starts_after_consec_loss_limit() -> Result<()> { let Some((dev, trainer)) = build_trainer() else { return Ok(()) }; let stream = dev.cuda_stream()?.clone(); - // Pre-condition: consec=0, cooldown_remaining=0, limit=10, duration=500. sync(&trainer); assert_eq!(trainer.read_isv_host(RL_CONSEC_LOSS_COUNT_INDEX), 0.0); assert_eq!(trainer.read_isv_host(RL_COOLDOWN_REMAINING_STEPS_INDEX), 0.0); - // 10 done-losses (one per batch element). DD limit not breached - // (per-loss ≈ -10, total -100 ≫ -3500), so we isolate the consec arm. - let rewards_d = upload_f32(&stream, &[-10.0_f32; 10])?; - let dones_d = upload_f32(&stream, &[1.0_f32; 10])?; - trainer.launch_rl_cmdp_constraints_check(&rewards_d, &dones_d, 10)?; + // 10 sequential done-losses on the same account. Per-loss -$10, total + // -$100 ≫ -$3500 limit, so we isolate the consec-streak arm. + let rewards_d = upload_f32(&stream, &[-10.0])?; + let dones_d = upload_f32(&stream, &[1.0])?; + for _ in 0..10 { + trainer.launch_rl_cmdp_constraints_check(&rewards_d, &dones_d, 1)?; + } sync(&trainer); let cooldown = trainer.read_isv_host(RL_COOLDOWN_REMAINING_STEPS_INDEX); let consec = trainer.read_isv_host(RL_CONSEC_LOSS_COUNT_INDEX); + // The 10th launch trips the limit. Kernel sets per_batch_cooldown[0]=duration + // then decrement *also runs this step* — so summary reads (duration - 1). + // Spec is "limit hit → start cooldown" but the decrement happens above + // the streak-check in the kernel ordering, so a single launch nets to + // exactly `duration` (no decrement yet because cooldown was 0 entering + // the launch). Confirm against `COOLDOWN_DURATION` directly. assert!( (cooldown - COOLDOWN_DURATION).abs() < 1e-3, - "cooldown_remaining must snap to duration={COOLDOWN_DURATION}; got {cooldown}" + "cooldown_remaining must snap to duration={COOLDOWN_DURATION} on the 10th loss; got {cooldown}" ); assert_eq!(consec, 0.0, "consec_loss_count must reset on limit hit"); eprintln!( - "G2 OK — 10 losses in single step → cooldown set to {cooldown}, consec reset to {consec}" + "G2 OK — 10 consecutive losses on one account → cooldown = {cooldown}, streak reset to {consec}" ); Ok(()) } -// G3 — Cooldown decrements one step per launch (independent of consec arm). -// We bypass the consec gate by sending one win plus a pre-seeded -// cooldown_remaining and verify exactly one decrement per launch. +// G3 — Cooldown decrements one step per launch on the affected account. +// Seed cooldown_remaining_per_batch[0]=50, run 3 idle launches, expect 47. #[test] #[ignore = "requires CUDA (MlDevice::cuda(0))"] fn g3_cmdp_cooldown_decrements_per_step() -> Result<()> { - let Some((dev, trainer)) = build_trainer() else { return Ok(()) }; + let Some((dev, mut trainer)) = build_trainer() else { return Ok(()) }; let stream = dev.cuda_stream()?.clone(); - trainer.isv_mapped.write_record(RL_COOLDOWN_REMAINING_STEPS_INDEX, 50.0); + // Seed the per-batch cooldown directly (the ISV summary slot is + // kernel-OUT only, not an input). + write_pb(&stream, &mut trainer.cooldown_remaining_per_batch_d, &[50.0])?; - // Three idle launches with a single break-even close (r=0, done=1) — - // neither increments nor resets consec, just exercises the decrement. + // Three idle launches with a break-even close — exercises the + // decrement without touching the streak counter. let rewards_d = upload_f32(&stream, &[0.0])?; let dones_d = upload_f32(&stream, &[1.0])?; for _ in 0..3 { @@ -168,15 +193,16 @@ fn g3_cmdp_cooldown_decrements_per_step() -> Result<()> { Ok(()) } -// G4 — A done-win resets consec_loss_count to 0 (canonical streak break). -// Seed consec=5, feed one win → counter must reset before the limit check. +// G4 — A done-win resets that account's consec_loss streak to 0. +// Seed consec_loss_per_batch[0]=5, feed one win → counter resets before +// the limit check. #[test] #[ignore = "requires CUDA (MlDevice::cuda(0))"] fn g4_cmdp_win_resets_consec_loss_counter() -> Result<()> { - let Some((dev, trainer)) = build_trainer() else { return Ok(()) }; + let Some((dev, mut trainer)) = build_trainer() else { return Ok(()) }; let stream = dev.cuda_stream()?.clone(); - trainer.isv_mapped.write_record(RL_CONSEC_LOSS_COUNT_INDEX, 5.0); + write_pb(&stream, &mut trainer.consec_loss_per_batch_d, &[5.0])?; let rewards_d = upload_f32(&stream, &[25.0])?; let dones_d = upload_f32(&stream, &[1.0])?; trainer.launch_rl_cmdp_constraints_check(&rewards_d, &dones_d, 1)?;