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

feat(sp22-vnext): Phase F-3c — HEALTH_DIAG snap + console line for K=3 CE EMA

Browse Source 
Completes the F-3 observability story. Smoke runs now print
aux_trade_outcome_ce_ema every epoch in the standard HEALTH_DIAG output
— no ISV inspector required.

Changes:

health_diag.rs:
- New field aux_trade_outcome_ce: f32 appended at END of
  HealthDiagSnapshot per "Field order is stable; adding fields
  appends to the end" doc rule. Existing fields' byte offsets
  preserved → no kernel-side word offset re-validation.
- snapshot_size_is_stable pin: 149 × 4 → 150 × 4 = 600 bytes.

health_diag_kernel.cu:
- New WORD_AUX_TRADE_OUTCOME_CE = 149 (appended at end).
- WORD_TOTAL = 150 (was 149); static_assert bumped.
- New kernel arg aux_trade_outcome_ce_idx appended after
  moe_lambda_eff_idx.
- New mirror write after the existing MoE mirrors. Stream-implicit
  ordering: aux_outcome_ce_ema_update (F-3b) fires before
  health_diag_isv_mirror, so the read picks up the just-updated EMA.

gpu_health_diag.rs:
- launch_isv_mirror gets new aux_trade_outcome_ce_idx: i32 arg.

gpu_dqn_trainer.rs:
- launch_health_diag_isv_mirror passes
  AUX_TRADE_OUTCOME_CE_EMA_INDEX as the new arg.

training_loop.rs:
- Per-epoch metrics push appends ("aux_trade_outcome_ce_ema",
  ISV[538]) to the standard out vec. Console / CSV automatically
  includes the new column.

End-to-end F-3 chain now closed:
  K=3 fwd → aux_to_loss_scalar_buf → aux_outcome_ce_ema_update
  → ISV[538] → health_diag_isv_mirror → snap.aux_trade_outcome_ce
  → training_loop metrics → console.

Operator sees CE every epoch:
- Cold-start: 0.000
- After bootstrap: ~1.098 (= ln(3))
- After training: ideally 0.5-0.7 (head learning)

Phase F end-to-end ready. The vNext stack has:
- Full GPU kernel chain (A2-A5 + D + plan-conditioning)
- Full Rust wireup (B0-B4, B4b-1/2, C-1/C-2, B5b)
- Real labels reaching trainer
- K=3 → policy via state slots + atom-shift
- End-to-end CE observability

Verification:
- cargo check -p ml clean.
- cargo test -p ml --lib → 1016/0 green (incl. bumped
  snapshot_size_is_stable byte-size pin test).

Audit: docs/dqn-wire-up-audit.md Phase F-3c section.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-14 12:22:43 +02:00
parent 02479c885d
commit cdd3dc6edb
6 changed files with 112 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files

5
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
View File

@@ -18970,6 +18970,7 @@ impl GpuDqnTrainer {
"launch_health_diag_isv_mirror: health_diag orchestrator not constructed".into()
)),
};
use crate::cuda_pipeline::sp22_isv_slots::AUX_TRADE_OUTCOME_CE_EMA_INDEX;
hd.launch_isv_mirror(
&self.stream,
self.isv_signals_dev_ptr,
@@ -18988,6 +18989,10 @@ impl GpuDqnTrainer {
MOE_EXPERT_UTIL_EMA_BASE as i32,
MOE_GATE_ENTROPY_EMA_INDEX as i32,
MOE_LAMBDA_EFF_INDEX as i32,
// SP22 H6 vNext Phase F-3c (2026-05-14): K=3 trade-outcome
// head CE EMA slot. Kernel reads ISV[538] and mirrors into
// snap.aux_trade_outcome_ce (appended field at word 149).
AUX_TRADE_OUTCOME_CE_EMA_INDEX as i32,
)
}

6
crates/ml/src/cuda_pipeline/gpu_health_diag.rs
View File

@@ -157,6 +157,9 @@ impl GpuHealthDiag {
moe_expert_util_base_idx: i32,
moe_gate_entropy_idx: i32,
moe_lambda_eff_idx: i32,
// SP22 H6 vNext Phase F-3c (2026-05-14): K=3 trade-outcome head
// CE EMA ISV slot index (= AUX_TRADE_OUTCOME_CE_EMA_INDEX = 538).
aux_trade_outcome_ce_idx: i32,
) -> Result<(), MLError> {
let snap_dev_ptr: u64 = self.snapshot.dev_ptr();
unsafe {
@@ -181,6 +184,9 @@ impl GpuHealthDiag {
.arg(&moe_expert_util_base_idx)
.arg(&moe_gate_entropy_idx)
.arg(&moe_lambda_eff_idx)
// SP22 H6 vNext Phase F-3c (2026-05-14): K=3 head's CE
// EMA ISV slot index (= 538).
.arg(&aux_trade_outcome_ce_idx)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (1, 1, 1),

21
crates/ml/src/cuda_pipeline/health_diag.rs
View File

@@ -269,6 +269,20 @@ pub struct HealthDiagSnapshot {
/// aggregator scripts) read the same 12-bin layout that already lives
/// in `MonitoringSummary`.
pub action_counts: [u32; HEALTH_DIAG_ACTION_BINS],
// ── SP22 H6 vNext Phase F-3c (2026-05-14): trade-outcome head CE EMA ────
/// K=3 trade-outcome aux head's batch-mean sparse cross-entropy EMA,
/// mirrored from `ISV[AUX_TRADE_OUTCOME_CE_EMA_INDEX=538]` by an
/// extended `health_diag_isv_mirror` kernel write. Cold-start sentinel
/// 0.0; Pearl A bootstrap replaces with the first non-zero loss
/// observation. Healthy learning: `~1.098 = ln(3)` early →
/// monotonic-ish decrease toward 0.5-0.7 over epochs. Pinned at
/// `~1.098` indicates the head can't learn (label noise / under-
/// capacity / outcomes genuinely unconditional). **Appended at end
/// of struct** per "Field order is stable" doc rule — existing
/// fields' offsets preserved so the kernel's pre-vNext word writes
/// don't need re-validation.
pub aux_trade_outcome_ce: f32,
}
impl Default for HealthDiagSnapshot {
@@ -400,11 +414,14 @@ mod tests {
// = 8+8+7+2+1+10+4+5+9+2+6+6+6+3+3+4+7+3+16+4+4+6+3+10+12 = 149 fields
// SP13 B1.0 (2026-05-05): aux block dropped from 4 → 3 floats
// (retired `aux_label_scale` together with ISV[117]).
let expected_bytes = 149 * 4;
// SP22 H6 vNext Phase F-3c (2026-05-14): +1 field at end —
// `aux_trade_outcome_ce: f32` (K=3 head CE EMA mirror). Total
// now 150 × 4 = 600 bytes. Matches kernel WORD_TOTAL = 150.
let expected_bytes = 150 * 4;
assert_eq!(
std::mem::size_of::<HealthDiagSnapshot>(),
expected_bytes,
"HealthDiagSnapshot must be {} bytes (149 × 4); update test if struct changes",
"HealthDiagSnapshot must be {} bytes (150 × 4); update test if struct changes",
expected_bytes,
);
}

23
crates/ml/src/cuda_pipeline/health_diag_kernel.cu
View File

@@ -197,10 +197,15 @@
/* Action-count histogram (12 words, [137..149), all u32). */
#define WORD_ACTION_COUNTS_BASE 137
#define WORD_TOTAL 149
/* SP22 H6 vNext Phase F-3c (2026-05-14): K=3 trade-outcome head CE EMA,
* appended at end of struct per "Field order is stable" doc rule. Word
* 149 (= old WORD_TOTAL). Snap layout total bumps to 150. */
#define WORD_AUX_TRADE_OUTCOME_CE 149
#define WORD_TOTAL 150
/* Pin the layout: matches `snapshot_size_is_stable` on the Rust side. */
static_assert(WORD_TOTAL == 149,
static_assert(WORD_TOTAL == 150,
"HealthDiagSnapshot layout drift — update kernel WORD_* and the "
"Rust struct in lockstep, then bump snapshot_size_is_stable.");
@@ -264,7 +269,12 @@ extern "C" __global__ void health_diag_isv_mirror(
* argument dropped from kernel signature. */
int moe_expert_util_base_idx, /* ISV[118], 8 contiguous slots */
int moe_gate_entropy_idx, /* ISV[126] */
int moe_lambda_eff_idx /* ISV[128] */
int moe_lambda_eff_idx, /* ISV[128] */
/* SP22 H6 vNext Phase F-3c (2026-05-14): K=3 trade-outcome head's
* sparse-CE EMA. Source: ISV[AUX_TRADE_OUTCOME_CE_EMA_INDEX=538]
* (producer: aux_outcome_ce_ema_update kernel, single-thread direct-
* to-ISV EMA with Pearl A bootstrap, α=0.05). */
int aux_trade_outcome_ce_idx /* ISV[538] */
) {
/* Single-block, single-thread contract. Mirrors aux_heads_loss_ema's
* grid guard: the work is ~22 scalar EMA reads + writes. */
@@ -309,6 +319,13 @@ extern "C" __global__ void health_diag_isv_mirror(
snap[WORD_AUX_MOE_ENT] = isv[moe_gate_entropy_idx];
snap[WORD_AUX_MOE_LAMBDA_EFF] = isv[moe_lambda_eff_idx];
/* SP22 H6 vNext Phase F-3c (2026-05-14): K=3 trade-outcome head's
* sparse-CE EMA mirrored from ISV[538] into the appended struct
* field at word 149. The producer (`aux_outcome_ce_ema_update`)
* fires every step before this mirror runs (stream-implicit
* ordering), so the read picks up the just-updated EMA. */
snap[WORD_AUX_TRADE_OUTCOME_CE] = isv[aux_trade_outcome_ce_idx];
/* No __threadfence_system() here — the finalise kernel (Phase 2E)
* issues a single fence covering all writes from the family. */
}

11
crates/ml/src/trainers/dqn/trainer/training_loop.rs
View File

@@ -10766,6 +10766,17 @@ impl DQNTrainer {
"aux_regime_ce_ema".to_owned(),
trainer.read_isv_signal_at(AUX_REGIME_CE_EMA_INDEX) as f64,
));
// SP22 H6 vNext Phase F-3c (2026-05-14): K=3 trade-outcome
// head's batch-mean sparse-CE EMA. Healthy learning: ~ln(3)
// ≈ 1.098 early → decrease toward 0.5-0.7. Pinned at ln(3)
// for many epochs indicates the head can't learn.
{
use crate::cuda_pipeline::sp22_isv_slots::AUX_TRADE_OUTCOME_CE_EMA_INDEX;
out.push((
"aux_trade_outcome_ce_ema".to_owned(),
trainer.read_isv_signal_at(AUX_TRADE_OUTCOME_CE_EMA_INDEX) as f64,
));
}
out.push((
"isv_reward_popart_ema".to_owned(),
trainer.read_isv_signal_at(REWARD_POPART_EMA_INDEX) as f64,

51
docs/dqn-wire-up-audit.md
View File

@@ -18294,3 +18294,54 @@ Completes the Phase F-3 observability chain. The kernel + ISV slot landed in F-3
- Smoke observability: F-3 + F-3b ✓
**Phase F deployment** (next): push branch + `argo-train.sh` smoke run. Decisive test of the spec's hypothesis. Spec falsification criterion (per the original spec doc): if smoke produces WR ≈ 0.4346 ± 0.005, the trade-outcome hypothesis is falsified; the conclusion would be "WR is data-determined at this bar resolution + horizon, not a model architecture issue", and future work explores different data sampling / different prediction target.
#### Phase F-3c — HEALTH_DIAG snap + console line for K=3 CE EMA (2026-05-14)
Completes the observability story. Smoke runs now print `aux_trade_outcome_ce_ema` every epoch in the standard HEALTH_DIAG output — no ISV inspector required.
**Changes**:
`health_diag.rs` (struct):
- New field `aux_trade_outcome_ce: f32` appended at END of `HealthDiagSnapshot` per the documented "Field order is stable; adding fields appends to the end" rule. Existing fields' byte offsets preserved → no kernel-side word offset re-validation needed.
- `snapshot_size_is_stable` test pin bumped 149 × 4 → 150 × 4 = 600 bytes.
`health_diag_kernel.cu`:
- New `WORD_AUX_TRADE_OUTCOME_CE = 149` (appended at end).
- `WORD_TOTAL = 150` (was 149); static_assert updated.
- New kernel arg `aux_trade_outcome_ce_idx` appended after `moe_lambda_eff_idx`.
- New mirror write `snap[WORD_AUX_TRADE_OUTCOME_CE] = isv[aux_trade_outcome_ce_idx]` after the existing MoE mirrors. Stream-implicit ordering: the `aux_outcome_ce_ema_update` producer (Phase F-3b) fires before `health_diag_isv_mirror` in the per-step training graph, so the read picks up the just-updated EMA.
`gpu_health_diag.rs`:
- `launch_isv_mirror` Rust launcher gets new `aux_trade_outcome_ce_idx: i32` arg threaded through.
`gpu_dqn_trainer.rs`:
- `launch_health_diag_isv_mirror` passes `AUX_TRADE_OUTCOME_CE_EMA_INDEX as i32` as the new arg.
`training_loop.rs`:
- Per-epoch metrics push appends `("aux_trade_outcome_ce_ema", ISV[538])` to the standard `out` vec. The downstream metrics output (CSV, console, etc.) automatically includes the new column.
**End-to-end Phase F-3 chain now closed**:
```
K=3 forward (D + B5b) → aux_to_loss_scalar_buf[0]
→ aux_outcome_ce_ema_update (F-3a kernel) → ISV[538]
→ health_diag_isv_mirror (F-3c kernel ext) → snap.aux_trade_outcome_ce
→ training_loop metrics push (F-3c Rust ext) → console output
```
Operator now sees CE every epoch:
- Cold-start: `aux_trade_outcome_ce_ema=0.000`
- After bootstrap: `aux_trade_outcome_ce_ema≈1.098` (= ln(3))
- After training: ideally `aux_trade_outcome_ce_ema≈0.5-0.7`
**Verification**:
- `cargo check -p ml` clean.
- `cargo test -p ml --lib` → **1016/0 green** (including the bumped `snapshot_size_is_stable` byte-size pin test).
**Phase F end-to-end ready**. The trade-outcome vNext stack now has:
- Full GPU kernel chain (A2-A5 + D + plan-conditioning)
- Full Rust wireup (B0-B4, B4b-1/2, C-1/C-2, B5b)
- Real labels reaching the trainer (B4b-2 direct-path gather)
- K=3 → policy via TWO paths (state slots + atom-shift)
- End-to-end CE observability (F-3 + F-3b + F-3c)
Spec-falsification criterion (per docs/plans/2026-05-14-sp22-h6-vNext-trade-outcome-aux.md): if the smoke produces WR ≈ 0.4346 ± 0.005 AND CE drops meaningfully (showing the head IS learning the labels), the spec is FALSIFIED — trade-outcome prediction is learnable but doesn't influence policy effectively. If CE pinned AND WR pinned, the spec is INCONCLUSIVE — label noise or under-capacity, head couldn't learn the signal so we don't know if the signal would have helped.