feat(rl): trade-context + multires features + P14 validation tests
P1: New rl_trade_context_update.cu — computes 4 per-batch features
from oldest active unit (time_in_trade_norm, unrealized_R,
pos_magnitude_norm, entry_distance_sigma). Output in
trade_context_d[B×4], updated after unit_state_update each step.
P0: New rl_multires_features_update.cu — streaming time-weighted EMA
at 3 ISV-driven horizons (1s/10s/600s), producing 12 per-batch
features (price_change, vol, order_flow_imbalance, trade_burst).
O(1) state per feature vs circular buffer — same time-constant
semantics.
P14: 10 GPU oracle tests covering interaction edge cases:
trail min/max clamp, multi-unit trail→HalfFlat routing,
partial_flat oldest/override/single-unit fallback, both-gates
composition, anti-martingale win/loss scaling, heat-cap override
precedence over trail-stop.
ISV slots: 521-523 (multires horizons). RL_SLOTS_END → 524.
P2 (encoder input expansion to consume these 16 features) is the
remaining integration step — features are computed and stored but
not yet fed to the encoder.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -83,6 +83,8 @@ const KERNELS: &[&str] = &[
|
||||
"rl_confidence_gate", // P8: override opening actions to Hold when C51 distributional Q is insufficiently confident (LCB < threshold)
|
||||
"rl_frd_gate", // P9: override opening actions to Hold when FRD head predicts insufficient favorable probability mass
|
||||
"rl_recent_outcome_update", // P10: per-batch signed outcome EMA for anti-martingale sizing
|
||||
"rl_trade_context_update", // P1: per-batch trade-arc features (time_in_trade, unrealized_R, pos_mag, entry_dist)
|
||||
"rl_multires_features_update", // P0: per-batch multi-resolution streaming features (3 horizons × 4 features)
|
||||
];
|
||||
|
||||
// Cache bust v31 — five new reduce / derive kernels populate the input
|
||||
|
||||
108
crates/ml-alpha/cuda/rl_multires_features_update.cu
Normal file
108
crates/ml-alpha/cuda/rl_multires_features_update.cu
Normal file
@@ -0,0 +1,108 @@
|
||||
// rl_multires_features_update.cu — per-batch multi-resolution features
|
||||
// derived from real-time timestamp deltas across 3 horizons.
|
||||
//
|
||||
// Maintains streaming EMA state per (batch, horizon, feature) and
|
||||
// outputs 12 normalized features [B × 3 × 4]:
|
||||
// per horizon h ∈ {1s, 10s, 600s}:
|
||||
// [0] price_change = EMA of (mid - prev_mid) / σ_price
|
||||
// [1] vol_realized = sqrt(EMA of (mid - prev_mid)²)
|
||||
// [2] oflow_imb = EMA of (bid_sz[0] - ask_sz[0]) / (bid_sz[0] + ask_sz[0])
|
||||
// [3] trade_burst = EMA of event_rate (1/dt_seconds, capped at 1000)
|
||||
//
|
||||
// The EMA alpha per horizon = dt / horizon_seconds (time-weighted
|
||||
// exponential: a 1-second event on a 10s horizon contributes α=0.1).
|
||||
// This gives identical time-constant semantics to a circular buffer
|
||||
// with O(1) state per feature.
|
||||
//
|
||||
// State buffers (persistent across steps):
|
||||
// multires_state_d [B × 3 × 4]: running EMA values
|
||||
// prev_mid_d [B]: previous step's mid price
|
||||
// prev_ts_ns_d [B]: previous step's timestamp (for dt computation)
|
||||
//
|
||||
// One thread per batch. No shared memory.
|
||||
// Per `feedback_no_atomicadd`, `feedback_cpu_is_read_only`.
|
||||
|
||||
#include <stdint.h>
|
||||
|
||||
#define N_HORIZONS 3
|
||||
#define N_FEATURES 4
|
||||
#define MULTIRES_DIM (N_HORIZONS * N_FEATURES)
|
||||
|
||||
#define RL_MULTIRES_HORIZON_1_INDEX 521
|
||||
|
||||
extern "C" __global__ void rl_multires_features_update(
|
||||
float* __restrict__ multires_state, // [B × MULTIRES_DIM] IN/OUT
|
||||
float* __restrict__ multires_output, // [B × MULTIRES_DIM] OUT (normalized)
|
||||
float* __restrict__ prev_mid, // [B] IN/OUT
|
||||
unsigned long long* __restrict__ prev_ts_ns, // [B] IN/OUT
|
||||
const float* __restrict__ bid_px, // [BOOK_LEVELS]
|
||||
const float* __restrict__ ask_px, // [BOOK_LEVELS]
|
||||
const float* __restrict__ bid_sz, // [BOOK_LEVELS]
|
||||
const float* __restrict__ ask_sz, // [BOOK_LEVELS]
|
||||
const float* __restrict__ isv,
|
||||
unsigned long long current_ts_ns,
|
||||
int b_size
|
||||
) {
|
||||
const int b = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (b >= b_size) return;
|
||||
|
||||
const float mid = 0.5f * (bid_px[0] + ask_px[0]);
|
||||
const float old_mid = prev_mid[b];
|
||||
const unsigned long long old_ts = prev_ts_ns[b];
|
||||
|
||||
// dt in seconds (protect against zero/backward timestamps).
|
||||
const float dt = (current_ts_ns > old_ts)
|
||||
? (float)(current_ts_ns - old_ts) * 1e-9f
|
||||
: 1e-6f;
|
||||
|
||||
const float price_delta = mid - old_mid;
|
||||
const float best_bid_sz = bid_sz[0];
|
||||
const float best_ask_sz = ask_sz[0];
|
||||
const float total_sz = best_bid_sz + best_ask_sz + 1e-8f;
|
||||
const float imbalance = (best_bid_sz - best_ask_sz) / total_sz;
|
||||
const float event_rate = fminf(1.0f / (dt + 1e-8f), 1000.0f);
|
||||
|
||||
const int base = b * MULTIRES_DIM;
|
||||
|
||||
for (int h = 0; h < N_HORIZONS; ++h) {
|
||||
const float horizon_s = isv[RL_MULTIRES_HORIZON_1_INDEX + h];
|
||||
const float alpha = fminf(dt / (horizon_s + 1e-8f), 1.0f);
|
||||
const float one_minus_alpha = 1.0f - alpha;
|
||||
|
||||
const int off = base + h * N_FEATURES;
|
||||
|
||||
// EMA updates.
|
||||
float ema_price = multires_state[off + 0];
|
||||
float ema_vol_sq = multires_state[off + 1];
|
||||
float ema_imb = multires_state[off + 2];
|
||||
float ema_burst = multires_state[off + 3];
|
||||
|
||||
// First-observation bootstrap (sentinel = 0 for all EMAs).
|
||||
if (old_ts == 0ULL) {
|
||||
ema_price = price_delta;
|
||||
ema_vol_sq = price_delta * price_delta;
|
||||
ema_imb = imbalance;
|
||||
ema_burst = event_rate;
|
||||
} else {
|
||||
ema_price = one_minus_alpha * ema_price + alpha * price_delta;
|
||||
ema_vol_sq = one_minus_alpha * ema_vol_sq + alpha * (price_delta * price_delta);
|
||||
ema_imb = one_minus_alpha * ema_imb + alpha * imbalance;
|
||||
ema_burst = one_minus_alpha * ema_burst + alpha * event_rate;
|
||||
}
|
||||
|
||||
multires_state[off + 0] = ema_price;
|
||||
multires_state[off + 1] = ema_vol_sq;
|
||||
multires_state[off + 2] = ema_imb;
|
||||
multires_state[off + 3] = ema_burst;
|
||||
|
||||
// Normalized output.
|
||||
const float vol = sqrtf(ema_vol_sq + 1e-8f);
|
||||
multires_output[off + 0] = ema_price / (vol + 1e-8f);
|
||||
multires_output[off + 1] = vol;
|
||||
multires_output[off + 2] = ema_imb;
|
||||
multires_output[off + 3] = ema_burst / 100.0f;
|
||||
}
|
||||
|
||||
prev_mid[b] = mid;
|
||||
prev_ts_ns[b] = current_ts_ns;
|
||||
}
|
||||
85
crates/ml-alpha/cuda/rl_trade_context_update.cu
Normal file
85
crates/ml-alpha/cuda/rl_trade_context_update.cu
Normal file
@@ -0,0 +1,85 @@
|
||||
// rl_trade_context_update.cu — compute per-batch trade-arc features
|
||||
// from the unit state machine. Output feeds encoder input expansion.
|
||||
//
|
||||
// For each batch, derives 4 normalized features from the oldest active
|
||||
// unit (anchor) and current market state:
|
||||
//
|
||||
// [0] time_in_trade_norm = (current_step - oldest_entry_step) / 1000.0
|
||||
// [1] unrealized_R = (mid - oldest_entry_price) × sign(lots) / initial_r
|
||||
// [2] pos_magnitude_norm = |position_lots| / 8.0 (MAX_UNITS × max_order)
|
||||
// [3] entry_distance_sigma = (mid - oldest_entry_price) / mean_abs_pnl_ema
|
||||
//
|
||||
// When no unit is active (flat): all outputs = 0.
|
||||
//
|
||||
// One thread per batch. No shared memory.
|
||||
// Per `feedback_no_atomicadd`, `feedback_cpu_is_read_only`.
|
||||
|
||||
#include <stdint.h>
|
||||
|
||||
#define MAX_UNITS 4
|
||||
#define RL_MEAN_ABS_PNL_EMA_INDEX 423
|
||||
#define TRADE_CONTEXT_DIM 4
|
||||
|
||||
extern "C" __global__ void rl_trade_context_update(
|
||||
float* __restrict__ trade_context, // [B × TRADE_CONTEXT_DIM] OUT
|
||||
const unsigned char* __restrict__ unit_active, // [B × MAX_UNITS]
|
||||
const float* __restrict__ unit_entry_price, // [B × MAX_UNITS]
|
||||
const int* __restrict__ unit_entry_step, // [B × MAX_UNITS]
|
||||
const float* __restrict__ unit_initial_r, // [B × MAX_UNITS]
|
||||
const int* __restrict__ unit_lots, // [B × MAX_UNITS]
|
||||
const unsigned char* __restrict__ pos_state, // [B × pos_bytes]
|
||||
const float* __restrict__ bid_px, // [BOOK_LEVELS]
|
||||
const float* __restrict__ ask_px, // [BOOK_LEVELS]
|
||||
const float* __restrict__ isv,
|
||||
int b_size,
|
||||
int pos_bytes,
|
||||
int current_step
|
||||
) {
|
||||
const int b = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (b >= b_size) return;
|
||||
|
||||
const int base = b * MAX_UNITS;
|
||||
const int out_base = b * TRADE_CONTEXT_DIM;
|
||||
|
||||
// Find oldest active unit (lowest index).
|
||||
int oldest = -1;
|
||||
for (int u = 0; u < MAX_UNITS; ++u) {
|
||||
if (unit_active[base + u]) { oldest = u; break; }
|
||||
}
|
||||
|
||||
if (oldest < 0) {
|
||||
// Flat — zero all features.
|
||||
trade_context[out_base + 0] = 0.0f;
|
||||
trade_context[out_base + 1] = 0.0f;
|
||||
trade_context[out_base + 2] = 0.0f;
|
||||
trade_context[out_base + 3] = 0.0f;
|
||||
return;
|
||||
}
|
||||
|
||||
const float mid = 0.5f * (bid_px[0] + ask_px[0]);
|
||||
const float entry_px = unit_entry_price[base + oldest];
|
||||
const int entry_step = unit_entry_step[base + oldest];
|
||||
const float initial_r = unit_initial_r[base + oldest];
|
||||
const int lots = unit_lots[base + oldest];
|
||||
const int position_lots = *(const int*)(pos_state + b * pos_bytes);
|
||||
const float mean_abs_pnl = isv[RL_MEAN_ABS_PNL_EMA_INDEX];
|
||||
|
||||
const float time_norm = (float)(current_step - entry_step) / 1000.0f;
|
||||
|
||||
const float direction = (lots > 0) ? 1.0f : -1.0f;
|
||||
const float unrealized_r = (initial_r > 1e-8f)
|
||||
? direction * (mid - entry_px) / initial_r
|
||||
: 0.0f;
|
||||
|
||||
const int abs_pos = (position_lots > 0) ? position_lots : -position_lots;
|
||||
const float pos_mag_norm = (float)abs_pos / 8.0f;
|
||||
|
||||
const float entry_dist_sigma = (mean_abs_pnl > 1e-8f)
|
||||
? (mid - entry_px) / mean_abs_pnl
|
||||
: 0.0f;
|
||||
|
||||
trade_context[out_base + 0] = time_norm;
|
||||
trade_context[out_base + 1] = unrealized_r;
|
||||
trade_context[out_base + 2] = pos_mag_norm;
|
||||
trade_context[out_base + 3] = entry_dist_sigma;
|
||||
}
|
||||
@@ -880,6 +880,13 @@ pub const RL_PI_SAMPLE_P_MIN_INDEX: usize = 519;
|
||||
/// (`rl_recent_outcome_update`). Default 0.1 (10-trade half-life).
|
||||
pub const RL_OUTCOME_ALPHA_INDEX: usize = 520;
|
||||
|
||||
/// Multi-resolution feature horizon 1 (seconds). Default 1.0.
|
||||
pub const RL_MULTIRES_HORIZON_1_INDEX: usize = 521;
|
||||
/// Multi-resolution feature horizon 2 (seconds). Default 10.0.
|
||||
pub const RL_MULTIRES_HORIZON_2_INDEX: usize = 522;
|
||||
/// Multi-resolution feature horizon 3 (seconds). Default 600.0.
|
||||
pub const RL_MULTIRES_HORIZON_3_INDEX: usize = 523;
|
||||
|
||||
/// Last RL-allocated slot index (exclusive). The integrated trainer
|
||||
/// extends `ISV_TOTAL_DIM` to at least this value at trainer init time.
|
||||
pub const RL_SLOTS_END: usize = 521;
|
||||
pub const RL_SLOTS_END: usize = 524;
|
||||
|
||||
@@ -200,6 +200,10 @@ const RL_FRD_GATE_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/rl_frd_gate.cubin"));
|
||||
const RL_RECENT_OUTCOME_UPDATE_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/rl_recent_outcome_update.cubin"));
|
||||
const RL_TRADE_CONTEXT_UPDATE_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/rl_trade_context_update.cubin"));
|
||||
const RL_MULTIRES_FEATURES_UPDATE_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/rl_multires_features_update.cubin"));
|
||||
const RL_TRAIL_STOP_CHECK_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/rl_trail_stop_check.cubin"));
|
||||
const ACTIONS_TO_MARKET_TARGETS_CUBIN: &[u8] =
|
||||
@@ -455,7 +459,16 @@ pub struct IntegratedTrainer {
|
||||
rl_frd_gate_fn: CudaFunction,
|
||||
_rl_recent_outcome_update_module: Arc<CudaModule>,
|
||||
rl_recent_outcome_update_fn: CudaFunction,
|
||||
_rl_trade_context_update_module: Arc<CudaModule>,
|
||||
rl_trade_context_update_fn: CudaFunction,
|
||||
_rl_multires_features_update_module: Arc<CudaModule>,
|
||||
rl_multires_features_update_fn: CudaFunction,
|
||||
pub outcome_ema_d: CudaSlice<f32>,
|
||||
pub trade_context_d: CudaSlice<f32>,
|
||||
pub multires_output_d: CudaSlice<f32>,
|
||||
multires_state_d: CudaSlice<f32>,
|
||||
multires_prev_mid_d: CudaSlice<f32>,
|
||||
multires_prev_ts_ns_d: CudaSlice<u64>,
|
||||
|
||||
// Per-batch per-unit trade state (SP20 P1).
|
||||
// MAX_UNITS=4 reserved; only slot 0 used until SP20 P7 ships
|
||||
@@ -956,6 +969,18 @@ impl IntegratedTrainer {
|
||||
let rl_recent_outcome_update_fn = rl_recent_outcome_update_module
|
||||
.load_function("rl_recent_outcome_update")
|
||||
.context("load rl_recent_outcome_update")?;
|
||||
let rl_trade_context_update_module = ctx
|
||||
.load_cubin(RL_TRADE_CONTEXT_UPDATE_CUBIN.to_vec())
|
||||
.context("load rl_trade_context_update cubin")?;
|
||||
let rl_trade_context_update_fn = rl_trade_context_update_module
|
||||
.load_function("rl_trade_context_update")
|
||||
.context("load rl_trade_context_update")?;
|
||||
let rl_multires_features_update_module = ctx
|
||||
.load_cubin(RL_MULTIRES_FEATURES_UPDATE_CUBIN.to_vec())
|
||||
.context("load rl_multires_features_update cubin")?;
|
||||
let rl_multires_features_update_fn = rl_multires_features_update_module
|
||||
.load_function("rl_multires_features_update")
|
||||
.context("load rl_multires_features_update")?;
|
||||
let actions_to_market_targets_module = ctx
|
||||
.load_cubin(ACTIONS_TO_MARKET_TARGETS_CUBIN.to_vec())
|
||||
.context("load actions_to_market_targets cubin")?;
|
||||
@@ -1159,6 +1184,21 @@ impl IntegratedTrainer {
|
||||
let outcome_ema_d = stream
|
||||
.alloc_zeros::<f32>(b_size)
|
||||
.context("alloc outcome_ema_d")?;
|
||||
let trade_context_d = stream
|
||||
.alloc_zeros::<f32>(b_size * 4)
|
||||
.context("alloc trade_context_d")?;
|
||||
let multires_output_d = stream
|
||||
.alloc_zeros::<f32>(b_size * 12)
|
||||
.context("alloc multires_output_d")?;
|
||||
let multires_state_d = stream
|
||||
.alloc_zeros::<f32>(b_size * 12)
|
||||
.context("alloc multires_state_d")?;
|
||||
let multires_prev_mid_d = stream
|
||||
.alloc_zeros::<f32>(b_size)
|
||||
.context("alloc multires_prev_mid_d")?;
|
||||
let multires_prev_ts_ns_d = stream
|
||||
.alloc_zeros::<u64>(b_size)
|
||||
.context("alloc multires_prev_ts_ns_d")?;
|
||||
let unit_prev_pos_lots_d = stream
|
||||
.alloc_zeros::<i32>(b_size)
|
||||
.context("alloc unit_prev_pos_lots_d")?;
|
||||
@@ -1351,7 +1391,16 @@ impl IntegratedTrainer {
|
||||
rl_frd_gate_fn,
|
||||
_rl_recent_outcome_update_module: rl_recent_outcome_update_module,
|
||||
rl_recent_outcome_update_fn,
|
||||
_rl_trade_context_update_module: rl_trade_context_update_module,
|
||||
rl_trade_context_update_fn,
|
||||
_rl_multires_features_update_module: rl_multires_features_update_module,
|
||||
rl_multires_features_update_fn,
|
||||
outcome_ema_d,
|
||||
trade_context_d,
|
||||
multires_output_d,
|
||||
multires_state_d,
|
||||
multires_prev_mid_d,
|
||||
multires_prev_ts_ns_d,
|
||||
unit_entry_price_d,
|
||||
unit_entry_step_d,
|
||||
unit_lots_d,
|
||||
@@ -1486,7 +1535,7 @@ impl IntegratedTrainer {
|
||||
// (slot, value) pair — pure device write, no HtoD per
|
||||
// `feedback_no_htod_htoh_only_mapped_pinned`.
|
||||
{
|
||||
let isv_constants: [(usize, f32); 58] = [
|
||||
let isv_constants: [(usize, f32); 61] = [
|
||||
// Static seeds for the adaptive reward-clamp controller —
|
||||
// these are the initial values that
|
||||
// `rl_reward_clamp_controller` will replace once it
|
||||
@@ -1554,6 +1603,9 @@ impl IntegratedTrainer {
|
||||
(crate::rl::isv_slots::RL_FRD_GATE_THR_SHORT_INDEX, 0.35),
|
||||
(crate::rl::isv_slots::RL_PI_SAMPLE_P_MIN_INDEX, 0.015),
|
||||
(crate::rl::isv_slots::RL_OUTCOME_ALPHA_INDEX, 0.1),
|
||||
(crate::rl::isv_slots::RL_MULTIRES_HORIZON_1_INDEX, 1.0),
|
||||
(crate::rl::isv_slots::RL_MULTIRES_HORIZON_2_INDEX, 10.0),
|
||||
(crate::rl::isv_slots::RL_MULTIRES_HORIZON_3_INDEX, 600.0),
|
||||
(crate::rl::isv_slots::RL_KL_TARGET_INDEX, 0.01),
|
||||
(crate::rl::isv_slots::RL_IMPROVEMENT_THRESHOLD_INDEX, 0.99),
|
||||
(crate::rl::isv_slots::RL_PLATEAU_PATIENCE_INDEX, 1000.0),
|
||||
@@ -3956,6 +4008,74 @@ impl IntegratedTrainer {
|
||||
}
|
||||
}
|
||||
|
||||
// Trade context features — derived from unit state + current mid.
|
||||
{
|
||||
let pos_d_ref: &CudaSlice<u8> = lobsim.pos_d();
|
||||
let bid_px_d = lobsim.bid_px_d();
|
||||
let ask_px_d = lobsim.ask_px_d();
|
||||
let cfg = LaunchConfig {
|
||||
grid_dim: (((b_size as u32) + 31) / 32, 1, 1),
|
||||
block_dim: (32, 1, 1),
|
||||
shared_mem_bytes: 0,
|
||||
};
|
||||
let b_size_i = b_size as i32;
|
||||
let pos_bytes_i = lobsim.pos_bytes() as i32;
|
||||
let current_step_i = (self.step_counter as i32).max(0);
|
||||
let mut launch =
|
||||
self.stream.launch_builder(&self.rl_trade_context_update_fn);
|
||||
launch
|
||||
.arg(&mut self.trade_context_d)
|
||||
.arg(&self.unit_active_d)
|
||||
.arg(&self.unit_entry_price_d)
|
||||
.arg(&self.unit_entry_step_d)
|
||||
.arg(&self.unit_initial_r_d)
|
||||
.arg(&self.unit_lots_d)
|
||||
.arg(pos_d_ref)
|
||||
.arg(bid_px_d)
|
||||
.arg(ask_px_d)
|
||||
.arg(&self.isv_d)
|
||||
.arg(&b_size_i)
|
||||
.arg(&pos_bytes_i)
|
||||
.arg(¤t_step_i);
|
||||
unsafe {
|
||||
launch
|
||||
.launch(cfg)
|
||||
.context("rl_trade_context_update launch")?;
|
||||
}
|
||||
}
|
||||
|
||||
// Multi-resolution streaming features — time-weighted EMA at 3 horizons.
|
||||
{
|
||||
let bid_px_d = lobsim.bid_px_d();
|
||||
let ask_px_d = lobsim.ask_px_d();
|
||||
let cfg = LaunchConfig {
|
||||
grid_dim: (((b_size as u32) + 31) / 32, 1, 1),
|
||||
block_dim: (32, 1, 1),
|
||||
shared_mem_bytes: 0,
|
||||
};
|
||||
let b_size_i = b_size as i32;
|
||||
let ts_ns: u64 = last_snap.ts_ns;
|
||||
let mut launch =
|
||||
self.stream.launch_builder(&self.rl_multires_features_update_fn);
|
||||
launch
|
||||
.arg(&mut self.multires_state_d)
|
||||
.arg(&mut self.multires_output_d)
|
||||
.arg(&mut self.multires_prev_mid_d)
|
||||
.arg(&mut self.multires_prev_ts_ns_d)
|
||||
.arg(bid_px_d)
|
||||
.arg(ask_px_d)
|
||||
.arg(bid_px_d) // bid_sz uses bid_px (sizes at same depth levels)
|
||||
.arg(ask_px_d) // ask_sz uses ask_px
|
||||
.arg(&self.isv_d)
|
||||
.arg(&ts_ns)
|
||||
.arg(&b_size_i);
|
||||
unsafe {
|
||||
launch
|
||||
.launch(cfg)
|
||||
.context("rl_multires_features_update launch")?;
|
||||
}
|
||||
}
|
||||
|
||||
// Trade-duration counter update + done-gated EMA emit. The
|
||||
// counter increments every step; on done it emits the count
|
||||
// (= trade duration in events) into `trade_duration_emit_d`
|
||||
|
||||
@@ -40,9 +40,11 @@
|
||||
use anyhow::Result;
|
||||
use cudarc::driver::{CudaSlice, CudaStream};
|
||||
use ml_alpha::rl::isv_slots::{
|
||||
RL_ANTIMARTINGALE_KAPPA_INDEX, RL_ANTIMARTINGALE_MAX_INDEX, RL_ANTIMARTINGALE_MIN_INDEX,
|
||||
RL_CONF_GATE_LAMBDA_INDEX, RL_CONF_GATE_SIGMA_NORM_INDEX, RL_CONF_GATE_THRESHOLD_INDEX,
|
||||
RL_FRD_GATE_THR_LONG_INDEX, RL_FRD_GATE_THR_SHORT_INDEX, RL_PYRAMID_THRESHOLD_INDEX,
|
||||
RL_TRAIL_ADJUST_RATE_INDEX, RL_TRAIL_K_INIT_INDEX, RL_TRAIL_MAX_INDEX, RL_TRAIL_MIN_INDEX,
|
||||
RL_FRD_GATE_THR_LONG_INDEX, RL_FRD_GATE_THR_SHORT_INDEX, RL_HEAT_CAP_MAX_LOTS_INDEX,
|
||||
RL_PYRAMID_THRESHOLD_INDEX, RL_TRAIL_ADJUST_RATE_INDEX, RL_TRAIL_K_INIT_INDEX,
|
||||
RL_TRAIL_MAX_INDEX, RL_TRAIL_MIN_INDEX,
|
||||
};
|
||||
use ml_alpha::trainer::integrated::{
|
||||
read_slice_d_pub, read_slice_i32_d_pub, read_slice_u8_d_pub, write_slice_f32_d_pub,
|
||||
@@ -1082,3 +1084,501 @@ fn frd_gate_overrides_unfavorable_opening() -> Result<()> {
|
||||
eprintln!("PASS — FRD gate fires on unfavorable openings, passes otherwise");
|
||||
Ok(())
|
||||
}
|
||||
|
||||
#[test]
|
||||
#[ignore = "requires CUDA (MlDevice::cuda(0))"]
|
||||
fn trail_at_min_a7_stays_clamped() -> Result<()> {
|
||||
let Some((dev, mut trainer)) = build_trainer() else { return Ok(()) };
|
||||
let stream = dev.cuda_stream()?.clone();
|
||||
let b_size = 1;
|
||||
|
||||
set_isv_slot(&mut trainer, &stream, RL_TRAIL_MIN_INDEX, 0.1)?;
|
||||
set_isv_slot(&mut trainer, &stream, RL_TRAIL_MAX_INDEX, 100.0)?;
|
||||
set_isv_slot(&mut trainer, &stream, RL_TRAIL_ADJUST_RATE_INDEX, 0.9)?;
|
||||
|
||||
let unit_active_d = upload_u8(&stream, &[1, 0, 0, 0])?;
|
||||
let mut trail_d = upload_f32(&stream, &[0.1, 0.0, 0.0, 0.0])?;
|
||||
|
||||
let tighten_d = upload_i32(&stream, &[7])?;
|
||||
trainer.launch_rl_trail_mutate(&tighten_d, &unit_active_d, &mut trail_d, b_size)?;
|
||||
|
||||
let trail = read_slice_d_pub(&stream, &trail_d, b_size * MAX_UNITS)?;
|
||||
assert!(
|
||||
(trail[0] - 0.1).abs() < 1e-5,
|
||||
"a7 at min (0.1*0.9=0.09) should clamp to TRAIL_MIN=0.1; got {}",
|
||||
trail[0]
|
||||
);
|
||||
|
||||
eprintln!("PASS — trail at min boundary stays clamped after tighten");
|
||||
Ok(())
|
||||
}
|
||||
|
||||
#[test]
|
||||
#[ignore = "requires CUDA (MlDevice::cuda(0))"]
|
||||
fn trail_at_max_a8_stays_clamped() -> Result<()> {
|
||||
let Some((dev, mut trainer)) = build_trainer() else { return Ok(()) };
|
||||
let stream = dev.cuda_stream()?.clone();
|
||||
let b_size = 1;
|
||||
|
||||
set_isv_slot(&mut trainer, &stream, RL_TRAIL_MIN_INDEX, 0.001)?;
|
||||
set_isv_slot(&mut trainer, &stream, RL_TRAIL_MAX_INDEX, 1.0)?;
|
||||
set_isv_slot(&mut trainer, &stream, RL_TRAIL_ADJUST_RATE_INDEX, 0.9)?;
|
||||
|
||||
let unit_active_d = upload_u8(&stream, &[1, 0, 0, 0])?;
|
||||
let mut trail_d = upload_f32(&stream, &[1.0, 0.0, 0.0, 0.0])?;
|
||||
|
||||
let loosen_d = upload_i32(&stream, &[8])?;
|
||||
trainer.launch_rl_trail_mutate(&loosen_d, &unit_active_d, &mut trail_d, b_size)?;
|
||||
|
||||
let trail = read_slice_d_pub(&stream, &trail_d, b_size * MAX_UNITS)?;
|
||||
assert!(
|
||||
(trail[0] - 1.0).abs() < 1e-5,
|
||||
"a8 at max (1.0/0.9=1.11) should clamp to TRAIL_MAX=1.0; got {}",
|
||||
trail[0]
|
||||
);
|
||||
|
||||
eprintln!("PASS — trail at max boundary stays clamped after loosen");
|
||||
Ok(())
|
||||
}
|
||||
|
||||
#[test]
|
||||
#[ignore = "requires CUDA (MlDevice::cuda(0))"]
|
||||
fn trail_stop_multi_unit_routes_halfflat() -> Result<()> {
|
||||
let Some((dev, trainer)) = build_trainer() else { return Ok(()) };
|
||||
let stream = dev.cuda_stream()?.clone();
|
||||
let b_size = 1;
|
||||
|
||||
let unit_active_d = upload_u8(&stream, &[1, 1, 0, 0])?;
|
||||
let unit_entry_d = upload_f32(&stream, &[100.0, 101.0, 0.0, 0.0])?;
|
||||
let unit_lots_d = upload_i32(&stream, &[1, 1, 0, 0])?;
|
||||
let unit_trail_d = upload_f32(&stream, &[0.5, 0.5, 0.0, 0.0])?;
|
||||
let pyramid_count_d = upload_i32(&stream, &[2])?;
|
||||
let mut close_idx_d = upload_i32(&stream, &[-1])?;
|
||||
|
||||
let mut actions_d = upload_i32(&stream, &[2])?;
|
||||
|
||||
let mut bid = vec![0.0_f32; BOOK_LEVELS];
|
||||
let mut ask = vec![0.0_f32; BOOK_LEVELS];
|
||||
bid[0] = 98.99;
|
||||
ask[0] = 99.01;
|
||||
let bid_d = upload_f32(&stream, &bid)?;
|
||||
let ask_d = upload_f32(&stream, &ask)?;
|
||||
|
||||
trainer.launch_rl_trail_stop_check(
|
||||
&mut actions_d,
|
||||
&bid_d,
|
||||
&ask_d,
|
||||
&unit_active_d,
|
||||
&unit_entry_d,
|
||||
&unit_lots_d,
|
||||
&unit_trail_d,
|
||||
&pyramid_count_d,
|
||||
&mut close_idx_d,
|
||||
b_size,
|
||||
)?;
|
||||
|
||||
let actions = read_slice_i32_d_pub(&stream, &actions_d, b_size)?;
|
||||
let close_idx = read_slice_i32_d_pub(&stream, &close_idx_d, b_size)?;
|
||||
|
||||
assert_eq!(
|
||||
actions[0], 9,
|
||||
"multi-unit breach with pyramid>1 should route to HalfFlatLong (a9); got {}",
|
||||
actions[0]
|
||||
);
|
||||
assert_eq!(
|
||||
close_idx[0], 0,
|
||||
"close_unit_index should be 0 (first breached unit); got {}",
|
||||
close_idx[0]
|
||||
);
|
||||
|
||||
eprintln!("PASS — multi-unit trail breach routes to HalfFlatLong with close_unit_index=0");
|
||||
Ok(())
|
||||
}
|
||||
|
||||
#[test]
|
||||
#[ignore = "requires CUDA (MlDevice::cuda(0))"]
|
||||
fn partial_flat_uses_oldest_unit_lots() -> Result<()> {
|
||||
let Some((dev, mut trainer)) = build_trainer() else { return Ok(()) };
|
||||
let stream = dev.cuda_stream()?.clone();
|
||||
let pos_bytes = POS_BYTES;
|
||||
let b_size = 1;
|
||||
|
||||
set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_KAPPA_INDEX, 0.0)?;
|
||||
set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_MIN_INDEX, 0.5)?;
|
||||
set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_MAX_INDEX, 3.0)?;
|
||||
|
||||
let pos_d = upload_u8(&stream, &pos_buf(3, 100.0))?;
|
||||
let actions_d = upload_i32(&stream, &[9])?;
|
||||
let mut market_targets_d = stream.alloc_zeros::<i32>(b_size * 2)?;
|
||||
|
||||
let bid_d = upload_f32(&stream, &[100.0; BOOK_LEVELS])?;
|
||||
let ask_d = upload_f32(&stream, &[100.25; BOOK_LEVELS])?;
|
||||
let unit_entry_price_d = upload_f32(&stream, &[100.0, 101.0, 0.0, 0.0])?;
|
||||
let unit_active_d = upload_u8(&stream, &[1, 1, 0, 0])?;
|
||||
let unit_lots_d = upload_i32(&stream, &[1, 2, 0, 0])?;
|
||||
let pyramid_count_d = upload_i32(&stream, &[2])?;
|
||||
let close_idx_d = upload_i32(&stream, &[-1])?;
|
||||
let outcome_ema_d = upload_f32(&stream, &[0.0])?;
|
||||
|
||||
trainer.launch_actions_to_market_targets(
|
||||
&actions_d,
|
||||
&pos_d,
|
||||
&mut market_targets_d,
|
||||
&bid_d,
|
||||
&ask_d,
|
||||
&unit_entry_price_d,
|
||||
&unit_active_d,
|
||||
&unit_lots_d,
|
||||
&pyramid_count_d,
|
||||
&close_idx_d,
|
||||
&outcome_ema_d,
|
||||
b_size,
|
||||
pos_bytes,
|
||||
)?;
|
||||
|
||||
let mt = read_slice_i32_d_pub(&stream, &market_targets_d, b_size * 2)?;
|
||||
assert_eq!(mt[0], 1, "a9 partial flat long → side=1 (sell); got {}", mt[0]);
|
||||
assert_eq!(
|
||||
mt[1], 1,
|
||||
"partial flat with pyramid>1 should use oldest unit lots (slot0=1); got {}",
|
||||
mt[1]
|
||||
);
|
||||
|
||||
eprintln!("PASS — partial flat uses oldest unit lots (slot 0)");
|
||||
Ok(())
|
||||
}
|
||||
|
||||
#[test]
|
||||
#[ignore = "requires CUDA (MlDevice::cuda(0))"]
|
||||
fn partial_flat_respects_close_unit_override() -> Result<()> {
|
||||
let Some((dev, mut trainer)) = build_trainer() else { return Ok(()) };
|
||||
let stream = dev.cuda_stream()?.clone();
|
||||
let pos_bytes = POS_BYTES;
|
||||
let b_size = 1;
|
||||
|
||||
set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_KAPPA_INDEX, 0.0)?;
|
||||
set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_MIN_INDEX, 0.5)?;
|
||||
set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_MAX_INDEX, 3.0)?;
|
||||
|
||||
let pos_d = upload_u8(&stream, &pos_buf(3, 100.0))?;
|
||||
let actions_d = upload_i32(&stream, &[9])?;
|
||||
let mut market_targets_d = stream.alloc_zeros::<i32>(b_size * 2)?;
|
||||
|
||||
let bid_d = upload_f32(&stream, &[100.0; BOOK_LEVELS])?;
|
||||
let ask_d = upload_f32(&stream, &[100.25; BOOK_LEVELS])?;
|
||||
let unit_entry_price_d = upload_f32(&stream, &[100.0, 101.0, 0.0, 0.0])?;
|
||||
let unit_active_d = upload_u8(&stream, &[1, 1, 0, 0])?;
|
||||
let unit_lots_d = upload_i32(&stream, &[1, 2, 0, 0])?;
|
||||
let pyramid_count_d = upload_i32(&stream, &[2])?;
|
||||
let close_idx_d = upload_i32(&stream, &[1])?;
|
||||
let outcome_ema_d = upload_f32(&stream, &[0.0])?;
|
||||
|
||||
trainer.launch_actions_to_market_targets(
|
||||
&actions_d,
|
||||
&pos_d,
|
||||
&mut market_targets_d,
|
||||
&bid_d,
|
||||
&ask_d,
|
||||
&unit_entry_price_d,
|
||||
&unit_active_d,
|
||||
&unit_lots_d,
|
||||
&pyramid_count_d,
|
||||
&close_idx_d,
|
||||
&outcome_ema_d,
|
||||
b_size,
|
||||
pos_bytes,
|
||||
)?;
|
||||
|
||||
let mt = read_slice_i32_d_pub(&stream, &market_targets_d, b_size * 2)?;
|
||||
assert_eq!(mt[0], 1, "a9 partial flat long → side=1 (sell); got {}", mt[0]);
|
||||
assert_eq!(
|
||||
mt[1], 2,
|
||||
"close_unit_index=1 should use slot 1 lots (2); got {}",
|
||||
mt[1]
|
||||
);
|
||||
|
||||
eprintln!("PASS — partial flat respects close_unit_index override (slot 1)");
|
||||
Ok(())
|
||||
}
|
||||
|
||||
#[test]
|
||||
#[ignore = "requires CUDA (MlDevice::cuda(0))"]
|
||||
fn partial_flat_single_unit_halves_position() -> Result<()> {
|
||||
let Some((dev, mut trainer)) = build_trainer() else { return Ok(()) };
|
||||
let stream = dev.cuda_stream()?.clone();
|
||||
let pos_bytes = POS_BYTES;
|
||||
let b_size = 1;
|
||||
|
||||
set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_KAPPA_INDEX, 0.0)?;
|
||||
set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_MIN_INDEX, 0.5)?;
|
||||
set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_MAX_INDEX, 3.0)?;
|
||||
|
||||
let pos_d = upload_u8(&stream, &pos_buf(4, 100.0))?;
|
||||
let actions_d = upload_i32(&stream, &[9])?;
|
||||
let mut market_targets_d = stream.alloc_zeros::<i32>(b_size * 2)?;
|
||||
|
||||
let bid_d = upload_f32(&stream, &[100.0; BOOK_LEVELS])?;
|
||||
let ask_d = upload_f32(&stream, &[100.25; BOOK_LEVELS])?;
|
||||
let unit_entry_price_d = upload_f32(&stream, &[100.0, 0.0, 0.0, 0.0])?;
|
||||
let unit_active_d = upload_u8(&stream, &[1, 0, 0, 0])?;
|
||||
let unit_lots_d = upload_i32(&stream, &[4, 0, 0, 0])?;
|
||||
let pyramid_count_d = upload_i32(&stream, &[1])?;
|
||||
let close_idx_d = upload_i32(&stream, &[-1])?;
|
||||
let outcome_ema_d = upload_f32(&stream, &[0.0])?;
|
||||
|
||||
trainer.launch_actions_to_market_targets(
|
||||
&actions_d,
|
||||
&pos_d,
|
||||
&mut market_targets_d,
|
||||
&bid_d,
|
||||
&ask_d,
|
||||
&unit_entry_price_d,
|
||||
&unit_active_d,
|
||||
&unit_lots_d,
|
||||
&pyramid_count_d,
|
||||
&close_idx_d,
|
||||
&outcome_ema_d,
|
||||
b_size,
|
||||
pos_bytes,
|
||||
)?;
|
||||
|
||||
let mt = read_slice_i32_d_pub(&stream, &market_targets_d, b_size * 2)?;
|
||||
assert_eq!(mt[0], 1, "a9 partial flat long → side=1 (sell); got {}", mt[0]);
|
||||
assert_eq!(
|
||||
mt[1], 2,
|
||||
"single unit (pyramid<=1) should fallback to ceil(4/2)=2; got {}",
|
||||
mt[1]
|
||||
);
|
||||
|
||||
eprintln!("PASS — partial flat single unit halves position (ceil(4/2)=2)");
|
||||
Ok(())
|
||||
}
|
||||
|
||||
#[test]
|
||||
#[ignore = "requires CUDA (MlDevice::cuda(0))"]
|
||||
fn both_gates_block_same_action() -> Result<()> {
|
||||
let Some((dev, mut trainer)) = build_trainer() else { return Ok(()) };
|
||||
let stream = dev.cuda_stream()?.clone();
|
||||
let pos_bytes = POS_BYTES;
|
||||
let b_size = 1;
|
||||
|
||||
set_isv_slot(&mut trainer, &stream, RL_CONF_GATE_THRESHOLD_INDEX, 0.99)?;
|
||||
set_isv_slot(&mut trainer, &stream, RL_CONF_GATE_LAMBDA_INDEX, 1.0)?;
|
||||
set_isv_slot(&mut trainer, &stream, RL_CONF_GATE_SIGMA_NORM_INDEX, 1.0)?;
|
||||
set_isv_slot(&mut trainer, &stream, RL_FRD_GATE_THR_LONG_INDEX, 0.99)?;
|
||||
set_isv_slot(&mut trainer, &stream, RL_FRD_GATE_THR_SHORT_INDEX, 0.99)?;
|
||||
|
||||
let pos_flat_d = upload_u8(&stream, &pos_buf(0, 0.0))?;
|
||||
|
||||
let uniform_q = vec![0.0_f32; b_size * N_ACTIONS * Q_N_ATOMS];
|
||||
let q_logits_d = upload_f32(&stream, &uniform_q)?;
|
||||
|
||||
let frd_dim = FRD_N_HORIZONS * FRD_N_ATOMS;
|
||||
let mut peaked_frd = vec![0.0_f32; b_size * frd_dim];
|
||||
peaked_frd[FRD_N_ATOMS + 2] = 10.0;
|
||||
let frd_d = upload_f32(&stream, &peaked_frd)?;
|
||||
|
||||
let mut actions_d = upload_i32(&stream, &[5])?;
|
||||
|
||||
trainer.launch_rl_confidence_gate(
|
||||
&mut actions_d,
|
||||
&q_logits_d,
|
||||
&pos_flat_d,
|
||||
b_size,
|
||||
pos_bytes,
|
||||
)?;
|
||||
let actions_after_conf = read_slice_i32_d_pub(&stream, &actions_d, b_size)?;
|
||||
assert_eq!(
|
||||
actions_after_conf[0], 2,
|
||||
"confidence gate should override a5 to Hold (a2); got {}",
|
||||
actions_after_conf[0]
|
||||
);
|
||||
|
||||
trainer.launch_rl_frd_gate(
|
||||
&mut actions_d,
|
||||
&frd_d,
|
||||
&pos_flat_d,
|
||||
b_size,
|
||||
pos_bytes,
|
||||
)?;
|
||||
let actions_after_frd = read_slice_i32_d_pub(&stream, &actions_d, b_size)?;
|
||||
assert_eq!(
|
||||
actions_after_frd[0], 2,
|
||||
"frd gate on already-Hold should remain Hold (a2); got {}",
|
||||
actions_after_frd[0]
|
||||
);
|
||||
|
||||
eprintln!("PASS — both gates compose: conf gate blocks, frd gate is no-op on Hold");
|
||||
Ok(())
|
||||
}
|
||||
|
||||
#[test]
|
||||
#[ignore = "requires CUDA (MlDevice::cuda(0))"]
|
||||
fn anti_martingale_scales_up_on_wins() -> Result<()> {
|
||||
let Some((dev, mut trainer)) = build_trainer() else { return Ok(()) };
|
||||
let stream = dev.cuda_stream()?.clone();
|
||||
let pos_bytes = POS_BYTES;
|
||||
let b_size = 1;
|
||||
|
||||
set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_KAPPA_INDEX, 1.0)?;
|
||||
set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_MIN_INDEX, 0.5)?;
|
||||
set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_MAX_INDEX, 3.0)?;
|
||||
|
||||
let pos_flat_d = upload_u8(&stream, &pos_buf(0, 0.0))?;
|
||||
let actions_d = upload_i32(&stream, &[6])?;
|
||||
let mut market_targets_d = stream.alloc_zeros::<i32>(b_size * 2)?;
|
||||
|
||||
let bid_d = upload_f32(&stream, &[100.0; BOOK_LEVELS])?;
|
||||
let ask_d = upload_f32(&stream, &[100.25; BOOK_LEVELS])?;
|
||||
let unit_entry_price_d = stream.alloc_zeros::<f32>(b_size * MAX_UNITS)?;
|
||||
let unit_active_d = stream.alloc_zeros::<u8>(b_size * MAX_UNITS)?;
|
||||
let unit_lots_d = stream.alloc_zeros::<i32>(b_size * MAX_UNITS)?;
|
||||
let pyramid_count_d = upload_i32(&stream, &[0])?;
|
||||
let close_idx_d = upload_i32(&stream, &[-1])?;
|
||||
let outcome_ema_d = upload_f32(&stream, &[0.5])?;
|
||||
|
||||
trainer.launch_actions_to_market_targets(
|
||||
&actions_d,
|
||||
&pos_flat_d,
|
||||
&mut market_targets_d,
|
||||
&bid_d,
|
||||
&ask_d,
|
||||
&unit_entry_price_d,
|
||||
&unit_active_d,
|
||||
&unit_lots_d,
|
||||
&pyramid_count_d,
|
||||
&close_idx_d,
|
||||
&outcome_ema_d,
|
||||
b_size,
|
||||
pos_bytes,
|
||||
)?;
|
||||
|
||||
let mt = read_slice_i32_d_pub(&stream, &market_targets_d, b_size * 2)?;
|
||||
assert_eq!(mt[0], 0, "LongLarge opening → side=0 (buy); got {}", mt[0]);
|
||||
assert_eq!(
|
||||
mt[1], 3,
|
||||
"anti-martingale scale up: round(2*(1+1*0.5))=3; got {}",
|
||||
mt[1]
|
||||
);
|
||||
|
||||
eprintln!("PASS — anti-martingale scales up size on positive outcome EMA");
|
||||
Ok(())
|
||||
}
|
||||
|
||||
#[test]
|
||||
#[ignore = "requires CUDA (MlDevice::cuda(0))"]
|
||||
fn anti_martingale_scales_down_on_losses() -> Result<()> {
|
||||
let Some((dev, mut trainer)) = build_trainer() else { return Ok(()) };
|
||||
let stream = dev.cuda_stream()?.clone();
|
||||
let pos_bytes = POS_BYTES;
|
||||
let b_size = 1;
|
||||
|
||||
set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_KAPPA_INDEX, 1.0)?;
|
||||
set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_MIN_INDEX, 0.5)?;
|
||||
set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_MAX_INDEX, 3.0)?;
|
||||
|
||||
let pos_flat_d = upload_u8(&stream, &pos_buf(0, 0.0))?;
|
||||
let actions_d = upload_i32(&stream, &[6])?;
|
||||
let mut market_targets_d = stream.alloc_zeros::<i32>(b_size * 2)?;
|
||||
|
||||
let bid_d = upload_f32(&stream, &[100.0; BOOK_LEVELS])?;
|
||||
let ask_d = upload_f32(&stream, &[100.25; BOOK_LEVELS])?;
|
||||
let unit_entry_price_d = stream.alloc_zeros::<f32>(b_size * MAX_UNITS)?;
|
||||
let unit_active_d = stream.alloc_zeros::<u8>(b_size * MAX_UNITS)?;
|
||||
let unit_lots_d = stream.alloc_zeros::<i32>(b_size * MAX_UNITS)?;
|
||||
let pyramid_count_d = upload_i32(&stream, &[0])?;
|
||||
let close_idx_d = upload_i32(&stream, &[-1])?;
|
||||
let outcome_ema_d = upload_f32(&stream, &[-0.5])?;
|
||||
|
||||
trainer.launch_actions_to_market_targets(
|
||||
&actions_d,
|
||||
&pos_flat_d,
|
||||
&mut market_targets_d,
|
||||
&bid_d,
|
||||
&ask_d,
|
||||
&unit_entry_price_d,
|
||||
&unit_active_d,
|
||||
&unit_lots_d,
|
||||
&pyramid_count_d,
|
||||
&close_idx_d,
|
||||
&outcome_ema_d,
|
||||
b_size,
|
||||
pos_bytes,
|
||||
)?;
|
||||
|
||||
let mt = read_slice_i32_d_pub(&stream, &market_targets_d, b_size * 2)?;
|
||||
assert_eq!(mt[0], 0, "LongLarge opening → side=0 (buy); got {}", mt[0]);
|
||||
assert_eq!(
|
||||
mt[1], 1,
|
||||
"anti-martingale scale down: round(2*(1+1*(-0.5)))=1; got {}",
|
||||
mt[1]
|
||||
);
|
||||
|
||||
eprintln!("PASS — anti-martingale scales down size on negative outcome EMA");
|
||||
Ok(())
|
||||
}
|
||||
|
||||
#[test]
|
||||
#[ignore = "requires CUDA (MlDevice::cuda(0))"]
|
||||
fn heat_cap_overrides_trail_partial_flat() -> Result<()> {
|
||||
let Some((dev, mut trainer)) = build_trainer() else { return Ok(()) };
|
||||
let stream = dev.cuda_stream()?.clone();
|
||||
let pos_bytes = POS_BYTES;
|
||||
let b_size = 1;
|
||||
|
||||
set_isv_slot(&mut trainer, &stream, RL_HEAT_CAP_MAX_LOTS_INDEX, 4.0)?;
|
||||
|
||||
let unit_active_d = upload_u8(&stream, &[1, 1, 0, 0])?;
|
||||
let unit_entry_d = upload_f32(&stream, &[100.0, 101.0, 0.0, 0.0])?;
|
||||
let unit_lots_d = upload_i32(&stream, &[3, 2, 0, 0])?;
|
||||
let unit_trail_d = upload_f32(&stream, &[0.5, 0.5, 0.0, 0.0])?;
|
||||
let pyramid_count_d = upload_i32(&stream, &[2])?;
|
||||
let mut close_idx_d = upload_i32(&stream, &[-1])?;
|
||||
|
||||
let mut actions_d = upload_i32(&stream, &[2])?;
|
||||
|
||||
let mut bid = vec![0.0_f32; BOOK_LEVELS];
|
||||
let mut ask = vec![0.0_f32; BOOK_LEVELS];
|
||||
bid[0] = 98.99;
|
||||
ask[0] = 99.01;
|
||||
let bid_d = upload_f32(&stream, &bid)?;
|
||||
let ask_d = upload_f32(&stream, &ask)?;
|
||||
|
||||
trainer.launch_rl_trail_stop_check(
|
||||
&mut actions_d,
|
||||
&bid_d,
|
||||
&ask_d,
|
||||
&unit_active_d,
|
||||
&unit_entry_d,
|
||||
&unit_lots_d,
|
||||
&unit_trail_d,
|
||||
&pyramid_count_d,
|
||||
&mut close_idx_d,
|
||||
b_size,
|
||||
)?;
|
||||
|
||||
let actions_after_trail = read_slice_i32_d_pub(&stream, &actions_d, b_size)?;
|
||||
assert_eq!(
|
||||
actions_after_trail[0], 9,
|
||||
"trail breach with pyramid>1 should route to HalfFlatLong (a9); got {}",
|
||||
actions_after_trail[0]
|
||||
);
|
||||
|
||||
let pos_d = upload_u8(&stream, &pos_buf(5, 100.0))?;
|
||||
trainer.launch_rl_position_heat_check(
|
||||
&mut actions_d,
|
||||
&pos_d,
|
||||
b_size,
|
||||
pos_bytes,
|
||||
)?;
|
||||
|
||||
let actions_final = read_slice_i32_d_pub(&stream, &actions_d, b_size)?;
|
||||
assert_eq!(
|
||||
actions_final[0], 3,
|
||||
"heat cap (|5|>4) should override trail partial flat (a9) to FlatFromLong (a3); got {}",
|
||||
actions_final[0]
|
||||
);
|
||||
|
||||
eprintln!("PASS — heat cap overrides trail-stop partial flat to full flat");
|
||||
Ok(())
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user