diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index 377a2facc..71b6649f3 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -20,11 +20,12 @@ const KERNELS: &[&str] = &[ "layer_norm", // Phase 1: trunk pre-CfC normalisation ]; -// Cache bust v7 (2026-05-17): Phase 1.7 GRN heads — appended -// multi_horizon_heads_grn_fwd_batched + _bwd_batched kernels to -// cuda/multi_horizon_heads.cu (TFT gated residual: 2-layer MLP body + -// GLU gate + skip-projection from trunk). Old cubins don't have the -// new symbols. Force fresh nvcc compile. +// Cache bust v8 (2026-05-17): Phase 2C TGN Δt Fourier features — +// snap_feature_assemble.cu now emits 40 floats per snapshot (was 32) +// with 8 sin/cos features encoding Δt at log-spaced periods. Old +// cubins write only 32 floats; FEATURE_DIM bumped to 40 in Rust. +// Force fresh nvcc compile so the kernel write size matches the +// buffer allocations. fn main() { println!("cargo:rerun-if-changed=build.rs"); diff --git a/crates/ml-alpha/cuda/snap_feature_assemble.cu b/crates/ml-alpha/cuda/snap_feature_assemble.cu index 346ae9974..55dfc281b 100644 --- a/crates/ml-alpha/cuda/snap_feature_assemble.cu +++ b/crates/ml-alpha/cuda/snap_feature_assemble.cu @@ -1,7 +1,8 @@ // snap_feature_assemble.cu // -// Reads one MBP-10 snapshot in struct-of-arrays layout, emits the 32-dim -// feature vector documented in spec Section 2. +// Reads one MBP-10 snapshot in struct-of-arrays layout, emits the 40-dim +// feature vector documented in spec Section 2 (bumped from 32 → 40 in +// Phase 2C to add TGN-style Δt Fourier features). // // Layout enforced by tests: // out[0] = (mid - prev_mid) / tick_size [O(1), tick-normalized return] @@ -17,21 +18,53 @@ // out[20..26]= regime[0..6] (loader-precomputed EMA cascade — gives // model multi-minute context unreachable // inside the K-snapshot BPTT window) -// out[26..32]= reserved, zero +// out[26..32]= reserved, zero (future macro context) +// out[32..40]= TGN-style Δt Fourier features: +// (cos(ω_k · Δt_ns), sin(ω_k · Δt_ns))_{k=0..3} +// with log-spaced periods [60s, 6s, 600ms, 60ms]. +// Gives the model explicit Δt encoding that's +// discriminative across scales — critical with +// decision-stride > 1 where the gap between K-positions +// becomes irregular. Range: [-1, 1] for each slot. // // 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). +// parameters; Fourier features are intrinsically [-1, 1]. 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. +#define SNAP_FEATURE_DIM 40 + __device__ __forceinline__ float signed_log1p(float x) { return copysignf(log1pf(fabsf(x)), x); } +// Angular frequencies for TGN Δt Fourier encoding (rad / ns). +// Periods are log-spaced: 60s, 6s, 600ms, 60ms — span sub-second +// microstructure to minute-scale macro context. 2π / period_ns. +__device__ __constant__ float SNAP_DT_OMEGAS[4] = { + 1.0471975511965977e-10f, // 2π / 60e9 ns (60s) + 1.0471975511965977e-9f, // 2π / 6e9 ns (6s) + 1.0471975511965977e-8f, // 2π / 600e6 ns (600ms) + 1.0471975511965977e-7f, // 2π / 60e6 ns (60ms) +}; + +__device__ __forceinline__ void emit_dt_fourier(float* o_32, long dt_ns) { + // Convert to float once — at the longest period (60s) the angle for + // Δt = 60s is exactly 2π, so float precision is fine within ±100s. + const float dt = (float)dt_ns; + #pragma unroll + for (int k = 0; k < 4; ++k) { + const float theta = SNAP_DT_OMEGAS[k] * dt; + o_32[2 * k] = cosf(theta); + o_32[2 * k + 1] = sinf(theta); + } +} + extern "C" __global__ void snap_feature_assemble( const float* __restrict__ bid_px, // [10] const float* __restrict__ bid_sz, // [10] @@ -46,7 +79,7 @@ extern "C" __global__ void snap_feature_assemble( long ts_ns, long prev_ts_ns, float tick_size, - float* __restrict__ out // [32] + float* __restrict__ out // [40] ) { if (threadIdx.x != 0 || blockIdx.x != 0) return; @@ -64,13 +97,21 @@ extern "C" __global__ void snap_feature_assemble( out[17] = log1pf((float) trade_count); out[18] = signed_log1p(trade_signed_vol); - const float dt_ms = (float)(ts_ns - prev_ts_ns) * 1e-6f; + const long dt_ns = ts_ns - prev_ts_ns; + const float dt_ms = (float)dt_ns * 1e-6f; out[19] = log1pf(fmaxf(dt_ms, 0.0f)); // Regime context (loader-precomputed EMA cascade) — six slots. for (int i = 0; i < 6; ++i) out[20 + i] = regime[i]; + // Reserved-zero slots [26..32] kept for future macro context. for (int i = 26; i < 32; ++i) out[i] = 0.0f; + + // TGN Δt Fourier features [32..40] — 4 log-spaced periods, + // (cos, sin) pair each. Δt clamped to >= 0 to handle the rare + // out-of-order timestamp without producing nonsense angles. + const long dt_ns_clamped = dt_ns >= 0 ? dt_ns : 0; + emit_dt_fourier(out + 32, dt_ns_clamped); } @@ -101,7 +142,7 @@ extern "C" __global__ void snap_feature_assemble_batched( const long* __restrict__ prev_ts_ns, // [N] float tick_size, int N, - float* __restrict__ out // [N, 32] + float* __restrict__ out // [N, 40] ) { int n = blockIdx.x * blockDim.x + threadIdx.x; if (n >= N) return; @@ -113,7 +154,7 @@ extern "C" __global__ void snap_feature_assemble_batched( const float* pbs = prev_bid_sz + n * 10; const float* pas = prev_ask_sz + n * 10; const float* rg = regime + n * 6; - float* o = out + n * 32; + float* o = out + n * SNAP_FEATURE_DIM; const float mid = 0.5f * (bx[0] + ax[0]); const float pm = prev_mid[n]; @@ -130,9 +171,14 @@ extern "C" __global__ void snap_feature_assemble_batched( o[17] = log1pf((float) trade_count[n]); o[18] = signed_log1p(trade_signed_vol[n]); - const float dt_ms = (float)(ts_ns[n] - prev_ts_ns[n]) * 1e-6f; + const long dt_ns = ts_ns[n] - prev_ts_ns[n]; + const float dt_ms = (float)dt_ns * 1e-6f; o[19] = log1pf(fmaxf(dt_ms, 0.0f)); for (int i = 0; i < 6; ++i) o[20 + i] = rg[i]; for (int i = 26; i < 32; ++i) o[i] = 0.0f; + + // TGN Δt Fourier features at slots [32..40]. + const long dt_ns_clamped = dt_ns >= 0 ? dt_ns : 0; + emit_dt_fourier(o + 32, dt_ns_clamped); } diff --git a/crates/ml-alpha/src/cfc/snap_features.rs b/crates/ml-alpha/src/cfc/snap_features.rs index b3b56d95d..033147a5f 100644 --- a/crates/ml-alpha/src/cfc/snap_features.rs +++ b/crates/ml-alpha/src/cfc/snap_features.rs @@ -1,8 +1,11 @@ //! `snap_feature_assemble` kernel binding. //! -//! Per spec Section 2, builds a 32-dim feature vector from one MBP-10 -//! snapshot. Production code calls this inside the captured Graph A -//! (Task 11); the standalone helper exists for bit-equiv tests. +//! Per spec Section 2 + Phase 2C amendment (2026-05-17), builds a 40-dim +//! feature vector from one MBP-10 snapshot. Slots [0..32] match the +//! original layout; slots [32..40] hold TGN-style Δt Fourier features +//! (4 log-spaced periods × sin+cos). Production code calls this inside +//! the captured Graph A (Task 11); the standalone helper exists for +//! bit-equiv tests. use std::sync::Arc; @@ -15,7 +18,7 @@ use crate::pinned_mem::MappedF32Buffer; const KERNEL_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/snap_feature_assemble.cubin")); pub const ES_TICK_SIZE: f32 = 0.25; -pub const FEATURE_DIM: usize = 32; +pub const FEATURE_DIM: usize = 40; /// Number of regime-context features pre-computed by the loader and /// piped into snap_features slots `out[20..20+REGIME_DIM]`. These give diff --git a/crates/ml-alpha/tests/snap_feature_bit_equiv.rs b/crates/ml-alpha/tests/snap_feature_bit_equiv.rs index 14f9524ab..6855b47fc 100644 --- a/crates/ml-alpha/tests/snap_feature_bit_equiv.rs +++ b/crates/ml-alpha/tests/snap_feature_bit_equiv.rs @@ -93,11 +93,57 @@ fn reserved_slots_are_zero() { let dev = test_device(); let input = synthetic_input(); let gpu = snap_feature_assemble_gpu(&dev, &input).expect("gpu"); - for i in 20..FEATURE_DIM { - assert_eq!(gpu[i], 0.0, "reserved slot {i} must be 0, got {}", gpu[i]); + // Slots [20..26] are regime features. The synthetic input sets + // regime: [0.0; 6], so these come out zero by virtue of input. + // Slots [26..32] are the genuine reserved-zero block. + // Slots [32..40] are TGN Δt Fourier features — non-zero whenever + // Δt > 0 (synthetic has Δt = 500ms, so they're non-zero). + for i in 20..32 { + assert_eq!(gpu[i], 0.0, "reserved/regime slot {i} must be 0 with zero regime, got {}", gpu[i]); } } +#[test] +fn dt_fourier_features_are_bounded() { + let dev = test_device(); + let input = synthetic_input(); + let gpu = snap_feature_assemble_gpu(&dev, &input).expect("gpu"); + // Slots [32..40] are (cos, sin) pairs at 4 frequencies — must all + // be in [-1, 1]. + for i in 32..FEATURE_DIM { + assert!(gpu[i].abs() <= 1.0 + 1e-6, "Fourier slot {i} out of [-1,1]: {}", gpu[i]); + } + // cos²+sin² = 1 for each frequency pair. + for k in 0..4 { + let c = gpu[32 + 2 * k]; + let s = gpu[32 + 2 * k + 1]; + assert_relative_eq!(c * c + s * s, 1.0, epsilon = 1e-5); + } +} + +#[test] +fn dt_fourier_discriminates_scales() { + let dev = test_device(); + // Δt = 1ms — short-period frequencies (60ms, 600ms) should differ + // from the long-period ones (60s, 6s) more than the latter pair. + let mut input_1ms = synthetic_input(); + input_1ms.prev_ts_ns = input_1ms.ts_ns - 1_000_000; + // Δt = 1s — long-period frequencies should now be near-zero phase. + let mut input_1s = synthetic_input(); + input_1s.prev_ts_ns = input_1s.ts_ns - 1_000_000_000; + let gpu_1ms = snap_feature_assemble_gpu(&dev, &input_1ms).expect("gpu"); + let gpu_1s = snap_feature_assemble_gpu(&dev, &input_1s).expect("gpu"); + + // L2 distance over Fourier features [32..40] must be substantial — + // these should be visibly different Δt regimes. + let mut l2 = 0.0_f32; + for i in 32..FEATURE_DIM { + let d = gpu_1ms[i] - gpu_1s[i]; + l2 += d * d; + } + assert!(l2.sqrt() > 0.1, "Fourier features should differ at L2 > 0.1 across 1ms vs 1s, got {}", l2.sqrt()); +} + #[test] fn zero_prev_mid_yields_zero_tick_return() { let dev = test_device();