|
| 1 | +"""GB10/sm_121 F0 tensor-core precompute probe (§26). |
| 2 | +
|
| 3 | +Builds the NEW chunk_precompute_fwd_cuda_prim (cb + summary_states T.gemm) |
| 4 | +at prod cfg and: |
| 5 | + 1. dumps generated CUDA, greps mma.sync (REAL sm_120 MMA, not relabeled serial). |
| 6 | + 2. times new F0 per-call (warm median + min). |
| 7 | + 3. parity of cb + summary_states + dA_cumsum vs: |
| 8 | + (a) the SERIAL F0 reference (eager precompute math), |
| 9 | + (b) the SERIAL Metal F0 prim built on CUDA target == NO, the serial prim is |
| 10 | + Metal-only; instead we compare against the eager-precompute fp32 reference |
| 11 | + AND re-build the *serial* prim source for byte check separately. |
| 12 | + over ALL elements (max|abs|), gate fp16 < 5e-4. |
| 13 | + 4. F0->F1->F2 chain timing with the NEW F0. |
| 14 | +
|
| 15 | +RULE #1: every kernel goes through its ONE builder; failures SURFACED. |
| 16 | +""" |
| 17 | +import sys, time, traceback, re |
| 18 | +try: |
| 19 | + from cppmega_mlx._gb10_nvrtc_env import ensure_nvrtc_builtins_path |
| 20 | + ensure_nvrtc_builtins_path() |
| 21 | +except Exception as _e: |
| 22 | + print(f"[nvrtc-env] unavailable: {_e!r}") |
| 23 | + |
| 24 | +import torch |
| 25 | +from einops import rearrange, repeat |
| 26 | +DEV = "cuda" |
| 27 | + |
| 28 | + |
| 29 | +def eager_precompute_f0(C, Bmat, x, A, dt, chunk_size): |
| 30 | + """SERIAL F0 reference (fp32): cb, dA_cumsum, summary_states. |
| 31 | + Mirrors chunk_precompute math exactly (the contract anchor).""" |
| 32 | + batch, seqlen, ngroups, dstate = C.shape |
| 33 | + _, _, nheads, headdim = x.shape |
| 34 | + nchunks = seqlen // chunk_size |
| 35 | + h = nheads // ngroups |
| 36 | + Cf, Bf, xf, Af, dtf = C.float(), Bmat.float(), x.float(), A.float(), dt.float() |
| 37 | + a = Af.view(1, 1, nheads) * dtf |
| 38 | + a_c = rearrange(a, "b (c l) hh -> b hh c l", c=nchunks) |
| 39 | + dA_cumsum = torch.cumsum(a_c, dim=-1) # (b,h,c,l) |
| 40 | + Cc = rearrange(Cf, "b (c l) g n -> b c l g n", c=nchunks) |
| 41 | + Bc = rearrange(Bf, "b (c s) g n -> b c s g n", c=nchunks) |
| 42 | + cb = torch.einsum("bclgn,bcsgn->bcgls", Cc, Bc) # (b,c,g,l,s) |
| 43 | + Bexp = Bc.repeat_interleave(h, dim=3) # (b,c,s,hh,n) |
| 44 | + xexp = rearrange(xf, "b (c s) hh p -> b c s hh p", c=nchunks) |
| 45 | + dtc = rearrange(dtf, "b (c s) hh -> b hh c s", c=nchunks) |
| 46 | + decay_states = torch.exp(dA_cumsum[:, :, :, -1:] - dA_cumsum) # (b,h,c,s) |
| 47 | + # summary_states[b,c,hh,p,n] = sum_s decay*dt * x * B |
| 48 | + states = torch.einsum("bhcs,bhcs,bcshp,bcshn->bchpn", |
| 49 | + decay_states, dtc, xexp, Bexp) |
| 50 | + return (cb.float().contiguous(), |
| 51 | + dA_cumsum.float().contiguous(), |
| 52 | + states.float().contiguous()) |
| 53 | + |
| 54 | + |
| 55 | +def _time(fn, n=30): |
| 56 | + for _ in range(5): |
| 57 | + fn() |
| 58 | + torch.cuda.synchronize() |
| 59 | + ts = [] |
| 60 | + for _ in range(n): |
| 61 | + torch.cuda.synchronize() |
| 62 | + t0 = time.perf_counter() |
| 63 | + fn() |
| 64 | + torch.cuda.synchronize() |
| 65 | + ts.append((time.perf_counter() - t0) * 1e3) |
| 66 | + ts.sort() |
| 67 | + return ts[len(ts)//2], ts[0] |
| 68 | + |
| 69 | + |
| 70 | +def run_cfg(batch, seqlen, chunk, ngroups, nheads, headdim, dstate): |
| 71 | + from cppmega_mlx.nn._tilelang.mamba3_chunked_precompute_core import ( |
| 72 | + build_chunk_precompute_metal, build_inter_chunk_recur_metal, |
| 73 | + chunk_precompute_fwd_cuda_prim, chunk_precompute_fwd_metal_prim, |
| 74 | + ) |
| 75 | + from cppmega_mlx.nn._tilelang.mamba3_chunked_scan_core import ( |
| 76 | + build_chunk_scan_combine_metal, |
| 77 | + ) |
| 78 | + nchunks = seqlen // chunk |
| 79 | + print(f"\n=== cfg b={batch} S={seqlen} c={chunk} g={ngroups} H={nheads} " |
| 80 | + f"P={headdim} N={dstate} nchunks={nchunks} ===") |
| 81 | + |
| 82 | + torch.manual_seed(0) |
| 83 | + C = (torch.randn(batch, seqlen, ngroups, dstate, device=DEV) * 0.1).half() |
| 84 | + Bmat = (torch.randn(batch, seqlen, ngroups, dstate, device=DEV) * 0.1).half() |
| 85 | + x = (torch.randn(batch, seqlen, nheads, headdim, device=DEV) * 0.1).half() |
| 86 | + A = -torch.rand(nheads, device=DEV).half() |
| 87 | + dt = (torch.rand(batch, seqlen, nheads, device=DEV) * 0.05).half() |
| 88 | + D = torch.randn(nheads, device=DEV).half() |
| 89 | + h0 = (torch.randn(batch, nheads, headdim, dstate, device=DEV) * 0.1).float() |
| 90 | + |
| 91 | + # ---- SERIAL F0 reference (fp32 einsum on the SAME fp16 inputs) ---- |
| 92 | + cb_ref, dac_ref, ss_ref = eager_precompute_f0( |
| 93 | + C, Bmat, x, A, dt, chunk) |
| 94 | + |
| 95 | + # =================== NEW CUDA F0 (tensor-core) =================== |
| 96 | + try: |
| 97 | + k0 = build_chunk_precompute_metal( |
| 98 | + batch, seqlen, chunk, ngroups, nheads, headdim, dstate, target="cuda") |
| 99 | + except Exception: |
| 100 | + print("[F0-CUDA] COMPILE FAILED:") |
| 101 | + traceback.print_exc() |
| 102 | + return False |
| 103 | + print("[F0-CUDA] COMPILE ok") |
| 104 | + |
| 105 | + # ---- dump CUDA source, grep mma.sync ---- |
| 106 | + try: |
| 107 | + src = k0.get_kernel_source() |
| 108 | + # tl::mma_sync is the codegen wrapper for the real PTX |
| 109 | + # mma.sync.aligned.m16n8k16 (CuTe SM80_16x8x16_F32F16F16F32_TN). |
| 110 | + n_mma = src.count("tl::mma_sync") + src.count("mma.sync") |
| 111 | + n_ldm = src.count("ptx_ldmatrix") |
| 112 | + n_serial_dot = len(re.findall(r"for \(int n = 0; n < %d" % dstate, src)) |
| 113 | + print(f"[F0-CUDA] codegen: tl::mma_sync x{n_mma} ldmatrix x{n_ldm} " |
| 114 | + f"leftover_serial_N_dot x{n_serial_dot} len={len(src)}B") |
| 115 | + with open("/tmp/f0_cuda_src.cu", "w") as f: |
| 116 | + f.write(src) |
| 117 | + print("[F0-CUDA] source -> /tmp/f0_cuda_src.cu") |
| 118 | + except Exception: |
| 119 | + print("[F0-CUDA] source dump FAILED:") |
| 120 | + traceback.print_exc() |
| 121 | + n_mma = -1 |
| 122 | + |
| 123 | + cb = torch.zeros(batch, nchunks, ngroups, chunk, chunk, device=DEV, dtype=torch.float16) |
| 124 | + dA_cumsum = torch.zeros(batch, nheads, nchunks, chunk, device=DEV, dtype=torch.float16) |
| 125 | + summary_states = torch.zeros(batch, nchunks, nheads, headdim, dstate, device=DEV, dtype=torch.float32) |
| 126 | + |
| 127 | + def _run_f0(): |
| 128 | + k0(x.contiguous(), Bmat.contiguous(), C.contiguous(), A.contiguous(), |
| 129 | + dt.contiguous(), cb, dA_cumsum, summary_states) |
| 130 | + try: |
| 131 | + _run_f0(); torch.cuda.synchronize() |
| 132 | + except Exception: |
| 133 | + print("[F0-CUDA] RUN FAILED:") |
| 134 | + traceback.print_exc() |
| 135 | + return False |
| 136 | + print("[F0-CUDA] RUN ok (LAUNCHES at N=%d)" % dstate) |
| 137 | + |
| 138 | + # ---- PARITY (all elements) of new F0 vs serial F0 reference ---- |
| 139 | + nan0 = bool(torch.isnan(cb).any()) or bool(torch.isnan(summary_states).any()) \ |
| 140 | + or bool(torch.isnan(dA_cumsum).any()) |
| 141 | + cb_err = float((cb.float().cpu() - cb_ref.cpu()).abs().max()) |
| 142 | + dac_err = float((dA_cumsum.float().cpu() - dac_ref.cpu()).abs().max()) |
| 143 | + ss_err = float((summary_states.float().cpu() - ss_ref.cpu()).abs().max()) |
| 144 | + print(f"[F0-CUDA vs SERIAL-ref] cb max|abs|={cb_err:.3e} " |
| 145 | + f"dA_cumsum={dac_err:.3e} summary_states={ss_err:.3e} NaN={nan0}") |
| 146 | + |
| 147 | + f0_med, f0_min = _time(_run_f0) |
| 148 | + print(f"[F0-CUDA] TIMING median={f0_med:.4f} ms/call min={f0_min:.4f} ms") |
| 149 | + |
| 150 | + # =================== SERIAL METAL F0 prim built on CUDA (the 24.57ms baseline) =================== |
| 151 | + # Build the *serial* metal prim explicitly on the CUDA target so we get the |
| 152 | + # apples-to-apples 24.57ms serial F0 on the SAME gb10 device. |
| 153 | + f0_serial_med = None |
| 154 | + cb_s_err = ss_s_err = None |
| 155 | + try: |
| 156 | + import tilelang |
| 157 | + from cppmega_mlx.nn._tilelang.mamba3_chunked_scan_core import _resolve_chunked_compile_target |
| 158 | + prim_s = chunk_precompute_fwd_metal_prim( |
| 159 | + batch, seqlen, chunk, ngroups, nheads, headdim, dstate) |
| 160 | + k0s = tilelang.compile(prim_s, out_idx=[5, 6, 7], |
| 161 | + target=_resolve_chunked_compile_target("cuda")) |
| 162 | + cb_s = torch.zeros_like(cb); dac_s = torch.zeros_like(dA_cumsum) |
| 163 | + ss_s = torch.zeros_like(summary_states) |
| 164 | + def _run_f0s(): |
| 165 | + k0s(x.contiguous(), Bmat.contiguous(), C.contiguous(), A.contiguous(), |
| 166 | + dt.contiguous(), cb_s, dac_s, ss_s) |
| 167 | + _run_f0s(); torch.cuda.synchronize() |
| 168 | + print("[F0-SERIAL on CUDA] RUN ok") |
| 169 | + f0_serial_med, f0_serial_min = _time(_run_f0s) |
| 170 | + print(f"[F0-SERIAL on CUDA] TIMING median={f0_serial_med:.4f} ms min={f0_serial_min:.4f} ms") |
| 171 | + # parity new-CUDA vs serial-prim (the tightest gate: same algorithm, fp16) |
| 172 | + cb_s_err = float((cb.float() - cb_s.float()).abs().max()) |
| 173 | + ss_s_err = float((summary_states.float() - ss_s.float()).abs().max()) |
| 174 | + dac_s_err = float((dA_cumsum.float() - dac_s.float()).abs().max()) |
| 175 | + print(f"[F0-CUDA vs F0-SERIAL prim] cb={cb_s_err:.3e} " |
| 176 | + f"summary_states={ss_s_err:.3e} dA_cumsum={dac_s_err:.3e}") |
| 177 | + except Exception: |
| 178 | + print("[F0-SERIAL on CUDA] build/run FAILED (serial prim may be Metal-only):") |
| 179 | + traceback.print_exc() |
| 180 | + |
| 181 | + # =================== F0->F1->F2 chain timing with NEW F0 =================== |
| 182 | + try: |
| 183 | + k1 = build_inter_chunk_recur_metal( |
| 184 | + batch, seqlen, chunk, ngroups, nheads, headdim, dstate, target="cuda") |
| 185 | + k2 = build_chunk_scan_combine_metal( |
| 186 | + batch, seqlen, chunk, ngroups, nheads, headdim, dstate, target="cuda") |
| 187 | + except Exception: |
| 188 | + print("[CHAIN] F1/F2 COMPILE FAILED:") |
| 189 | + traceback.print_exc() |
| 190 | + return False |
| 191 | + prev_states = torch.zeros(batch, nchunks, nheads, headdim, dstate, device=DEV, dtype=torch.float32) |
| 192 | + final_state = torch.zeros(batch, nheads, headdim, dstate, device=DEV, dtype=torch.float32) |
| 193 | + out_u = torch.zeros(batch, seqlen, nheads, headdim, device=DEV, dtype=torch.float16) |
| 194 | + dt_k = rearrange(dt, "b (c s) hh -> b hh c s", c=nchunks).contiguous() |
| 195 | + |
| 196 | + def _run_f1(): |
| 197 | + k1(summary_states.contiguous(), dA_cumsum.contiguous(), h0.contiguous(), |
| 198 | + prev_states, final_state) |
| 199 | + def _run_f2(): |
| 200 | + k2(cb.contiguous(), x.contiguous(), dt_k.contiguous(), dA_cumsum.contiguous(), |
| 201 | + C.contiguous(), prev_states.contiguous(), D.contiguous(), out_u) |
| 202 | + try: |
| 203 | + _run_f1(); torch.cuda.synchronize() |
| 204 | + _run_f2(); torch.cuda.synchronize() |
| 205 | + except Exception: |
| 206 | + print("[CHAIN] F1/F2 RUN FAILED:") |
| 207 | + traceback.print_exc() |
| 208 | + return False |
| 209 | + f1_med, f1_min = _time(_run_f1) |
| 210 | + f2_med, f2_min = _time(_run_f2) |
| 211 | + chain = f0_med + f1_med + f2_med |
| 212 | + print(f"[CHAIN] F0={f0_med:.4f} F1={f1_med:.4f} F2={f2_med:.4f} " |
| 213 | + f"SUM={chain:.4f} ms (medians)") |
| 214 | + |
| 215 | + gate = 5e-4 |
| 216 | + parity_ok = (cb_err < gate and ss_err < gate and dac_err < 1e-2 and not nan0) |
| 217 | + # NOTE: dA_cumsum is fp16-stored so its gate is looser (cumsum magnitude), |
| 218 | + # report it but key the gate on cb/summary_states (the GEMM outputs). |
| 219 | + print(f"\n[SUMMARY] new_F0={f0_med:.4f} ms " |
| 220 | + f"serial_F0={'%.4f'%f0_serial_med if f0_serial_med else 'NA'} ms " |
| 221 | + f"baseline_24.57ms chain={chain:.4f} ms") |
| 222 | + print(f"[SUMMARY] parity cb={cb_err:.2e} ss={ss_err:.2e} -> " |
| 223 | + f"{'PASS' if parity_ok else 'FAIL'} (gate {gate})") |
| 224 | + print(f"[CSV] new_f0={f0_med:.4f} new_f0_min={f0_min:.4f} " |
| 225 | + f"serial_f0={f0_serial_med if f0_serial_med else -1:.4f} " |
| 226 | + f"f1={f1_med:.4f} f2={f2_med:.4f} chain={chain:.4f} " |
| 227 | + f"cb_err={cb_err:.3e} ss_err={ss_err:.3e} dac_err={dac_err:.3e} " |
| 228 | + f"cb_vs_serialprim={cb_s_err if cb_s_err else -1:.3e} " |
| 229 | + f"ss_vs_serialprim={ss_s_err if ss_s_err else -1:.3e} " |
| 230 | + f"mma_sync={n_mma} parity={'PASS' if parity_ok else 'FAIL'}") |
| 231 | + return parity_ok |
| 232 | + |
| 233 | + |
| 234 | +def main(): |
| 235 | + print("=== F0 tensor-core precompute probe (gb10 sm_121) ===") |
| 236 | + print("torch", torch.__version__, "cuda", torch.cuda.is_available(), |
| 237 | + torch.cuda.get_device_name(0) if torch.cuda.is_available() else "NONE") |
| 238 | + if not torch.cuda.is_available(): |
| 239 | + print("NO CUDA"); sys.exit(2) |
| 240 | + cfg = dict(batch=1, seqlen=4096, chunk=64, ngroups=8, nheads=112, |
| 241 | + headdim=64, dstate=64) |
| 242 | + ok = False |
| 243 | + try: |
| 244 | + ok = run_cfg(**cfg) |
| 245 | + except Exception: |
| 246 | + print("CFG CRASHED:"); traceback.print_exc() |
| 247 | + print("\nF0 PARITY PASS" if ok else "\nF0 PARITY/RUN FAILURE") |
| 248 | + sys.exit(0 if ok else 1) |
| 249 | + |
| 250 | + |
| 251 | +if __name__ == "__main__": |
| 252 | + main() |
0 commit comments