---
id: pr-deepgemm-304
repo: deepseek-ai/DeepGEMM
pr: 304
title: '[Public release 26/04] Introducing Mega MoE, FP4 Indexer and other features/fixes'
author: LyricZhao
date: '2026-04-17'
url: https://github.com/deepseek-ai/DeepGEMM/pull/304
source_category: upstream-code
architectures:
- sm100
- sm90
tags:
- gemm
- moe
- fused-kernel
- nvfp4
- fp8
- fp4
- block-scale
- tcgen05
- tmem
techniques:
- kernel-fusion
- fine-grained-quantization
- communication-overlap
hardware_features:
- tcgen05
- tmem
- nvfp4
- fp4
- fp8
- block-scale
kernel_types:
- gemm
- moe
- grouped-gemm
- fused-kernel
languages:
- cuda-cpp
- ptx
captured_at: '2026-04-27'
status: merged
merge_sha: 7f2a703ed51ac1f7af07f5e1453b2d3267d37d50
inclusion_reason: Major DeepGEMM public release with Mega MoE FP8xFP4 fused mega-kernel and FP4 Indexer; substantial SM100 kernel feature surface
changed_paths:
- CMakeLists.txt
- README.md
- csrc/apis/attention.hpp
- csrc/apis/einsum.hpp
- csrc/apis/gemm.hpp
- csrc/apis/hyperconnection.hpp
- csrc/apis/layout.hpp
- csrc/apis/mega.hpp
- csrc/apis/runtime.hpp
- csrc/indexing/main.cu
- csrc/jit/cache.hpp
- csrc/jit/compiler.hpp
- csrc/jit/device_runtime.hpp
- csrc/jit/handle.hpp
- csrc/jit/include_parser.hpp
- csrc/jit/kernel_runtime.hpp
- csrc/jit_kernels/heuristics/common.hpp
- csrc/jit_kernels/heuristics/config.hpp
- csrc/jit_kernels/heuristics/mega_moe.hpp
- csrc/jit_kernels/heuristics/runtime.hpp
- csrc/jit_kernels/heuristics/sm100.hpp
- csrc/jit_kernels/heuristics/sm90.hpp
- csrc/jit_kernels/heuristics/utils.hpp
- csrc/jit_kernels/impls/epilogue.hpp
- csrc/jit_kernels/impls/runtime_utils.hpp
- csrc/jit_kernels/impls/sm100_bf16_gemm.hpp
- csrc/jit_kernels/impls/sm100_bmk_bnk_mn.hpp
- csrc/jit_kernels/impls/sm100_fp8_fp4_gemm_1d1d.hpp
- csrc/jit_kernels/impls/sm100_fp8_fp4_mega_moe.hpp
- csrc/jit_kernels/impls/sm100_tf32_hc_prenorm_gemm.hpp
- csrc/jit_kernels/impls/sm90_bf16_gemm.hpp
- csrc/jit_kernels/impls/sm90_bmk_bnk_mn.hpp
- csrc/jit_kernels/impls/sm90_fp8_gemm_1d1d.hpp
- csrc/jit_kernels/impls/sm90_fp8_gemm_1d2d.hpp
- csrc/jit_kernels/impls/sm90_tf32_hc_prenorm_gemm.hpp
- csrc/jit_kernels/impls/smxx_clean_logits.hpp
- csrc/jit_kernels/impls/smxx_cublaslt.hpp
- csrc/jit_kernels/impls/smxx_fp8_fp4_mqa_logits.hpp
- csrc/jit_kernels/impls/smxx_fp8_fp4_paged_mqa_logits.hpp
- csrc/jit_kernels/impls/smxx_layout.hpp
- csrc/python_api.cpp
- csrc/utils/exception.hpp
- csrc/utils/hash.hpp
- csrc/utils/layout.hpp
- csrc/utils/math.hpp
- csrc/utils/system.hpp
- deep_gemm/__init__.py
- deep_gemm/include/deep_gemm/comm/barrier.cuh
- deep_gemm/include/deep_gemm/common/compile.cuh
- deep_gemm/include/deep_gemm/common/cute_tie.cuh
- deep_gemm/include/deep_gemm/common/exception.cuh
- deep_gemm/include/deep_gemm/common/math.cuh
- deep_gemm/include/deep_gemm/common/tma_copy.cuh
- deep_gemm/include/deep_gemm/common/types.cuh
- deep_gemm/include/deep_gemm/common/utils.cuh
- deep_gemm/include/deep_gemm/epilogue/sm100_store_cd.cuh
- deep_gemm/include/deep_gemm/epilogue/sm100_store_cd_swap_ab.cuh
- deep_gemm/include/deep_gemm/epilogue/transform.cuh
- deep_gemm/include/deep_gemm/impls/sm100_bf16_gemm.cuh
- deep_gemm/include/deep_gemm/impls/sm100_bmk_bnk_mn.cuh
- deep_gemm/include/deep_gemm/impls/sm100_fp4_mqa_logits.cuh
- deep_gemm/include/deep_gemm/impls/sm100_fp4_paged_mqa_logits.cuh
- deep_gemm/include/deep_gemm/impls/sm100_fp8_fp4_gemm_1d1d.cuh
- deep_gemm/include/deep_gemm/impls/sm100_fp8_fp4_mega_moe.cuh
- deep_gemm/include/deep_gemm/impls/sm100_fp8_mqa_logits.cuh
- deep_gemm/include/deep_gemm/impls/sm100_fp8_paged_mqa_logits.cuh
- deep_gemm/include/deep_gemm/impls/sm100_tf32_hc_prenorm_gemm.cuh
- deep_gemm/include/deep_gemm/impls/sm90_bf16_gemm.cuh
- deep_gemm/include/deep_gemm/impls/sm90_bmk_bnk_mn.cuh
- deep_gemm/include/deep_gemm/impls/sm90_fp8_gemm_1d1d.cuh
- deep_gemm/include/deep_gemm/impls/sm90_fp8_gemm_1d2d.cuh
- deep_gemm/include/deep_gemm/impls/sm90_fp8_mqa_logits.cuh
- deep_gemm/include/deep_gemm/impls/sm90_fp8_paged_mqa_logits.cuh
- deep_gemm/include/deep_gemm/impls/sm90_tf32_hc_prenorm_gemm.cuh
- deep_gemm/include/deep_gemm/impls/smxx_clean_logits.cuh
- deep_gemm/include/deep_gemm/impls/smxx_layout.cuh
- deep_gemm/include/deep_gemm/layout/mega_moe.cuh
- deep_gemm/include/deep_gemm/layout/sym_buffer.cuh
- deep_gemm/include/deep_gemm/mma/sm100.cuh
- deep_gemm/include/deep_gemm/mma/sm90.cuh
- deep_gemm/include/deep_gemm/ptx/ld_st.cuh
- deep_gemm/include/deep_gemm/ptx/tcgen05.cuh
- deep_gemm/include/deep_gemm/ptx/tma.cuh
- deep_gemm/include/deep_gemm/ptx/utils.cuh
- deep_gemm/include/deep_gemm/ptx/wgmma.cuh
- deep_gemm/include/deep_gemm/scheduler/gemm.cuh
- deep_gemm/include/deep_gemm/scheduler/mega_moe.cuh
- deep_gemm/include/deep_gemm/scheduler/paged_mqa_logits.cuh
- deep_gemm/legacy/a_fused_m_grouped_gemm.py
- deep_gemm/legacy/b_fused_k_grouped_gemm.py
- deep_gemm/mega/__init__.py
- deep_gemm/testing/bench.py
- deep_gemm/utils/__init__.py
- deep_gemm/utils/dist.py
- deep_gemm/utils/layout.py
- deep_gemm/utils/math.py
- setup.py
- tests/generators.py
- tests/test_attention.py
- tests/test_bf16.py
- tests/test_einsum.py
- tests/test_fp8_fp4.py
- tests/test_layout.py
- tests/test_lazy_init.py
- tests/test_mega_moe.py
- tests/test_sanitizer.py
- third-party/tilelang_ops/__init__.py
- third-party/tilelang_ops/swiglu_apply_weight_to_fp8.py
- third-party/tilelang_ops/utils.py
artifact_dir: evidence/pull-bundles/deepgemm/gh-304
head_ref_oid: 211d2678d9d178de7a77fbfefe7310e0f4f8821c
---
# [Public release 26/04] Introducing Mega MoE, FP4 Indexer and other features/fixes

- Repository: `deepseek-ai/DeepGEMM`
- PR: [304](https://github.com/deepseek-ai/DeepGEMM/pull/304)
- Decision: `Major DeepGEMM public release with Mega MoE FP8xFP4 fused mega-kernel and FP4 Indexer; substantial SM100 kernel feature surface`
- Reason: Major DeepGEMM public release with Mega MoE FP8xFP4 fused mega-kernel and FP4 Indexer; substantial SM100 kernel feature surface

## Changed Paths

- `CMakeLists.txt`
- `README.md`
- `csrc/apis/attention.hpp`
- `csrc/apis/einsum.hpp`
- `csrc/apis/gemm.hpp`
- `csrc/apis/hyperconnection.hpp`
- `csrc/apis/layout.hpp`
- `csrc/apis/mega.hpp`
- `csrc/apis/runtime.hpp`
- `csrc/indexing/main.cu`
- `csrc/jit/cache.hpp`
- `csrc/jit/compiler.hpp`
- `csrc/jit/device_runtime.hpp`
- `csrc/jit/handle.hpp`
- `csrc/jit/include_parser.hpp`
- `csrc/jit/kernel_runtime.hpp`
- `csrc/jit_kernels/heuristics/common.hpp`
- `csrc/jit_kernels/heuristics/config.hpp`
- `csrc/jit_kernels/heuristics/mega_moe.hpp`
- `csrc/jit_kernels/heuristics/runtime.hpp`
- `csrc/jit_kernels/heuristics/sm100.hpp`
- `csrc/jit_kernels/heuristics/sm90.hpp`
- `csrc/jit_kernels/heuristics/utils.hpp`
- `csrc/jit_kernels/impls/epilogue.hpp`
- `csrc/jit_kernels/impls/runtime_utils.hpp`
- `csrc/jit_kernels/impls/sm100_bf16_gemm.hpp`
- `csrc/jit_kernels/impls/sm100_bmk_bnk_mn.hpp`
- `csrc/jit_kernels/impls/sm100_fp8_fp4_gemm_1d1d.hpp`
- `csrc/jit_kernels/impls/sm100_fp8_fp4_mega_moe.hpp`
- `csrc/jit_kernels/impls/sm100_tf32_hc_prenorm_gemm.hpp`
- `csrc/jit_kernels/impls/sm90_bf16_gemm.hpp`
- `csrc/jit_kernels/impls/sm90_bmk_bnk_mn.hpp`
- `csrc/jit_kernels/impls/sm90_fp8_gemm_1d1d.hpp`
- `csrc/jit_kernels/impls/sm90_fp8_gemm_1d2d.hpp`
- `csrc/jit_kernels/impls/sm90_tf32_hc_prenorm_gemm.hpp`
- `csrc/jit_kernels/impls/smxx_clean_logits.hpp`
- `csrc/jit_kernels/impls/smxx_cublaslt.hpp`
- `csrc/jit_kernels/impls/smxx_fp8_fp4_mqa_logits.hpp`
- `csrc/jit_kernels/impls/smxx_fp8_fp4_paged_mqa_logits.hpp`
- `csrc/jit_kernels/impls/smxx_layout.hpp`
- `csrc/python_api.cpp`
- `csrc/utils/exception.hpp`
- `csrc/utils/hash.hpp`
- `csrc/utils/layout.hpp`
- `csrc/utils/math.hpp`
- `csrc/utils/system.hpp`
- `deep_gemm/__init__.py`
- `deep_gemm/include/deep_gemm/comm/barrier.cuh`
- `deep_gemm/include/deep_gemm/common/compile.cuh`
- `deep_gemm/include/deep_gemm/common/cute_tie.cuh`
- `deep_gemm/include/deep_gemm/common/exception.cuh`
- `deep_gemm/include/deep_gemm/common/math.cuh`
- `deep_gemm/include/deep_gemm/common/tma_copy.cuh`
- `deep_gemm/include/deep_gemm/common/types.cuh`
- `deep_gemm/include/deep_gemm/common/utils.cuh`
- `deep_gemm/include/deep_gemm/epilogue/sm100_store_cd.cuh`
- `deep_gemm/include/deep_gemm/epilogue/sm100_store_cd_swap_ab.cuh`
- `deep_gemm/include/deep_gemm/epilogue/transform.cuh`
- `deep_gemm/include/deep_gemm/impls/sm100_bf16_gemm.cuh`
- `deep_gemm/include/deep_gemm/impls/sm100_bmk_bnk_mn.cuh`
- `deep_gemm/include/deep_gemm/impls/sm100_fp4_mqa_logits.cuh`
- `deep_gemm/include/deep_gemm/impls/sm100_fp4_paged_mqa_logits.cuh`
- `deep_gemm/include/deep_gemm/impls/sm100_fp8_fp4_gemm_1d1d.cuh`
- `deep_gemm/include/deep_gemm/impls/sm100_fp8_fp4_mega_moe.cuh`
- `deep_gemm/include/deep_gemm/impls/sm100_fp8_mqa_logits.cuh`
- `deep_gemm/include/deep_gemm/impls/sm100_fp8_paged_mqa_logits.cuh`
- `deep_gemm/include/deep_gemm/impls/sm100_tf32_hc_prenorm_gemm.cuh`
- `deep_gemm/include/deep_gemm/impls/sm90_bf16_gemm.cuh`
- `deep_gemm/include/deep_gemm/impls/sm90_bmk_bnk_mn.cuh`
- `deep_gemm/include/deep_gemm/impls/sm90_fp8_gemm_1d1d.cuh`
- `deep_gemm/include/deep_gemm/impls/sm90_fp8_gemm_1d2d.cuh`
- `deep_gemm/include/deep_gemm/impls/sm90_fp8_mqa_logits.cuh`
- `deep_gemm/include/deep_gemm/impls/sm90_fp8_paged_mqa_logits.cuh`
- `deep_gemm/include/deep_gemm/impls/sm90_tf32_hc_prenorm_gemm.cuh`
- `deep_gemm/include/deep_gemm/impls/smxx_clean_logits.cuh`
- `deep_gemm/include/deep_gemm/impls/smxx_layout.cuh`
- `deep_gemm/include/deep_gemm/layout/mega_moe.cuh`
- `deep_gemm/include/deep_gemm/layout/sym_buffer.cuh`
- `deep_gemm/include/deep_gemm/mma/sm100.cuh`
- `deep_gemm/include/deep_gemm/mma/sm90.cuh`

## Evidence Bundle

`evidence/pull-bundles/deepgemm/gh-304`
