Skip to content

[Apple Silicon] [MLX] Fuse MoE combine multiply-reduce into one Metal kernel - #29804

Open
jlee5814 wants to merge 24 commits into
sgl-project:mainfrom
jlee5814:mlx-fused-moe-combine
Open

jlee5814 wants to merge 24 commits into
sgl-project:mainfrom
jlee5814:mlx-fused-moe-combine

Conversation

@jlee5814

@jlee5814 jlee5814 commented Jul 1, 2026

Copy link
Copy Markdown
Contributor

Motivation

The reference MoE combine, (y * scores[..., None]).sum(axis=-2), runs a broadcast multiply and a reduction with a transient materialized between them. Fusing both into one Metal kernel saves a full read of y, the transient round trip, and one dispatch per MoE layer.

This PR began as the kernel plus its verification suite. Review exposed a real coverage bug in the eligibility guard (rank 3 assumption against a rank 4 production call site), which the fix generalized. The performance question required measured serving numbers, which required an integration path. The abstraction ask captured the JIT declaration pattern this kernel now shares with fused_swiglu.py. A founder-level review of the abstraction and its test suite followed; this revision applies that review's fixes.

Modifications

Guard generalization and bf16 scores

  • SGLang never flattens batch and sequence upstream, so every real combine call site is rank 4 y with rank 3 scores (prefill [1, L, 4, 2048], batched decode [R, 1, 4, 2048]). The original rank 3 guard would have fallen back on every real forward pass; caught in review. Both target models run bf16 activations, and the original _SCORES_DTYPES rejected bf16 scores.
  • The contract is now scores.shape == y.shape[:-1] with rank at least 3. Leading dims flatten into the kernel's row dimension before dispatch and the output reshapes back. bf16 scores are accepted.

Default off integration

  • SGLANG_MLX_FUSE_MOE_COMBINE = EnvBool(False) in environ.py, plus patch_moe_combine_with_fused, a subclass swap on Qwen2MoeSparseMoeBlock and Qwen3MoeSparseMoeBlock, gated in _load_model beside the existing SwiGLU patch. Per call eligibility stays inside fused_combine, which falls back inline, so patching never changes results. Merging changes nothing for any user; the flag flip is a separate, evidence backed decision.

metal_jit abstraction

  • hardware_backend/mlx/metal_jit.py: kernel files declare source once via the @metal_jit.kernel(...) decorator on a MetalJitOp subclass. The base class owns a can_fuse to dispatch_fused or dispatch_fallback template with typed NotFusable routing. Compile failures degrade via a memoized guard mirroring jit_kernel/moe_fused_gate.py's can_use_moe_fused_gate pattern (catch, log once, memoize, fall back) instead of crashing the forward pass.
  • MetalJitKernel still owns the per dtype key compile cache and warm dedupe underneath, unchanged in behavior.
  • Both kernels migrate as the first consumers. Name templates keep kernel names byte identical, so profiler labels and shader cache keys do not change; the SwiGLU warm set folds into warm_once.
  • Patch registration stays in model_runner. The env flags' one reader should remain the load path, and the patch functions encode mlx_lm model class knowledge, not Metal JIT knowledge.

can_fuse is now abstract, per review

  • A founder-level review found FusedGateQmvSiluMulKernel was a live example of the risk in an optimistic base default: it inherited can_fuse returning True unconditionally instead of declaring its own eligibility, relying entirely on dispatch_fused raising NotFusable before any observable effect to stay safe. MetalJitOp.can_fuse now raises NotImplementedError (metal_jit.py:123-125, matching this codebase's existing abstract-method idiom of a bare raise NotImplementedError, used already by dispatch_fused/dispatch_fallback on the same class; no abc.ABC is used anywhere in this tree). The class docstring states the rule directly: every op declares its own predicate, the base declares none.
  • FusedMoeCombineKernel.can_fuse already had a real structural predicate (fused_combine.py:100-112); unchanged.
  • FusedGateQmvSiluMulKernel.can_fuse (new, fused_swiglu.py:246-263) delegates to the same _check_eligibility regime check dispatch_fused/dispatch_fallback already share, so all three stay in sync by construction. This does not change the op's observable behavior: for structurally ineligible inputs, dispatch() now short-circuits straight to dispatch_fallback instead of first attempting dispatch_fused (which raised the identical diagnostic anyway); for eligible inputs whose Metal compile fails, can_fuse cannot see that and the routing is unchanged. SwiGLU's asymmetry is preserved exactly: dispatch_fallback still always raises the diagnostic chain rather than computing, since its real fallback needs wrapper layer context the op does not receive; only its predicate became explicit.
  • Three test_metal_jit.py stub classes that exercised dispatch_fused/dispatch_fallback routing without overriding can_fuse (relying on the removed default) now declare an explicit can_fuse returning True, so they keep testing what they were written to test.

Tolerance test fix, per review

  • The same review found the bf16 1 ULP, zero headroom bound in test_fused_combine.py could reach tens of thousands of "ULP" at production scale. Root cause, found this session: the huge values are a metric artifact, not a kernel regression. Near zero, fp16/bf16 use absolute (subnormal-adjacent) steps far finer than the ULP ceiling's relative-error assumption, and the two paths can even land on opposite sides of zero (e.g. 0.0 vs -1e-7), where the bit-pattern-distance metric's own documented same-sign assumption breaks down. allclose(atol=tol) already bounds these elements correctly.
  • Fix: the integer-ULP ceiling is now checked only on elements at or above a 1e-3 ground-truth-magnitude floor (_ULP_FLOOR, test_fused_combine.py), where the bit-pattern-distance metric is meaningful. A live 100-seed-per-shape sweep across every shape in _SHAPES plus the real production prefill shape (1, 268, 4, 2048) found max ULP at or above the floor was 0-1 in every case, for both dtypes; the ceilings (_ROUNDED_MAX_ULP) are now 3 for both fp16 and bf16, roughly 3x headroom over that observed max. The check is not deleted or loosened past what the evidence supports.

Accuracy Tests

  • test_fused_combine.py: 102 of 102 matrix cases fire the fused path. fp16 and bf16 scores are bit exact against the fp32 reference; fp32 scores are asserted within narrowing tolerance and a floor-masked integer-ULP ceiling (see Modifications). [hardware verified, this run] the updated bound was demonstrated seed robust: test_scores_fp32_correctly_rounded and the production-shape case in test_leading_dims_match_reference were run across 7 shifted seed offsets (0 through 60000) on this hardware, all passing, with observed ULP at or above the floor staying in {0, 1} throughout. Patch idempotency on real, small dimension mlx_lm blocks is a no op with bit identical outputs.
  • test_metal_jit.py: cache identity, name formatting, warm dedupe, decorator validation, and the dispatch template's routing contract (fused, fallback, NotFusable propagation, now including can_fuse as a required override), including the compile failure memoization guard.
  • test_fused_swiglu.py: regime diagnostics through the wrapper, dtype membership fallback (not just dtype agreement), warm skip logging, warmup enumeration, and (new) the explicit can_fuse predicate.
  • 47 tests across the three files at head (17 + 10 + 20, pytest --collect-only) [hardware verified, this run]. Unchanged from before this revision: the fixes changed test bodies and one helper signature, not test count.
  • [hardware verified, this run] Full suite: 45 passed, 2 skipped (opt-in, HF-checkpoint-gated, not a hardware skip), 0 failed, both before and after this revision's changes.
  • [hardware verified, this run] The three NotFusable safety sub claims from the prior review (forced Metal compile failure on an eligible input: no raise to caller, bit exact output, no partial/corrupted cache state) were re-run against the real fused_combine() entry point after the can_fuse change and still hold; the equivalent check on FusedGateQmvSiluMulKernel's own compile-failure path (now gated by its own can_fuse) also holds.

Speed Tests and Profiling

Original end to end measurement [author reported, not reverified end to end], mlx-community/Qwen1.5-MoE-A2.7B-Chat-4bit on a 24 GB Apple Silicon machine, interleaved on/off server launches, warmup discarded, temperature 0, medians with min and max:

config decode bs=1 tok/s per step ms decode bs=4 tok/s prefill L=268 tok/s
unfused 131.02 [130.61, 131.67] 7.632 201.20 [199.13, 203.11] 758.4 [756.0, 815.8]*
fused 131.80 [131.64, 132.11] 7.587 203.56 [201.82, 205.55] 818.2 [809.8, 820.1]

* One unfused launch at 9.7 GB free excluded as memory pressure (732.7 tok/s).

Decode bs=1 is +0.6 percent, decode bs=4 is +1.2 percent, prefill is +0.0 to +1.5 percent depending on which launches are compared.

Isolated op measurement, decode bs=1 [hardware verified, this run]: re-ran the fused-vs-fallback dispatch check on this hardware after the can_fuse change. dispatch() still routes through exactly one metal_jit.get("fused_moe_combine", ...) call on the fused path, output still bit exact against an independent reference, and the fused path still measures faster than the fallback (paired timing this run: +2.4%, a simpler protocol than the prior review's bootstrapped interleaved A/B, so not a direct comparison to the fuller table below).

Prior isolated op measurement [prior review, not rerun this session], Apple M4 Pro, 24 GB, bootstrap 95 percent confidence interval:

shape fused (us) fallback (us) speedup 95% CI excludes zero
decode bs=1 17.87 18.80 +9.6% yes
decode bs=4 18.39 19.54 +7.4% yes
prefill L=268 34.83 190.94 +82.2% yes

The two end-to-end/isolated tables are not directly comparable: the isolated measurement strips out every other per token cost (attention, routing, the expert GEMMs, sampling, KV bookkeeping), so it upper bounds what the combine op itself could contribute, while the end-to-end measurement is what actually reaches a user after dilution by the rest of the forward pass. The metal_jit registry promotion and the can_fuse abstraction change are both same-process refactors with no dispatch-count or numerics change on the fused path, confirmed this run; neither table's numbers are affected by them.

Known Limitations

  • warmup_specs declares each op's supported variant set and has no consumer yet. The working warm path (SwiGLU only; combine pays its first dispatch JIT cost today) warms instantiated shapes derived from the live model. aot.py reconciles the two in a follow up PR.
  • The patch time can_fuse(switch_mlp) free function in fused_swiglu.py is unrelated to MetalJitOp.can_fuse; a rename is follow up material. This is a naming collision only; both FusedMoeCombineKernel and FusedGateQmvSiluMulKernel now declare their own MetalJitOp.can_fuse explicitly (see Modifications), so the base class no longer has an optimistic default for a future op to silently inherit.
  • SwiGLU's dispatch_fallback intentionally raises rather than computes, since its real fallback needs wrapper layer context the op does not receive. Unifying that is follow up scope.
  • Nothing counts fallback occurrences on either op (no logging on FusedMoeCombineKernel's fallback path, a once-per-process warning on SwiGLU's). This was flagged by review as a real but lower-leverage gap for this pass and is not addressed here.
  • The register_mlx_ci addition to ci_register.py this branch carries is now byte identical to what [Apple Silicon] [CI] Move the MLX lane to the check-changes + pr-gate composite #30121 independently merged to main; a real merge auto-resolves it with no manual edit, confirmed this run (git hash-object matches on both sides). Attempting the full merge this run surfaced a separate, unrelated conflict: an upstream commit (#32448, merged after this branch forked) relocated test_fused_swiglu.py to test/registered/ and rewrote its CI-harness registration, colliding with this branch's own, differently-evolved copy of that file at its old path. That conflict is outside what the prior review named and was not resolved unilaterally in this pass; it is reported here as the actual remaining blocker to a clean merge, and needs a codeowner call on how to reconcile the two files' content, not just their location.
  • Mergeable state at time of writing needs the test_fused_swiglu.py relocation above resolved before this branch can merge cleanly against current main; the ci_register.py duplication that previously blocked it is confirmed moot.

Checklist

Review and Merge Process

  1. Ping Merge Oncalls to start the process. See the PR Merge Process.
  2. Get approvals from CODEOWNERS (yeahdongcn for this path) and any other reviewers.
  3. Trigger CI tests with comments or contact authorized users to do so.
  4. After green CI and required approvals, ask Merge Oncalls or a maintainer with write permission to merge.

CI States

Latest PR Test (Base): ❌ Run #31957166728
Latest PR Test (Extra): ❌ Run #31957166641

@jlee5814
jlee5814 requested a review from yeahdongcn as a code owner July 1, 2026 04:26
@gemini-code-assist

Copy link
Copy Markdown
Contributor

Warning

You have reached your daily quota limit. Please wait up to 24 hours and I will start processing your requests again!

@chatgpt-codex-connector chatgpt-codex-connector Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: 975232d23f

ℹ️ 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".

Comment thread python/sglang/srt/hardware_backend/mlx/moe/fused_combine.py Outdated
@jlee5814 jlee5814 changed the title [Apple Silicon][MLX] Fuse MoE combine multiply-reduce into one Metal kernel [Apple Silicon] [MLX] Fuse MoE combine multiply-reduce into one Metal kernel Jul 1, 2026
jlee5814 added 3 commits July 4, 2026 15:12
…kernel

Fuses out[b,h] = sum_k y[b,k,h] * scores[b,k] into a single mx.fast.metal_kernel dispatch, fp32 product and fp32 accumulate narrowed to y.dtype on write. Scores carry an independent Metal dtype (TS), so the fp16-y + fp32-scores production combo fuses instead of falling back. Bit-exact vs fp32 reference for low-precision operands across 17 shapes, correctly-rounded within 2 ULP for fp32 scores. Adds test_fused_combine.py: 68 cases, fused path asserted fired, negative control for the fallback gate.
…gate

The fallback returned MLX's promoted dtype, fp32 for the production fp16-y + fp32-scores combo, while the fused path narrows to y.dtype, so the eligibility gate alone changed the output contract (caught by Codex review). The fallback now narrows to y.dtype, and a negative control with fp32 scores asserts the fallback's output dtype.

can_fuse also rejects zero-size B / TOP_K / H, which previously passed the gate and failed Metal compilation (TOP_K=0 emits a zero-length array, H=0 a constexpr division by zero); three controls cover them. The observed ULP maxima for fp32 scores (2 fp16, 1 bf16) are now asserted ceilings instead of report-only numbers, so the accuracy table is a contract.
Condenses both module docstrings, drops the section headers, and cuts comments that restate adjacent code, keeping the ones carrying non-obvious facts: geometry arithmetic, the scores-pointer dtype behavior, gate rationale, and the ULP encoding argument. Function docstrings collapse to one line except the fused_combine numerical contract. No code change: the AST is identical outside docstrings and a comment inside the Metal source string.

@yeahdongcn yeahdongcn left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What is the e2e perf gain?

_OUTPUTS_PER_TG = _THREADS_PER_TG * _N_READS # 256


_KERNEL_SOURCE = r"""

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Could we introduce some abstraction here for declaring JIT kernels on Apple Silicon? Since we already have python/sglang/srt/hardware_backend/mlx/moe/fused_swiglu.py, it seems we're starting to establish common patterns for things like registration, warmup, and related initialization. It would be good to capture those patterns so future JIT kernels can follow a consistent approach.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Added metal_jit.py in ab47947.

MetalJitKernel owns the source, per dtype lazy compile, warmup dedupe, and name mangling. Both kernels migrate as first consumers; kernel names stay byte identical. Patching stays in model_runner: the patch functions encode mlx-lm model knowledge, not Metal JIT knowledge. Tested in test_metal_jit.py, plus double patch is a no op.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Registry follow-up landed as well; see the thread on metal_jit.py for details. Declaring a new JIT kernel is now source plus one registration call.


def can_fuse(y: mx.array, scores: mx.array) -> bool:
"""Cheap structural check: does this combine match the fast-path regime?"""
if y.ndim != 3 or scores.ndim != 2:

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

can_fuse() only accepts y.ndim == 3 and scores.ndim == 2, why? The mlx-lm Qwen3-MoE combine site has y shaped like [batch, seq, top_k, hidden] and scores shaped like [batch, seq, top_k] before y = (y * scores[..., None]).sum(axis=-2).

@jlee5814 jlee5814 Jul 6, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Fixed in a34b78.

SGLang never flattens upstream; the guard was written against a shape that never arrives. Contract is now scores.shape == y.shape[:-1], rank >= 3, leading dims flattened; bf16 bit exact. 102 matrix cases pass, including the serving shapes.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What is the e2e perf gain?

Measured on Qwen1.5-MoE-A2.7B-Chat-4bit, eight interleaved launch pairs, medians [min, max]. Decode bs=1: 131.02 [130.61, 131.67] vs 131.80 [131.64, 132.11] tok/s, +0.6%, one saved dispatch plus transient per MoE layer. bs=4: +1.2%, spreads overlap. Prefill L=268: +0.0 to +1.5% memory matched, but the unfused baseline hit a discrete slow mode (756 to 758 tok/s, 4 of 8 launches) that fused never entered in 8 of 8. Flag ships default off.

return str(dtype).replace("mlx.core.", "").replace(".", "_")


class MetalJitKernel:

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

MetalJitKernel is a useful step, but I do not think it fully removes the repeated pattern yet. Each fused kernel still has to declare a module-level _KERNEL = MetalJitKernel(...) and keep a local _get_kernel() wrapper mostly for caching/test seams.

Since we are starting to add more MLX JIT kernels, can we make metal_jit.py own a small registry/factory instead? For example, kernels could be registered by name/template and looked up through one shared API that owns dtype-key caching, name formatting, warmup bookkeeping, and the test seam. Then fused_combine.py and fused_swiglu.py would only define the kernel source plus call the registry, instead of each file carrying the same _KERNEL/_get_kernel pattern.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Confirmed. A registry would remove the per-file _KERNEL + _get_kernel pair. One thing not visible from the diff: the wrapper was the dispatch tests' spy seam, not leftover pattern. The registry subsumes it: tests patch the registry lookup, a cleaner single seam. Will add the name keyed registry, migrate both kernels to source plus registration, and move the seam.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Resolved in 4e1201f.

Kernels reduce to source plus registration, with the registry owning dtype caching, naming, warmup and the test seam. d684d3d adds the dispatch template on top per the follow-up sketch.

@@ -0,0 +1,101 @@
"""Unit tests for the MetalJitKernel JIT declaration surface.

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Although these tests have been added, could we also make sure they're actually executed by the CI in #30121?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, model free and in stage A's explicit list, so they run on every PR once #30121 lands.

Code owner review on sgl-project#29804 asked for one shared API instead of each
kernel file holding a private _KERNEL instance and a _get_kernel spy
wrapper. register()/get()/warm_once() replace both per file, and
MetalJitKernel's caching and warmup carry over unchanged underneath.
Kernel name templates are untouched, so shader cache keys and profiler
labels do not move. The dispatch spy seam for tests moves from each
file's local _get_kernel to the shared metal_jit.get.
"""


metal_jit.register(

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

How about something like this?

@metal_jit.kernel(
    name="fused_moe_combine",
    name_template="fused_moe_combine_y{0}_s{1}",
    input_names=["y", "scores"],
    output_names=["out"],
)
class FusedMoeCombineKernel(MetalJitOp):
    source = _KERNEL_SOURCE

    def dispatch(...): ...

    def warmup_specs(model) -> Iterable[WarmupSpec]: ...

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@metal_jit.kernel plus MetalJitOp only, SwiGLU migrated, register deleted. warmup_specs stubbed per your request, real specs in the aot.py follow up. Sources byte identical, tests pass.

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I just DMed you on Slack. Maybe we can have a quick discussion?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

SwiGLU's can_fuse is newly written (the module level one is a patch time check), and its fallback re-raises the existing diagnostics since the nn path needs wrapper context. Deferred: docstrings, can_patch rename, NotFusable.

jlee5814 added 4 commits July 7, 2026 21:08
Declarative registration surface proposed by yeahdongcn on PR sgl-project#29804:
subclasses set source and implement dispatch, the @metal_jit.kernel(...)
decorator validates and registers them. WarmupSpec and the default
warmup_specs hook are placeholders for the aot.py policy layer in a
follow up PR. register() stays as the transitional entry point until
both kernel modules migrate.
…eview

FusedMoeCombineKernel subclasses MetalJitOp with source, the can_fuse
guard, and dispatch (geometry unchanged); fused_combine stays as a thin
module level wrapper so no call sites outside this file change. Registry
tests register through the decorator; new tests cover decorator
validation, the missing source error, the warmup_specs default, and the
resolved kernel name for a representative dtype pair.
FusedGateQmvSiluMulKernel carries the Path B source and dispatch
(guards and geometry unchanged); fused_gate_qmv_silu_mul stays as a
thin module level wrapper so patch and warm paths are untouched. With
both kernels on the decorator, metal_jit.register has no callers, so
registration now lives only in the decorator and the API has one shape.
Comment on lines +223 to +234
try:
from mlx_lm.models.qwen2_moe import Qwen2MoeSparseMoeBlock

targets.append((Qwen2MoeSparseMoeBlock, _fused_qwen2_moe_call))
except ImportError:
pass
try:
from mlx_lm.models.qwen3_moe import Qwen3MoeSparseMoeBlock

targets.append((Qwen3MoeSparseMoeBlock, _fused_qwen3_moe_call))
except ImportError:
pass

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

  for module_name, class_name, fused_call in (
      ("mlx_lm.models.qwen2_moe", "Qwen2MoeSparseMoeBlock", _fused_qwen2_moe_call),
      ("mlx_lm.models.qwen3_moe", "Qwen3MoeSparseMoeBlock", _fused_qwen3_moe_call),
  ):
      try:
          module = importlib.import_module(module_name)
      except ImportError:
          continue
      targets.append((getattr(module, class_name), fused_call))

Looks better.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This has been adopted.


source = _KERNEL_SOURCE

@staticmethod

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It seems not necessary.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This has been dropped.

# Registered with the CPU suite (runtime no-op marker, parsed via AST). On
# non-Apple-Silicon CI runners the whole TestCase skips via the @skipUnless
# guard below, so this is the harmless "yes this test exists" registry signal.
register_cpu_ci(est_time=90, suite="base-a-test-cpu")

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Should we have the new register_mlx_ci?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Fixed.

Both files carry register_mlx_ci(suite="stage-a-unit-test-mlx") alongside register_cpu_ci per the lane convention. The marker is backported verbatim from #30121.

def warmup_specs(self, model) -> Iterable[WarmupSpec]:
"""Precompile keys aot.py should warm for ``model``.

Hook for the AOT policy layer (follow up PR) to enumerate per model

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Can we do it in this PR?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done.

Both ops enumerate WarmupSpecs from declared dtype support; aot.py consumes them in the follow up. Also fixed a gap this surfaced: eligibility checked dtype agreement but not membership, so matched fp32 could reach the unvalidated kernel; such dtypes now fall back.

…its own predicate

Founder review found FusedGateQmvSiluMulKernel a live counterexample to the
optimistic can_fuse default: it inherited an unconditional True instead of
declaring its own eligibility, relying entirely on dispatch_fused raising
NotFusable before any observable effect. can_fuse now raises
NotImplementedError on the base (matching this codebase's existing
raise-NotImplementedError abstract-method idiom, no abc.ABC), and
FusedGateQmvSiluMulKernel gets an explicit predicate delegating to the same
_check_eligibility regime check dispatch_fused/dispatch_fallback already
share. FusedMoeCombineKernel already had a real predicate and is unchanged.
Behavior is unchanged on both ops, including SwiGLU's asymmetric fallback
that always raises rather than computes; only the predicate became explicit.
Three test_metal_jit.py stub classes that exercised dispatch routing without
overriding can_fuse now declare one explicitly, so they keep testing what
they were written to test.
…o noise

Founder review found the bf16 1 ULP, zero-headroom bound in
test_scores_fp32_correctly_rounded fails at production prefill scale and
isn't seed robust, while the primary allclose gate never failed. Root cause,
found this session with a live 100-seed-per-shape sweep: the huge ULP values
(up to tens of thousands) are a metric artifact, not a kernel regression.
Near zero, fp16/bf16 space successive values by an absolute step far finer
than the ULP ceiling's relative-error assumption, and the two paths can land
on opposite sides of zero, where the bit-pattern-distance metric's own
documented same-sign assumption breaks down. allclose(atol=tol) already
bounds these elements correctly.

_low_precision_ulp now takes an optional (gt, floor) pair to exclude
elements below a 1e-3 ground-truth-magnitude floor from the max. Above that
floor, the same sweep (every _SHAPES entry plus the real production prefill
shape) found max ULP was 0-1 in every case for both dtypes, so
_ROUNDED_MAX_ULP moves from {fp16: 2, bf16: 1} (no headroom) to {fp16: 3,
bf16: 3} (roughly 3x headroom over the observed max). The check is not
deleted or loosened past what the evidence supports. Demonstrated seed
robust by running the affected tests across 7 shifted seed offsets, all
passing (reports/work-order-run.md).
…erge

test_fused_gate_qmv_silu_mul_matches_unfused lost its skip guard in the
upstream/main merge that relocated this file (87f2a6e); without it the
test hard-fails with KeyError on any host lacking SGLANG_MLX_TEST_MODEL
instead of skipping, unlike its sibling test_patched_switchglu_matches_unpatched.
Verified via a differential pytest run against the pre-relocation content at
7528970: identical 17 node ids both sides, 15 passed/2 skipped restored.
test/registered/unit/README.md requires registered unit tests to inherit
CustomTestCase (from sglang.test.test_utils) rather than plain
unittest.TestCase, a rule yeahdongcn has enforced before (sgl-project#32115). Both
files predate the relocation and were on bare unittest.TestCase; swap is
mechanical since CustomTestCase only adds a safe setUpClass/tearDownClass
wrapper. test_fused_swiglu.py uses pytest-native monkeypatch/caplog
fixtures throughout and is left as is; converting it away from those would
be a real rewrite, not a mechanical fix, so it is reported instead of
changed. Full MLX suite re-run after this change: 220 passed, 4 skipped,
0 failed.
… relocation merge

Upstream's relocation (580b1ac, sgl-project#32448) added `import sys` and
`if __name__ == "__main__": sys.exit(pytest.main([__file__, "-v"]))` when it
moved this file into test/registered/. The merge that pulled that commit into
this branch (87f2a6e) kept this branch's own local edits to the file at
its old path and dropped the __main__ addition -- the second casualty of that
merge alongside the @requires_model gate fixed in 2c62ea3.

Without it, ci_register.py's collect_tests() raises ValueError at collection
time under `test/run_suite.py --hw mlx --suite stage-a-unit-test-mlx`
(`python3 file.py -f`), which is the actual invocation pr-test-mlx.yml's
stage-a-unit-test-mlx job uses, not pytest. Bare pytest runs stayed green
throughout because pytest collection never checks for this entry point, so
no node-id diff, pass count, or ULP sweep could have surfaced it.

Re-verified through the lane's own command after the fix: all 17 files in
the suite pass, test_fused_swiglu.py reports 15 passed, 2 skipped -- matching
the pytest baseline exactly.
@jlee5814

Copy link
Copy Markdown
Contributor Author

/tag-and-rerun-ci

The module-level can_fuse(switch_mlp) function decides at patch time
whether a SwitchGLU module is eligible for fusion, a distinct concept
from MetalJitOp.can_fuse's per-op dispatch predicate that it happened
to share a name with. Rename removes the collision; the two class-level
can_fuse methods (FusedMoeCombineKernel, FusedGateQmvSiluMulKernel) and
fused_combine.py's own module-level can_fuse alias are unaffected.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants