[GH-ISSUE #235] [Question] Contiguous Grouped Gemm for CUDA Graph #135

New issue

Open

opened 2026-10-26 21:13:23 +00:00 by bsakaguchi · 2 comments

bsakaguchi commented

2026-10-26 21:13:23 +00:00

Owner

Originally created by @yuhyao on GitHub (Nov 28, 2025).
Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/235

Hi DeepSeek team,

I noticed that DeepEP provides an optional num_worst_tokens argument for CUDA Graph usage. Based on that, I assume contiguous Grouped GEMM is also intended to support CUDA Graph. However, it seems there is no input parameter to supply a valid m for the actual workload. While we can set m_indices to -1 for padded tokens, performance degrades significantly because is_computation_valid only skips the compute blocks, whereas the TMA loads are still issued for those padded regions.

Did I misunderstand something here, or is there a plan to address this limitation?

Thanks!

Originally created by @yuhyao on GitHub (Nov 28, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/235 Hi DeepSeek team, I noticed that DeepEP provides an optional `num_worst_tokens` argument for CUDA Graph usage. Based on that, I assume contiguous Grouped GEMM is also intended to support CUDA Graph. However, it seems there is no input parameter to supply a valid m for the actual workload. While we can set `m_indices` to -1 for padded tokens, performance degrades significantly because `is_computation_valid` only skips the compute blocks, whereas the TMA loads are still issued for those padded regions. Did I misunderstand something here, or is there a plan to address this limitation? Thanks!

bsakaguchi commented

2026-10-26 21:13:24 +00:00

Author

Owner

@LyricZhao commented on GitHub (Dec 5, 2025):

num_worst_tokens is supported by the community (maybe with their own GEMM kernels). In DeepSeek, we never use num_worst_tokens (prefill without CUDA graph).

We are releasing EPv2 soon, normal/low-latency layout will be merged into a single contiguous layout then. DeepGEMM will also provide the masked/non-masked API.

@LyricZhao commented on GitHub (Dec 5, 2025): `num_worst_tokens` is supported by the community (maybe with their own GEMM kernels). In DeepSeek, we never use `num_worst_tokens` (prefill without CUDA graph). We are releasing EPv2 soon, normal/low-latency layout will be merged into a single contiguous layout then. DeepGEMM will also provide the masked/non-masked API.

bsakaguchi commented

2026-10-26 21:13:24 +00:00

Author

Owner

@yuhyao commented on GitHub (Dec 8, 2025):

num_worst_tokens is supported by the community (maybe with their own GEMM kernels). In DeepSeek, we never use num_worst_tokens (prefill without CUDA graph).

We are releasing EPv2 soon, normal/low-latency layout will be merged into a single contiguous layout then. DeepGEMM will also provide the masked/non-masked API.

Thanks for the reply. It’s exciting to hear that DeepEP v2 is coming soon!
Please allow me to double-check my understanding: the low-latency mode in DeepEP v2 will also use contiguous layout, so that the input and output buffers for grouped GEMM consumes less memory compared to the current low-latency layout when number of experts per rank is larger than num_experts_per_tok?

@yuhyao commented on GitHub (Dec 8, 2025): > `num_worst_tokens` is supported by the community (maybe with their own GEMM kernels). In DeepSeek, we never use `num_worst_tokens` (prefill without CUDA graph). > > We are releasing EPv2 soon, normal/low-latency layout will be merged into a single contiguous layout then. DeepGEMM will also provide the masked/non-masked API. Thanks for the reply. It’s exciting to hear that DeepEP v2 is coming soon! Please allow me to double-check my understanding: the low-latency mode in DeepEP v2 will also use contiguous layout, so that the input and output buffers for grouped GEMM consumes less memory compared to the current low-latency layout when number of experts per rank is larger than `num_experts_per_tok`?

bsakaguchi referenced this issue

2026-10-26 21:15:28 +00:00

[PR #135] [MERGED] Fix import-time error: "undefined symbol: _ZN3c104cuda20getCurrentCUDAStreamEa" #261