[blackwell] sm100: QKV projection + RoPE Triton kernel

[Andrewxu313]

Parent: #204 — Blackwell Phase 2 kernel port
Branch: tairan/blackwell-02-05-qkv-rope-triton
Spec: blackwell-kernel-port-v1.md § Sub-task 5

What

Write batchgen_kernels/triton/qkv_proj_rope.py:

• tl.dot BF16 matmul [M,K] x [K,N]
• Fused split epilogue: writes directly to Q_out [M,q_size], K_out [M,kv_size], V_out [M,kv_size]
• Optional fused RoPE rotation on Q and K tiles (HAS_ROPE: tl.constexpr)

F.linear is NOT acceptable — it returns a single tensor, losing the fused split + RoPE.

Wire into batchgen/attention/fused_kernels/qkv_wgmma.py for arch == 'sm100'.

Acceptance

• Q/K/V split: max_abs_err < 1e-2 vs F.linear + split reference
• RoPE: Q/K match apply_rotary_pos_emb reference within 1e-2
• MLA prefill E2E on B200: no shape errors