From 2289fa062a2e1546f5bb0e7831da85a72c6ac293 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 17 May 2026 00:33:17 +0200 Subject: [PATCH] feat(ml-alpha): normalize snap features to O(1)-O(10) scale MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit First L40S run (alpha-perception-s6hqv, commit 586d1e782) trained flat at chance: train_loss=0.6953, val_loss=0.6943, AUCs all near 0.50 across 5 epochs and 40k gradient steps. Synthetic overfit on the same trainer hit 0.59→0.19 in 250 steps, so wiring was sound. The signal-killer was feature scale: raw size deltas were ±100, dt_ms could exceed 1e4, and log-returns sat at ~1e-4. Mamba2's W_in projection saturates on those extremes, gradient bleeds out. Now in the kernel: out[0] = (mid - prev_mid) / tick_size [tick-return] out[12..17]= sgn(d) * log1p(|d|) [signed-log OFI] out[18] = sgn(v) * log1p(|v|) [signed-log vol] out[19] = log1p(max(dt_ms, 0)) [log dt] No tuned constants — tick_size is a market quantity; log1p and signed-log are monotone bounded transforms. Bit-equiv tests updated to assert the new closed-form expressions (still GPU-only, no CPU oracle). Co-Authored-By: Claude Opus 4.7 --- crates/ml-alpha/cuda/snap_feature_assemble.cu | 38 ++++++++++++------- .../ml-alpha/tests/snap_feature_bit_equiv.rs | 34 ++++++++++------- 2 files changed, 45 insertions(+), 27 deletions(-) diff --git a/crates/ml-alpha/cuda/snap_feature_assemble.cu b/crates/ml-alpha/cuda/snap_feature_assemble.cu index 6fac2d08b..b5ed16a77 100644 --- a/crates/ml-alpha/cuda/snap_feature_assemble.cu +++ b/crates/ml-alpha/cuda/snap_feature_assemble.cu @@ -3,20 +3,32 @@ // Reads one MBP-10 snapshot in struct-of-arrays layout, emits the 32-dim // feature vector documented in spec Section 2. // -// Layout enforced bit-for-bit by tests: -// out[0] = mid log-return since prev snapshot -// out[1] = L1 spread in ticks -// out[2..7] = bid-side per-level log-size (L1..L5) -// out[7..12] = ask-side per-level log-size (L1..L5) -// out[12..17]= per-level OFI (bid_delta - ask_delta) -// out[17] = log(trade_count + 1) -// out[18] = trade_signed_vol -// out[19] = dt_ms (snapshot inter-arrival time) +// Layout enforced by tests: +// out[0] = (mid - prev_mid) / tick_size [O(1), tick-normalized return] +// out[1] = (ask[0] - bid[0]) / tick_size [O(1), spread in ticks] +// out[2..7] = log1p(bid_sz[0..5]) [0..~7, log size] +// out[7..12] = log1p(ask_sz[0..5]) [0..~7] +// out[12..17]= sgn(d) * log1p(|d|), d = bid_delta - ask_delta +// [signed-log OFI, ±~7] +// out[17] = log1p(trade_count) [0..~5] +// out[18] = sgn(v) * log1p(|v|), v = trade_signed_vol +// [signed-log, ±~7] +// out[19] = log1p(max(dt_ms, 0)) [0..~9] // out[20..32]= reserved, zero // +// All transforms are monotone, bounded-rate, dimensionally clean — +// tick_size is a market constant; log1p / signed-log carry no tuned +// parameters. Brings every feature to comparable O(1)-O(10) scale so +// Mamba2's W_in projection doesn't saturate on the unbounded features +// (raw bid/ask size deltas were ±1e2, dt_ms could exceed 1e4). +// // Single-thread per-snapshot kernel; the call site launches with one // block of one thread per snapshot. +__device__ __forceinline__ float signed_log1p(float x) { + return copysignf(log1pf(fabsf(x)), x); +} + extern "C" __global__ void snap_feature_assemble( const float* __restrict__ bid_px, // [10] const float* __restrict__ bid_sz, // [10] @@ -35,7 +47,7 @@ extern "C" __global__ void snap_feature_assemble( if (threadIdx.x != 0 || blockIdx.x != 0) return; const float mid = 0.5f * (bid_px[0] + ask_px[0]); - out[0] = (mid > 0.0f && prev_mid > 0.0f) ? logf(mid / prev_mid) : 0.0f; + out[0] = (mid > 0.0f && prev_mid > 0.0f) ? (mid - prev_mid) / tick_size : 0.0f; out[1] = (ask_px[0] - bid_px[0]) / tick_size; for (int i = 0; i < 5; ++i) { @@ -43,13 +55,13 @@ extern "C" __global__ void snap_feature_assemble( out[7 + i] = log1pf(ask_sz[i]); const float bid_delta = bid_sz[i] - prev_bid_sz[i]; const float ask_delta = ask_sz[i] - prev_ask_sz[i]; - out[12 + i] = bid_delta - ask_delta; + out[12 + i] = signed_log1p(bid_delta - ask_delta); } out[17] = log1pf((float) trade_count); - out[18] = trade_signed_vol; + out[18] = signed_log1p(trade_signed_vol); const float dt_ms = (float)(ts_ns - prev_ts_ns) * 1e-6f; - out[19] = dt_ms; + out[19] = log1pf(fmaxf(dt_ms, 0.0f)); for (int i = 20; i < 32; ++i) out[i] = 0.0f; } diff --git a/crates/ml-alpha/tests/snap_feature_bit_equiv.rs b/crates/ml-alpha/tests/snap_feature_bit_equiv.rs index c885f7464..bafd51b81 100644 --- a/crates/ml-alpha/tests/snap_feature_bit_equiv.rs +++ b/crates/ml-alpha/tests/snap_feature_bit_equiv.rs @@ -46,17 +46,19 @@ fn spread_in_ticks_matches_synthetic() { } #[test] -fn mid_log_return_sign_tracks_input() { +fn mid_tick_return_sign_tracks_input() { let dev = test_device(); - // mid=5500.125; prev_mid=5499.875 → positive log-return + // mid=5500.125; prev_mid=5499.875 → positive tick-return (+1 tick) let input_up = synthetic_input(); let gpu_up = snap_feature_assemble_gpu(&dev, &input_up).expect("gpu"); - assert!(gpu_up[0] > 0.0, "mid > prev_mid must give positive log-return, got {}", gpu_up[0]); + assert!(gpu_up[0] > 0.0, "mid > prev_mid must give positive tick-return, got {}", gpu_up[0]); + // O(1) — must be within a few ticks for a one-tick price move. + assert!(gpu_up[0].abs() < 10.0, "tick-return out of expected range: {}", gpu_up[0]); let mut input_down = synthetic_input(); input_down.prev_mid = 5501.0; let gpu_down = snap_feature_assemble_gpu(&dev, &input_down).expect("gpu"); - assert!(gpu_down[0] < 0.0, "mid < prev_mid must give negative log-return, got {}", gpu_down[0]); + assert!(gpu_down[0] < 0.0, "mid < prev_mid must give negative tick-return, got {}", gpu_down[0]); } #[test] @@ -64,11 +66,13 @@ fn ofi_sign_tracks_bid_minus_ask_size() { let dev = test_device(); let input = synthetic_input(); let gpu = snap_feature_assemble_gpu(&dev, &input).expect("gpu"); - // prev sizes init to 0, so OFI[i] = bid_sz[i] - ask_sz[i]. - // bid_sz[0]=12, ask_sz[0]=10 → OFI[0] = 2. - assert_relative_eq!(gpu[12], input.bid_sz[0] - input.ask_sz[0], epsilon = 1e-6); - // Level 4: bid_sz[4]=16, ask_sz[4]=16 → OFI = 0. - assert_relative_eq!(gpu[16], input.bid_sz[4] - input.ask_sz[4], epsilon = 1e-6); + // prev sizes init to 0, so OFI[i] = bid_sz[i] - ask_sz[i] after signed-log compression. + // bid_sz[0]=12, ask_sz[0]=10 → raw d=2 → signed_log1p(2) = log1p(2) ≈ 1.0986. + let d0 = input.bid_sz[0] - input.ask_sz[0]; + assert_relative_eq!(gpu[12], d0.signum() * (1.0 + d0.abs()).ln(), epsilon = 1e-5); + // Level 4: bid_sz[4]=16, ask_sz[4]=16 → OFI = 0 → signed_log1p(0) = 0. + let d4 = input.bid_sz[4] - input.ask_sz[4]; + assert_relative_eq!(gpu[16], d4.signum() * (1.0 + d4.abs()).ln(), epsilon = 1e-5); } #[test] @@ -94,19 +98,21 @@ fn reserved_slots_are_zero() { } #[test] -fn zero_prev_mid_yields_zero_log_return() { +fn zero_prev_mid_yields_zero_tick_return() { let dev = test_device(); let mut input = synthetic_input(); input.prev_mid = 0.0; let gpu = snap_feature_assemble_gpu(&dev, &input).expect("gpu"); - assert_eq!(gpu[0], 0.0, "log-return must be 0 when prev_mid is 0"); + assert_eq!(gpu[0], 0.0, "tick-return must be 0 when prev_mid is 0"); } #[test] -fn dt_in_ms_matches_input_nanoseconds() { +fn dt_is_log_compressed() { let dev = test_device(); let input = synthetic_input(); let gpu = snap_feature_assemble_gpu(&dev, &input).expect("gpu"); - let expected_ms = (input.ts_ns - input.prev_ts_ns) as f32 * 1e-6; - assert_relative_eq!(gpu[19], expected_ms, epsilon = 1e-3); + let dt_ms = (input.ts_ns - input.prev_ts_ns) as f32 * 1e-6; + let expected = (1.0 + dt_ms.max(0.0)).ln(); + assert_relative_eq!(gpu[19], expected, epsilon = 1e-5); + assert!(gpu[19] >= 0.0, "log-compressed dt must be non-negative, got {}", gpu[19]); }