diff --git a/crates/ml-alpha/cuda/snap_feature_assemble.cu b/crates/ml-alpha/cuda/snap_feature_assemble.cu index 49dbd6255..55dfc281b 100644 --- a/crates/ml-alpha/cuda/snap_feature_assemble.cu +++ b/crates/ml-alpha/cuda/snap_feature_assemble.cu @@ -18,15 +18,7 @@ // out[20..26]= regime[0..6] (loader-precomputed EMA cascade — gives // model multi-minute context unreachable // inside the K-snapshot BPTT window) -// out[26] = sgn(c) * log1p(|c|), c = trade_count delta -// [≥ 0 within file; ~0..5] -// SECOND exposure of count-delta to restore -// the Phase 1+2+3 redundancy that the C16 -// tick-rule swap broke (slot [18] used to -// equal slot [17] for non-negative deltas; -// now it's an uncorrelated signal — see -// `pearl_input_feature_redundancy_loadbearing`). -// out[27..32]= reserved, zero (future macro context) +// 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]. @@ -112,18 +104,8 @@ extern "C" __global__ void snap_feature_assemble( // Regime context (loader-precomputed EMA cascade) — six slots. for (int i = 0; i < 6; ++i) out[20 + i] = regime[i]; - // out[26] = signed_log1p(trade_count) — second exposure of the dense - // count-delta signal. Restores the Phase 1+2+3 redundancy where - // slots [17] and [18] were both log1p(count_delta); we kept the new - // tick-rule signal at [18] but the encoder lost the 2× W_in capacity - // allocation. Empirically (7.8M ES snaps, Q1+Q2 2024) tick-rule is - // uncorrelated with count-delta (r=0.0002) and misses 44% of trade - // events, so it can't substitute. Keeping both: dense, exact - // count_delta at [17] AND [26]; sparse, noisy, signed tick-rule - // at [18]. Slot was previously zero-reserved, so LayerNorm and - // Mamba2 W_in dimensions are unchanged. - out[26] = signed_log1p((float) trade_count); - for (int i = 27; i < 32; ++i) out[i] = 0.0f; + // 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 @@ -194,10 +176,7 @@ extern "C" __global__ void snap_feature_assemble_batched( o[19] = log1pf(fmaxf(dt_ms, 0.0f)); for (int i = 0; i < 6; ++i) o[20 + i] = rg[i]; - // Second exposure of count-delta — restores Phase 1+2+3 redundancy - // (see single-snapshot kernel comment for the empirical case). - o[26] = signed_log1p((float) trade_count[n]); - for (int i = 27; i < 32; ++i) o[i] = 0.0f; + 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; diff --git a/crates/ml-alpha/tests/snap_feature_bit_equiv.rs b/crates/ml-alpha/tests/snap_feature_bit_equiv.rs index c0f24e814..6855b47fc 100644 --- a/crates/ml-alpha/tests/snap_feature_bit_equiv.rs +++ b/crates/ml-alpha/tests/snap_feature_bit_equiv.rs @@ -95,21 +95,12 @@ fn reserved_slots_are_zero() { let gpu = snap_feature_assemble_gpu(&dev, &input).expect("gpu"); // Slots [20..26] are regime features. The synthetic input sets // regime: [0.0; 6], so these come out zero by virtue of input. - for i in 20..26 { - assert_eq!(gpu[i], 0.0, "regime slot {i} must be 0 with zero regime, got {}", gpu[i]); - } - // Slot [26] = signed_log1p(trade_count) — second exposure of the - // count-delta signal (added 2026-05-18 to restore Phase 1+2+3 - // redundancy after the C16 tick-rule swap). Synthetic input has - // trade_count = 7, so slot 26 = log(8) ≈ 2.0794. - let expected_slot_26 = (1.0f32 + 7.0f32).ln(); - assert_relative_eq!(gpu[26], expected_slot_26, epsilon = 1e-5); - // Slots [27..32] remain reserved-zero. - for i in 27..32 { - assert_eq!(gpu[i], 0.0, "reserved slot {i} must be 0, got {}", gpu[i]); - } + // 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). + // Δ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]