From 26ce7ba690932ee7dfa4404e44702c32faf9c941 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Wed, 13 May 2026 21:55:34 +0200 Subject: [PATCH] =?UTF-8?q?feat(sp22-vNext):=20Phase=20A2=20=E2=80=94=20tr?= =?UTF-8?q?ade-outcome=20label=20producer=20kernel?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit First foundation kernel for the H6 vNext trade-outcome aux head. Per-env classification at trade-close events into K=3 outcomes: - 0 = Profit (pnl_vs_target >= 1.0) - 1 = Stop (pnl_vs_stop >= 1.0) - 2 = Timeout (neither threshold hit) Sparse labels — most bars get -1 mask. Priority: Profit > Stop > Timeout. Pure per-env map; no atomicAdd, no reduction. Launch: grid(ceil(N/256)), block(256). Registered in build.rs. Cubin compiles clean. Currently dead code — launcher wireup comes in Phase A3+ commits. See docs/plans/2026-05-14-sp22-h6-vNext-trade-outcome-aux.md for the full vNext architecture. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 10 ++ .../trade_outcome_label_kernel.cu | 129 ++++++++++++++++++ docs/dqn-wire-up-audit.md | 28 ++++ 3 files changed, 167 insertions(+) create mode 100644 crates/ml/src/cuda_pipeline/trade_outcome_label_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index b72cf6c90..d5ed341f1 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -792,6 +792,16 @@ fn main() { // producer chain has fresh labels every rollout step. Same per-thread // O(1) map; same lookahead window semantics; same skip sentinel. "aux_sign_label_per_step_kernel.cu", + // SP22 H6 vNext (2026-05-13): trade-outcome label producer kernel. + // Replaces the per-bar binary direction label with a per-trade-close + // 3-way outcome label {Profit, Stop, Timeout}. Sparse labels + // (only at trade-close events) but aligned with the system's + // actual multi-bar trade decision horizon. See + // `docs/plans/2026-05-14-sp22-h6-vNext-trade-outcome-aux.md` for + // the full vNext architecture spec. This kernel is the FOUNDATION + // (Phase A2) — additional kernels for aux forward / loss / + // backward + 12-weight W atom-shift land in subsequent commits. + "trade_outcome_label_kernel.cu", // SP22 H6 (2026-05-12): per-env p_up extractor that copies // `aux_softmax[env, 1]` → `prev_aux_dir_prob[env]` after each aux // forward. The cache buffer is read by `experience_state_gather` diff --git a/crates/ml/src/cuda_pipeline/trade_outcome_label_kernel.cu b/crates/ml/src/cuda_pipeline/trade_outcome_label_kernel.cu new file mode 100644 index 000000000..64550b042 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/trade_outcome_label_kernel.cu @@ -0,0 +1,129 @@ +// crates/ml/src/cuda_pipeline/trade_outcome_label_kernel.cu +// +// SP22 H6 vNext (2026-05-13) — trade-outcome label producer. +// +// Classifies the trade that just CLOSED at this (env, step) into one of three +// outcomes: Profit, Stop, or Timeout. Labels are sparse — only emitted on +// bars where `trade_close_per_sample[i] != 0`. All other bars get the mask +// sentinel `-1` so the downstream aux-trade-outcome CE loss skips them. +// +// Replaces `aux_sign_label_per_step_kernel.cu`'s per-bar binary direction +// label with a per-trade-close 3-way outcome label that ALIGNS with the +// system's actual decision horizon (multi-bar trade lifecycle). +// +// Outcome classification +// ────────────────────── +// At trade-close, the kernel reads two saved per-env scratch values +// produced by `experience_env_step`'s segment_complete branch: +// - `pnl_vs_target_at_close[i]` = unrealized / (profit_target × equity_at_entry) +// - `pnl_vs_stop_at_close[i]` = -unrealized / (stop_loss × equity_at_entry) +// +// Both are computed in the canonical PLAN_ISV slot semantics (see +// `state_layout.cuh` PLAN_ISV_PNL_VS_TARGET=1, PLAN_ISV_PNL_VS_STOP=2). +// Values ≥ 1.0 mean the threshold was hit. +// +// Classification (priority: Profit > Stop > Timeout): +// - if pnl_vs_target >= 1.0 → 0 (Profit) — trade hit profit_target +// - elif pnl_vs_stop >= 1.0 → 1 (Stop) — trade hit stop_loss +// - else → 2 (Timeout) — neither threshold hit at close +// +// "Timeout" covers two cases: hold_time reached target_bars without hitting +// either threshold, OR voluntary policy exit (Flat/reverse action) before +// either threshold. Both are signal-rich: "this trade closed marginally — +// not a clear win or loss". Distinguishing the two would require additional +// state and likely doesn't add useful signal vs the policy's existing +// hold-time observation in plan_isv[PROGRESS]. +// +// Sparse labels +// ───────────── +// Most bars (~99% in a typical rollout) have no trade close. Those bars get +// `out_labels[i] = -1` (mask). The downstream aux loss kernel +// (`aux_trade_outcome_loss_reduce`) sums only over valid labels (label != -1) +// and divides by valid count, so gradient magnitude is normalized to +// per-trade-close, not per-bar. This concentrates the aux head's learning +// signal where it matters: at the moments the policy is asked to decide. +// +// Discipline +// ────────── +// - Pure map; no atomicAdd, no reduction. Per-(env) parallelism. +// - GPU-only data per `feedback_cpu_is_read_only.md`. Inputs are mapped-pinned +// or device-allocated (see launcher); outputs are device-allocated. +// - No HtoD per `feedback_no_htod_htoh_only_mapped_pinned.md`. +// +// Launch +// ────── +// grid = (ceil(n_episodes / 256)), block = (256). One thread per env. +// +// Inputs +// ────── +// trade_close_per_sample [N*L] — int flag at out_off = i * L + t indicating +// trade-close event at this (env, step). Produced by +// `experience_env_step` per (env, step). For per-step producer use, +// only the current step's slice [i in 0..n_episodes] is read at +// offset `i * L + t`. +// pnl_vs_target_at_close [n_episodes] f32 — saved per-env at close from +// plan_isv[PLAN_ISV_PNL_VS_TARGET=1]. Sentinel 0.0 when no trade closed. +// pnl_vs_stop_at_close [n_episodes] f32 — saved per-env at close from +// plan_isv[PLAN_ISV_PNL_VS_STOP=2]. Sentinel 0.0 when no trade closed. +// +// Outputs +// ─────── +// out_labels [n_episodes] i32 in {-1, 0, 1, 2} — trade outcome label per env. + +#include + +#define TRADE_OUTCOME_PROFIT 0 +#define TRADE_OUTCOME_STOP 1 +#define TRADE_OUTCOME_TIMEOUT 2 +#define TRADE_OUTCOME_MASK -1 + +extern "C" __global__ void trade_outcome_label_kernel( + /* Per-(env, step) trade-close flag. Read at `i * L + t` for current + * rollout step. Non-zero means a trade just closed for this env. */ + const int* __restrict__ trade_close_per_sample, + /* Per-env P&L vs profit-target at close. Computed in env_step's + * segment_complete branch (PLAN_ISV_PNL_VS_TARGET semantics: + * unrealized / (profit_target × equity_at_entry)). Value ≥ 1.0 + * means the profit threshold was hit. Sentinel 0.0 when no trade + * closed this step. */ + const float* __restrict__ pnl_vs_target_at_close, + /* Per-env P&L vs stop-loss at close (PLAN_ISV_PNL_VS_STOP semantics: + * -unrealized / (stop_loss × equity_at_entry)). Value ≥ 1.0 means + * the stop threshold was hit. Sentinel 0.0 when no trade closed. */ + const float* __restrict__ pnl_vs_stop_at_close, + /* Current rollout step (0..timesteps). Used to compute out_off for + * reading trade_close_per_sample. */ + int t, + /* Number of envs in the rollout slice. */ + int n_episodes, + /* Buffer stride: timesteps per env in the per-sample buffer layout. */ + int L, + /* Per-env i32 output. Written -1 when no trade close, else 0/1/2. */ + int* __restrict__ out_labels +) { + int env = blockIdx.x * blockDim.x + threadIdx.x; + if (env >= n_episodes) return; + + /* Default: mask. Most bars have no trade close. */ + int label = TRADE_OUTCOME_MASK; + + /* Layout: trade_close_per_sample is [N, L] row-major. For env i at step + * t, the slot is at `i * L + t`. */ + long long out_off = (long long)env * (long long)L + (long long)t; + int closed = trade_close_per_sample[out_off]; + + if (closed != 0) { + float pvt = pnl_vs_target_at_close[env]; + float pvs = pnl_vs_stop_at_close[env]; + /* Priority: Profit > Stop > Timeout. If both target and stop + * were touched within the same close (rare edge case — trades + * usually exit on the first threshold hit), prefer Profit + * classification — the policy saw the profit side first if + * the trade succeeded. */ + if (pvt >= 1.0f) label = TRADE_OUTCOME_PROFIT; + else if (pvs >= 1.0f) label = TRADE_OUTCOME_STOP; + else label = TRADE_OUTCOME_TIMEOUT; + } + + out_labels[env] = label; +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 61fa7e10c..d5da0231a 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -17631,3 +17631,31 @@ Cargo check clean. - ~2-3 days focused engineering, reuses ~80% of current infrastructure Cargo check clean. + +#### Phase A2 — Trade-outcome label producer kernel (2026-05-13) + +First foundation kernel for the SP22 H6 vNext trade-outcome aux head proposal. + +**`trade_outcome_label_kernel.cu` (NEW)**: +- Per-env per-step classification kernel +- Reads `trade_close_per_sample[i*L+t]` (existing trade-close flag), `pnl_vs_target_at_close[env]`, `pnl_vs_stop_at_close[env]` (saved per-env at segment_complete) +- Writes K=3 outcome label `{0=Profit, 1=Stop, 2=Timeout}` plus mask sentinel `-1` for bars without trade close +- Priority: Profit > Stop > Timeout (rare edge case where both thresholds touched within single close prefers Profit) +- Pure per-env map; no atomicAdd, no reduction. Launch: grid(ceil(n_episodes/256)), block(256). +- Registered in `build.rs::kernels_with_common`. Cubin builds clean. + +**Status**: kernel compiled and registered as dead code (no launcher wireup yet). This is the FOUNDATION for the trade-outcome aux head — the new label signal source that aligns with the system's multi-bar trade decision horizon. + +**Remaining vNext work**: +- Phase A3: aux trade-outcome forward kernel (or extend existing K=2 head to K=3 mode) +- Phase A4: aux trade-outcome loss reduce kernel +- Phase A5: aux trade-outcome backward kernel +- Phase B: plan-params concat input (h_s2_aux + plan_params) +- Phase C: 3-slot state assembly (state[121..124]) +- Phase D: 12-weight W atom-shift (4 actions × 3 outcomes) +- Phase E: dW backward + Adam for 12-weight W +- Phase F: validation smoke + +Per the vNext spec, total remaining effort ~2-3 days. This commit lands the cornerstone label producer cleanly so subsequent phases have a stable foundation. + +**Save-for-backward** for the label producer's inputs (`pnl_vs_target_at_close`, `pnl_vs_stop_at_close`): need to add 2 new mapped-pinned buffers populated by `experience_env_step`'s segment_complete branch. Deferred to Phase A3 wireup commit.