Remove dead warpspec/dqonce kernels, add deployment analysis docs

Remove the warp-specialized and dequant-once grouped GEMM kernels
from ops.cu — both were correct but slower than the baseline on Ada
due to register pressure from multiple accumulator sets. Also remove
the unused get_num_sms() helper. See moe-kernel-spec.md for the full
post-mortem.

Add deployment-summary.md with per-kernel performance tables at
M=1/4/64+, workload-weighted vLLM analysis, and memory savings.
Add moe-kernel-spec.md documenting the MoE optimization attempts
and hybrid dequant+cuBLAS BMM benchmarks. Update kbit-kernel-spec.md
with grouped scalar GEMV section and current optimization priorities.

226/226 tests pass.

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

Author: TimDettmers

benchmarks/bench_fp16_moe_sweep.py

@@ -0,0 +1,42 @@
+"""fp16 BMM baseline for MoE shapes across wide M range.
+
+Uses CUDA events (accurate for fp16 bmm which has no Python overhead).
+"""
+import torch
+
+NUM_EXPERTS = 8
+WARMUP = 50
+ITERS = 200
+
+dev = torch.device("cuda")
+
+shapes = [
+    ("moe_gu", 2048, 512),
+    ("moe_dn",  512, 2048),
+]
+
+m_vals = [1, 2, 4, 8, 16, 32, 64, 128, 256, 512, 1024, 2048, 4096]
+
+print(f"{'shape':<8} {'M':>5} {'fp16_us':>8}")
+print("-" * 24)
+
+for name, K_dim, N in shapes:
+    for M in m_vals:
+        A = torch.randn(NUM_EXPERTS, M, K_dim, dtype=torch.float16, device=dev)
+        B = torch.randn(NUM_EXPERTS, K_dim, N, dtype=torch.float16, device=dev)
+
+        fn = lambda: torch.bmm(A, B)
+        for _ in range(WARMUP):
+            fn()
+        torch.cuda.synchronize()
+
+        start = torch.cuda.Event(enable_timing=True)
+        end = torch.cuda.Event(enable_timing=True)
+        start.record()
+        for _ in range(ITERS):
+            fn()
+        end.record()
+        torch.cuda.synchronize()
+        t = start.elapsed_time(end) / ITERS * 1000  # us
+        print(f"{name:<8} {M:>5} {t:>8.1f}")
+    print()

benchmarks/ncu_moe_sweep.py

@@ -0,0 +1,65 @@
+"""NCU driver for MoE grouped MMA sweep across wide M range.
+
+Only k=4, but all power-of-2 M values from 1 to 4096.
+Usage: ncu --kernel-name "kbit_grouped_gemm_prod" --metrics gpu__time_duration.avg python benchmarks/ncu_moe_sweep.py
+"""
+import os, sys, torch
+
+for p in [".", ".."]:
+    if os.path.isdir(os.path.join(p, "bitsandbytes")):
+        sys.path.insert(0, os.path.abspath(p))
+        break
+
+import bitsandbytes
+from bitsandbytes.functional import create_normal_float_codebook
+
+NUM_EXPERTS = 8
+K_BITS = 4
+WARMUP = 3
+PROFILED = 5
+
+dev = torch.device("cuda")
+codebook = create_normal_float_codebook(K_BITS, device=dev)
+
+shapes = [
+    ("moe_gu", 2048, 512),
+    ("moe_dn",  512, 2048),
+]
+
+m_vals = [1, 2, 4, 8, 16, 32, 64, 128, 256, 512, 1024, 2048, 4096]
+
+# Pre-quantize
+moe_data = {}
+for name, K_dim, N in shapes:
+    packed_list, absmax_list = [], []
+    for _ in range(NUM_EXPERTS):
+        W = torch.randn(K_dim * N, device=dev, dtype=torch.float32)
+        pf, af = torch.ops.bitsandbytes.quantize_kbit(W, codebook, K_BITS)
+        pt, at = torch.ops.bitsandbytes.repack_kbit(pf, af, K_dim, N, K_BITS)
+        packed_list.append(pt)
+        absmax_list.append(at)
+    B_packed_all = torch.cat(packed_list, dim=0)
+    B_absmax_all = torch.cat(absmax_list, dim=0)
+    moe_data[name] = (K_dim, N, B_packed_all, B_absmax_all)
+
+# Print config to stderr
+print(f"shapes={[s[0] for s in shapes]} k={K_BITS} M={m_vals} W={WARMUP} P={PROFILED}", file=sys.stderr)
+
+for name, K_dim, N in shapes:
+    K_dim, N, B_packed_all, B_absmax_all = moe_data[name]
+    for M in m_vals:
+        total_tokens = M * NUM_EXPERTS
+        A_concat = torch.randn(total_tokens, K_dim, dtype=torch.float16, device=dev)
+        offsets = list(range(0, total_tokens + 1, M))
+        expert_offsets = torch.tensor(offsets, dtype=torch.int32, device=dev)
+
+        fn = lambda: torch.ops.bitsandbytes.kbit_grouped_gemm(
+            A_concat, B_packed_all, B_absmax_all, codebook,
+            expert_offsets, K_dim, N, K_BITS, NUM_EXPERTS, M)
+
+        for _ in range(WARMUP):
+            fn()
+        torch.cuda.synchronize()
+        for _ in range(PROFILED):
+            fn()
+        torch.cuda.synchronize()

benchmarks/ncu_single_moe.py

@@ -0,0 +1,44 @@
+"""Single MoE kernel invocation for detailed NCU profiling."""
+import torch, sys
+sys.path.insert(0, ".")
+import bitsandbytes
+from bitsandbytes.functional import quantize_kbit, create_normal_float_codebook
+
+torch.manual_seed(42)
+k, K_dim, N, num_experts, M = 4, 2048, 512, 8, 512
+codebook = create_normal_float_codebook(k, device="cuda")
+
+packed_list, absmax_list = [], []
+for e in range(num_experts):
+    W = torch.randn(N, K_dim, dtype=torch.float16, device="cuda")
+    packed, absmax, _ = quantize_kbit(W, k, codebook=codebook, absmax_format="fp32")
+    packed_list.append(packed)
+    absmax_list.append(absmax)
+
+B_packed_list, B_absmax_list = [], []
+for e in range(num_experts):
+    bp, ba = torch.ops.bitsandbytes.repack_kbit(packed_list[e], absmax_list[e], K_dim, N, k)
+    B_packed_list.append(bp)
+    B_absmax_list.append(ba)
+
+B_packed_all = torch.cat(B_packed_list)
+B_absmax_all = torch.cat(B_absmax_list)
+
+A_list, offsets = [], [0]
+for e in range(num_experts):
+    A = torch.randn(M, K_dim, dtype=torch.float16, device="cuda")
+    A_list.append(A)
+    offsets.append(offsets[-1] + M)
+A_concat = torch.cat(A_list, dim=0)
+eo = torch.tensor(offsets, dtype=torch.int32, device="cuda")
+
+# Warmup
+for _ in range(3):
+    C = torch.ops.bitsandbytes.kbit_grouped_gemm(
+        A_concat, B_packed_all, B_absmax_all, codebook, eo, K_dim, N, k, num_experts, M)
+torch.cuda.synchronize()
+
+# Profiled call
+C = torch.ops.bitsandbytes.kbit_grouped_gemm(
+    A_concat, B_packed_all, B_absmax_all, codebook, eo, K_dim, N, k, num_experts, M)
+torch.cuda.synchronize()

csrc/ops.cu

@@ -2507,6 +2507,11 @@ __global__ void kbit_grouped_gemm_prod(
     } // end persistent work loop
 }
 
+// [REMOVED: Warp-specialized and dequant-once grouped GEMM kernels.
+//  Both were correct but slower than the baseline on Ada (sm_89) due to
+//  register pressure from multiple accumulator sets. See moe-kernel-spec.md
+//  for the full analysis. Code removed in dead-code cleanup.]
+
 // Grouped GEMM launcher — supports TILE_N=64/128 and auto k_splits
 template <int K, int MB, int TN, typename scalar_t>
 static void kbitGroupedGemmProdLaunch(
@@ -2605,16 +2610,6 @@ void kbitGroupedGemmProd(
     }
 }
 
-// Cached SM count to avoid repeated cudaGetDevice/cudaDeviceGetAttribute calls
-static int cached_num_sms = 0;
-static int get_num_sms() {
-    if (cached_num_sms == 0) {
-        int dev;
-        cudaGetDevice(&dev);
-        cudaDeviceGetAttribute(&cached_num_sms, cudaDevAttrMultiProcessorCount, dev);
-    }
-    return cached_num_sms;
-}
 
 // ===================================================================
 // Scalar GEMV kernel: C[M,N] = A[M,K_dim] * W_kbit^T  (M=1..4)