feat(ml-alpha): TGN Δt Fourier features in snap_feature_assemble (Phase 2C)

Bumps FEATURE_DIM 32→40. Slots [32..40] now carry 8 TGN-style Fourier
features encoding the elapsed time Δt = ts_ns - prev_ts_ns:
  (cos(ω_k · Δt_ns), sin(ω_k · Δt_ns))_{k=0..3}
at log-spaced periods [60s, 6s, 600ms, 60ms].

This gives Mamba2's input vector explicit Δt encoding that's
discriminative across temporal scales — particularly important once
decision-stride > 1 (Phase 2A) lands and the gap between K-positions
becomes irregular. Without these features the model has no way to
distinguish "1ms gap" from "1s gap" between consecutive K-positions.

Slots [0..32] unchanged (bit-equivalent for the first 32 features).
Reserved-zero slots [26..32] kept for future macro context. Slots
[20..26] still hold the loader-precomputed EMA regime cascade.

Frequencies stored in __constant__ memory (SNAP_DT_OMEGAS[4]) — small
table, broadcast read pattern, no register pressure. Frequency
selection rationale (one per log-decade):
  60s    — minute-scale macro session context
  6s     — 10s-scale liquidity windows
  600ms  — sub-second microstructure
  60ms   — tick-cluster spacing

Δt clamped to >= 0 so the rare out-of-order timestamp doesn't produce
nonsense angles. Each (cos, sin) pair satisfies cos²+sin² = 1
(verified by new test `dt_fourier_features_are_bounded`).

New tests in snap_feature_bit_equiv.rs:
  - dt_fourier_features_are_bounded: |slot| <= 1 + cos²+sin² == 1
  - dt_fourier_discriminates_scales: Δt=1ms vs Δt=1s produce L2-distinct
    Fourier vectors (>0.1)
  - reserved_slots_are_zero updated to check [20..32] (regime + reserved)
    instead of [20..FEATURE_DIM]

All 8 perception_overfit smokes still pass (synthetic stride=1 and
stride=4 both converge 0.32 → 0.0000) — proves the wider FEATURE_DIM=40
input doesn't break the Mamba2+LN+GRN chain.

build.rs cache-bust → v8.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>

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");

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

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

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();