diff --git a/docs/benchmark_results/fused-moe-decode-dff-sweep-gb10-2026-07-03.csv b/docs/benchmark_results/fused-moe-decode-dff-sweep-gb10-2026-07-03.csv new file mode 100644 index 000000000..5d9547f6b --- /dev/null +++ b/docs/benchmark_results/fused-moe-decode-dff-sweep-gb10-2026-07-03.csv @@ -0,0 +1,15 @@ +# Fused decode-MoE vs gather_qmm Dff crossover sweep (#626). +# Host: GB10 (DGX Spark), sm_121, CUDA 13.0. MLX pin e9463bb. +# Harness: mlxcel-bench-decode, prompt "Hello, how are you today?", 100 decode +# tokens after a 20-token in-process warmup, median of 3 invocations per side. +# Fallback side forced with MLXCEL_FUSED_MOE_MAX_DFF=1 (gather_qmm); +# fused side forced with MLXCEL_FUSED_MOE_MAX_DFF=20000 (fused two-kernel path). +# decode tok/s. NA = model does not run in the text decode bench harness. +model,dff,quant_mode,fallback_run1,fallback_run2,fallback_run3,fallback_median,fused_run1,fused_run2,fused_run3,fused_median,fused_over_fallback +qwen3.5-35b-a3b-4bit,512,affine,NA,NA,NA,NA,NA,NA,NA,NA,NA +qwen3-30b-a3b-4bit,768,affine,91.00,91.66,90.44,91.00,88.43,91.91,91.34,91.34,1.0037 +lfm2-8b-a1b-4bit,1792,affine,140.65,139.82,140.35,140.35,157.36,157.80,157.76,157.76,1.1240 +gpt-oss-20b-mxfp4,2880,mxfp4,78.93,78.59,77.10,78.59,78.00,78.90,76.57,78.00,0.9925 +phi-3.5-moe-4bit,6400,affine,49.44,51.96,51.09,51.09,53.62,53.41,53.78,53.62,1.0495 +llama-4-scout-17b-4bit,8192,affine,21.23,21.23,21.28,21.23,21.00,21.27,21.08,21.08,0.9929 +mixtral-8x7b-4bit,14336,affine,28.15,28.23,28.30,28.23,27.88,27.89,27.90,27.89,0.9880 diff --git a/docs/benchmark_results/fused-moe-decode-kernel-design.md b/docs/benchmark_results/fused-moe-decode-kernel-design.md index 9d4784a9d..48465a7ae 100644 --- a/docs/benchmark_results/fused-moe-decode-kernel-design.md +++ b/docs/benchmark_results/fused-moe-decode-kernel-design.md @@ -168,7 +168,7 @@ falls back to the proven `gather_qmm` / `SwitchGLU` path automatically. Set | `MLXCEL_FUSED_MOE` | unset (on) | On by default. Set to `0` (also `false`/`off`/`no`, case-insensitive) to force the proven `gather_qmm` / `SwitchGLU` path; any other value, or leaving it unset, keeps the kernel on. | | `MLXCEL_FUSED_MOE_SGY` | 8 | Simdgroups per threadgroup (one output row each). Tune per hardware. | | `MLXCEL_FUSED_MOE_RELU2` | unset (off) | Enable the squared-ReLU (fc1/relu²/fc2) fused path used by nemotron-class experts. Correct but measured performance-neutral on nemotron-h; kept for a future MoE-dominated squared-ReLU model. | -| `MLXCEL_FUSED_MOE_MAX_DFF` | 4096 | Expert-intermediate (Dff) upper bound. Above it, `forward_fused_kernel` declines and the caller falls back to `gather_qmm`. The fused path wins only while `gather_qmm` underutilizes the GPU (small experts); for large experts `gather_qmm` already saturates and the extra dispatch plus global-memory activation staging is a net loss. M1 Ultra measurements: Dff 704..2560 gain +3.5% to +15.4%, Dff 6400 (phi-3.5-moe) loses 5.9%, Dff 14336 (mixtral) loses 21%. On CUDA (GB10) the crossover is much higher, ~13-14k (Dff 768 +60%, 1792 +13%, 6400 +2%, 14336 -2%), so 4096 is conservative there but kept as the shared default. The break-even is hardware-dependent, so the bound is tunable. | +| `MLXCEL_FUSED_MOE_MAX_DFF` | 4096 (Metal) / 8192 (CUDA) | Expert-intermediate (Dff) upper bound. Above it, `forward_fused_kernel` declines and the caller falls back to `gather_qmm`. The fused path wins only while `gather_qmm` underutilizes the GPU (small experts); for large experts `gather_qmm` already saturates and the extra dispatch plus global-memory activation staging is a net loss. The break-even is backend-dependent, so the default is chosen from the live backend (`mlx::core::metal::is_available()`). M1 Ultra measurements: Dff 704..2560 gain +3.5% to +15.4%, Dff 6400 (phi-3.5-moe) loses 5.9%, Dff 14336 (mixtral) loses 21%, so Metal keeps 4096. On CUDA the crossover is higher; re-measured on GB10 under MLX pin e9463bb (#626) it collapsed from the old ~13-14k to ~8000 (Dff 6400 +5%, 8192 break-even, 14336 -1.2%; see the 2026-07-03 addendum), so the CUDA default is 8192. An explicit value overrides both backends. | ### Measured decode gains (M1 Ultra, `MLXCEL_FUSED_MOE=1`) @@ -250,11 +250,13 @@ Forcing the fused path past the cap (`MLXCEL_FUSED_MOE_MAX_DFF=20000`): | 6400 | phi-3.5-moe | +2% | | 14336 | mixtral | −2% | -The CUDA crossover is ~13–14k (vs Metal's ~4096, where phi-3.5-moe already loses -5.9%). The 4096 cap is therefore conservative on CUDA but kept as the shared -default: the meaningful wins are all below it, and the 4096–14k range gains at -most ~2% (phi-3.5-moe) while mixtral (14336) slightly prefers `gather_qmm`. CUDA -users with mid-size experts can raise `MLXCEL_FUSED_MOE_MAX_DFF`. +The CUDA crossover measured here (#319) was ~13–14k (vs Metal's ~4096, where +phi-3.5-moe already loses 5.9%). At that time 4096 was kept as the shared default +even on CUDA. That was re-measured after the #625 MLX pin bump (which moved +`gather_gemm` to JIT and sped the fallback): the CUDA crossover collapsed to +~8000 and the default is now backend-aware, 4096 on Metal and 8192 on CUDA. See +the 2026-07-03 addendum below for the new sweep; CUDA users with larger experts +can still raise or lower `MLXCEL_FUSED_MOE_MAX_DFF`. **Parity caveat.** The kernel accumulates the GEMV in f32 with a different reduction order than `gather_qmm`'s tiling, so it is within the f16 jitter @@ -282,8 +284,9 @@ SwiGLU activation, text-only decode path), phimoe (Phi-3.5-MoE; migrated from its local `SwitchGLU`/`SwitchLinear` to the shared ones; checkpoints pre-stacked under `block_sparse_moe.switch_mlp.{gate,up,down}_proj`; `sanitize_weights` still handles the unstacked `experts.{i}.w1/w2/w3` layout for community checkpoints; the -expert intermediate is 6400, above `MLXCEL_FUSED_MOE_MAX_DFF`, so like mixtral the -kernel declines and decode stays on `gather_qmm`), olmoe (OLMoE; migrated from +expert intermediate is 6400: below the CUDA default (8192) so the kernel dispatches +on CUDA at a measured +5%, above the Metal default (4096) so it declines to +`gather_qmm` on Metal; see the 2026-07-03 addendum), olmoe (OLMoE; migrated from its local `SwitchGLU`/`SwitchLinear` to the shared ones; SwiGLU, softmax-routed with optional `norm_topk_prob`, weights stacked under `switch_mlp.{gate,up,down}_proj` or joined from the per-expert `experts.{i}` layout by `sanitize_weights`; expert @@ -301,3 +304,47 @@ to `gather_qmm`, generating coherent output with no crash or OOM; fused-path throughput and output-parity validation remain blocked on a fitting 4-bit or 8-bit minimax checkpoint). nemotron-h's MoE runs through the separate C++ `fused_moe_forward` and is wired behind `MLXCEL_FUSED_MOE_RELU2` only. + +### 2026-07-03 addendum: GB10 crossover re-measured under MLX pin e9463bb (#626) + +The #625 MLX pin bump to `e9463bb` moved `gather_gemm` to a JIT path, changing +the `gather_qmm` fallback cost, so the CUDA crossover was re-measured on GB10 (DGX +Spark, sm_121, CUDA 13.0). Harness: `mlxcel-bench-decode`, prompt "Hello, how are +you today?", 100 decode tokens after a 20-token warmup, median of 3. Each side is +forced with the env var: `MLXCEL_FUSED_MOE_MAX_DFF=1` selects the `gather_qmm` +fallback, `=20000` selects the fused kernel. decode tok/s: + +| Dff | model | gather_qmm | fused | fused/fallback | +|----:|-------|-----------:|------:|---------------:| +| 768 | qwen3-30b-a3b | 91.00 | 91.34 | 1.00 | +| 1792 | lfm2-8b-a1b | 140.35 | 157.76 | **1.12** | +| 2880 | gpt-oss-20b (mxfp4) | 78.59 | 78.00 | 0.99 | +| 6400 | phi-3.5-moe | 51.09 | 53.62 | **1.05** | +| 8192 | llama-4-scout-17b | 21.23 | 21.08 | 0.99 | +| 14336 | mixtral-8x7b | 28.23 | 27.89 | 0.99 | + +Findings: + +- **The crossover moved down from ~13-14k to ~8000, but stays well above Metal's + 4096.** The fused path clearly wins at Dff 6400 (phi-3.5-moe +5%) and below + (lfm2 +12%), is break-even at 8192 (llama-4-scout, within run-to-run noise: one + of three fused runs matched the fallback), and loses at 14336 (mixtral -1.2%). + Interpolating phi (+5%) and llama-4-scout (break-even) puts the ratio=1.0 + crossover at ~8000. +- **The old +55% qwen3-30b-a3b win is gone**, not because the fused kernel + regressed but because the JIT `gather_gemm` made the fallback far faster for the + pathological 128-expert config (old fallback 58.2 -> now 91.0 tok/s). The fused + path itself is unchanged there (~91). Low-expert configs are unaffected: lfm2 is + still +12%, matching the #319 measurement. +- **gpt-oss-20b is a control**: it is `mxfp4`, which `forward_fused_kernel` + declines (non-affine), so the env var is a no-op and both sides match (0.99). + qwen3.5-35b-a3b (Dff 512, multimodal `Qwen3_5MoeForConditionalGeneration`) does + not run in the text decode bench harness and is below 4096 regardless. + +**Default set to 8192 on CUDA** (the break-even boundary, rounded down from the +14336 regression), keeping Metal at 4096. Models with an expert intermediate in +`(4096, 8192]` now take the fused kernel by default on CUDA: phi-3.5-moe (Dff +6400, 51.09 -> 53.62 tok/s, +5%) and llama-4-scout-17b (Dff 8192, 21.23 -> 21.08 +tok/s, break-even). Mixtral (Dff 14336) stays on `gather_qmm`. The env var +`MLXCEL_FUSED_MOE_MAX_DFF` still overrides both backends. Raw per-run data: +`fused-moe-decode-dff-sweep-gb10-2026-07-03.csv`. diff --git a/docs/environment-variables.md b/docs/environment-variables.md index b67def5a6..62baf778d 100644 --- a/docs/environment-variables.md +++ b/docs/environment-variables.md @@ -208,7 +208,7 @@ recommended as normal deployment settings. | `MLXCEL_PIPELINE_GRANULARITY` | `off`, `layer`, `block:N` | `off` | Inserts layer-boundary async-eval hints for pipeline experiments. | | `MLXCEL_FUSED_MOE` | `0`/`false`/`off`/`no` disable; any other value or unset enables | on | Fused single-token decode-MoE kernel (#268), on by default since #282 (Metal) and #319 (CUDA, via `mx.fast.cuda_kernel`); validated on M1 Ultra, M5, and GB10. Set to `0` to force the proven `gather_qmm`/`SwitchGLU` path. Active for qwen3_moe, qwen3_next, dots.llm1, gemma4, qwen2_moe, mixtral, phimoe, lfm2, qwen3_vl_moe, and olmoe decode. | | `MLXCEL_FUSED_MOE_SGY` | `1`-`32` | `8` | Simdgroups (Metal) / warps-per-block (CUDA) per threadgroup for the fused decode-MoE kernel; tune per hardware. | -| `MLXCEL_FUSED_MOE_MAX_DFF` | positive int | `4096` | Expert-intermediate (Dff) upper bound for the fused path; above it the caller falls back to `gather_qmm`. The fused path wins only while `gather_qmm` underutilizes the GPU (small experts). The break-even is hardware-dependent: ~4096 on M1 Ultra, ~13-14k on GB10/CUDA. 4096 is the conservative shared default; raise it on CUDA for mid-size experts. | +| `MLXCEL_FUSED_MOE_MAX_DFF` | positive int | `4096` (Metal) / `8192` (CUDA) | Expert-intermediate (Dff) upper bound for the fused path; above it the caller falls back to `gather_qmm`. The fused path wins only while `gather_qmm` underutilizes the GPU (small experts), so the break-even is backend-dependent and the default is chosen from the live backend: `4096` on Metal (M1 Ultra tuning) and `8192` on CUDA (GB10 re-measured under MLX pin e9463bb, #626; fused wins through Dff 6400 and is break-even at 8192). An explicit value overrides the default on both backends: lower it to force `gather_qmm` sooner, raise it (e.g. `20000`) to force the fused kernel on larger experts such as mixtral (Dff 14336, where it is a slight net loss). | | `MLXCEL_FUSED_MOE_RELU2` | presence enables | off | Enables the squared-ReLU fused MoE path for nemotron-class experts; performance-neutral on nemotron-h, kept for a future MoE-dominated squared-ReLU model. | | `MLXCEL_FUSED_QK_NORM` | `1`/`true`/`on`/`yes` enable; any other value or unset disables | off (opt-in) | Fused single-token QKV projection + Q/K RMSNorm + RoPE kernel (#326) for Qwen3 and Qwen3-MoE decode. Opt-in: set to `1` to enable. Matches the graph path within RMS < 5e-3 (the reduction is over the transpose-invariant head_dim axis), but greedy temp-0 is not byte-identical over long generation; on CUDA the graph path is itself non-deterministic run-to-run from GPU FP-reduction order, while the fused path is deterministic, so its output stays inside the graph baseline's own envelope. The kernel cuts Rust/C++ FFI crossings rather than MLX op count, so it does not speed up the GPU/bandwidth-bound decode loop: on M1 Ultra it measured 1 to 3.4% slower (qwen3-0.6b 275 vs 284, qwen3-8b 82.3 vs 83.2 tok/s); on GB10/CUDA (SM 12.1) it is also slower (qwen3-0.6b 0.96x, qwen3-8b ~1.0x, qwen3-30b-a3b 0.92x fused/graph; see `docs/benchmark_results/fused-qk-norm-decode-gb10.md`), so there is no per-backend win. Ships as a reusable shared primitive for the deferred QK-norm families and stays opt-in (default off) on every measured backend (M1 Ultra, M5 Max, GB10/CUDA), mirroring the opt-in `MLXCEL_FUSED_MOE_RELU2`. Active only when `l == 1` (decode) and weights are quantized. | | `MLXCEL_FUSED_XIELU` | `0`/`false`/`off`/`no` disable; any other value or unset enables | on | Fused single-launch Metal xIELU kernel for the Apertus MLP activation (#409), on by default since the M5 Max validation. `MLP::forward` routes through one Metal dispatch covering the ~11 elementwise ops in `apertus_xielu` (square, minimum, expm1, where, and neighbors) instead of the per-op graph. Greedy temp-0 decode is byte-identical to the elementwise path on Apple Silicon: every intermediate stays in the input dtype (bf16) and the kernel reproduces MLX's `expm1f` exactly. Measured decode speedup on M1 Ultra (+2.7%, Apertus-8B 83.4 to 85.7 tok/s) and M5 Max (+1.9%, 112.0 to 114.2 tok/s), with no regression. Set to `0` to force the elementwise path. On non-Metal back-ends the FFI falls back to an equivalent elementwise graph, so the flag is safe to set everywhere. Apertus only; no other model family is affected. | diff --git a/src/lib/mlxcel-core/cpp/mlx_cxx_bridge.h b/src/lib/mlxcel-core/cpp/mlx_cxx_bridge.h index b003e146f..6cabb5083 100644 --- a/src/lib/mlxcel-core/cpp/mlx_cxx_bridge.h +++ b/src/lib/mlxcel-core/cpp/mlx_cxx_bridge.h @@ -1558,6 +1558,11 @@ void fused_gated_delta_decode_step( // Check if GatedDeltaNet Metal kernel is available (Metal GPU only) bool gated_delta_kernel_available(); +// True when the MLX Metal backend is available at runtime (macOS Apple +// Silicon); false on CUDA-only and CPU-only builds. Mirrors the +// metal::is_available() gate used to pick the Metal vs CUDA kernel port. #626. +bool metal_is_available(); + // Start/stop Metal GPU capture. Requires the process to run under // `MTL_CAPTURE_ENABLED=1`; otherwise Metal drops the capture silently. void metal_start_capture(rust::Str path); diff --git a/src/lib/mlxcel-core/cpp/mlx_cxx_kernels.cpp b/src/lib/mlxcel-core/cpp/mlx_cxx_kernels.cpp index 8fbbf6465..5a70b1bc9 100644 --- a/src/lib/mlxcel-core/cpp/mlx_cxx_kernels.cpp +++ b/src/lib/mlxcel-core/cpp/mlx_cxx_kernels.cpp @@ -769,6 +769,19 @@ bool gated_delta_kernel_available() { #endif } +// True when the MLX Metal backend is available at runtime. Mirrors the +// metal::is_available() gate used to pick the CUDA vs Metal kernel port +// (run_fused_moe_two_kernel, bitlinear_matmul) so Rust callers can choose a +// backend-specific default without a compile-time cfg. False on CUDA-only and +// CPU-only builds. See issue #626. +bool metal_is_available() { +#ifdef __APPLE__ + return mlx::core::metal::is_available(); +#else + return false; +#endif +} + // Start a GPU trace capture and stop an active one. Mirrors the Python // `mx.metal.start_capture` / `mx.metal.stop_capture` API so a mlxcel // decode run can emit the same `.gputrace` files that `mlx-lm` produces diff --git a/src/lib/mlxcel-core/src/lib.rs b/src/lib/mlxcel-core/src/lib.rs index 7e5e53d1e..1ee13aa96 100644 --- a/src/lib/mlxcel-core/src/lib.rs +++ b/src/lib/mlxcel-core/src/lib.rs @@ -1801,6 +1801,13 @@ mod ffi { /// Check whether the current default device is GPU fn is_gpu_available() -> bool; + /// True when the MLX Metal backend is available at runtime (macOS + /// Apple Silicon). False on CUDA-only and CPU-only builds. Mirrors the + /// `mlx::core::metal::is_available()` gate that picks the Metal vs CUDA + /// kernel port, so Rust callers can choose a backend-specific default + /// without a compile-time cfg. See issue #626. + fn metal_is_available() -> bool; + /// Fused sampling: temperature + top-k + top-p + min-p + categorical /// in a single C++ call to minimize FFI round-trips. /// Input: 2D logits [batch, vocab] (already sliced, penalties applied) diff --git a/src/models/switch_layers.rs b/src/models/switch_layers.rs index 79a6b7f57..67814b858 100644 --- a/src/models/switch_layers.rs +++ b/src/models/switch_layers.rs @@ -61,6 +61,42 @@ pub(crate) fn fused_moe_enabled_from(value: Option<&str>) -> bool { } } +/// Metal default for the fused decode-MoE expert-intermediate (Dff) upper bound. +/// +/// The fused two-kernel path wins only while `gather_qmm` underutilizes the GPU +/// (small experts); above this bound the caller falls back to `gather_qmm`. +/// Tuned on M1 Ultra, where phi-3.5-moe (Dff 6400) already regresses. See +/// `docs/benchmark_results/fused-moe-decode-kernel-design.md`. +const FUSED_MOE_MAX_DFF_METAL: i32 = 4096; + +/// CUDA default for the same bound. At batch=1 `gather_qmm` leaves far more of +/// the GPU idle than on Metal, so the fused path keeps winning to a larger +/// expert size. Re-measured on GB10 (DGX Spark, sm_121) under MLX pin e9463bb +/// (#626, after the #625 pin bump moved `gather_gemm` to JIT and sped the +/// fallback): fused wins through Dff 6400 (phi-3.5-moe +5%, lfm2 +12%), is +/// break-even at 8192 (llama-4-scout, within run-to-run noise) and loses at +/// 14336 (mixtral -1.2%), so the crossover interpolates to ~8000. Set to the +/// 8192 break-even boundary, which stays well below the 14336 regression while +/// capturing the phi-3.5-moe / llama-4-scout mid-size experts that the old 4096 +/// default silently declined. The env var overrides this on both backends. +const FUSED_MOE_MAX_DFF_CUDA: i32 = 8192; + +/// Resolve the fused-MoE Dff upper bound. An explicit positive +/// `MLXCEL_FUSED_MOE_MAX_DFF` wins on both backends; otherwise the default is +/// backend-specific (Metal keeps the conservative 4096, CUDA uses the higher +/// measured crossover). Split out as a pure function so the backend-aware +/// default is unit-testable without mutating process-global env or querying the +/// live device; `metal_available` mirrors `mlx::core::metal::is_available()`. +pub(crate) fn fused_moe_max_dff_from(env: Option<&str>, metal_available: bool) -> i32 { + env.and_then(|s| s.parse::().ok()) + .filter(|v| *v > 0) + .unwrap_or(if metal_available { + FUSED_MOE_MAX_DFF_METAL + } else { + FUSED_MOE_MAX_DFF_CUDA + }) +} + /// Per-expert 3D linear layer (falls back to gather_mm for non-quantized models) /// Supports affine, mxfp4, nvfp4, and mxfp8 quantization modes. pub enum SwitchLinear { @@ -394,16 +430,14 @@ impl SwitchGLU { let din = gw_shape[2] * (32 / gate.bits); // Large experts: gather_qmm already saturates the GPU, so the two-kernel // fused path's all-cores advantage disappears and its extra dispatch + - // global-memory activation staging becomes a net loss. Measured on M1 - // Ultra: wins for small experts (Dff 704..2560: +3.5% to +15.4%), loses - // for large ones (phi-3.5-moe Dff 6400: -5.9%, mixtral Dff 14336: -21%). - // Decline above the break-even so the caller falls back to gather_qmm. - // The break-even is hardware-dependent, so the bound is env-tunable. - let max_dff = std::env::var("MLXCEL_FUSED_MOE_MAX_DFF") - .ok() - .and_then(|s| s.parse::().ok()) - .filter(|v| *v > 0) - .unwrap_or(4096); + // global-memory activation staging becomes a net loss. The break-even is + // backend-dependent (Metal ~4096, CUDA ~8192; see FUSED_MOE_MAX_DFF_*), + // so the default follows the live backend and the env var overrides both. + // Decline above the bound so the caller falls back to gather_qmm. + let max_dff = fused_moe_max_dff_from( + std::env::var("MLXCEL_FUSED_MOE_MAX_DFF").ok().as_deref(), + mlxcel_core::metal_is_available(), + ); if dff > max_dff { return None; } @@ -659,6 +693,32 @@ mod tests { } } + #[test] + fn fused_moe_max_dff_default_is_backend_aware_and_env_overrides() { + // Unset: Metal keeps the conservative 4096; CUDA (no Metal) uses the + // higher measured GB10 crossover (#626). + assert_eq!(fused_moe_max_dff_from(None, true), FUSED_MOE_MAX_DFF_METAL); + assert_eq!(fused_moe_max_dff_from(None, false), FUSED_MOE_MAX_DFF_CUDA); + // The CUDA default must exceed Metal's, which is the whole point of #626. + assert!(FUSED_MOE_MAX_DFF_CUDA > FUSED_MOE_MAX_DFF_METAL); + // An explicit positive value overrides the default on both backends. + for metal in [true, false] { + assert_eq!(fused_moe_max_dff_from(Some("16384"), metal), 16384); + assert_eq!(fused_moe_max_dff_from(Some("1"), metal), 1); + } + // Non-positive or unparseable values fall back to the backend default. + for bad in ["0", "-1", "abc", "", " 4096"] { + assert_eq!( + fused_moe_max_dff_from(Some(bad), true), + FUSED_MOE_MAX_DFF_METAL + ); + assert_eq!( + fused_moe_max_dff_from(Some(bad), false), + FUSED_MOE_MAX_DFF_CUDA + ); + } + } + #[test] fn switch_linear_stacks_individual_quantized_experts() { // Per-expert affine 4-bit layout (Qwen1.5-MoE / Qwen2-MoE checkpoints):