Wave 3a half of the val-cost-streams refactor (3b host-side wire-up
follows). Atomically migrates the kernel-side contracts; 5 launchers
remain orphan transiently awaiting 3b production callers.
Baseline kernels (1.4):
- 4 baseline_*_kernel signatures gain 'out: float*' parameter writing
per-window [mean, std, raw_sharpe] (matches 1.1.b sharpe_per_bar shape)
- ISV writes to slots 409, 410, 412, 416 removed entirely
(per-window output is correct for WindowMetrics consumption;
ISV-scalar writes were spec scaffolding for a single-fold-aggregate
version that 1.4.b's per-window contract supersedes)
- 4 ISV slot constants removed from sp15_isv_slots.rs
- state_reset_registry: NO entries to remove (verified via grep —
the 4 slots never had registry entries / dispatch arms in the first
place; they were single-fold-aggregate scalars defaulted at every
fold start by the constructor-write that initialises the ISV bus).
Task 4 from the dispatch is a no-op; the
every_fold_and_soft_reset_entry_has_dispatch_arm regression test
continues to pass unchanged.
- 4 oracle tests migrated to output-buffer assertion
- layout_fingerprint_seed string updated (4 retired entries removed,
4 trunk-shared entries retained; layout-break-class change)
New action_decoding_helpers.cuh:
- Extracts factored_action_to_dir_idx + factored_action_to_position
__device__ helpers (the latter is a higher-level position state-
machine helper not previously available)
- Mirrors trade_physics.cuh::decode_direction_4b semantics exactly so
on-policy and counterfactual paths agree on factored-action meaning
- Single source of truth for action→direction→position mapping;
consumers #include the header
New position_history_derivation_kernel.cu (post-loop derivation for
cost_net_sharpe consumer in Wave 3b):
- Reads actions_history_buf, reconstructs per-bar position_history
(-1/0/+1), side_ind (1.0 on position change), rt_ind (1.0 on
transition-to-flat from non-flat) via sequential walk (single
block per window, no atomicAdd per feedback_no_atomicadd)
- New launcher launch_sp15_position_history_derivation in
gpu_dqn_trainer.rs
- New cubin manifest entry in build.rs
- 1 oracle test covering 8-bar Short→Hold→Long→Hold→Flat→Long→Flat→
Short sequence; expected position/side_ind/rt_ind triples match
hand-computed values
Atomic per feedback_no_partial_refactor for the ISV-contract change
(every consumer of slots 409/410/412/416 migrated in this commit; their
consumers were the 4 oracle tests, all migrated). The orphan launcher
transient state for the 5 baselines + derivation kernel is explicitly
the 3a/3b split point — production callers land in 3b's
GpuBacktestEvaluator::new constructor signature change.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
e968f4ded9