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]); }