feat(sp21): T3.3 ISV defrost — hold_reward_ema (atomic Phase 3.2 wireup)
Closes the Phase 3.2 forward-reference loop in the SP20 aggregator.
Previously `out_inputs->per_bar_hold_reward = 0.0f` was hardcoded; the
per-bar Hold opportunity-cost producer existed
(experience_kernels.cu:3823 — `per_bar_opp_cost = -aux_conf × cost_scale`)
and wrote to `hold_baseline_buffer` and `r_micro` directly, but never
reached HOLD_REWARD_EMA. Result: HOLD_REWARD_EMA frozen at sentinel 0.0
across all observed epochs; the SP20 reward centering loop
(`r_micro += per_bar_opp_cost - HOLD_REWARD_EMA`) stayed uncentered,
biasing the policy's reward signal away from zero-mean.
T3.3 wireup mirrors the T2.2 alpha refactor: a per-env scratch buffer
that the producer writes at every step (alongside the existing
`hold_baseline_buffer` write), and the aggregator reads at the same
step. Sums opp_cost over Hold-direction envs only; emits the mean as
`per_bar_hold_reward`. The HOLD_REWARD_EMA's gate (`hold_fraction > 0.5f`
from T3.2) preserves the strict-majority semantic from before.
Atomic across producer site, kernel signature, aggregator, launcher,
collector, and tests (per feedback_no_partial_refactor):
- experience_kernels.cu — new `float* per_bar_opp_cost_per_env`
kernel arg, NULL-tolerant write at line ~3823.
- sp20_aggregate_inputs_kernel.cu — new arg, 6th shmem stripe
(`sh_opp_cost_sum`), per-thread Hold-gated accumulation, tree
reduction extended to 6 stripes, output `per_bar_hold_reward =
opp_cost_sum / hold_count` when hold_count > 0 else 0.
- sp20_aggregate_inputs.rs — launcher signature, dynamic_shmem_bytes
5→6 stripes, doc table, internal test.
- gpu_experience_collector.rs — new `per_bar_opp_cost_per_env:
CudaSlice<f32>` field, alloc, struct construction, kernel-arg
pass at both experience_env_step and sp20_aggregate_inputs
launches (raw_ptr).
- sp20_aggregate_inputs_test.rs — helper renamed
`run_kernel_with_is_win_and_opp_cost`, 4 existing sites pass
NULL, 2 NEW oracle tests verifying real producer + NULL fallback.
- sp20_phase1_4_wireup_test.rs — NULL fallback at the wireup site;
HOLD_REWARD_EMA-stays-at-zero assertion remains valid.
Verification:
- cargo check -p ml --tests: passes (warnings only)
- cargo test -p ml --test sp20_aggregate_inputs_test --features cuda
-- --ignored: 12/12 GPU oracle tests pass on RTX 3050 Ti, including
both new T3.3 tests (per_bar_hold_reward_means_over_hold_envs_only,
null_per_bar_opp_cost_emits_zero).
- cargo test -p ml --test sp20_phase1_4_wireup_test --features cuda
-- --ignored: 2/2 pass under the new NULL-tolerant contract.
Plan reference: docs/plans/2026-05-10-sp21-train-eval-coherence-isv-defrost.md
Tier 3 status: T3.1 ✓, T3.2 ✓, T3.3 ✓ (this commit), T3.4 withdrawn,
T3.5 cascade-pending, T3.6 withdrawn.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -2212,6 +2212,17 @@ extern "C" __global__ void experience_env_step(
|
||||
* 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,
|
||||
/* SP21 T3.3 (2026-05-10) — per-env per-bar Hold opportunity-cost.
|
||||
* `per_bar_opp_cost_per_env` ← [N] f32: written at the per-bar
|
||||
* Hold opp-cost computation site (line ~3823) alongside the
|
||||
* existing `hold_baseline_buffer` write. Each slot is
|
||||
* `-aux_conf × cost_scale ∈ [-0.25, 0]` for the env's current
|
||||
* step. Consumer: sp20_aggregate_inputs_kernel sums over Hold-
|
||||
* direction envs and emits the mean as
|
||||
* `out_inputs->per_bar_hold_reward` → HOLD_REWARD_EMA. NULL-
|
||||
* tolerant for tests (consumer falls back to 0.0f matching the
|
||||
* pre-T3.3 Phase 3.2-forward-reference contract). */
|
||||
float* __restrict__ per_bar_opp_cost_per_env,
|
||||
/* SP20 Phase 3 Task 3.2 (2026-05-10) — Hold opportunity-cost dual
|
||||
* emission (Component 2) per spec §4.2.
|
||||
*
|
||||
@@ -3822,6 +3833,16 @@ extern "C" __global__ void experience_env_step(
|
||||
: 0.0f;
|
||||
const float per_bar_opp_cost = -aux_conf_i * cost_scale;
|
||||
|
||||
/* SP21 T3.3 (2026-05-10): write the per-env per-bar opp_cost
|
||||
* to the new per_bar_opp_cost_per_env buffer (consumed by
|
||||
* sp20_aggregate_inputs_kernel for the HOLD_REWARD_EMA
|
||||
* input). NULL-tolerant — when the buffer isn't wired
|
||||
* (oracle tests / pre-T3.3 callers), the slot stays
|
||||
* untouched. Race-free per-i write (one thread per env). */
|
||||
if (per_bar_opp_cost_per_env != NULL) {
|
||||
per_bar_opp_cost_per_env[i] = per_bar_opp_cost;
|
||||
}
|
||||
|
||||
/* Path 2 — always written when buffer wired. */
|
||||
if (has_hold_buf) {
|
||||
const long long row_off =
|
||||
|
||||
@@ -743,6 +743,17 @@ pub struct GpuExperienceCollector {
|
||||
/// `step_ret > 0` predicate when NULL for test scaffolds).
|
||||
pub(crate) is_win_per_env: CudaSlice<i32>, // [alloc_episodes]
|
||||
|
||||
/// SP21 T3.3 (2026-05-10): per-env per-bar Hold opportunity-cost
|
||||
/// scratch. Written every step by `experience_env_step` as
|
||||
/// `per_bar_opp_cost = -aux_conf × cost_scale ∈ [-0.25, 0]`.
|
||||
/// Consumer: `sp20_aggregate_inputs_kernel` reads `[env]` slot only
|
||||
/// for envs in `dir == DIR_HOLD` and emits the mean as
|
||||
/// `out_inputs->per_bar_hold_reward`. NULL-tolerant for tests
|
||||
/// (consumer falls back to `0.0f` matching pre-T3.3 Phase 3.2-
|
||||
/// forward-reference contract). Defrosts HOLD_REWARD_EMA — was
|
||||
/// stuck at 0 because the aggregator hardcoded 0.0f for this slot.
|
||||
pub(crate) per_bar_opp_cost_per_env: CudaSlice<f32>, // [alloc_episodes]
|
||||
|
||||
/// SP20 Phase 3 Task 3.2 (2026-05-10): per-env circular Hold
|
||||
/// opportunity-cost buffer (Component 2) per spec §4.2.
|
||||
///
|
||||
@@ -1893,6 +1904,19 @@ impl GpuExperienceCollector {
|
||||
"sp20: alloc is_win_per_env ({alloc_episodes} i32): {e}"
|
||||
)))?;
|
||||
|
||||
// SP21 T3.3 (2026-05-10): per-env per-bar Hold opportunity-cost
|
||||
// scratch. Sentinel 0.0; written by experience_env_step every
|
||||
// step as `per_bar_opp_cost = -aux_conf × cost_scale ∈ [-0.25, 0]`.
|
||||
// Consumed by sp20_aggregate_inputs_kernel for envs in DIR_HOLD,
|
||||
// emitted as mean → SP20EmaInputs.per_bar_hold_reward → drives
|
||||
// HOLD_REWARD_EMA (which was previously hardcoded 0 by the Phase
|
||||
// 3.2 forward-reference contract). Sized [alloc_episodes].
|
||||
let per_bar_opp_cost_per_env = stream
|
||||
.alloc_zeros::<f32>(alloc_episodes)
|
||||
.map_err(|e| MLError::ModelError(format!(
|
||||
"sp21: alloc per_bar_opp_cost_per_env ({alloc_episodes} f32): {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
|
||||
@@ -2954,6 +2978,7 @@ impl GpuExperienceCollector {
|
||||
label_at_open_per_env,
|
||||
alpha_per_env,
|
||||
is_win_per_env,
|
||||
per_bar_opp_cost_per_env,
|
||||
hold_baseline_buffer,
|
||||
states_out,
|
||||
actions_out,
|
||||
@@ -6396,6 +6421,16 @@ impl GpuExperienceCollector {
|
||||
// tick at close bars is dominated by tx_cost deduction
|
||||
// → systematically negative even on winning trades).
|
||||
.arg(&mut self.is_win_per_env)
|
||||
// SP21 T3.3 (2026-05-10): per-env per-bar Hold
|
||||
// opportunity-cost scratch buffer. Producer:
|
||||
// experience_env_step writes per_bar_opp_cost = -aux_conf
|
||||
// × cost_scale at line ~3823 alongside the existing
|
||||
// hold_baseline_buffer write. Consumer:
|
||||
// sp20_aggregate_inputs_kernel sums over Hold-direction
|
||||
// envs and emits mean → SP20EmaInputs.per_bar_hold_reward
|
||||
// → HOLD_REWARD_EMA (defrosts the previously-hardcoded
|
||||
// Phase 3.2 forward-reference 0.0f).
|
||||
.arg(&mut self.per_bar_opp_cost_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:
|
||||
@@ -6733,6 +6768,13 @@ impl GpuExperienceCollector {
|
||||
// `step_ret > 0` predicate that pinned WR_EMA at 0 in
|
||||
// production.
|
||||
let is_win_per_env_dev = self.is_win_per_env.raw_ptr();
|
||||
// SP21 T3.3 (2026-05-10): per-env per-bar Hold opp-cost
|
||||
// pointer. Populated every step by `experience_env_step`
|
||||
// (launched immediately above on the same stream — stream-
|
||||
// implicit producer→consumer ordering). Replaces the
|
||||
// hardcoded `out_inputs->per_bar_hold_reward = 0.0f`
|
||||
// placeholder atomically per `feedback_no_partial_refactor`.
|
||||
let per_bar_opp_cost_per_env_dev = self.per_bar_opp_cost_per_env.raw_ptr();
|
||||
unsafe {
|
||||
launch_sp20_aggregate_inputs(
|
||||
&self.stream,
|
||||
@@ -6747,6 +6789,7 @@ impl GpuExperienceCollector {
|
||||
aux_dir_acc_dev,
|
||||
alpha_per_env_dev,
|
||||
is_win_per_env_dev,
|
||||
per_bar_opp_cost_per_env_dev,
|
||||
n_envs_i32_sp20,
|
||||
dir_divisor,
|
||||
hold_action_id,
|
||||
|
||||
@@ -18,7 +18,7 @@
|
||||
//! | `trade_duration` | `round(mean(hold_at_exit) over closed envs)` if `is_close == 1`, else `0` |
|
||||
//! | `hold_fraction` | `hold_count / n_envs ∈ [0, 1]` (SP21 T3.2: was binary `action_is_hold` strict-majority vote) |
|
||||
//! | `alpha` | `mean(alpha_per_env) over closed envs` if `is_close == 1`, else `0` — **SP20 Phase 2 Task 2.2** (R_event - hold_baseline). |
|
||||
//! | `per_bar_hold_reward` | `0.0` — **Phase 3.2 forward reference** (-aux_conf * cost_scale) |
|
||||
//! | `per_bar_hold_reward` | `mean(per_bar_opp_cost over Hold-direction envs)` when `hold_count > 0`, else `0.0` (SP21 T3.3, was Phase 3.2 forward reference) |
|
||||
//! | `aux_logits_p50` | `aux_logits_p50_dev[0]` (sp20_stats output) |
|
||||
//! | `aux_logits_std` | `aux_logits_std_dev[0]` (sp20_stats output) |
|
||||
//! | `aux_dir_acc` | `aux_dir_acc_dev[0]` (existing aux_dir_acc_reduce_kernel output [0]) |
|
||||
@@ -67,13 +67,14 @@
|
||||
pub const SP20_AGGREGATE_BLOCK: u32 = 64;
|
||||
|
||||
/// Compute the dynamic shared-memory bytes the aggregation kernel
|
||||
/// needs for the configured block size: 5 stripes × `bdim` × 4 bytes
|
||||
/// (3 i32 stripes for closed/wins/hold counts + 2 f32 stripes for the
|
||||
/// hold_at_exit sum and the SP20 Phase 2 Task 2.2 alpha sum).
|
||||
/// needs for the configured block size: 6 stripes × `bdim` × 4 bytes
|
||||
/// (3 i32 stripes for closed/wins/hold counts + 3 f32 stripes for the
|
||||
/// hold_at_exit sum, the SP20 Phase 2 Task 2.2 alpha sum, and the
|
||||
/// SP21 T3.3 per-bar Hold opp-cost sum).
|
||||
#[must_use]
|
||||
pub fn dynamic_shmem_bytes() -> u32 {
|
||||
SP20_AGGREGATE_BLOCK
|
||||
.checked_mul(5)
|
||||
.checked_mul(6)
|
||||
.and_then(|x| x.checked_mul(std::mem::size_of::<i32>() as u32))
|
||||
.expect("SP20 aggregate: shmem bytes overflow")
|
||||
}
|
||||
@@ -144,6 +145,15 @@ pub unsafe fn launch_sp20_aggregate_inputs(
|
||||
// the per-bar `step_ret` at close bars is dominated by tx-cost so
|
||||
// the legacy predicate was essentially always false.
|
||||
is_win_per_env_dev: u64,
|
||||
// SP21 T3.3 (2026-05-10): per-env per-bar Hold opportunity-cost
|
||||
// device pointer. `[n_envs]` f32 buffer written every step by
|
||||
// `experience_env_step` as `per_bar_opp_cost = -aux_conf × cost_scale`.
|
||||
// The aggregator reads it only for Hold-direction envs and emits
|
||||
// mean opp-cost when `hold_count > 0`. NULL-tolerant — when `0`
|
||||
// the kernel emits `out_inputs->per_bar_hold_reward = 0.0f`,
|
||||
// matching the pre-T3.3 Phase 3.2-forward-reference contract used
|
||||
// by oracle tests that don't wire the producer.
|
||||
per_bar_opp_cost_per_env_dev: u64,
|
||||
n_envs: i32,
|
||||
dir_divisor: i32,
|
||||
hold_action_id: i32,
|
||||
@@ -191,6 +201,7 @@ pub unsafe fn launch_sp20_aggregate_inputs(
|
||||
.arg(&aux_dir_acc_dev)
|
||||
.arg(&alpha_per_env_dev)
|
||||
.arg(&is_win_per_env_dev)
|
||||
.arg(&per_bar_opp_cost_per_env_dev)
|
||||
.arg(&n_envs)
|
||||
.arg(&dir_divisor)
|
||||
.arg(&hold_action_id)
|
||||
@@ -218,7 +229,7 @@ mod tests {
|
||||
fn shmem_bytes_match_layout() {
|
||||
// SP20 Phase 2 Task 2.2: 5 stripes × bdim × 4 bytes (added the
|
||||
// alpha-sum stripe atomically with the per-env alpha producer).
|
||||
let expected = SP20_AGGREGATE_BLOCK as usize * 5 * 4;
|
||||
let expected = SP20_AGGREGATE_BLOCK as usize * 6 * 4;
|
||||
assert_eq!(dynamic_shmem_bytes() as usize, expected);
|
||||
}
|
||||
|
||||
|
||||
@@ -180,6 +180,22 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel(
|
||||
* with the per-env buffer populated by `experience_env_step`'s
|
||||
* segment_complete branch as `(segment_return > 0.0f) ? 1 : 0`. */
|
||||
const int* __restrict__ is_win_per_env, /* [n_envs] or NULL */
|
||||
/* SP21 T3.3 (2026-05-10) — per-env per-bar Hold opportunity-cost
|
||||
* input. `[n_envs]` f32 device buffer; written every step by
|
||||
* `experience_env_step` as `per_bar_opp_cost = -aux_conf × cost_scale ∈
|
||||
* [-0.25, 0]`. The aggregator reads `per_bar_opp_cost_per_env[env]`
|
||||
* only for envs in `dir == hold_action_id` (Hold-direction envs) so
|
||||
* non-Hold envs' opp_costs don't contaminate the average. The
|
||||
* downstream HOLD_REWARD_EMA's gate (`hold_fraction > 0.5f` at the
|
||||
* EMA kernel) ensures the EMA fires only when majority Hold —
|
||||
* preserving the strict-majority semantic of the pre-T3.2 binary
|
||||
* `action_is_hold` field.
|
||||
*
|
||||
* NULL-tolerant — when NULL the kernel emits
|
||||
* `out_inputs->per_bar_hold_reward = 0.0f` (matches the pre-T3.3
|
||||
* Phase 3.2-forward-reference contract used by aggregate-kernel
|
||||
* oracle tests that don't wire the per-bar producer). */
|
||||
const float* __restrict__ per_bar_opp_cost_per_env, /* [n_envs] or NULL */
|
||||
int n_envs,
|
||||
/* Action decoding parameters (production packed factored action
|
||||
* encoding). `dir_divisor = NUM_MAGNITUDES * NUM_ORD * NUM_URG`,
|
||||
@@ -197,19 +213,20 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel(
|
||||
const int tid = threadIdx.x;
|
||||
const int bdim = blockDim.x;
|
||||
|
||||
/* Dynamic shared memory: 5 stripes of bdim × 4 bytes each (4 stripes
|
||||
* pre-Task-2.2 + 1 stripe for the SP20 Phase 2 Task 2.2 alpha sum).
|
||||
/* Dynamic shared memory: 6 stripes of bdim × 4 bytes each (5 stripes
|
||||
* pre-T3.3 + 1 stripe for the SP21 T3.3 per-bar Hold opp-cost sum).
|
||||
* We use `int` storage for the 3 count reductions and `float` storage
|
||||
* for the hold-at-exit and alpha sums (alias-compatible because both
|
||||
* are 4 bytes wide; we cast between views as needed). */
|
||||
* for the hold-at-exit, alpha, and opp-cost sums (alias-compatible
|
||||
* because both are 4 bytes wide; we cast between views as needed). */
|
||||
extern __shared__ int sh_int[];
|
||||
/* Layout: [closed_count | wins_count | hold_count |
|
||||
* hold_at_exit_sum_f32 | alpha_sum_f32] */
|
||||
* hold_at_exit_sum_f32 | alpha_sum_f32 | opp_cost_sum_f32] */
|
||||
int* sh_closed_count = &sh_int[0];
|
||||
int* sh_wins_count = &sh_int[bdim];
|
||||
int* sh_hold_count = &sh_int[2 * bdim];
|
||||
float* sh_hold_at_exit_sum = (float*)&sh_int[3 * bdim];
|
||||
float* sh_alpha_sum = (float*)&sh_int[4 * bdim];
|
||||
float* sh_opp_cost_sum = (float*)&sh_int[5 * bdim];
|
||||
|
||||
/* Per-thread strided accumulation over [0, n_envs). */
|
||||
int local_closed_count = 0;
|
||||
@@ -217,6 +234,7 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel(
|
||||
int local_hold_count = 0;
|
||||
float local_hold_at_exit_sum = 0.0f;
|
||||
float local_alpha_sum = 0.0f;
|
||||
float local_opp_cost_sum = 0.0f;
|
||||
|
||||
for (int env = tid; env < n_envs; env += bdim) {
|
||||
const int tc = trade_close_per_env[env * env_stride];
|
||||
@@ -262,7 +280,17 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel(
|
||||
local_alpha_sum += alpha_per_env[env];
|
||||
}
|
||||
}
|
||||
if (dir == hold_action_id) local_hold_count += 1;
|
||||
if (dir == hold_action_id) {
|
||||
local_hold_count += 1;
|
||||
/* SP21 T3.3 (2026-05-10): accumulate per-bar Hold opp-cost
|
||||
* for Hold-direction envs only. NULL-tolerant — when the
|
||||
* producer isn't wired (oracle tests / pre-T3.3 callers),
|
||||
* the accumulator stays at 0 → `out_inputs->per_bar_hold_reward
|
||||
* = 0.0f`, matching the Phase 3.2 forward-reference contract. */
|
||||
if (per_bar_opp_cost_per_env != NULL) {
|
||||
local_opp_cost_sum += per_bar_opp_cost_per_env[env];
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
sh_closed_count[tid] = local_closed_count;
|
||||
@@ -270,10 +298,11 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel(
|
||||
sh_hold_count[tid] = local_hold_count;
|
||||
sh_hold_at_exit_sum[tid] = local_hold_at_exit_sum;
|
||||
sh_alpha_sum[tid] = local_alpha_sum;
|
||||
sh_opp_cost_sum[tid] = local_opp_cost_sum;
|
||||
__syncthreads();
|
||||
|
||||
/* Tree reduction across `bdim`. Standard log2(bdim) pattern. We
|
||||
* reduce all five stripes in one loop body — single barrier per
|
||||
* reduce all six stripes in one loop body — single barrier per
|
||||
* pass keeps the cost equal to a single reduction. */
|
||||
for (int s = bdim / 2; s > 0; s >>= 1) {
|
||||
if (tid < s) {
|
||||
@@ -282,6 +311,7 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel(
|
||||
sh_hold_count[tid] += sh_hold_count[tid + s];
|
||||
sh_hold_at_exit_sum[tid] += sh_hold_at_exit_sum[tid + s];
|
||||
sh_alpha_sum[tid] += sh_alpha_sum[tid + s];
|
||||
sh_opp_cost_sum[tid] += sh_opp_cost_sum[tid + s];
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
@@ -293,6 +323,7 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel(
|
||||
const int hold_count = sh_hold_count[0];
|
||||
const float hold_at_exit_sum_f = sh_hold_at_exit_sum[0];
|
||||
const float alpha_sum_f = sh_alpha_sum[0];
|
||||
const float opp_cost_sum_f = sh_opp_cost_sum[0];
|
||||
|
||||
const int is_close_out = (closed_count > 0) ? 1 : 0;
|
||||
|
||||
@@ -366,7 +397,21 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel(
|
||||
out_inputs->trade_duration = trade_duration_out;
|
||||
out_inputs->hold_fraction = hold_fraction_out; /* SP21 T3.2 */
|
||||
out_inputs->alpha = alpha_out; /* Task 2.2 — was 0.0f */
|
||||
out_inputs->per_bar_hold_reward = 0.0f; /* Phase 3.2 fwd ref */
|
||||
/* SP21 T3.3 (2026-05-10): per_bar_hold_reward = mean per-bar Hold
|
||||
* opp-cost over Hold-direction envs at this step. Replaces the
|
||||
* Phase 3.2 forward-reference `0.0f` placeholder. The accumulator
|
||||
* `opp_cost_sum_f` only summed over envs with `dir == hold_action_id`
|
||||
* (per-thread gate above), so dividing by `hold_count` gives the
|
||||
* Hold-conditional mean. When no envs are in Hold OR the producer
|
||||
* is NULL (oracle tests), the sum is 0 and we emit 0.0f matching
|
||||
* pre-T3.3 contract. The downstream HOLD_REWARD_EMA's gate
|
||||
* (`hold_fraction > 0.5f` at the EMA kernel) ensures the EMA fires
|
||||
* only when majority Hold — preserving strict-majority semantic. */
|
||||
float per_bar_hold_reward_out = 0.0f;
|
||||
if (hold_count > 0) {
|
||||
per_bar_hold_reward_out = opp_cost_sum_f / (float)hold_count;
|
||||
}
|
||||
out_inputs->per_bar_hold_reward = per_bar_hold_reward_out;
|
||||
out_inputs->aux_logits_p50 = p50_in;
|
||||
out_inputs->aux_logits_std = std_in;
|
||||
out_inputs->aux_dir_acc = dir_acc_in;
|
||||
|
||||
@@ -149,6 +149,7 @@ mod gpu {
|
||||
dir_buf.dev_ptr,
|
||||
/* alpha_per_env_dev = */ 0, // NULL — test scaffold without SP20 reward producer
|
||||
/* is_win_per_env_dev = */ 0, // NULL — falls back to `sr > 0` legacy predicate
|
||||
/* per_bar_opp_cost_per_env_dev = */ 0, // SP21 T3.3 NULL — emits per_bar_hold_reward=0
|
||||
/* n_envs = */ n as i32,
|
||||
DIR_DIVISOR,
|
||||
HOLD_ACTION_ID,
|
||||
@@ -218,6 +219,7 @@ mod gpu {
|
||||
dir_buf.dev_ptr,
|
||||
alpha_buf.dev_ptr, // SP20 Phase 2 Task 2.2 alpha producer
|
||||
/* is_win_per_env_dev = */ 0, // NULL — alpha-only variant; legacy `sr > 0` predicate
|
||||
/* per_bar_opp_cost_per_env_dev = */ 0, // SP21 T3.3 NULL — emits per_bar_hold_reward=0
|
||||
/* n_envs = */ n as i32,
|
||||
DIR_DIVISOR,
|
||||
HOLD_ACTION_ID,
|
||||
@@ -237,12 +239,17 @@ mod gpu {
|
||||
/// per-bar `step_ret > 0` predicate at close (which is dominated by
|
||||
/// tx-cost so almost always false at close bars, pinning WR_EMA at
|
||||
/// 0 in production).
|
||||
fn run_kernel_with_is_win(
|
||||
///
|
||||
/// SP21 T3.3 (2026-05-10): also accepts `per_bar_opp_cost_per_env`
|
||||
/// to exercise the new Hold opportunity-cost producer path. Pass
|
||||
/// an empty slice to use the NULL fallback (matches pre-T3.3 contract).
|
||||
fn run_kernel_with_is_win_and_opp_cost(
|
||||
trade_close: &[i32],
|
||||
step_ret: &[f32],
|
||||
hold_at_exit: &[f32],
|
||||
actions: &[i32],
|
||||
is_win_per_env: &[i32],
|
||||
per_bar_opp_cost_per_env: &[f32],
|
||||
aux_logits_p50: f32,
|
||||
aux_logits_std: f32,
|
||||
aux_dir_acc: f32,
|
||||
@@ -267,6 +274,22 @@ mod gpu {
|
||||
let win_buf = unsafe { MappedI32Buffer::new(n) }.expect("alloc is_win");
|
||||
win_buf.write_from_slice(is_win_per_env);
|
||||
|
||||
// SP21 T3.3 (2026-05-10): per-env per-bar Hold opp-cost producer
|
||||
// buffer. When the caller passes an empty slice, we pass NULL
|
||||
// (preserving pre-T3.3 contract). When non-empty, we wire the
|
||||
// producer so the aggregator emits a non-zero per_bar_hold_reward
|
||||
// for tests that exercise the new T3.3 path.
|
||||
let opp_buf_opt = if per_bar_opp_cost_per_env.is_empty() {
|
||||
None
|
||||
} else {
|
||||
assert_eq!(per_bar_opp_cost_per_env.len(), n,
|
||||
"per_bar_opp_cost_per_env len must match n_envs (or be empty for NULL)");
|
||||
let buf = unsafe { MappedF32Buffer::new(n) }.expect("alloc opp_cost");
|
||||
buf.write_from_slice(per_bar_opp_cost_per_env);
|
||||
Some(buf)
|
||||
};
|
||||
let opp_buf_dev = opp_buf_opt.as_ref().map_or(0_u64, |b| b.dev_ptr);
|
||||
|
||||
let p50_buf = unsafe { MappedF32Buffer::new(1) }.expect("alloc p50");
|
||||
p50_buf.write_from_slice(&[aux_logits_p50]);
|
||||
let std_buf = unsafe { MappedF32Buffer::new(1) }.expect("alloc std");
|
||||
@@ -292,6 +315,7 @@ mod gpu {
|
||||
dir_buf.dev_ptr,
|
||||
/* alpha_per_env_dev = */ 0,
|
||||
win_buf.dev_ptr, // SP20 Phase 2 Task 2.2-fix is_win producer
|
||||
opp_buf_dev, // SP21 T3.3 — wired iff per_bar_opp_cost_per_env non-empty
|
||||
/* n_envs = */ n as i32,
|
||||
DIR_DIVISOR,
|
||||
HOLD_ACTION_ID,
|
||||
@@ -566,9 +590,10 @@ mod gpu {
|
||||
// and accumulates `wins_count = 3`.
|
||||
let is_win_per_env = vec![1_i32, 1, 1, 0];
|
||||
|
||||
let (ic, iw, _, _, _, _, _, _, _) = run_kernel_with_is_win(
|
||||
let (ic, iw, _, _, _, _, _, _, _) = run_kernel_with_is_win_and_opp_cost(
|
||||
&trade_close, &step_ret, &hold_at_exit, &actions,
|
||||
&is_win_per_env, 0.0, 0.0, 0.0,
|
||||
&is_win_per_env, &[], // SP21 T3.3 NULL: opp-cost producer not exercised here
|
||||
0.0, 0.0, 0.0,
|
||||
);
|
||||
|
||||
assert_eq!(ic, 1, "is_close = 1 (4 envs closed)");
|
||||
@@ -615,4 +640,91 @@ mod gpu {
|
||||
fraction not binary majority. got iw={iw}",
|
||||
);
|
||||
}
|
||||
|
||||
/// SP21 T3.3 (2026-05-10) — per_bar_hold_reward producer wired.
|
||||
///
|
||||
/// Constructs 4 envs, 3 in DIR_HOLD with synthetic per-bar opp-costs
|
||||
/// `[-0.10, -0.20, -0.05, _]` (1 env in Long has its slot ignored —
|
||||
/// only Hold-direction envs contribute to the mean). Expected:
|
||||
/// `per_bar_hold_reward = mean(-0.10, -0.20, -0.05) = -0.1167`.
|
||||
/// Pins the contract that the aggregator (a) sums opp_cost only
|
||||
/// over Hold-dir envs, (b) divides by hold_count (NOT n_envs), and
|
||||
/// (c) emits the mean as a real signal (not 0.0 placeholder).
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn per_bar_hold_reward_means_over_hold_envs_only() {
|
||||
let trade_close = vec![0_i32; 4]; // no closes this step
|
||||
let step_ret = vec![0.0_f32; 4];
|
||||
let hold_at_exit = vec![0.0_f32; 4];
|
||||
// 3 Hold envs (DIR_HOLD=1), 1 Long env (DIR_LONG=2).
|
||||
let actions = vec![
|
||||
encode_dir(HOLD_ACTION_ID),
|
||||
encode_dir(HOLD_ACTION_ID),
|
||||
encode_dir(HOLD_ACTION_ID),
|
||||
encode_dir(2),
|
||||
];
|
||||
let is_win_per_env = vec![0_i32; 4]; // no wins (no closes)
|
||||
// Hold envs have opp_costs; Long env's slot is read by the
|
||||
// kernel but the per-thread accumulator is gated on
|
||||
// `dir == hold_action_id`, so its value (here -1.0 sentinel)
|
||||
// is correctly excluded from the sum.
|
||||
let per_bar_opp_cost = vec![-0.10_f32, -0.20, -0.05, -1.0];
|
||||
|
||||
let (_, _, _, ah, _, hr, _, _, _) = run_kernel_with_is_win_and_opp_cost(
|
||||
&trade_close, &step_ret, &hold_at_exit, &actions,
|
||||
&is_win_per_env, &per_bar_opp_cost,
|
||||
0.0, 0.0, 0.0,
|
||||
);
|
||||
|
||||
// hold_fraction = 3/4 = 0.75 (T3.2 contract).
|
||||
assert!(
|
||||
(ah - 0.75).abs() < 1e-6,
|
||||
"hold_fraction = 3/4 = 0.75 exact; got {ah}",
|
||||
);
|
||||
// per_bar_hold_reward = mean over Hold-dir envs only.
|
||||
// = (-0.10 + -0.20 + -0.05) / 3 = -0.35 / 3 ≈ -0.11667.
|
||||
// The Long env's -1.0 slot MUST NOT contribute (gated out).
|
||||
let expected = (-0.10_f32 - 0.20 - 0.05) / 3.0;
|
||||
assert!(
|
||||
(hr - expected).abs() < 1e-6,
|
||||
"per_bar_hold_reward = mean over Hold-dir envs = {expected:.6}; \
|
||||
got {hr}. The Long env's -1.0 sentinel MUST be excluded from \
|
||||
the sum (per-thread gate on `dir == hold_action_id`).",
|
||||
);
|
||||
}
|
||||
|
||||
/// SP21 T3.3 (2026-05-10) — NULL producer fallback contract.
|
||||
///
|
||||
/// When `per_bar_opp_cost_per_env` is NULL (oracle scaffolds without
|
||||
/// the new producer), the aggregator MUST emit
|
||||
/// `per_bar_hold_reward = 0.0f` regardless of hold_count or any
|
||||
/// other state. Preserves the pre-T3.3 Phase 3.2 forward-reference
|
||||
/// placeholder behavior bit-identically for backward compat.
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn null_per_bar_opp_cost_emits_zero() {
|
||||
let trade_close = vec![0_i32; 4];
|
||||
let step_ret = vec![0.0_f32; 4];
|
||||
let hold_at_exit = vec![0.0_f32; 4];
|
||||
// All 4 in Hold — hold_count = 4, hold_fraction = 1.0.
|
||||
let actions = vec![encode_dir(HOLD_ACTION_ID); 4];
|
||||
let is_win_per_env = vec![0_i32; 4];
|
||||
|
||||
let (_, _, _, ah, _, hr, _, _, _) = run_kernel_with_is_win_and_opp_cost(
|
||||
&trade_close, &step_ret, &hold_at_exit, &actions,
|
||||
&is_win_per_env, &[], // NULL — no opp-cost producer wired
|
||||
0.0, 0.0, 0.0,
|
||||
);
|
||||
|
||||
assert!(
|
||||
(ah - 1.0).abs() < 1e-6,
|
||||
"hold_fraction = 4/4 = 1.0; got {ah}",
|
||||
);
|
||||
assert_eq!(
|
||||
hr, 0.0,
|
||||
"NULL per_bar_opp_cost_per_env MUST emit per_bar_hold_reward=0.0 \
|
||||
(pre-T3.3 Phase 3.2 forward-reference contract preserved); \
|
||||
got {hr}",
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -178,6 +178,10 @@ mod gpu {
|
||||
// (e.g. WR_EMA = 1.0 with one
|
||||
// winning trade at step_ret > 0)
|
||||
// remain valid.
|
||||
/* per_bar_opp_cost_per_env_dev = */ 0, // SP21 T3.3 — wireup test stays
|
||||
// NULL; consumer falls back to
|
||||
// per_bar_hold_reward = 0.0f
|
||||
// matching pre-T3.3 contract.
|
||||
n_envs as i32,
|
||||
DIR_DIVISOR,
|
||||
HOLD_ACTION_ID,
|
||||
|
||||
@@ -2,6 +2,85 @@
|
||||
|
||||
**Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7.
|
||||
|
||||
## 2026-05-10 — SP21 T3.3: hold_reward_ema defrost (atomic Phase 3.2 producer wireup)
|
||||
|
||||
**Branch:** `sp20-aux-h-fixed` (off `1790a31b6`).
|
||||
**Commits:** 1 atomic (this commit).
|
||||
|
||||
Closes the Phase 3.2 forward-reference loop in the SP20 aggregator.
|
||||
Previously `out_inputs->per_bar_hold_reward = 0.0f` was hardcoded as a
|
||||
deferred Phase 3.2 placeholder; the per-bar Hold opportunity-cost
|
||||
producer existed (`experience_kernels.cu:3823` — `per_bar_opp_cost =
|
||||
-aux_conf × cost_scale`) and wrote to `hold_baseline_buffer` and
|
||||
`r_micro` directly, but never reached HOLD_REWARD_EMA. Result:
|
||||
HOLD_REWARD_EMA frozen at sentinel 0.0 across all observed epochs,
|
||||
the SP20 reward centering loop (`r_micro += per_bar_opp_cost -
|
||||
HOLD_REWARD_EMA`) stayed uncentered, biasing the policy's reward
|
||||
signal away from zero-mean.
|
||||
|
||||
**T3.3 wireup mirrors the T2.2 alpha refactor**: a per-env scratch
|
||||
buffer that the producer writes at every step (alongside the existing
|
||||
`hold_baseline_buffer` write), and the aggregator reads at the same
|
||||
step. Sums opp_cost over Hold-direction envs only; emits the mean as
|
||||
`per_bar_hold_reward`. The HOLD_REWARD_EMA's gate (`hold_fraction >
|
||||
0.5f` from T3.2) preserves the strict-majority semantic from before
|
||||
the binary→fractional refactor.
|
||||
|
||||
**Affected files (atomic per `feedback_no_partial_refactor`):**
|
||||
|
||||
- `crates/ml/src/cuda_pipeline/experience_kernels.cu` — new
|
||||
`float* per_bar_opp_cost_per_env` kernel arg; per-env producer
|
||||
write at the existing per-bar opp-cost computation site (line
|
||||
~3823) alongside the `hold_baseline_buffer` write.
|
||||
- `crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu` —
|
||||
new kernel arg, 6th shmem stripe (`sh_opp_cost_sum`), per-thread
|
||||
Hold-gated accumulation (`if (dir == hold_action_id)
|
||||
local_opp_cost_sum += per_bar_opp_cost_per_env[env]`), tree
|
||||
reduction extended to 6 stripes, output computation
|
||||
(`per_bar_hold_reward = opp_cost_sum / hold_count`).
|
||||
- `crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs` — launcher
|
||||
signature accepts `per_bar_opp_cost_per_env_dev: u64` (NULL = `0`),
|
||||
`dynamic_shmem_bytes()` 5→6 stripes, doc table, internal test.
|
||||
- `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` — new
|
||||
`pub(crate) per_bar_opp_cost_per_env: CudaSlice<f32>` field, alloc
|
||||
in constructor, struct-construction line, kernel-arg pass at both
|
||||
the experience_env_step launch and the sp20_aggregate_inputs
|
||||
launch (raw_ptr).
|
||||
- `crates/ml/tests/sp20_aggregate_inputs_test.rs` — `run_kernel_with_
|
||||
is_win` renamed `run_kernel_with_is_win_and_opp_cost`; 4 existing
|
||||
call sites pass NULL; **2 new GPU oracle tests** verifying
|
||||
(a) `per_bar_hold_reward = mean over Hold-dir envs` with synthetic
|
||||
opp_costs and a Long env whose slot must be excluded; (b) NULL
|
||||
producer fallback emits 0.0 preserving the pre-T3.3 contract.
|
||||
- `crates/ml/tests/sp20_phase1_4_wireup_test.rs` — NULL fallback at
|
||||
the wireup test's `launch_sp20_aggregate_inputs` site (the
|
||||
`alpha_and_hold_reward_emas_stay_at_zero_with_null_producers`
|
||||
assertion remains valid).
|
||||
|
||||
**Verification:**
|
||||
|
||||
- `cargo check -p ml --tests` — passes (warnings only).
|
||||
- `cargo test -p ml --test sp20_aggregate_inputs_test --features cuda
|
||||
-- --ignored` — **12/12 GPU oracle tests pass on RTX 3050 Ti**,
|
||||
including both new T3.3 tests (`per_bar_hold_reward_means_over_
|
||||
hold_envs_only`, `null_per_bar_opp_cost_emits_zero`).
|
||||
- `cargo test -p ml --test sp20_phase1_4_wireup_test --features cuda
|
||||
-- --ignored` — 2/2 pass; HOLD_REWARD_EMA assertion holds under
|
||||
the new NULL-tolerant contract.
|
||||
|
||||
**Pearls applied:**
|
||||
|
||||
- `pearl_first_observation_bootstrap` — HOLD_REWARD_EMA's existing
|
||||
bootstrap path (sentinel 0.0 → first observation) still applies;
|
||||
T3.3 only changes WHAT the observation is (real opp_cost mean
|
||||
instead of hardcoded 0).
|
||||
- `feedback_no_partial_refactor` — kernel signature, launcher,
|
||||
collector, tests all migrated atomically in this single commit.
|
||||
- `pearl_blend_formulas_must_have_permanent_floor` — the consumer
|
||||
side at `experience_kernels.cu:3838` (`r_micro += per_bar_opp_cost
|
||||
- HOLD_REWARD_EMA`) is now correctly centered; previously the
|
||||
EMA=0 made the centering a no-op.
|
||||
|
||||
## 2026-05-10 — SP21 Tier-3 foundation: T3.1+T3.2 ISV defrost (wr_ema, hold_pct_ema)
|
||||
|
||||
**Branch:** `sp20-aux-h-fixed` (off `6df44e4c6`).
|
||||
|
||||
@@ -132,9 +132,9 @@ Each ISV slot below stayed constant across all 3 observed epochs of `d7bj7`. The
|
||||
| T3.1 | `wr_ema` | 0.0000 (frozen) | EMA of segment-close win rate ∈ [0, 1] | `pearl_per_bar_vs_segment_pnl_signal_mismatch` patch added `is_win_per_env` buffer; verify producer is reading it AND writing to `wr_ema_idx` |
|
||||
| T3.2 | `hold_pct_ema` | 0.0000 (frozen) | EMA of fraction of bars where policy holds an active position | Likely an EMA producer reading from action_counts or position_active_mask |
|
||||
| T3.3 | `hold_reward_ema` | 0.0000 (frozen) | EMA of opportunity-cost reward magnitude per bar | Likely populated alongside hold_pct |
|
||||
| T3.4 | `loss_cap` | -1.0000 (sentinel) | Should be replaced with adaptive value after first observation per Pearl-A | First-observation bootstrap kernel — sentinel-detection guard may have wrong condition |
|
||||
| ~~T3.4~~ | ~~`loss_cap`~~ | ~~-1.0000 (sentinel)~~ | **WITHDRAWN 2026-05-10 (not a bug)**: investigation shows the value `-1.0` is NOT a sentinel — it's the LOSS_CAP controller's correct output for `WR_EMA < 0.50` (`-1.0 - clamp((wr_ema-0.50)/0.05, 0, 1)`). The actual sentinel per `state_reset_registry` is `0.0` (FoldReset bootstrap). After T3.1 fix, WR_EMA moves to actual val WR ≈ 0.46, but LOSS_CAP stays at -1.0 because WR is still below the 0.50 ramp threshold. **Will move to -2.0 only when WR_EMA crosses 0.55** (linear ramp in [0.50, 0.55]). GPU test `loss_cap_ramp_boundaries` verified the controller logic is correct (passed in our run). | **No fix needed at this slot.** Same shape as T3.6 — downstream of upstream WR signal being below the controller's design threshold. Will defrost automatically when WR climbs above 0.50. |
|
||||
| T3.5 | `hold_cost_scale` | 0.0100 (at min floor) | Controller-driven within [min, max] bounds | Controller reads `hold_pct_ema` (which is also frozen at 0) — chain failure: T3.2 must fix first |
|
||||
| T3.6 | `aux_conf_threshold` | 0.0100 (at min floor) | Controller drives based on aux confidence distribution | Producer kernel reads aux confidence histogram; check if histogram producer is firing |
|
||||
| ~~T3.6~~ | ~~`aux_conf_threshold`~~ | ~~0.0100 (at min floor)~~ | **WITHDRAWN 2026-05-10 (not a bug)**: investigation shows the controller at `sp20_controllers_compute_kernel.cu:152-157` (`clamp(aux_dir_acc_ema - 0.50, 0.01, 0.20)`) is *working correctly*. With observed `aux_dir_acc ≈ 0.46` (below random in d7bj7 logs), raw = -0.04, controller clamps at floor 0.01. Chain `aux_dir_acc_reduce_kernel → aggregator → EMA blend → controller` is intact and the EMA tracks the actual ~0.46 value. The slot will defrost on its own when aux improves above 0.51. **Downstream symptom of the upstream aux signal problem, NOT an independent producer-chain bug.** | **No fix needed at this slot.** Upstream problem (aux_dir_acc < 0.51) addressed by upstream feature/architecture work — separate from SP21 scope. |
|
||||
|
||||
**Strategy**: T3.1 fixed first as canonical example (it has a recent pearl + recent commit `64bbbe418` claiming the fix). Verify the patch actually propagated to the ISV slot, not just to `is_win_per_env` buffer. Once T3.1 is verified working, use the same pattern (find producer → check input availability → check write-to-isv → check sentinel guard) for T3.2-T3.6 in parallel.
|
||||
|
||||
|
||||
Reference in New Issue
Block a user