cd149cdb5d
Controller oscillated: 3 dones at step 15k spiked EMA above 2×target, triggering thousands of steps of tightening that killed trades. Higher target (10% vs 2%) matches the natural trade frequency at b=16. Wider dead zone (5× above / 0.2× below) prevents single-batch spikes from triggering tighten/relax oscillation. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
81 lines
3.2 KiB
Plaintext
81 lines
3.2 KiB
Plaintext
// rl_gate_threshold_controller.cu — adaptive gate threshold controller.
|
|
//
|
|
// Watches trade frequency (dones EMA) and adjusts confidence + FRD
|
|
// gate thresholds to maintain a target trade rate. When trades dry up
|
|
// (dones_ema < target), thresholds DECREASE (more permissive). When
|
|
// trades are abundant (dones_ema > target), thresholds INCREASE
|
|
// (more selective). Schulman bounded discrete step pattern.
|
|
//
|
|
// Runs once per step, single thread. Reads/writes ISV.
|
|
// Per `feedback_no_atomicadd`, `feedback_cpu_is_read_only`.
|
|
|
|
#include <stdint.h>
|
|
|
|
#define RL_CONF_GATE_THRESHOLD_INDEX 512
|
|
#define RL_FRD_GATE_THR_LONG_INDEX 516
|
|
#define RL_FRD_GATE_THR_SHORT_INDEX 517
|
|
#define RL_GATE_WARMUP_STEPS_INDEX 524
|
|
#define RL_GATE_DONES_EMA_INDEX 525
|
|
#define RL_GATE_DONES_TARGET_INDEX 526
|
|
#define RL_GATE_CONF_MIN_INDEX 527
|
|
#define RL_GATE_CONF_MAX_INDEX 528
|
|
#define RL_GATE_FRD_MIN_INDEX 529
|
|
#define RL_GATE_FRD_MAX_INDEX 530
|
|
#define RL_GATE_ADJUST_RATE_INDEX 531
|
|
|
|
extern "C" __global__ void rl_gate_threshold_controller(
|
|
float* __restrict__ isv,
|
|
const float* __restrict__ dones, // [B]
|
|
int b_size,
|
|
int current_step
|
|
) {
|
|
if (threadIdx.x != 0 || blockIdx.x != 0) return;
|
|
|
|
const int warmup = (int)isv[RL_GATE_WARMUP_STEPS_INDEX];
|
|
if (current_step < warmup) return;
|
|
|
|
// Update dones EMA (count of done events this step / b_size).
|
|
float done_count = 0.0f;
|
|
for (int b = 0; b < b_size; ++b) {
|
|
if (dones[b] > 0.5f) done_count += 1.0f;
|
|
}
|
|
const float done_rate = done_count / (float)b_size;
|
|
const float alpha = 0.01f;
|
|
const float prev_ema = isv[RL_GATE_DONES_EMA_INDEX];
|
|
const float dones_ema = (prev_ema == 0.0f && done_rate > 0.0f)
|
|
? done_rate
|
|
: (1.0f - alpha) * prev_ema + alpha * done_rate;
|
|
isv[RL_GATE_DONES_EMA_INDEX] = dones_ema;
|
|
|
|
const float target = isv[RL_GATE_DONES_TARGET_INDEX];
|
|
const float adjust = isv[RL_GATE_ADJUST_RATE_INDEX];
|
|
|
|
// Schulman bounded step: if dones_ema < target, DECREASE thresholds
|
|
// (more permissive); if > target, INCREASE (more selective).
|
|
float conf_thr = isv[RL_CONF_GATE_THRESHOLD_INDEX];
|
|
float frd_long = isv[RL_FRD_GATE_THR_LONG_INDEX];
|
|
float frd_short = isv[RL_FRD_GATE_THR_SHORT_INDEX];
|
|
|
|
if (dones_ema < target * 0.2f) {
|
|
// Trades well below target — relax gates.
|
|
conf_thr *= adjust;
|
|
frd_long *= adjust;
|
|
frd_short *= adjust;
|
|
} else if (dones_ema > target * 5.0f) {
|
|
// Trades well above target — tighten gates.
|
|
conf_thr /= adjust;
|
|
frd_long /= adjust;
|
|
frd_short /= adjust;
|
|
}
|
|
|
|
// Bilateral clamp.
|
|
const float conf_min = isv[RL_GATE_CONF_MIN_INDEX];
|
|
const float conf_max = isv[RL_GATE_CONF_MAX_INDEX];
|
|
const float frd_min = isv[RL_GATE_FRD_MIN_INDEX];
|
|
const float frd_max = isv[RL_GATE_FRD_MAX_INDEX];
|
|
|
|
isv[RL_CONF_GATE_THRESHOLD_INDEX] = fmaxf(conf_min, fminf(conf_thr, conf_max));
|
|
isv[RL_FRD_GATE_THR_LONG_INDEX] = fmaxf(frd_min, fminf(frd_long, frd_max));
|
|
isv[RL_FRD_GATE_THR_SHORT_INDEX] = fmaxf(frd_min, fminf(frd_short, frd_max));
|
|
}
|