diff --git a/.isort.cfg b/.isort.cfg
index 83509515..f34d926e 100644
--- a/.isort.cfg
+++ b/.isort.cfg
@@ -1,3 +1,4 @@
 [settings]
 profile=black
 known_first_party=sglang
+known_third_party=sgl_kernel
diff --git a/benchmark/bench_fused_moe_fp8.py b/benchmark/bench_fused_moe_fp8.py
new file mode 100644
index 00000000..e3a7ecfe
--- /dev/null
+++ b/benchmark/bench_fused_moe_fp8.py
@@ -0,0 +1,139 @@
+"""Small opt-in benchmark for the FP8 split-activation MoE path."""
+
+import argparse
+import ctypes
+import importlib.util
+import statistics
+from pathlib import Path
+
+import torch
+
+# Xe2 executes these FP8 E4M3 weights with BF16 activations through W8A16.
+
+
+def _preload_fp8_instances():
+    spec = importlib.util.find_spec("sgl_kernel")
+    if spec is None or spec.submodule_search_locations is None:
+        return
+    for package_dir in spec.submodule_search_locations:
+        for path in sorted(
+            Path(package_dir).glob("libsgl-ops-sycl-GroupGemmFp8Xe20_inst*.so")
+        ):
+            ctypes.CDLL(str(path), mode=ctypes.RTLD_GLOBAL)
+
+
+_preload_fp8_instances()
+from sgl_kernel import fused_experts
+
+PROFILES = {
+    "qwen-tp4": (3584, 1280, 8, 8),
+    "deepseek-tp8": (7168, 512, 8, 8),
+    "qwen35-tp4": (2048, 128, 16, 10),
+}
+TOKENS = [1, 32, 2048]
+DEFAULT_WARMUP = 20
+DEFAULT_REPETITIONS = 30
+DEFAULT_INNER_REPETITIONS = 10
+
+
+def _make_case(hidden, intermediate, experts, topk, tokens):
+    torch.manual_seed(0)
+    torch.xpu.manual_seed_all(0)
+    hidden_states = torch.randn((tokens, hidden), dtype=torch.bfloat16, device="xpu")
+    w1_bf16 = torch.randn(
+        (experts, 2 * intermediate, hidden), dtype=torch.bfloat16, device="xpu"
+    )
+    w2_bf16 = torch.randn(
+        (experts, hidden, intermediate), dtype=torch.bfloat16, device="xpu"
+    )
+    w1 = w1_bf16.to(torch.float8_e4m3fn)
+    w2 = w2_bf16.to(torch.float8_e4m3fn)
+    del w1_bf16, w2_bf16
+    w1_scale = torch.ones(
+        (experts, 2 * intermediate // 128, hidden // 128),
+        dtype=torch.float32,
+        device="xpu",
+    )
+    w2_scale = torch.ones(
+        (experts, hidden // 128, intermediate // 128),
+        dtype=torch.float32,
+        device="xpu",
+    )
+    topk_ids = (
+        torch.arange(tokens * topk, device="xpu", dtype=torch.int32).reshape(
+            tokens, topk
+        )
+        % experts
+    )
+    topk_weights = torch.full(
+        (tokens, topk), 1.0 / topk, dtype=torch.float32, device="xpu"
+    )
+    return hidden_states, w1, w2, topk_weights, topk_ids, w1_scale, w2_scale
+
+
+def _run_case(profile_name, tokens, warmup, repetitions, inner_repetitions):
+    hidden, intermediate, experts, topk = PROFILES[profile_name]
+    args = _make_case(hidden, intermediate, experts, topk, tokens)
+
+    def run():
+        fused_experts(
+            *args[:5],
+            activation="silu",
+            use_fp8_w8a8=True,
+            w1_scale=args[5],
+            w2_scale=args[6],
+        )
+
+    for _ in range(warmup):
+        run()
+    torch.xpu.synchronize()
+    samples = []
+    for _ in range(repetitions):
+        start = torch.xpu.Event(enable_timing=True)
+        end = torch.xpu.Event(enable_timing=True)
+        start.record()
+        for _ in range(inner_repetitions):
+            run()
+        end.record()
+        torch.xpu.synchronize()
+        samples.append(start.elapsed_time(end) / inner_repetitions)
+    samples.sort()
+    median = statistics.median(samples)
+    del args
+    torch.xpu.empty_cache()
+    return hidden, intermediate, experts, topk, median
+
+
+def main():
+    parser = argparse.ArgumentParser()
+    parser.add_argument("--tokens", type=int, nargs="+", default=TOKENS)
+    parser.add_argument("--warmup", type=int, default=DEFAULT_WARMUP)
+    parser.add_argument("--repetitions", type=int, default=DEFAULT_REPETITIONS)
+    parser.add_argument(
+        "--inner-repetitions", type=int, default=DEFAULT_INNER_REPETITIONS
+    )
+    args = parser.parse_args()
+    print("profile tokens hidden intermediate experts topk median_ms", flush=True)
+    for profile_name in PROFILES:
+        for tokens in args.tokens:
+            hidden, intermediate, experts, topk, median = _run_case(
+                profile_name,
+                tokens,
+                args.warmup,
+                args.repetitions,
+                args.inner_repetitions,
+            )
+            print(
+                profile_name,
+                tokens,
+                hidden,
+                intermediate,
+                experts,
+                topk,
+                f"{median:.3f}",
+                flush=True,
+            )
+
+
+if __name__ == "__main__":
+    main()
diff --git a/benchmark/bench_moe_fp8_w8a16_grouped_gemm.py b/benchmark/bench_moe_fp8_w8a16_grouped_gemm.py
new file mode 100644
index 00000000..59ffdb66
--- /dev/null
+++ b/benchmark/bench_moe_fp8_w8a16_grouped_gemm.py
@@ -0,0 +1,191 @@
+"""Benchmark the Xe2 FP8 W8A16 grouped GEMM against matched references.
+
+This is the op-level companion to ``bench_fused_moe_fp8.py``, following the
+same split as the W4A16 benchmarks. It compares scalar SGL/vLLM kernels on
+identical logical FP8 weights and reports the SGL 128x128 block-scale path as
+a separate quantization contract.
+
+Run:
+  python benchmark/bench_moe_fp8_w8a16_grouped_gemm.py
+  SGL_MOE_BENCH_FULL_SHAPES=1 python benchmark/bench_moe_fp8_w8a16_grouped_gemm.py
+"""
+
+import os
+
+import sgl_kernel  # noqa: F401 - registers torch.ops.sgl_kernel
+import torch
+import triton
+
+
+def _import_vllm_grouped_gemm():
+    import vllm_xpu_kernels._moe_C  # noqa: F401
+    import vllm_xpu_kernels._xpu_C  # noqa: F401
+    from vllm_xpu_kernels.fused_moe_interface import cutlass_grouped_gemm_xe2
+
+    return cutlass_grouped_gemm_xe2
+
+
+try:
+    _vllm_grouped_gemm = _import_vllm_grouped_gemm()
+    VLLM_AVAILABLE = True
+except Exception as exc:
+    _vllm_grouped_gemm = None
+    VLLM_AVAILABLE = False
+    print(f"[vLLM provider disabled: {type(exc).__name__}: {exc}]")
+
+
+FP8_MAX = 448.0
+BLOCK_SIZE = 128
+QUICK_SHAPES = [
+    # (experts, avg_m, N, K): controlled scalar dispatch boundaries.
+    (8, 4, 1024, 1024),
+    (8, 16, 1024, 1024),
+    (8, 64, 1024, 2048),
+    (8, 129, 1024, 1024),
+]
+FULL_SHAPES = QUICK_SHAPES + [
+    # One-rank model GEMM1/GEMM2 shapes used during optimization.
+    (8, 64, 2560, 3584),
+    (8, 64, 3584, 1280),
+    (8, 64, 1024, 7168),
+    (8, 64, 7168, 512),
+    (16, 64, 256, 2048),
+    (16, 64, 2048, 128),
+]
+BENCH_SHAPES = (
+    FULL_SHAPES if os.environ.get("SGL_MOE_BENCH_FULL_SHAPES") == "1" else QUICK_SHAPES
+)
+
+
+def _quantize_scalar(weight):
+    scale = weight.float().abs().amax().clamp_min(1e-12) / FP8_MAX
+    quantized = (
+        (weight.float() / scale).clamp(-FP8_MAX, FP8_MAX).to(torch.float8_e4m3fn)
+    )
+    return quantized, scale
+
+
+def _quantize_block(weight):
+    experts, rows, columns = weight.shape
+    assert rows % BLOCK_SIZE == 0 and columns % BLOCK_SIZE == 0
+    blocked = weight.float().reshape(
+        experts,
+        rows // BLOCK_SIZE,
+        BLOCK_SIZE,
+        columns // BLOCK_SIZE,
+        BLOCK_SIZE,
+    )
+    scales = blocked.abs().amax((2, 4), keepdim=True).clamp_min(1e-12) / FP8_MAX
+    quantized = (blocked / scales).clamp(-FP8_MAX, FP8_MAX).to(torch.float8_e4m3fn)
+    return (
+        quantized.reshape_as(weight),
+        scales.reshape(experts, rows // BLOCK_SIZE, columns // BLOCK_SIZE),
+    )
+
+
+def _make_inputs(experts, avg_m, gemm_n, gemm_k, provider):
+    torch.manual_seed(0)
+    torch.xpu.manual_seed_all(0)
+    total_m = experts * avg_m
+    activations = (
+        torch.randn((total_m, gemm_k), device="xpu", dtype=torch.bfloat16) / 16
+    )
+    weight_bf16 = (
+        torch.randn((experts, gemm_n, gemm_k), device="xpu", dtype=torch.bfloat16) / 16
+    )
+    rows_per_expert = torch.full((experts,), avg_m, device="xpu", dtype=torch.int32)
+    output = torch.empty((total_m, gemm_n), device="xpu", dtype=torch.bfloat16)
+
+    if provider == "sgl_block":
+        weights, scales = _quantize_block(weight_bf16)
+    else:
+        weights, scale = _quantize_scalar(weight_bf16)
+        scales = torch.full(
+            (experts, 1), scale.item(), device="xpu", dtype=torch.float32
+        )
+    del weight_bf16
+
+    if provider == "vllm_scalar":
+        weights = weights.transpose(-1, -2).contiguous()
+        scales = scales.flatten()
+
+    return activations, weights, scales, rows_per_expert, output
+
+
+def _run_sgl(inputs, experts):
+    activations, weights, scales, rows_per_expert, output = inputs
+    torch.ops.sgl_kernel.moe_grouped_mm_nt_xe20_fp8_w8a16(
+        output,
+        activations,
+        weights,
+        scales,
+        None,
+        rows_per_expert,
+        experts,
+    )
+    return output
+
+
+def _run_vllm(inputs, experts, gemm_n, gemm_k):
+    activations, weights, scales, rows_per_expert, output = inputs
+    _vllm_grouped_gemm(
+        activations,
+        weights,
+        scales,
+        None,
+        output,
+        rows_per_expert,
+        gemm_n,
+        gemm_k,
+        experts,
+    )
+    return output
+
+
+PROVIDERS = ["sgl_scalar", "sgl_block"]
+PROVIDER_NAMES = ["SGL scalar W8A16", "SGL 128x128 block W8A16"]
+STYLES = [("green", "-"), ("blue", "-")]
+if VLLM_AVAILABLE:
+    PROVIDERS.insert(1, "vllm_scalar")
+    PROVIDER_NAMES.insert(1, "vLLM scalar W8A16")
+    STYLES.insert(1, ("red", "-"))
+
+
+@triton.testing.perf_report(
+    triton.testing.Benchmark(
+        x_names=["experts", "avg_m", "gemm_n", "gemm_k"],
+        x_vals=BENCH_SHAPES,
+        line_arg="provider",
+        line_vals=PROVIDERS,
+        line_names=PROVIDER_NAMES,
+        styles=STYLES,
+        ylabel="Time (ms)",
+        plot_name="moe-fp8-w8a16-grouped-gemm",
+        args={},
+    )
+)
+def benchmark(experts, avg_m, gemm_n, gemm_k, provider):
+    inputs = _make_inputs(experts, avg_m, gemm_n, gemm_k, provider)
+    if provider == "vllm_scalar":
+        run = lambda: _run_vllm(inputs, experts, gemm_n, gemm_k)
+    else:
+        run = lambda: _run_sgl(inputs, experts)
+    return triton.testing.do_bench(run, warmup=100, rep=300)
+
+
+def _check_scalar_reference():
+    if not VLLM_AVAILABLE:
+        return
+    shape = QUICK_SHAPES[1]
+    sgl_inputs = _make_inputs(*shape, "sgl_scalar")
+    vllm_inputs = _make_inputs(*shape, "vllm_scalar")
+    sgl_output = _run_sgl(sgl_inputs, shape[0])
+    vllm_output = _run_vllm(vllm_inputs, shape[0], shape[2], shape[3])
+    torch.xpu.synchronize()
+    torch.testing.assert_close(sgl_output, vllm_output, rtol=5e-2, atol=5e-2)
+    print("[correctness] SGL/vLLM scalar outputs match", flush=True)
+
+
+if __name__ == "__main__":
+    _check_scalar_reference()
+    benchmark.run(print_data=True)
diff --git a/include/sgl_kernel_ops.h b/include/sgl_kernel_ops.h
index b5b75b23..c5bbb3a2 100644
--- a/include/sgl_kernel_ops.h
+++ b/include/sgl_kernel_ops.h
@@ -712,6 +712,18 @@ void moe_grouped_mm_nt_xe20_w4a16(
     bool is_int4,
     const int64_t group_size);
 
+// FP8 weight-only MoE grouped GEMM. Activations are BF16, weights are FP8
+// E4M3, and weight_scales is [E, 1]/[E, 2] for per-expert scalar scales or
+// [E, ceil(N/128), K/128] for 128x128 block scales.
+void moe_grouped_mm_nt_xe20_fp8_w8a16(
+    torch::Tensor& output,
+    const torch::Tensor& activations,
+    const torch::Tensor& weights,
+    const torch::Tensor& weight_scales,
+    const std::optional<at::Tensor>& bias,
+    const torch::Tensor& total_rows_for_experts,
+    const int64_t n_experts);
+
 void prepare_moe_input(
     const torch::Tensor& topk_ids,
     torch::Tensor& expert_offsets,
@@ -724,6 +736,13 @@ void prepare_moe_input(
     const int64_t n,
     const int64_t k);
 
+void prepare_moe_input_small(
+    const torch::Tensor& input,
+    const torch::Tensor& topk_ids,
+    torch::Tensor& expert_counts,
+    torch::Tensor& output_permutation,
+    torch::Tensor& output);
+
 void ep_moe_pre_reorder(
     torch::Tensor input,
     torch::Tensor gateup_input,
diff --git a/python/sgl_kernel/moe.py b/python/sgl_kernel/moe.py
index a9afb56e..03035482 100755
--- a/python/sgl_kernel/moe.py
+++ b/python/sgl_kernel/moe.py
@@ -370,7 +370,13 @@ def cutlass_fp4_group_mm(
 
 
 _MOE_WS_HEADROOM = 1.1
+_MOE_SMALL_PREPARE_MAX_ROUTES = 64
+_MOE_SMALL_PREPARE_MAX_TOPK = 16
+_MOE_SMALL_PREPARE_MAX_ELEMENTS = 81920
 _moe_ws_cache: Dict[Tuple[str, torch.device], torch.Tensor] = {}
+_moe_ws_view_cache: Dict[
+    Tuple[str, torch.device], Tuple[torch.Tensor, tuple, torch.Tensor]
+] = {}
 
 
 def _get_moe_ws(
@@ -397,7 +403,62 @@ def _get_moe_ws(
         new_numel = max(numel, int(numel * _MOE_WS_HEADROOM))
         cur = torch.empty(new_numel, dtype=dtype, device=device)
         _moe_ws_cache[key] = cur
-    return cur.narrow(0, 0, numel).view(shape)
+    cached_view = _moe_ws_view_cache.get(key)
+    shape = tuple(shape)
+    if cached_view is None or cached_view[0] is not cur or cached_view[1] != shape:
+        view = cur.narrow(0, 0, numel).view(shape)
+        _moe_ws_view_cache[key] = (cur, shape, view)
+        return view
+    return cached_view[2]
+
+
+def _should_use_small_moe_prepare(
+    num_tokens: int, topk: int, hidden_dims: int, num_experts: int
+) -> bool:
+    routed_rows = num_tokens * topk
+    return (
+        1 <= topk <= _MOE_SMALL_PREPARE_MAX_TOPK
+        and routed_rows <= min(num_experts, _MOE_SMALL_PREPARE_MAX_ROUTES)
+        and routed_rows * hidden_dims <= _MOE_SMALL_PREPARE_MAX_ELEMENTS
+    )
+
+
+def _validate_fp8_weight_scale(
+    scale: torch.Tensor,
+    weights: torch.Tensor,
+    name: str,
+    allow_scalar: bool,
+) -> None:
+    """Validate an FP8 expert scale tensor against its physical weight shape."""
+    assert scale.dtype == torch.float32, f"{name} must be float32"
+    assert scale.ndim in (
+        (2, 3) if allow_scalar else (3,)
+    ), f"{name} must be 3D block scales or 2D scalar scales"
+    assert scale.shape[0] == weights.shape[0], (
+        f"{name} expert dimension {scale.shape[0]} must match weights "
+        f"expert dimension {weights.shape[0]}"
+    )
+    if scale.ndim == 2:
+        assert allow_scalar, f"{name} scalar scales are not supported for this FP8 path"
+        expected_columns = 2 if name == "w1_scale" else 1
+        assert scale.shape[1] in (1, expected_columns), (
+            f"{name} scalar scale shape must be [E, 1] or "
+            f"[E, {expected_columns}], got {tuple(scale.shape)}"
+        )
+        return
+
+    expected_shape = (
+        weights.shape[0],
+        (weights.shape[1] + 127) // 128,
+        (weights.shape[2] + 127) // 128,
+    )
+    assert tuple(scale.shape) == expected_shape, (
+        f"{name} block scales must have shape [E, ceil(N/128), ceil(K/128)] "
+        f"={expected_shape}, got {tuple(scale.shape)}"
+    )
+    assert (
+        weights.shape[2] % 128 == 0
+    ), f"{name} block scales require K divisible by 128, got K={weights.shape[2]}"
 
 
 def fused_experts(
@@ -442,8 +503,10 @@ def fused_experts(
     - b2 (Optional[torch.Tensor]): Optional bias for w2.
     - inplace (bool): If True, perform operations in-place to save memory. Defaults to False.
     - activation (str): The activation function to use ('silu' or 'gelu'). Defaults to 'silu'.
-    - use_fp8_w8a8 (bool): If True, use fp8 arithmetic to compute the inner
-        products for w1 and w2. Defaults to False.
+    - use_fp8_w8a8 (bool): If True, use FP8 E4M3 expert weights from a W8A8
+        checkpoint. Xe2 currently has no native FP8-A MoE kernel, so this
+        path falls back to BF16 activations and the W8A16 operator. Defaults
+        to False.
     - use_mxfp4_w4a16 (bool): If True, w1 and w2 are in MXFP4 packed format
         (int8 or uint8, two E2M1 nibbles per byte) with corresponding E8M0 block
         scales supplied via w1_scale and w2_scale. Scales may be represented
@@ -483,12 +546,14 @@ def fused_experts(
         support, applied to the intermediate activation before GEMM2 to
         match the K-dim sort applied to w2 at weight-load time. Only valid
         with use_int4_w4a16=True.
-    - a1_scale (Optional[torch.Tensor]): Optional scale to be used for
-        a1.
-    - a2_scale (Optional[torch.Tensor]): Optional scale to be used for
-        a2.
-    - block_shape: (Optional[List[int]]): Optional block size for block-wise
-        quantization.
+    - a1_scale (Optional[torch.Tensor]): Reserved for a future prequantized
+        FP8 activation input. It is currently rejected because the Xe2
+        fallback consumes BF16 activations.
+    - a2_scale (Optional[torch.Tensor]): Reserved for a future prequantized
+        FP8 activation input. It is currently rejected because the Xe2
+        fallback consumes BF16 activations.
+    - block_shape: (Optional[List[int]]): Weight block size metadata. FP8
+        block scales must use [128, 128]; the value is validated when supplied.
     - no_combine (bool): If True, skip the combine step. Defaults to False.
     - routed_scaling_factor (Optional[float]): Optional scaling factor for routed tokens, used by Llama4 only.
     - gemm1_alpha (Optional[float]): Optional gemm1_alpha for the activation
@@ -502,15 +567,50 @@ def fused_experts(
     - torch.Tensor: The output tensor after applying the MoE layer.
     """
 
-    assert use_fp8_w8a8 is False, "current MoE does not support use_fp8_w8a8"
-    assert a1_scale is None, "current MoE does not support a1_scale"
-    assert a2_scale is None, "current MoE does not support a2_scale"
-    assert block_shape is None, "current MoE does not support block_shape"
-    assert activation in (
-        "silu",
-        "gelu",
-        "relu2",
-    ), f"Only silu, gelu and relu2 are supported but got {activation}"
+    use_fp8_weight = use_fp8_w8a8
+    assert a1_scale is None, (
+        "prequantized FP8 activation input is not supported: " "a1_scale must be None"
+    )
+    assert a2_scale is None, (
+        "prequantized FP8 activation input is not supported: " "a2_scale must be None"
+    )
+    if block_shape is not None:
+        assert use_fp8_weight, "block_shape is only supported for FP8 MoE paths"
+        assert list(block_shape) == [
+            128,
+            128,
+        ], "FP8 MoE currently supports only block_shape=[128, 128]"
+    if use_fp8_weight:
+        assert activation in ("silu", "gelu", "relu2"), (
+            "FP8 MoE supports silu, gelu, relu2, GPT-OSS SwiGLU, and "
+            "DeepSeek-V4 clamped SwiGLU only"
+        )
+        assert (
+            w1_g_idx_perm is None and w2_g_idx_perm is None
+        ), "w1_g_idx_perm/w2_g_idx_perm are only supported by the INT4 W4A16 path"
+        if activation == "gelu" or activation == "relu2":
+            assert (
+                gemm1_alpha is None and gemm1_limit is None and swiglu_limit is None
+            ), f"{activation} cannot be combined with a SwiGLU alpha or clamp"
+        elif gemm1_alpha is not None:
+            assert gemm1_limit is not None and swiglu_limit is None, (
+                "GPT-OSS SwiGLU requires gemm1_alpha and gemm1_limit, "
+                "and cannot use swiglu_limit"
+            )
+        elif swiglu_limit is not None:
+            assert (
+                swiglu_limit == 10 and gemm1_limit is None and gemm1_alpha is None
+            ), "FP8 DeepSeek-V4 SwiGLU currently requires swiglu_limit=10"
+        else:
+            assert (
+                gemm1_limit is None
+            ), "gemm1_limit requires gemm1_alpha for GPT-OSS SwiGLU"
+    else:
+        assert activation in (
+            "silu",
+            "gelu",
+            "relu2",
+        ), f"Only silu, gelu and relu2 are supported but got {activation}"
 
     # Unified 4-bit W4A16 MoE (mxfp4 or int4). Weights are packed int8/uint8
     # [E, N, K/2]; scales are [E, N, K/group_size] N-outer. For mxfp4 the
@@ -525,6 +625,9 @@ def fused_experts(
     assert not (
         use_mxfp4_w4a16 and use_int4_w4a16
     ), "use_mxfp4_w4a16 and use_int4_w4a16 are mutually exclusive"
+    assert not (
+        use_4bit_w4a16 and use_fp8_weight
+    ), "4-bit W4A16 and FP8 paths are mutually exclusive"
     if use_4bit_w4a16:
         assert (
             w1.dtype == torch.int8 or w1.dtype == torch.uint8
@@ -557,8 +660,9 @@ def fused_experts(
                     w2_zp.dtype == w2_scale.dtype and w2_zp.shape == w2_scale.shape
                 ), "w2_zp must have the same dtype and shape as w2_scale"
     else:
-        assert w1_scale is None, "w1_scale is only supported for 4-bit W4A16 MoE"
-        assert w2_scale is None, "w2_scale is only supported for 4-bit W4A16 MoE"
+        if not use_fp8_weight:
+            assert w1_scale is None, "w1_scale is only supported for 4-bit W4A16 MoE"
+            assert w2_scale is None, "w2_scale is only supported for 4-bit W4A16 MoE"
         assert (
             w1_zp is None and w2_zp is None
         ), "w1_zp/w2_zp are only supported for 4-bit W4A16 MoE"
@@ -571,6 +675,25 @@ def fused_experts(
         assert (
             use_int4_w4a16
         ), "w1_g_idx_perm/w2_g_idx_perm only apply to use_int4_w4a16"
+    elif use_fp8_weight:
+        assert (
+            w1.dtype == torch.float8_e4m3fn
+        ), "FP8 weight-only MoE requires w1 to be float8_e4m3fn"
+        assert (
+            w2.dtype == torch.float8_e4m3fn
+        ), "FP8 weight-only MoE requires w2 to be float8_e4m3fn"
+        assert (
+            w1_scale is not None and w2_scale is not None
+        ), "w1_scale/w2_scale must be provided for FP8 weight-only MoE"
+        assert (
+            w1.is_contiguous() and w2.is_contiguous()
+        ), "FP8 weight-only MoE requires contiguous expert weights"
+        _validate_fp8_weight_scale(w1_scale, w1, "w1_scale", allow_scalar=True)
+        _validate_fp8_weight_scale(w2_scale, w2, "w2_scale", allow_scalar=True)
+        assert (
+            w1_scale.ndim == w2_scale.ndim
+        ), "w1_scale and w2_scale must use the same scalar or block layout"
+        assert hidden_states.dtype == torch.bfloat16, "hidden_states must be bfloat16"
     if b1 is not None:
         assert (
             b1.dtype == torch.bfloat16 or b1.dtype == torch.float32
@@ -635,42 +758,57 @@ def fused_experts(
     else:
         out_hidden_states = torch.empty_like(hidden_states)
 
-    topk_ids = topk_ids.int() if topk_ids.dtype == torch.long else topk_ids
     expert_offsets = _get_moe_ws(
         "expert_offsets", (E,), torch.int32, hidden_states.device
     )
-    problem_sizes1 = _get_moe_ws(
-        "problem_sizes1", (E, 3), torch.int32, hidden_states.device
+    use_small_prepare = (
+        _should_use_small_moe_prepare(M, TopK, hidden_dims, E) and use_fp8_weight
     )
-    problem_sizes2 = _get_moe_ws(
-        "problem_sizes2", (E, 3), torch.int32, hidden_states.device
-    )
-    a_map = _get_moe_ws("a_map", (topk_ids.numel(),), torch.int32, hidden_states.device)
     c_map = _get_moe_ws("c_map", (topk_ids.numel(),), torch.int32, hidden_states.device)
-    torch.ops.sgl_kernel.prepare_moe_input.default(
-        topk_ids,
-        expert_offsets,
-        None,
-        problem_sizes1,
-        problem_sizes2,
-        a_map,
-        c_map,
-        E,
-        hidden_dims,
-        TopK,
-    )
     input_A_shuffle = _get_moe_ws(
         "input_A_shuffle",
         (num_tokens * TopK, K),
         hidden_states.dtype,
         hidden_states.device,
     )
-    # Use scatter_tokens_to_experts (IPEX MoEScatter style):
-    # 1 WG per source token, reads sequentially, scatters to TopK destinations,
-    # with coalesced reads and data reuse.
-    torch.ops.sgl_kernel.scatter_tokens_to_experts.default(
-        hidden_states, c_map, input_A_shuffle
-    )
+    if use_small_prepare:
+        # TODO: Support strided topk_ids in the prepare kernels and remove this copy.
+        topk_ids_small = topk_ids if topk_ids.is_contiguous() else topk_ids.contiguous()
+        torch.ops.sgl_kernel.prepare_moe_input_small.default(
+            hidden_states, topk_ids_small, expert_offsets, c_map, input_A_shuffle
+        )
+    else:
+        if topk_ids.dtype == torch.long:
+            topk_ids_int = _get_moe_ws(
+                "topk_ids_int", topk_ids.shape, torch.int32, hidden_states.device
+            )
+            topk_ids_int.copy_(topk_ids)
+        else:
+            topk_ids_int = topk_ids
+        problem_sizes1 = _get_moe_ws(
+            "problem_sizes1", (E, 3), torch.int32, hidden_states.device
+        )
+        problem_sizes2 = _get_moe_ws(
+            "problem_sizes2", (E, 3), torch.int32, hidden_states.device
+        )
+        a_map = _get_moe_ws(
+            "a_map", (topk_ids.numel(),), torch.int32, hidden_states.device
+        )
+        torch.ops.sgl_kernel.prepare_moe_input.default(
+            topk_ids_int,
+            expert_offsets,
+            None,
+            problem_sizes1,
+            problem_sizes2,
+            a_map,
+            c_map,
+            E,
+            hidden_dims,
+            TopK,
+        )
+        torch.ops.sgl_kernel.scatter_tokens_to_experts.default(
+            hidden_states, c_map, input_A_shuffle
+        )
     if w1_g_idx_perm is not None:
         # GPTQ desc_act/g_idx: reorder each expert's activation slice to
         # match the K-dim sort applied to w1 at weight-load time.
@@ -685,6 +823,94 @@ def fused_experts(
         hidden_states.device,
     )
 
+    if use_fp8_weight:
+        if activation == "gelu":
+            activation_type = 1
+        elif activation == "relu2":
+            activation_type = 3
+        elif activation != "silu":
+            raise ValueError(
+                f"FP8 W8A16 Xe2 path does not support activation={activation!r}; "
+                "supported activations are 'silu', 'gelu', and 'relu2'"
+            )
+        elif gemm1_alpha is not None:
+            if gemm1_limit is None:
+                raise AssertionError(
+                    "gemm1_limit must be provided when gemm1_alpha is set for swiglu for GPT-OSS"
+                )
+            activation_type = 2
+        elif swiglu_limit is not None:
+            activation_type = 4
+            gemm1_limit = float(swiglu_limit)
+        else:
+            activation_type = 0
+
+        gemm1_output_width = N if activation_type == 3 else 2 * N
+        intermediate_cache1 = _get_moe_ws(
+            "intermediate_cache1",
+            (M * TopK, gemm1_output_width),
+            hidden_states.dtype,
+            hidden_states.device,
+        )
+        torch.ops.sgl_kernel.moe_grouped_mm_nt_xe20_fp8_w8a16(
+            intermediate_cache1,
+            input_A_shuffle,
+            w1,
+            w1_scale,
+            b1,
+            expert_offsets,
+            E,
+        )
+
+        if activation_type == 2:
+            intermediate_cache2 = torch.ops.sgl_kernel.swiglu_gpt_oss_sigmoid_alpha(
+                intermediate_cache1, gemm1_alpha, gemm1_limit
+            )
+        else:
+            intermediate_cache2 = _get_moe_ws(
+                "intermediate_cache2",
+                (M * TopK, N),
+                hidden_states.dtype,
+                hidden_states.device,
+            )
+            if activation_type == 0:
+                torch.ops.sgl_kernel.silu_and_mul(
+                    intermediate_cache2, intermediate_cache1
+                )
+            elif activation_type == 4:
+                torch.ops.sgl_kernel.silu_and_mul_clamp(
+                    intermediate_cache2, intermediate_cache1, swiglu_limit
+                )
+            elif activation_type == 1:
+                torch.ops.sgl_kernel.gelu_tanh_and_mul(
+                    intermediate_cache2, intermediate_cache1
+                )
+            elif activation_type == 3:
+                torch.clamp_min(intermediate_cache1, 0, out=intermediate_cache2)
+                torch.square(intermediate_cache2, out=intermediate_cache2)
+            else:
+                raise AssertionError(
+                    f"unsupported FP8 activation type: {activation_type}"
+                )
+
+        torch.ops.sgl_kernel.moe_grouped_mm_nt_xe20_fp8_w8a16(
+            intermediate_cache3,
+            intermediate_cache2,
+            w2,
+            w2_scale,
+            b2,
+            expert_offsets,
+            E,
+        )
+
+        rsf = 1.0
+        if routed_scaling_factor is not None:
+            rsf = routed_scaling_factor
+        torch.ops.sgl_kernel.apply_shuffle_mul_sum.default(
+            intermediate_cache3, out_hidden_states, c_map, rsf, topk_weights
+        )
+        return out_hidden_states
+
     # 0=silu, 1=gelu, 2=swiglu (silu with alpha/limit clamping for gpt-oss),
     # 3=relu2, 4=swiglu_deepseek_v4 (clamp gate/up then plain silu * up).
     if activation == "silu":
@@ -740,12 +966,6 @@ def fused_experts(
             hidden_states.dtype,
             hidden_states.device,
         )
-        intermediate_cache2 = _get_moe_ws(
-            "intermediate_cache2",
-            (M * TopK, N),
-            hidden_states.dtype,
-            hidden_states.device,
-        )
         # GEMM1: B = w1 (gate+up).
         if use_4bit_w4a16:
             torch.ops.sgl_kernel.moe_grouped_mm_nt_xe20_w4a16(
@@ -773,22 +993,32 @@ def fused_experts(
                 gemm1_alpha=float(gemm1_alpha) if gemm1_alpha is not None else 1.702,
                 gemm1_limit=float(gemm1_limit) if gemm1_limit is not None else 7.0,
             )
-        if activation_type == 0:
-            torch.ops.sgl_kernel.silu_and_mul(intermediate_cache2, intermediate_cache1)
-        elif activation_type == 4:
-            torch.ops.sgl_kernel.silu_and_mul_clamp(
-                intermediate_cache2, intermediate_cache1, swiglu_limit
-            )
-        elif activation_type == 1:
-            torch.ops.sgl_kernel.gelu_tanh_and_mul(
-                intermediate_cache2, intermediate_cache1
-            )
-        elif activation_type == 2:
+        if activation_type == 2:
             intermediate_cache2 = torch.ops.sgl_kernel.swiglu_gpt_oss_sigmoid_alpha(
                 intermediate_cache1, gemm1_alpha, gemm1_limit
             )
-        elif activation_type == 3:
-            intermediate_cache2 = torch.square(torch.relu(intermediate_cache1))
+        else:
+            intermediate_cache2 = _get_moe_ws(
+                "intermediate_cache2",
+                (M * TopK, N),
+                hidden_states.dtype,
+                hidden_states.device,
+            )
+            if activation_type == 0:
+                torch.ops.sgl_kernel.silu_and_mul(
+                    intermediate_cache2, intermediate_cache1
+                )
+            elif activation_type == 4:
+                torch.ops.sgl_kernel.silu_and_mul_clamp(
+                    intermediate_cache2, intermediate_cache1, swiglu_limit
+                )
+            elif activation_type == 1:
+                torch.ops.sgl_kernel.gelu_tanh_and_mul(
+                    intermediate_cache2, intermediate_cache1
+                )
+            elif activation_type == 3:
+                torch.clamp_min(intermediate_cache1, 0, out=intermediate_cache2)
+                torch.square(intermediate_cache2, out=intermediate_cache2)
         if w2_g_idx_perm is not None:
             # GPTQ desc_act/g_idx: reorder each expert's activation slice to
             # match the K-dim sort applied to w2 at weight-load time.
diff --git a/src/BuildOnLinux.cmake b/src/BuildOnLinux.cmake
index ff3e2f76..e6b920dd 100644
--- a/src/BuildOnLinux.cmake
+++ b/src/BuildOnLinux.cmake
@@ -127,6 +127,7 @@ set(XE20_OFFLINE_COMPILER_FLAGS "${XE20_OFFLINE_COMPILER_AOT_OPTIONS}${SYCL_OFFL
 set(SGL_XE20_BUNDLE_PREFIXES
   "GroupGemmXe20_inst_"
   "GroupGemmW4A16Xe20_inst_"
+  "GroupGemmFp8Xe20_inst_"
   "xe_fmha_fwd_decode_page_"
   "xe_fmha_fwd_decode_nopage_"
   "xe_fmha_fwd_split_decode_page_"
@@ -228,6 +229,19 @@ endif()
 if(USE_MOE AND USE_SYCL_JIT AND TARGET sgl-ops-sycl-GroupGemmW4A16Xe20)
   target_link_libraries(sgl-ops-sycl-GroupGemmW4A16Xe20 PRIVATE sgl_jit)
 endif()
+# The FP8 W8A16 grouped GEMM dispatcher follows the same JIT path.
+if(USE_MOE AND USE_SYCL_JIT AND TARGET sgl-ops-sycl-GroupGemmFp8W8A16Xe20)
+  target_link_libraries(sgl-ops-sycl-GroupGemmFp8W8A16Xe20 PRIVATE sgl_jit)
+endif()
+# In AOT builds, keep the bundled FP8 instances as an explicit dependency of
+# their dispatcher so --as-needed does not drop them from the runtime graph.
+if(USE_MOE AND NOT USE_SYCL_JIT
+   AND TARGET sgl-ops-sycl-GroupGemmFp8W8A16Xe20
+   AND TARGET sgl-ops-sycl-GroupGemmFp8Xe20_inst)
+  target_link_libraries(
+    sgl-ops-sycl-GroupGemmFp8W8A16Xe20
+    PUBLIC sgl-ops-sycl-GroupGemmFp8Xe20_inst)
+endif()
 # The GDN chunk delta-rule dispatch (chunk_gated_delta_rule.cpp) calls the JIT engine.
 if(USE_FMHA AND USE_SYCL_JIT AND TARGET sgl-ops-sycl-chunk_gated_delta_rule)
   target_link_libraries(sgl-ops-sycl-chunk_gated_delta_rule PRIVATE sgl_jit)
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index 7059fd15..3837832e 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -114,7 +114,7 @@ if(NOT USE_MLA)
     list(FILTER device_cpp EXCLUDE REGEX "/mla_(decode|prefill|sparse_decode)\\.cpp$")
 endif()
 if(NOT USE_MOE)
-    list(FILTER device_cpp EXCLUDE REGEX "/(GroupGemm(Xe20|W4A16Xe20)|MoE(Align|PrepareInputs|Sum|_fused_gate|_sum_reduce))\\.cpp$")
+    list(FILTER device_cpp EXCLUDE REGEX "/(GroupGemm(Xe20|W4A16Xe20|Fp8W8A16Xe20)|MoE(Align|PrepareInputs|Sum|_fused_gate|_sum_reduce))\\.cpp$")
 endif()
 
 # BMG only
@@ -157,6 +157,7 @@ if(USE_MOE)
     if(NOT USE_SYCL_JIT)
         include(${CMAKE_CURRENT_SOURCE_DIR}/GroupGemmXe20.cmake)
         include(${CMAKE_CURRENT_SOURCE_DIR}/GroupGemmW4A16Xe20.cmake)
+        include(${CMAKE_CURRENT_SOURCE_DIR}/GroupGemmFp8W8A16Xe20.cmake)
     endif()
 endif()
 
diff --git a/src/GroupGemmFp8W8A16Xe20.cmake b/src/GroupGemmFp8W8A16Xe20.cmake
new file mode 100644
index 00000000..6fe7906d
--- /dev/null
+++ b/src/GroupGemmFp8W8A16Xe20.cmake
@@ -0,0 +1,30 @@
+set(GROUP_GEMM_FP8_XE20_TEMPLATE "${CMAKE_CURRENT_SOURCE_DIR}/sycl/GroupGemmFp8W8A16Xe20LauncherInstance.cpp.in")
+set(GROUP_GEMM_FP8_XE20_GEN_DIR "${CMAKE_CURRENT_BINARY_DIR}/generated/group_gemm_fp8_xe20")
+set(GROUP_GEMM_FP8_XE20_INST_SRCS)
+file(MAKE_DIRECTORY ${GROUP_GEMM_FP8_XE20_GEN_DIR})
+
+function(add_group_gemm_fp8_xe20_inst TILE_M TILE_N TILE_K SG_SHAPE SG_STRIDE SCALE_COUNTS)
+    set(TILE "Shape<${TILE_M}, ${TILE_N}, ${TILE_K}>")
+    set(SGLAYOUT "Layout<Shape<${SG_SHAPE}>, Stride<${SG_STRIDE}>>")
+    foreach(with_bias false true)
+        set(WITH_BIAS ${with_bias})
+        foreach(scale_count IN LISTS SCALE_COUNTS)
+            set(SCALE_COUNT ${scale_count})
+            set(SCALE_GEN_SRC
+                "${GROUP_GEMM_FP8_XE20_GEN_DIR}/GroupGemmFp8Xe20_inst_${TILE_M}_${TILE_N}_${TILE_K}_b${WITH_BIAS}_s${SCALE_COUNT}.cpp")
+            configure_file(${GROUP_GEMM_FP8_XE20_TEMPLATE} ${SCALE_GEN_SRC} @ONLY)
+            list(APPEND GROUP_GEMM_FP8_XE20_INST_SRCS ${SCALE_GEN_SRC})
+        endforeach()
+    endforeach()
+    set(GROUP_GEMM_FP8_XE20_INST_SRCS ${GROUP_GEMM_FP8_XE20_INST_SRCS} PARENT_SCOPE)
+endfunction()
+
+# W8A16 scalar/block menu: 16x64x32 and 32x64x32 for small/medium M,
+# 64x64x32 for medium-M long-K scalar GEMMs, plus 128x128x16 for
+# large-M scalar GEMMs. Activation is external.
+add_group_gemm_fp8_xe20_inst("_16" "_64" "_32" "_1, _4, _1" "_4, _1, _0" "1;2;3")
+add_group_gemm_fp8_xe20_inst("_32" "_64" "_32" "_1, _4, _1" "_4, _1, _0" "1;2;3")
+add_group_gemm_fp8_xe20_inst("_64" "_64" "_32" "_2, _4, _1" "_4, _1, _0" "2")
+add_group_gemm_fp8_xe20_inst("_128" "_128" "_16" "_4, _2, _1" "_2, _1, _0" "1;2;3")
+
+list(APPEND ATen_XPU_SYCL_XE20 ${GROUP_GEMM_FP8_XE20_INST_SRCS})
diff --git a/src/jit/moe_jit.cpp b/src/jit/moe_jit.cpp
index 21ca54a4..9d2a6151 100644
--- a/src/jit/moe_jit.cpp
+++ b/src/jit/moe_jit.cpp
@@ -253,5 +253,120 @@ bool w4a16_grouped_gemm_launch(
   return true;
 }
 
+// ---------------------------------------------------------------------------
+// FP8 W8A16 grouped GEMM.
+// ---------------------------------------------------------------------------
+
+namespace {
+
+using Fp8W8A16Fn = void (*)(
+    void*, const void*, const void*, const void*, const void*, void*, int, int, const int*, int, int*, int, bool, bool);
+
+struct Fp8TileCfg {
+  const char* tile;
+  const char* sglayout;
+};
+
+constexpr Fp8TileCfg kFp8Tiles[] = {
+    {"Shape<_16, _64, _32>", "Layout<Shape<_1, _4, _1>, Stride<_4, _1, _0>>"},
+    {"Shape<_32, _64, _32>", "Layout<Shape<_1, _4, _1>, Stride<_4, _1, _0>>"},
+    {"Shape<_64, _64, _32>", "Layout<Shape<_2, _4, _1>, Stride<_4, _1, _0>>"},
+    {"Shape<_128, _128, _16>", "Layout<Shape<_4, _2, _1>, Stride<_2, _1, _0>>"},
+};
+
+int fp8_tile_id(int avg_m, int gemm_k, int gemm_n, int scale_count) {
+  if (scale_count == 3) {
+    if (avg_m <= 4) return 0;
+    if (avg_m >= 1024 || (avg_m > 128 && gemm_k >= 512 && gemm_n >= 512)) return 3;
+    return 1;
+  }
+  if (avg_m <= 8) return 0;
+  if (scale_count == 2 && avg_m > 32 && avg_m <= 128 && gemm_k >= 2048) return 2;
+  if (avg_m <= 32 || (scale_count == 2 && gemm_k >= 4096 && avg_m <= 512)) return 1;
+  return 3;
+}
+
+uint64_t pack_fp8_w8a16_key(int tile_id, int scale_count, bool with_bias, int arch) {
+  uint64_t key = static_cast<uint64_t>(arch) & 0xFF;
+  key = (key << 8) | (static_cast<uint64_t>(tile_id) & 0xFF);
+  key = (key << 8) | (static_cast<uint64_t>(scale_count) & 0xFF);
+  key = (key << 1) | (with_bias ? 1u : 0u);
+  return key;
+}
+
+jit::JitFnCache<Fp8W8A16Fn> g_fp8_w8a16_fns("FP8 W8A16 grouped GEMM");
+
+Fp8W8A16Fn resolve_fp8_w8a16(int tile_id, int scale_count, bool with_bias, int arch, std::string* err) {
+  const uint64_t key = pack_fp8_w8a16_key(tile_id, scale_count, with_bias, arch);
+  auto build = [&](std::string* build_err) -> void* {
+    const jit::JitConfig& cfg = jit::default_config();
+    if (!cfg.valid) {
+      *build_err = "unavailable: " + cfg.error;
+      return nullptr;
+    }
+    if (cfg.src_root.empty()) {
+      *build_err = "source template root not resolved";
+      return nullptr;
+    }
+
+    jit::CompileSpec spec;
+    spec.template_path = cfg.src_root + "/sycl/GroupGemmFp8W8A16Xe20LauncherInstance.cpp.in";
+    spec.subs["TILE"] = kFp8Tiles[tile_id].tile;
+    spec.subs["SGLAYOUT"] = kFp8Tiles[tile_id].sglayout;
+    spec.subs["WITH_BIAS"] = with_bias ? "true" : "false";
+    spec.subs["SCALE_COUNT"] = std::to_string(scale_count);
+    const jit::ArchSpec arch_spec = jit::arch_spec(static_cast<jit::Arch>(arch), "-DSGL_FP8_W8A16_JIT_ENTRY");
+    spec.extra_flags = arch_spec.extra_flags;
+    spec.target = arch_spec.target;
+    spec.entry_symbol = "sgl_moe_fp8_w8a16_entry";
+    spec.name = std::string("group_gemm_fp8_w8a16_t") + std::to_string(tile_id) + "_s" + std::to_string(scale_count) +
+                "_b" + (with_bias ? "1" : "0") + "_" + arch_spec.suffix;
+    return jit::get_or_compile(spec, cfg, build_err);
+  };
+  return g_fp8_w8a16_fns.get(key, build, err);
+}
+
+}  // namespace
+
+bool fp8_w8a16_grouped_gemm_launch(
+    int avg_m,
+    int scale_count,
+    bool with_bias,
+    void* queue,
+    const void* activations,
+    const void* weights,
+    const void* weight_scales,
+    const void* bias,
+    void* outputs,
+    int gemm_n,
+    int gemm_k,
+    const int* rows_per_expert,
+    int num_experts,
+    int* workspace,
+    int ld_b,
+    bool weight_scale_blocked,
+    bool static_scheduler,
+    int arch,
+    std::string* err) {
+  const int tile_id = fp8_tile_id(avg_m, gemm_k, gemm_n, scale_count);
+  Fp8W8A16Fn fn = resolve_fp8_w8a16(tile_id, scale_count, with_bias, arch, err);
+  if (!fn) return false;
+  fn(queue,
+     activations,
+     weights,
+     weight_scales,
+     bias,
+     outputs,
+     gemm_n,
+     gemm_k,
+     rows_per_expert,
+     num_experts,
+     workspace,
+     ld_b,
+     weight_scale_blocked,
+     static_scheduler);
+  return true;
+}
+
 }  // namespace moe_jit
 }  // namespace sgl
diff --git a/src/jit/moe_jit.h b/src/jit/moe_jit.h
index b075ce98..cde6509e 100644
--- a/src/jit/moe_jit.h
+++ b/src/jit/moe_jit.h
@@ -61,5 +61,29 @@ bool w4a16_grouped_gemm_launch(
     int arch = 0,  // sgl::jit::Arch code (0=BMG/Xe20, 1=XE3P/Xe35)
     std::string* err = nullptr);
 
+// Launch the FP8-weight/BF16-activation grouped GEMM. The tile is selected
+// from the routed M, GEMM shape, and scalar/block scale layout exactly as in
+// GroupGemmFp8W8A16Xe20.cpp.
+bool fp8_w8a16_grouped_gemm_launch(
+    int avg_m,
+    int scale_count,
+    bool with_bias,
+    void* queue,
+    const void* activations,
+    const void* weights,
+    const void* weight_scales,
+    const void* bias,
+    void* outputs,
+    int gemm_n,
+    int gemm_k,
+    const int* rows_per_expert,
+    int num_experts,
+    int* workspace,
+    int ld_b,
+    bool weight_scale_blocked,
+    bool static_scheduler,
+    int arch = 0,
+    std::string* err = nullptr);
+
 }  // namespace moe_jit
 }  // namespace sgl
diff --git a/src/sycl/GroupGemmFp8W8A16Xe20.cpp b/src/sycl/GroupGemmFp8W8A16Xe20.cpp
new file mode 100644
index 00000000..2be32db0
--- /dev/null
+++ b/src/sycl/GroupGemmFp8W8A16Xe20.cpp
@@ -0,0 +1,257 @@
+#define SYCL_INTEL_TARGET 20
+
+#include <ATen/ATen.h>
+#include <c10/xpu/XPUStream.h>
+#include <torch/all.h>
+
+#include <cute/tensor.hpp>
+#include <map>
+
+#include "Utils.h"
+#include "cutlass/gemm/device/gemm_universal_adapter.h"
+#include "cutlass/gemm/group_array_problem_shape.hpp"
+#include "kernels/moe/xe20/fp8/moe_kernel.hpp"
+#include "sgl_kernel_export.h"
+#ifdef USE_MOE_JIT
+#include "jit/moe_jit.h"
+#endif
+
+using namespace cute;
+
+template <typename Tile, typename SGLayout, bool WithBias, int WeightScaleCount>
+__attribute__((visibility("default"))) void Xe20MoEGEMMFp8W8A16Launcher(
+    sycl::queue q,
+    const void* activations,
+    const void* weights,
+    const void* weight_scales,
+    const void* bias,
+    void* outputs,
+    const int gemm_n,
+    const int gemm_k,
+    const int* num_rows_per_expert_device,
+    const int num_experts,
+    int* workspace,
+    int ld_b,
+    bool weight_scale_blocked,
+    bool static_scheduler);
+
+using Tile_16_64_32 = Shape<_16, _64, _32>;
+using Tile_32_64_32 = Shape<_32, _64, _32>;
+using Tile_64_64_32 = Shape<_64, _64, _32>;
+using Tile_128_128_16 = Shape<_128, _128, _16>;
+
+using SG_1_4_1 = Layout<Shape<_1, _4, _1>, Stride<_4, _1, _0>>;
+using SG_2_4_1 = Layout<Shape<_2, _4, _1>, Stride<_4, _1, _0>>;
+using SG_4_2_1 = Layout<Shape<_4, _2, _1>, Stride<_2, _1, _0>>;
+
+#define DECLARE_XE20_MOE_FP8_W8A16_EXTERN(Tile, SGLayout, WithBias, ScaleCount)           \
+  extern template void Xe20MoEGEMMFp8W8A16Launcher<Tile, SGLayout, WithBias, ScaleCount>( \
+      sycl::queue,                                                                        \
+      const void*,                                                                        \
+      const void*,                                                                        \
+      const void*,                                                                        \
+      const void*,                                                                        \
+      void*,                                                                              \
+      const int,                                                                          \
+      const int,                                                                          \
+      const int*,                                                                         \
+      const int,                                                                          \
+      int*,                                                                               \
+      int,                                                                                \
+      bool,                                                                               \
+      bool);
+
+#define DECLARE_XE20_MOE_FP8_W8A16_BIAS_VARIANTS(Tile, SGLayout, ScaleCount) \
+  DECLARE_XE20_MOE_FP8_W8A16_EXTERN(Tile, SGLayout, false, ScaleCount)       \
+  DECLARE_XE20_MOE_FP8_W8A16_EXTERN(Tile, SGLayout, true, ScaleCount)
+
+#define DECLARE_XE20_MOE_FP8_W8A16_ALL_SCALE_VARIANTS(Tile, SGLayout) \
+  DECLARE_XE20_MOE_FP8_W8A16_BIAS_VARIANTS(Tile, SGLayout, 1)         \
+  DECLARE_XE20_MOE_FP8_W8A16_BIAS_VARIANTS(Tile, SGLayout, 2)         \
+  DECLARE_XE20_MOE_FP8_W8A16_BIAS_VARIANTS(Tile, SGLayout, 3)
+
+DECLARE_XE20_MOE_FP8_W8A16_ALL_SCALE_VARIANTS(Tile_16_64_32, SG_1_4_1)
+DECLARE_XE20_MOE_FP8_W8A16_ALL_SCALE_VARIANTS(Tile_32_64_32, SG_1_4_1)
+DECLARE_XE20_MOE_FP8_W8A16_BIAS_VARIANTS(Tile_64_64_32, SG_2_4_1, 2)
+DECLARE_XE20_MOE_FP8_W8A16_ALL_SCALE_VARIANTS(Tile_128_128_16, SG_4_2_1)
+
+#undef DECLARE_XE20_MOE_FP8_W8A16_ALL_SCALE_VARIANTS
+#undef DECLARE_XE20_MOE_FP8_W8A16_BIAS_VARIANTS
+#undef DECLARE_XE20_MOE_FP8_W8A16_EXTERN
+
+#define LAUNCH_MOE_FP8_W8A16(ScaleCount, WithBiasVal, ...)           \
+  Xe20MoEGEMMFp8W8A16Launcher<__VA_ARGS__, WithBiasVal, ScaleCount>( \
+      queue,                                                         \
+      activations.data_ptr(),                                        \
+      weights.data_ptr(),                                            \
+      weight_scales.data_ptr(),                                      \
+      bias_ptr,                                                      \
+      output.data_ptr(),                                             \
+      gemm_n,                                                        \
+      gemm_k,                                                        \
+      total_rows_for_experts.data_ptr<int>(),                        \
+      n_experts,                                                     \
+      atomic_buffer.data_ptr<int>(),                                 \
+      ld_b,                                                          \
+      weight_scale_blocked,                                          \
+      static_scheduler)
+
+#define DISPATCH_MOE_FP8_W8A16_BLOCK_TILES(WithBiasVal)                            \
+  do {                                                                             \
+    if (avg_m <= 4) {                                                              \
+      LAUNCH_MOE_FP8_W8A16(3, WithBiasVal, Tile_16_64_32, SG_1_4_1);               \
+    } else if (avg_m >= 1024 || (avg_m > 128 && gemm_k >= 512 && gemm_n >= 512)) { \
+      LAUNCH_MOE_FP8_W8A16(3, WithBiasVal, Tile_128_128_16, SG_4_2_1);             \
+    } else {                                                                       \
+      LAUNCH_MOE_FP8_W8A16(3, WithBiasVal, Tile_32_64_32, SG_1_4_1);               \
+    }                                                                              \
+  } while (0)
+
+#define DISPATCH_MOE_FP8_W8A16_SCALAR_TILES(ScaleCount, WithBiasVal)                 \
+  do {                                                                               \
+    if (avg_m <= 8) {                                                                \
+      LAUNCH_MOE_FP8_W8A16(ScaleCount, WithBiasVal, Tile_16_64_32, SG_1_4_1);        \
+    } else if (ScaleCount == 2 && avg_m > 32 && avg_m <= 128 && gemm_k >= 2048) {    \
+      LAUNCH_MOE_FP8_W8A16(ScaleCount, WithBiasVal, Tile_64_64_32, SG_2_4_1);        \
+    } else if (avg_m <= 32 || (ScaleCount == 2 && gemm_k >= 4096 && avg_m <= 512)) { \
+      LAUNCH_MOE_FP8_W8A16(ScaleCount, WithBiasVal, Tile_32_64_32, SG_1_4_1);        \
+    } else {                                                                         \
+      LAUNCH_MOE_FP8_W8A16(ScaleCount, WithBiasVal, Tile_128_128_16, SG_4_2_1);      \
+    }                                                                                \
+  } while (0)
+
+SGL_KERNEL_EXPORT void moe_grouped_mm_nt_xe20_fp8_w8a16(
+    torch::Tensor& output,
+    const torch::Tensor& activations,
+    const torch::Tensor& weights,
+    const torch::Tensor& weight_scales,
+    const std::optional<at::Tensor>& bias,
+    const torch::Tensor& total_rows_for_experts,
+    const int64_t n_experts) {
+  CHECK_INPUT(output);
+  CHECK_INPUT(activations);
+  CHECK_INPUT(weights);
+  CHECK_INPUT(weight_scales);
+  CHECK_INPUT(total_rows_for_experts);
+  TORCH_CHECK(output.device() == activations.device(), "output must be on the same device as activations");
+  TORCH_CHECK(weights.device() == activations.device(), "weights must be on the same device as activations");
+  TORCH_CHECK(
+      weight_scales.device() == activations.device(), "weight_scales must be on the same device as activations");
+  TORCH_CHECK(
+      total_rows_for_experts.device() == activations.device(),
+      "total_rows_for_experts must be on the same device as activations");
+  if (bias.has_value()) {
+    const auto& bias_tensor = *bias;
+    CHECK_INPUT(bias_tensor);
+    TORCH_CHECK(bias_tensor.device() == activations.device(), "bias must be on the same device as activations");
+    TORCH_CHECK(bias_tensor.scalar_type() == at::kFloat, "bias must be float32");
+    TORCH_CHECK(bias_tensor.dim() == 2, "bias must be 2D [E, N]");
+  }
+  TORCH_CHECK(activations.scalar_type() == at::ScalarType::BFloat16, "W8A16 activations must be bfloat16");
+  TORCH_CHECK(weights.scalar_type() == at::ScalarType::Float8_e4m3fn, "W8A16 weights must be float8_e4m3fn");
+  TORCH_CHECK(weight_scales.scalar_type() == at::kFloat, "W8A16 weight scales must be float32");
+  TORCH_CHECK(output.scalar_type() == at::ScalarType::BFloat16, "W8A16 output must be bfloat16");
+  TORCH_CHECK(
+      weight_scales.dim() == 2 || weight_scales.dim() == 3,
+      "W8A16 weight scales must be [E, 1]/[E, 2] or [E, N/128, K/128]");
+  if (weight_scales.dim() == 2) {
+    TORCH_CHECK(weight_scales.size(1) == 1 || weight_scales.size(1) == 2, "W8A16 scale count must be 1 or 2");
+  }
+  TORCH_CHECK(n_experts > 0 && n_experts % 8 == 0, "n_experts must be a positive multiple of 8");
+  TORCH_CHECK(activations.dim() == 2, "W8A16 activations must be 2D [M_total, K]");
+  TORCH_CHECK(weights.dim() == 3, "W8A16 weights must be 3D [E, N, K]");
+  TORCH_CHECK(output.dim() == 2, "W8A16 output must be 2D [M_total, N]");
+  TORCH_CHECK(weights.size(0) == n_experts, "weights expert dimension mismatch");
+  TORCH_CHECK(weight_scales.size(0) == n_experts, "weight scales expert dimension mismatch");
+  TORCH_CHECK(
+      total_rows_for_experts.dim() == 1 && total_rows_for_experts.size(0) == n_experts, "rows_for_experts must be [E]");
+  TORCH_CHECK(total_rows_for_experts.scalar_type() == at::ScalarType::Int, "rows_for_experts must be int32");
+  TORCH_CHECK(weights.size(2) == activations.size(1), "W8A16 K dimension mismatch");
+  TORCH_CHECK(
+      activations.is_contiguous() && weights.is_contiguous() && weight_scales.is_contiguous(),
+      "W8A16 tensors must be contiguous");
+  TORCH_CHECK(weights.size(1) % 64 == 0 && weights.size(2) % 32 == 0, "W8A16 N must be divisible by 64 and K by 32");
+
+  int total_m = static_cast<int>(activations.size(0));
+  int gemm_k = static_cast<int>(activations.size(1));
+  int gemm_n = static_cast<int>(weights.size(1));
+  int avg_m = total_m / static_cast<int>(n_experts);
+  int ld_b = static_cast<int>(weights.stride(1));
+  int scale_count = weight_scales.dim() == 2 ? static_cast<int>(weight_scales.size(1)) : 3;
+  bool weight_scale_blocked = weight_scales.dim() == 3;
+  bool static_scheduler = total_m <= n_experts || (!weight_scale_blocked && gemm_k <= 128);
+  if (weight_scale_blocked) {
+    TORCH_CHECK(weight_scales.size(1) == (gemm_n + 127) / 128, "W8A16 block scale N dimension must be ceil(N/128)");
+    TORCH_CHECK(
+        gemm_k % 128 == 0 && weight_scales.size(2) == gemm_k / 128, "W8A16 block scale K dimension must be K/128");
+  }
+  TORCH_CHECK(output.size(0) == total_m, "output rows must equal M_total");
+  TORCH_CHECK(output.size(1) == gemm_n, "output must have the same columns as weights");
+  if (bias.has_value()) {
+    TORCH_CHECK(bias->size(0) == n_experts && bias->size(1) == gemm_n, "bias shape must be [E, N]");
+  }
+  auto stream = at::xpu::getCurrentXPUStream();
+  auto queue = stream.queue();
+  using StreamKey = std::pair<c10::DeviceIndex, c10::StreamId>;
+  thread_local std::map<StreamKey, at::Tensor> atomic_buffers;
+  auto [buffer_it, inserted] = atomic_buffers.try_emplace(StreamKey{stream.device_index(), stream.id()}, at::Tensor{});
+  at::Tensor& atomic_buffer = buffer_it->second;
+  if (inserted || !atomic_buffer.defined()) {
+    atomic_buffer = at::empty({1}, activations.options().dtype(at::kInt));
+  }
+  if (!static_scheduler) {
+    queue.memset(atomic_buffer.data_ptr(), 0, sizeof(int32_t));
+  }
+  bool with_bias = bias.has_value();
+  void* bias_ptr = with_bias ? bias->data_ptr() : nullptr;
+
+#ifdef USE_MOE_JIT
+  std::string jit_err;
+  TORCH_CHECK(
+      sgl::moe_jit::fp8_w8a16_grouped_gemm_launch(
+          avg_m,
+          scale_count,
+          with_bias,
+          &queue,
+          activations.data_ptr(),
+          weights.data_ptr(),
+          weight_scales.data_ptr(),
+          bias_ptr,
+          output.data_ptr(),
+          gemm_n,
+          gemm_k,
+          total_rows_for_experts.data_ptr<int>(),
+          static_cast<int>(n_experts),
+          atomic_buffer.data_ptr<int>(),
+          ld_b,
+          weight_scale_blocked,
+          static_scheduler,
+          jit_arch_code(),
+          &jit_err),
+      jit_err);
+#else
+  if (scale_count == 3) {
+    if (with_bias) {
+      DISPATCH_MOE_FP8_W8A16_BLOCK_TILES(true);
+    } else {
+      DISPATCH_MOE_FP8_W8A16_BLOCK_TILES(false);
+    }
+  } else if (scale_count == 1) {
+    if (with_bias) {
+      DISPATCH_MOE_FP8_W8A16_SCALAR_TILES(1, true);
+    } else {
+      DISPATCH_MOE_FP8_W8A16_SCALAR_TILES(1, false);
+    }
+  } else if (with_bias) {
+    DISPATCH_MOE_FP8_W8A16_SCALAR_TILES(2, true);
+  } else {
+    DISPATCH_MOE_FP8_W8A16_SCALAR_TILES(2, false);
+  }
+#endif
+}
+
+#undef DISPATCH_MOE_FP8_W8A16_BLOCK_TILES
+#undef DISPATCH_MOE_FP8_W8A16_SCALAR_TILES
+#undef LAUNCH_MOE_FP8_W8A16
+
+#undef SYCL_INTEL_TARGET
diff --git a/src/sycl/GroupGemmFp8W8A16Xe20LauncherInstance.cpp.in b/src/sycl/GroupGemmFp8W8A16Xe20LauncherInstance.cpp.in
new file mode 100644
index 00000000..2f02b839
--- /dev/null
+++ b/src/sycl/GroupGemmFp8W8A16Xe20LauncherInstance.cpp.in
@@ -0,0 +1,139 @@
+#define SYCL_INTEL_TARGET 20
+
+#include <ATen/ATen.h>
+#include <c10/xpu/XPUStream.h>
+#include <torch/all.h>
+
+#include <cute/tensor.hpp>
+
+#include "sycl/Utils.h"
+#include "cutlass/gemm/device/gemm_universal_adapter.h"
+#include "cutlass/gemm/group_array_problem_shape.hpp"
+#include "sycl/kernels/moe/xe20/fp8/moe_kernel.hpp"
+
+using namespace cute;
+
+template <typename, typename, typename, typename, typename, typename, bool, int>
+class GemmXe20Fp8W8A16Name;
+
+template <typename Tile, typename SGLayout, bool WithBias, int WeightScaleCount>
+void Xe20MoEGEMMFp8W8A16Launcher(
+    sycl::queue q,
+    const void* activations,
+    const void* weights,
+    const void* weight_scales,
+    const void* bias,
+    void* outputs,
+    const int gemm_n,
+    const int gemm_k,
+    const int* num_rows_per_expert_device,
+    const int num_experts,
+    int* workspace,
+    int ld_b_param,
+    bool weight_scale_blocked,
+    bool static_scheduler) {
+  using ElementA = cutlass::bfloat16_t;
+  using ElementD = cutlass::bfloat16_t;
+  auto make_dummy_tensor = [&](auto val, auto stride) {
+    return make_tensor(make_gmem_ptr(&val), make_layout(repeat<rank_v<decltype(stride)>>(1), stride));
+  };
+  auto make_dummy_bias = [&](auto val) {
+    return make_tensor(make_gmem_ptr(&val), make_layout(Shape<int>{}, Stride<_1>{}));
+  };
+  using TensorA = decltype(make_dummy_tensor(ElementA{}, Stride<int, _1>{}));
+  using TensorBPacked = decltype(make_dummy_tensor(cutlass::float_e4m3_t{}, Stride<int, _1>{}));
+  using TensorD = decltype(make_dummy_tensor(ElementD{}, Stride<int, _1>{}));
+  using TensorBias = decltype(make_dummy_bias(float{}));
+  using MMA = typename TiledMMAHelper<
+      MMA_Atom<XE_DPAS_TT<8, float, cutlass::bfloat16_t>>, Layout<Tile>, SGLayout>::TiledMMA;
+  auto mma = MMA{};
+  int sm_count = cutlass::KernelHardwareInfo::query_device_multiprocessor_count(0);
+  auto max_threads_per_workgroup = size(mma);
+  static constexpr int MaxThreadsPerSM = 512;
+  TORCH_CHECK(
+      MaxThreadsPerSM % max_threads_per_workgroup == 0,
+      "MaxThreadsPerSM must be divisible by MaxThreadsPerWorkgroup");
+  sycl::range<3> local(1, 1, max_threads_per_workgroup);
+  sycl::range<3> global(1, sm_count * MaxThreadsPerSM / max_threads_per_workgroup, 1);
+  namespace syclex = sycl::ext::oneapi::experimental;
+  namespace intelex = sycl::ext::intel::experimental;
+  syclex::properties kernel_props{syclex::sub_group_size<16>, intelex::grf_size<256>};
+  using Kernel = MoE_FP8::MoEGEMMFp8Weight<
+      Tile,
+      SGLayout,
+      TensorA,
+      TensorBPacked,
+      TensorD,
+      TensorBias,
+      MMA,
+      WithBias,
+      WeightScaleCount != 3,
+      WeightScaleCount == 3,
+      WeightScaleCount == 1>;
+  typename Kernel::Params params{
+      static_cast<const uint8_t*>(activations),
+      static_cast<const uint8_t*>(weights),
+      static_cast<const float*>(weight_scales),
+      static_cast<const float*>(bias),
+      static_cast<ElementD*>(outputs),
+      num_rows_per_expert_device,
+      gemm_n,
+      gemm_k,
+      num_experts,
+      workspace,
+      mma,
+      static_cast<int32_t>(ld_b_param),
+      weight_scale_blocked,
+      static_scheduler};
+  q.submit([&](sycl::handler& h) {
+    sycl::local_accessor<int32_t, 1> local_mem(sycl::range<1>(1), h);
+    h.parallel_for<GemmXe20Fp8W8A16Name<
+        Tile, SGLayout, TensorA, TensorBPacked, TensorD, ElementD, WithBias, WeightScaleCount>>(
+        sycl::nd_range<3>(global * local, local), kernel_props, [=](sycl::nd_item<3> item) {
+          int32_t* slm_mem =
+              static_cast<int32_t*>(local_mem.template get_multi_ptr<sycl::access::decorated::no>().get());
+          Kernel{}(params, item, slm_mem);
+        });
+  });
+}
+
+template __attribute__((visibility("default"))) void Xe20MoEGEMMFp8W8A16Launcher<
+    @TILE@, @SGLAYOUT@, @WITH_BIAS@, @SCALE_COUNT@>(
+    sycl::queue, const void*, const void*, const void*, const void*, void*, const int, const int,
+    const int*, const int, int*, int, bool, bool);
+
+  #ifdef SGL_FP8_W8A16_JIT_ENTRY
+  extern "C" __attribute__((visibility("default"))) void sgl_moe_fp8_w8a16_entry(
+    void* queue_ptr,
+    const void* activations,
+    const void* weights,
+    const void* weight_scales,
+    const void* bias,
+    void* outputs,
+    int gemm_n,
+    int gemm_k,
+    const int* rows_per_expert,
+    int num_experts,
+    int* workspace,
+    int ld_b,
+    bool weight_scale_blocked,
+    bool static_scheduler) {
+    Xe20MoEGEMMFp8W8A16Launcher<@TILE@, @SGLAYOUT@, @WITH_BIAS@, @SCALE_COUNT@>(
+      *static_cast<sycl::queue*>(queue_ptr),
+      activations,
+      weights,
+      weight_scales,
+      bias,
+      outputs,
+      gemm_n,
+      gemm_k,
+      rows_per_expert,
+      num_experts,
+      workspace,
+      ld_b,
+      weight_scale_blocked,
+      static_scheduler);
+  }
+  #endif
+
+#undef SYCL_INTEL_TARGET
diff --git a/src/sycl/MoEPrepareInputs.cpp b/src/sycl/MoEPrepareInputs.cpp
index 42ad7472..edf49bf6 100644
--- a/src/sycl/MoEPrepareInputs.cpp
+++ b/src/sycl/MoEPrepareInputs.cpp
@@ -3,7 +3,10 @@
 #include <c10/xpu/XPUStream.h>
 #include <torch/all.h>
 
+#include <algorithm>
+#include <limits>
 #include <sycl/sycl.hpp>
+#include <unordered_map>
 
 #include "SYCLHelpers.h"
 #include "Utils.h"
@@ -496,6 +499,250 @@ SGL_KERNEL_EXPORT void prepare_moe_input(
   return;
 }
 
+template <typename IndexType, typename ScalarT>
+struct PrepareMoeInputSmall : public __SYCL_KER_CONFIG_CONVENTION__ {
+  // TODO: Add benchmarked WG/max-route specializations when this path is
+  // enabled beyond BMG; select among static variants using device capabilities.
+  static constexpr int WGSize = 256;
+  static constexpr int MaxRoutes = 64;
+  static constexpr int ElementsPerVector = 8;
+  static constexpr int RequiredSubGroupSize = 16;
+
+  static_assert(WGSize % RequiredSubGroupSize == 0);
+  static_assert(MaxRoutes <= WGSize);
+  static_assert(ElementsPerVector * sizeof(ScalarT) == 16);
+
+  PrepareMoeInputSmall(
+      const ScalarT* input,
+      const IndexType* topk_ids,
+      int32_t* expert_counts,
+      int32_t* output_permutation,
+      ScalarT* output,
+      int32_t num_experts,
+      int32_t input_rows,
+      int32_t topk,
+      int32_t hidden_dim)
+      : input_(input),
+        topk_ids_(topk_ids),
+        expert_counts_(expert_counts),
+        output_permutation_(output_permutation),
+        output_(output),
+        num_experts_(num_experts),
+        input_rows_(input_rows),
+        topk_(topk),
+        hidden_dim_(hidden_dim) {}
+
+  void sycl_ker_config_convention(sycl::handler& cgh) {
+    route_positions_ = sycl::local_accessor<int32_t, 1>(MaxRoutes, cgh);
+    local_counts_ = sycl::local_accessor<int32_t, 1>(num_experts_, cgh);
+  }
+
+  [[sycl::reqd_sub_group_size(RequiredSubGroupSize)]] void operator()(sycl::nd_item<1> item) const {
+    int local_id = item.get_local_linear_id();
+    for (int expert = local_id; expert < num_experts_; expert += WGSize) {
+      expert_counts_[expert] = 0;
+      local_counts_[expert] = 0;
+    }
+    sycl::group_barrier(item.get_group());
+
+    if (input_rows_ == 1) {
+      if (local_id == 0) {
+        int32_t order[16];
+        for (int rank = 0; rank < topk_; ++rank) {
+          order[rank] = rank;
+        }
+        for (int rank = 1; rank < topk_; ++rank) {
+          int32_t current = order[rank];
+          int insert_at = rank;
+          while (insert_at > 0 && topk_ids_[order[insert_at - 1]] > topk_ids_[current]) {
+            order[insert_at] = order[insert_at - 1];
+            --insert_at;
+          }
+          order[insert_at] = current;
+        }
+        for (int destination = 0; destination < topk_; ++destination) {
+          int32_t route = order[destination];
+          ++expert_counts_[static_cast<int32_t>(topk_ids_[route])];
+          route_positions_[route] = destination;
+          output_permutation_[route] = destination;
+        }
+      }
+      sycl::group_barrier(item.get_group());
+
+      using Vector = sycl::vec<uint16_t, ElementsPerVector>;
+      auto input_vectors = reinterpret_cast<const Vector*>(input_);
+      auto output_vectors = reinterpret_cast<Vector*>(output_);
+      int vector_count = hidden_dim_ % ElementsPerVector == 0 ? hidden_dim_ / ElementsPerVector : 0;
+      for (int vector_id = local_id; vector_id < vector_count; vector_id += WGSize) {
+        Vector value = input_vectors[vector_id];
+        for (int route = 0; route < topk_; ++route) {
+          output_vectors[route_positions_[route] * vector_count + vector_id] = value;
+        }
+      }
+      for (int column = vector_count * ElementsPerVector + local_id; column < hidden_dim_; column += WGSize) {
+        ScalarT value = input_[column];
+        for (int route = 0; route < topk_; ++route) {
+          output_[route_positions_[route] * hidden_dim_ + column] = value;
+        }
+      }
+      return;
+    }
+
+    int route_count = topk_ * input_rows_;
+    if (local_id < route_count) {
+      int32_t expert = static_cast<int32_t>(topk_ids_[local_id]);
+      sycl::atomic_ref<
+          int32_t,
+          sycl::memory_order::relaxed,
+          sycl::memory_scope::work_group,
+          sycl::access::address_space::local_space>
+          count(local_counts_[expert]);
+      count.fetch_add(1);
+    }
+    sycl::group_barrier(item.get_group());
+
+    if (local_id == 0) {
+      int32_t offset = 0;
+      for (int expert = 0; expert < num_experts_; ++expert) {
+        int32_t count = local_counts_[expert];
+        expert_counts_[expert] = count;
+        local_counts_[expert] = offset;
+        offset += count;
+      }
+    }
+    sycl::group_barrier(item.get_group());
+
+    if (local_id < route_count) {
+      int32_t expert = static_cast<int32_t>(topk_ids_[local_id]);
+      int32_t destination = local_counts_[expert];
+      for (int route = 0; route < local_id; ++route) {
+        destination += static_cast<int32_t>(topk_ids_[route]) == expert;
+      }
+      route_positions_[local_id] = destination;
+      output_permutation_[local_id] = destination;
+    }
+    sycl::group_barrier(item.get_group());
+
+    using Vector = sycl::vec<uint16_t, ElementsPerVector>;
+    auto input_vectors = reinterpret_cast<const Vector*>(input_);
+    auto output_vectors = reinterpret_cast<Vector*>(output_);
+    int vector_count = hidden_dim_ % ElementsPerVector == 0 ? hidden_dim_ / ElementsPerVector : 0;
+    int vector_tasks = route_count * vector_count;
+    for (int task = local_id; task < vector_tasks; task += WGSize) {
+      int route = task / vector_count;
+      int vector_id = task % vector_count;
+      int source_row = route / topk_;
+      output_vectors[route_positions_[route] * vector_count + vector_id] =
+          input_vectors[source_row * vector_count + vector_id];
+    }
+    int tail_start = vector_count * ElementsPerVector;
+    int tail_size = hidden_dim_ - tail_start;
+    int tail_tasks = route_count * tail_size;
+    for (int task = local_id; task < tail_tasks; task += WGSize) {
+      int route = task / tail_size;
+      int column = tail_start + task % tail_size;
+      int source_row = route / topk_;
+      output_[route_positions_[route] * hidden_dim_ + column] = input_[source_row * hidden_dim_ + column];
+    }
+  }
+
+  const ScalarT* input_;
+  const IndexType* topk_ids_;
+  int32_t* expert_counts_;
+  int32_t* output_permutation_;
+  ScalarT* output_;
+  int32_t num_experts_;
+  int32_t input_rows_;
+  int32_t topk_;
+  int32_t hidden_dim_;
+  mutable sycl::local_accessor<int32_t, 1> route_positions_;
+  mutable sycl::local_accessor<int32_t, 1> local_counts_;
+};
+
+template <typename Kernel>
+size_t prepare_moe_input_small_local_memory_capacity(at::DeviceIndex device_index) {
+  static thread_local std::unordered_map<at::DeviceIndex, size_t> capacity_by_device;
+  auto cached = capacity_by_device.find(device_index);
+  if (cached != capacity_by_device.end()) {
+    return cached->second;
+  }
+
+  auto* properties = at::xpu::getDeviceProperties(device_index);
+  TORCH_CHECK(
+      std::find(properties->sub_group_sizes.begin(), properties->sub_group_sizes.end(), Kernel::RequiredSubGroupSize) !=
+          properties->sub_group_sizes.end(),
+      "prepare_moe_input_small requires subgroup size ",
+      Kernel::RequiredSubGroupSize);
+  TORCH_CHECK(
+      dpcppMaxWorkGroupSize<Kernel>(device_index) >= Kernel::WGSize,
+      "prepare_moe_input_small requires work-group size ",
+      Kernel::WGSize);
+  return capacity_by_device.emplace(device_index, properties->local_mem_size).first->second;
+}
+
+SGL_KERNEL_EXPORT void prepare_moe_input_small(
+    const torch::Tensor& input,
+    const torch::Tensor& topk_ids,
+    torch::Tensor& expert_counts,
+    torch::Tensor& output_permutation,
+    torch::Tensor& output) {
+  TORCH_CHECK(
+      input.is_xpu() && input.dim() == 2 && input.is_contiguous(), "input must be contiguous XPU [rows, hidden_dim]");
+  TORCH_CHECK(
+      topk_ids.is_xpu() && topk_ids.dim() == 2 && topk_ids.is_contiguous() && topk_ids.size(0) == input.size(0),
+      "topk_ids must be contiguous XPU [rows, topk] with rows matching input");
+  TORCH_CHECK(topk_ids.size(1) > 0 && topk_ids.size(1) <= 16, "topk must be in [1, 16]");
+  TORCH_CHECK(
+      topk_ids.numel() <= (PrepareMoeInputSmall<int32_t, c10::BFloat16>::MaxRoutes), "routed rows must be <= 64");
+  TORCH_CHECK(input.scalar_type() == at::kBFloat16, "input must be bfloat16");
+  TORCH_CHECK(
+      expert_counts.is_xpu() && expert_counts.dim() == 1 && expert_counts.numel() > 0 && expert_counts.is_contiguous(),
+      "expert_counts must be a non-empty contiguous XPU vector");
+  TORCH_CHECK(
+      output_permutation.is_xpu() && output_permutation.dim() == 1 && output_permutation.is_contiguous(),
+      "output_permutation must be a contiguous XPU vector");
+  TORCH_CHECK(
+      output.is_xpu() && output.dim() == 2 && output.is_contiguous(),
+      "output must be contiguous XPU [routes, hidden_dim]");
+  TORCH_CHECK(
+      input.device() == topk_ids.device() && input.device() == expert_counts.device() &&
+          input.device() == output_permutation.device() && input.device() == output.device(),
+      "all tensors must be on the same device");
+  TORCH_CHECK(expert_counts.scalar_type() == at::kInt, "expert_counts must be int32");
+  TORCH_CHECK(output_permutation.scalar_type() == at::kInt, "output_permutation must be int32");
+  TORCH_CHECK(output.scalar_type() == input.scalar_type(), "output dtype must match input");
+  TORCH_CHECK(output.size(0) == topk_ids.numel() && output.size(1) == input.size(1), "output shape mismatch");
+  TORCH_CHECK(output_permutation.numel() == topk_ids.numel(), "output_permutation shape mismatch");
+  TORCH_CHECK(
+      expert_counts.numel() <= std::numeric_limits<int32_t>::max() &&
+          input.size(0) <= std::numeric_limits<int32_t>::max() && input.size(1) <= std::numeric_limits<int32_t>::max(),
+      "expert count and input dimensions must fit in int32");
+
+  auto queue = at::xpu::getCurrentXPUStream().queue();
+  AT_DISPATCH_INDEX_TYPES(topk_ids.scalar_type(), "prepare_moe_input_small", [&] {
+    using Kernel = PrepareMoeInputSmall<index_t, c10::BFloat16>;
+    const size_t local_memory_capacity = prepare_moe_input_small_local_memory_capacity<Kernel>(input.device().index());
+    const size_t required_local_memory =
+        (static_cast<size_t>(Kernel::MaxRoutes) + static_cast<size_t>(expert_counts.numel())) * sizeof(int32_t);
+    TORCH_CHECK(
+        required_local_memory <= local_memory_capacity,
+        "prepare_moe_input_small requires ",
+        required_local_memory,
+        " bytes of local memory");
+    Kernel task(
+        input.const_data_ptr<c10::BFloat16>(),
+        topk_ids.const_data_ptr<index_t>(),
+        expert_counts.mutable_data_ptr<int32_t>(),
+        output_permutation.mutable_data_ptr<int32_t>(),
+        output.mutable_data_ptr<c10::BFloat16>(),
+        static_cast<int32_t>(expert_counts.numel()),
+        static_cast<int32_t>(input.size(0)),
+        static_cast<int32_t>(topk_ids.size(1)),
+        static_cast<int32_t>(input.size(1)));
+    sycl_kernel_submit(Kernel::WGSize, Kernel::WGSize, queue, task);
+  });
+}
+
 // Scatter kernel: 1 WG per source token, reads token once, scatters to topk destinations.
 // Equivalent to IPEX MoEScatter but uses precomputed src2dst_map (c_map / output_permutation).
 template <typename T>
diff --git a/src/sycl/kernels/moe/xe20/fp8/moe_kernel.hpp b/src/sycl/kernels/moe/xe20/fp8/moe_kernel.hpp
new file mode 100644
index 00000000..c9e74f7f
--- /dev/null
+++ b/src/sycl/kernels/moe/xe20/fp8/moe_kernel.hpp
@@ -0,0 +1,199 @@
+/***************************************************************************************************
+ * Copyright (C) 2025 Intel Corporation, All rights reserved.
+ * SPDX-License-Identifier: BSD-3-Clause
+ **************************************************************************************************/
+
+// FP8 E4M3 weight, BF16 activation grouped GEMM for Xe2 (BMG). Activation is
+// applied externally between GEMM1 and GEMM2.
+
+#pragma once
+
+#include "../common/block_2d_copy_d.hpp"
+#include "../w4a16/gemm_xe2.hpp"
+#include "cute/tensor.hpp"
+#include "cutlass/cutlass.h"
+#include "cutlass/float8.h"
+#include "cutlass/gemm/gemm.h"
+#include "cutlass/gemm/group_array_problem_shape.hpp"
+#include "cutlass/gemm/kernel/tile_scheduler.hpp"
+#include "cutlass/kernel_hardware_info.hpp"
+#include "cutlass/platform/platform.h"
+#include "cutlass/util/packed_stride.hpp"
+#include "moe_mainloop.hpp"
+
+#pragma clang diagnostic ignored "-Wpass-failed"
+#pragma clang diagnostic ignored "-Wdeprecated-declarations"
+
+namespace MoE_FP8 {
+using namespace cute;
+
+template <
+    typename TileShape,
+    typename SubgroupLayout,
+    typename TensorA,
+    typename TensorBPacked,
+    typename TensorD,
+    typename TensorBias,
+    typename TiledMMA,
+    bool WithBias,
+    bool WeightScalePerExpert = false,
+    bool WeightScaleBlocked = false,
+    bool SingleWeightScale = false>
+class MoEGEMMFp8Weight {
+ public:
+  using ElementA = cutlass::bfloat16_t;
+  using ElementD = cutlass::bfloat16_t;
+  using TiledCopyA = decltype(make_block_2d_copy_A(TiledMMA{}, TensorA{}));
+  using TiledCopyBPacked = decltype(make_block_2d_copy_B(TiledMMA{}, TensorBPacked{}));
+  using TiledCopyD = decltype(moe_xe20::make_moe_block_2d_copy_D<void>(TiledMMA{}, TensorD{}));
+  using SGPerWG = decltype(product(take<1, 4>(shape(typename TiledMMA::ThrLayoutVMNK{}))));
+
+  constexpr static int Stages = 3;
+  using MainloopDispatchPolicy = MoE_FP8::XeDefault<Stages>;
+  using CollectiveMainloop = MoEMainloopFp8Weight<
+      MainloopDispatchPolicy,
+      TiledCopyA,
+      TiledCopyBPacked,
+      TiledCopyD,
+      TensorA,
+      TensorBPacked,
+      TensorD,
+      TensorBias,
+      TiledMMA,
+      WithBias,
+      WeightScalePerExpert,
+      WeightScaleBlocked>;
+
+  struct Params {
+    const uint8_t* Activations;    // [M_total, K] bf16 bytes
+    const uint8_t* PackedWeights;  // [num_experts, N, K] fp8 e4m3 raw bytes
+    const float* WeightScales;     // Per-expert scalar or [E, ceil(N/128), K/128] block scales
+    const float* Bias;
+    ElementD* Outputs;
+    const int32_t* M_per_group;
+    const int32_t N;
+    const int32_t K;
+    const int32_t num_experts;
+    int32_t* workspace;
+    TiledMMA mma;
+    int32_t ld_b;
+    bool weight_scale_blocked = false;
+    bool static_scheduler = false;
+  };
+
+  auto make_A_tensor(uint8_t* ptr_A, int M, int K) {
+    auto* bf16_ptr = reinterpret_cast<cutlass::bfloat16_t*>(ptr_A);
+    return make_tensor(make_gmem_ptr(bf16_ptr), make_layout(make_shape(M, K), make_stride(K, _1{})));
+  }
+
+  auto make_B_tensors(uint8_t* ptr_B, int N, int K, int ld_b) {
+    auto* e4m3_ptr = reinterpret_cast<cutlass::float_e4m3_t*>(ptr_B);
+    auto B = make_tensor(make_gmem_ptr(e4m3_ptr), make_layout(make_shape(N, K), make_stride(ld_b, _1{})));
+    return B;
+  }
+
+  // Per-expert weight-scale pointers + row stride. Mirrors MXFP4's
+  // make_scale_ptrs exactly (same gate/up split conventions), just with
+  // K_scale = K / FP8_GROUP_SIZE_K instead of K / MXFP4_GROUP_SIZE.
+  auto make_Bias_tensors(float* ptr_Bias, int N) {
+    return make_tensor(make_gmem_ptr<float>(ptr_Bias), make_layout(make_shape(N), make_stride(_1{})));
+  }
+
+  auto make_D_tensors(ElementD* ptr_D, int pre_rows, int M, int N) {
+    return make_tensor(
+        make_gmem_ptr<ElementD>(ptr_D + pre_rows * N), make_layout(make_shape(M, N), make_stride(N, _1{})));
+  }
+
+  CUTLASS_DEVICE void operator()(Params const& params, sycl::nd_item<3> item, int32_t* slm_mem) {
+    auto N = params.N;
+    auto K = params.K;
+    auto M_per_group = params.M_per_group;
+    auto num_experts = params.num_experts;
+    auto mma = params.mma;
+    auto workspace = params.workspace;
+
+    auto wg_tile = mma.tile_mnk();
+    auto wg_tile_m = get<0>(wg_tile);
+    auto wg_tile_n = get<1>(wg_tile);
+
+    int group_id = item.get_group_linear_id();
+    int N_pad = ceil_div(N, wg_tile_n) * wg_tile_n;
+    int group_m_id = (group_id * wg_tile_n) / N_pad;
+    int group_range = item.get_group_range(1);
+    int32_t thr_id = int32_t(item.get_local_linear_id());
+
+    const int64_t K_scale = K / FP8_GROUP_SIZE_K;
+    int64_t scale_n =
+        WeightScalePerExpert ? (SingleWeightScale ? 1 : 2) : (params.weight_scale_blocked ? (N + 127) / 128 : N);
+
+    int pre_rows = 0;
+    int pre_tiles = 0;
+    for (int i = 0; i < num_experts; ++i) {
+      int M = M_per_group[i];
+      int cumsum_rows_for_experts = M + pre_rows;
+      int cumsum_tiles_for_experts = (M + wg_tile_m - 1) / wg_tile_m + pre_tiles;
+
+      if (group_m_id >= cumsum_tiles_for_experts) {
+        pre_rows = cumsum_rows_for_experts;
+        pre_tiles = cumsum_tiles_for_experts;
+        continue;
+      }
+
+      int expert_id = i;
+      int ld_b = params.ld_b;
+      int64_t B_offset = static_cast<int64_t>(expert_id) * static_cast<int64_t>(N) * static_cast<int64_t>(ld_b);
+      int64_t S_offset = static_cast<int64_t>(expert_id) * scale_n * (WeightScalePerExpert ? 1 : K_scale);
+
+      uint8_t* ptr_A_curr_batch = const_cast<uint8_t*>(params.Activations) + pre_rows * K * sizeof(ElementA);
+      uint8_t* ptr_B_curr_batch = const_cast<uint8_t*>(params.PackedWeights) + B_offset;
+      float* ptr_S_curr_batch = const_cast<float*>(params.WeightScales) + S_offset;
+      float* ptr_Bias_curr_batch = nullptr;
+      if constexpr (WithBias) {
+        ptr_Bias_curr_batch = const_cast<float*>(params.Bias) + expert_id * N;
+      }
+
+      auto A_tensor = make_A_tensor(ptr_A_curr_batch, M, K);
+      auto B_tensor = make_B_tensors(ptr_B_curr_batch, N, K, ld_b);
+      auto D_tensor = make_D_tensors(params.Outputs, pre_rows, M, N);
+      auto Bias_tensor = make_Bias_tensors(ptr_Bias_curr_batch, N);
+
+      while (group_m_id < cumsum_tiles_for_experts) {
+        int n_coord = (group_id * wg_tile_n) % N_pad / wg_tile_n;
+        int m_coord = (group_m_id - pre_tiles);
+
+        auto tile_coord = make_coord(m_coord, n_coord, _, 0);
+        if constexpr (SingleWeightScale) {
+          moe_w4a16::xe_gemm<void, void, void>(
+              A_tensor, B_tensor, ptr_S_curr_batch, ptr_Bias_curr_batch, D_tensor, tile_coord, mma);
+        } else {
+          CollectiveMainloop mainloop;
+          mainloop(
+              A_tensor,
+              B_tensor,
+              ptr_S_curr_batch,
+              WeightScalePerExpert ? scale_n : K_scale,
+              D_tensor,
+              tile_coord,
+              mma,
+              thr_id,
+              Bias_tensor,
+              N);
+        }
+
+        if (params.static_scheduler) {
+          group_id += group_range;
+        } else {
+          if (thr_id == 0) {
+            slm_mem[0] = cutlass::atomicAdd(workspace, 1);
+          }
+          item.barrier(sycl::access::fence_space::local_space);
+          group_id = group_range + slm_mem[0];
+        }
+        group_m_id = (group_id * wg_tile_n) / N_pad;
+      }
+      pre_rows = cumsum_rows_for_experts;
+      pre_tiles = cumsum_tiles_for_experts;
+    }
+  };
+};
+}  // namespace MoE_FP8
diff --git a/src/sycl/kernels/moe/xe20/fp8/moe_mainloop.hpp b/src/sycl/kernels/moe/xe20/fp8/moe_mainloop.hpp
new file mode 100644
index 00000000..869ba7ca
--- /dev/null
+++ b/src/sycl/kernels/moe/xe20/fp8/moe_mainloop.hpp
@@ -0,0 +1,386 @@
+/***************************************************************************************************
+ * Copyright (C) 2025 Intel Corporation, All rights reserved.
+ * SPDX-License-Identifier: BSD-3-Clause
+ **************************************************************************************************/
+
+// FP8 E4M3 weight, BF16 activation grouped-GEMM mainloop for Xe2 (BMG).
+// Weight scales are either one scalar per expert/projection or one value per
+// 128x128 weight block.
+
+#pragma once
+
+#include <cute/tensor.hpp>
+#include <cute/util/compat.hpp>
+#include <sycl/ext/intel/experimental/grf_size_properties.hpp>
+#include <sycl/sycl.hpp>
+
+#include "cutlass/float8.h"
+#include "cutlass/half.h"
+#include "cutlass/kernel_hardware_info.h"
+#include "cutlass/platform/platform.h"
+#include "cutlass/tensor_ref.h"
+#include "cutlass/util/GPU_Clock.hpp"
+#include "cutlass/util/reference/device/gemm_complex.h"
+#include "cutlass/util/reference/device/tensor_compare.h"
+#include "cutlass/util/reference/host/tensor_fill.h"
+#include "cutlass/util/sycl_event_manager.hpp"
+#include "sycl/SYCLHelpers.h"
+
+#pragma clang diagnostic ignored "-Wpass-failed"
+#pragma clang diagnostic ignored "-Wdeprecated-declarations"
+
+namespace MoE_FP8 {
+
+using namespace cute;
+
+static constexpr int FP8_GROUP_SIZE_K = 128;
+
+// Number of work-items per subgroup on Xe (SIMD lane count).
+static constexpr int SUBGROUP_SIZE = 16;
+
+template <typename Element>
+CUTE_DEVICE Element apply_bf16_weight_scale(Element value, float scale) {
+  static_assert(cute::is_same_v<Element, cutlass::bfloat16_t> || cute::is_same_v<Element, cutlass::half_t>);
+  uint16_t bits = sycl::bit_cast<uint16_t>(value);
+#if defined(__SYCL_DEVICE_ONLY__) && defined(SYCL_INTEL_TARGET)
+  if constexpr (cute::is_same_v<Element, cutlass::bfloat16_t>) {
+    asm("{\n"
+        ".decl Z_BF16 v_type=G type=BF num_elts=16 alias=<%0,0>\n"
+        ".decl Y_FP32 v_type=G type=F num_elts=16 alias=<%1,0>\n"
+        "mul (M1, 16) Z_BF16(0,0)<1> Z_BF16(0,0)<1;1,0> Y_FP32(0,0)<1;1,0>\n"
+        "}\n"
+        : "+rw"(bits)
+        : "rw"(scale));
+  } else {
+    asm("{\n"
+        ".decl Z_FP16 v_type=G type=HF num_elts=16 alias=<%0,0>\n"
+        ".decl Y_FP32 v_type=G type=F num_elts=16 alias=<%1,0>\n"
+        "mul (M1, 16) Z_FP16(0,0)<1> Z_FP16(0,0)<1;1,0> Y_FP32(0,0)<1;1,0>\n"
+        "}\n"
+        : "+rw"(bits)
+        : "rw"(scale));
+  }
+#else
+  return Element(static_cast<float>(value) * scale);
+#endif
+  return sycl::bit_cast<Element>(bits);
+}
+
+template <int Stages>
+class XeDefault {};
+
+template <
+    class DispatchPolicy_,
+    class TiledCopyA_,
+    class TiledCopyBPacked_,
+    class TiledCopyD_,
+    class ATensor_,
+    class BPackedTensor_,
+    class DTensor_,
+    class BiasTensor_,
+    class TiledMMA_,
+    bool WithBias,
+    bool WeightScalePerExpert = false,
+    bool WeightScaleBlocked = false>
+struct MoEMainloopFp8Weight {
+  static_assert(cutlass::detail::dependent_false<DispatchPolicy_>, "Could not find a mainloop specialization.");
+};
+
+template <
+    int Stages,
+    class TiledCopyA_,
+    class TiledCopyBPacked_,
+    class TiledCopyD_,
+    class ATensor_,
+    class BPackedTensor_,
+    class DTensor_,
+    class BiasTensor_,
+    class TiledMMA_,
+    bool WithBias,
+    bool WeightScalePerExpert,
+    bool WeightScaleBlocked>
+struct MoEMainloopFp8Weight<
+    XeDefault<Stages>,
+    TiledCopyA_,
+    TiledCopyBPacked_,
+    TiledCopyD_,
+    ATensor_,
+    BPackedTensor_,
+    DTensor_,
+    BiasTensor_,
+    TiledMMA_,
+    WithBias,
+    WeightScalePerExpert,
+    WeightScaleBlocked> {
+  using TiledMMA = TiledMMA_;
+  using TiledCopyA = TiledCopyA_;
+  using TiledCopyBPacked = TiledCopyBPacked_;
+  using TiledCopyD = TiledCopyD_;
+  using ATensor = ATensor_;
+  using BPackedTensor = BPackedTensor_;
+  using DTensor = DTensor_;
+  using BiasTensor = BiasTensor_;
+
+  MoEMainloopFp8Weight() {}
+
+  template <typename Coord>
+  CUTLASS_DEVICE void run_w8a16_block(
+      ATensor& A,
+      BPackedTensor& Bp,
+      const float* w_scale_gmem,
+      int w_scale_row_stride,
+      DTensor& D,
+      Coord blk_coord,
+      TiledMMA mma,
+      int thr_id,
+      BiasTensor Bias,
+      int gemm_n) {
+    auto wg_m = get<0>(blk_coord);
+    auto wg_n = get<1>(blk_coord);
+    auto wg_tile = mma.tile_mnk();
+    auto wg_coord = make_coord(wg_m, wg_n, 0);
+    constexpr int BLK_M = get<0>(decltype(wg_tile){});
+    constexpr int BLK_N = get<1>(decltype(wg_tile){});
+    constexpr int BLK_K = get<2>(decltype(wg_tile){});
+    constexpr int ATOM_M_V = get<1>(typename TiledMMA::ThrLayoutVMNK{}.shape());
+    constexpr int ATOM_N_V = get<2>(typename TiledMMA::ThrLayoutVMNK{}.shape());
+    constexpr int SG_M = BLK_M / ATOM_M_V;
+    constexpr int SG_N = BLK_N / ATOM_N_V;
+    constexpr int N_ATOMS = SG_N / SUBGROUP_SIZE;
+    constexpr int RELOAD_CADENCE = FP8_GROUP_SIZE_K / BLK_K;
+    static_assert(
+        RELOAD_CADENCE == 4 || RELOAD_CADENCE == 8,
+        "W8A16 block fast path supports K tiles of 32 or 16 per scale group");
+    static_assert(BLK_N <= 128, "W8A16 block fast path requires one N tile per weight-scale block");
+
+    Tensor cA = make_identity_tensor(A.shape());
+    Tensor cBp = make_identity_tensor(Bp.shape());
+    Tensor cD = make_identity_tensor(D.shape());
+    Tensor gA = local_tile(cA, select<0, 2>(wg_tile), make_coord(wg_m, _));
+    Tensor gBp = local_tile(cBp, select<1, 2>(wg_tile), make_coord(wg_n, _));
+    Tensor gD = local_tile(cD, wg_tile, wg_coord, Step<_1, _1, X>{});
+
+    TiledCopyA tiled_copy_a{A};
+    TiledCopyBPacked tiled_copy_b{Bp};
+    TiledCopyD tiled_copy_d{D};
+    auto thr_copy_a = tiled_copy_a.get_slice(thr_id);
+    auto thr_copy_b = tiled_copy_b.get_slice(thr_id);
+    auto thr_copy_d = tiled_copy_d.get_slice(thr_id);
+    auto thr_mma = mma.get_slice(thr_id);
+    auto tAgA = thr_copy_a.partition_S(gA);
+    auto tBgBp = thr_copy_b.partition_S(gBp);
+    using CopyAFragment = decltype(thr_copy_a.partition_sg_fragment_D(gA(_, _, 0)));
+    using CopyBFragment = decltype(thr_copy_b.partition_sg_fragment_D(gBp(_, _, 0)));
+    using MmaAFragment = decltype(thr_mma.partition_sg_fragment_A(gA(_, _, 0)));
+    using MmaBFragment = decltype(thr_mma.partition_sg_fragment_B(gBp(_, _, 0)));
+    CopyAFragment tArA_packed = thr_copy_a.partition_sg_fragment_D(gA(_, _, 0));
+    CopyBFragment tBrB_packed = thr_copy_b.partition_sg_fragment_D(gBp(_, _, 0));
+    MmaAFragment tSrA = thr_mma.partition_sg_fragment_A(gA(_, _, 0));
+    MmaBFragment tSrB = thr_mma.partition_sg_fragment_B(gBp(_, _, 0));
+    SubgroupTensor tCrC = thr_mma.partition_sg_fragment_C(gD);
+    cute::clear(tCrC);
+
+    auto prefetch_a = make_block_2d_prefetch(tiled_copy_a);
+    auto prefetch_b = make_block_2d_prefetch(tiled_copy_b);
+    auto pAgA = prefetch_a.get_slice(thr_id).partition_S(gA);
+    auto pBgBp = prefetch_b.get_slice(thr_id).partition_S(gBp);
+    constexpr SPIRVScope barrier_scope = ScopeWorkgroup;
+    const int k_tile_count = ceil_div(shape<1>(A), BLK_K);
+    const int full_group_count = k_tile_count / RELOAD_CADENCE;
+    CUTE_UNROLL
+    for (int prefetch_k = 0; prefetch_k < Stages; ++prefetch_k) {
+      if (prefetch_k < k_tile_count) {
+        prefetch(prefetch_a, pAgA(_, _, _, prefetch_k));
+        prefetch(prefetch_b, pBgBp(_, _, _, prefetch_k));
+      }
+    }
+
+    float w_scale = 1.0f;
+    const int scale_n = (wg_n * BLK_N) / 128;
+    auto load_group_scale = [&](int group) { w_scale = w_scale_gmem[scale_n * w_scale_row_stride + group]; };
+    load_group_scale(0);
+
+    auto run_k_tile = [&](int k_tile) {
+      barrier_arrive(barrier_scope);
+      copy(tiled_copy_a, tAgA(_, _, _, k_tile), tArA_packed);
+      copy(tiled_copy_b, tBgBp(_, _, _, k_tile), tBrB_packed);
+      const int prefetch_idx = k_tile + Stages;
+      if (prefetch_idx < k_tile_count) {
+        prefetch(prefetch_a, pAgA(_, _, _, prefetch_idx));
+        prefetch(prefetch_b, pBgBp(_, _, _, prefetch_idx));
+      }
+      reorder(tArA_packed, tSrA);
+      reorder(tBrB_packed, tSrB);
+      CUTLASS_PRAGMA_UNROLL
+      for (int i = 0; i < tSrB.size(); ++i) {
+        tSrB(i) = apply_bf16_weight_scale(tSrB(i), w_scale);
+      }
+      cute::gemm(mma, tSrA, tSrB, tCrC);
+      barrier_wait(barrier_scope);
+    };
+
+    for (int group = 0; group < full_group_count; ++group) {
+      CUTE_UNROLL
+      for (int group_offset = 0; group_offset < RELOAD_CADENCE; ++group_offset) {
+        run_k_tile(group * RELOAD_CADENCE + group_offset);
+      }
+      if (group + 1 < full_group_count) {
+        load_group_scale(group + 1);
+      }
+    }
+
+    if constexpr (WithBias) {
+      add_bias<SG_M, SG_N, BLK_N>(Bias, tCrC, wg_n, thr_id, gemm_n);
+    }
+    SubgroupTensor tCrD = thr_copy_d.partition_sg_fragment_S(gD);
+    Tensor tCgD = thr_copy_d.partition_D(gD);
+    reorder(tCrC, tCrD);
+    copy(tiled_copy_d, tCrD, tCgD);
+  }
+
+  template <typename Coord>
+  CUTLASS_DEVICE void run_w8a16_scalar(
+      ATensor& A,
+      BPackedTensor& Bp,
+      const float* w_scale_gmem,
+      int weight_scale_count,
+      DTensor& D,
+      Coord blk_coord,
+      TiledMMA mma,
+      int thr_id,
+      BiasTensor Bias,
+      int gemm_n) {
+    auto wg_m = get<0>(blk_coord);
+    auto wg_n = get<1>(blk_coord);
+    auto wg_tile = mma.tile_mnk();
+    auto wg_coord = make_coord(wg_m, wg_n, 0);
+    constexpr int BLK_M = get<0>(decltype(wg_tile){});
+    constexpr int BLK_N = get<1>(decltype(wg_tile){});
+    constexpr int ATOM_M_V = get<1>(typename TiledMMA::ThrLayoutVMNK{}.shape());
+    constexpr int ATOM_N_V = get<2>(typename TiledMMA::ThrLayoutVMNK{}.shape());
+    constexpr int SG_M = BLK_M / ATOM_M_V;
+
+    Tensor cA = make_identity_tensor(A.shape());
+    Tensor cBp = make_identity_tensor(Bp.shape());
+    Tensor cD = make_identity_tensor(D.shape());
+    Tensor gA = local_tile(cA, select<0, 2>(wg_tile), make_coord(wg_m, _));
+    Tensor gBp = local_tile(cBp, select<1, 2>(wg_tile), make_coord(wg_n, _));
+    Tensor gD = local_tile(cD, wg_tile, wg_coord, Step<_1, _1, X>{});
+
+    TiledCopyA tiled_copy_a{A};
+    TiledCopyBPacked tiled_copy_b{Bp};
+    auto thr_copy_a = tiled_copy_a.get_slice(thr_id);
+    auto thr_copy_b = tiled_copy_b.get_slice(thr_id);
+    auto thr_mma = mma.get_slice(thr_id);
+    auto tAgA = thr_copy_a.partition_S(gA);
+    auto tBgBp = thr_copy_b.partition_S(gBp);
+    using CopyAFragment = decltype(thr_copy_a.partition_sg_fragment_D(gA(_, _, 0)));
+    using CopyBFragment = decltype(thr_copy_b.partition_sg_fragment_D(gBp(_, _, 0)));
+    using MmaAFragment = decltype(thr_mma.partition_sg_fragment_A(gA(_, _, 0)));
+    using MmaBFragment = decltype(thr_mma.partition_sg_fragment_B(gBp(_, _, 0)));
+    CopyAFragment tArA_packed = thr_copy_a.partition_sg_fragment_D(gA(_, _, 0));
+    CopyBFragment tBrB_packed = thr_copy_b.partition_sg_fragment_D(gBp(_, _, 0));
+    MmaAFragment tSrA = thr_mma.partition_sg_fragment_A(gA(_, _, 0));
+    MmaBFragment tSrB = thr_mma.partition_sg_fragment_B(gBp(_, _, 0));
+    SubgroupTensor tCrC = thr_mma.partition_sg_fragment_C(gD);
+    cute::clear(tCrC);
+
+    auto prefetch_a = make_block_2d_prefetch(tiled_copy_a);
+    auto prefetch_b = make_block_2d_prefetch(tiled_copy_b);
+    auto pAgA = prefetch_a.get_slice(thr_id).partition_S(gA);
+    auto pBgBp = prefetch_b.get_slice(thr_id).partition_S(gBp);
+    constexpr SPIRVScope barrier_scope = ScopeWorkgroup;
+    const int k_tile_count = ceil_div(shape<1>(A), get<2>(wg_tile));
+    int k_tile_prefetch = 0;
+
+    CUTE_UNROLL
+    for (; k_tile_prefetch < Stages; ++k_tile_prefetch) {
+      if (k_tile_prefetch < k_tile_count) {
+        prefetch(prefetch_a, pAgA(_, _, _, k_tile_prefetch));
+        prefetch(prefetch_b, pBgBp(_, _, _, k_tile_prefetch));
+      }
+    }
+
+    for (int k_tile = 0; k_tile < k_tile_count; ++k_tile, ++k_tile_prefetch) {
+      barrier_arrive(barrier_scope);
+
+      copy(tiled_copy_a, tAgA(_, _, _, k_tile), tArA_packed);
+      copy(tiled_copy_b, tBgBp(_, _, _, k_tile), tBrB_packed);
+
+      if (k_tile_prefetch < k_tile_count) {
+        prefetch(prefetch_a, pAgA(_, _, _, k_tile_prefetch));
+        prefetch(prefetch_b, pBgBp(_, _, _, k_tile_prefetch));
+      }
+
+      reorder(tArA_packed, tSrA);
+      reorder(tBrB_packed, tSrB);
+      cute::gemm(mma, tSrA, tSrB, tCrC);
+
+      barrier_wait(barrier_scope);
+    }
+
+    constexpr int SG_N = BLK_N / ATOM_N_V;
+    auto sg_local_n_coord = cutlass::get_sub_group_id() % ATOM_N_V;
+    int sg_local_id = cutlass::get_sub_group_local_id();
+    constexpr int sg_local_range = 16;
+    int n_tile_start = wg_n * BLK_N;
+    int n_sg_start = sg_local_n_coord * SG_N;
+    CUTLASS_PRAGMA_UNROLL
+    for (int sn = 0; sn < SG_N / sg_local_range; ++sn) {
+      int global_n = n_tile_start + n_sg_start + sn * sg_local_range + sg_local_id;
+      float weight_scale = w_scale_gmem[weight_scale_count == 2 && global_n >= gemm_n / 2 ? 1 : 0];
+      CUTLASS_PRAGMA_UNROLL
+      for (int sm = 0; sm < SG_M; ++sm) {
+        tCrC(sn * SG_M + sm) *= weight_scale;
+      }
+    }
+    if constexpr (WithBias) {
+      add_bias<SG_M, SG_N, BLK_N>(Bias, tCrC, wg_n, thr_id, gemm_n);
+    }
+    TiledCopyD tiled_copy_d{D};
+    auto thr_copy_d = tiled_copy_d.get_slice(thr_id);
+    SubgroupTensor tCrD = thr_copy_d.partition_sg_fragment_S(gD);
+    Tensor tCgD = thr_copy_d.partition_D(gD);
+    reorder(tCrC, tCrD);
+    copy(tiled_copy_d, tCrD, tCgD);
+  }
+
+  template <typename Coord>
+  CUTLASS_DEVICE void operator()(
+      ATensor& A,
+      BPackedTensor& Bp,
+      const float* w_scale_gmem,
+      int w_scale_row_stride,
+      DTensor& D,
+      Coord blk_coord,
+      TiledMMA mma,
+      int thr_id,
+      BiasTensor Bias,
+      int gemm_n) {
+    if constexpr (WeightScalePerExpert) {
+      run_w8a16_scalar(A, Bp, w_scale_gmem, w_scale_row_stride, D, blk_coord, mma, thr_id, Bias, gemm_n);
+    } else {
+      run_w8a16_block(A, Bp, w_scale_gmem, w_scale_row_stride, D, blk_coord, mma, thr_id, Bias, gemm_n);
+    }
+  }
+
+  template <int SG_M, int SG_N, int BLK_N, typename tCrC_t>
+  void add_bias(const BiasTensor& Bias, tCrC_t& tCrC, int wg_n, int thr_id, int gemm_n) {
+    static constexpr auto ATOM_N = get<2>(typename TiledMMA::ThrLayoutVMNK{}.shape());
+    constexpr int N_ATOMS = SG_N / SUBGROUP_SIZE;
+    int sg_local_n_coord = (thr_id / SUBGROUP_SIZE) % ATOM_N;
+    int lane = thr_id % SUBGROUP_SIZE;
+
+    CUTLASS_PRAGMA_UNROLL
+    for (int na = 0; na < N_ATOMS; ++na) {
+      int n = wg_n * BLK_N + sg_local_n_coord * SG_N + na * SUBGROUP_SIZE + lane;
+      float bias = (n < gemm_n) ? static_cast<float>(Bias(n)) : 0.0f;
+      CUTLASS_PRAGMA_UNROLL
+      for (int sm = 0; sm < SG_M; ++sm) {
+        tCrC(na * SG_M + sm) += bias;
+      }
+    }
+  }
+};
+
+}  // namespace MoE_FP8
diff --git a/src/torch_extension_sycl.cc b/src/torch_extension_sycl.cc
index 2c58255f..c49220f5 100644
--- a/src/torch_extension_sycl.cc
+++ b/src/torch_extension_sycl.cc
@@ -216,11 +216,20 @@ TORCH_LIBRARY_FRAGMENT(sgl_kernel, m) {
       "Tensor? zeros, Tensor? bias, Tensor rows_per_expert, int n_experts, bool is_int4, int group_size) -> ()");
   m.impl("moe_grouped_mm_nt_xe20_w4a16", torch::kXPU, &moe_grouped_mm_nt_xe20_w4a16);
 
+  m.def(
+      "moe_grouped_mm_nt_xe20_fp8_w8a16(Tensor! output, Tensor activations, Tensor weights, "
+      "Tensor weight_scales, Tensor? bias, Tensor total_rows_for_experts, int n_experts) -> ()");
+  m.impl("moe_grouped_mm_nt_xe20_fp8_w8a16", torch::kXPU, &moe_grouped_mm_nt_xe20_fp8_w8a16);
+
   m.def(
       "prepare_moe_input(Tensor topk_ids, Tensor! expert_offsets, Tensor? blockscale_offsets, Tensor! problem_sizes1,"
       " Tensor! problem_sizes2, Tensor! input_permutation, Tensor! output_permutation, int num_experts, int n, int k)"
       " -> ()");
   m.impl("prepare_moe_input", torch::kXPU, &prepare_moe_input);
+  m.def(
+      "prepare_moe_input_small(Tensor input, Tensor topk_ids, Tensor! expert_counts, Tensor! output_permutation, "
+      "Tensor! output) -> ()");
+  m.impl("prepare_moe_input_small", torch::kXPU, &prepare_moe_input_small);
   m.def("scatter_tokens_to_experts(Tensor input, Tensor src2dst_map, Tensor! output) -> ()");
   m.impl("scatter_tokens_to_experts", torch::kXPU, &scatter_tokens_to_experts);
   m.def(
diff --git a/tests/test_moe_fused_experts_workspace.py b/tests/test_moe_fused_experts_workspace.py
index b9041ca6..60b98960 100644
--- a/tests/test_moe_fused_experts_workspace.py
+++ b/tests/test_moe_fused_experts_workspace.py
@@ -20,7 +20,12 @@
 import pytest
 import torch
 from sgl_kernel import fused_experts
-from sgl_kernel.moe import _MOE_WS_HEADROOM, _get_moe_ws, _moe_ws_cache
+from sgl_kernel.moe import (
+    _MOE_WS_HEADROOM,
+    _get_moe_ws,
+    _moe_ws_cache,
+    _moe_ws_view_cache,
+)
 from test_moe_gemm import create_random_xpu_tensor, torch_naive_moe
 
 pytestmark = pytest.mark.skipif(
@@ -41,8 +46,10 @@ def _xpu_device() -> torch.device:
 @pytest.fixture(autouse=True)
 def clear_moe_workspace_cache():
     _moe_ws_cache.clear()
+    _moe_ws_view_cache.clear()
     yield
     _moe_ws_cache.clear()
+    _moe_ws_view_cache.clear()
 
 
 def _cache_key(name: str):
@@ -65,13 +72,14 @@ def test_get_moe_ws_first_call_allocates_with_headroom():
 
 
 def test_get_moe_ws_reuses_same_storage_for_equal_or_smaller_shape():
-    _get_moe_ws("foo", (100,), torch.float32, _xpu_device())
+    first_view = _get_moe_ws("foo", (100,), torch.float32, _xpu_device())
     key = _cache_key("foo")
     ptr1 = _moe_ws_cache[key].data_ptr()
     numel1 = _moe_ws_cache[key].numel()
 
     # Same shape: must reuse the exact same underlying buffer.
-    _get_moe_ws("foo", (100,), torch.float32, _xpu_device())
+    second_view = _get_moe_ws("foo", (100,), torch.float32, _xpu_device())
+    assert second_view is first_view
     assert _moe_ws_cache[key].data_ptr() == ptr1
     assert _moe_ws_cache[key].numel() == numel1
 
@@ -79,6 +87,7 @@ def test_get_moe_ws_reuses_same_storage_for_equal_or_smaller_shape():
     t3 = _get_moe_ws("foo", (10,), torch.float32, _xpu_device())
     assert _moe_ws_cache[key].data_ptr() == ptr1
     assert t3.shape == (10,)
+    assert _get_moe_ws("foo", (10,), torch.float32, _xpu_device()) is t3
 
 
 def test_get_moe_ws_grows_for_larger_shape():
diff --git a/tests/test_moe_gemm.py b/tests/test_moe_gemm.py
index c00b807d..33fa3266 100644
--- a/tests/test_moe_gemm.py
+++ b/tests/test_moe_gemm.py
@@ -915,5 +915,772 @@ def test_moe_grouped_mm_nt_xe20_w4a16_mxfp4_input_dtypes(weight_dtype, scale_dty
     )
 
 
+# ---------------------------------------------------------------------------
+# FP8 (E4M3) W8A16 expert-weight helpers
+# ---------------------------------------------------------------------------
+#
+# Quantization is done with plain torch.float8_e4m3fn casts (a numeric cast,
+# not a call into the kernel-under-test), so these helpers are independent of
+# sgl_kernel and can run entirely on CPU. This mirrors the MXFP4 test's
+# philosophy: both the kernel and the reference see the *same rounded*
+# values, so any remaining numerical difference is GEMM arithmetic noise
+# (FP32 reference accumulation vs. the kernel's BF16 DPAS path), not
+# quantization error.
+
+FP8_E4M3_MAX = 448.0
+FP8_BLOCK_SIZE = 128  # matches FP8_GROUP_SIZE_K in moe_mainloop.hpp
+
+
+def _make_fp8_api_validation_inputs(hidden_dtype=torch.bfloat16):
+    """Create CPU tensors that reach fused_experts' FP8 contract checks."""
+    num_experts = 8
+    hidden_size = intermediate_size = 128
+    return dict(
+        hidden_states=torch.zeros((1, hidden_size), dtype=hidden_dtype),
+        w1=torch.zeros(
+            (num_experts, 2 * intermediate_size, hidden_size),
+            dtype=torch.float8_e4m3fn,
+        ),
+        w2=torch.zeros(
+            (num_experts, hidden_size, intermediate_size),
+            dtype=torch.float8_e4m3fn,
+        ),
+        topk_weights=torch.ones((1, 1), dtype=torch.float32),
+        topk_ids=torch.zeros((1, 1), dtype=torch.long),
+        w1_scale=torch.ones(
+            (num_experts, 2, 1),
+            dtype=torch.float32,
+        ),
+        w2_scale=torch.ones(
+            (num_experts, 1, 1),
+            dtype=torch.float32,
+        ),
+    )
+
+
+def test_fp8_moe_requires_bfloat16_activation():
+    inputs = _make_fp8_api_validation_inputs(hidden_dtype=torch.float16)
+    with pytest.raises(AssertionError, match="hidden_states must be bfloat16"):
+        fused_experts(**inputs, use_fp8_w8a8=True)
+
+
+def test_fp8_moe_rejects_external_activation_scales():
+    inputs = _make_fp8_api_validation_inputs()
+    with pytest.raises(AssertionError, match="prequantized FP8 activation input"):
+        fused_experts(**inputs, use_fp8_w8a8=True, a1_scale=torch.ones(1, 1))
+
+
+def test_fp8_moe_accepts_only_128_by_128_block_metadata():
+    inputs = _make_fp8_api_validation_inputs()
+    with pytest.raises(AssertionError, match=r"block_shape=\[128, 128\]"):
+        fused_experts(**inputs, use_fp8_w8a8=True, block_shape=[64, 128])
+
+
+def test_fp8_moe_rejects_non_block_w8a8_scales():
+    inputs = _make_fp8_api_validation_inputs()
+    inputs["w1_scale"] = torch.ones((8, 256, 1), dtype=torch.float32)
+    with pytest.raises(AssertionError, match="w1_scale block scales"):
+        fused_experts(**inputs, use_fp8_w8a8=True)
+
+
+def _quant_dequant_fp8_block(w: torch.Tensor, block_size: int = FP8_BLOCK_SIZE):
+    """2-D block (e.g. DeepSeek-style 128x128) fp8 e4m3 quantize + dequantize
+    for a 3-D expert weight tensor [E, N, K]. N and K must be multiples of
+    block_size.
+
+    Returns (scale [E, N/block_size, K/block_size] fp32, q_fp8 [E, N, K], dequantized tensor
+    in w's original dtype).
+    """
+    E, N, K = w.shape
+    assert (
+        N % block_size == 0 and K % block_size == 0
+    ), f"N={N} and K={K} must both be multiples of block_size={block_size}"
+    w_f32 = w.float().reshape(
+        E, N // block_size, block_size, K // block_size, block_size
+    )
+    amax = w_f32.abs().amax(dim=(2, 4), keepdim=True).clamp(min=1e-12)
+    scale = amax / FP8_E4M3_MAX
+    q = (w_f32 / scale).clamp(-FP8_E4M3_MAX, FP8_E4M3_MAX).to(torch.float8_e4m3fn)
+    dq = (q.float() * scale).reshape(E, N, K).to(w.dtype)
+    return scale.reshape(E, N // block_size, K // block_size), q.reshape(E, N, K), dq
+
+
+def torch_naive_moe_fp8_w8a16(
+    a,
+    w1_dq,
+    w2_dq,
+    topk_ids,
+    topk_weight,
+    topk,
+    b1,
+    b2,
+    routed_scaling_factor=None,
+    activation="silu",
+    gemm1_alpha=None,
+    gemm1_limit=None,
+    swiglu_limit=None,
+):
+    """Reference for the Xe2 FP8-weight fallback path.
+
+    ``use_fp8_w8a8=True`` identifies the checkpoint's FP8 weight format, but
+    Xe2 executes the MoE with BF16 activations through the W8A16 operator.
+    ``w1_dq`` and ``w2_dq`` must already be the dequantized FP8-rounded
+    weights so the reference includes weight quantization noise.
+    """
+    B, D = a.shape
+    a_rep = a.view(B, -1, D).repeat(1, topk, 1).reshape(-1, D)
+
+    out = torch.zeros(B * topk, w2_dq.shape[1], dtype=a.dtype, device=a.device)
+    topk_weight_flat = topk_weight.reshape(-1)
+    topk_ids_flat = topk_ids.reshape(-1)
+    b1 = (
+        b1
+        if b1 is not None
+        else torch.zeros(w1_dq.shape[:2], dtype=torch.float32, device=a.device)
+    )
+    b2 = (
+        b2
+        if b2 is not None
+        else torch.zeros(w2_dq.shape[:2], dtype=torch.float32, device=a.device)
+    )
+
+    for i in range(w1_dq.shape[0]):
+        mask = topk_ids_flat == i
+        if mask.sum():
+            gemm1 = (a_rep[mask].float() @ w1_dq[i].float().transpose(0, 1)) + b1[
+                i
+            ].float()
+            if gemm1_alpha is not None:
+                tmp = swiglu_gpt_oss_sigmoid_alpha(
+                    gemm1.to(a.dtype), gemm1_alpha, gemm1_limit
+                )
+            elif swiglu_limit is not None:
+                gate, up = gemm1.to(a.dtype).chunk(2, dim=-1)
+                gate = torch.minimum(gate, torch.tensor(swiglu_limit, dtype=gate.dtype))
+                up = torch.clamp(up, -swiglu_limit, swiglu_limit)
+                tmp = (F.silu(gate.float()) * up.float()).to(a.dtype)
+            elif activation == "silu":
+                act_fn = F.silu
+            elif activation == "relu2":
+                tmp = torch.square(torch.relu(gemm1.to(a.dtype)))
+            elif activation == "gelu":
+                act_fn = lambda x: F.gelu(x, approximate="tanh")
+            else:
+                raise AssertionError(f"unsupported test activation: {activation}")
+            if activation != "relu2" and gemm1_alpha is None and swiglu_limit is None:
+                tmp = apply_act_and_mul(gemm1.to(a.dtype), act_fn)
+            gemm2 = (tmp.float() @ w2_dq[i].float().transpose(0, 1)) + b2[i].float()
+            out[mask] = gemm2.to(a.dtype)
+
+    result = (
+        out.view(B, topk, w2_dq.shape[1]) * topk_weight.view(B, topk, 1).to(out.dtype)
+    ).sum(dim=1)
+    if routed_scaling_factor is not None:
+        result = result * routed_scaling_factor
+    return result
+
+
+# ---------------------------------------------------------------------------
+# FP8 W8A16 expert-weight tests (fused_experts level)
+# ---------------------------------------------------------------------------
+
+
+@pytest.mark.parametrize(
+    "num_tokens,topk,num_experts,hidden_size,intermediate_size,with_bias",
+    [
+        (1, 1, 8, 256, 256, False),  # avg_m <= 4: Tile16
+        (32, 2, 8, 256, 256, True),  # avg_m <= 128: Tile32, with bias
+        (1024, 2, 8, 512, 512, False),  # avg_m > 128: Tile128
+    ],
+)
+def test_moe_gemm_fp8_w8a16_block_weights(
+    num_tokens,
+    topk,
+    num_experts,
+    hidden_size,
+    intermediate_size,
+    with_bias,
+):
+    """Test fused_experts with FP8 E4M3 weights and 128x128 block scales.
+
+    Weights are block-quantized (128x128, DeepSeek-style) then dequantized
+    for the reference so both paths see identical fp8-rounded weights; the
+    reference (torch_naive_moe_fp8_w8a16) dequantizes the FP8-rounded weights
+    and keeps activations in BF16, matching the Xe2 W8A16 fallback.
+    Any remaining numerical difference should be GEMM arithmetic noise only.
+
+    The cases cover each block-scale tile tier and runtime-optional bias.
+    """
+    from sgl_kernel.moe import _moe_ws_cache, _moe_ws_view_cache
+
+    _moe_ws_cache.clear()
+    _moe_ws_view_cache.clear()
+    torch.xpu.synchronize()
+    torch.xpu.empty_cache()
+    torch.manual_seed(0)
+    torch.xpu.manual_seed_all(0)
+
+    rtol, atol = 1e-1, 5e-2
+
+    a = create_random_cpu_tensor((num_tokens, hidden_size), torch.bfloat16)
+    # w1: gate+up projection [E, 2*I, H]; w2: down projection [E, H, I].
+    w1_bf16 = create_random_cpu_tensor(
+        (num_experts, 2 * intermediate_size, hidden_size), torch.bfloat16
+    )
+    w2_bf16 = create_random_cpu_tensor(
+        (num_experts, hidden_size, intermediate_size), torch.bfloat16
+    )
+
+    b1, b2 = None, None
+    if with_bias:
+        b1 = create_random_cpu_tensor(
+            (num_experts, 2 * intermediate_size), torch.float32, std=0.005
+        )
+        b2 = create_random_cpu_tensor(
+            (num_experts, hidden_size), torch.float32, std=0.005
+        )
+
+    score = torch.randn([num_tokens, num_experts], dtype=torch.bfloat16)
+    score = torch.softmax(score, dim=-1, dtype=torch.float32)
+    topk_weight, topk_ids = torch.topk(score, topk)
+
+    w1_scale, w1_fp8, w1_dq = _quant_dequant_fp8_block(w1_bf16)
+    w2_scale, w2_fp8, w2_dq = _quant_dequant_fp8_block(w2_bf16)
+
+    torch_output = torch_naive_moe_fp8_w8a16(
+        a,
+        w1_dq,
+        w2_dq,
+        topk_ids,
+        topk_weight,
+        topk,
+        b1,
+        b2,
+    )
+
+    device = "xpu"
+    sglang_output = fused_experts(
+        a.to(device),
+        w1_fp8.to(device),
+        w2_fp8.to(device),
+        topk_weight.to(device),
+        topk_ids.to(device),
+        b1.to(device) if b1 is not None else None,
+        b2.to(device) if b2 is not None else None,
+        activation="silu",
+        use_fp8_w8a8=True,
+        w1_scale=w1_scale.to(device),
+        w2_scale=w2_scale.to(device),
+    )
+
+    torch.xpu.synchronize()
+    torch.testing.assert_close(
+        torch_output, sglang_output.to("cpu"), rtol=rtol, atol=atol
+    )
+    _moe_ws_cache.clear()
+    _moe_ws_view_cache.clear()
+    torch.xpu.synchronize()
+    torch.xpu.empty_cache()
+
+
+def test_moe_gemm_fp8_gelu_split_activation():
+    """Keep one small GELU case covering FP8's external activation wiring."""
+    torch.manual_seed(1)
+    torch.xpu.manual_seed_all(1)
+    num_tokens, topk, num_experts = 17, 2, 8
+    hidden_size = intermediate_size = 256
+
+    a = create_random_cpu_tensor((num_tokens, hidden_size), torch.bfloat16)
+    w1_bf16 = create_random_cpu_tensor(
+        (num_experts, 2 * intermediate_size, hidden_size), torch.bfloat16
+    )
+    w2_bf16 = create_random_cpu_tensor(
+        (num_experts, hidden_size, intermediate_size), torch.bfloat16
+    )
+    score = torch.randn((num_tokens, num_experts), dtype=torch.bfloat16)
+    topk_weight, topk_ids = torch.topk(
+        torch.softmax(score, dim=-1, dtype=torch.float32), topk
+    )
+    w1_scale, w1_fp8, w1_dq = _quant_dequant_fp8_block(w1_bf16)
+    w2_scale, w2_fp8, w2_dq = _quant_dequant_fp8_block(w2_bf16)
+
+    torch_output = torch_naive_moe_fp8_w8a16(
+        a,
+        w1_dq,
+        w2_dq,
+        topk_ids,
+        topk_weight,
+        topk,
+        None,
+        None,
+        activation="gelu",
+    )
+    sglang_output = fused_experts(
+        a.to("xpu"),
+        w1_fp8.to("xpu"),
+        w2_fp8.to("xpu"),
+        topk_weight.to("xpu"),
+        topk_ids.to("xpu"),
+        activation="gelu",
+        w1_scale=w1_scale.to("xpu"),
+        w2_scale=w2_scale.to("xpu"),
+        use_fp8_w8a8=True,
+    )
+    torch.testing.assert_close(
+        torch_output, sglang_output.to("cpu"), rtol=1e-1, atol=1e-2
+    )
+
+
+def test_moe_gemm_fp8_relu2():
+    """Cover FP8's non-gated ReLU2 path with a single-width GEMM1 output."""
+    torch.manual_seed(2)
+    torch.xpu.manual_seed_all(2)
+    num_tokens, topk, num_experts = 17, 2, 8
+    hidden_size = intermediate_size = 256
+
+    a = create_random_cpu_tensor((num_tokens, hidden_size), torch.bfloat16)
+    w1_bf16 = create_random_cpu_tensor(
+        (num_experts, intermediate_size, hidden_size), torch.bfloat16
+    )
+    w2_bf16 = create_random_cpu_tensor(
+        (num_experts, hidden_size, intermediate_size), torch.bfloat16
+    )
+    b1 = create_random_cpu_tensor(
+        (num_experts, intermediate_size), torch.float32, std=0.005
+    )
+    b2 = create_random_cpu_tensor((num_experts, hidden_size), torch.float32, std=0.005)
+    score = torch.randn((num_tokens, num_experts), dtype=torch.bfloat16)
+    topk_weight, topk_ids = torch.topk(
+        torch.softmax(score, dim=-1, dtype=torch.float32), topk
+    )
+    w1_scale, w1_fp8, w1_dq = _quant_dequant_fp8_block(w1_bf16)
+    w2_scale, w2_fp8, w2_dq = _quant_dequant_fp8_block(w2_bf16)
+
+    torch_output = torch_naive_moe_fp8_w8a16(
+        a,
+        w1_dq,
+        w2_dq,
+        topk_ids,
+        topk_weight,
+        topk,
+        b1,
+        b2,
+        activation="relu2",
+    )
+    sglang_output = fused_experts(
+        a.to("xpu"),
+        w1_fp8.to("xpu"),
+        w2_fp8.to("xpu"),
+        topk_weight.to("xpu"),
+        topk_ids.to("xpu"),
+        b1.to("xpu"),
+        b2.to("xpu"),
+        activation="relu2",
+        w1_scale=w1_scale.to("xpu"),
+        w2_scale=w2_scale.to("xpu"),
+        use_fp8_w8a8=True,
+    )
+    torch.testing.assert_close(
+        torch_output, sglang_output.to("cpu"), rtol=1e-1, atol=1e-2
+    )
+
+
+@pytest.mark.parametrize("activation_variant", ["gpt_oss", "deepseek_v4"])
+def test_moe_gemm_fp8_swiglu_variants(activation_variant):
+    torch.manual_seed(3)
+    torch.xpu.manual_seed_all(3)
+    num_tokens, topk, num_experts = 17, 2, 8
+    hidden_size = intermediate_size = 256
+    a = create_random_cpu_tensor((num_tokens, hidden_size), torch.bfloat16)
+    w1_bf16 = create_random_cpu_tensor(
+        (num_experts, 2 * intermediate_size, hidden_size), torch.bfloat16
+    )
+    w2_bf16 = create_random_cpu_tensor(
+        (num_experts, hidden_size, intermediate_size), torch.bfloat16
+    )
+    score = torch.randn((num_tokens, num_experts), dtype=torch.bfloat16)
+    topk_weight, topk_ids = torch.topk(
+        torch.softmax(score, dim=-1, dtype=torch.float32), topk
+    )
+    w1_scale, w1_fp8, w1_dq = _quant_dequant_fp8_block(w1_bf16)
+    w2_scale, w2_fp8, w2_dq = _quant_dequant_fp8_block(w2_bf16)
+    activation_kwargs = (
+        {"gemm1_alpha": SWIGLU_ALPHA, "gemm1_limit": SWIGLU_LIMIT}
+        if activation_variant == "gpt_oss"
+        else {"swiglu_limit": 10}
+    )
+    torch_output = torch_naive_moe_fp8_w8a16(
+        a,
+        w1_dq,
+        w2_dq,
+        topk_ids,
+        topk_weight,
+        topk,
+        None,
+        None,
+        **activation_kwargs,
+    )
+    sglang_output = fused_experts(
+        a.to("xpu"),
+        w1_fp8.to("xpu"),
+        w2_fp8.to("xpu"),
+        topk_weight.to("xpu"),
+        topk_ids.to("xpu"),
+        activation="silu",
+        w1_scale=w1_scale.to("xpu"),
+        w2_scale=w2_scale.to("xpu"),
+        use_fp8_w8a8=True,
+        **activation_kwargs,
+    )
+    torch.testing.assert_close(
+        torch_output, sglang_output.to("cpu"), rtol=1e-1, atol=1e-2
+    )
+
+
+def test_moe_gemm_fp8_w8a8_flag_falls_back_for_scalar_scales():
+    """The W8A8 checkpoint flag must use W8A16 for scalar FP8 scales on Xe2."""
+    torch.manual_seed(4)
+    torch.xpu.manual_seed_all(4)
+    num_tokens, topk, num_experts = 17, 2, 8
+    hidden_size = intermediate_size = 128
+
+    a = create_random_cpu_tensor((num_tokens, hidden_size), torch.bfloat16)
+    w1_bf16 = create_random_cpu_tensor(
+        (num_experts, 2 * intermediate_size, hidden_size), torch.bfloat16
+    )
+    w2_bf16 = create_random_cpu_tensor(
+        (num_experts, hidden_size, intermediate_size), torch.bfloat16
+    )
+    score = torch.randn((num_tokens, num_experts), dtype=torch.bfloat16)
+    topk_weight, topk_ids = torch.topk(
+        torch.softmax(score, dim=-1, dtype=torch.float32), topk
+    )
+
+    w1_scale = (
+        w1_bf16.float().abs().amax((1, 2), keepdim=True).clamp_min(1e-12) / FP8_E4M3_MAX
+    )
+    w2_scale = (
+        w2_bf16.float().abs().amax((1, 2), keepdim=True).clamp_min(1e-12) / FP8_E4M3_MAX
+    )
+    w1_fp8 = (w1_bf16.float() / w1_scale).to(torch.float8_e4m3fn)
+    w2_fp8 = (w2_bf16.float() / w2_scale).to(torch.float8_e4m3fn)
+    w1_dq = (w1_fp8.float() * w1_scale).to(torch.bfloat16)
+    w2_dq = (w2_fp8.float() * w2_scale).to(torch.bfloat16)
+    w1_scale = w1_scale.view(num_experts, 1).repeat(1, 2).contiguous()
+    w2_scale = w2_scale.view(num_experts, 1).contiguous()
+
+    torch_output = torch_naive_moe(
+        a, w1_dq, w2_dq, topk_ids, topk_weight, topk, None, None
+    )
+    sglang_output = fused_experts(
+        a.to("xpu"),
+        w1_fp8.to("xpu"),
+        w2_fp8.to("xpu"),
+        topk_weight.to("xpu"),
+        topk_ids.to("xpu"),
+        activation="silu",
+        use_fp8_w8a8=True,
+        w1_scale=w1_scale.to("xpu"),
+        w2_scale=w2_scale.to("xpu"),
+    )
+    torch.testing.assert_close(
+        torch_output, sglang_output.to("cpu"), rtol=1e-1, atol=5e-2
+    )
+
+
+def test_moe_gemm_fp8_w8a16_distinct_w1_scalar_scales():
+    torch.manual_seed(5)
+    torch.xpu.manual_seed_all(5)
+    num_tokens, topk, num_experts = 17, 2, 8
+    hidden_size = intermediate_size = 128
+
+    a = create_random_cpu_tensor((num_tokens, hidden_size), torch.bfloat16)
+    w1_gate = create_random_cpu_tensor(
+        (num_experts, intermediate_size, hidden_size), torch.bfloat16
+    )
+    w1_up = (
+        create_random_cpu_tensor(
+            (num_experts, intermediate_size, hidden_size), torch.bfloat16
+        )
+        * 0.125
+    )
+    w2_bf16 = create_random_cpu_tensor(
+        (num_experts, hidden_size, intermediate_size), torch.bfloat16
+    )
+    score = torch.randn((num_tokens, num_experts), dtype=torch.bfloat16)
+    topk_weight, topk_ids = torch.topk(
+        torch.softmax(score, dim=-1, dtype=torch.float32), topk
+    )
+
+    gate_scale = (
+        w1_gate.float().abs().amax((1, 2), keepdim=True).clamp_min(1e-12) / FP8_E4M3_MAX
+    )
+    up_scale = (
+        w1_up.float().abs().amax((1, 2), keepdim=True).clamp_min(1e-12) / FP8_E4M3_MAX
+    )
+    w2_scale = (
+        w2_bf16.float().abs().amax((1, 2), keepdim=True).clamp_min(1e-12) / FP8_E4M3_MAX
+    )
+    gate_fp8 = (w1_gate.float() / gate_scale).to(torch.float8_e4m3fn)
+    up_fp8 = (w1_up.float() / up_scale).to(torch.float8_e4m3fn)
+    w2_fp8 = (w2_bf16.float() / w2_scale).to(torch.float8_e4m3fn)
+    w1_fp8 = torch.cat((gate_fp8, up_fp8), dim=1).contiguous()
+    w1_dq = torch.cat(
+        (
+            (gate_fp8.float() * gate_scale).to(torch.bfloat16),
+            (up_fp8.float() * up_scale).to(torch.bfloat16),
+        ),
+        dim=1,
+    )
+    w2_dq = (w2_fp8.float() * w2_scale).to(torch.bfloat16)
+    w1_scale = torch.cat((gate_scale, up_scale), dim=2).view(num_experts, 2)
+    w2_scale = w2_scale.view(num_experts, 1)
+
+    torch_output = torch_naive_moe(
+        a, w1_dq, w2_dq, topk_ids, topk_weight, topk, None, None
+    )
+    sglang_output = fused_experts(
+        a.to("xpu"),
+        w1_fp8.to("xpu"),
+        w2_fp8.to("xpu"),
+        topk_weight.to("xpu"),
+        topk_ids.to("xpu"),
+        activation="silu",
+        use_fp8_w8a8=True,
+        w1_scale=w1_scale.to("xpu"),
+        w2_scale=w2_scale.to("xpu"),
+    )
+    torch.testing.assert_close(
+        torch_output, sglang_output.to("cpu"), rtol=1e-1, atol=5e-2
+    )
+
+
+@pytest.mark.parametrize(
+    ("rows_per_expert", "intermediate_size"), [(2, 128), (16, 128), (40, 192)]
+)
+def test_moe_grouped_mm_fp8_w8a16_distinct_two_scales(
+    rows_per_expert, intermediate_size
+):
+    torch.manual_seed(6)
+    torch.xpu.manual_seed_all(6)
+    num_experts, hidden_size = 8, 128
+    total_rows = num_experts * rows_per_expert
+
+    activations = torch.randn((total_rows, hidden_size), dtype=torch.bfloat16)
+    gate = (
+        torch.randn((num_experts, intermediate_size, hidden_size), dtype=torch.bfloat16)
+        * 0.25
+    )
+    up = (
+        torch.randn((num_experts, intermediate_size, hidden_size), dtype=torch.bfloat16)
+        * 0.03125
+    )
+    gate_scale = (
+        gate.float().abs().amax((1, 2), keepdim=True).clamp_min(1e-12) / FP8_E4M3_MAX
+    )
+    up_scale = (
+        up.float().abs().amax((1, 2), keepdim=True).clamp_min(1e-12) / FP8_E4M3_MAX
+    )
+    gate_fp8 = (gate.float() / gate_scale).to(torch.float8_e4m3fn)
+    up_fp8 = (up.float() / up_scale).to(torch.float8_e4m3fn)
+    weights = torch.cat((gate_fp8, up_fp8), dim=1).contiguous()
+    scales = torch.cat((gate_scale, up_scale), dim=2).view(num_experts, 2)
+    dequantized_weights = torch.cat(
+        (
+            gate_fp8.float() * gate_scale,
+            up_fp8.float() * up_scale,
+        ),
+        dim=1,
+    )
+    reference = torch.cat(
+        [
+            activations[
+                expert * rows_per_expert : (expert + 1) * rows_per_expert
+            ].float()
+            @ dequantized_weights[expert].transpose(0, 1)
+            for expert in range(num_experts)
+        ],
+        dim=0,
+    ).to(torch.bfloat16)
+
+    output = torch.empty(
+        (total_rows, 2 * intermediate_size), device="xpu", dtype=torch.bfloat16
+    )
+    torch.ops.sgl_kernel.moe_grouped_mm_nt_xe20_fp8_w8a16(
+        output,
+        activations.to("xpu"),
+        weights.to("xpu"),
+        scales.to("xpu"),
+        None,
+        torch.full((num_experts,), rows_per_expert, device="xpu", dtype=torch.int32),
+        num_experts,
+    )
+    torch.testing.assert_close(reference, output.cpu(), rtol=5e-2, atol=5e-2)
+
+
+@pytest.mark.parametrize(
+    "num_tokens,topk,num_experts,hidden_size,intermediate_size",
+    [
+        (1, 2, 8, 128, 128),  # avg_m <= 8: Tile16
+        (128, 2, 8, 256, 256),  # avg_m <= 32: Tile32
+        (256, 2, 8, 2048, 256),  # GEMM1 medium-M long-K: Tile64
+        (512, 2, 8, 256, 256),  # avg_m <= 128, short-K: Tile128
+        (1024, 2, 8, 256, 256),  # avg_m > 128: Tile128
+    ],
+)
+def test_moe_gemm_fp8_w8a16(
+    num_tokens, topk, num_experts, hidden_size, intermediate_size
+):
+    torch.manual_seed(0)
+    torch.xpu.manual_seed_all(0)
+
+    a = create_random_cpu_tensor((num_tokens, hidden_size), torch.bfloat16)
+    w1_bf16 = create_random_cpu_tensor(
+        (num_experts, 2 * intermediate_size, hidden_size), torch.bfloat16
+    )
+    w2_bf16 = create_random_cpu_tensor(
+        (num_experts, hidden_size, intermediate_size), torch.bfloat16
+    )
+    score = torch.randn((num_tokens, num_experts), dtype=torch.bfloat16)
+    topk_weight, topk_ids = torch.topk(
+        torch.softmax(score, dim=-1, dtype=torch.float32), topk
+    )
+    if num_tokens == 1:
+        topk_ids = torch.cat((topk_ids, topk_ids), dim=1)[:, ::2]
+        assert not topk_ids.is_contiguous()
+
+    s1 = (
+        w1_bf16.float().abs().amax((1, 2), keepdim=True).clamp_min(1e-12) / FP8_E4M3_MAX
+    )
+    s2 = (
+        w2_bf16.float().abs().amax((1, 2), keepdim=True).clamp_min(1e-12) / FP8_E4M3_MAX
+    )
+    w1_fp8 = (w1_bf16.float() / s1).to(torch.float8_e4m3fn)
+    w2_fp8 = (w2_bf16.float() / s2).to(torch.float8_e4m3fn)
+    w1_dq = (w1_fp8.float() * s1).to(torch.bfloat16)
+    w2_dq = (w2_fp8.float() * s2).to(torch.bfloat16)
+
+    w1_scale = s1.view(num_experts, 1).repeat(1, 2).contiguous()
+    w2_scale = s2.view(num_experts, 1).contiguous()
+
+    torch_output = torch_naive_moe(
+        a,
+        w1_dq,
+        w2_dq,
+        topk_ids,
+        topk_weight,
+        topk,
+        None,
+        None,
+    )
+    sglang_output = fused_experts(
+        a.to("xpu"),
+        w1_fp8.to("xpu"),
+        w2_fp8.to("xpu"),
+        topk_weight.to("xpu"),
+        topk_ids.to("xpu"),
+        activation="silu",
+        use_fp8_w8a8=True,
+        w1_scale=w1_scale.to("xpu"),
+        w2_scale=w2_scale.to("xpu"),
+    )
+    torch.testing.assert_close(
+        torch_output, sglang_output.to("cpu"), rtol=1e-1, atol=5e-2
+    )
+
+
+@pytest.mark.parametrize("num_tokens", [1, 16, 128, 512])
+def test_moe_gemm_fp8_w8a16_block_scales(num_tokens):
+    """Cover FP8 W8A16 with 128x128 weight block scales."""
+    torch.manual_seed(3)
+    torch.xpu.manual_seed_all(3)
+    num_experts, topk = 8, 2
+    hidden_size = intermediate_size = 256
+
+    a = create_random_cpu_tensor((num_tokens, hidden_size), torch.bfloat16)
+    w1_bf16 = create_random_cpu_tensor(
+        (num_experts, 2 * intermediate_size, hidden_size), torch.bfloat16
+    )
+    w2_bf16 = create_random_cpu_tensor(
+        (num_experts, hidden_size, intermediate_size), torch.bfloat16
+    )
+    score = torch.randn((num_tokens, num_experts), dtype=torch.bfloat16)
+    topk_weight, topk_ids = torch.topk(
+        torch.softmax(score, dim=-1, dtype=torch.float32), topk
+    )
+
+    w1_scale, w1_fp8, w1_dq = _quant_dequant_fp8_block(w1_bf16)
+    w2_scale, w2_fp8, w2_dq = _quant_dequant_fp8_block(w2_bf16)
+    torch_output = torch_naive_moe(
+        a, w1_dq, w2_dq, topk_ids, topk_weight, topk, None, None
+    )
+    sglang_output = fused_experts(
+        a.to("xpu"),
+        w1_fp8.to("xpu"),
+        w2_fp8.to("xpu"),
+        topk_weight.to("xpu"),
+        topk_ids.to("xpu"),
+        activation="silu",
+        use_fp8_w8a8=True,
+        w1_scale=w1_scale.to("xpu"),
+        w2_scale=w2_scale.to("xpu"),
+    )
+    torch.testing.assert_close(
+        torch_output, sglang_output.to("cpu"), rtol=1e-1, atol=5e-2
+    )
+
+
+@pytest.mark.parametrize(
+    "rows_per_expert", [(0, 1, 2, 3, 4, 5, 7, 10), (8, 9, 11, 13, 16, 19, 23, 29)]
+)
+def test_moe_grouped_mm_fp8_w8a16_heterogeneous_block_scales(rows_per_expert):
+    torch.manual_seed(7)
+    torch.xpu.manual_seed_all(7)
+    num_experts = 8
+    gemm_n = gemm_k = 256
+    total_rows = sum(rows_per_expert)
+
+    activations = torch.randn((total_rows, gemm_k), dtype=torch.bfloat16)
+    block_values = torch.randn(
+        (num_experts, 2, FP8_BLOCK_SIZE, 2, FP8_BLOCK_SIZE),
+        dtype=torch.float32,
+    )
+    expert = torch.arange(num_experts, dtype=torch.int32).view(-1, 1, 1)
+    n_block = torch.arange(2, dtype=torch.int32).view(1, -1, 1)
+    k_block = torch.arange(2, dtype=torch.int32).view(1, 1, -1)
+    exponents = ((expert * 3 + n_block * 2 + k_block) % 9 - 4).float()
+    block_values *= torch.pow(2.0, exponents).view(num_experts, 2, 1, 2, 1)
+    weights_bf16 = block_values.reshape(num_experts, gemm_n, gemm_k).to(torch.bfloat16)
+    scales, weights_fp8, weights_dq = _quant_dequant_fp8_block(weights_bf16)
+
+    reference_parts = []
+    row_start = 0
+    for expert_id, rows in enumerate(rows_per_expert):
+        row_end = row_start + rows
+        if rows:
+            reference_parts.append(
+                activations[row_start:row_end].float()
+                @ weights_dq[expert_id].float().transpose(0, 1)
+            )
+        row_start = row_end
+    reference = torch.cat(reference_parts, dim=0).to(torch.bfloat16)
+
+    output = torch.empty((total_rows, gemm_n), device="xpu", dtype=torch.bfloat16)
+    torch.ops.sgl_kernel.moe_grouped_mm_nt_xe20_fp8_w8a16(
+        output,
+        activations.to("xpu"),
+        weights_fp8.to("xpu"),
+        scales.to("xpu"),
+        None,
+        torch.tensor(rows_per_expert, device="xpu", dtype=torch.int32),
+        num_experts,
+    )
+    actual = output.cpu().float()
+    expected = reference.float()
+    error = (actual - expected).abs()
+    assert error.mean() <= expected.abs().mean() * 1e-2
+    assert error.max() <= expected.abs().max() * 2e-2
+
+
 if __name__ == "__main__":
     sys.exit(pytest.main([__file__]))
diff --git a/tests/test_moe_prepare_input.py b/tests/test_moe_prepare_input.py
index 0acfcd9d..3ea1680c 100644
--- a/tests/test_moe_prepare_input.py
+++ b/tests/test_moe_prepare_input.py
@@ -8,6 +8,102 @@
     prepare_moe_input,
     scatter_tokens_to_experts,
 )
+from sgl_kernel.moe import _should_use_small_moe_prepare
+
+
+@pytest.mark.parametrize(
+    "num_tokens,top_k,hidden_dim,num_experts,expected",
+    [
+        (1, 0, 2048, 64, False),
+        (10, 4, 2048, 64, True),
+        (1, 16, 2048, 64, True),
+        (1, 17, 2048, 64, False),
+        (11, 4, 2048, 64, False),
+        (16, 4, 1024, 64, True),
+        (17, 4, 1024, 128, False),
+        (8, 8, 1024, 32, False),
+    ],
+)
+def test_should_use_small_moe_prepare(
+    num_tokens, top_k, hidden_dim, num_experts, expected
+):
+    assert (
+        _should_use_small_moe_prepare(num_tokens, top_k, hidden_dim, num_experts)
+        is expected
+    )
+
+
+@pytest.mark.parametrize(
+    "num_tokens,top_k,hidden_dim,index_dtype,routing",
+    [
+        (1, 16, 15, torch.int32, "random"),
+        (1, 8, 16, torch.int64, "all_same"),
+        (3, 3, 33, torch.int64, "random"),
+        (5, 7, 16, torch.int32, "skewed"),
+        (8, 8, 33, torch.int32, "all_same"),
+    ],
+)
+def test_prepare_moe_input_small(num_tokens, top_k, hidden_dim, index_dtype, routing):
+    num_experts = 32
+    torch.manual_seed(41)
+    input_tensor = torch.randn(
+        num_tokens, hidden_dim, dtype=torch.bfloat16, device="xpu"
+    )
+    if routing == "random":
+        topk_ids = torch.stack(
+            [torch.randperm(num_experts)[:top_k] for _ in range(num_tokens)]
+        )
+    elif routing == "skewed":
+        topk_ids = torch.arange(top_k).repeat(num_tokens, 1)
+    else:
+        topk_ids = torch.zeros((num_tokens, top_k), dtype=torch.int64)
+    topk_ids = topk_ids.to(device="xpu", dtype=index_dtype)
+    expert_counts = torch.full((num_experts,), 123, dtype=torch.int32, device="xpu")
+    output_permutation = torch.empty(
+        num_tokens * top_k, dtype=torch.int32, device="xpu"
+    )
+    output = torch.empty(
+        num_tokens * top_k, hidden_dim, dtype=torch.bfloat16, device="xpu"
+    )
+
+    torch.ops.sgl_kernel.prepare_moe_input_small.default(
+        input_tensor, topk_ids, expert_counts, output_permutation, output
+    )
+
+    flat_ids = topk_ids.cpu().flatten()
+    sorted_routes = torch.argsort(flat_ids, stable=True)
+    expected_permutation = torch.empty_like(sorted_routes, dtype=torch.int32)
+    expected_permutation[sorted_routes] = torch.arange(
+        sorted_routes.numel(), dtype=torch.int32
+    )
+    expected_output = input_tensor.cpu()[sorted_routes // top_k]
+    expected_counts = torch.bincount(flat_ids, minlength=num_experts).to(torch.int32)
+
+    torch.testing.assert_close(expert_counts.cpu(), expected_counts)
+    torch.testing.assert_close(output_permutation.cpu(), expected_permutation)
+    torch.testing.assert_close(output.cpu(), expected_output)
+
+
+@pytest.mark.parametrize(
+    "num_tokens,top_k,error",
+    [(1, 17, "topk must be in"), (13, 5, "routed rows must be")],
+)
+def test_prepare_moe_input_small_rejects_capacity_overflow(num_tokens, top_k, error):
+    hidden_dim = 16
+    input_tensor = torch.empty(
+        (num_tokens, hidden_dim), dtype=torch.bfloat16, device="xpu"
+    )
+    topk_ids = torch.zeros((num_tokens, top_k), dtype=torch.int32, device="xpu")
+    route_count = topk_ids.numel()
+
+    with pytest.raises(RuntimeError, match=error):
+        torch.ops.sgl_kernel.prepare_moe_input_small.default(
+            input_tensor,
+            topk_ids,
+            torch.empty(32, dtype=torch.int32, device="xpu"),
+            torch.empty(route_count, dtype=torch.int32, device="xpu"),
+            torch.empty((route_count, hidden_dim), dtype=torch.bfloat16, device="xpu"),
+        )
 
 
 @pytest.mark.parametrize("num_tokens", [1, 2, 5, 16, 64, 128, 224, 1024])