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

[Bignonia7]

Scope

Adds only csrc/cuda/moe/mxfp8_mxfp4_grouped_gemm_wgmma.cu (forward and backward CuTe WGMMA compute, expert-offset pointer prep, validation, and launch wrappers). Packed E2M1 weights use separate shared-memory staging, then layout-aware writes to the GMMA B tile.

Review status

• This one-file PR is for implementation review. The target dsv4-p5-dev branch does not yet wire this file into setup.py or csrc/ops.cpp, so this PR is not independently buildable or runnable.
• The isolated 4a forward CUTLASS type probe was reported to compile; it does not compile or validate this compute file.
• The compute file has not been compiled with nvcc or run on SM90. Numeric correctness, racecheck, and performance remain unverified.
• The current compute path uses synchronous global-to-shared loads and static tile-prefix scheduling. It does not yet use the planned packed TMA pipeline or persistent grouped adapter.

Checks

• Confirmed the PR diff contains exactly one .cu file.
• git diff --check passes.

Follow-up validation: build with KERNEL_ALIGN_MOE_SM90=1, wire the bindings, run P5 forward/backward comparisons and SM90 racecheck, then benchmark against the strict path.

Summary by CodeRabbit

• New Features
  • Added a high-performance grouped mixture-of-experts matrix multiplication path for H100 GPUs.
  • Supports MXFP8 activations and MXFP4 weights for forward computations.
  • Supports corresponding backward input-gradient calculations with FP32 accumulation.
  • Added validation for supported shapes, data types, memory layout, alignment, and expert ranges.

[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: 926e8fc0-7701-4e8f-b1f5-25909c8cb880

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

✨ Finishing Touches

🧪 Generate unit tests (beta)

• Create a new PR

---

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