From 62897104635e2ab14eb2659f26dc4a1eb67e1313 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 10 May 2026 09:05:46 +0200 Subject: [PATCH] merge(sp20): restore is_win_per_env across struct/launcher/registry after Phase 3 cherry-pick MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Cherry-picking Phase 3 (a8731dda4 + 84929f419) onto Phase 2 fix tip (13ce78385) with --strategy-option=theirs dropped is_win_per_env arg + field + registration the wr_ema fix added. Manual restoration: - experience_kernels.cu: re-added is_win_per_env arg to experience_env_step - gpu_experience_collector.rs: re-added struct field, alloc, ctor entry, launcher arg - state_reset_registry.rs: re-added FoldReset RegistryEntry - training_loop.rs: re-added named-reset dispatch arm - docs/dqn-wire-up-audit.md: documented the merge restoration Test sp20_is_win_per_env_registered_fold_reset passes. Workspace + examples compile clean. wr_ema runtime mystery (fix wired, but production still shows wr_ema=0) unchanged — debug per 13ce78385 scaffold pending. Co-Authored-By: Claude Opus 4.7 (1M context) --- .../src/cuda_pipeline/experience_kernels.cu | 7 +++++ .../cuda_pipeline/gpu_experience_collector.rs | 29 +++++++++++++++++++ .../src/trainers/dqn/state_reset_registry.rs | 8 +++++ .../src/trainers/dqn/trainer/training_loop.rs | 11 +++++++ docs/dqn-wire-up-audit.md | 4 +++ 5 files changed, 59 insertions(+) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index d2150cd16..6c259b37f 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -2205,6 +2205,13 @@ extern "C" __global__ void experience_env_step( float* __restrict__ alpha_per_env, int loss_cap_idx, int alpha_ema_idx, + /* SP20 Phase 2 Task 2.2-fix (2026-05-10) — segment-level win flag. + * `is_win_per_env` ← [N] i32: written at `is_close && segment_ + * complete && segment_hold_time > 0.0f` with + * `(segment_return > 0.0f) ? 1 : 0`. Consumer: + * sp20_aggregate_inputs_kernel reads this in place of the broken + * per-bar `step_ret > 0` predicate. NULL-tolerant for tests. */ + int* __restrict__ is_win_per_env, /* SP20 Phase 3 Task 3.2 (2026-05-10) — Hold opportunity-cost dual * emission (Component 2) per spec §4.2. * diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 712c0b68d..4ed32ea84 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -735,6 +735,14 @@ pub struct GpuExperienceCollector { /// don't leak into the EMA. pub(crate) alpha_per_env: CudaSlice, // [alloc_episodes] + /// SP20 Phase 2 Task 2.2-fix (2026-05-10): per-env segment-level win + /// flag scratch. Companion to `alpha_per_env` — same producer site + /// (`segment_complete` branch in experience_kernels.cu), same fold- + /// reset semantics, same NULL-tolerant consumer (sp20_aggregate_ + /// inputs_kernel reads when non-NULL, falls back to legacy per-bar + /// `step_ret > 0` predicate when NULL for test scaffolds). + pub(crate) is_win_per_env: CudaSlice, // [alloc_episodes] + /// SP20 Phase 3 Task 3.2 (2026-05-10): per-env circular Hold /// opportunity-cost buffer (Component 2) per spec §4.2. /// @@ -1873,6 +1881,18 @@ impl GpuExperienceCollector { "sp20: alloc alpha_per_env ({alloc_episodes} f32): {e}" )))?; + // SP20 Phase 2 Task 2.2-fix (2026-05-10): per-env segment-level + // win flag (1 iff segment_return > 0 at trade close, else 0). + // Sentinel = 0; written by experience_env_step at the same + // segment_complete site as alpha_per_env. Consumed by the + // aggregation kernel via the new is_win_per_env arg replacing + // the broken per-bar step_ret predicate. Sized [alloc_episodes]. + let is_win_per_env = stream + .alloc_zeros::(alloc_episodes) + .map_err(|e| MLError::ModelError(format!( + "sp20: alloc is_win_per_env ({alloc_episodes} i32): {e}" + )))?; + // SP20 Phase 3 Task 3.2 (2026-05-10): per-env Hold opportunity- // cost circular buffer. Layout `[alloc_episodes × // HOLD_BASELINE_BUFFER_SIZE = 30]` f32 row-major. Sentinel 0.0 @@ -2933,6 +2953,7 @@ impl GpuExperienceCollector { episode_starts_buf, label_at_open_per_env, alpha_per_env, + is_win_per_env, hold_baseline_buffer, states_out, actions_out, @@ -6367,6 +6388,14 @@ impl GpuExperienceCollector { use crate::cuda_pipeline::sp14_isv_slots::ALPHA_EMA_INDEX; ALPHA_EMA_INDEX as i32 }) + // SP20 Phase 2 Task 2.2-fix (2026-05-10): segment-level + // win flag i32 buffer. Written at segment_complete with + // `(segment_return > 0.0f) ? 1 : 0`. Replaces the broken + // per-bar `step_ret > 0` predicate in sp20_aggregate_ + // inputs_kernel's wins_count accumulation (the per-bar + // tick at close bars is dominated by tx_cost deduction + // → systematically negative even on winning trades). + .arg(&mut self.is_win_per_env) // SP20 Phase 3 Task 3.2 (2026-05-10): Hold // opportunity-cost dual emission (Component 2). Six // new args at the end of the kernel signature: diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 1cfd22407..870de5c46 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -2132,6 +2132,14 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "GpuExperienceCollector.alpha_per_env [alloc_episodes] f32 device-resident scratch — SP20 Phase 2 Task 2.2 (2026-05-09) per-env trade-close alpha (= R_event - hold_baseline) per spec §4.1. Phase 2 emits `hold_baseline = 0.0f` as a Phase 3.2 forward-reference placeholder; Phase 3.2's Hold opportunity-cost dual emission lands the real `hold_baseline_buffer.sum_over_trade_range()` consumer. Producer: `experience_env_step` `is_close` (segment_complete) branch — writes `alpha = sp20_compute_event_reward(segment_return, label_at_open_per_env[env], ISV[LOSS_CAP_INDEX])`. Consumer: `sp20_aggregate_inputs_kernel` (Path C aggregation) — sums `alpha_per_env[env]` over closed envs and divides by closed-count to populate `EmaInputs.alpha`. FoldReset sentinel 0.0 across all `alloc_episodes` slots — leftover alpha from the previous fold's last trade-close would corrupt the new fold's first ALPHA_EMA observation. The sentinel-0 reads as `alpha == 0` for non-close steps; the aggregation kernel reads only when `trade_close_per_env[env] != 0` so the FoldReset sentinel never reaches the EMA. Reset path: `cudarc::driver::CudaStream::memset_zeros` on the `CudaSlice` (device-resident, no host buffer to fill). Replaces the Phase 1.4 hardcoded `out_inputs->alpha = 0.0f` placeholder atomically per `feedback_no_partial_refactor` (Task 2.2 lands buffer + kernel-arg + aggregation-kernel-update + 2 placeholder pin tests in one commit).", }, + // ── SP20 Phase 2 Task 2.2-fix (2026-05-10): per-env segment- + // level win flag scratch. Companion to alpha_per_env, same + // segment_complete producer site, same FoldReset semantics. + RegistryEntry { + name: "is_win_per_env", + category: ResetCategory::FoldReset, + description: "GpuExperienceCollector.is_win_per_env [alloc_episodes] i32 device-resident scratch — SP20 Phase 2 Task 2.2-fix (2026-05-10) per-env segment-level win flag (1 iff segment_return > 0.0f at trade close, else 0). Producer: `experience_env_step` `is_close && segment_complete && segment_hold_time > 0.0f` branch — writes alongside `alpha_per_env[i] = alpha`. Consumer: `sp20_aggregate_inputs_kernel` reads `is_win_per_env[env]` to compute `wins_count` (replacing the broken per-bar `step_ret > 0` predicate; per-bar step_ret at close bars is dominated by tx_cost deduction → systematically negative even on winning trades, pinning WR_EMA at 0). FoldReset sentinel 0 across all `alloc_episodes` slots — leftover win flag from the previous fold's last trade-close would corrupt the new fold's first WR_EMA observation. The sentinel-0 reads as is_win=0 (not-won) for non-close steps; the aggregation kernel reads only when `trade_close_per_env[env] != 0` so the FoldReset sentinel never reaches the EMA. Reset path: `cudarc::driver::CudaStream::memset_zeros` on the `CudaSlice` (device-resident, no host buffer to fill).", + }, // ── SP20 Phase 3 Task 3.2 (2026-05-10): per-env Hold // opportunity-cost circular buffer. Component 2 producer/ // consumer per spec §4.2. Written at every per-bar diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 8a1483c87..6d3ed6f72 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -9567,6 +9567,17 @@ impl DQNTrainer { )))?; } } + // SP20 Phase 2 Task 2.2-fix (2026-05-10): per-env segment- + // level win flag scratch. Companion to alpha_per_env above. + "is_win_per_env" => { + if let Some(ref mut collector) = self.gpu_experience_collector { + let stream = collector.stream().clone(); + stream.memset_zeros(&mut collector.is_win_per_env) + .map_err(|e| crate::MLError::ModelError(format!( + "is_win_per_env memset_zeros: {e}" + )))?; + } + } // SP20 Phase 3 Task 3.2 (2026-05-10): per-env Hold // opportunity-cost circular buffer (Component 2). Same // device-resident `CudaSlice` reset pattern as diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 31675a62b..4b74074a6 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,10 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +## 2026-05-10 — SP20 Phase 2 fix + Phase 3 merge: restore is_win_per_env + +Cherry-picking Phase 3 commits onto Phase 2 fix tip with `--strategy-option=theirs` dropped the `is_win_per_env` kernel arg + struct field + reset registration that the wr_ema fix branch added (Phase 3 had its own version of the `experience_env_step` signature without that arg). Manual restoration: kernel signature, struct field, alloc, launcher arg, registry entry, reset dispatch arm. Test `sp20_is_win_per_env_registered_fold_reset` passes. + ## 2026-05-10 — SP20 Phase 3 Task 3.4: replay-buffer schema for per-bar aux_conf (atomic) Lands the producer→ring→consumer schema plumbing for per-bar