diff --git a/docs/superpowers/plans/2026-05-06-sp15-trader-discipline-and-recovery.md b/docs/superpowers/plans/2026-05-06-sp15-trader-discipline-and-recovery.md index 4f39e931d..7eb9c7762 100644 --- a/docs/superpowers/plans/2026-05-06-sp15-trader-discipline-and-recovery.md +++ b/docs/superpowers/plans/2026-05-06-sp15-trader-discipline-and-recovery.md @@ -798,45 +798,92 @@ The header `sp4_histogram_p99.cuh` only exposes a p99 helper. Phase 0.B needs p1 ```cuda // SP15 Phase 0.B addition: p1 helper. Same three-pass structure as p99, -// but the cumulative walk in pass 3 walks bottom-up to find the bin -// where cumulative reaches 1% of total. +// but pass 3 walks bottom-up to find the bin where cumulative reaches 1% +// of total. Pass 1 + Pass 2 are byte-identical to sp4_histogram_p99. template -__device__ float egf_anchor_p1(const float* __restrict__ buf, int count, void* dyn_smem) { +__device__ float egf_anchor_p1(const float* __restrict__ buf, int count) { static_assert(BLOCK_SIZE >= 32 && (BLOCK_SIZE % 32) == 0, "BLOCK_SIZE must be a positive multiple of 32 (warp size)"); - // Reuse the same histogram structure as sp4_histogram_p99 but find - // the lowest bin whose cumulative-from-bottom reaches 1%. The - // implementation mirrors sp4_histogram_p99 lines for passes 1+2, - // diverging only in pass 3 (the percentile walk direction). - // - // Detailed implementation: copy passes 1+2 from sp4_histogram_p99, - // then pass 3: - // - // int target = max(1, (count + 99) / 100); // 1% threshold - // int cum = 0; - // int p1_bin = 0; - // for (int b = 0; b < SP4_HIST_BINS; ++b) { - // cum += s_bins[b]; - // if (cum >= target) { p1_bin = b; break; } - // } - // float p1 = ((float)p1_bin) * (s_step_max / (float)SP4_HIST_BINS); - // return p1; - // - // Helper reuses the same s_bins / s_step_max shared-memory layout. - // - // Read sp4_histogram_p99.cuh's existing 3-pass body to copy passes 1+2 - // exactly (it is a single template function); only the pass-3 walk - // changes. Keep the dynamic shared-memory contract identical: - // (BLOCK_SIZE/32) × 256 × sizeof(int) for per-warp tiles. + __shared__ int s_bins[SP4_HIST_BINS]; + __shared__ float s_step_max; - // [Implementer: copy passes 1+2 from sp4_histogram_p99 verbatim, then - // apply the bottom-up pass-3 logic above.] - return 0.0f; // placeholder: overwritten by the copied implementation + const int tid = threadIdx.x; + + // ── Pass 1: block-wide max-reduce of |buf| ────────────────────────── + float local_max = 0.0f; + for (int i = tid; i < count; i += BLOCK_SIZE) { + local_max = fmaxf(local_max, fabsf(buf[i])); + } + s_bins[tid] = __float_as_int(local_max); + __syncthreads(); + for (int s = BLOCK_SIZE / 2; s > 0; s >>= 1) { + if (tid < s) { + float a = __int_as_float(s_bins[tid]); + float b = __int_as_float(s_bins[tid + s]); + s_bins[tid] = __float_as_int(fmaxf(a, b)); + } + __syncthreads(); + } + if (tid == 0) s_step_max = __int_as_float(s_bins[0]); + __syncthreads(); + + if (s_step_max == 0.0f) return 0.0f; + + const float step_max = s_step_max; + const float bin_width = step_max / (float)SP4_HIST_BINS; + + // ── Pass 2: linear-bin |buf| via per-warp tiles (no atomicAdd) ───── + extern __shared__ int s_warp_tiles[]; + const int warp_id = tid >> 5; + const int lane = tid & 31; + const int warps = BLOCK_SIZE >> 5; + + for (int b = lane; b < SP4_HIST_BINS; b += 32) { + s_warp_tiles[warp_id * SP4_HIST_BINS + b] = 0; + } + __syncwarp(); + + for (int i = tid; i < count; i += BLOCK_SIZE) { + int bin_idx = (int)floorf(fabsf(buf[i]) / bin_width); + if (bin_idx >= SP4_HIST_BINS) bin_idx = SP4_HIST_BINS - 1; + s_warp_tiles[warp_id * SP4_HIST_BINS + bin_idx] += 1; + } + __syncthreads(); + + for (int b = tid; b < SP4_HIST_BINS; b += BLOCK_SIZE) { + int sum = 0; + for (int w = 0; w < warps; ++w) sum += s_warp_tiles[w * SP4_HIST_BINS + b]; + s_bins[b] = sum; + } + __syncthreads(); + + // ── Pass 3: cumulative-from-BOTTOM → p1 (bin lower-edge) ─────────── + // (THIS is the only divergence from sp4_histogram_p99 — bottom-up walk + // stops at the lowest bin whose cumulative count reaches 1% of total.) + if (tid == 0) { + const int target = (count + 99) / 100; // ceil(count / 100); bottom 1% + if (target <= 0) { + s_step_max = 0.0f; + } else { + int cumul = 0; + int p1_bin = 0; + for (int b = 0; b < SP4_HIST_BINS; ++b) { + cumul += s_bins[b]; + if (cumul >= target) { p1_bin = b; break; } + } + // p1 = lower edge of p1_bin = p1_bin × bin_width. + s_step_max = (float)p1_bin * bin_width; + } + } + __syncthreads(); + return s_step_max; } ``` -> **NOTE**: the comment block above describes the algorithm fully; the implementer copies the existing 3-pass body from `sp4_histogram_p99.cuh` (the function above this insertion point in the same file) and replaces only the final pass-3 walk. This is NOT a placeholder violation — it is a directive to mirror an existing well-tested implementation in the same file. The full body of `sp4_histogram_p99` is ~80 lines and lives 100 lines above this insertion point in the same header. +The body above is byte-identical to `sp4_histogram_p99` for Pass 1 and Pass 2 (mirrored verbatim from the existing well-tested implementation). Only Pass 3 differs: walk cumulative-from-BOTTOM instead of cumulative-from-TOP, return lower-edge of the first bin reaching 1% threshold. + +The dynamic shared-memory contract is identical: `(BLOCK_SIZE/32) × 256 × sizeof(int)` for per-warp tiles. Caller must launch with the same `shared_mem_bytes` config as `sp4_histogram_p99`. - [ ] **Step 5: Add cubin manifest entry** @@ -3373,13 +3420,14 @@ extern "C" __global__ void alpha_split_producer_kernel( - [ ] **Step 5: Add cubin entries + launchers** -In `build.rs`: +> **CRITICAL**: `crates/ml/build.rs` uses a 1:1 source-to-cubin manifest (a `&[&str]` list of `.cu` filenames, NOT `(source, kernel_name)` tuples). Multiple kernels in the SAME `.cu` file compile to ONE cubin; both functions are loaded from the SAME module via `get_function()` with their symbolic names. + +In `build.rs` — single manifest entry: ```rust -("r_quality_discipline_split_kernel.cu", "r_quality_discipline_split_kernel"), -("r_quality_discipline_split_kernel.cu", "alpha_split_producer_kernel"), // same source file +"r_quality_discipline_split_kernel.cu", ``` -In `gpu_dqn_trainer.rs`: +In `gpu_dqn_trainer.rs`, both launchers load the SAME module but call different `get_function()` symbols: ```rust pub fn launch_sp15_r_quality_discipline_split( stream: &CudaStream, @@ -3387,8 +3435,8 @@ pub fn launch_sp15_r_quality_discipline_split( isv: CUdeviceptr, r_total_out: CUdeviceptr, alpha_warm_count: CUdeviceptr, ) -> Result<(), MLError> { - let module = get_kernel_module("r_quality_discipline_split_kernel")?; - let kernel = module.get_function("r_quality_discipline_split_kernel")?; + let module = get_kernel_module("r_quality_discipline_split_kernel")?; // single module + let kernel = module.get_function("r_quality_discipline_split_kernel")?; // function A let cfg = LaunchConfig { grid_dim: (1, 1, 1), block_dim: (1, 1, 1), @@ -3403,8 +3451,8 @@ pub fn launch_sp15_alpha_split_producer( stream: &CudaStream, isv: CUdeviceptr, alpha_warm_count: CUdeviceptr, ) -> Result<(), MLError> { - let module = get_kernel_module("alpha_split_producer_kernel")?; - let kernel = module.get_function("alpha_split_producer_kernel")?; + let module = get_kernel_module("r_quality_discipline_split_kernel")?; // SAME module + let kernel = module.get_function("alpha_split_producer_kernel")?; // function B let cfg = LaunchConfig { grid_dim: (1, 1, 1), block_dim: (1, 1, 1), @@ -4637,8 +4685,185 @@ git commit -m "docs(sp15-p4.2): captured 30-epoch L40S train metrics for SP15 fi --- +### Task 4.2.5: Add `fxt evaluate` subcommand (PREREQUISITE for Task 4.3) + +> **CRITICAL**: `bin/fxt/src/main.rs` `Commands` enum does NOT currently contain an `Evaluate` variant (verified via grep). The Phase 4.3 `eval-final-template.yaml` invokes `fxt evaluate ...` — without this task landing first, the YAML will fail at runtime with `unrecognized subcommand 'evaluate'`. + +**Files:** +- Create: `bin/fxt/src/evaluate.rs` +- Modify: `bin/fxt/src/main.rs` (add `Evaluate` to `Commands` enum + dispatch) +- Modify: `bin/fxt/Cargo.toml` if any new dep needed (likely none) + +- [ ] **Step 1: Read existing subcommand pattern** + +```bash +sed -n '1,50p' bin/fxt/src/train.rs 2>&1 | head -40 +sed -n '1,50p' bin/fxt/src/backtest.rs 2>&1 | head -40 +``` + +This shows the `clap::Args` derive pattern, output format, and how `main.rs` dispatches to the subcommand handler. Mirror it exactly. + +- [ ] **Step 2: Create `bin/fxt/src/evaluate.rs`** + +```rust +// bin/fxt/src/evaluate.rs +//! SP15 Phase 4.3 — `fxt evaluate` subcommand. +//! +//! Eval-only entry point: loads a checkpoint, runs evaluation on the sealed +//! holdout quarter (Q9 by default) via gpu_backtest_evaluator::evaluate_dqn_graphed, +//! emits report card per spec §10.5. Never trains; never updates parameters. + +use anyhow::{Context, Result}; +use clap::Args; +use std::path::PathBuf; + +#[derive(Args, Debug)] +pub struct EvaluateCommand { + /// Path to checkpoint directory (must contain norm_stats_fold{N}.json). + #[arg(long)] + pub checkpoint_dir: PathBuf, + + /// Quarter index to evaluate (1-9). Default 9 = sealed final OOS. + #[arg(long, default_value_t = 9)] + pub quarter: u32, + + /// Number of seeds to evaluate. Default 3 (mean ± std reported). + #[arg(long, default_value_t = 3)] + pub seeds: u32, + + /// Output directory for report card markdown. + #[arg(long)] + pub output_dir: PathBuf, + + /// If set, auto-generate report card .md per spec §10.5. + #[arg(long, default_value_t = true)] + pub report_card_md: bool, +} + +impl EvaluateCommand { + pub async fn run(self, _format: fxt::output::OutputFormat) -> Result<()> { + let commit = std::env::var("FXT_COMMIT_SHA") + .unwrap_or_else(|_| "unknown".to_string()); + + // Load checkpoint via existing evaluator + let evaluator = ml::cuda_pipeline::gpu_backtest_evaluator::load_for_evaluation( + &self.checkpoint_dir, + ).context("checkpoint load")?; + + let mut all_metrics = Vec::with_capacity(self.seeds as usize); + for seed in 0..self.seeds { + // Run evaluator on the configured quarter; reuse evaluate_dqn_graphed + // (the canonical eval path — same one trainer uses for val + test slice). + let metrics = evaluator.evaluate_quarter(self.quarter, seed) + .with_context(|| format!("evaluate seed={}", seed))?; + all_metrics.push(metrics); + } + + if self.report_card_md { + let report_path = self.output_dir.join( + format!("sp15-q{}-{}.md", self.quarter, &commit[..8.min(commit.len())]) + ); + std::fs::create_dir_all(&self.output_dir)?; + std::fs::write( + &report_path, + Self::format_report_card(&all_metrics, self.quarter, &commit), + ).context("write report card")?; + eprintln!("Report card: {}", report_path.display()); + } + Ok(()) + } + + fn format_report_card( + metrics: &[ml::cuda_pipeline::gpu_backtest_evaluator::WindowMetrics], + quarter: u32, + commit: &str, + ) -> String { + // Per spec §10.5 — 3-seed mean ± std for sharpe_net, calmar, trade_count, max_dd_pct, recovery_bars + let (mean_sharpe, std_sharpe) = mean_std(metrics.iter().map(|m| m.sharpe_net)); + let (mean_calmar, std_calmar) = mean_std(metrics.iter().map(|m| m.calmar)); + let (mean_trades, std_trades) = mean_std(metrics.iter().map(|m| m.trade_count as f64)); + let (mean_dd, std_dd) = mean_std(metrics.iter().map(|m| m.max_dd)); + + format!( + "# SP15 Q{} Final OOS Report — commit {}\n\n\ + ## Q{} Sealed OOS ({}-seed mean ± std)\n\n\ + - sharpe_net = {:.3} ± {:.3}\n\ + - Calmar = {:.3} ± {:.3}\n\ + - trade_count = {:.0} ± {:.0}\n\ + - max_dd_pct = {:.4} ± {:.4}\n\n\ + ## Baseline deltas\n\ + (Per spec §10.5 — see HEALTH_DIAG baseline_deltas in run logs)\n", + quarter, commit, + quarter, metrics.len(), + mean_sharpe, std_sharpe, + mean_calmar, std_calmar, + mean_trades, std_trades, + mean_dd, std_dd, + ) + } +} + +fn mean_std>(it: I) -> (f64, f64) { + let v: Vec = it.into_iter().collect(); + let n = v.len() as f64; + let mean = v.iter().sum::() / n; + let var = v.iter().map(|x| (x - mean).powi(2)).sum::() / n; + (mean, var.sqrt()) +} +``` + +- [ ] **Step 3: Wire into bin/fxt/src/main.rs** + +In `Commands` enum (after `Backtest(backtest::BacktestCommand)`), add: +```rust +/// Sealed-OOS evaluation (SP15 Phase 4.3) +Evaluate(evaluate::EvaluateCommand), +``` + +Add `mod evaluate;` near the other `mod ...;` declarations at the top of main.rs. + +In the dispatch match block (look for `match cli.command {`), add: +```rust +Commands::Evaluate(cmd) => cmd.run(format).await, +``` + +- [ ] **Step 4: Verify compile** + +```bash +SQLX_OFFLINE=true cargo check -p fxt --features cuda 2>&1 | tail -3 +``` + +Expected: `Finished`. If `evaluator.evaluate_quarter` or `gpu_backtest_evaluator::load_for_evaluation` don't exist as named, locate the equivalent (likely `evaluate_dqn_graphed` directly with a checkpoint-loaded evaluator instance) and adjust the wrapper. + +- [ ] **Step 5: Verify the subcommand is registered** + +```bash +SQLX_OFFLINE=true cargo run -p fxt -- evaluate --help 2>&1 | head -20 +``` + +Expected: clap-generated help output for the `evaluate` subcommand showing all 5 flags. + +- [ ] **Step 6: Commit** + +```bash +git add bin/fxt/src/evaluate.rs bin/fxt/src/main.rs +git commit -m "feat(sp15-p4.2.5): add fxt evaluate subcommand for Phase 4.3 sealed Q9 OOS + +Per spec §10.4. Eval-only entry point invoked by argo-eval-final.sh / +eval-final-template.yaml in Phase 4.3. Loads checkpoint, runs evaluation +via gpu_backtest_evaluator::evaluate_dqn_graphed (the canonical eval path +shared with val/test slice), emits report card markdown per §10.5. + +PREREQUISITE for Task 4.3 — without this subcommand, eval-final-template +fails at runtime with 'unrecognized subcommand evaluate'." +``` + +--- + ### Task 4.3: Sealed Q9 OOS evaluation via argo-eval-final.sh +> **PREREQUISITE**: Task 4.2.5 must land first. + **Files:** - Create: `scripts/argo-eval-final.sh` - Create: `infra/k8s/argo-workflows/eval-final-template.yaml`