Spec: docs/superpowers/specs/2026-05-28-state-conditional-q-synthesis.md
§3 Sub-phase 2.1.
Replaces the distributional Bellman bootstrap
(fused_select_and_project_bellman) with the Dueling-DQN scalar
baseline target
y(s_t, a_t) = r_t + γⁿ × V_target(s_{t+1})
advantage_scalar = y − V(s_t)
projected via two-bin linear interpolation onto the action's atom
support. The Q-head's atom logits are now interpreted as the
advantage distribution A_dist(s, a); the CE loss between
softmax(A_dist(s_t, a_t)) and the projected target trains A to
model deviation from the scalar V baseline.
Changes:
- NEW crates/ml-alpha/cuda/bellman_scalar_target_project.cu —
device-side scalar Bellman target + per-action atom projection.
Grid (B, 1, 1), Block (Q_N_ATOMS=21, 1, 1). No atomicAdd, no
cross-thread accumulation: scalar advantage produces a delta on
exactly two atoms so each thread writes its own atom value.
- crates/ml-alpha/build.rs registers the new kernel.
- crates/ml-alpha/src/rl/dqn.rs:
* Loads bellman_scalar_target_project from its own cubin.
* Adds DqnHead::bellman_scalar_target_project wrapper.
* Adds host-side composed_q_value(v_scalar, a_logits,
atom_supports) -> Vec<f32> implementing the Dueling
composition Q = V + A - mean_a(A) for test oracles.
- crates/ml-alpha/src/trainer/integrated.rs:
* Allocates sampled_v_curr_d + sampled_v_target_tp1_d for
V evaluated on PER-sampled hidden states.
* step_synthetic_body Bellman build: replaces forward_target +
fused_select_and_project_bellman with V forwards on sampled
buffers + the new scalar Bellman projection (single source of
truth — same swap applied to dqn_replay_step for k>0 path).
* Test accessors isv_host_slice_mut_for_test,
isv_dev_ptr_for_test, dqn_head_ref so kernel-oracle unit
tests can drive the new kernel without a full mega-graph.
- NEW crates/ml-alpha/tests/q_projection_unit.rs — pre-impl oracle:
projects y=0 to center atom, y=V_MAX to last, y=V_MIN to first,
arbitrary in-range cases recover EV within Δz/2, verifies V_t
enters as a subtraction. Plus pure host-side check of
composed_q_value.
All Phase 2.1 unit tests pass (5 CUDA --ignored + 2 host-side).
Smoke status: BLOCKED on pre-existing NaN abort at step 4 affecting
HEAD (c52282fb4 + f11bab542 + 7e38e46e6) — l_pi=NaN, l_aux=NaN at
step 4 reproduce IDENTICALLY with and without Phase 2.1 changes.
Per feedback_stop_on_anomaly, the pre-existing NaN must be diagnosed
before the Phase 2.1 smoke gate (l_v<2.0, l_q<3.0,
action_entropy>1.5 for 500 steps) can be validated. The Phase 2.1
changes themselves are structurally sound — unit tests confirm
correct projection and the smoke loss trajectories at steps 0-3 are
within ~1% of HEAD (l_q 2.9484 vs 2.9926, l_v 6.3293 vs 6.3291).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Latent bug surfaced by f11bab542 test audit. `h_mag_per_bucket_kernel`
launched with block_dim=32 (single warp) but BUCKET_DIMS terciles are
[43, 43, 42] — channels 32..bdim were silently dropped from the sum,
producing under-counted Controller D dead-bucket signal in production
(perception.rs::h_mag_cfg).
Fix: grow block_dim to 128 (smallest pow2 ≥ MAX_BUCKET_DIM=96), expand
shared mem from `__shared__ float sdata[32]` to `sdata[128]`, change
reduction stride from 16→32→64 (single-warp shuffle equivalent) to
64→32→16→8→4→2→1 (multi-warp block-tree). No atomicAdd, no warp
divergence in tail lanes (uniform predicate tid < bdim).
Changes:
- bucket_transition_kernels.cu: kernel block-tree reduction over 128
lanes, shared mem 128 floats, launch comment updated
- perception.rs::h_mag_cfg: block_dim 32→128, shared 32*4→128*4
- bucket_transition_kernels.rs test: launch config matches new contract
+ bug-surfaced comment replaced with fix-landed reference
Test: tests/bucket_transition_kernels::h_mag_per_bucket_kernel_computes
_mean_abs_per_bucket now passes (was deliberately failing in f11bab542
to surface this bug per feedback_no_todo_fixme — tests assert
invariants not observed behavior).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Updates 13 test files that were locked to stale kernel contracts.
Each change verified against the actual kernel/code (cuda/*.cu and
crates/ml-alpha/src/*) — not normalised to broken behavior:
Controller bootstrap values (isv_bootstrap.rs,
r5_controllers_and_soft_update.rs):
* γ → 0.995 — `rl_gamma_controller.cu` GAMMA_MIN raised per
`pearl_edge_lives_at_wave_timescale_not_tick` (long-horizon credit
assignment on surfer-philosophy timescale)
* entropy_coef → 0.5 — `rl_entropy_coef_controller.cu` COEF_MAX
bootstrap at deficit=1.0 per
`pearl_pi_actor_collapses_without_entropy_floor` (cold-start
emergency exploration)
* per_α → 0.4 — R9 derive-from-input pattern at kurt_excess=0
* τ structurally pinned at TAU_BOOTSTRAP=0.005 — TAU_MAX (slot 573)
equals bootstrap, controller has zero headroom upward (intentional
per slot doc-comment)
Topology / horizon consolidation (bucket_transition_kernels.rs,
cfc_step_per_branch.rs, heads_block_diagonal.rs, heads_bit_equiv.rs,
bce_grad_finite_diff.rs):
* N_HORIZONS: 5 → 3 (tercile layout: cut points [43, 86], dims
[43, 43, 42]) — kernel-side `bucket_transition_kernels.cu` /
`bce_loss_multi_horizon.cu` / `heads.rs` were consolidated
* MAX_BUCKET_DIM: 28 → 96 (kernel #define MAX_BUCKET_DIM = 96)
* Per-test asserts mirror tercile assignment via shared
`tercile_bucket(c)` helper
Done-gating semantics (r3_ema_advantage.rs):
* `compute_advantage_return.cu` line 30: `advantages[b] = done *
(ret - vt)` — V learns from trade closes only, π trained via
target-Q distillation + SAC entropy. Test now two-cohorts
(done=1 sees -k, done=0 sees 0)
Anti-collapse mechanisms (r4_action_kernels.rs):
* Thompson kernel applies ISV-driven epsilon-greedy floor (slot 588
= 0.05) per `pearl_pi_actor_collapses_without_entropy_floor`.
Wins/100 expectation widened to (1 − 0.05) × 100 ± 3σ ≈ ≥88
Confidence gate (trade_management_kernels.rs):
* `rl_confidence_gate.cu` does NOT check `lots != 0 → skip` — the
position-conditional escape was moved to the Phase 2.3 state-
mandatory action mask (separate kernel). Case 3 now asserts gate
fires regardless of position
Smoothness controller (smoothness_lambda_controller_invariants.rs):
* TARGET_K_RATIO[2] = 0.1 (sqrt-scaling fix) — kernel comment at
`smoothness_lambda_controller.cu:38-45` explains the migration
from geometric to sqrt anchoring
PER buffer migration (r7d_per_wiring.rs):
* CPU `trainer.replay` field replaced by GPU-resident
`trainer.gpu_replay.replay_len_d` (graph-safe push). Test reads
the counter via the canonical mapped-pinned pattern
(`MappedI32Buffer` + `memcpy_dtod_async` + sync) per
`feedback_no_htod_htoh_only_mapped_pinned`. Asserts mechanism
invariants (monotone growth, bounded by MAX_INJECTS_PER_STEP,
saturates at capacity) — looser than the old "len == step"
because HER + n-step push counts vary per step
EMA-input sentinel-zero loops (isv_bootstrap.rs,
r5_controllers_and_soft_update.rs):
* Switched from a sprawling negative exception list ("any slot in
[417, RL_SLOTS_END) except these N constants") to a positive
enumeration of the ~7 documented EMA-input slots. New ISV
constants no longer break the test by being added to the slot
range
Stale fixture (perception_forward_golden.bin):
* Old golden was captured with probs.len() = 160; current model
produces 96. Regenerated. Test logic captures-on-first-run +
compares-on-subsequent-runs is unchanged
Latent kernel bug surfaced (bucket_transition_kernels.rs::
h_mag_per_bucket_kernel_computes_mean_abs_per_bucket):
* `h_mag_per_bucket_kernel` launches with block_dim=32 but
BUCKET_DIMS=[43, 43, 42] — only the first 32 channels contribute
to the sum while the divisor uses the full bdim. Production
launch site `perception.rs::h_mag_cfg` uses the same block_dim,
so Controller D's dead-bucket signal is silently under-counted.
Test asserts the IDEAL behavior (full-bucket mean) and DELIBERATELY
FAILS until the kernel is fixed — per
`pearl_tests_must_prove_not_lock_observations`
Implements §3 Sub-phase 2.0 of the state-conditional Q synthesis spec
(docs/superpowers/specs/2026-05-28-state-conditional-q-synthesis.md).
Adds a done-gated trade-magnitude envelope that clamps the V regression
target before MSE, bounding `(v_pred - target)²` regardless of
fat-tail trade-close spikes from `returns = r + γ(1-done)·v_tp1`. The
envelope decouples from the C51 atom span (slots 484/485) so it can
tighten below the atom span floor as the trade-magnitude EMA narrows.
Wiring:
* New ISV slots 593-596: V_MAGNITUDE_EMA, V_TARGET_MAX, V_TARGET_MIN,
V_TARGET_K. RL_SLOTS_END = 597.
* New kernel `rl_v_target_envelope_update.cu` — done-gated Wiener-α
EMA of `max(|reward| | done)`; derives `V_max_eff = +k × ema,
V_min_eff = -k × ema`. Sparse-aware (skip update on no-done steps),
first-observation bootstrap, single-thread single-block.
* `v_head_bwd.cu` — clamp `r_target` to `[V_min_eff, V_max_eff]`
before computing MSE. Updates `ValueHead::backward` Rust signature
to take the ISV device pointer.
* Trainer integration: single helper `launch_rl_v_target_envelope_update`
is invoked from all three reward-pipeline paths (public
`launch_apply_reward_scale`, `step_with_lobsim_reward_and_train` —
two inline sites) — single source of truth per
`feedback_single_source_of_truth_no_duplicates`.
* Bootstrap ISV slots: V_max=+200, V_min=-200, k=3.0 (matches
dd049d9a4 baseline avg trade tail per spec). Bumps ISV constant
array from 115 → 118 entries.
V plateau LR decay (also called out in the spec) was already wired
end-to-end via `rl_lr_controller` (slots 433/434/435/438) — no
duplication needed.
Smoke validation:
* Real-data 1000-step run (`alpha_rl_train --n-backtests 128`):
max l_v = 0.6024 throughout 1000 steps (well under the spec's
5.0 ceiling).
* Synthetic Rust integration test
(`tests/v_head_stabilization_smoke.rs`): 1000 steps batch=1,
cold-start window 100 steps, post-warmup max l_v = 0.0035.
Layer 3 of the four-layer defense per pearl_wr_above_50_with_negative_pnl_loss_cutting.
The wr-targeting controller (Layer 4) successfully drives wr=0.55+ but
PnL spirals negative because Q can only learn exit timing from sparse
close-event rewards. Without per-step exit signal during open trades,
Q can't learn "close losing trade early".
Fix: add continuous drawdown penalty during open trades.
- rl_fused_reward_pipeline.cu PHASE 5 shaping:
if (current_lots != 0 && unrealized < 0)
r += unrealized * rate * reward_scale
- Penalty-only (no symmetric reward for unrealized gain) — avoids
exposure-positive bias per pearl_event_driven_reward_density_alignment
- Reads unrealized_pnl from PosFlat offset 24 (added this session)
ISV slot 592 RL_DRAWDOWN_PENALTY_RATE_INDEX, bootstrap 0.001.
Conservative rate to avoid "Q learns trading is punished" pathology
(observed when rate was 0.01 in earlier session).
Local smoke 1000 steps (b=128):
- wr trajectory: 0.31 → 0.46 (Layer 4 controller still working)
- pnl: -$105k (vs -$5M without layer 3 at cluster scale)
- Completes clean, no NaN
The per-step penalty magnitude (~1e-5 to 1e-3 scaled units at rate=0.001)
provides Q with the exit gradient it was missing. Cluster validation
will show if surfer + positive PnL emerges at 5k+ steps.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The adaptive WIN/LOSS controller (rl_reward_clamp_controller.cu:303-309)
was driving LOSS away from the static 3.0 bootstrap toward 1.0 floor
via ratio × WIN. With L/W trade ratio near 1.0, LOSS ≈ WIN (symmetric
clamp), Q loses loss-aversion, and wr drops to 0.30 trend-follower.
Reverted to static bounds matching dd049d9a4 (wr=0.567 baseline):
- WIN remains at bootstrap 1.0
- LOSS remains at bootstrap 3.0
- Controllers still observe margin/ratio for diagnostics but don't write
Per pearl_loss_clamp_controls_entropy_stability: LOSS=3.0 gives stable
entropy + high wr; the entropy collapse concern is addressed separately
by Thompson sampling floor (kept).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The max_hold force-close used event-timestamp nanoseconds, but MBP-10
events fire at variable rates. At dense periods (~100 events/sec), a
60s cap = 6000 events = 6000 training steps — effectively never fires.
Observed: alpha-rl-final-f0 step 1500+ → dones=0 (model "hold forever"
attractor), pnl growing from unrealized drift, not real trades.
Fix: switch to step-based max_hold.
- Added entry_step: u32 to PosFlat at offset 28 (struct 28 → 32 bytes)
- order_match.cu + resting_orders.cu record entry_step on open/flip
- Replaced ns-based check with step-based: current_step - entry_step
- max_hold_steps_d replaces max_hold_ns_d in LobSimCuda
- alpha_rl_train.rs sets max_hold_steps=100 (matches γ-derived min_hold)
- New fill_u32.cu kernel; deleted fill_u64.cu (no remaining consumers)
The mega-graph captures self.isv_dev_ptr + 548*4 as current_step pointer
so rl_increment_step's in-graph ISV[548] mutation drives the lobsim
clock live during fast-path replay.
Local smoke 1000 steps (b=32, per=4096):
- 750 trades, dones fire consistently every 100-step window
- avg_hold stabilizes at 35 (well within 100-step cap)
- No "never close" collapse
- All architectural fixes preserved
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
dqn_distributional_q.cu:285 atomicAdds ce/B into ss_q_loss (kernel-side
B divide, matching ppo_clipped_surrogate.cu:269-271's b_inv pattern).
Host previously divided by B AGAIN at integrated.rs:3944 and :5196 —
true mean CE (≈ 3.0 for 21-atom softmax) appeared as 3.0/1024 ≈ 0.003.
The "baseline l_q=1.45 at step 500" was actually a missing-zero-out
artifact: ss_q_loss accumulated 500 × 0.003 = 1.46 across steps before
bf8887400 fixed the zero-out, exposing the underlying double-divide.
Q learning has been working correctly the whole time — only the
diagnostic display was wrong.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
cuBLAS GEMM calls inside CUDA-graph capture regions cause
CUDA_ERROR_STREAM_CAPTURE_INVALIDATED on first use of new (m,n,k)
shapes. The mega-graph was silently failing → fast-path replay was
a NO-OP → sps stayed at ~10.
Phase 1 — DQN distributional Q head:
- New crates/ml-alpha/cuda/dqn_q_head_fwd_bwd.cu (3 kernels)
- Removed cuBLAS field + gemm_f32 helper from dqn.rs
Phase 2 — IQN ensemble heads:
- New crates/ml-alpha/cuda/rl_iqn_matmul.cu (3 kernels)
- Removed cuBLAS field from iqn.rs
Phase 3 — Mamba2 SKIPPED: workspace pre-warms during 65+ eager
warmup steps before mega-graph capture.
Local smoke (RTX 3050 Ti, b=128, 500 steps):
- Mega-graph captures cleanly at step 67
- 8.9 sps avg, 11 sps peak post-capture (GPU-bound on mobile GPU)
- l_q=0.024 (healthy), l_pi rising, V converging
- Production L40S should see full mega-graph speedup now
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Two perf fixes for mega-graph throughput.
1. Inline ema_update_per_step(q_divergence) — eliminates the per-step
cuMemAllocAsync from `.clone()` on ema_input_scratch_d. With cudarc
0.19.3 async alloc enabled (commit 29d21b324), every CudaSlice clone
becomes a `cuMemAllocAsync` graph node captured in the mega-graph.
Same pattern as the kl_pi inline at line 4762 (commented "avoids
.clone()" but never applied to the q_divergence sibling 60 lines
later).
2. Bump cli.diag_every default 1 → 10. Each diag sync_and_swap event-
syncs ~15 DtoD copies on the diag stream (~3-10ms). At 1/step that
caps throughput at ~100-300 sps. The cluster runs at 10 by default
(300+ sps) and can dial to 100 for max perf or 1 for full diag
fidelity.
Local RTX 3050 Ti perf baseline (b=128, per_capacity=8192, n_steps=500):
sustained 8 sps with mega-graph engaged. GPU is the bottleneck on this
hardware (133ms/step measured via cuGraphLaunch host time post-queue-
fill — RTX 3050 Ti SM count + memory bandwidth, not host overhead).
Expect L40S/H100 to hit the original commit-message claim of 300+ sps
unchanged since this commit only removes host-side artifacts.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Two fixes in the integrated trainer hot path.
1. l_pi / l_q monotonic accumulation (root cause: 90f178ae9 perf removal).
The PPO surrogate fwd kernel atomicAdds into ss_pi_loss /
ss_pi_loss_entropy, and dqn_distributional_q backward atomicAdds into
ss_q_loss. Per the original `BUG FIX` comment in dqn_replay_step,
these MUST be zeroed each step or they read as the running sum across
all steps.
90f178ae9 (remove K-loop memsets) correctly removed the per-K-iter
zeros for buffers that are OVERWRITTEN by backward kernels and
reduce_axis0, but the comment in the dead code branch flagged the
loss-scalar atomics as a known exception. The exception was lost
when the if-false block ceased to execute.
Smoke before fix at step 499: l_q=10.96 l_pi=1384.46
Smoke after fix at step 499: l_q=0.024 l_pi=6.65 (both bounded).
v_loss writes via reduce_axis0_mapped (single-writer overwrite) and
does not need zeroing.
2. LobPtrs dead-code refactor (root cause: 9c1b70ede merge resolution).
The 9c1b70ede mega-graph commit introduced LobPtrs to cache
pos_d().raw_ptr() / bid_px_d() / ask_px_d() / pos_bytes() once per
step instead of dispatching through the RlLobBackend vtable at every
raw_launch call site. The struct merged in but the 35 consumer sites
did not because of structural divergence between phase-a's
integrated.rs and clean-clamp's. Per `feedback_no_hiding`, wire it
up or delete — wired across step_with_lobsim,
step_with_lobsim_reward_and_train, and step_with_lobsim_gpu_body.
Two sites kept the old `lobsim.pos_book_and_market_targets_mut()`
pattern because that method returns disjoint borrows including a
`&mut CudaSlice<i32>` for market_targets — incompatible with the
immutable lob snapshot.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
CLI flags: --resume-from <path> and --checkpoint-every <n> (default 5000).
Resume loads checkpoint and continues from saved step. Save writes
rolling checkpoints (keeps last 2, deletes older).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Binary format: magic + version + step + sections. Sections include
ISV (585 f32), 26 device buffers (DQN/policy/value/IQN/FRD/NoisyNet/
outcome weights + targets), 20 AdamW optimizer states, and encoder
(delegated to CfcTrunk). All device transfers use mapped-pinned DtoD.
load_checkpoint invalidates CUDA graphs to force re-capture.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Serializes m, v moments + step_count + hyperparams to a binary writer.
Uses mapped-pinned DtoD (not raw memcpy_htod/dtoh) per
feedback_no_htod_htoh_only_mapped_pinned.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
ml-backtesting/build.rs reads FOXHUNT_CUDA_ARCH (defaults sm_86),
ml-alpha/build.rs reads CUDA_COMPUTE_CAP. Template only set the
latter, so lobsim cubins compiled for sm_86 on H100 (sm_90) →
CUDA_ERROR_NO_BINARY_FOR_GPU at runtime.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The graph was captured at step 2 with PER nearly empty (2048/65536
entries). PER fills at step 64 (65536/1024). The captured graph baked
empty-PER kernel behavior — once PER filled, sps dropped from 338→7.
Now: warmup runs eagerly for per_capacity/b_size + 2 steps (66 steps).
Capture happens at step 67 when PER is full and all kernel behaviors
are in steady state.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
At mega-graph speeds (329 sps), the diag sync_and_swap blocks for
100ms+ because the DtoD copies haven't finished in the 3ms step time.
Gate sync+snapshot to every 10th step (or log/checkpoint boundaries).
The DiagFrame still sends every step using stale staging data — the
background writer drops most frames anyway via try_send(1).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The per-step while loop computing (grid, start, nodes_at_level) for
each tree level was host-side work inside the mega-graph replay path.
Now precomputed at init: 16 entries for capacity=65536, 192 bytes.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The mega-graph ran at 338 sps for 60 steps then degraded to 7 sps.
Root cause: PER fills at step 64 (65536/1024), exposing three bugs:
1. rl_per_tree_rebuild: used __threadfence() as a barrier (it's only
a memory fence). Race condition benign when tree mostly zeros,
corrupts sum-tree when full. Fix: multi-kernel (one per tree level),
stream ordering provides the barrier. Graph-compatible.
2. rl_per_push_flush: volatile spin with 131k threads saturating SMs.
Fix: split into prefix_sum (Grid=1) + coalesced write (Grid=b_size,
Block=128). No volatile spin.
3. rl_per_sample: Block=(1) doing 256 sequential random reads per
sample. Fix: Block=(128) with coalesced cooperative gathers. 128×
less memory traffic.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Captures the ENTIRE per-step training pipeline into ONE cuGraphLaunch,
eliminating ~150 individual kernel launches and ~143ms of host-side
Rust overhead per step.
Phase 1: Lobsim → raw_launch
- apply_snapshot_from_device inlined as 4× raw_memcpy_dtod + raw_launch
- step_fill_from_market_targets inlined as 2× raw_launch
- LobSimRawPtrs trait method caches 30+ device pointers
Phase 2: LR controller → GPU kernel
- New rl_lr_from_mapped_pinned.cu reads losses from device pointers
- New adamw_step_isv_lr kernel reads LR from ISV instead of scalar arg
- All 12 Adam .step() calls use step_isv_lr() in mega-graph mode
Phase 3: PER → main stream
- mega_graph_single_stream flag routes PER to self.raw_stream
- Cross-stream events skipped, K forced to 1
Phase 4: Mega-graph capture/replay
- Three-state machine: warmup → capture → replay
- enable_mega_graph() propagates to perception trainer
- Perception sub-graph guards prevent sub-captures inside mega-capture
- grad_h_accumulate_scaled_isv reads lambda from ISV on-device
Validated: 500 steps, 68-76 sps sustained, no NaN, all losses finite.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
nsys profiling identified two hot-path bottlenecks:
1. sync_training_event() — cuEventSynchronize blocked host 5-11ms
EVERY step (87.9% of CUDA API wall time). Removed: LR controller
uses EMA smoothing, 2-step deferred mapped-pinned reads are fine.
2. 24 raw_memset_d8_zero calls per K-loop iteration — zeroed scratch
buffers that backward kernels overwrite completely. At K_max=4:
72 memsets/step, 26% of CUDA API time. Disabled via if-false gate.
Expected combined impact: ~2s saved per 200 steps → measurable sps
improvement at b=1024.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The 94KB per-step JSON record (130 fields, per-batch unit arrays)
took ~150ms CPU to build+serialize, blocking the GPU pipeline that
only needs ~13ms. Training was CPU-bound at 6.2 sps.
New architecture: training loop snapshots DiagFrame data (~0.1ms
memcpy from mapped-pinned to owned Vecs), sends via non-blocking
try_send on sync_channel(1). Background writer thread builds JSON
and writes to BufWriter at its own pace. Drops frames under
backpressure — acceptable for diagnostics.
Training loop per-step: 13ms GPU + 0.1ms snapshot = ~13.1ms.
Expected: ~50-77 sps at b=1024 on L40S (was 6.2).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
nsys profile showed gpu_gather_bce_labels consumed 95.6% of GPU time:
5 separate kernel launches per step, each random-accessing a 61M-element
array. The snapshot gather kernel (gpu_sample_and_gather) already
computes the same global_idx — fusing the label reads into the same
pass eliminates 5 launches with zero extra random access cost.
Before: 6 random-access kernels (1 snapshot + 5 labels) = 13ms/step
After: 1 random-access kernel (snapshot + labels fused) = ~2.5ms/step
Expected speedup: 6 sps → 30-50+ sps at b=1024.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
cudarc's CudaSlice::Drop checked has_async_alloc (always false) and
fell back to stream.synchronize() + free_sync() on every drop. This
blocked the host for ~12.6ms per sync × 4.7 drops/step = ~59ms/step
of pure idle waiting (35.8% of wall time per nsys).
With has_async_alloc=true, drops use cuMemFreeAsync (non-blocking,
same-stream ordering). CUDA 12.4 on sm_86+ fully supports this.
Combined with the background diag writer and fused label gather,
the training loop should now be GPU-bound at ~13ms/step → ~77 sps.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
DQN head (dqn.rs) and Mamba2 encoder (mamba2_block.rs) were forcing
CUBLAS_COMPUTE_32F — pure FP32 scalar, bypassing Tensor Cores entirely.
TF32 (19-bit mantissa) gives 2-3× SGEMM throughput on L40S/H100 with
negligible precision loss (well within RL gradient noise).
8+ SGEMMs per step now use Tensor Cores: DQN forward/backward (4) +
Mamba2 W_in/W_a/W_b/W_out projections (4+).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Previous Phase 2 fix only blocked explicit close actions (FlatL/FlatS/
HalfFlatL/HalfFlatS). But position reversal via direct directional
actions (ShortHard while long, LongHard while short) can also reduce
|position| by transitioning through zero, firing done events that
bypass the min_hold gate.
Local smoke comparison @ step 499:
Without reversal block: avg_hold=15.7, trades=3660
With reversal block: avg_hold=36.8, trades=1474
The block doubles avg_hold and halves trade frequency, consistent with
enforcing wave-scale operation. Per pearl_edge_lives_at_wave_timescale_not_tick.
Logic by position sign:
- pos > 0 (long): block actions {0,1,3,9} — short reversal + close + half-close
- pos < 0 (short): block actions {4,5,6,10} — close + long reversal + half-close
- pos == 0 (flat): no blocking — opening allowed
Heat-cap exemption preserved BEFORE the block — margin protection
overrides patience (safety wins).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Phase 2: γ-derived min-hold bootstrap
- rl_min_hold_check.cu rewritten to compute hold = current_step -
oldest_active_unit.entry_step, NOT steps_since_done (which counts
time since last close, including flat periods — wrong signal).
- Bootstrap value changed from hardcoded 100.0 → round(-ln(2)/-ln(γ))
= 138 at γ=0.995. Wave-scale operation per
pearl_edge_lives_at_wave_timescale_not_tick.
- Two launch sites updated to push unit_active_d + unit_entry_step_d
(matches new kernel signature).
Phase 3: Realistic published ES futures transaction costs
- cost_per_lot_per_side = $0.82 (CME $0.31 + NFA $0.02 + AMP $0.49)
- Set in alpha_rl_train.rs after LobSimCuda::new
- Half-spread cost is already modeled by lobsim's bid-ask matching.
Phase 4 audit (no code changes): Mbp10Raw struct has no forward-looking
fields. FRD/BCE labels only flow to aux supervised heads, not encoder
state. No data pipeline leakage detected.
Local smoke (RTX 3050 Ti, b=128, 500 steps):
Build + run clean. PnL drops from $560k → $64k under realistic costs
(8.8× reduction). avg_hold metric reports 15.7 but this is "avg
time between consecutive close events" (broken metric — includes
flat time between trades), not the actual position hold time.
KNOWN LIMITATION: min_hold blocks explicit FlatL/FlatS but reversal
actions (e.g., ShortHard while long) can also close positions via
transition through zero in the lobsim. The block doesn't catch these.
Phase 2.1 follow-up may add reversal-direction blocking when hold <
min_hold. Cluster validation will show if this materially affects
hold dynamics under full data.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Yesterday's fix `trade_context[1] = fminf(0.0f, unrealized_R)` created a
structural 12× W/L bias: the agent saw losses fast (closed them), couldn't
see wins (rode blindly until max_hold/session). The reproducibility across
folds (f0=f1=f2 at step 500) was the SAME exploit pattern in each window,
not real edge.
This replaces the asymmetric clip with drawdown-from-peak:
trade_context[1] = unrealized_R - peak(unrealized_R) (always ≤ 0)
The peak tracks per-active-unit and resets on slot activation. Symmetric:
fires on losing positions (monotone drawdown from entry) AND on
winning-then-retracing positions (drawdown from high-water mark even
while still net-positive).
Sentinel-zero bootstrap per pearl_first_observation_bootstrap: peak==0
on freshly opened/added/reversed slot; first observation directly
replaces peak with unrealized_R.
Implementation:
- New CudaSlice unit_peak_unrealized_r_d [B×MAX_UNITS], allocated zero
- rl_unit_state_update + rl_fused_reward_pipeline write 0.0f sentinel
on each OPEN/REVERSE/PYRAMID-ADD activation
- rl_trade_context_update reads/updates peak inline, outputs drawdown
Local smoke (RTX 3050 Ti, b=128, 500 steps):
- avg_hold rose from 17.6 → 26-32 steps (wave-scale-ier)
- W/L magnitude ratio is no longer pinned at 12× — varies 0.67-27 across
the run, sometimes L>W (genuinely symmetric signal)
- No crash, training trajectory healthy
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Two architectural fixes addressing the "never close" overfit attractor
that caused walk-forward fold 0 to converge to dones=0 by step 2000.
1. Asymmetric clip on unrealized_R in trade_context (value leak fix):
trade_context[1] = fminf(0.0f, unrealized_R)
The model previously saw unrealized_R as a state feature. Q learned
"high unrealized = high V(state)" → Q(close) < V(hold) → never close.
Now Q only sees the LOSS side: losing positions visible (cut losses),
winning positions invisible (close decision driven by general policy,
not state-conditioned on profit). Realized PnL on close still
teaches "take profits" via Bellman backup.
2. Enable max_hold_ns = 60s (environment constraint):
LobSimCuda::max_hold_ns_d was alloc_zeros (disabled). The lobsim
has session-gap force-close (>1 hour gaps) but no per-trade timeout,
letting the agent hold positions indefinitely within sessions.
Now alpha_rl_train sets max_hold to 60s via new upload_max_hold_ns
method backed by fill_u64.cu kernel (device-side, complies with
feedback_no_htod_htoh_only_mapped_pinned).
These are environment/architectural fixes, NOT reward shaping hacks.
Together they address why the model COULD learn "never close" and
ensure dones signal density for Q-learning.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Standard PPO practice missing in this codebase. Without it, lifting
the WIN clamp from 1.0 to 167× (yesterday's fix) expanded |A| by
167× → gradients 167× → KL 167× → ε saturated MAX → unbounded
policy drift → overfit (alpha-rl-wf-f0 step 2000: l_pi=105k, dones=0).
Implementation:
- New kernel compute_advantage_rms.cu: single-block warp-shuffle
reduce computes sqrt(mean(A²)) over the batch each step
- New ISV slot 589 RL_ADVANTAGE_BATCH_RMS_INDEX
- ppo_clipped_surrogate fwd/bwd divide A by RMS (with floor bootstrap)
- Dispatched after compute_advantage_return in step_with_lobsim
Scale-only normalization (RMS, not z-score) preserves A=0 on
non-done steps → no spurious PPO gradient on Hold actions where
no trade closed.
This is architectural correctness, not a quickfix. Fights the
ROOT CAUSE (reward range expansion breaking LR calibration), not
the SYMPTOM (ε saturation, l_pi explosion).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Diag analysis of alpha-rl-pifix-4h5fd (500 steps, 80MB JSONL) revealed:
clip_rate_ema = 99.99% across the entire run
pos_scaled_max_ema = 30-116 (30× the WIN=1.0 clamp)
MARGIN saturated at MAX=5.0 (controller wants WIN wider)
reward_clamp_win STUCK at 1.0 for 500 steps
Root cause: rl_reward_clamp_controller.cu had `(void) margin;` at
Step 4 — the MARGIN was computed but discarded. WIN was hardcoded
to the 1.0 bootstrap from an old G.2 feedback-loop fix. With 99.99%
of rewards clipped to ±1, Q saw saturated signal and couldn't
distinguish small from huge wins.
Fix: write WIN = MARGIN × pos_max_ema (the controller's intent),
and LOSS = RATIO × WIN (preserves loss-aversion asymmetry). The
G.2 feedback concern is addressed by the slow atom-span EWMA
(α=0.001, half-life 700 steps) at Step 5 — atom span damps the
WIN→atom→Q loop sufficiently.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
PPO surrogate uses atomicAdd to accumulate into these buffers, but
they were NEVER reset between steps. After 10k steps, l_pi reported
1.27e8 — the SUM of all per-step losses, not the per-step mean.
This explains the apparent l_pi explosion in alpha-rl-floor-zwlk4
(step 9999: l_pi=127580728). The actual per-step l_pi is ~12k/step
(typical PPO surrogate magnitude with the current ratio clamp).
ss_q_loss was correctly zeroed; ss_v_loss uses reduce_axis0_mapped
which overwrites. Only the PPO accumulators leaked.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Two fixes that decouple entropy stability from loss aversion:
1. Thompson sampling probability floor (RL_THOMPSON_FLOOR_INDEX=588,
bootstrap 0.05): 5% of steps pick uniform random action. Prevents
any action from reaching π=0 (δ-function attractor). ISV-driven.
Per pearl_pi_actor_collapses_without_entropy_floor.
2. LOSS bootstrap restored to 3.0 (was 1.5). LOSS=1.5 caused entropy
collapse to 0.94 despite max SAC. LOSS=3.0 keeps entropy stable at
1.87 (proven over 58k steps). The done-gated EMAs (slots 585/586)
will adapt LOSS toward ~1.45 AFTER entropy stabilizes.
Together: entropy stays stable (LOSS=3.0 + 5% floor) AND PnL improves
(adaptive LOSS lowers toward real L/W ratio after warmup).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Added unrealized_pnl field to PosFlat (offset 24, 24→28 bytes).
Computed every step in both fill kernels (order_match + resting_orders)
as (mid - vwap_entry) × direction × |lots|. Zeroed when flat.
This gives Q real-time visibility into open-position drawdown through
the trade_context features. The confidence gate's margin floor now
checks total equity (realized + unrealized) instead of just realized.
Units: price-units × lots (same as realized_pnl). USD = value × $50.
14 files updated: struct def (Rust + CUDA), both fill kernels,
margin floor consumer, layout docs in 6 kernels, tests.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Broker hard constraint: $35k capital, ES maintenance margin ~$14.4k.
If cumulative realized_pnl drops below -$20k, ALL trading actions
are overridden to Hold (if flat) or FlatL/FlatS (if positioned).
ISV-driven at slot 587 (RL_MARGIN_FLOOR_USD_INDEX, bootstrap 20000).
Fires FIRST in the confidence gate — before warmup check, before
exploration slots, before everything. The broker will liquidate at
margin call; this prevents the model from reaching that point.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
LOSS=3.0 bootstrap means Q tolerates 3× losses during warmup before
the done-gated EMAs adapt. This accumulates -$2M before the controller
kicks in. Bootstrap 1.5 = moderate loss aversion from step 0.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The old approach used pre-clamp reward tail maxes (neg_ema/pos_ema ≈
0.07) → LOSS hit floor at 1.0 → no loss aversion → wr dropped to 0.44.
New approach: the clamp controller loops over done-step rewards,
separates wins from losses, maintains two ISV EMAs:
- RL_DONE_WIN_MAGNITUDE_EMA_INDEX (585) — avg winning trade magnitude
- RL_DONE_LOSS_MAGNITUDE_EMA_INDEX (586) — avg losing trade magnitude
LOSS = clamp(1.0, loss_ema/win_ema × 1.1, 3.0). With observed L/W
ratio of 1.32, LOSS should settle at ~1.45 — preserves loss aversion
while preventing the 3× tolerance that caused massive losses.
On top of dd049d9a4 baseline (wr=0.567). Only this one change.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
On top of dd049d9a4 baseline (wr=0.567). Single change: LOSS clamp
adapts from observed loss/win magnitude ratio instead of static 3.0.
LOSS = clamp(1.0, ratio × 1.1, 3.0). WIN stays structural at 1.0.
Gives Q accurate loss magnitude perception without over-allocating
resolution to rare catastrophic tails. Should improve L/W from 1.32
toward break-even (1.28) while preserving wr≥0.55.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The clamp bounds were bootstrapped at ±0.5 (matching the old ±0.5 atom
span) instead of the intended asymmetric WIN=1.0, LOSS=3.0. This
caused: (1) apply_reward_scale clamped at ±0.5 instead of [-3,+1],
(2) the C51 atom span EWMA target = max(1.0, 0.5) = 1.0 → V_MIN
never moved because target min(-1.0, -0.5) = -1.0 = bootstrap.
Now: WIN=1.0 (positive reward ceiling), LOSS=3.0 (loss-aversion
asymmetry per pearl_audit_unboundedness). Atom span will EWMA from
[-1, +1] toward [-3, +1] over ~2100 steps (α=0.001).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Same issue as apply_reward_scale: the clamp controller and atom
updater were wired in the OLD step path (step_with_lobsim_reward_and_train)
but not in the GPU path (step_with_lobsim_gpu_body). My earlier fix
inlined apply_reward_scale but missed the two kernels that follow it.
Without this, the C51 atom span EWMA never fires — V_MIN/V_MAX stay
frozen at bootstrap [-1.0, +1.0] despite the EWMA code being correct.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
G.2 disabled atom span adaptation to break a positive feedback loop.
With clamp bounds now FROZEN, re-enabling span adaptation is safe.
The span anchors on the CLAMP bounds (WIN=1.0, LOSS=3.0) — the
structural ceiling on post-clamp rewards that Q actually sees. The
earlier version incorrectly used pre-clamp EMAs (pos_ema ≈ 50) which
would blow atoms to [-50, +50] with Δ_z=5.
V_MAX EWMAs from 1.0 toward WIN=1.0 (stays put).
V_MIN EWMAs from -1.0 toward -LOSS=-3.0 (slowly widens to cover the
loss tail). Final asymmetric span [-3, +1] with Δ_z=0.2 gives Q
full resolution across the clamped reward range.
α=0.001 (half-life ~700 steps), floor ±1.0.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
G.2 audit froze the atom span and said "atom span tracks the SEED
clamp range" — but the seed was ±0.5 while the clamp WIN=1.0. With
reward_scale keeping typical scaled rewards at ±0.6-0.9, half the
reward range was projected to edge atoms (binary resolution above
V_MAX=0.5).
Now: V_MIN=-1.0, V_MAX=+1.0 in both Q_V_MIN/Q_V_MAX constants and
ISV bootstrap. Δ_z = 2.0/20 = 0.1, giving ~6 atoms for a typical
scaled reward of 0.6. Q can now distinguish "good trade" from "great
trade" instead of just "positive" vs "negative."
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
launch_apply_reward_scale was defined but NEVER CALLED from the step
body. Raw PnL ($100-1000) went directly into C51 Bellman with atom
span [-0.5, +0.5]. Every reward projected to edge atoms → binary Q
resolution → Thompson sampling couldn't distinguish actions → wr stuck
at 0.36.
The full reward chain: apply_reward_scale → reward_clamp_controller →
atom_support_update was all wired (kernels loaded, methods written)
but the entry point was orphaned. The clamp controller and atom
updater ran but saw zero input (pos_max_ema=0, neg_max_ema=0).
Now: fused_reward_pipeline → apply_reward_scale → (existing) clamp
controller → atom_support_update. Rewards are scaled to fit the C51
atom span, the span adapts to reward magnitudes, and Q gets proper
distributional resolution.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The gate created a self-reinforcing Hold trap: gate forces Hold → Q
learns Hold is best → conf=0 for trading actions → gate blocks
everything → Hold=100%. Even max SAC (α=2.0, τ=50) couldn't break it.
Two fixes:
1. Exploration slots: the first `b_size × (1 - max_hold_frac)` batch
elements are NEVER gated. At max_hold_frac=0.85 with b=1024, 154
slots always use the Thompson-sampled action. This guarantees Q
sees trading outcomes and can learn their value.
2. Symmetric threshold decay: the adaptive threshold had 100× asymmetry
(10% raise vs 0.1% lower). When Hold exceeded target, the threshold
barely decreased. Now both directions use THRESHOLD_ADJUST_RATE=1.1.
New ISV slot: RL_CONF_GATE_MAX_HOLD_FRAC_INDEX=584 (bootstrap 0.85).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Root cause: confidence gate decouples policy from actions. Policy
softmax stays high-entropy (~2.4) while the gate forces Hold, producing
action entropy ~0.7. SAC read policy entropy EMA (slot 420), saw
"above target", and kept LOWERING τ — the exact opposite of what was
needed.
New kernel `action_entropy_per_step` computes H(action_histogram) from
the POST-gate actions buffer and writes to ISV slot 583
(RL_ACTION_ENTROPY_EMA_INDEX). SAC co-tuning in rl_q_pi_distill_grad
now reads this slot. Launched OUTSIDE CUDA Graph capture (after
confidence gate + FRD gate) in both training and prefill paths.
Kernel design: 11 threads (N_ACTIONS), each thread counts its action
across all b_size elements. Thread 0 computes entropy from the
histogram and updates the EMA. No atomics.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Two bugs caused τ to collapse to floor (1.0) instead of ramping up:
1. Single-sample entropy: SAC auto-tune computed s_entropy from block 0
only (one batch element). At b=1024 this is noise. Now reads
RL_ENTROPY_OBSERVED_EMA_INDEX (slot 420) — the batch-average entropy
EMA written by ema_update_per_step via tree reduction.
2. Symmetric step rate: SAC_ALPHA_STEP=0.001 was identical for ramp and
decay. Entropy collapsed faster than the controller could recover.
Now asymmetric: SAC_STEP_RAMP=0.01 (100× faster when entropy is below
target) vs SAC_STEP_DECAY=0.0001. Per pearl_asymmetric_controller_decay.
Also adds sac_alpha + sac_entropy_target to JSONL isv_config for
diagnostic visibility.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
