Add SM120 support

[Snowfall99]

Summary

Adds an SM120 path so MPK runs on consumer/workstation Blackwell (e.g. RTX Pro 6000, cc120). Previously MPK had no SM120 support:

• linear_layer / linear_with_residual_layer hit assert False, "Unsupported compute capability: 120".
• The shared-memory ladder misclassified SM120 into the >= 90 (Hopper/datacenter-Blackwell) bucket and requested ~207–220KB of dynamic smem, but SM120 only has ~99KB usable per block. This over-request is the root cause of the persistent-kernel scheduler hangs seen on this hardware.

Changes

• runtime_header.h: add an explicit MPK_TARGET_CC == 120 rung (99KB) ahead of the >= 90 check in both shared-memory ladders.
• hopper/utils.cuh: wg_sync now emits portable bar.sync PTX for SM120 instead of falling through to the brkpt trap.
• persistent_kernel.py: rmsnorm_layer bounds the Hopper kernel dispatch to 90 <= cc < 120 (previously any cc >= 90 incorrectly picked the Hopper-specific TMA/WGMMA kernel); linear_layer / linear_with_residual_layer route SM120 to the ampere task set (no TMA/WGMMA required there).

All SM120 gates use exact == 120 rather than >= 120, so a future higher compute capability still falls through to the loud assert False / correct datacenter rung instead of silently reusing the SM120-specific path.

Test results

Verified on real RTX Pro 6000 (Blackwell Workstation, cc120) hardware. Both tests previously failed/hung on this GPU. Compiled with -DMPK_TARGET_CC=120; the megakernel now requests smem size: 98304 (96KB, under the ~99KB limit — the fix that resolves the hang).

test_diamond_fork_join_testmode.py

[MPK INIT] Total tasks: 1546, Total events: 5
diff A = 0.00390625, B = 0.0009765625, C = 0.0009765625, D = 0.001953125
PASSED: diamond_fork_join test_mode produces correct output

test_qwen3_mlp_testmode.py (three stages, bf16 tolerance)

Test: Gate+Up linear only (dense Qwen3 MLP)
  Max absolute diff: 0.007812
  PASSED: gate+up linear layer produces correct output

Test: Gate+Up linear + SiLU-Mul
  Max absolute diff: 0.015625
  PASSED: gate+up + silu_mul produces correct output

Test: Full MLP pipeline (gate+up → silu_mul → down+residual)
  B=8, hidden=4096, intermediate=2048
  Max absolute diff: 0.015625
  PASSED: full MLP pipeline produces correct output

Out of scope

No ampere/generic fallback kernel exists to route these to (would need new kernel work): MoE layers, FP8 group GEMM, split-KV attention — all Hopper/datacenter-Blackwell-only today.

[Snowfall99]

@xinhaoc Would you mind reviewing this when you have a chance?