[GH-ISSUE #356] Does Mega MoE support RTX 50-series GPUs, and can it work in a PCIe-only setup without NVLink? #170

New issue

Open

opened 2026-10-26 21:14:06 +00:00 by bsakaguchi · 1 comment

bsakaguchi commented

2026-10-26 21:14:06 +00:00

Owner

Originally created by @wukong1992 on GitHub (Jun 10, 2026).
Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/356

Does Mega MoE support RTX 50-series GPUs, and can it work in a PCIe-only setup without NVLink?

Originally created by @wukong1992 on GitHub (Jun 10, 2026). Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/356 Does Mega MoE support RTX 50-series GPUs, and can it work in a PCIe-only setup without NVLink?

bsakaguchi commented

2026-10-26 21:14:09 +00:00

Author

Owner

@aganhui commented on GitHub (Sep 11, 2026):

Short answer as of nv_dev today: no, Mega MoE does not run on RTX 50-series (SM120) yet — it is SM100-only (B200).

The dispatch in csrc/apis/mega.hpp hard-requires arch_major == 10 (see the fp8_fp4_mega_moe / fp4_fp4_mega_moe / bf16_mega_moe entry points), and the underlying kernels are sm100_fp8_fp4_mega_moe.cuh / sm100_bf16_mega_moe.cuh — they use SM100's UMMA/tcgen05 machinery and tensor-memory features that consumer Blackwell does not expose. Mega MoE landed for SM100 first (April, #304); the SM90 port is still in flight (#323, with #383 merged earlier this month); an SM120 port has not been started.

On the NVLink/PCIe question: NVLink should not matter for DeepGEMM itself. DeepGEMM provides the per-GPU MoE compute kernels; the inter-GPU communication lives in the serving stack around it (e.g., DeepEP, which is where NVLink vs PCIe matters). So whenever SM120 Mega MoE arrives, a PCIe-only setup should be fine at the DeepGEMM level — the effective constraint is memory capacity and bandwidth of the card, not the interconnect.

What SM120 users do get today (nv_dev): the fp8/fp4 GEMM family (normal + m-grouped contiguous/masked + k-grouped), bf16 GEMM, and the fp8/fp4 MQA-logits kernels (dense + paged). Note if you are running on a 5090: the fp8 MQA-logits path has a swizzle bug for head_dim 32/64 that is being fixed in #379/#434, so builds from older nv_dev revisions may return wrong logits for those head dims.

@aganhui commented on GitHub (Sep 11, 2026): Short answer as of nv_dev today: **no, Mega MoE does not run on RTX 50-series (SM120) yet — it is SM100-only (B200).** The dispatch in `csrc/apis/mega.hpp` hard-requires `arch_major == 10` (see the `fp8_fp4_mega_moe` / `fp4_fp4_mega_moe` / `bf16_mega_moe` entry points), and the underlying kernels are `sm100_fp8_fp4_mega_moe.cuh` / `sm100_bf16_mega_moe.cuh` — they use SM100's UMMA/tcgen05 machinery and tensor-memory features that consumer Blackwell does not expose. Mega MoE landed for SM100 first (April, #304); the SM90 port is still in flight (#323, with #383 merged earlier this month); an SM120 port has not been started. On the NVLink/PCIe question: NVLink should not matter for DeepGEMM itself. DeepGEMM provides the per-GPU MoE compute kernels; the inter-GPU communication lives in the serving stack around it (e.g., DeepEP, which is where NVLink vs PCIe matters). So whenever SM120 Mega MoE arrives, a PCIe-only setup should be fine at the DeepGEMM level — the effective constraint is memory capacity and bandwidth of the card, not the interconnect. What SM120 users do get today (nv_dev): the fp8/fp4 GEMM family (normal + m-grouped contiguous/masked + k-grouped), bf16 GEMM, and the fp8/fp4 MQA-logits kernels (dense + paged). Note if you are running on a 5090: the fp8 MQA-logits path has a swizzle bug for head_dim 32/64 that is being fixed in #379/#434, so builds from older nv_dev revisions may return wrong logits for those head dims.

No labels

pull-request

No milestone

No project

No assignees

1 participant

Notifications

Due date

The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set.

Reference

deepseek-ai/DeepGEMM#170

No description provided.

Rows
Columns