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

fix(rl): CMDP gates per-batch (one independent session per b)

Browse Source 
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).
This commit is contained in:
jgrusewski
2026-05-30 22:06:22 +02:00
parent 6e0f568160
commit 39efacf77d
4 changed files with 180 additions and 79 deletions
Download Patch File Download Diff File Expand all files Collapse all files

17
crates/ml-alpha/cuda/actions_to_market_targets.cu
View File

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

104
crates/ml-alpha/cuda/rl_cmdp_constraints_check.cu
View File

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

42
crates/ml-alpha/src/trainer/integrated.rs
View File

@@ -743,6 +743,16 @@ pub struct IntegratedTrainer {
rl_write_u64_fn: CudaFunction,
ts_ns_d: CudaSlice<u64>,
pub outcome_ema_d: CudaSlice<f32>,
/// 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<f32>,
pub session_dd_triggered_per_batch_d: CudaSlice<f32>,
pub consec_loss_per_batch_d: CudaSlice<f32>,
pub cooldown_remaining_per_batch_d: CudaSlice<f32>,
pub trade_context_d: CudaSlice<f32>,
pub multires_output_d: CudaSlice<f32>,
multires_state_d: CudaSlice<f32>,
@@ -1983,6 +1993,21 @@ impl IntegratedTrainer {
let outcome_ema_d = stream
.alloc_zeros::<f32>(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::<f32>(b_size)
.context("alloc session_pnl_per_batch_d")?;
let session_dd_triggered_per_batch_d = stream
.alloc_zeros::<f32>(b_size)
.context("alloc session_dd_triggered_per_batch_d")?;
let consec_loss_per_batch_d = stream
.alloc_zeros::<f32>(b_size)
.context("alloc consec_loss_per_batch_d")?;
let cooldown_remaining_per_batch_d = stream
.alloc_zeros::<f32>(b_size)
.context("alloc cooldown_remaining_per_batch_d")?;
let trade_context_d = stream
.alloc_zeros::<f32>(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();

96
crates/ml-alpha/tests/risk_stack_invariants.rs
View File

@@ -70,88 +70,113 @@ fn upload_f32(stream: &Arc<CudaStream>, host: &[f32]) -> Result<CudaSlice<f32>>
// 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<CudaStream>,
dst: &mut CudaSlice<f32>,
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)?;