Add an opt-in fused gate+up grouped GEMM for AutoEP experts by yh0903 · Pull Request #8733 · deepspeedai/DeepSpeed

yh0903 · 2026-10-03T08:21:38Z

AutoEP's grouped experts run the gate (w1) and up (w3) projections of the same routed rows as two grouped GEMMs. In the backward pass, they compute two input-gradient GEMMs and add the two BF16 results. This adds an opt-in expert_parallel.gate_up_impl="fused", which runs both projections as one torch._grouped_mm over the two weights, concatenated on each call.

Change

GroupedExperts(gate_up_impl="fused") concatenates w1 and w3 per call and runs one grouped GEMM with 2 * hidden_dim output columns.
- The backward pass therefore has one input-gradient GEMM, which accumulates both projections in FP32, and no separate addition.
The activation reads the two halves of the packed output in place:
- swiglu uses a new Triton kernel, swiglu_packed, with the same FP32 math as the existing swiglu kernel.
- The other registered forms (geglu_tanh, swiglu_clamped, swiglu_oai) run their eager expression on views of the halves.
- New forms can register a packed_fused_fn.
The parameters keep their separate w1/w3 layout. Checkpoints, universal checkpoints, expert tensor parallelism and optimizer state, including Muon's per-expert treatment, are unchanged.
The default, "separate", keeps the existing path.
"fused" is rejected rather than ignored when the experts would not run torch._grouped_mm:
- use_grouped_mm=false, at config validation and in GroupedExperts;
- devices that select the Triton grouped GEMM, unless disable_triton_grouped_mm=true.
Docs: autoep.rst and the config reference.

Numerical and autograd contract

The forward expression and parameter-gradient semantics are unchanged. Exact output/weight-gradient equality is an observation in the tested H100 cases, not a guarantee across GEMM shapes, implementations or hardware. The input gradient can differ because it is accumulated in one FP32 GEMM rather than as two rounded BF16 results added. The FP64 regression now uses the same represented BF16 weights as the tested paths, rather than the original pre-quantization FP32 weights; its existing relative-L2 bound (fused error <= 1.05 * separate error) is unchanged. This is not a universal accuracy-improvement or bitwise-equivalence claim.

Cross-GEMM output and parameter-gradient assertions now use the existing grouped-GEMM BF16 policy (rtol=0.02, atol=0.02 * max(abs(reference))) rather than zero tolerance. The skewed/empty-expert forward and smooth-form parameter gradients use the independent FP64 reference. Clamped-form parameter gradients use an independent dense BF16 reference, because rounding a projection across a clamp boundary changes its derivative discontinuously; the bound is not enlarged to absorb those discontinuities. The input-gradient relative-L2 comparison is unchanged. The multi-step test uses the same policy for comparisons across algorithms, while saved/reloaded parameters, optimizer momentum tensors and parameter-group configuration must still match exactly. Exact parity checks between the two implementations of the same elementwise SwiGLU kernel remain unchanged.

The packed SwiGLU Triton path has first-order gradients only: create_graph=True raises rather than silently dropping its second derivative. Forward and backward select the input's CUDA device and its current stream, restoring the caller's current device. The eager expressions for the other activation forms and the default separate path are unchanged.

Tests

CPU, TestFusedGateUp in tests/unit/v1/moe/test_autoep_expert_activation.py:
- packed versus separate halves for every registered form, forward and backward;
- only swiglu has a packed kernel;
- the default is unchanged;
- unsupported settings are rejected (bad value, use_grouped_mm=false, Triton grouped GEMM path), at config validation and in GroupedExperts;
- the config reaches the replaced experts, whose parameters remain w1, w2, w3.
Earlier GPU evidence, on one H100 (305 passed): test_swiglu_triton.py, test_autoep_expert_activation.py and test_autoep_unit.py.
- swiglu_packed matches the existing swiglu kernel's values exactly, forward and backward (rtol=0, atol=0, not a signed-zero bit-pattern assertion). Covered: FP32/FP16/BF16, odd widths, 3-D inputs, empty inputs, non-contiguous input, and an odd last dimension (rejected).
- The fused path joins test_grouped_gemm_paths_match_the_for_loop for every form.
- New test_matches_separate_with_skewed_and_empty_experts, for every form, with an empty expert and unaligned rows:
  - output and w1/w2/w3 gradients match exactly in the tested shapes;
  - the original relative-L2 FP64 error comparison is retained with the corrected represented-weight reference described above.
- Not part of this change: the 2 failures are TestMiniMaxM3Block::test_replaced_block_matches_hf_block. That preset requires Transformers >= 5.15.0, and the image has 5.12.0.
The shipped GroupedExperts, one H100, Qwen3-30B-A3B per-rank expert shape (131,072 routed rows, hidden 2048, expert hidden 768, 16 local experts, BF16):

Per call (ms) Forward, no grad Forward Backward Total

"separate" 2.149 2.292 4.857 9.298

"fused" 2.224 2.236 4.321 8.781

Output and all weight gradients matched exactly; the input gradient differed by 2.9e-3 relative.
End to end, Qwen3-30B-A3B on 8 H100s:
- Setup: expert parallel size 8, micro-batch 4, 16 accumulation steps, sequence length 4096, DeepEP, reentrant activation checkpointing, BF16 parameters and Adam state, 30 steps with 10 warmup.
- Measured with an equivalent benchmark patch on our benchmark base: the same concatenation and the same packed SwiGLU kernel.
- Separate 22,018.8 / 22,018.5 ms against fused 21,601.2 / 21,608.4 ms, in two pairs in the same job: -414 ms (-1.9%).
- Step-1 loss identical. Later losses agree within same-arm repeat noise.
- Peak allocated memory unchanged; reserved memory about 0.4 GB higher.

Review follow-up validation

Current-master targeted validation on NVIDIA H100 80GB HBM3, torch 2.10/CUDA 13.0, Triton 3.6.0 and transformers 5.12.0: 83 passed, 0 skipped (the full 63-case SwiGLU file plus 20 TestFusedGateUp cases). The SwiGLU file also passed 63/63 with no skips under pytest --forked. Cross-device cases used two GPUs; the lifecycle cases use the actual native grouped GEMM.
With only the same new tests applied to the prior runtime, all four focused guards fail as intended: both FP16/BF16 cross-device cases reproduce Triton's inaccessible-pointer error, and both FP32/BF16 higher-order cases fail because create_graph=True was not rejected. This verifies the actual root causes, not an unrelated collection or test-API failure.
Regressions cover strided packed inputs and upstream gradients, two-device/current-stream selection and caller restoration, explicit higher-order rejection, the corrected FP64 reference, and two SGD updates with two accumulation micro-batches per step. All four activations run with and without reentrant recomputation; a model/optimizer state-dict round trip between updates keeps the separate w1/w2/w3 names and shapes and checks the next update against the independent separate path.
The SwiGLU test file imports its runtime in the worker, not during collection, so its Triton driver probe is deferred until after pytest --forked starts the worker.
The follow-up changes launch-device selection and unsupported-autograd handling, not the GEMM implementation, kernel arithmetic, launch configuration, default or parameter layout. The earlier local and E2E performance evidence is reused without retiming; the historical percentages are not new measurements of these guards.

AI review follow-up

Changed tests and documentation only; the runtime is byte-identical to the previous validated head.
The numerical-contract comment is addressed by the reference/tolerance policy above, plus a regression that accepts a single BF16 representable-step change but rejects a clearly incorrect value. On NVIDIA H100 80GB HBM3, torch 2.10/CUDA 13.0, Triton 3.6.0 and transformers 5.12.0, the focused fused-gate-up/grouped-GEMM suite passed 33/33 with no skips.
A synthetic one-ULP perturbation of the fused down-GEMM output passed all 12/12 reference/lifecycle cases with the contract-level checks; the old exact checks failed all 12/12 under the same perturbation. An independently injected 25% output/gradient error was rejected by all four reference cases. These are small test controls, not new GEMM-backend or performance measurements; they demonstrate that legitimate rounding is allowed without accepting an incorrect computation.
The DCO comment was checked against the actual GitHub PR commit list and local ancestry: every PR commit has the configured author's Signed-off-by trailer, and DCO passes. The SHA cited in that comment is not in the PR and is not found by GitHub's commit API, so no commit amendment or history rewrite is needed.

CI and scope

DCO, test collection and changed-files pre-commit passed on the published follow-up. Full GPU CI and infrastructure-smoke jobs were skipped on the PR event; they are not reported as full GPU CI acceptance. The targeted H100 runs above provide the device evidence for this follow-up.

This branch inherits recent upstream master. The backward-overlap child (#8739) remains a separate stacked PR and is not included here; no recompute-skip, buffer-slice, row-weighting-activation or communication optimization is added by this PR.

The new optimizer round trip is a focused SGD test, not an independent qualification of every optimizer, expert-TP/universal-checkpoint combination, ZeRO stage or accelerator. The E2E evidence covers one model/configuration on H100, not long-run convergence or a universal speedup.

AutoEP's grouped experts run the gate and up projections of the same input as two grouped GEMMs, and in the backward pass add the two input-gradient GEMM results. expert_parallel.gate_up_impl="fused" concatenates w1 and w3 on each call and runs one torch._grouped_mm with twice the output columns. A single input-gradient GEMM then accumulates both projections in FP32. The activation reads the two halves of the packed output in place: swiglu through a new Triton kernel, swiglu_packed, with the same FP32 math as the existing swiglu kernel; the other registered forms through their eager expression on views. Parameters keep the separate w1/w3 layout, so checkpoints, expert tensor parallelism and optimizer state are unchanged. The default, "separate", keeps the existing path. "fused" is rejected when the experts would not run torch._grouped_mm. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Signed-off-by: yh0903 <helloyu0903@gmail.com>

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Signed-off-by: yh0903 <helloyu0903@gmail.com>

Select the packed input device for forward and backward while preserving its current stream and the caller device. Reject create_graph through the new packed kernel instead of silently dropping its second derivative. Defer SwiGLU test driver imports until the worker starts. Compare FP64 input-gradient references against the represented BF16 weights, without loosening the existing error bound. Cover all activations over two updates with GAS two, reentrant recomputation and model/optimizer checkpoint reload. Keep kernel arithmetic, GEMM logic, separate parameter layout and opt-in default unchanged, and scope numerical/performance statements to the validated cases. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Signed-off-by: yh0903 <helloyu0903@gmail.com>

chatgpt-codex-connector · 2026-10-10T00:01:06Z

Codex Review Summary

This comment shows the latest Codex review activity on this pull request.

Review	Status	Commit	Review trigger
📝 Code Review	✅ Completed 2026-10-10T00:06:36.351014Z	`ed88d3d`	Draft marked ready

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

Open a pull request for review
Mark a draft as ready
Comment "@codex review" or "@codex security review".

Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings.

chatgpt-codex-connector

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: ed88d3dfd9

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

Open a pull request for review
Mark a draft as ready
Comment "@codex review".

If Codex has suggestions, it will comment; otherwise it will react with 👍.

Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".

chatgpt-codex-connector · 2026-10-10T00:06:33Z

    config.score_apply = param_dict.get("score_apply", "auto")
    config.combine_impl = param_dict.get("combine_impl", "auto")
    config.row_weighting_impl = param_dict.get("row_weighting_impl", "auto")
+    config.gate_up_impl = param_dict.get("gate_up_impl", "separate")


Add the required Signed-off-by trailer

This is a one-parent non-merge commit, but git log -1 --format='%B' 8213255561f7700869cafa528429eec5546669e3 contains no Signed-off-by: trailer. Add the trailer using the configured author identity; otherwise the commit cannot satisfy the repository's DCO/CI requirement.

AGENTS.md reference: AGENTS.md:L8-L9

Useful? React with 👍 / 👎.

chatgpt-codex-connector · 2026-10-10T00:06:33Z

+        torch.testing.assert_close(out_fused, out_separate, rtol=0, atol=0)
+        for name in ("w1", "w2", "w3"):
+            torch.testing.assert_close(getattr(fused, name).grad, getattr(separate, name).grad, rtol=0, atol=0)


Relax parity checks to the documented numerical contract

On a CUDA architecture or library version where the single wide GEMM and the two narrow GEMMs select different valid algorithms, _run_experts_grouped_mm_fused_gate_up explicitly permits rounding differences, but these zero-tolerance output and weight-gradient assertions fail before the FP64 error comparison. The update/checkpoint test repeats the same assumption, so compare against the independent FP64 reference or use contract-level tolerances rather than requiring bitwise parity across GEMM shapes.

AGENTS.md reference: AGENTS.md:L31-L34

Useful? React with 👍 / 👎.

Reuse the existing BF16 grouped-GEMM rounding bound and compare against independent references instead of requiring cross-algorithm exact parity. Keep the input-gradient relative-L2 bound and exact checkpoint state/config restoration assertions. Clamped parameter gradients use dense BF16 references to preserve their derivative branches without enlarging the tolerance. Exercise permitted one-ULP rounding and rejection of incorrect values. Runtime, kernels, opt-in default and parameter layout remain unchanged. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Signed-off-by: yh0903 <helloyu0903@gmail.com>

tohtana

Looks good to me, thank you!

@tohtana

…peedai#8652) # AutoEP: fused Triton kernel for the clamped expert activations  **Branch:** `peng-add-fuse-fn-clamped-form-v1`, rebased on `master` at `5f266e26d` (2026-10-15). The expert activation registry it registers into is merged (deepspeedai#8650), and so is the fused gate+up projection (deepspeedai#8733), which this PR now also covers. **Part of:** deepspeedai#8672 (Q4 2026 roadmap), "Mixture-of-Expert (MoE) Training Scaling and Kernel Development": "AutoEP Performance Optimization & Parallel Folding (AutoEP + AutoTP)" (PoC @tohtana @pengdurice). Follows the merged expert activation registry (deepspeedai#8650) and MiniMax-M3 preset (deepspeedai#8651). ## Summary The clamped expert activations `swiglu_clamped` (DeepSeek-V4) and `swiglu_oai` (GPT-OSS, MiniMax-M3) run as plain PyTorch expressions that keep about four `[rows, ffn]` temporaries for backward. Plain `swiglu` has fused Triton kernels that keep only `gate` and `up`. This PR adds the same kernels for the clamped forms, for both expert layouts: - `gate_up_impl="separate"` (two grouped GEMMs): `swiglu_clamped(gate, up, ...)`. - `gate_up_impl="fused"` (one grouped GEMM, deepspeedai#8733): `swiglu_clamped_packed(gate_up, ...)`, which reads the two halves of the packed GEMM output in place, as master's `swiglu_packed` does. At the busiest rank of a MiniMax-M3 run that ran out of memory (1,045,081 rows x 3,072, bf16), one `swiglu_oai` call goes from 41.86 GiB extra peak memory and 77.8 ms to 17.94 GiB and 11.9 ms, the same as plain `swiglu`. ## Duplicate check Open PRs and issues searched on 2026-10-12 for "swiglu", "clamp", "expert activation" and "gpt-oss": no other work on fused clamped activations. deepspeedai#8748 (another author, open) adds a different kernel, `swiglu_weighted` (routing weight inside the activation, plain `swiglu` only), to the same two files (`swiglu_triton.py` and the registry in `ep_experts.py`). The two do not overlap in function; whichever merges second needs a small rebase. ## Problem Each clamped form is two clamps, a sigmoid, three products and an add in plain PyTorch, each a full `[rows, ffn]` temporary, most of them saved for backward. On an expert-parallel rank this is where the memory goes, because AutoEP has no capacity limit and real routers send some ranks several times the mean number of rows. MiniMax-M3 at 65,536 tokens with expert parallel 32 on 4 nodes of 8 H200 ran out of memory there: the busiest rank received 1,045,081 rows (3.99 times its own 262,144) and failed inside `apply_expert_activation` on one 11.96 GiB temporary, with 129.50 GiB already in use. (Setup: all 60 layers, real checkpoint, `minimax_m3` preset, ZeRO-3 with the optimizer state on the host, activation checkpointing; torch 2.8.0+cu126, transformers 5.15.1, DeepSpeed master `4a5856c25` plus the then-open registry PR.) Over the 57 MoE layers, the busiest of 32 ranks received a median 4.20 times the mean number of rows (maximum 7.84). ## The change | File | What | |---|---| | `deepspeed/ops/triton_ops/swiglu_triton.py` | `_swiglu_clamped_fwd_kernel` / `_swiglu_clamped_bwd_kernel`, modelled on master's swiglu kernels: int64 offsets, float32 math, one rounding on store, `OAI: tl.constexpr` selects the form, `PACKED: tl.constexpr` reads `gate` and `up` from one `[rows, 2 * INTER]` tensor with master's packed indexing. The clamp gradient passes at the limit, as `torch.clamp`'s does. Public `swiglu_clamped(gate, up, alpha, limit, oai)` and `swiglu_clamped_packed(gate_up, alpha, limit, oai)`; plain-PyTorch fallback without Triton or on the CPU. The packed backward raises for `create_graph=True`, as master's packed swiglu does | | `deepspeed/moe/ep_experts.py` | registers the kernels as `fused_fn` and `packed_fused_fn` of `swiglu_clamped` and `swiglu_oai`. Nothing else changes: the grouped-GEMM expert paths reach them through `apply_expert_activation(..., fused=True)` and `apply_packed_expert_activation`; the for-loop reference path keeps `fused=False` | | `docs/_pages/config-json.md` | one sentence: which forms run a fused kernel | | `tests/unit/v1/ops/triton_ops/test_swiglu_triton.py`, `tests/unit/v1/moe/test_autoep_expert_activation.py` | see Tests | Size: 264 added and 2 removed lines under `deepspeed/` (232 added in `swiglu_triton.py`, 32 in `ep_experts.py`). ## Measured One H200, torch 2.8.0+cu126, Triton 3.4.0, through `apply_expert_activation` / `apply_packed_expert_activation` (`fused=True`). 1,045,081 rows x 3,072 in bf16 (one `[rows, ffn]` tensor = 5.98 GiB), alpha 1.702, limit 7.0, inputs at 6 x standard normal so most values are past a clamp. Extra peak = `max_memory_allocated` during forward + backward minus the memory allocated before (inputs and upstream gradient). Kept = memory still allocated after the forward, minus the output. Time = median of 5 forward + backward passes after 3 warm-ups. | Form | Layout | master `1c862f7df`: extra peak / kept / time | this PR: extra peak / kept / time | |---|---|---|---| | `swiglu` (existing kernels, for comparison) | separate | 17.94 GiB / 0.00 GiB / 11.9 ms | 17.94 GiB / 0.00 GiB / 11.9 ms | | `swiglu` | packed | 17.94 / 0.00 / 11.9 ms | 17.94 / 0.00 / 11.9 ms | | `swiglu_clamped` | separate | 35.88 / 17.94 / 51.6 ms | **17.94 / 0.00 / 11.9 ms** | | `swiglu_clamped` | packed | 35.88 / 17.94 / 92.1 ms | **17.94 / 0.00 / 13.1 ms** | | `swiglu_oai` | separate | 41.86 / 23.92 / 77.8 ms | **17.94 / 0.00 / 11.9 ms** | | `swiglu_oai` | packed | 41.86 / 23.92 / 118.3 ms | **17.94 / 0.00 / 13.2 ms** | Per MoE layer on that rank, `swiglu_oai` keeps 23.92 GiB less for backward and runs 6.5 times faster (separate) or 9.0 times faster (packed). The packed clamped kernels are 1.2-1.3 ms slower than master's packed `swiglu` kernel, which does less arithmetic per element. Bench script: `expert_activation_bench.py` (TODO(user): attach as a gist or drop this line). ## Tests Command, from `tests/`, on 1 node of 8 H200 (143,771 MiB each, driver 580.126.09), torch 2.8.0+cu126, Triton 3.4.0, transformers 5.15.1, pytest 9.1.1: ```bash DEEPSPEED_RUN_H100_TESTS=1 python -m pytest unit/v1/ops/triton_ops/test_swiglu_triton.py \ unit/v1/moe/test_autoep_expert_activation.py -v \ -k "not (skewed_and_empty or accumulated_updates_and_checkpoint)" ``` Result: **204 passed**, 12 deselected, in 31 s. On master `5f266e26d` (2026-10-15): **204 passed**, 12 deselected, in 34 s on torch 2.8.0; the two empty-expert tests from deepspeedai#8733 on torch 2.9.1: 12 passed on this branch and 12 passed on master. The two deselected tests are master's (deepspeedai#8733). On torch 2.8.0+cu126 on H200 they do not finish, on master as well as on this branch: they pass an empty expert group (`counts = [37, 0, 291, 72]` and `[16, 0, 32, 16]`) to `torch._grouped_mm`, whose backward never returns for some shapes when a group is empty on that torch version (GPU at 100% with no progress; known from MiniMax-M3 work, one-GPU reproduction available). AutoEP itself never passes an empty group, because `generate_permute_indices` gives every local expert at least 8 rows. On torch 2.8.0 master's `test_matches_separate_with_skewed_and_empty_experts[swiglu]` did not finish in 300 s; on torch 2.9.1+cu128 (Triton 3.5.1) both tests pass, **12 passed** on this branch and 12 passed on master. New and changed tests: | Test | What it would catch | |---|---| | `test_clamped_forward_matches_reference` (2 forms x 3 dtypes x 6 shapes) | a wrong output: compared with the expression evaluated in float32 and rounded once; inputs at scale 4 so about 8% are past a clamp; the kernel must also be no farther from float32 than the eager low-precision evaluation | | `test_clamped_backward_matches_reference` (2 x 3 x 5) | wrong input gradients, the same way, with a random upstream gradient | | `test_clamped_alpha_and_limit_are_honoured` | alpha or limit ignored; a nonzero gradient where an input was clamped | | `test_clamped_gradient_at_the_limit_is_inclusive` | a gradient that stops at `gate == limit` or `up == -limit`, where `torch.clamp`'s passes | | `test_clamped_empty_input`, `test_clamped_non_contiguous_input`, `test_clamped_shape_and_dtype_mismatch_rejected` | the edge cases master's swiglu tests cover | | `test_clamped_packed_matches_the_separate_kernel_exactly` (2 forms x 2 dtypes x 4 shapes, including 3-D and empty) | a packed path that reads the wrong half or row: output and gradient bit-equal to the separate kernel | | `test_fused_kernel_matches_plain_pytorch` (every form with a `fused_fn`) | `apply_expert_activation(fused=True)` differing from the plain form evaluated in float32 and rounded once (bf16, 256 x 3,072) | | `test_packed_kernels_are_registered` (was `test_only_swiglu_has_a_packed_kernel`) | the fused gate+up path silently running the clamped forms in plain PyTorch | | `test_grouped_gemm_paths_match_the_for_loop`, `test_packed_activation_matches_the_separate_halves` (master's, now running the new kernels) | the grouped-GEMM paths differing from the for-loop reference path | `pre-commit run --files` passes on every changed file. ## Compatibility The change is an elementwise kernel behind AutoEP's existing activation registry; it does not depend on the ZeRO stage, offload or parallel layout. | | Validated | Not validated | |---|---|---| | Precision | float32, float16, bf16 (kernel tests); bf16 through the expert paths | | | Expert layouts | `gate_up_impl` `"separate"` and `"fused"`; `torch._grouped_mm` and Triton grouped-GEMM paths (unit tests) | | | End to end | | no training job has run with this kernel. The MiniMax-M3 runs on 4 nodes used the plain PyTorch form (the out-of-memory case above was avoided with expert parallel 4) | | Accelerators | CUDA (H200); CPU fallback (plain PyTorch) | ROCm, XPU and other Triton backends. The packed path sets the device with `torch.cuda.device` (`#ignore-cuda`), as master's packed swiglu does | ## AI assistance This PR was written with AI assistance (Claude). 🤖 Generated with [Claude Code](https://claude.com/claude-code) Signed-off-by: pengdurice <pengduhit@gmail.com>

This was referenced Oct 4, 2026

Overlap AutoEP's expert weight gradients with the DeepEP backward combine #8739

Open

AutoEP: opt-in skip of the recomputed DeepEP combine #8745

Open

yh0903 and others added 2 commits October 9, 2026 16:24

Sync upstream master for fused gate-up review

0c0720b

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Signed-off-by: yh0903 <helloyu0903@gmail.com>

yh0903 marked this pull request as ready for review October 10, 2026 00:00

yh0903 requested review from hwchen2017, loadams, tjruwase and tohtana as code owners October 10, 2026 00:00

chatgpt-codex-connector Bot reviewed Oct 10, 2026

View reviewed changes

tohtana approved these changes Oct 12, 2026

View reviewed changes

tohtana enabled auto-merge October 12, 2026 05:39

tohtana added this pull request to the merge queue Oct 12, 2026

Merged via the queue into deepspeedai:master with commit 3a656d3 Oct 12, 2026
15 checks passed

pengdurice mentioned this pull request Oct 16, 2026

AutoEP: fused Triton kernel for the clamped expert activations #8652

Merged

Navigation Menu

Repository navigation

Add an opt-in fused gate+up grouped GEMM for AutoEP experts - #8733

Add an opt-in fused gate+up grouped GEMM for AutoEP experts#8733
tohtana merged 4 commits into
deepspeedai:masterfrom
yh0903:yh0903/autoep-fused-gate-up

yh0903 commented Oct 3, 2026 •

edited

Loading

Uh oh!

chatgpt-codex-connector Bot commented Oct 10, 2026 •

edited

Loading

Uh oh!

chatgpt-codex-connector Bot left a comment

Uh oh!

chatgpt-codex-connector Bot Oct 10, 2026

Uh oh!

chatgpt-codex-connector Bot Oct 10, 2026

Uh oh!

tohtana left a comment

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Per call (ms)	Forward, no grad	Forward	Backward	Total
`"separate"`	2.149	2.292	4.857	9.298
`"fused"`	2.224	2.236	4.321	8.781

Conversation

yh0903 commented Oct 3, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Change

Numerical and autograd contract

Tests

Review follow-up validation

AI review follow-up

CI and scope

Uh oh!

chatgpt-codex-connector Bot commented Oct 10, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codex Review Summary

Uh oh!

chatgpt-codex-connector Bot left a comment

Choose a reason for hiding this comment

💡 Codex Review

Uh oh!

chatgpt-codex-connector Bot Oct 10, 2026

Choose a reason for hiding this comment

Uh oh!

chatgpt-codex-connector Bot Oct 10, 2026

Choose a reason for hiding this comment

Uh oh!

tohtana left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

yh0903 commented Oct 3, 2026 •

edited

Loading

chatgpt-codex-connector Bot commented Oct 10, 2026 •

edited

Loading