plan(sp15): fix 3 critical compile-breakers from third review
Targeted fixes per user direction (option 3 of "fix critical only, accept rest as execution-time gaps"): 1. build.rs cubin manifest is 1:1 source-to-cubin (verified: list of .cu filenames, NOT (source, name) tuples). Fixed Phase 3.1 Step 5 to use single manifest entry; both kernels load from SAME module via get_function() with their symbolic names. 2. egf_anchor_p1 helper in sp4_histogram_p99.cuh: replaced the `return 0.0f` stub with the full ~80-line body. Pass 1 + Pass 2 byte-identical to sp4_histogram_p99 (mirrored verbatim from the existing sibling). Pass 3 walks cumulative-from-bottom for p1 instead of cumulative-from-top for p99. Returns lower-edge of the first bin reaching 1% threshold. 3. Added Task 4.2.5: `fxt evaluate` subcommand (PREREQUISITE for 4.3). Verified that bin/fxt/src/main.rs Commands enum has no Evaluate variant. Without this task, eval-final-template.yaml fails at runtime. Full subcommand shown: bin/fxt/src/evaluate.rs with 5 flags (--checkpoint-dir, --quarter, --seeds, --output-dir, --report-card-md), report card markdown emitter per spec §10.5. 5 IMPORTANT and 2 NIT issues from review remain documented as execution-time friction; subagent-driven-development's spec-compliance reviewer between tasks is expected to catch them per-task. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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 <int BLOCK_SIZE>
|
||||
__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<I: IntoIterator<Item=f64>>(it: I) -> (f64, f64) {
|
||||
let v: Vec<f64> = it.into_iter().collect();
|
||||
let n = v.len() as f64;
|
||||
let mean = v.iter().sum::<f64>() / n;
|
||||
let var = v.iter().map(|x| (x - mean).powi(2)).sum::<f64>() / 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`
|
||||
|
||||
Reference in New Issue
Block a user