[DSv4][P5-4] MXFP8×MXFP4 Routed Expert grouped GEMM: wgmma triton kernel

[Bignonia7]

Add the experimental SM90 MXFP8 x MXFP4 grouped GEMM implemented in Triton with WGMMA, together with its SM90 PTX/numeric probe and an end-to-end prep/compute benchmark. The kernel is gated and intentionally not registered as a provider until SM90 numeric/PTX/sanitizer checks pass.

Files

• rl_engine/moe/triton_grouped_gemm_wgmma.py — SM90 MXFP8×MXFP4 routed-expert grouped GEMM (Triton WGMMA) with GPU-side scheduling: per-expert element offsets + exact M-tile prefix; a fixed program count traverses the prefix without host-side device reads; recursive GPU scan for large E (no CPU fallback); checked entry points synchronize to report invalid content; validate_contents=False only for caller-guaranteed-valid offsets/E8M0 codes. Not registered as a provider until SM90 numeric/PTX/sanitizer gates pass.
• scripts/probe_p5_triton_wgmma.py — SM90 gate: asserts wgmma.mma_async is present in the compiled PTX and compares GPU output against the serial FP32 oracle. Deliberately separate from provider registration.
• benchmarks/benchmark_triton_grouped_wgmma.py — experimental SM90 prep/compute end-to-end timing; run the probe and sanitizers first; timing success is not a numeric gate; all inputs stay packed with explicit GPU/host latency labels.

[coderabbitai]

Important

Review skipped

Auto reviews are disabled on base/target branches other than the default branch.

Please check the settings in the CodeRabbit UI or the .coderabbit.yaml file in this repository. To trigger a single review, invoke the @coderabbitai review command.

⚙️ Run configuration

Configuration used: defaults

Review profile: CHILL

Plan: Advanced

Run ID: de691f3b-3c97-4a3d-adee-a646d6435e02

You can disable this status message by setting the reviews.review_status to false in the CodeRabbit configuration file.

Use the checkbox below for a quick retry:

• 🔍 Trigger review

---

_{Comment @coderabbitai help to get the list of available commands.}