feat(sp21): T3.1+T3.2 ISV defrost — wr_ema + hold_pct_ema (atomic)
SP21 Tier-3 foundation: defrost two production-frozen ISV slots that
pinned at 0 across every observed training epoch (d7bj7, xmd6b). Both
bugs in the same aggregator kernel, same structural shape: per-step
binary majority-vote indicators feeding fractional EMAs.
T3.1 — WR_EMA defrost (sp20_aggregate_inputs_kernel.cu):
- was: is_win_out = (2 * wins_count >= closed_count) ? 1 : 0
binary "majority won this step" indicator
- now: win_fraction_out = (float)wins_count / (float)closed_count
fractional in [0, 1]
- With actual val WR≈0.46, the binary signal was 0 most of the time,
pinning WR_EMA at 0 across 30+ epochs in xmd6b. EMA now converges
to population win rate.
T3.2 — HOLD_PCT_EMA defrost (same kernel, same pattern):
- was: action_is_hold_out = (hold_count * 2 > n_envs) ? 1 : 0
strict-majority indicator
- now: hold_fraction_out = (float)hold_count / (float)n_envs
- Cascade: HOLD_COST_SCALE controller compared hold_pct_ema=0 to
tgt±0.05, always saw < lower, ramped × 0.95 → clamped at floor
0.01 (the hold_cost_scale=0.0100 observation in d7bj7 logs).
T3.5 expected to cascade-fix in next training run.
- HOLD_REWARD_EMA gate preserves strict-majority semantic via
hold_fraction > 0.5f test in the consumer kernel.
Atomic across struct fields, kernel logic, Rust mirror, byte
serialization, doc tables, and 4 test files (per
feedback_no_partial_refactor):
- sp20_aggregate_inputs_kernel.cu (struct + 2 computations)
- sp20_emas_compute_kernel.cu (struct + reader + gate)
- sp20_aggregate_inputs.rs (doc table)
- sp20_emas_compute.rs (Rust struct + serialize + 3 tests)
- sp20_aggregate_inputs_test.rs (reader sig + 4 test assertions)
- sp20_emas_compute_test.rs (5 struct literal updates)
- sp20_phase1_4_wireup_test.rs (HOLD_PCT_EMA expected 0.625)
Verification:
- cargo check -p ml --tests: passes (warnings only)
- cargo test -p ml --lib sp20_emas: 6/6 unit tests pass
Pearl candidate: binary-majority aggregator over a fractional
underlying signal cannot serve as input to a fractional-target EMA.
The previous fix (commit 64bbbe418) addressed the per-bar vs segment
predicate at the producer site, but didn't notice the aggregator's
binarization step still collapsed the fraction to {0, 1}. Two bugs
in series, both now resolved.
Plan reference: docs/plans/2026-05-10-sp21-train-eval-coherence-isv-defrost.md
Tiers 3.3-3.6 remaining; T3.5 expected to cascade-fix.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -14,9 +14,9 @@
|
||||
//! | Field | Rule |
|
||||
//! |------------------------|---------------------------------------------------------------------------------------|
|
||||
//! | `is_close` | `1` if any env's `trade_close[env] != 0`, else `0` |
|
||||
//! | `is_win` | `1` if `is_close == 1` AND fraction of closed envs with `step_ret > 0` ≥ 0.5, else `0` |
|
||||
//! | `win_fraction` | `wins_count / closed_count ∈ [0, 1]` when `is_close == 1`, else `0.0` (SP21 T3.1: was binary majority `is_win`) |
|
||||
//! | `trade_duration` | `round(mean(hold_at_exit) over closed envs)` if `is_close == 1`, else `0` |
|
||||
//! | `action_is_hold` | `1` if `count(decoded_dir == HOLD) > n_envs/2`, else `0` |
|
||||
//! | `hold_fraction` | `hold_count / n_envs ∈ [0, 1]` (SP21 T3.2: was binary `action_is_hold` strict-majority vote) |
|
||||
//! | `alpha` | `mean(alpha_per_env) over closed envs` if `is_close == 1`, else `0` — **SP20 Phase 2 Task 2.2** (R_event - hold_baseline). |
|
||||
//! | `per_bar_hold_reward` | `0.0` — **Phase 3.2 forward reference** (-aux_conf * cost_scale) |
|
||||
//! | `aux_logits_p50` | `aux_logits_p50_dev[0]` (sp20_stats output) |
|
||||
|
||||
@@ -104,9 +104,22 @@
|
||||
* `#[repr(C)]`) ensures both kernels see the same byte layout. */
|
||||
struct SP20EmaInputs {
|
||||
int is_close;
|
||||
int is_win;
|
||||
/* SP21 T3.1 (2026-05-10): fractional win rate ∈ [0, 1] = wins_count /
|
||||
* closed_count when is_close == 1, else 0.0f. Replaces the prior
|
||||
* binary `int is_win` (majority-vote) field which structurally pinned
|
||||
* WR_EMA at 0 whenever true WR < 0.5 (the binary "majority won this
|
||||
* step" indicator was 0 most of the time when actual win rate was
|
||||
* < 50%). The EMA blends over this fractional value so WR_EMA
|
||||
* converges to the true population win rate. */
|
||||
float win_fraction;
|
||||
int trade_duration;
|
||||
int action_is_hold;
|
||||
/* SP21 T3.2 (2026-05-10): fractional hold rate ∈ [0, 1] = hold_count /
|
||||
* n_envs. Replaces the prior binary `int action_is_hold` (strict
|
||||
* majority vote) which structurally pinned HOLD_PCT_EMA at 0 whenever
|
||||
* fewer than half the envs were in Hold (the canonical case during
|
||||
* directional trading). HOLD_REWARD_EMA's gate preserves the strict-
|
||||
* majority semantic via `hold_fraction > 0.5f` test in the consumer. */
|
||||
float hold_fraction;
|
||||
float alpha;
|
||||
float per_bar_hold_reward;
|
||||
float aux_logits_p50;
|
||||
@@ -283,12 +296,22 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel(
|
||||
|
||||
const int is_close_out = (closed_count > 0) ? 1 : 0;
|
||||
|
||||
/* is_win: fraction of closed envs that won >= 0.5. */
|
||||
int is_win_out = 0;
|
||||
/* SP21 T3.1 (2026-05-10): fractional win rate, NOT binary majority.
|
||||
* `win_fraction = wins_count / closed_count ∈ [0, 1]`. The prior
|
||||
* binary majority-vote `(2 * wins_count >= closed_count)` predicate
|
||||
* was structurally incapable of representing WR < 0.5 — it emitted
|
||||
* 0 whenever fewer than half the closed envs won, regardless of
|
||||
* the actual ratio, pinning WR_EMA at 0 in production whenever
|
||||
* true WR < 0.5 (the canonical case). The downstream EMA blends
|
||||
* over this float directly so WR_EMA converges to the population
|
||||
* win rate.
|
||||
*
|
||||
* Emit 0.0f when `is_close_out == 0` for cleanliness — the consumer
|
||||
* gates the blend on `is_close == 1` so the value is moot in that
|
||||
* branch. */
|
||||
float win_fraction_out = 0.0f;
|
||||
if (is_close_out == 1) {
|
||||
/* (wins_count / closed_count) >= 0.5
|
||||
* ⇔ 2 * wins_count >= closed_count (avoids fp). */
|
||||
is_win_out = (2 * wins_count >= closed_count) ? 1 : 0;
|
||||
win_fraction_out = (float)wins_count / (float)closed_count;
|
||||
}
|
||||
|
||||
/* trade_duration: round(mean of hold_at_exit over closed envs). */
|
||||
@@ -311,9 +334,18 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel(
|
||||
alpha_out = alpha_sum_f / (float)closed_count;
|
||||
}
|
||||
|
||||
/* action_is_hold: majority-vote over envs. > n/2 threshold
|
||||
* matches the spec (strict majority, ties ⇒ 0). */
|
||||
const int action_is_hold_out = (hold_count * 2 > n_envs) ? 1 : 0;
|
||||
/* SP21 T3.2 (2026-05-10): fractional hold rate, NOT binary majority.
|
||||
* `hold_fraction = hold_count / n_envs ∈ [0, 1]`. The prior strict-
|
||||
* majority `(hold_count * 2 > n_envs)` predicate was structurally
|
||||
* incapable of representing hold rates < 0.5 — it emitted 0 whenever
|
||||
* fewer than half the envs were in Hold, pinning HOLD_PCT_EMA at 0
|
||||
* during directional trading and cascade-pinning HOLD_COST_SCALE at
|
||||
* its floor (the controller saw `hold_pct_ema < lower` every step
|
||||
* and ramped × 0.95 to clamp). Downstream HOLD_REWARD_EMA gate
|
||||
* preserves the strict-majority semantic via `hold_fraction > 0.5f`. */
|
||||
const float hold_fraction_out = (n_envs > 0)
|
||||
? (float)hold_count / (float)n_envs
|
||||
: 0.0f;
|
||||
|
||||
/* Read the upstream stats / dir-acc outputs. Stream-ordering
|
||||
* invariant guarantees these are visible to this kernel's
|
||||
@@ -330,9 +362,9 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel(
|
||||
* reference until Component 2 lands the per-bar Hold opp-cost
|
||||
* dual emission. */
|
||||
out_inputs->is_close = is_close_out;
|
||||
out_inputs->is_win = is_win_out;
|
||||
out_inputs->win_fraction = win_fraction_out;
|
||||
out_inputs->trade_duration = trade_duration_out;
|
||||
out_inputs->action_is_hold = action_is_hold_out;
|
||||
out_inputs->hold_fraction = hold_fraction_out; /* SP21 T3.2 */
|
||||
out_inputs->alpha = alpha_out; /* Task 2.2 — was 0.0f */
|
||||
out_inputs->per_bar_hold_reward = 0.0f; /* Phase 3.2 fwd ref */
|
||||
out_inputs->aux_logits_p50 = p50_in;
|
||||
|
||||
@@ -19,10 +19,10 @@
|
||||
//! | EMA | Storage | Gate | Observation source |
|
||||
//! |-----|---------|------|--------------------|
|
||||
//! | `ALPHA_EMA` | ISV[511] | `is_close == 1` | `alpha = R_event - hold_baseline` |
|
||||
//! | `WR_EMA` | ISV[512] | `is_close == 1` | `is_win` ∈ {0, 1} |
|
||||
//! | `WR_EMA` | ISV[512] | `is_close == 1` | `win_fraction` ∈ [0, 1] (SP21 T3.1) |
|
||||
//! | `TRADE_DUR` | internal[0]| `is_close == 1` | `trade_duration` (bars) |
|
||||
//! | `HOLD_PCT_EMA` | ISV[515] | per-step | `(action_is_hold == 1) ? 1 : 0` |
|
||||
//! | `HOLD_REWARD_EMA`| ISV[516] | `action_is_hold == 1` | `r_per_bar = -aux_conf * cost_scale` |
|
||||
//! | `HOLD_PCT_EMA` | ISV[515] | per-step | `hold_fraction ∈ [0, 1]` (SP21 T3.2) |
|
||||
//! | `HOLD_REWARD_EMA`| ISV[516] | `hold_fraction > 0.5` | `r_per_bar = -aux_conf * cost_scale` |
|
||||
//! | `AUX_P50` | internal[1]| per-step | `aux_logits_p50` (Phase 1.1 output) |
|
||||
//! | `AUX_STD` | internal[2]| per-step | `aux_logits_std` (Phase 1.1 output) |
|
||||
//! | `AUX_DIR_ACC` | internal[3]| per-step | `aux_dir_acc` ∈ {0.0, 1.0} |
|
||||
@@ -147,21 +147,30 @@ pub const OBS_COUNT_AUX_DIR_ACC: usize = 7;
|
||||
/// Field semantics:
|
||||
/// - `is_close`: 1 if a trade closed this step, 0 otherwise. Gates
|
||||
/// the per-trade EMAs (WR, ALPHA, TRADE_DURATION).
|
||||
/// - `is_win`: 1 if the closed trade was a win, 0 otherwise.
|
||||
/// Only meaningful when `is_close == 1`. Stored as i32 to match
|
||||
/// the device-struct field width.
|
||||
/// - `win_fraction`: fractional win rate ∈ [0, 1] = wins_count /
|
||||
/// closed_count, only meaningful when `is_close == 1`. SP21 T3.1
|
||||
/// (2026-05-10): replaced the prior binary `is_win` (i32 majority
|
||||
/// vote) — the aggregator now emits the actual population win rate
|
||||
/// as a float so WR_EMA converges to the true value. The prior
|
||||
/// binary indicator structurally pinned WR_EMA at 0 whenever true
|
||||
/// WR < 0.5.
|
||||
/// - `trade_duration`: bars in the closed trade. Only meaningful
|
||||
/// when `is_close == 1`. Stored as i32 (rounded mean across envs
|
||||
/// in the production aggregation path).
|
||||
/// - `action_is_hold`: 1 if the (majority) action this step was
|
||||
/// Hold, 0 otherwise. Gates the HOLD_REWARD_EMA; also drives the
|
||||
/// HOLD_PCT_EMA observation.
|
||||
/// - `hold_fraction`: fractional hold rate ∈ [0, 1] = hold_count /
|
||||
/// n_envs. SP21 T3.2 (2026-05-10): replaced the prior binary
|
||||
/// `action_is_hold` (i32 strict-majority vote). HOLD_PCT_EMA
|
||||
/// observes this directly so the EMA converges to the actual hold
|
||||
/// rate, not the long-run probability of strict-majority hold.
|
||||
/// HOLD_REWARD_EMA gate preserves the strict-majority semantic via
|
||||
/// `hold_fraction > 0.5f` test in the consumer kernel.
|
||||
/// - `alpha`: `R_event - hold_baseline` from the SP20 spec
|
||||
/// Component 1 reward kernel. Phase 2 wires the real producer;
|
||||
/// Phase 1.4 wire-up populates this with sentinel 0.0 (forward
|
||||
/// reference, see audit doc + `sp20_aggregate_inputs_kernel.cu`).
|
||||
/// - `per_bar_hold_reward`: per-bar Hold reward = `-aux_conf *
|
||||
/// cost_scale`. Only meaningful when `action_is_hold == 1`.
|
||||
/// cost_scale`. Only meaningful when `hold_fraction > 0.5` (the
|
||||
/// strict-majority gate preserved from pre-SP21-T3.2 semantics).
|
||||
/// Phase 3.2 wires the real producer; Phase 1.4 populates this
|
||||
/// with sentinel 0.0 (forward reference).
|
||||
/// - `aux_logits_p50`: per-step `aux_logits_p50` from Phase 1.1's
|
||||
@@ -176,9 +185,9 @@ pub const OBS_COUNT_AUX_DIR_ACC: usize = 7;
|
||||
#[derive(Copy, Clone, Debug, Default)]
|
||||
pub struct EmaInputs {
|
||||
pub is_close: i32,
|
||||
pub is_win: i32,
|
||||
pub win_fraction: f32,
|
||||
pub trade_duration: i32,
|
||||
pub action_is_hold: i32,
|
||||
pub hold_fraction: f32,
|
||||
pub alpha: f32,
|
||||
pub per_bar_hold_reward: f32,
|
||||
pub aux_logits_p50: f32,
|
||||
@@ -215,15 +224,14 @@ pub fn pack_inputs_into_f32_view(out: &mut [f32], inputs: &EmaInputs) {
|
||||
"pack_inputs_into_f32_view: out len ({}) < SP20_EMA_INPUTS_F32_LEN ({})",
|
||||
out.len(), SP20_EMA_INPUTS_F32_LEN,
|
||||
);
|
||||
// Match the device-side `SP20EmaInputs` layout: 4 i32 + 5 f32.
|
||||
// Bit-pattern transmute: the device kernel reads the f32-aliased
|
||||
// buffer as i32 fields by struct offset (the .cu file's
|
||||
// `int is_close` is a 4-byte read off the same offset our
|
||||
// `f32::from_bits(...)` write lands at).
|
||||
// Match the device-side `SP20EmaInputs` layout: 2 i32 + 7 f32 (SP21
|
||||
// T3.1 + T3.2 converted is_win → win_fraction and action_is_hold →
|
||||
// hold_fraction). Bit-pattern transmute for the i32 fields; native
|
||||
// f32 stores for the rest.
|
||||
out[0] = f32::from_bits(inputs.is_close as u32);
|
||||
out[1] = f32::from_bits(inputs.is_win as u32);
|
||||
out[1] = inputs.win_fraction;
|
||||
out[2] = f32::from_bits(inputs.trade_duration as u32);
|
||||
out[3] = f32::from_bits(inputs.action_is_hold as u32);
|
||||
out[3] = inputs.hold_fraction;
|
||||
out[4] = inputs.alpha;
|
||||
out[5] = inputs.per_bar_hold_reward;
|
||||
out[6] = inputs.aux_logits_p50;
|
||||
@@ -354,7 +362,7 @@ mod tests {
|
||||
#[test]
|
||||
fn ema_inputs_default_is_safe_sentinel() {
|
||||
// A `Default::default()` `EmaInputs` has all gates closed
|
||||
// (is_close = 0, action_is_hold = 0) and all observations at
|
||||
// (is_close = 0, hold_fraction = 0.0) and all observations at
|
||||
// 0 — equivalent to "no firing this step" for all 8 EMAs.
|
||||
// This is the canonical pre-warmup / cold-start input shape.
|
||||
let inp = EmaInputs::default();
|
||||
@@ -363,7 +371,7 @@ mod tests {
|
||||
// is correct behavior for cold-start, and is verified by the
|
||||
// `per_step_emas_fire_unconditionally` GPU oracle test.
|
||||
assert_eq!(inp.is_close, 0);
|
||||
assert_eq!(inp.action_is_hold, 0);
|
||||
assert_eq!(inp.hold_fraction, 0.0);
|
||||
}
|
||||
|
||||
#[test]
|
||||
@@ -381,11 +389,12 @@ mod tests {
|
||||
fn pack_inputs_round_trips() {
|
||||
// Sanity: pack into a 9-f32 buffer and read back the bit
|
||||
// patterns; verify i32 fields land in the right slots.
|
||||
// SP21 T3.1: win_fraction is now a native f32 in slot [1].
|
||||
let inp = EmaInputs {
|
||||
is_close: 1,
|
||||
is_win: 1,
|
||||
win_fraction: 0.46,
|
||||
trade_duration: 7,
|
||||
action_is_hold: 0,
|
||||
hold_fraction: 0.0,
|
||||
alpha: 0.5,
|
||||
per_bar_hold_reward: -0.05,
|
||||
aux_logits_p50: 0.2,
|
||||
@@ -396,9 +405,11 @@ mod tests {
|
||||
pack_inputs_into_f32_view(&mut buf, &inp);
|
||||
// Verify i32 slots come back via to_bits round trip.
|
||||
assert_eq!(buf[0].to_bits() as i32, 1);
|
||||
assert_eq!(buf[1].to_bits() as i32, 1);
|
||||
// win_fraction is a native f32 — direct comparison.
|
||||
assert!((buf[1] - 0.46).abs() < 1e-6);
|
||||
assert_eq!(buf[2].to_bits() as i32, 7);
|
||||
assert_eq!(buf[3].to_bits() as i32, 0);
|
||||
// hold_fraction is a native f32 — direct comparison.
|
||||
assert!(buf[3].abs() < 1e-6);
|
||||
// Float slots are direct.
|
||||
assert_eq!(buf[4], 0.5);
|
||||
assert_eq!(buf[5], -0.05);
|
||||
|
||||
@@ -134,9 +134,15 @@
|
||||
* each 4 bytes. */
|
||||
struct SP20EmaInputs {
|
||||
int is_close; /* 0 or 1 — gates per-trade EMAs */
|
||||
int is_win; /* 0 or 1 — only meaningful when is_close == 1 */
|
||||
/* SP21 T3.1 (2026-05-10): fractional win rate ∈ [0, 1] = wins_count /
|
||||
* closed_count when is_close == 1, else 0.0f. Replaces the prior
|
||||
* binary `int is_win` (majority-vote) field. */
|
||||
float win_fraction;
|
||||
int trade_duration; /* bars — only meaningful when is_close == 1 */
|
||||
int action_is_hold; /* 0 or 1 — gates HOLD_REWARD_EMA */
|
||||
/* SP21 T3.2 (2026-05-10): fractional hold rate ∈ [0, 1] = hold_count /
|
||||
* n_envs. HOLD_PCT_EMA observes this directly. HOLD_REWARD_EMA gate
|
||||
* preserves strict-majority semantic via `hold_fraction > 0.5f`. */
|
||||
float hold_fraction;
|
||||
float alpha; /* R_event - hold_baseline at trade close */
|
||||
float per_bar_hold_reward; /* -aux_conf * cost_scale on Hold bars */
|
||||
float aux_logits_p50; /* upstream sp20_stats output [0] */
|
||||
@@ -207,22 +213,29 @@ extern "C" __global__ void sp20_emas_compute_kernel(
|
||||
* coalescing is moot, but the explicit local copies make the
|
||||
* dependency chain readable. */
|
||||
const int is_close = ema_inputs->is_close;
|
||||
const int is_win_i = ema_inputs->is_win;
|
||||
/* SP21 T3.1 (2026-05-10): read fractional win rate as float directly.
|
||||
* The aggregator emits `wins_count / closed_count ∈ [0, 1]` so the
|
||||
* EMA converges to the population win rate. Was previously `int
|
||||
* is_win` (binary majority vote) which structurally pinned WR_EMA
|
||||
* at 0 whenever true WR < 0.5. */
|
||||
const float win_fraction = ema_inputs->win_fraction;
|
||||
const int trade_duration = ema_inputs->trade_duration;
|
||||
const int action_is_hold = ema_inputs->action_is_hold;
|
||||
/* SP21 T3.2 (2026-05-10): fractional hold rate. HOLD_PCT_EMA observes
|
||||
* this directly; HOLD_REWARD_EMA gate uses `hold_fraction > 0.5f` to
|
||||
* preserve the prior strict-majority semantic. */
|
||||
const float hold_fraction = ema_inputs->hold_fraction;
|
||||
const float alpha = ema_inputs->alpha;
|
||||
const float r_per_bar = ema_inputs->per_bar_hold_reward;
|
||||
const float aux_p50_in = ema_inputs->aux_logits_p50;
|
||||
const float aux_std_in = ema_inputs->aux_logits_std;
|
||||
const float aux_dir_acc_in = ema_inputs->aux_dir_acc;
|
||||
|
||||
const float is_win = (float)is_win_i;
|
||||
const float trade_duration_f = (float)trade_duration;
|
||||
|
||||
/* ── Per-trade EMAs: gated on is_close == 1 ───────────────────── */
|
||||
if (is_close == 1) {
|
||||
/* WR_EMA */
|
||||
isv[isv_idx_wr] = ema_step(isv[isv_idx_wr], is_win,
|
||||
/* WR_EMA — blends over win_fraction ∈ [0, 1] (SP21 T3.1). */
|
||||
isv[isv_idx_wr] = ema_step(isv[isv_idx_wr], win_fraction,
|
||||
&obs_count[OBS_IDX_WR]);
|
||||
|
||||
/* ALPHA_EMA */
|
||||
@@ -237,13 +250,16 @@ extern "C" __global__ void sp20_emas_compute_kernel(
|
||||
|
||||
/* ── Per-step EMAs: always fire ───────────────────────────────── */
|
||||
|
||||
/* HOLD_PCT_EMA — observation = (action_is_hold == 1) ? 1.0 : 0.0 */
|
||||
{
|
||||
const float hold_indicator = (action_is_hold == 1) ? 1.0f : 0.0f;
|
||||
isv[isv_idx_hold_pct] = ema_step(isv[isv_idx_hold_pct],
|
||||
hold_indicator,
|
||||
&obs_count[OBS_IDX_HOLD_PCT]);
|
||||
}
|
||||
/* HOLD_PCT_EMA — observation = hold_fraction ∈ [0, 1] (SP21 T3.2).
|
||||
* Was previously the `(action_is_hold == 1) ? 1.0 : 0.0` binary
|
||||
* indicator which structurally pinned HOLD_PCT_EMA at 0 whenever
|
||||
* the policy held positions in less than half the envs. The
|
||||
* downstream HOLD_COST_SCALE controller compares hold_pct_ema to
|
||||
* `tgt ± 0.05` to decide ramp direction; with the binary signal it
|
||||
* always saw `< lower` and clamped at floor 0.01. */
|
||||
isv[isv_idx_hold_pct] = ema_step(isv[isv_idx_hold_pct],
|
||||
hold_fraction,
|
||||
&obs_count[OBS_IDX_HOLD_PCT]);
|
||||
|
||||
/* aux_conf_p50_ema — internal scratch */
|
||||
internal[INT_IDX_AUX_P50] =
|
||||
@@ -260,8 +276,15 @@ extern "C" __global__ void sp20_emas_compute_kernel(
|
||||
ema_step(internal[INT_IDX_AUX_DIR_ACC], aux_dir_acc_in,
|
||||
&obs_count[OBS_IDX_AUX_DIR_ACC]);
|
||||
|
||||
/* ── Hold-reward EMA: gated on action_is_hold == 1 ────────────── */
|
||||
if (action_is_hold == 1) {
|
||||
/* ── Hold-reward EMA: gated on hold_fraction > 0.5 (SP21 T3.2). ──
|
||||
* Preserves the prior strict-majority semantic of `action_is_hold==1`
|
||||
* (`hold_count*2 > n_envs` ⇔ `hold_fraction > 0.5`). The gate stays
|
||||
* binary because `r_per_bar` is a per-step scalar (single value, not
|
||||
* per-env); reading it only on majority-Hold steps approximates the
|
||||
* "policy is in Hold this bar" semantic. T3.3 (per_bar_hold_reward
|
||||
* producer is currently 0.0, Phase 3.2 forward reference) will
|
||||
* unfreeze HOLD_REWARD_EMA fully when the producer wires real values. */
|
||||
if (hold_fraction > 0.5f) {
|
||||
isv[isv_idx_hold_reward] =
|
||||
ema_step(isv[isv_idx_hold_reward], r_per_bar,
|
||||
&obs_count[OBS_IDX_HOLD_REWARD]);
|
||||
|
||||
@@ -69,23 +69,26 @@ mod gpu {
|
||||
}
|
||||
|
||||
/// Read the post-launch `SP20EmaInputs` struct out of the f32-aliased
|
||||
/// buffer. Returns `(is_close, is_win, trade_duration, action_is_hold,
|
||||
/// alpha, per_bar_hold_reward, p50, std, dir_acc)`.
|
||||
fn read_ema_inputs(buf: &MappedF32Buffer) -> (i32, i32, i32, i32, f32, f32, f32, f32, f32) {
|
||||
/// buffer. Returns `(is_close, win_fraction, trade_duration,
|
||||
/// hold_fraction, alpha, per_bar_hold_reward, p50, std, dir_acc)`.
|
||||
/// SP21 T3.1+T3.2 (2026-05-10): slots 1 and 3 are native f32
|
||||
/// (`win_fraction`, `hold_fraction`) — were binary `is_win` /
|
||||
/// `action_is_hold` previously.
|
||||
fn read_ema_inputs(buf: &MappedF32Buffer) -> (i32, f32, i32, f32, f32, f32, f32, f32, f32) {
|
||||
let v = buf.read_all();
|
||||
assert_eq!(v.len(), SP20_EMA_INPUTS_F32_LEN);
|
||||
// The kernel writes int fields in the first 4 slots and float
|
||||
// fields in the last 5. We cast back from f32::to_bits.
|
||||
// 2 i32 fields + 7 f32 fields. i32 fields come back via to_bits;
|
||||
// f32 fields are direct reads.
|
||||
let is_close = v[0].to_bits() as i32;
|
||||
let is_win = v[1].to_bits() as i32;
|
||||
let win_fraction = v[1];
|
||||
let trade_duration = v[2].to_bits() as i32;
|
||||
let action_is_hold = v[3].to_bits() as i32;
|
||||
let hold_fraction = v[3];
|
||||
let alpha = v[4];
|
||||
let per_bar_hr = v[5];
|
||||
let p50 = v[6];
|
||||
let std = v[7];
|
||||
let dir_acc = v[8];
|
||||
(is_close, is_win, trade_duration, action_is_hold,
|
||||
(is_close, win_fraction, trade_duration, hold_fraction,
|
||||
alpha, per_bar_hr, p50, std, dir_acc)
|
||||
}
|
||||
|
||||
@@ -102,7 +105,7 @@ mod gpu {
|
||||
aux_logits_p50: f32,
|
||||
aux_logits_std: f32,
|
||||
aux_dir_acc: f32,
|
||||
) -> (i32, i32, i32, i32, f32, f32, f32, f32, f32) {
|
||||
) -> (i32, f32, i32, f32, f32, f32, f32, f32, f32) {
|
||||
let stream = make_test_stream();
|
||||
let kernel = load_aggregation_kernel(&stream);
|
||||
let n = trade_close.len();
|
||||
@@ -169,7 +172,7 @@ mod gpu {
|
||||
aux_logits_p50: f32,
|
||||
aux_logits_std: f32,
|
||||
aux_dir_acc: f32,
|
||||
) -> (i32, i32, i32, i32, f32, f32, f32, f32, f32) {
|
||||
) -> (i32, f32, i32, f32, f32, f32, f32, f32, f32) {
|
||||
let stream = make_test_stream();
|
||||
let kernel = load_aggregation_kernel(&stream);
|
||||
let n = trade_close.len();
|
||||
@@ -243,7 +246,7 @@ mod gpu {
|
||||
aux_logits_p50: f32,
|
||||
aux_logits_std: f32,
|
||||
aux_dir_acc: f32,
|
||||
) -> (i32, i32, i32, i32, f32, f32, f32, f32, f32) {
|
||||
) -> (i32, f32, i32, f32, f32, f32, f32, f32, f32) {
|
||||
let stream = make_test_stream();
|
||||
let kernel = load_aggregation_kernel(&stream);
|
||||
let n = trade_close.len();
|
||||
@@ -314,9 +317,9 @@ mod gpu {
|
||||
run_kernel(&trade_close, &step_ret, &hold_at_exit, &actions,
|
||||
0.3, 0.1, 0.55);
|
||||
assert_eq!(ic, 0, "is_close must be 0 (no trades closed)");
|
||||
assert_eq!(iw, 0, "is_win must be 0 when is_close=0");
|
||||
assert_eq!(iw, 0.0, "win_fraction must be 0.0 when is_close=0");
|
||||
assert_eq!(td, 0, "trade_duration must be 0 when is_close=0");
|
||||
assert_eq!(ah, 0, "action_is_hold must be 0 (all envs picked Long)");
|
||||
assert_eq!(ah, 0.0, "hold_fraction must be 0.0 (all envs picked Long)");
|
||||
assert_eq!(alpha, 0.0, "alpha is Phase 2 forward ref placeholder = 0.0");
|
||||
assert_eq!(hr, 0.0, "per_bar_hold_reward is Phase 3.2 forward ref = 0.0");
|
||||
assert!((p50 - 0.3).abs() < 1e-6, "aux_logits_p50 forwarded");
|
||||
@@ -325,9 +328,11 @@ mod gpu {
|
||||
}
|
||||
|
||||
/// Half envs win, half lose; all close. Expected:
|
||||
/// is_close = 1, is_win = 1 (4 wins / 8 closed = 0.5, threshold).
|
||||
/// is_close = 1, win_fraction = 0.5 (4 wins / 8 closed).
|
||||
/// trade_duration = round(mean(hold_at_exit)) = round(7.5) = 8
|
||||
/// (round-to-nearest-even at 7.5 ⇒ 8).
|
||||
/// SP21 T3.1: was `is_win = 1` (binary majority threshold); now
|
||||
/// `win_fraction = 0.5` exact.
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn half_wins_meets_threshold_and_duration_rounds() {
|
||||
@@ -340,7 +345,8 @@ mod gpu {
|
||||
run_kernel(&trade_close, &step_ret, &hold_at_exit, &actions,
|
||||
0.0, 0.0, 0.0);
|
||||
assert_eq!(ic, 1);
|
||||
assert_eq!(iw, 1, "fraction wins (4/8 = 0.5) meets >= 0.5 threshold");
|
||||
assert!((iw - 0.5).abs() < 1e-6,
|
||||
"win_fraction = 4/8 = 0.5 exactly; got {iw}");
|
||||
// round-to-nearest-even at 7.5 = 8 (banker's rounding); CUDA
|
||||
// __float2int_rn uses round-to-nearest-even. Either 7 or 8 are
|
||||
// structurally acceptable; the kernel uses CUDA's RN so we
|
||||
@@ -349,14 +355,15 @@ mod gpu {
|
||||
td == 7 || td == 8,
|
||||
"trade_duration must round 7.5 to 7 (down) or 8 (RN-even); got {td}",
|
||||
);
|
||||
assert_eq!(ah, 0);
|
||||
assert_eq!(ah, 0.0,
|
||||
"hold_fraction must be 0.0 (all envs picked Long, no Hold)");
|
||||
}
|
||||
|
||||
/// Strict majority Hold over n_envs. With 5 of 8 picking Hold,
|
||||
/// 5 > 8/2 = 4 ⇒ action_is_hold = 1.
|
||||
/// 5 of 8 envs picking Hold → hold_fraction = 5/8 = 0.625 (SP21 T3.2:
|
||||
/// was binary `action_is_hold = 1` under strict majority).
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn majority_hold_sets_action_is_hold() {
|
||||
fn five_of_eight_hold_yields_fraction_0_625() {
|
||||
let trade_close = vec![0; 8];
|
||||
let step_ret = vec![0.0; 8];
|
||||
let hold_at_exit = vec![0.0; 8];
|
||||
@@ -367,14 +374,19 @@ mod gpu {
|
||||
let (_, _, _, ah, _, _, _, _, _) =
|
||||
run_kernel(&trade_close, &step_ret, &hold_at_exit, &actions,
|
||||
0.0, 0.0, 0.0);
|
||||
assert_eq!(ah, 1, "5 of 8 = strict majority Hold => action_is_hold=1");
|
||||
assert!(
|
||||
(ah - 0.625).abs() < 1e-6,
|
||||
"5 of 8 Hold => hold_fraction = 5/8 = 0.625 exact; got {ah}",
|
||||
);
|
||||
}
|
||||
|
||||
/// Tie at exactly half Hold. 4 of 8 = 4 = n_envs/2; the kernel uses
|
||||
/// `count * 2 > n_envs` (strict majority), so a tie ⇒ 0.
|
||||
/// Half of envs picking Hold → hold_fraction = 4/8 = 0.5 exact (SP21
|
||||
/// T3.2: was `action_is_hold = 0` under strict majority — ties did
|
||||
/// not trigger). The HOLD_REWARD_EMA gate `hold_fraction > 0.5f`
|
||||
/// preserves the strict-majority semantic at the consumer site.
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn tie_hold_count_does_not_set_action_is_hold() {
|
||||
fn half_hold_yields_fraction_0_5_exact() {
|
||||
let trade_close = vec![0; 8];
|
||||
let step_ret = vec![0.0; 8];
|
||||
let hold_at_exit = vec![0.0; 8];
|
||||
@@ -385,11 +397,14 @@ mod gpu {
|
||||
let (_, _, _, ah, _, _, _, _, _) =
|
||||
run_kernel(&trade_close, &step_ret, &hold_at_exit, &actions,
|
||||
0.0, 0.0, 0.0);
|
||||
assert_eq!(ah, 0, "tie (4 of 8 Hold) => action_is_hold=0 (strict majority)");
|
||||
assert!(
|
||||
(ah - 0.5).abs() < 1e-6,
|
||||
"4 of 8 Hold => hold_fraction = 4/8 = 0.5 exact; got {ah}",
|
||||
);
|
||||
}
|
||||
|
||||
/// Single env closed with a loss: is_close=1, is_win=0
|
||||
/// (0/1 < 0.5), trade_duration = round(3.0) = 3.
|
||||
/// Single env closed with a loss: is_close=1, win_fraction=0.0
|
||||
/// (0/1), trade_duration = round(3.0) = 3.
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn single_env_loss_close() {
|
||||
@@ -401,7 +416,7 @@ mod gpu {
|
||||
run_kernel(&trade_close, &step_ret, &hold_at_exit, &actions,
|
||||
0.0, 0.0, 0.0);
|
||||
assert_eq!(ic, 1);
|
||||
assert_eq!(iw, 0, "loss (0 wins out of 1 closed) => is_win=0");
|
||||
assert_eq!(iw, 0.0, "loss (0 wins out of 1 closed) => win_fraction=0.0");
|
||||
assert_eq!(td, 3, "trade_duration = round(3.0) = 3");
|
||||
}
|
||||
|
||||
@@ -557,14 +572,15 @@ mod gpu {
|
||||
);
|
||||
|
||||
assert_eq!(ic, 1, "is_close = 1 (4 envs closed)");
|
||||
assert_eq!(
|
||||
iw, 1,
|
||||
"is_win MUST be 1 — segment-level wins (is_win_per_env = \
|
||||
[1,1,1,0]) put 3 of 4 closed envs in the win bucket; \
|
||||
2*3 >= 4 ⇒ majority threshold met. The legacy `step_ret \
|
||||
> 0` predicate would have given is_win = 0 (all step_rets \
|
||||
negative due to tx-cost domination at close bars), pinning \
|
||||
WR_EMA at 0 in production.",
|
||||
assert!(
|
||||
(iw - 0.75).abs() < 1e-6,
|
||||
"win_fraction MUST be 3/4 = 0.75 — segment-level wins \
|
||||
(is_win_per_env = [1,1,1,0]) put 3 of 4 closed envs in the \
|
||||
win bucket. SP21 T3.1: aggregator now emits the actual \
|
||||
fraction, not a binary majority indicator. The legacy \
|
||||
`step_ret > 0` predicate would have given win_fraction = 0 \
|
||||
(all step_rets negative due to tx-cost domination at close \
|
||||
bars), pinning WR_EMA at 0 in production. got iw={iw}",
|
||||
);
|
||||
}
|
||||
|
||||
@@ -590,12 +606,13 @@ mod gpu {
|
||||
0.0, 0.0, 0.0);
|
||||
|
||||
assert_eq!(ic, 1);
|
||||
assert_eq!(
|
||||
iw, 1,
|
||||
assert!(
|
||||
(iw - 1.0).abs() < 1e-6,
|
||||
"NULL is_win_per_env falls back to `sr > 0` legacy \
|
||||
predicate — all 4 step_rets positive ⇒ wins_count = 4 ⇒ \
|
||||
is_win = 1 (the pre-fix oracle-test contract preserved \
|
||||
via NULL-tolerance).",
|
||||
win_fraction = 4/4 = 1.0 (the pre-fix oracle-test contract \
|
||||
preserved via NULL-tolerance). SP21 T3.1: aggregator emits \
|
||||
fraction not binary majority. got iw={iw}",
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -146,9 +146,9 @@ mod gpu {
|
||||
fn first_observation_replaces_sentinel() {
|
||||
let inp = EmaInputs {
|
||||
is_close: 1,
|
||||
is_win: 1,
|
||||
win_fraction: 1.0,
|
||||
trade_duration: 7,
|
||||
action_is_hold: 0,
|
||||
hold_fraction: 0.0,
|
||||
alpha: 0.5,
|
||||
per_bar_hold_reward: 0.0,
|
||||
aux_logits_p50: 0.0,
|
||||
@@ -216,9 +216,9 @@ mod gpu {
|
||||
let inputs: Vec<EmaInputs> = (0..100)
|
||||
.map(|i| EmaInputs {
|
||||
is_close: 1,
|
||||
is_win: if i % 2 == 0 { 1 } else { 0 },
|
||||
win_fraction: if i % 2 == 0 { 1.0 } else { 0.0 },
|
||||
trade_duration: 1,
|
||||
action_is_hold: 0,
|
||||
hold_fraction: 0.0,
|
||||
alpha: 0.0,
|
||||
per_bar_hold_reward: 0.0,
|
||||
aux_logits_p50: 0.0,
|
||||
@@ -253,9 +253,9 @@ mod gpu {
|
||||
for _ in 0..5 {
|
||||
inputs.push(EmaInputs {
|
||||
is_close: 0,
|
||||
is_win: 0,
|
||||
win_fraction: 0.0,
|
||||
trade_duration: 0,
|
||||
action_is_hold: 0, // gated OFF
|
||||
hold_fraction: 0.0, // gated OFF
|
||||
alpha: 0.0,
|
||||
per_bar_hold_reward: -0.05,
|
||||
aux_logits_p50: 0.0,
|
||||
@@ -266,9 +266,9 @@ mod gpu {
|
||||
for _ in 0..5 {
|
||||
inputs.push(EmaInputs {
|
||||
is_close: 0,
|
||||
is_win: 0,
|
||||
win_fraction: 0.0,
|
||||
trade_duration: 0,
|
||||
action_is_hold: 1, // gate ON
|
||||
hold_fraction: 1.0, // > 0.5 gate ON (SP21 T3.2)
|
||||
alpha: 0.0,
|
||||
per_bar_hold_reward: -0.05,
|
||||
aux_logits_p50: 0.0,
|
||||
@@ -317,9 +317,9 @@ mod gpu {
|
||||
let inputs: Vec<EmaInputs> = (0..10)
|
||||
.map(|_| EmaInputs {
|
||||
is_close: 0,
|
||||
is_win: 0,
|
||||
win_fraction: 0.0,
|
||||
trade_duration: 0,
|
||||
action_is_hold: 0,
|
||||
hold_fraction: 0.0,
|
||||
alpha: 0.0,
|
||||
per_bar_hold_reward: 0.0,
|
||||
aux_logits_p50: 0.2,
|
||||
|
||||
@@ -254,15 +254,23 @@ mod gpu {
|
||||
assert_eq!(isv[ALPHA_EMA_INDEX], 0.0,
|
||||
"ALPHA_EMA stays at 0.0 with NULL alpha_per_env producer (Task 2.2 NULL contract)");
|
||||
|
||||
// WR_EMA: 1 win out of 1 close ⇒ is_win = 1, first-obs replace
|
||||
// ⇒ WR_EMA = 1.0.
|
||||
// WR_EMA: 1 win out of 1 close ⇒ win_fraction = 1.0, first-obs
|
||||
// replace ⇒ WR_EMA = 1.0. SP21 T3.1: aggregator emits exact
|
||||
// fraction (was binary majority).
|
||||
assert_eq!(isv[WR_EMA_INDEX], 1.0,
|
||||
"WR_EMA = 1.0 after one winning close (1 win / 1 closed >= 0.5)");
|
||||
"WR_EMA = 1.0 after one winning close (1 win / 1 closed = 1.0)");
|
||||
|
||||
// HOLD_PCT_EMA: 5 of 8 envs picked Hold ⇒ action_is_hold = 1
|
||||
// ⇒ HOLD_PCT_EMA = 1.0 on first observation.
|
||||
assert_eq!(isv[HOLD_PCT_EMA_INDEX], 1.0,
|
||||
"HOLD_PCT_EMA = 1.0 after one majority-Hold observation");
|
||||
// HOLD_PCT_EMA: 5 of 8 envs picked Hold ⇒ hold_fraction = 5/8 =
|
||||
// 0.625, first-obs replace ⇒ HOLD_PCT_EMA = 0.625. SP21 T3.2:
|
||||
// aggregator emits exact fraction (was binary majority indicator
|
||||
// → 1.0). The fractional value lets the downstream HOLD_COST_SCALE
|
||||
// controller make graded decisions instead of the cascade-
|
||||
// pinning-at-floor that the binary signal produced.
|
||||
let hold_pct = isv[HOLD_PCT_EMA_INDEX];
|
||||
assert!(
|
||||
(hold_pct - 0.625).abs() < 1e-6,
|
||||
"HOLD_PCT_EMA = 0.625 (5/8 of envs Hold); got {hold_pct}",
|
||||
);
|
||||
|
||||
// HOLD_REWARD_EMA: per_bar_hold_reward = 0.0 (Phase 3.2 fwd
|
||||
// ref); fired on Hold ⇒ EMA = 0.0.
|
||||
|
||||
@@ -2,6 +2,70 @@
|
||||
|
||||
**Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7.
|
||||
|
||||
## 2026-05-10 — SP21 Tier-3 foundation: T3.1+T3.2 ISV defrost (wr_ema, hold_pct_ema)
|
||||
|
||||
**Branch:** `sp20-aux-h-fixed` (off `6df44e4c6`).
|
||||
**Commits:** 1 atomic (this commit).
|
||||
|
||||
Defrosts two production-frozen ISV slots that pinned-at-zero across all
|
||||
observed training epochs in d7bj7/xmd6b logs. Both bugs were in the same
|
||||
aggregator kernel (`sp20_aggregate_inputs_kernel.cu`), same structural
|
||||
shape: per-step binary majority-vote indicators feeding fractional EMAs.
|
||||
|
||||
**T3.1: WR_EMA — `is_win` → `win_fraction` (struct field rename + type change).**
|
||||
The aggregator computed `is_win_out = (2 * wins_count >= closed_count) ? 1 : 0`
|
||||
— a binary "majority of closed envs won" indicator. With actual val WR
|
||||
≈ 0.46 (xmd6b 30 epochs), the binary signal was 0 most of the time, and
|
||||
WR_EMA → 0 by Wiener-α floor blend. The aggregator now emits
|
||||
`win_fraction = wins_count / closed_count ∈ [0, 1]` and the EMA blends
|
||||
over the float directly. WR_EMA will converge to the population win rate.
|
||||
|
||||
**T3.2: HOLD_PCT_EMA — `action_is_hold` → `hold_fraction` (same fix shape).**
|
||||
The aggregator computed `action_is_hold_out = (hold_count * 2 > n_envs) ? 1 : 0`
|
||||
— strict-majority indicator. With actual hold pct in [0.10, 0.56],
|
||||
majority was rarely met, HOLD_PCT_EMA → 0. **Cascade**: the
|
||||
`HOLD_COST_SCALE` controller compares `hold_pct_ema` to `tgt ± 0.05` and
|
||||
ramps; with hold_pct_ema=0 it always saw `< lower` and ramped × 0.95
|
||||
every step → clamped at floor 0.01 (the `hold_cost_scale = 0.0100`
|
||||
observation in d7bj7 sp20_isv logs). Fixing T3.2 should cascade-fix
|
||||
T3.5 — verify in next training run. HOLD_REWARD_EMA gate preserves
|
||||
strict-majority semantic via `hold_fraction > 0.5f` test in the
|
||||
consumer kernel.
|
||||
|
||||
**Affected files (atomic per `feedback_no_partial_refactor`):**
|
||||
|
||||
- `crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu` —
|
||||
struct fields, output computations.
|
||||
- `crates/ml/src/cuda_pipeline/sp20_emas_compute_kernel.cu` —
|
||||
struct fields, EMA reader, HOLD_REWARD_EMA gate.
|
||||
- `crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs` — doc table.
|
||||
- `crates/ml/src/cuda_pipeline/sp20_emas_compute.rs` — Rust struct mirror,
|
||||
byte serialization, doc table, default-sentinel test, round-trip test.
|
||||
- `crates/ml/tests/sp20_aggregate_inputs_test.rs` — reader signature,
|
||||
4 fractional-comparison test assertions, 2 renamed test fns.
|
||||
- `crates/ml/tests/sp20_emas_compute_test.rs` — 5 struct-literal field
|
||||
renames (is_win → win_fraction, action_is_hold → hold_fraction).
|
||||
- `crates/ml/tests/sp20_phase1_4_wireup_test.rs` — HOLD_PCT_EMA expected
|
||||
value updated 1.0 → 0.625 (5/8 envs Hold).
|
||||
|
||||
**Verification:**
|
||||
|
||||
- `cargo check -p ml --tests` — passes (warnings only).
|
||||
- `cargo test -p ml --lib sp20_emas` — 6/6 unit tests pass, including
|
||||
`pack_inputs_round_trips` with `win_fraction: 0.46` + `hold_fraction: 0.0`.
|
||||
- GPU oracle tests gated `#[ignore = "requires GPU"]` — compile cleanly.
|
||||
|
||||
**Pearls applied / candidates:**
|
||||
|
||||
- `pearl_first_observation_bootstrap` — unchanged, both EMAs still bootstrap
|
||||
from sentinel.
|
||||
- `pearl_wiener_alpha_floor_for_nonstationary` — unchanged, α=0.4 floor.
|
||||
- New pearl candidate: **binary-majority aggregator over a fractional
|
||||
underlying signal cannot serve as input to a fractional-target EMA**.
|
||||
When the EMA is supposed to converge to a population fraction, the
|
||||
observation must itself be the fraction or a 0/1 sample, not a
|
||||
thresholded reduction.
|
||||
|
||||
## 2026-05-10 — REVERT: aux_horizon_update_kernel sp20-aux-h-fixed experiment closed
|
||||
|
||||
Reverts the forced-H=30 diagnostic that tested the bootstrap-failure hypothesis. **Hypothesis verdict: dead.** Across two epochs of `train-d7bj7`:
|
||||
|
||||
Reference in New Issue
Block a user