Revert "arch(ml-alpha): restore count-delta redundancy at feature slot [26]"

This reverts commit 008f65d894.

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;

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]