gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

feat(rl): ISV-driven minimum hold time — hard constraint on churning

Browse Source 
Overrides closing actions (a3/a4/a9/a10) to Hold when
steps_since_done < RL_MIN_HOLD_STEPS_INDEX (slot 536, default 20).
The agent CANNOT exit before the minimum — forced to ride the wave.

Trail stops still fire regardless (safety overrides patience) —
if the market moves against the position past the trail distance,
the stop-loss exits even within the min-hold window.

Pipeline order: action selection → confidence gate → FRD gate →
min_hold_check → trail_mutate → trail_stop → heat_cap →
actions_to_market_targets

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-25 12:27:47 +02:00
parent 60c714c8d1
commit 7c38f339cd
4 changed files with 89 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1
crates/ml-alpha/build.rs
View File

@@ -88,6 +88,7 @@ const KERNELS: &[&str] = &[
"rl_encoder_context_broadcast", // P2: broadcast per-batch context (16 dims) into encoder input [B,K,56] at cols 40-55
"rl_gate_threshold_controller", // adaptive gate thresholds from trade frequency (dones EMA)
"rl_reward_shaping", // surfer-philosophy: entry cost + short-hold penalty + hold bonus
"rl_min_hold_check", // hard minimum hold time — override close actions to Hold before N steps
];
// Cache bust v31 — five new reduce / derive kernels populate the input

45
crates/ml-alpha/cuda/rl_min_hold_check.cu Normal file
View File

@@ -0,0 +1,45 @@
// rl_min_hold_check.cu — ISV-driven minimum hold time.
//
// Overrides closing actions (a3/a4/a9/a10) to Hold (a2) when the
// current position has been held for fewer than min_hold steps.
// Hard constraint — the agent cannot exit before the minimum.
//
// Runs AFTER action selection, BEFORE trail_stop_check (which can
// still force-exit on genuine stop-loss breaches regardless of
// hold time — safety overrides patience).
//
// Per `feedback_no_atomicadd`: per-batch element-wise.
// Per `feedback_cpu_is_read_only`: pure device-side.
#include <stdint.h>
#define ACTION_HOLD 2
#define ACTION_FLAT_FROM_LONG 3
#define ACTION_FLAT_FROM_SHORT 4
#define ACTION_HALF_FLAT_LONG 9
#define ACTION_HALF_FLAT_SHORT 10
#define RL_MIN_HOLD_STEPS_INDEX 536
extern "C" __global__ void rl_min_hold_check(
int* __restrict__ actions, // [B] IN/OUT
const int* __restrict__ steps_since_done, // [B]
const float* __restrict__ isv,
int b_size
) {
const int b = blockIdx.x * blockDim.x + threadIdx.x;
if (b >= b_size) return;
const int min_hold = (int)isv[RL_MIN_HOLD_STEPS_INDEX];
if (min_hold <= 0) return;
const int hold = steps_since_done[b];
if (hold >= min_hold) return;
const int action = actions[b];
if (action == ACTION_FLAT_FROM_LONG ||
action == ACTION_FLAT_FROM_SHORT ||
action == ACTION_HALF_FLAT_LONG ||
action == ACTION_HALF_FLAT_SHORT) {
actions[b] = ACTION_HOLD;
}
}

7
crates/ml-alpha/src/rl/isv_slots.rs
View File

@@ -926,6 +926,11 @@ pub const RL_SHORT_HOLD_PENALTY_INDEX: usize = 534;
/// Scales with sqrt(hold_time). Default: 0.50.
pub const RL_HOLD_BONUS_INDEX: usize = 535;
/// Hard minimum hold time (steps). Closing actions are overridden to
/// Hold when steps_since_done < this value. Trail stops still fire
/// regardless (safety overrides patience). Default: 20.
pub const RL_MIN_HOLD_STEPS_INDEX: usize = 536;
/// 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 = 536;
pub const RL_SLOTS_END: usize = 537;

39
crates/ml-alpha/src/trainer/integrated.rs
View File

@@ -204,6 +204,8 @@ const RL_GATE_THRESHOLD_CONTROLLER_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/rl_gate_threshold_controller.cubin"));
const RL_REWARD_SHAPING_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/rl_reward_shaping.cubin"));
const RL_MIN_HOLD_CHECK_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/rl_min_hold_check.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] =
@@ -467,6 +469,8 @@ pub struct IntegratedTrainer {
rl_gate_threshold_controller_fn: CudaFunction,
_rl_reward_shaping_module: Arc<CudaModule>,
rl_reward_shaping_fn: CudaFunction,
_rl_min_hold_check_module: Arc<CudaModule>,
rl_min_hold_check_fn: CudaFunction,
_rl_trade_context_update_module: Arc<CudaModule>,
rl_trade_context_update_fn: CudaFunction,
_rl_multires_features_update_module: Arc<CudaModule>,
@@ -990,6 +994,12 @@ impl IntegratedTrainer {
let rl_reward_shaping_fn = rl_reward_shaping_module
.load_function("rl_reward_shaping")
.context("load rl_reward_shaping")?;
let rl_min_hold_check_module = ctx
.load_cubin(RL_MIN_HOLD_CHECK_CUBIN.to_vec())
.context("load rl_min_hold_check cubin")?;
let rl_min_hold_check_fn = rl_min_hold_check_module
.load_function("rl_min_hold_check")
.context("load rl_min_hold_check")?;
let rl_trade_context_update_module = ctx
.load_cubin(RL_TRADE_CONTEXT_UPDATE_CUBIN.to_vec())
.context("load rl_trade_context_update cubin")?;
@@ -1419,6 +1429,8 @@ impl IntegratedTrainer {
rl_gate_threshold_controller_fn,
_rl_reward_shaping_module: rl_reward_shaping_module,
rl_reward_shaping_fn,
_rl_min_hold_check_module: rl_min_hold_check_module,
rl_min_hold_check_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,
@@ -1564,7 +1576,7 @@ impl IntegratedTrainer {
// (slot, value) pair — pure device write, no HtoD per
// `feedback_no_htod_htoh_only_mapped_pinned`.
{
let isv_constants: [(usize, f32); 72] = [
let isv_constants: [(usize, f32); 73] = [
// Static seeds for the adaptive reward-clamp controller —
// these are the initial values that
// `rl_reward_clamp_controller` will replace once it
@@ -1643,6 +1655,7 @@ impl IntegratedTrainer {
(crate::rl::isv_slots::RL_SHORT_HOLD_MIN_STEPS_INDEX, 20.0),
(crate::rl::isv_slots::RL_SHORT_HOLD_PENALTY_INDEX, 0.5),
(crate::rl::isv_slots::RL_HOLD_BONUS_INDEX, 0.50),
(crate::rl::isv_slots::RL_MIN_HOLD_STEPS_INDEX, 20.0),
(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),
@@ -3875,7 +3888,29 @@ impl IntegratedTrainer {
let pos_bytes_i = lobsim.pos_bytes() as i32;
let b_size_i = b_size as i32;
// ── SP20 P1+P5 trail mutate (a7/a8) — BEFORE actions_to_market_targets.
// Min hold check — override closing actions to Hold when
// hold time < ISV-driven minimum. Trail stops still fire
// (they run AFTER this, safety overrides patience).
{
let cfg = LaunchConfig {
grid_dim: (((b_size as u32) + 31) / 32, 1, 1),
block_dim: (32, 1, 1),
shared_mem_bytes: 0,
};
let mut launch = self.stream.launch_builder(&self.rl_min_hold_check_fn);
launch
.arg(&mut self.actions_d)
.arg(&self.steps_since_done_d)
.arg(&self.isv_d)
.arg(&b_size_i);
unsafe {
launch
.launch(cfg)
.context("rl_min_hold_check launch")?;
}
}
// Trail mutate (a7/a8) — BEFORE actions_to_market_targets.
// Mutates unit_trail_distance for active units if the chosen action
// is TrailTighten/Loosen. Non-trail actions pass through.
{