Replace static moe_lambda=0.01 with a GPU-driven controller that scales
λ inversely with observed gate-entropy: λ_eff = λ_floor + λ_max_extra ×
max(0, ent_target − ent_ema)/ent_target, where ent_ema is the existing
ISV[126] producer (Phase 2 task 2.4 wire-up).
Motivation: the L40S validation run train-multi-seed-fgg9x exposed
single-expert specialization in fold 1 (expert 4 → 81% util, 7/8
experts < 5%) under static λ=0.01. The load-balance push was too weak
to prevent winner-take-all once the gate had a strong preference.
Adaptive λ rises proportional to the entropy deficit, restoring
diversity pressure when the gate concentrates.
Per pearl_blend_formulas_must_have_permanent_floor: λ_floor remains a
permanent minimum (= old static 0.01) so the controller never weakens
below baseline. Per feedback_isv_for_adaptive_bounds: signal flows
through ISV[128] (new MOE_LAMBDA_EFF_INDEX); consumer kernel reads
the slot at runtime; no DtoH on hot path.
Wiring (feedback_no_partial_refactor):
• new kernel moe_lambda_eff_update (single-block cold-path cadence,
matches kelly_cap_update / cql_alpha_seed_update precedent)
• new ISV slot 128, ISV_TOTAL_DIM 127 → 129
• layout_fingerprint_seed updated (schema_hash bumps; PVC fxcache will
auto-regen on next deploy)
• moe_load_balance_loss kernel takes (isv_signals, isv_lambda_eff_idx)
instead of float lambda — matches mag_concat_qdir's ISV-read pattern
• config: pub moe_lambda field replaced with moe_lambda_floor (0.01),
moe_lambda_max_extra (0.09), moe_entropy_target_frac (0.7)
• state-reset registry entry — ISV[128] resets to floor at fold boundary
• HEALTH_DIAG aux_moe line gains λ_eff field for observability
• hyperopt adapter (DQNHyperparameters) migrated to new knobs
• constructor bootstraps ISV[128]=floor + ISV[118..127)=uniform 1/K +
ISV[126]=ln(K) so cold-start matches legacy byte-for-byte
Smoke validates: magnitude_distribution still passes/fails identically
to HEAD (eq=0.586 same as pre-change eq=0.592 — unrelated Kelly cap
behaviour per project_magnitude_eval_collapse_kelly_capped.md).
HEALTH_DIAG fold 3 confirms controller live: ent decays 1.381 → 0.746
across epochs 7-19, λ_eff rises 0.0146 → 0.0539 monotonically.
Unit test moe_lambda_eff_update_writes_correct_values exercises 5
regimes (above-target / at-target / half-collapse / full-collapse /
deeply-above-target) on the GPU kernel — all pass within 1e-5.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
compute-sanitizer caught 16 OOB errors in magma_sgemmEx_kernel
<f,f,f,1,0,6,4,6,3,4>+0xe90 surfacing as mamba2_scan_projected_bwd
LAUNCH_FAILED. Pre-existing latent, exposed when q_var_buf_trainer
allocator reshuffle (fa92cb8db) moved subsequent buffer pointers.
Root cause: DuelingWeightSet::from_flat_buffer and
BranchingWeightSet::from_flat_buffer in gpu_weights.rs used hard-coded
layout indices [0..11] / [12..23] from before the GRN trunk expansion
(Plan 4 Task 2c.3a). The expansion inserted 9 GRN tensors at indices
1..12, shifting every value/branch tensor by +9. The 12-index sequential
walk in from_flat_buffer silently slid every dueling-weight pointer
forward into adjacent tensors:
- online_dueling.w_v1 → w_residual_h_s1 (sized SH1*s1_input_dim, not
VALUE_H*SH2)
- online_dueling.b_v1 → gamma_h_s1 (sized SH1=64 floats=256 bytes,
not VALUE_H=128 floats=512 bytes)
- online_dueling.w_v2 → beta_h_s1; b_v2 → w_a_h_s2; w_a1 → b_a_h_s2; …
The pathological alias surfaced via the ensemble multi-head clone path
(forward_value_head_for_ensemble): the cuBLAS RELU_BIAS epilogue tried
to read bias[0..128] from b_v1's 64-float buffer, producing exactly the
observed 16 thread (0..7, 2..3) reads at 1..61 bytes past a 256-byte
allocation.
Production GEMMs never tripped this because they read from
params_buf via f32_weight_ptrs_from_base which has always used the
correct 163-tensor GRN layout. flatten/unflatten paths used
matched-stale self-copies (no-op when online_d.w_s1 == params_buf_ptr).
Only the ensemble clone (`clone_dueling_weights`) made independent
copies of the misaliased pointers and then handed them to the GEMM as
if they were value-head tensors.
Per feedback_no_partial_refactor: every consumer of the weight-set/
flat-buffer contract migrated in lockstep:
- New DUELING_FLAT_INDICES = [0,1,2,3, 13,14,15,16, 17,18,19,20]
and BRANCHING_FLAT_INDICES = [21..32] in gpu_weights.rs encode the
authoritative mapping from DWS/BWS slots to GRN-expanded layout
indices.
- DuelingWeightSet::from_flat_buffer + BranchingWeightSet::from_flat_buffer
rewritten to use these mappings with a full prefix-sum byte-offsets
table (matches f32_weight_ptrs_from_base byte layout).
- flatten_online_weights, unflatten_online_weights,
flatten_target_weights, unflatten_target_weights (gpu_dqn_trainer.rs)
rewritten to keyed [(ptr, layout_idx); 24] pairs and write at
byte_offsets[layout_idx] instead of sequential prefix-sum over
sizes[0..23]. The no-op zero-copy check (online_d.w_s1 == src_base)
is preserved because DUELING_FLAT_INDICES[0] == 0.
Sanitizer (RTX 3050 Ti, magnitude_distribution smoke):
magma_sgemmEx_kernel OOB count: 16 → 0
Non-sanitizer smoke completes all 20 epochs without LAUNCH_FAILED
(was crashing on epoch 1 prior to fix); MAG_DIST/EVAL_DIST results
reflect real model behavior (Q=0.349, H=0.298, F=0.353 train-mode).
The unrelated F_Full eval-cap assertion is the ongoing Kelly cap
issue (project_magnitude_eval_collapse_kelly_capped.md).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Smoke after var_scale floor (73fd49f4b) still showed eval Full = 0.000 with
intent = 0.911. Root cause: trade_physics.cuh::unified_env_step_core applied
Kelly cap using ONLY direction conviction (q_dir_gap / q_dir_abs_ref). At
smoke horizon with direction Q-values close (q_s≈q_h≈q_l≈q_f within ~0.07),
direction conviction was ~0.04 → safety = max(0.75, 0.04) = 0.75,
effective_kelly = max(0, 0.75) = 0.75, cap = 0.75 × max_pos × 0.75 =
0.5625 × max_pos → Half bucket pin (abs_pos < 0.75).
The magnitude branch was simultaneously highly confident (q_f / q_h ≈ 1.9,
mag conviction ≈ 0.5) — the policy strongly preferred Full. That signal
never reached the cap.
Thread per-sample magnitude conviction into the cap formula:
- new buffer magnitude_conviction_buf[N*L] (training) parallel to
conviction_buf, and chunked_magnitude_conviction_buf[chunk_len * n_windows]
(eval) parallel to chunked_conviction_buf
- experience_action_select writes (max(q_mag) − min(q_mag)) /
fmaxf(ISV[16], 1e-6); ISV[16] = Q_ABS_REF (magnitude branch's |Q| EMA,
recycled — no new ISV slot)
- scripted_policy_kernel writes 0.0 (no real magnitude preference; helper
composer falls back to direction signal)
- unified_env_step_core takes a parallel magnitude_conviction parameter
- combined: policy_conviction = max(dir_conv, mag_conv); cap formula
becomes safety = max(health_safety, policy_conviction); warmup_floor
uses policy_conviction
- all three env_step kernel call sites migrated together
(backtest_env_step, backtest_env_step_batch, experience_env_step) per
feedback_no_partial_refactor — no mixed state
Per pearl_blend_formulas_must_have_permanent_floor: max() composition,
not multiplicative blend. Per feedback_isv_for_adaptive_bounds: signal
flows from existing ISV[16] = Q_ABS_REF, no tuned constants.
Smoke verification (20-epoch RTX 3050 Ti, magnitude_distribution test):
- pre-fix: [EVAL_DIST] Q=0.586 H=0.414 F=0.000 (intent F=0.911)
- post-fix: [EVAL_DIST] Q=0.618 H=0.153 F=0.229 (intent F=0.911)
Intent unchanged (network already learned Full preference); the Kelly
cap no longer silences it. Both smoke assertions now pass:
ef >= 0.05 (gate) and eh + ef >= 0.30 (Task 2.2 H10).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
compute-sanitizer caught 4 OOB writes in compute_expected_q at threads
60..63 of block 0 (production: batch=64, total_actions=13). Stride
52 bytes/thread (= 13 floats), increasing 45 → 97 → 149 → 201 bytes
past a 4-byte-aligned neighbour allocation. The kernel writes
`q_variance[i*total_actions + a]` with stride total_actions=13 (4+3+3+3
factored layout) into a buffer that was still sized at b*12 from the
legacy 3+3+3+3 exposure layout. Same class of bug as `d625ca28e` (q_readback
12→total_actions): the direction branch grew from 3 → 4 actions but
several legacy 12-stride consumers were not migrated.
Per feedback_no_partial_refactor: q_var_buf_trainer's contract is a
shared one — every consumer migrates in lockstep. Five touch points:
• alloc q_var_buf_trainer: b*12 → b*total_actions
• compute_expected_q (writer): unchanged, already stride total_actions
• q_denoise_step kernel: add `var_stride` param, read var_q at stride
total_actions; FC layer dim D=12 unchanged (denoiser MLP weights stay)
• q_denoise_backward kernel: same — add `var_stride`, two read sites
• denoise_build_input kernel: add `var_stride` between D and schedule
• launch_loss_reduce: b_qvar = b*12 → b*total_actions for the third
c51_loss_reduce launch (ISV slot 3/4 mean reduction)
The denoiser's internal D=12 is preserved — its FC layers operate on the
12-slot exposure layout and only consume the first 12 variance entries
per sample. The 13th variance slot (last urgency action) is computed by
compute_expected_q but currently unused by the denoiser; epistemic_gate
already reads with stride total_actions and benefits from full per-action
variance now that the buffer is correctly sized.
This was a latent bug masked by allocator fortune — earlier mag_concat
corruption (`8bc6f1ccd`) was overwriting buffers first, so this OOB
landed in already-corrupted memory and produced no compute-sanitizer
report. With mag_concat fixed, the q_var stride mismatch became the
dominant remaining corruption source, surfacing as
`bias_grad_reduce_f32_p1: DriverError(CUDA_ERROR_LAUNCH_FAILED)` on the
next launch in production runs.
Sanitizer (RTX 3050 Ti, magnitude_distribution smoke):
compute_expected_q OOB count: 4 → 0
Note: the run still trips 16 OOB reads in `magma_sgemmEx_kernel` (cuBLAS
GEMM) under a different code path; those are pre-existing — they appear
only because q_var growing from b*12 to b*total_actions reshuffled
allocation addresses, exposing a separate latent cuBLAS leading-dim
mismatch that was previously aliasing into q_var's old footprint. That
finding is filed in the wire-up audit and is not in scope for this commit.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Smoke (with mag_concat off-by-one fixed and fxcache regenerated):
- training MAG_DIST: Q=0.287 H=0.301 F=0.412 (healthy diversity)
- INTENT at eval: Q=0.045 H=0.045 F=0.911 (network learned Full)
- realised at eval: Q=0.592 H=0.408 F=0.000 (Full silenced)
Even though q_f >> q_h >> q_q (network strongly prefers Full), eval
realised Full = 0.000. Position-sizing pipeline composes four
multiplicative shrinks at experience_kernels.cu:1615-1672:
effective_max_pos = max_position
× cvar_scale (line 1621)
× q_gap_conviction (line 1628, clamped [0.25,1])
× kelly_f (line 1649, only if >20 trades)
× var_scale (line 1671, = 1/(1+sqrt(var_q)))
Each term well-bounded individually but composing them silences the
policy at validation when var_q persists high (var_q ≈ 30 → var_scale
≈ 0.15; even with conviction = 1.0, kelly_f = 1.0, the compound 0.15
falls in the Quarter bucket [0, 0.375)). Network INTENT reaches the
target_position kernel correctly; var_scale strips it back out.
Same family as val-Flat-collapse (warm-branch Kelly = 0 from balanced
priors) — fix is the same pearl
(`pearl_blend_formulas_must_have_permanent_floor.md`):
var_scale = max(var_scale, q_gap_conviction)
The q_gap-derived conviction is already an adaptive ISV-coupled signal
(line 1628), already clamped to [0.25, 1.0]. Using it as a permanent
floor on var_scale lets the policy's magnitude intent reach the
realised position when conviction is high, regardless of variance.
No tuned constants — feedback_adaptive_not_tuned + feedback_isv_for_adaptive_bounds
both honoured by reusing the existing adaptive bound rather than
introducing a new threshold.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Smoke and localdev profiles previously used data_source = "ohlcv", divergent
from production (mbp10). Mismatch silently produced cache-key collisions in
the SHA256 hash path: smoke fxcache could not be loaded by production-shape
training without an explicit override. Local-dev fxcache regen also required
remembering to pass --data-source ohlcv to match.
Globalize mbp10 as the default everywhere it isn't deliberately overridden:
- config/training/dqn-smoketest.toml: data_source = "mbp10"
- config/training/dqn-localdev.toml: data_source = "mbp10"
- training_profile.rs: doc Default → "mbp10"
- trainers/dqn/config.rs: Default impl → "mbp10"
- hyperopt/adapters/dqn.rs: default → "mbp10"
- examples/precompute_features.rs: doc updated
- fxcache.rs / feature_cache.rs: discover_and_load + cache-key tests
use "mbp10" arguments
- docs/dqn-wire-up-audit.md: new entry per Invariant 7
Documentation strings retained "ohlcv" only where they document the two
available choices (config.rs:946, training_profile.rs:83).
Local fxcache regenerated to v6 mbp10:
test_data/feature-cache/13c0b086a975cc7e2384377a2cd0e97738c9410292fcfecb5807c29bf885cb48.fxcache
(175874 bars, 55 MB, OFI_DIM=32). Stale v5 ohlcv fxcache untracked
from git index (already gitignored post-79578bbaf).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
compute-sanitizer pinpointed mag_concat_qdir at experience_kernels.cu:3590
writing 260 floats/state (SH2 + b0_size where b0=4) into a buffer
allocated for 259 floats/state (SH2 + 3, legacy 3-direction layout).
Off-by-one corrupted the next row's first column on every write and
overran past the buffer end on the final row, surfacing as
CUDA_ERROR_ILLEGAL_ADDRESS in downstream kernels (denoise_bias_grad_p1,
cublasLt h_v matmul) on L40S production batch sizes.
The `+3` constant was overloaded:
- direction-conditioned (mag_concat, w_b1fc, w_gate_1) — incorrectly
hardcoded SH2+3 instead of SH2+branch_0_size when the kernel migrated
to 4-direction (S/H/L/F).
- OFI-conditioned (ord_concat, urg_concat, w_b2fc, w_b3fc, w_gate_2,
w_gate_3) — correctly SH2+3 for 3 OFI features per branch
(concat_ofi_features).
Migrated all direction-conditioned consumers in lockstep
(feedback_no_partial_refactor):
- gpu_dqn_trainer.rs: w_b1fc, w_gate_1 use shared_h2+branch_0_size
- gpu_dqn_trainer.rs: split mag_concat_dim (SH2+b0) from
ofi_concat_dim (SH2+3) for buffer alloc
- gpu_dqn_trainer.rs: accumulate_d_h_s2_from_concat takes src_stride
param so mag callers pass SH2+b0, ord/urg callers pass SH2+3
- batched_forward.rs: split mag_concat_dim (SH2+b0) from ofi_concat_dim
(SH2+3); strided_scatter dst_stride and fc_k now diverge between mag
(d==1) and ord/urg (d==2,3); add separate (SH2+3) GEMM cache shape
- batched_backward.rs: d==1 (magnitude) dX/dW dims use SH2+b0;
d==2/3 (order/urgency) keep SH2+3; add separate (SH2+3) GEMM cache
shape for OFI branches
- gradient_budget.rs: smoke test now allocates b1 with SH2+b0
and b2/b3 with SH2+3 (was buggy SH2+3 for all three)
- value_decoder.rs: doc updated
- docs/dqn-named-dims.md: new "Branch FC input strides" section
documenting the direction- vs OFI-conditioning invariant
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
L40S production training crashed in launch_ofi_embed_backward with
CUDA_ERROR_ILLEGAL_ADDRESS at the first denoise_bias_grad_p1 launch in
the adam_grad child graph (workflow train-multi-seed-tdjtr fold 0).
Root cause: denoise_bias_grad_p1/p2 CUfunction handles were shared
across two independently-captured child graphs (post_aux and adam_grad).
On Ada/Hopper, when partials buffers land in different VRAM arenas
(L40S B=16384 post-MoE memory layout), the driver's per-node handle
validation trips. Pre-existing latent bug — masked on smaller GPUs and
pre-MoE allocator states where buffers happened to be address-adjacent.
Fix: load a dedicated CUmodule for the OFI bias-grad reduction kernels,
separate from the denoise variant. Same kernel bytecode (EXPECTED_Q_CUBIN
contains both); each child graph gets its own CUfunction reference.
Mirrors the graph_safe_copy_kernel precedent at line ~11523:
"Hopper CUfunction isolation: each child graph needs its own CUmodule".
Constructor adds ~2ms of cubin-load time; zero runtime cost. Smoke
validates no regression at B=64.
Production validation deferred to follow-up L40S re-deploy.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Reads the most recent HEALTH_DIAG aux_moe line from the smoke log
and asserts: (1) no expert utilization < 2% (anti-collapse working),
(2) no expert > 80% (no snap-collapse to single expert), (3) gate
entropy < 0.9·ln(8) (peakier than uniform on average).
Failures are findings, not test infrastructure bugs — concrete numbers
reported per feedback_kill_runs_on_anomaly_quickly.md.
Spec: docs/superpowers/specs/2026-04-27-moe-regime-redesign-design.md §8.3.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Atomic cleanup per feedback_no_partial_refactor.md:
- Delete crates/ml-dqn/src/regime_conditional.rs and all
RegimeConditional* exports from lib.rs. regime_classifier.rs
(RegimeType, RegimeClassConfig) is kept — used by validation layer
and ml-regime-detection crate.
- Rewrite DQNAgentType to wrap DQN directly (no RegimeConditionalDQN
field, no 3-head delegation gymnastics).
- DQNAgentType::get_count_bonuses_branched no longer returns hardcoded
None — it delegates to DQN::get_count_bonuses_branched() and UCB
count bonuses reach the GPU action selector for the first time
(ghost feature from Phase 0 deferral). action.rs caller simplified
to direct destructure of fixed arrays, no Option matching.
- Drop 4 regime-threshold fields from DQNConfig (regime_adx_idx,
regime_cusum_idx, regime_adx_threshold, regime_cusum_threshold) +
matching Default and aggressive() builder entries.
- serialize_model rewritten to serialize single DQN branching network
directly (no trending__/ranging__/volatile__ prefix namespace).
- curriculum.rs import fixed: crate::dqn::regime_conditional::RegimeType
→ crate::dqn::RegimeType (re-exported from regime_classifier).
- Constructor drops RegimeConditionalDQN::new_on_device, calls
DQN::new_on_device directly.
- Stale doc comments in dqn.rs, moe.rs, regime_classifier.rs updated.
- Wire-up audit updated.
Phase 3 MoE (commit a52d99613) provides the regime-conditioned behavior
the legacy 3-head architecture pretended to do — gate sees ADX (40)
and CUSUM (41) as part of the full 128-dim state vector and learns its
own decomposition, strictly subsuming the threshold classifier.
Smoke: 3/3 folds passed (405s), gate util=0.119,0.119,0.169,... preserved,
all fold checkpoints written. Workspace: 0 errors across all crates,
tests, and examples.
Spec: docs/superpowers/specs/2026-04-27-moe-regime-redesign-design.md §7.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Phase 3 of the MoE regime redesign per
docs/superpowers/specs/2026-04-27-moe-regime-redesign-design.md.
Atomic forward+backward wire-up — replaces the existing h_s2 producer
in fused_training.rs with the gated expert mixture:
state[B, 128] -> shared GRN trunk -> h_s1[B, 256]
|
+-> 8 expert MLPs (256->64->256) -> expert_outputs[8, B, 256]
+-> gate (128->64->8 softmax) ----> gate[B, 8]
|
moe_mixture_forward -> h_s2[B, 256] -> branching heads + C51 + IQN
Backward: moe_mixture_backward (de_k = g · dh_s2) + moe_dgate_reduce
(dg = Σ_c e_k · dh_s2) + load-balance aux gradient + cuBLAS SGEMM
backward through gate + each expert's 2 linear layers. Adam optimizer
step now updates gate + 8 experts via params_buf.
Loss: λ · K · Σ_k (mean_b g[b,k])² added to total loss with λ from
hyperparams.moe_lambda (default 0.01).
Per-step ISV producer launch (moe_expert_util_ema_update) writes 8
utilization EMA + 1 gate-entropy EMA into ISV[118..127). Per-epoch
HEALTH_DIAG aux_moe line emits utilization vector + entropy live so
operators can see whether experts are differentiating or collapsing.
Smoke test: DONE. 3/3 folds, all checkpoints saved, 728s (12.1 min,
within 25-min budget). Gate differentiated by epoch 1: expert 2 rose
from 0.119 → 0.286 → 0.323 over fold 1-2 while others remained at
0.097-0.113. Gate entropy 1.611 at fold 2 epoch 4 < ln(8)=2.079.
val_loss finite across all 3 folds; average fold metric 22.4.
Per feedback_no_partial_refactor.md: all consumers of the h_s2 contract
(branching heads, IQN aux, attention focus, backward chain) migrate in
this single commit.
Per feedback_no_htod_htoh_only_mapped_pinned.md: no new HtoD/HtoH
introduced; gate softmax + expert outputs + mixture all GPU-resident,
ISV producer GPU-driven. Load-balance scalar uses cuMemAllocHost +
cuMemHostGetDevicePointer_v2 (mapped pinned), matching existing pattern.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Single-thread cold-path-cadence kernel writes 8 per-expert utilization
EMAs + 1 gate-entropy EMA into ISV slots [118..127), α=0.05. Same shape
as h_s2_rms_ema_update / aux_heads_loss_ema_update.
GPU-resident, CPU read-only per pearl_cold_path_no_exception_to_gpu_drives.md.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Two-step design (per-k contribution then K-element reduce) avoids
atomicAdd per feedback_no_atomicadd.md. Tests verify CPU reference
match and uniform-gate minimum (loss = λ).
Backward gradient lands in Phase 3 (wire-up).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Two-stage backward: moe_mixture_backward computes de_k = g · dh_s2 in
one kernel; moe_dgate_reduce computes dg via shmem reduction over c.
FD test verifies analytic backward matches numerical for B=2, K=4, C=32
within 1e-3 tolerance (perturbation 1e-3).
All test data flows via mapped pinned buffers (no HtoD/HtoH).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Moves MappedF32Buffer + MappedI32Buffer from distributional_q_tests.rs
local definitions to crates/ml/src/cuda_pipeline/mapped_pinned.rs so all
kernel test wrappers (Phase 0.F, upcoming MoE Phase 2 tests) share one
implementation.
Adds write_from_slice helper for direct host_ptr write (no memcpy)
per feedback_no_htod_htoh_only_mapped_pinned.md.
Test 0.F bit-identical post-move: sigma_C51, argmax, Thompson counts
unchanged.
Per spec docs/superpowers/specs/2026-04-27-moe-regime-redesign-design.md §6.4
and feedback_no_htod_htoh_only_mapped_pinned.md.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
T1.6 implementer correctly identified that the gate input dim is
ml_core::state_layout::STATE_DIM=128, not the literal 42 the spec/plan
incorrectly stated. The 42-dim figure was the bar-feature subset; the
actual state vector is 128-dim (42 features + portfolio + MTF + OFI
padded to 128 for cuBLAS alignment).
Updated spec §3 architecture diagram, §4.1 gate subnetwork description
+ parameter count (3,272 → 8,776), and plan header architecture line.
Implementation in commit 28c707f6a is correct; this commit just makes
the spec match the implementation.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Gate (4 tensors, [127..131)) + 8 experts × 4 tensors (32 tensors, [131..163))
appended to GpuDqnTrainer params_buf layout. NUM_WEIGHT_TENSORS 127 → 163.
Layout fingerprint hash recomputes — old checkpoints fail to load with
fingerprint-mismatch error per the no-fallback contract in the spec.
Gate tensors: gate_w1[STATE_DIM=128,64], gate_b1[64], gate_w2[64,8],
gate_b2[8]. Zero-init so g(s) = uniform 1/K at cold start.
Expert tensors (per expert k∈[0,8)): w1[SH2,BTN], b1[BTN], w2[BTN,SH2],
b2[SH2] where SH2=cfg.shared_h2=256, BTN=MOE_EXPERT_BOTTLENECK=64.
Xavier init on w1/w2; zero on biases. Total ~268k new params.
Adam state (m_buf/v_buf), gradient scratch, target_params_buf all extend
in lockstep — all sized from compute_total_params() which sums over the
full 163-tensor layout. No static sizes to update.
Tensors allocated but not yet wired into forward/backward — Phase 3
wires gate forward, expert forward, mixture replacement of h_s2, and
the corresponding backward chain.
No test assert updates required — no test hardcodes NUM_WEIGHT_TENSORS
or the layout fingerprint value.
Spec: docs/superpowers/specs/2026-04-27-moe-regime-redesign-design.md §6.1.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
8 slots for per-expert utilization EMA + 1 slot for gate-entropy EMA.
ISV_TOTAL_DIM 118 -> 127. Not yet consumed; producers + consumers land
in subsequent commits per the MoE plan.
Spec: docs/superpowers/specs/2026-04-27-moe-regime-redesign-design.md §4.6.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Anti-collapse load-balancing aux loss weight for the upcoming Mixture-of-
Experts redesign. Configurable via hyperopt; not a kernel constant.
Spec: docs/superpowers/specs/2026-04-27-moe-regime-redesign-design.md §4.3.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Atomic removal of 5 boolean feature flags from DQNConfig per
`feedback_no_feature_flags.md`, all of which gated dead or redundant code.
(use_iqn already removed in da632446c.)
- `use_dueling`, `use_distributional`, `use_noisy_nets`, `use_branching`:
pure-cosmetic always-on flags. Only metadata strings remained.
4 `if true /* always on */` blocks in dqn.rs collapsed to live arms;
dead else arms (legacy non-branching code paths, 5-flat ExposureLevel)
deleted.
- `use_count_bonus`: redundant with `count_bonus_coefficient` (the live
numeric kill-switch). Recording made unconditional; coefficient is the
single dial.
- `use_cvar_action_selection`: dead post-use_iqn cleanup (only consumers
were inside the deleted IQN arms). cvar_alpha retained for cuda_pipeline
CVaR loss usage.
count_bonus.rs API tightened: `bonuses_branched_f32(&self,
exp_out: &mut [f32], ord_out: &mut [f32], urg_out: &mut [f32])` write-into
API alongside the existing Vec<f64> form (kept for tests). Production
DQN::get_count_bonuses_branched() returns fixed `([f32; 4], [f32; 3],
[f32; 3])` arrays — allocation-free, f32 throughout, fed directly to GPU
action selector via stack-allocated buffers.
Test 0.F outputs bit-identical vs pre-cleanup (sigma_C51, argmax,
Thompson counts) — confirms legacy use_* arms were unreachable as
expected.
`docs/dqn-wire-up-audit.md` updated per Invariant 7.
Precondition for the MoE regime redesign per
`docs/superpowers/specs/2026-04-27-moe-regime-redesign-design.md`.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per feedback_no_htod_htoh_only_mapped_pinned.md (newly recorded): every
CPU<->GPU path in this redesign uses mapped pinned memory exclusively.
No cudaMemcpy HtoD, no Vec-to-Vec defensive copies, including in test
code. CPU is strictly read-only on the production surface.
Plan changes:
- New Task 2.0 promotes MappedF32Buffer / MappedI32Buffer from
distributional_q_tests.rs local definitions to a shared
crates/ml/src/cuda_pipeline/mapped_pinned.rs module so all kernel
test wrappers (Test 0.F, upcoming MoE tests) share one
implementation. Adds write_from_slice helper for direct host_ptr
write (no memcpy).
- Task 2.1 test wrapper rewritten to allocate mapped pinned buffers
+ write to host_ptr + read GPU-written output via host_ptr. No more
memcpy_stod / memcpy_dtov in test code.
Spec: new section 6.4 codifies the mapped-pinned-only constraint and
references the shared module + reference implementation.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
End-to-end investigation (2026-04-27) confirmed RegimeConditionalDQN is
vestigial decoration — 3 heads constructed at training start but only
trending_head ever receives gradient updates. GpuDqnTrainer (the actual
production GPU trainer) has zero references to RegimeType/regime
routing; experience replay inserts go to trending_head.memory only;
ranging_head and volatile_head stay at random init for the entire
training run. Several support APIs (get_count_bonuses_branched, config,
get_state_dim) hardcode-delegate to trending_head, ignoring the regime
split entirely. Per `feedback_no_hiding.md` (wire up or delete) and the
user's preference to fix not delete: the design wires regime
conditioning properly via Mixture-of-Experts replacing the vestigial
3-head architecture.
Pearl introduced and saved as `pearl_learned_gate_subsumes_handcoded.md`:
when the network already sees the heuristic's inputs, a learned gate
strictly subsumes any hand-coded discretization. This is the
load-bearing rationale — ADX/CUSUM are already at state indices 40/41,
so threshold-based regime classification is a strict information
bottleneck the gate can recover and improve on.
Design summary:
- Architecture: shared GRN trunk -> K=8 small expert MLPs (256->64->256
bottleneck per expert, ~33k params each) -> learned gating network
(state[42]->64->8 softmax) -> mixed h_s2 -> existing branching heads
+ C51 + IQN dual head. Soft full mixture (no top-k hardcoding); gate
emerges peaky or flat from data. Anti-collapse load-balancing aux
loss with default lambda=0.01 (configurable hyperparameter, not a
kernel constant) prevents init-noise-dominated single-expert lock-in
without forcing uniform utilization. User-confirmed signal:
"collapses don't recover well in this codebase".
- 9 new ISV slots (118-126: per-expert utilization EMA + gate entropy
EMA), GPU-driven producer per
`pearl_cold_path_no_exception_to_gpu_drives.md`.
- 3 new small CUDA kernels (moe_mixture_forward/backward,
moe_load_balance_loss) + 1 ISV producer; everything else is cuBLAS-
reusable. CUDA Graph capture compatible.
- Atomic deletion (no fallback): regime_conditional.rs (~700 LOC),
RegimeType enum, classify_from_features, RegimeMetrics,
RegimeClassConfig, 4 DQNConfig regime threshold fields, per-regime
3-file checkpoint format. DQNAgentType becomes thin wrapper over
single DQN. Old checkpoints fail loudly with layout-fingerprint
mismatch.
- 5-layer testing strategy (unit kernels, smoke, gate-differentiation
validation, L40S production validation with explicit kill criteria,
architecture-hash backward-incompat).
Out of scope (explicit): top-k routing, per-expert action heads,
hierarchical MoE, regime-conditional CountBonus/NoisySigma broadcast,
expert warm-start from existing trending checkpoint, CVaR action
selection on mixed C51 distribution.
Precondition: the in-progress use_* flag cleanup + count_bonus
[f32; N] refactor lands as its own commit before MoE implementation
begins, per `feedback_no_partial_refactor.md`.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
`use_iqn` is exactly the `use_/enable_` boolean banned by
`feedback_no_feature_flags.md`. It gated dead code: production training
runs IQN unconditionally through `cuda_pipeline/gpu_iqn_head.rs` +
`iqn_dual_head_kernel`, properly wired into the branching architecture
with the `FIXED_TAUS` 5-quantile schedule. Nothing in the cuda_pipeline
production path ever read `use_iqn` or `iqn_network`.
The legacy `iqn_network: Option<QuantileNetwork>` field on `DQN` was a
parallel CPU-side network from a pre-branching era, structurally
unreachable in production: every consumer was gated behind the
`if true /* use_branching: always on */` arm at `q_values_for_batch`,
so the IQN else-if at L1822 was dead code. Training optimised
`iqn_network` parameters in isolation; inference never read them. That's
the train/inference mismatch the L283 comment ("IQN trains base
q_network but inference uses dist_dueling network (zero gradients)")
was working around by **disabling** the feature instead of fixing the
inference path. Per `feedback_no_quickfixes.md` + `feedback_no_hiding.md`
the fix is to remove the dead path entirely.
Strip:
- `DQNConfig::use_iqn` field + parses + checkpoint hash + metadata.
`dqn.use_iqn` / `dqn.iqn_embedding_dim` / `dqn.iqn_num_quantiles` /
`dqn.iqn_kappa` from older checkpoints are silently dropped on load
(same pattern used for `use_dueling`). `iqn_lambda` stays — the
cuda_pipeline dual head consumes it as the IQN aux loss weight.
- 3 vestigial config fields (`iqn_num_quantiles`, `iqn_kappa`,
`iqn_embedding_dim`) — never read in production; kernel-side macros
(`IQN_NUM_QUANTILES = 5`, embed_dim 64, kappa 1.0) are the actual
config.
- 4 default builders (`Default`, `aggressive`, `conservative`,
`emergency_safe_defaults`) drop the 4 IQN-related fields each.
- `DQN::iqn_network` field + initialisation block in `new_with_stream`.
- 6 conditional gates in `select_action`, `select_action_with_confidence`,
`select_action_inference`, `q_values_for_batch` — all collapse to the
live (branching or standard-Q) arm.
- `DQN::get_state_embedding` (only consumed by deleted IQN paths).
- The entire `crates/ml-dqn/src/quantile_regression.rs` module (392 LOC)
+ its 2 lib.rs exports. Nothing outside `ml-dqn` ever imported it
(the `quantile_huber_loss` reference in `gpu_iqn_head.rs` is a CUDA
kernel name string, unrelated to this Rust module).
Downstream call sites:
- `crates/ml/src/trainers/dqn/{config,fused_training,trainer/constructor}.rs`:
drop `iqn_num_quantiles` / `iqn_embedding_dim` / `iqn_kappa` references
off `DQNConfig`; substitute kernel-fixed literals (64, 1.0) where
`GpuDqnTrainConfig` / `GpuIqnConfig` still expect them.
- `crates/ml/examples/evaluate_baseline.rs`: drop two `iqn_num_quantiles`
hyperparam reads (their `..DQNConfig::default()` fallbacks now stand
alone).
- `crates/ml/tests/dqn_action_collapse_fix_test.rs`: drop the
`assert!(!config.use_iqn, "...gradient dead zone")` and the explicit
`config.use_iqn = false` setter; the dead-zone pathology is now
structurally impossible.
- `crates/ml/tests/dqn_inference_test.rs`: drop `config.use_iqn = false`.
- `services/trading_service/src/services/dqn_model.rs`: drop the
`iqn={}` debug-log field.
- `crates/ml/src/trainers/dqn/distributional_q_tests.rs`: ship Test 0.F
(Plan A Task 8 #186) — converged-checkpoint extraction harness +
`MappedF32Buffer` (mapped pinned f32 mirror) + `compute_sigma_c51_test`
kernel handle. The structural assertions panic on the local 5-epoch
smoke checkpoint as designed (under-converged: sigma_C51 spread ~1.4%
across directions, P(active)=0.4330 >= 0.20). Docstring rewritten to
drop `use_iqn=false` framing and the legacy "Tier-B-prime" caveat;
Tier-A version is GPU-integration-only per
`feedback_no_cpu_forwards.md` (CPU is read-only).
`docs/dqn-wire-up-audit.md` updated per Invariant 7.
Build: `cargo check --workspace --tests` clean (0 errors).
Test: `cargo test -p ml --lib distributional_q_tests::test_0f
-- --ignored --nocapture` produces bit-identical sigma_C51 /
argmax / Thompson values vs pre-removal — confirms branching
forward path is the same code post-cleanup as pre-cleanup
(legacy `use_iqn` arms were unreachable, as expected).
Net: 12 files, +458 / -785 lines (327 LOC deleted).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Two-part fix for a class of bugs causing un-normalised features to
silently flow into training.
(1) data_source mismatch between precompute writer and trainer reader.
precompute_features.rs:214,633 hardcoded "ohlcv"
train_baseline_rl.rs:582 hardcoded "ohlcv"
config/training/dqn-production.toml: data_source = "mbp10"
Production runs the trainer with the production profile (data_source
= mbp10), but the actual cache lookup hardcoded "ohlcv". Smoke worked
by accident (smoke profile is also "ohlcv"). Any future profile with
a different data_source silently mismatches → cache MISS → DBN
fallback path.
Both call sites now hardcode "mbp10" (the canonical production data
source per CLAUDE.md). precompute_features adds a `--data-source`
CLI override for the rare case a smoke flow needs to regenerate
the local "ohlcv" fxcache; default is "mbp10".
(2) DBN-fallback path didn't normalise features.
precompute_features.rs:629 applies NormStats::normalize_batch on the
canonical fxcache write path. The fallback in train_baseline_rl.rs
(cache-miss → load DBN files → extract features → upload to GPU) did
NOT normalise. Any cache miss (data_source drift, schema-hash mismatch,
missing file) silently uploaded RAW features. Raw close prices
(~$5180 ES futures) flowed into next_states[:, 0]; the aux next-bar
head's label_scale EMA latched onto raw-price magnitude (~5443 vs
expected ~1.0 z-score); the shared trunk learned to predict next-bar
prices; the policy effectively traded with future-price knowledge →
train-h5gxb epoch-0 Sharpe = 141 with 0.32% max-drawdown over 214k
bars (impossibly good = oracle leak).
DBN fallback now applies the same z-score normalisation unconditionally
as defence-in-depth, so a future cache-miss cannot reintroduce raw
values into training.
Audit entry updated.
build.rs::emit_feature_schema_hash only hashed
src/features/extraction.rs, src/fxcache.rs, and
../ml-core/src/state_layout.rs. The z-score normalization step lives
in examples/precompute_features.rs:625-631 (added 2026-04-03 in
9f7c14978f) and was NOT covered by the hash. A fxcache written by an
older precompute build (raw features, no normalization) silently
passed today's validate() because every other field matched.
Empirical impact (L40S Argo train-f8h6q, 2026-04-27):
- PVC fxcache: stale, written pre-normalization → feature column 0
contains RAW CLOSE PRICES (~$5180 ES futures) instead of z-
normalized log-returns
- aux head reads next_states[:, 0] as its next-bar regression
label (gpu_dqn_trainer.rs:7758-7789)
- EMA label_scale climbed to 5420 (vs smoke 0.05) → shared trunk
learned to predict next-bar prices → policy effectively traded
with future-bar information
- epoch-0 Sharpe = 141.99 with 0.32% max-drawdown over 214k bars
— physically impossible; clear future-leak signature
Fix adds examples/precompute_features.rs to schema_sources. New hash
invalidates the stale PVC cache. Argo's ensure-fxcache step has a
regenerate-on-failure branch (infra/k8s/argo/train-template.yaml:
372-383) that auto-regens with current normalized precompute.
Generalises beyond this incident: any future change to feature
normalization, target ordering, or precompute post-processing now
bumps the hash and forces fxcache regen.
Audit entry updated.
Retracts the prior SUPERSEDED footer (commit 42ffd6aad). The technical
proposal still stands — Thompson sampling on C51+IQN distributions is
the canonical action selector for distributional RL (Bellemare 2017,
Dabney 2018). Phase 0 tests and the Aggregation Contract are sound
math/engineering regardless of the measurement-bug findings.
What changed is the URGENCY framing, not the validity. The val-Flat-
collapse / Short-collapse observations cited as motivating evidence
were partly distorted by three measurement bugs (a86fba2b1 + b8788511c)
in the diagnostic infrastructure. The "ship Thompson NOW because
val_dir_dist collapses to 80%+ Hold/Flat" narrative dissolves; the
"Thompson is the principled action selector for our distributional
model" narrative stands.
Sequencing: PAUSED pending evidence from a fresh L40S 30-epoch
baseline (train-f8h6q, 2026-04-27 12:20) on post-fix code. The
baseline is a bug-hunting expedition — kill on anomaly, diagnose,
fix, re-run per feedback_stop_on_anomaly.md. Once healthy baseline
established, Phase 2 ships as principled improvement with clean A/B
against the trustworthy post-fix metrics.
Plans B / C / D are PAUSED, not cancelled. Files remain in
docs/superpowers/plans/. Resumption gate: post-fix baseline run is
bug-free or all surfaced bugs are addressed.
The val-Flat-collapse / Short-collapse / C51 expected-Q bias hypothesis
that motivated the 4-plan distributional-RL Thompson rollout was
largely a measurement artefact in the diagnostic infrastructure, not
a real policy pathology. Three layered bugs in actions_history_buf
init + reader + mag_stats attribution conspired to inflate val_dir_dist
Short, inflate active_frac, and pin wr_h/wr_f to zero. After fixing
all three (commits a86fba2b1 + b8788511c), val_dir_dist matches
val_picked_dir_dist within ~5pp — no collapse, diverse picks.
Status footer added to the spec documenting:
- what was actually wrong (3 measurement bugs)
- what the post-fix data shows (mild passivity bias from early
training, not a structural collapse)
- what is preserved (Phase 0 unit tests as latent infrastructure
for any future distributional Q-head; Aggregation Contract
pearl as a sound engineering invariant)
- what is cancelled (Plans B / C / D — Phase 1 audit, Phase 2
Thompson integration, Phase 3 long verification — superseded)
- future revival condition (if fresh L40S 30-epoch on post-fix
code shows val_picked_dir_dist itself collapsing toward Hold/Flat,
reopen)
Plan files remain in docs/superpowers/plans/ as historical record.
experience_kernels.cu line 1916 binned action_mag_per_sample by
actual_mag_core at every step, including trade-close events. But
unified_env_step_core forces `actual_mag = 0 (Quarter)` whenever
actual_dir is Hold/Flat (trade_physics.cuh:772) and trade closes
always land in Hold/Flat state — so every Half/Full close was
attributed to the Quarter bin. close_counts[Half] and close_counts[Full]
were structurally pinned to 0, giving wr_h = wr_f = 0 across all
training runs.
Fix: introduce `seg_mag_bin = is_close ? pre_mag_bin : actual_mag_core`.
At close events bin by pre_mag_bin (the magnitude of the position
being closed); at non-close events keep actual_mag_core (current
realized magnitude). pre_mag_bin is always 0/1/2 at close events
since exiting/reversing requires prev_sign != 0 → pre_trade_position
!= 0 → pre_frac > 0.001 → pre_mag_bin in {0,1,2}.
Smoke verification (5-epoch local): wr_h/wr_f remain 0 because
var_scale (1/(1+sqrt(var_q))) shrinks effective_max_pos to 10-19% of
broker max at smoke maturity → even Long Full target lands at
abs_pos ≈ 0.15 < 0.375 → all positions decode as Quarter; no
Half/Full positions exist for the fix to attribute. This is the
expected structural consequence of the var_scale design (uncertain Q
→ smaller position, conservative). The fix is latent correctness:
in mature L40S 30+ epoch runs where var_q drops and var_scale
grows above 0.375, Half/Full positions become reachable and
wr_h/wr_f will reflect their real realized win rates.
Without the fix, even mature training would show wr_h = wr_f = 0
because of the Hold/Flat → Quarter close convention masking real
per-magnitude win rates. Audit entry updated.
After done_flags[w]=1 (capital floor breach), backtest_env_step
early-returns without writing actions_history_buf for remaining slots
in [done_step, max_len). The prior zero-init decoded those slots as
Short Quarter Market Normal (action 0) via `dir = 0/27 = 0` and
inflated val_dir_dist's Short bucket / active_frac to a measurement
artifact masking real model behaviour.
Two-part fix:
1. `gpu_backtest_evaluator.rs::reset_evaluation_state`: replace
`memset_zeros` for actions_history_buf with
`cuMemsetD32Async(0xFFFFFFFFu32)` writing -1 sentinel. The Rust
readers already filter `if a < 0 { continue; }` so unwritten slots
are skipped correctly post-fix.
2. `backtest_metrics_kernel.cu`: add `if (act < 0) continue;` after
reading actions_history. The reduce-side metrics
(buy_count/sell_count/hold_count → active_frac/dir_entropy +
bnd_* trade-boundary detection) now consistently skip unwritten
slots. step_returns at those slots are still zero-init (correct)
so summing them with r=0 is a no-op.
Empirical impact (local 3-fold × 5-epoch smoke, RTX 3050 Ti):
val_dir_dist Short: 81-84% → 13-29% (matches val_picked within 5pp)
active_frac: 87-91% → 31-50%
dir_entropy: 0.57 → 0.83-1.02
The pre-fix "val-Flat-collapse" / "Short-collapse" pathology that
motivated substantial subsequent investigation (incl. the 4-plan
distributional-RL Thompson rollout draft) was largely a measurement
artifact from this bug surfacing differently before vs after the
Kelly cap fix (`0c9d1ee39`). Pre-Kelly the Kelly cap clamped most
Long/Short → Flat → actions_history was densely written with Flat
encoding → 80% Flat reading (real Kelly pathology + small artifact).
Post-Kelly the picks survive but the poor smoke-trained model
breaches capital floor often → many unwritten Short slots → 83%
Short reading (pure artifact). With both fixes, val_dir_dist now
reflects real model behaviour.
Audit entry updated in docs/dqn-wire-up-audit.md.
Algorithmic property test (CPU). Confirms Thompson exploration discovers
KNOWN +0.005 edge in 100 iterations on a 1-state bandit, while
argmax-only training never updates Q[Long].
Setup revised from plan A draft (option 3 — production-realistic):
p_long initial = [0.10, 0.20, 0.40, 0.20, 0.10] (uniform, E=0, has σ)
p_flat initial = [0, 0, 1, 0, 0] (δ(v=0), deterministic)
Argmax with strict-> ties at E=0 → always picks Flat → never explores
Long → Q[Long] stays at 0, never discovers edge.
Thompson samples Long > 0 with P≈0.30 → ~30 effective updates → mean
drifts toward +0.005, crosses Q[Flat]=0 within budget.
Plan's prior draft (initial p_long with mean=-0.015 + p_flat=δ(0)) was
calibration-bound: Thompson drift was directionally correct (-0.015 →
-0.005) but didn't cross zero in 100 iters. Revised setup eliminates
the artificial initial bias and matches production reality more
closely (Flat = δ(0) by construction; Long starts spread from random
init, then accumulates true edge).
Stop condition: if Thompson e_long ≤ e_flat with this setup, the
hypothesis is genuinely wrong and reward shaping must change before
proceeding to Phase 2.
Observed (local RTX 3050 Ti, ~0.00s test wall, ~1.57s 5-test suite):
argmax : e_long=0.000000, e_flat=0.000000 (asserts e_long ≤ 0.001 OK)
thompson: e_long=0.003550, e_flat=0.000000 (asserts e_long > e_flat OK)
All 5 Phase 0 tests pass: 0.A bias-reproduces, 0.B inverse-CDF,
0.C IQN symmetry, 0.D Thompson-reverses, 0.E synthetic-edge.
- Use u32::div_ceil for grid-dim arithmetic (style)
- Tighten u8::try_from to validate direction < B0_SIZE=4 instead of
fits-in-u8: kernel-OOB write fails loudly at the first malformed
index, not silently in production.
No semantic changes.
Per feedback_gpu_cpu_roundtrip.md, the per-seed dtoh in
launch_thompson_direction (10k iterations for Test 0.A, 100k for Test
0.B) violated the no-dtoh-on-hot-or-tight-loop invariant. Replaces
the per-seed kernel with a batched kernel (one thread per seed) and
the dtoh wrapper with a MappedI32Buffer using cuMemHostAlloc(
DEVICEMAP|PORTABLE) — same pattern as gpu_training_guard.rs
MappedBuffer.
Kernel changes (thompson_test_kernel.cu):
- thompson_direction_test_batched: replaces thompson_direction_test;
one thread per seed, writes via mapped device pointer with
__threadfence_system() for PCIe coherence.
- argmax_eq_test, compute_sigma_c51/iqn_test: __threadfence_system()
added before kernel exit.
Test refactor (distributional_q_tests.rs):
- MappedI32Buffer helper (test-utility version of the production
f32-only MappedBuffer).
- Single batched launch per test instead of N launches.
- Tests 0.A and 0.B preserved assertions; runtime drops from ~1.86s
(Test 0.A) and ~3.8s (Test 0.B) to ~0.15s each on RTX 3050 Ti.
Dead code: launch_thompson_direction (single-seed) and the OnceLock
KernelSet single-launch wrappers deleted; orphan code per
feedback_wire_everything_up.md.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Originally specified by Plan A Task 4 with P(d=0)≈0.2. That assumed
d=1's sample of 0 would break the tie when d=0's C51 sample is also 0.
The actual `thompson_direction_test` kernel uses strict-> with
`best_d=0` initialised, so d=0 wins whenever its sample is ≥ 0
(atoms v=0 or v=+1).
Corrected expected: P(d=0) = p[v=0] + p[v=+1] = 0.9. Test still
verifies inverse-CDF correctness — wrong math would deviate this
proportion measurably.
Also applies clippy::erasing_op / clippy::identity_op cleanup (`0 *`,
`1 *` index expressions) to Test 0.A from Task 3 review. No behaviour
change in 0.A.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Adds distributional_q_tests.rs with launch_thompson_direction and
launch_argmax_eq wrappers around thompson_test_kernel.cu. Tests follow
in subsequent tasks.
Implements inverse-CDF over C51 atoms, uniform-τ interpolation over IQN
quantiles, and argmax of E[Q] for eval mode. Plus diagnostic σ kernels
(compute_sigma_c51_test, compute_sigma_iqn_test) used by Test 0.F.
Exercised only by Phase 0 tests in distributional_q_tests.rs; Phase 2
production kernel will integrate the same __device__ __forceinline__
math into experience_action_select. No production callers in this
commit.
Plan A Task 1 of docs/superpowers/plans/2026-04-27-distributional-rl-thompson-plan-A-phase-0.md.
User correctly identified that CPU mirror function tests don't test
the production GPU code path. A bug shared between mirror and kernel
(translated identically wrong) would slip through. Mirror tests + a
single GPU bridge test were a weak compromise.
GPU-direct testing strategy:
- All Phase 0 kernel-correctness tests (0.A, 0.B, 0.C, 0.D, 0.F):
launch tiny test-only kernels with the SAME math the Phase 2
production kernel will use; assert properties of the output.
- Test 0.E (synthetic edge discovery): stays CPU. It tests an
ALGORITHMIC PROPERTY of Thompson exploration (does it discover
edge if edge exists?), not a kernel correctness property.
- All Phase 2 unit tests (2.A-2.D): GPU-direct against the
modified production kernel.
- Phase 0.F (real checkpoint extraction): unchanged — already GPU.
Local development uses RTX 3050 GPU (per memory user_dev_environment.md).
CI runs --ignored flag to skip GPU tests on CPU-only runners.
Time budget: Phase 0 was 1-2 days (CPU mirror); now 2-3 days
(GPU-direct, includes kernel wrapper setup half-day).
Other delta:
- Phase 0 deliverable file renamed: distributional_q.rs ->
distributional_q_tests.rs (no mirror functions, just tests +
kernel wrappers).
- Phase 2 unit tests rephrased to launch production kernel rather
than compare against CPU mirror.
- L1 verification gate runtime: seconds -> minutes (GPU launch
overhead per test).
The user's intuition was right: testing production directly is the
honest approach. Mirror was an optimization that traded correctness
for speed; with local GPU available the optimization isn't needed.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Self-review identified 8 major + 8 medium + 15 minor issues. All fixed:
MAJORS:
M1 Test 0.A: clarified — compare argmax(E[Q]), Boltzmann(E[Q]),
and Thompson sampling distributions; assertion is on relative
ordering across all three.
M2 IQN quantile count: replaced hardcoded `5` with N_IQN_QUANTILES
constant (defined per task #147 fixed-quantile design).
M3 "Converged checkpoint" definition: ≥60 epochs trained AND
val_sharpe stabilised (no >10% change over last 10 epochs).
Cites prior 60-epoch validation runs (task #80, train-7rgqd).
M4 R3 reframed: replaced "no issue" handwave with explicit
by-design tradeoff acknowledgment + cost analysis. Wasted
exploration is the cost of finding out whether edge exists.
M5 Test 0.D σ_long=0.05 justified: chosen to match expected order
of magnitude given typical |return| ~ 50bps; Phase 0.F
validates against real checkpoint.
M6 rng_ctr post-increment: clarified — matches existing pattern
at experience_kernels.cu:858 (no behaviour change).
M7 train_active_frac instrumentation: NEW Phase 2 deliverable —
existing HEALTH_DIAG only has val_active_frac, but L3 verifies
training-time active_frac. Spec now explicitly adds this
~10-line metrics.rs change as a Phase 2 deliverable.
M8 eps_dir cleanup code-level detail: explicit reference to
experience_kernels.cu lines 814-865; remove eps_dir from both
static EPS_FLOOR clamp AND adaptive boost block; verify
variable can be removed from kernel signature via grep.
MEDIUMS:
Med1 Current C51/IQN combination: clarified that compute_expected_q
blends per training schedule; Phase 2 replaces with explicit
0.5*E_C51 + 0.5*E_IQN equal weighting; Phase 0.F verifies.
Med2 Eval mode phrasing: "eval mode already sets eps=0 in existing
kernel" — no semantic override, factually correct.
Med3 Magnitude σ claim: clarified — magnitude branch likely has σ
bias in OPPOSITE direction (Full has larger σ; UCB would
prefer Full and worsen saturation). Empirical verification
deferred. Phase 0.F should also report per-magnitude σ.
Med4 Hierarchical sampling claim corrected: it's not about
balancing 50/50 (already 50/50). It's about decoupling
cluster-best decisions; clarified.
Med5 n_atoms vs N_IQN_QUANTILES: clarified — n_atoms variable per
config (currently 51); N_IQN_QUANTILES fixed at 5.
Med6 Conviction code: removed pseudo-code; references existing
implementation at experience_kernels.cu:1091; provides
implementation hint for E[Q] reuse.
Med7 Q-target propagation: clarified — uses full distribution
(C51 atom projection / IQN quantile regression), not just
E[Q]. Thompson modifies action selection only.
Med8 References: added Thompson 1933 (original), Bellemare 2017
(C51), Dabney 2018 (IQN) for theoretical foundations.
MINORS:
Min1 Date updated to 2026-04-27.
Min2-3 Argmax monotonic /2 simplified out — argmax(a+b) =
argmax((a+b)/2). Code clarity improved.
Min4 P(argmax picks Long) = 0 deterministic; reframed assertion.
Min5 Test 0.F structural assertions added: σ_C51[FLAT] < 0.01 ×
σ_C51[LONG]; same for IQN; E[Q_FLAT] > E[Q_LONG]; argmax
picks FLAT; Thompson P(LONG)+P(SHORT) ≥ 0.20.
Min6 -INFINITY → CUDART_INF_F (CUDA convention).
Min7 dir_idx scope: comment notes it's declared earlier in kernel.
Min8 action_select args: explicit — three buffers exist on GPU
but not currently passed; new params, no new buffers.
Min9 Phase 0 time math: 5 hours tests + 1 hour enumeration + 2
hours 0.F + (3 hours runtime if checkpoint training needed,
runs in parallel). Honest budget.
Min10 "Two-stream" → "5-Layer Gate" header.
Min11 Plan 5 reference uses full path consistently.
Min12 Plan B time budget: explicit 1 day if pass; 2-5 days if bug.
Min13 active_frac: clarified Long+Short combined, not per direction.
Min14 train_active_frac: now in Phase 2 deliverables (see M7).
Min15 "20 mechanisms" → 21, with sub-counts in section headers.
Spec now 615 lines, comprehensive coverage of:
- Pearl + theoretical foundation
- Problem statement (with bias-might-be-correct caveat)
- Architecture (Thompson at training, argmax at eval)
- Train vs eval distinct selectors with behavior-change disclosure
- Direction-only scope with magnitude σ-bias warning
- Conviction stays E[Q]-based (no Kelly cap jitter)
- Interaction matrix: 21 mechanisms in 3 categories
- 6 v2 enhancements documented + deferred
- Phase 0/1/2/3 with tests, exit gates, time budgets
- 5-layer verification + train_active_frac instrumentation
- 8 risks with mitigations + 5 stop conditions
- What v1 doesn't touch (referencing interaction matrix)
- References (Thompson 1933, Bellemare 2017, Dabney 2018, etc.)
- Aggregation contract (project-wide pearl, enforced)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per-user direction: every existing mechanism in the DQN system MUST be
explicitly considered for interaction with Thompson, and Thompson itself
MUST be examined for system-specific improvements beyond vanilla.
INTERACTION MATRIX (3 categories, 20 mechanisms):
Category 1 — Compose with Thompson (no change required):
Counterfactual reward, B.2 novelty bonus, PopArt, Saboteur,
Curiosity, NoisyNets/VSN, Distillation, CQL, Polyak target EMA,
HER, PER, Replay warm-start, Multi-fold validation harness.
Category 2 — Trivially adapt to Thompson (one-line changes):
D7/N7 contrarian sign flip (negate the SAMPLE), cosine epsilon
schedule (still applies to mag/ord/urg), per-sample epsilon (IQL
expectile gap), adaptive Boltzmann tau (still applies to mag/ord/urg).
Category 3 — Take precedence over Thompson (hard constraints):
Plan-based action lock (Thompson sample discarded if plan active),
per-magnitude Kelly cap, trail stop, capital floor breach.
Critical insights from the audit:
1. NoisyNets is ALREADY a form of training-time Thompson at the
parameter level. Output-space Thompson stacks on top —
total exploration = parameter-space ⊗ output-space (multiplicative).
2. Curiosity is ORTHOGONAL to Thompson — Thompson explores actions
whose Q is uncertain; curiosity explores states whose dynamics
are uncertain. Both axes desirable; no conflict.
3. Plan lock takes precedence; same as currently with Boltzmann.
OUTSIDE-THE-BOX v2 ENHANCEMENTS (deferred to follow-up specs):
v2.1 Triple-source Thompson (C51 + IQN + Ensemble) — incorporate
the existing ensemble Q-head as 3rd uncertainty source.
v2.2 Persistent Thompson (anti-churn for HFT) — bias sampling
toward current direction, ISV-driven; reduces tx_cost from
Long/Short oscillation across bars.
v2.3 CVaR-aware eval (risk-adjusted deployment) — eval picks
argmax(E[Q] − λ·CVaR_α[Q]); risk-aware decision making for
production with real capital.
v2.4 Information-Directed Sampling (Russo & Van Roy 2014) — picks
action minimizing regret²/info_gain; more efficient than
vanilla Thompson when learning saturates.
v2.5 Hierarchical Thompson on (trade vs no-trade) → (which
direction) — addresses 50/50 structural advantage of no-trade.
v2.6 Composition with curiosity-driven exploration — explicit
coupling beyond reward-side composition.
Each v2 enhancement gets its own spec/plan when prioritised. Vanilla
Thompson is v1; ships first; verified independently.
ALSO FIXED (from earlier self-review):
- Pearl claim softened: only the C51 Hold/Flat bias is directly
attributed; other historic bugs had different mechanisms.
- TFT entry removed from contract table — TFT is a Variable
Selection Network (feature processor), not a Q-head. Replaced
with generic "Future Q-head additions" placeholder.
- Eval direction = argmax E[Q] explicitly flagged as a behavior
change from current Boltzmann-with-tau (val_dir_dist will be
more concentrated than current).
- Phase 0.E budget reduced (1-2 hours, not 1 day) — synthetic
bandit is ~50 lines of Rust, not full RL training loop.
- Phase 0 enumerates existing checkpoints before training new one.
- Architecture diagram parenthesis fixed.
- Conviction implementation note: compute E[Q] once, reuse for
conviction AND eval-mode argmax — no redundant computation.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
User correctly challenged the "band-aid removal" framing. All fixes
shipped during the val-Flat-collapse investigation addressed real bugs
at their respective layers and should be preserved:
- Kelly cap warm-branch (0c9d1ee39): post-decision physics layer.
Thompson-independent. KEEP.
- Train Return display + Sharpe annualization (non-tau parts of
7a3d88646): display/metric layer. Thompson-independent. KEEP.
- Direction Boltzmann tau-floor (tau part of 7a3d88646) +
adaptive eps_dir floor (d54b49efc): gates inside direction-branch
action selection. Phase 2 replaces direction-branch action
selection wholesale (eps-greedy + Boltzmann → Thompson), so these
direction-only code paths become structurally unreachable.
The latter two are NOT band-aids being removed because Thompson is
better. They are dead code being cleaned up because Thompson replaces
the surrounding mechanism. Magnitude/order/urgency branches keep their
existing eps-greedy + Boltzmann + tau-floor + EPS_FLOOR paths intact.
Reframed Phase 3 deliverable: "direction-branch dead-code cleanup"
with explicit rationale (per feedback_no_legacy_aliases.md and
feedback_no_partial_refactor.md). 0.5 day budget instead of 1.
Also clarified eval action selection: argmax of (E[Q_C51]+E[Q_IQN])/2
is correct. Bellman backup is a Q-learning UPDATE rule, not an
action-selection rule. Once Q is learned, optimal policy is greedy
argmax of learned Q. Online Bellman lookahead at eval would require
a forward model of market dynamics — not available, not standard.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
