Floor = max(|mean_Q| * 0.01, 1e-6). At Q-mean=0.005: floor=5e-5.
At Q-mean=1.0: floor=0.01. Fully adaptive, zero hardcoded constants.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
tau=Q_gap/3 gave 20:1 weight ratio (95%/4%/1%), making the target
essentially argmax. Once the online distribution matched this sharp
target, cross-entropy gradient vanished → Q-gap froze at 0.158,
val_Sharpe locked at 2.40 from epoch 22.
tau=Q_gap gives a CONSTANT e:1 ≈ 2.72:1 ratio (63%/23%/14% for 3
actions) regardless of Q-gap magnitude. The target always maintains
meaningful mixture → perpetual gradient pressure → Q-values keep
evolving. Scale-invariant: the softmax ratio depends only on relative
Q-differences, not their absolute magnitude.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
All 4 branches (dir, mag, order, urgency) now use Boltzmann sampling
in eval mode instead of argmax. argmax(softmax(Q/tau)) = argmax(Q) —
temperature is meaningless with argmax. Every state picked the same
winning action → identical trades → Sharpe=0.00 from epoch 8 onward.
Boltzmann sampling with Philox seed gives deterministic but
action-diverse evaluation. When Q-values differentiate, Boltzmann
naturally sharpens toward the best action.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Softmax weights used implicit tau=1, producing near-uniform weights
when Q-gap=0.06: exp(0.06)≈1.06 → weights [0.35, 0.33, 0.32].
The mixture averaged all distributions → zero gradient to differentiate
actions → Q-gap frozen → val_Sharpe locked at 0.00 after epoch 19.
Now uses tau = max(Q_gap/3, 0.001). With Q-gap=0.06 and tau=0.02:
exp(0.06/0.02) = exp(3) ≈ 20:1 weight ratio between best and worst.
The 0.001 floor means even tiny Q-differences produce sharp weights,
preventing the uniform-averaging fixed point.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
read_total_loss() calls cuStreamSynchronize which is illegal during
CUDA Graph stream capture. Was inside submit_aux_ops (captured in
graph_mega/graph_aux). Moved to Step 4 (conditional ops outside graph).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
branching_action_select exposure head used hard argmax while order/
urgency used Boltzmann. experience_action_select eval mode used pure
argmax for direction. Both replaced with Boltzmann softmax:
tau = max(Q_max - Q_min, 0.01)
When Q-values differentiate: converges to argmax. When flat: spreads
evenly. Eval mode takes argmax of softmax probs (deterministic).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Two fixes:
1. Width: half = max(10*q_gap, 3*q_std, 0.1) instead of fixed 3*q_std+1.0.
With q_gap=0.06 and 51 atoms, gives ~5 atoms of action resolution
(was 1.5 atoms with the 1.0 floor). No more catastrophic -40 dips
from atom underresolution.
2. Baseline: uses q_std_ema (proportional to natural Q oscillation)
instead of fixed 0.01. Normal Q-oscillation (±q_std) no longer
triggers aggressive alpha — only abnormal shifts do.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Readiness jumped from 0 to ~0.8 in one step when CV first improved,
causing per-sample support to snap from [-1,1] to V(s)-centered
discontinuously. Same adaptive alpha pattern as eval_v_range:
α = |error| / (|error| + 0.05), clamped [0.01, 0.3]. Smooth ramp-in.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
EMA init used == 0.0 sentinel and snapped to first non-zero Q-stats,
causing a discontinuous v_range transition → val_Sharpe=30 spike at
epoch 3 followed by -1 correction. Now seeds from neutral center
(mean=0, std=1 → v_range=[-4,4]) and the adaptive alpha tracks
to actual Q-stats smoothly over 2-3 steps.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
When Q-mean shifted from +0.01 to -0.01, fixed α=0.05 took ~20 updates
to catch up → misaligned backtest atoms → val_Sharpe dipped to -20.
Adaptive alpha: α = |error| / (|error| + 0.01), clamped [0.01, 0.5].
Tracks fast on Q-shifts (α→0.5), smooths when stable (α→0.01).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Fixed v_range from hyperparams (reward_scale/gamma=±240) was 4800x too
wide for Q-values near 0. All Q-values mapped to the same 2-3 center
atoms, making backtest action selection essentially random → val_Sharpe
jumped between -11 and +0.14 arbitrarily. Now uses the same
EMA-smoothed eval_v_range that the experience collector uses.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
reduce_current_q_stats returned PREVIOUS call's results via async
double-buffer. When the eval_v_range EMA converged, stale readback
created a limit cycle: Q-mean alternated between exactly -0.0190 and
-0.2102 for 20+ epochs. Now uses synchronous cuStreamSynchronize +
cuMemcpyDtoH (5µs cost every 50 steps). Removes dead
flush_q_stats_readback and q_stats_ready field.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Hard argmax picks actions based on noise when Q-values are flat
(all near zero). This causes policy collapse — the model locks into
a single action pattern after 1 epoch (Trades=56647 locked, action
distribution frozen). Expected SARSA-style softmax mixture:
p_target(z) = Σ_a softmax(E[Q(s',a)])[a] * p_target_a(z)
When Q-values differentiate, converges to argmax. When flat, averages
across actions → stable gradient signal → breaks policy collapse.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
update_eval_v_range snapped to current Q-stats every 50 steps. With
per-sample IQL support (the circular dependency fix), Q-stats fluctuate
between steps, causing the backtest atom placement to jump → val_Sharpe
oscillated between -42 and 0. EMA (α=0.05) smooths the transition,
giving the backtest consistent atoms across evaluations.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
compute_expected_q used eval_v_range to compute Q-values, then Q-stats
updated eval_v_range from those Q-values — circular dependency that
amplified any drift. Q-mean collapsed 0 → -0.69 in 10 epochs despite
mean_reward ≈ 0. Now uses per_sample_support from IQL (independently
trained V(s)), breaking the feedback loop.
Also removes 127 bf16 identity wrapper calls across 5 CUDA kernels
(c51_loss, epsilon_greedy, dqn_utility, experience, attention_backward).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
qvalue_stats_reduce and qvalue_divergence_check used bf16() identity
wrappers on native f32 data — legacy from the bf16 era. These
wrappers can cause subtle precision differences across GPU
architectures. Replaced with native fminf/fmaxf/__shfl_xor_sync.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Forward cached GEMMs (plain + RELU_BIAS) used AlgoGetIds which can
select graph-incompatible algorithms on H100 SM90. The DDQN target
forward pass uses these — biased target Q-values cause bootstrapping
collapse (Q-mean drifted from 0 to -0.69 over 10 epochs despite
mean_reward≈0). Heuristic selects graph-safe algorithms.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Consecutive counter reset on a single epoch above threshold, so the
trigger=5 was never reached when max_dd naturally bounced every 3-4
epochs. Now uses EMA of binary low-dd signal (α=0.15) with hysteresis:
activates at ratio>0.6 (sustained low dd), deactivates at ratio<0.4
(recovered). No fixed trigger count, no single-epoch reset.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Fixed threshold 0.5% MaxDD was calibrated for zero-gradient era.
With real gradients, max_dd=0.01-0.08% is always below 0.5%, so
adversarial activated at epoch 5 and never released — every epoch
trained under 3x spread, 2x tx_cost, 0.5x fill permanently.
Now uses EMA(max_dd) * 0.5 as threshold, self-calibrating to the
model's actual drawdown scale.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
domain_rand_episode_starts used constant seed (i, 0, 9999) producing
identical episode starting bars every epoch. With deterministic starts
and alternating weight updates, the OOS Sharpe oscillated perfectly
between +0.85 and -1.0. Adding epoch to the Philox hash key diversifies
starting positions across epochs.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Root causes are fixed (cuBLASLt heuristic, entropy inv_batch, overflow,
IQL OOB). Removes: run_buffer_diagnostics, run_recapture_diagnostics,
debug_buffer_norm_f32, diag_recapture_remaining field,
gradient_stage_diagnostics field + FOXHUNT_GRAD_DIAG env var.
Eliminates ~4 stream syncs per step for first 3 steps each fold.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The ungraphed path served its purpose — confirmed cuBLASLt AlgoGetIds
was selecting graph-incompatible algorithms on H100. Now that the
heuristic fix is in place, the diagnostic is dead code.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
MSE gradient kernel had hardcoded branch_scale=(d==1)?4.0:1.0 for
magnitude branch. Now uses the same IQL branch_scales [B,4] buffer
as c51_grad_kernel. Also fixes MSE entropy missing inv_batch.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Per-branch entropy boost in mse_grad_kernel added O(1) term without
1/B normalization. Same pattern as the C51 entropy fix.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The entropy regularization term in c51_grad_kernel was added without
the 1/B factor, making it O(1) per sample vs O(1/B) for the C51 loss
term. At batch_size=16384, entropy dominated by 16384x, causing
grad_norm=640 (should be ~33). Local batch=16: 5.5 → 0.29.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
d_logits clipping uses val_blocks(2048) + adv_blocks(2304) = 4352
partials, but buffer was sized for total_params(1298). 3054-slot
(12KB) overflow corrupted adjacent GPU memory on EVERY training step
inside graph_forward. This was captured in the CUDA Graph and
replayed every step — the actual root cause of zero grad_buf on H100.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
iql_gather_q_taken and iql_compute_advantage_weights indexed q_out
with factored action (0..80) into a [B,12] buffer — massive OOB.
compute-sanitizer found 2586 errors. Now decodes factored action
(dir*b1*b2*b3 + mag*b2*b3 + ord*b3 + urg) into 4 branch indices
and sums per-branch Q-values. Also fixes total_actions config
(was product 81, now sum 12). Sanitizer: 0 errors.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
AlgoGetIds picks algorithms by ID order. On H100 (SM90), the first
valid algo uses split-K/stream-K with internal workspace allocation
that silently produces zero output when replayed via CUDA Graph.
The heuristic selects graph-safe algorithms by design.
Also removes FOXHUNT_NO_GRAPH temp env var and raw DtoH diagnostic probe.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Temporarily sets FOXHUNT_NO_GRAPH=1 in train-best pods to run
forward+backward+Adam ungraphed. If grad_buf is non-zero ungraphed,
the root cause is CUDA Graph capture of cuBLASLt on H100 SM90.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Runs forward+backward+Adam ungraphed to isolate whether H100
grad_buf=0 is caused by graph capture or the cuBLAS operation itself.
NVIDIA docs confirm cublasLtMatmul inside graphs can silently fail
on Hopper when algorithm uses internal workspace allocation.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
debug_buffer_norm_f32 reuses grad_norm_partials (sized for total_params)
as scratch. d_adv_logits with batch_size*tba=589K elements needs 2304
blocks but partials only has 1137 — GPU memory overflow corrupting
adjacent allocations. On H100 this likely trashed grad_norm_buf,
explaining the persistent grad_norm=0.000000. Guard added to skip
oversized norms.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
All code now uses ptrs.grad_buf (u64) as single source of truth.
Vaccine rewritten to use launch_cublas_backward_to(scratch) instead
of std::mem::swap on CudaSlice which caused pointer divergence
between diagnostics and Adam/graphs. Locally verified: avg_grad=6.13.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
volatile float* bypasses GPU L2 cache for correct read-modify-write
across kernel launches. CPU reads directly via read_volatile at epoch
boundary. Removes dead ptr variable in update_eval_v_range.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
attention d_params_per_sample was B*P*4 = 1.6GB at batch_size=16384.
Caused CUDA_ERROR_OUT_OF_MEMORY on H100.
Same tile fix as IQL: process min(B,256) samples per tile, accumulate
d_params via +=. Buffer shrinks from 1.6GB to 26MB.
Also fixed: attn_weight_grad_reduce uses += for tile accumulation.
Also fixed: IQL VRAM log now shows tiled allocation size.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
At batch_size=16384, per-sample gradient buffer was B*P*4 = 1.7GB
per IQL trainer (3.4GB for dual tau). Caused OOM on H100 for attention.
Fix: tile backward+reduce into chunks of 256 samples. Same kernels,
launched multiple times with offset pointers. Per-sample buffer
shrinks to min(B,256)*P*4 = 27MB. 64x smaller.
Also fixed: weight_grad_reduce uses += (accumulate across tiles)
instead of = (overwrite). Backward kernel takes full_batch_size
param for correct 1/N gradient scaling.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Root cause: eval compute_expected_q used [-100,100] atom range. With 51
atoms that's delta_z=4.0 — all Q-values collapsed to ~0. Model weights
changed but eval actions didn't → identical val_Sharpe=-9.16 every epoch.
Fix: eval_v_range_buf initialized from config v_min/v_max, then updated
per-epoch from observed Q-stats (q_mean ± 3σ). Eval atoms now track
the actual Q-value distribution.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
GPU read-modify-write (acc_buf[0] += loss) requires device memory.
Mapped host memory doesn't support atomic-free += from GPU kernels
on all architectures — H100 showed avg_grad=0.000000.
Reverted to CudaSlice<f32> with DtoH at epoch boundary (gated by
FOXHUNT_GPU_SYNC_DIAG). Removed __threadfence_system from accumulator
writes (not needed for device memory).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Default: enabled (syncs ON for debugging).
Production: set FOXHUNT_GPU_SYNC_DIAG=0 to disable.
Gated:
- Experience collector progress logging sync (3x per collection)
- Training guard read_accumulators epoch-boundary sync
Not gated (required for correctness):
- Drop sync in experience collector (prevents VRAM leak)
- compute_q_stats DtoH (returns data to caller — needs mapped memory refactor)
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
evaluate_dqn_graphed now takes Option<&mut dyn QValueProvider>.
Training eval passes Some(fused_ctx), standalone eval uses closure path.
Removed dead evaluate_dqn non-graphed branch from evaluate_baseline.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Every constant now derives from data:
- cv_max decays at 0.999/step (1000-step half-life) — stale spikes
don't permanently anchor readiness
- p5 estimator inits from first batch's mean spread (not hardcoded 1.0)
- Frugal step uses 1/sqrt(step_count) — convergence guarantee
- Support floor fraction = 1/num_atoms (guarantees ≥1 atom resolution)
- Staleness decay_k = ln(20) — derived from "oldest = 5% weight" invariant
Zero manual tuning. Zero hardcoded constants. All self-calibrating.
50-epoch: fold 2 Sharpe 16.20, fold 3 Sharpe 6.42 at epoch 42.
No regression from making constants adaptive.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Single GPU-resident scalar: CV = sigma_adv / |mean_adv|.
When CV > 1 (noisy advantages), readiness < 1 → features suppressed.
When CV ≤ 1 (stable advantages), readiness = 1 → features fully active.
All 5 IQL downstream kernels blend toward neutral defaults at readiness=0:
- Per-sample C51 support: blends to [-1, 1] default
- PER modulation: blends to td_errors unmodified
- Branch scales: blends to uniform 0.25
- Advantage weights: blends to neutral 1.0
- Expectile gap epsilon: gap scaled by readiness
Zero new hyperparameters. Computed from existing adv_stats_buf.
Best Sharpe improved 11.39 → 14.23 on 10-epoch stability test.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Wired:
- modulate_td_errors: advantage-weighted PER priorities with staleness decay
- per_sample_epsilon: IQL expectile gap drives state-dependent exploration
in experience_action_select (NULL fallback for backtest evaluator)
- replay_write_cursor/capacity accessors for staleness computation
Removed:
- use_iql config flag (IQL is mandatory)
- rng_states from ALL 7 experience kernel signatures + Rust launchers
+ backtest evaluator (pure stateless Philox, zero LCG)
- Dead v_range comments cleaned
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
training_guard accumulator (loss_sum, grad_sum, step_count) now uses
MappedBuffer — GPU writes via device pointer, CPU reads via
read_volatile. Eliminates synchronous memcpy_dtoh at epoch boundary.
Also added __threadfence_system() after accumulator writes for
CPU visibility guarantee.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
