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>
Revises the C51-bias spec after deeper review surfaced 12 design gaps,
of which 4 were critical:
1. Train-only vs train+eval ambiguity — UCB at eval would conflate
"model recommends Long" with "model is uncertain about Long",
inflating reported edge. CRITICAL for trading where eval drives
real capital decisions.
2. Thompson sampling is more principled than UCB:
- parameter-free (no κ to tune)
- uses distribution directly without scalar reduction
- naturally explore-exploit balanced via distribution shape
3. c51_alpha is the wrong blend weight (it's C51-vs-MSE-warmup, not
C51-vs-IQN). Equal-weight average of C51 and IQN samples is the
structural choice — no tuned blend weight needed.
4. The bias might be CORRECT BEHAVIOUR — model rationally choosing
Flat when no edge has been discovered. Phase 0 must include a
synthetic-edge test (controlled MDP with KNOWN positive Long
expected value) to verify Thompson can discover edge if it exists.
Other gaps fixed:
- Eval at argmax E[Q] (not Boltzmann, not Thompson)
- Pearl wording broadened to cover ensembles + future methods
- Ensemble Q-head added to aggregation contract table
- Explicit caveat: NEVER extend Thompson to magnitude branch (would
worsen existing magnitude saturation)
- Phase 0.F uses CONVERGED checkpoint (≥30 epochs), not 2-epoch run
- L4 long smoke (30 epochs, ~1 hour) added — Thompson edge discovery
needs longer feedback loop than 5 epochs
- Phase 3 explicitly removes eps-floor + tau-floor band-aids
(Thompson replaces direction Boltzmann; band-aids become dead code)
- Conviction stays E[Q]-based, not sample-based (avoid Kelly cap
jitter from stochastic samples)
Architecture (Thompson only, no UCB):
TRAINING: dir_idx = argmax(0.5 × (sample_C51(d) + sample_IQN(d)))
magnitude/order/urgency: existing Boltzmann + ε-greedy
EVAL: dir_idx = argmax(0.5 × (E[Q_C51] + E[Q_IQN]))
magnitude/order/urgency: existing Boltzmann (eval mode)
Direction-branch ε-greedy + Boltzmann are REMOVED — Thompson is the
exploration mechanism. No new GPU buffers; existing C51 atoms + IQN
quantiles passed to action_select.
5-7 days active work across 4 sub-plans; each gets its own
writing-plans cycle.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Designs the structural fix for the C51 expected-Q Hold/Flat bias exposed
by the Kelly val-Flat-collapse fix. The bias is a manifestation of a
deeper project-wide pearl:
"Distributional RL aggregation discards uncertainty;
action selection must restore it."
Any value head representing Q as a distribution (atoms, quantiles,
ensembles) MUST expose both E[Q] and σ(Q) to action selection. Boltzmann
on E[Q] alone produces structural bias toward low-variance actions
regardless of expected payoff — the C51+IQN Flat-attractor is one
instance of this lost-information pattern.
Fix: extract σ(Q) from BOTH C51 atoms (closed form) and IQN quantiles
(IQR/1.349), blend by loss-time weight, feed Q_eff = E[Q] + κ·σ
(κ=1.0 structural identity) to direction-branch Boltzmann ONLY.
4-phase implementation:
Phase 0 — TDD hypothesis verification (Rust mirror functions + 5 unit
tests including GPU integration on real checkpoint)
Phase 1 — Audit existing reward levers (B.2, CF, PopArt, Q-target)
via 6 unit tests; fix any bugs found
Phase 2 — UCB integration: new compute_q_with_uncertainty kernel,
modified action_select, Rust orchestration, project-wide
aggregation contract in dqn-wire-up-audit.md
Phase 3 — Verification per Plan 5 Task 5 multi-seed × multi-fold
5-layer verification gate; 8 risks with mitigations; 4 stop conditions
that halt execution and force redesign.
Existing band-aid fixes (Kelly cap, eps-floor, tau-floor) stay — they
address symptoms at different layers. UCB adds the missing aggregation
step that was the common root across all the symptoms.
Direction-branch only — magnitude/order/urgency don't have the
Flat-attractor (atom-mass collapse asymmetry).
5-7 days active work across 4 sub-plans; each gets its own
writing-plans cycle.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Prior C51 bias fix (commit 7a3d88646: ISV-adaptive Boltzmann tau floor)
had no measurable effect on the post-training Hold/Flat collapse —
empirically confirmed in train-bscl2 epoch 2: val_dir_dist
[short=0.135 hold=0.358 long=0.142 flat=0.364], identical to the
pre-fix run [short=0.136 hold=0.348 long=0.148 flat=0.367].
Root cause: once Q-values reflect tx_cost-driven aversion, Boltzmann
correctly samples the biased Q distribution regardless of tau. Tau
adjustments protect cold-start exploration but can't combat learned
preferences. The 2% static eps_dir floor allows only 0.5% random
sampling per direction — too little to break the Q-value lock-in or
generate enough Long/Short experiences for edge discovery.
Fix:
if (ISV available) {
passive_pressure = clamp(0, 1, 1 − ISV[71]/max(ISV[72], 1e-4))
eps_dir = max(eps_dir, 0.5 × passive_pressure)
}
ISV[71] = TRADE_ATTEMPT_RATE_EMA (current Flat→Positioned rate, B.2 producer)
ISV[72] = TRADE_TARGET_RATE (target frozen at epoch 5 from measured EMA)
Feedback semantics:
- attempt_rate >= target → passive_pressure=0 → eps_dir at baseline 0.02
- attempt_rate = 0 (fully passive) → passive_pressure=1 → eps_dir floor=0.5
- intermediate → linear blend
The 0.5 ceiling is a structural blend point (half random / half policy)
— maximum exploration that still preserves directional Q-signal
propagation through the replay buffer. Not a tuned magnitude.
Cold-start safety: ISV[72] is 0 until epoch 5 freeze; with the 1e-4
target floor, passive_pressure clamps to ≈1 immediately, but eps_dir
also has a baseline 0.02 EPS_FLOOR so the formula's max() picks
whichever is larger. Once ISV[72] freezes, the feedback loop activates
properly.
Eval mode unaffected (eps logic gated behind !eval_mode).
Direction branch only — magnitude/order/urgency don't have the
Flat-attractor problem.
ISV-driven, no new ISV slots, no tuned constants per
feedback_isv_for_adaptive_bounds.md and feedback_adaptive_not_tuned.md.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
CONCERN 1 — train Return display (financials.rs)
Replace `mean_per_bar × bars_per_year` ("annualized arithmetic return")
with the actual cumulative compounded return: `prod(1+r) − 1` via
log-space sum for f64 stability. Step floor at 1+r >= 1e-10 so a -100%
bar gives -23 log contribution rather than -inf; result bounded above
-100%.
The old formula multiplied per-bar mean by ~98K, producing values like
Return=-51234% that LOOKED like portfolio collapse but were actually
just label artefact: -52 bps mean × 98,280 bars/year = -51,234%, while
the actual compounded return over the rollout was bounded and finite.
"Total Return" now means what a trader expects.
CONCERN 2 — val_Sharpe annualization source-of-truth (training_loop.rs +
gpu_backtest_evaluator.rs)
val_Sharpe was using `m.sharpe` from the kernel (annualised by
`sqrt(bars_per_day × trading_days_per_year)` at evaluator init), but
val_Sharpe_raw recomputed the divisor from `self.hyperparams.bars_per_day
* 252.0` in Rust. Empirically the two diverged by 5.4× (val_Sharpe=143.38
/ val_Sharpe_raw=0.0841 = 1704, expected sqrt(98280) = 313.5). Root cause
unclear without runtime instrumentation but the structural fix is to
read the EXACT factor the kernel applied: new
`GpuBacktestEvaluator::annualization_factor()` returns the f32 stored at
init time. Rust now divides by that, so the two numbers can never drift
regardless of how `bars_per_day` flows through the config layers.
Per `feedback_no_partial_refactor.md` — every consumer of a shared
contract must use the same source.
CONCERN 3 — C51 expected-Q Hold/Flat bias (experience_kernels.cu)
Direction-branch Boltzmann tau floor: replace static `0.01` with
ISV-adaptive `max(ISV[21], 0.01)` where ISV[21] is q_dir_abs_ref (EMA
of mean |Q| across direction bins, same signal that drives conviction).
The C51 expected-Q biases Flat above directional actions when the policy
has no measurable edge — Flat returns are exactly 0 (no position change),
collapsing C51's distribution to a delta at 0; directional returns
concentrate slightly below 0 from tx_cost without compensating edge,
giving E[Q_directional] < E[Q_flat] = 0. With Q-spread on the order
of 0.01-0.1 and the static 0.01 tau floor, exp(Q/tau) ratios approached
deterministic argmax, locking the policy into Hold/Flat within ~1 epoch
and preventing the exploration needed to discover real edge.
Validated empirically by the kelly-fix-multifold run (train-multi-seed-bs9m5):
val_dir_dist epoch 1 [S=.23 H=.17 L=.42 F=.18] (active 65%)
val_dir_dist epoch 2 [S=.14 H=.35 L=.15 F=.37] (active 29%)
That 35→72% Hold+Flat shift in one epoch is the C51 attractor in action.
The ISV-adaptive floor preserves relative spread for sampling: when
Q-magnitudes grow during training, the floor scales with them, so the
policy never collapses to deterministic argmax until q_range exceeds
the network's typical Q magnitude — i.e., until spread represents
real, scale-significant edge. Coherent with the existing conviction
formula (same ISV reference). No tuned constants per
`feedback_isv_for_adaptive_bounds.md` and `feedback_adaptive_not_tuned.md`.
Cold-start fallback retains 0.01 minimum so kernel doesn't divide by
zero pre-first-update. Once ISV producer fires (every epoch), adaptive
floor takes over.
7 financials unit tests pass. Workspace cargo check clean.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
VAL DIAGNOSTIC PROOF (train-4fpzx step=400):
pick=Long Full (target=0.66) prev_pos=0.0 → pos_post=0.0 actual_dir=Flat
trail=0, margin can't clip to 0 → only Kelly cap zeroed the target.
ROOT CAUSE:
effective_kelly = maturity * kelly_f + (1 - maturity) * warmup_floor
Cold-start fix (commit 2c97e0436) protected `warmup_floor` so it never
collapses to zero at maturity=0. But the warm branch was left exposed:
once maturity → 1 (>=10 completed trades), the blend collapses to
`kelly_f` alone, and `kelly_f = 0` is the natural state of
(payoff*win_rate - (1-win_rate)) / payoff
with balanced priors and small actual returns. Val environments with
pure per-bar P&L (no saboteur/shaping perturbations like training)
settle into this regime within ~10 trades — after which every non-Hold
target gets clamped to 0 deterministically. Identical bootstrap-deadlock
pattern to the IQN trunk SAXPY.
EVIDENCE:
val_picked_dir_dist [short=0.19 hold=0.20 long=0.39 flat=0.21] (kernel pick)
val_dir_dist [short=0.0001 hold=0.20 long=0.0000 flat=0.80] (post-physics)
100% of Long picks and ~99.95% of Short picks become actual_dir=Flat.
Hold passes through 1:1 (Hold skips margin/Kelly/trail in env_step).
VALDIAG step=400: act=77 (Long Full, target=0.66) prev=0.0 pos_post=0.0
-> confirms target zeroed before execute_trade; trail=0 rules out trail;
margin cap can't produce 0 with equity=$35K vs margin/contract=$17.9K.
FIX:
effective_kelly = max(kelly_f, warmup_floor)
The conviction-and-health-driven warmup_floor (in [0.5, 1.0]) becomes a
permanent minimum cap. `kelly_f` only takes over when the policy has
demonstrated enough edge to *exceed* the floor. Preserves design intent
(Kelly drives sizing once stats mature with real edge) while preventing
the bootstrap deadlock in environments where balanced trades naturally
yield kelly_f = 0.
All adaptive ISV-driven signals; no tuned constants.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
First 10 bars proved the env_step kernel produces actual_dir=Long for
Long picks correctly at cold start. So the gate that turns 100% of
Long/Short picks into actual_dir=Flat in val_dir_dist must activate
LATER in the window — likely past the first chunk boundary.
Extended sampling: first 10 bars, every 100 bars to 1000, every 1000
bars to 10000, every 10000 bars throughout. Also adds a "mismatch"
trigger that fires when picked dir is Long/Short but actual_dir is
not — every 137 bars to keep printf volume bounded. cash field added
to spot capital drain or anomalous accumulation.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Captures action_val + decoded dir/mag, max_position, prev_position,
position post-trade, actual_dir, actual_mag, trail_triggered,
conviction, health, value, max_equity inside backtest_env_step_batch.
Localises which step in unified_env_step_core zeros target_position
(or position) for non-Hold picks.
After tracing the math, every analytical candidate (margin cap, Kelly
cap with health-coupled warmup floor, trail stop) returns a positive
target at cold start with my fix. val_dir_dist still locks at 0% Long
while val_picked_dir_dist shows 17% Long — only direct kernel-side
observation can show what's actually clamping.
Removed once the gate is identified.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
