feat(sp22-vNext): Phase A2 — trade-outcome label producer kernel
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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`
|
||||
|
||||
129
crates/ml/src/cuda_pipeline/trade_outcome_label_kernel.cu
Normal file
129
crates/ml/src/cuda_pipeline/trade_outcome_label_kernel.cu
Normal file
@@ -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 <cuda_runtime.h>
|
||||
|
||||
#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;
|
||||
}
|
||||
@@ -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.
|
||||
|
||||
Reference in New Issue
Block a user