Remove sm_90 tile config (falls back to GB10 BM=32)

torch-ext/dg_w4_kernels/fused_moe_w4_v2.py

@@ -73,18 +73,10 @@ def moe_align_cuda(topk_ids, block_m, E):
                                                cap, nblk, E, BLOCK=1024, num_warps=4)
     return a_row, expert_ids, cap, inv
 
-# Per-arch launch configs (BM included: BM>=64 selects the wgmma lowering on sm_90).
+# Per-arch launch configs.
 # GB10 (sm_121) keeps the shipped originals byte-for-byte — records unaffected.
-# sm_90 tier swept 2026-07-23 on H100 SXM5 (w4_h100_sweep.py --wgmma): gates = numerics
-# (rel_mean<5e-3) + run-to-run SELF-DETERMINISM per T. Winner is rel~1.4e-05 vs the BM=32
-# path and self-det=True at every T; 0.88-1.39 ms/layer (297-469 GB/s), 1.76-2.7x over the
-# mma tier, ~4.6x over the untuned GB10 config on H100.
 _CC = torch.cuda.get_device_capability() if torch.cuda.is_available() else (12, 1)
-if _CC[0] == 9:                                     # Hopper (H100/H200): wgmma tier
-    BM_V2 = 64
-    G1 = dict(BN=64, num_warps=4, num_stages=2)
-    G2 = dict(BN=128, num_warps=4, num_stages=4)
-elif _CC[0] == 12 and _CC[1] == 0:                   # sm_120: RTX 5090 / Blackwell consumer
+if _CC[0] == 12 and _CC[1] == 0:                     # sm_120: RTX 5090 / Blackwell consumer
     # Swept + e2e-gated on a 5090 (torch 2.13/triton 3.7): BM=16 is 1.08x on the isolated layer and
     # +3.4%/+3.9% end-to-end on gsm8k/humaneval vs the BM=32 fallback, accuracy unchanged (48/50 both).
     # Smaller BM wins because M/expert is ~16 at canvas 256, so BM=32 tiles are half padding.