[Model] Support Qwen3.8-Flash-Next

.buildkite/test_areas/lm_eval.yaml

@@ -331,6 +331,42 @@ steps:
     - uv pip install --system 'gpt-oss[eval]==0.0.5'
     - pytest -s -v evals/gpt_oss/test_gpqa_correctness.py --config-list-file=configs/models-b200.txt
 
+- label: ":nvidia: (B200) Qwen3.8-Flash-Next-FP8 Accuracy Eval"
+  key: accuracy-eval-qwen3-8-flash-next-fp8-b200
+  timeout_in_minutes: 60
+  device: b200-k8s
+  optional: true
+  num_devices: 4
+  source_file_dependencies:
+  - csrc/libtorch_stable/gdn/
+  - tests/evals/qwen4_exp/
+  - vllm/model_executor/layers/mamba/
+  - vllm/models/qwen4_exp/
+  - vllm/transformers_utils/configs/qwen4_exp.py
+  - vllm/v1/attention/backends/short_conv_attn.py
+  - vllm/v1/spec_decode/qwen4_exp.py
+  commands:
+    - uv pip install --system 'evalscope==1.10.0'
+    - pytest -s -v evals/qwen4_exp/test_accuracy.py --config-list-file=configs/models-b200.txt
+
+- label: ":nvidia: (H200) Qwen3.8-Flash-Next-FP8 Accuracy Eval"
+  key: accuracy-eval-qwen3-8-flash-next-fp8-h200
+  timeout_in_minutes: 60
+  device: h200
+  optional: true
+  num_devices: 4
+  source_file_dependencies:
+  - csrc/libtorch_stable/gdn/
+  - tests/evals/qwen4_exp/
+  - vllm/model_executor/layers/mamba/
+  - vllm/models/qwen4_exp/
+  - vllm/transformers_utils/configs/qwen4_exp.py
+  - vllm/v1/attention/backends/short_conv_attn.py
+  - vllm/v1/spec_decode/qwen4_exp.py
+  commands:
+    - uv pip install --system 'evalscope==1.10.0'
+    - pytest -s -v evals/qwen4_exp/test_accuracy.py --config-list-file=configs/models-h200.txt
+
 - label: ":nvidia: (DGX) Spark GPQA Eval (GPT-OSS)"
   key: gpqa-eval-gpt-oss-spark
   timeout_in_minutes: 35

csrc/libtorch_stable/gdn/fused_gdn_decode_kernel.cu

@@ -4,6 +4,7 @@
  */
 
 #include <cstdint>
+#include <string>
 #include <cuda_bf16.h>
 #include <cuda_fp16.h>
 #include <cuda_runtime.h>
@@ -146,7 +147,7 @@ __device__ __forceinline__ Sum2 warp_reduce_sum_pair(float x, float y) {
   return {x, y};
 }
 
-template <typename StateT, int ValueHeadsPerKeyHead>
+template <typename StateT, int ValueHeadsPerKeyHead, bool SigmoidGate>
 __global__ __launch_bounds__(kThreads, 2) void gdn_decode_post_conv_mtp_kernel(
     const __nv_bfloat16* __restrict__ mixed_qkv,
     const __nv_bfloat16* __restrict__ a, const __nv_bfloat16* __restrict__ b,
@@ -354,9 +355,11 @@ __global__ __launch_bounds__(kThreads, 2) void gdn_decode_post_conv_mtp_kernel(
 #pragma unroll
     for (int i = 0; i < 4; ++i) {
       const int value = lane + i * 32;
-      const float gate = silu_fast(__bfloat162float(
+      const float gate_input = __bfloat162float(
           output_gate[static_cast<int64_t>(token) * strides.gate_row +
-                      value_head * kDimV + value]));
+                      value_head * kDimV + value]);
+      const float gate =
+          SigmoidGate ? sigmoid_fast(gate_input) : silu_fast(gate_input);
       const float weight =
           norm_weight_is_bf16
               ? __bfloat162float(
@@ -370,7 +373,7 @@ __global__ __launch_bounds__(kThreads, 2) void gdn_decode_post_conv_mtp_kernel(
   }
 }
 
-template <typename StateT, int ValueHeadsPerKeyHead>
+template <typename StateT, int ValueHeadsPerKeyHead, bool SigmoidGate>
 void launch_gdn_decode_post_conv_mtp(
     torch::stable::Tensor const& mixed_qkv, torch::stable::Tensor const& a_log,
     torch::stable::Tensor const& dt_bias,
@@ -395,7 +398,7 @@ void launch_gdn_decode_post_conv_mtp(
       get_current_cuda_stream(mixed_qkv.get_device_index());
   const int num_requests = static_cast<int>(state_indices.size(0));
   const dim3 grid(num_requests, num_value_heads);
-  gdn_decode_post_conv_mtp_kernel<StateT, ValueHeadsPerKeyHead>
+  gdn_decode_post_conv_mtp_kernel<StateT, ValueHeadsPerKeyHead, SigmoidGate>
       <<<grid, kThreads, 0, stream>>>(
           static_cast<const __nv_bfloat16*>(mixed_qkv.data_ptr()), a, b,
           static_cast<const float*>(a_log.data_ptr()), dt_bias.data_ptr(),
@@ -425,7 +428,7 @@ void fused_gdn_decode_post_conv_mtp(
     torch::stable::Tensor const& num_accepted_tokens,
     torch::stable::Tensor& state, torch::stable::Tensor const& output_gate,
     torch::stable::Tensor const& norm_weight, torch::stable::Tensor& out,
-    double scale, double norm_eps) {
+    double scale, double norm_eps, const std::string& output_gate_activation) {
   using torch::headeronly::ScalarType;
 
   STD_TORCH_CHECK(
@@ -466,6 +469,9 @@ void fused_gdn_decode_post_conv_mtp(
                   "norm_weight must be a CUDA float32 or bfloat16 tensor");
   STD_TORCH_CHECK(out.is_cuda() && out.scalar_type() == ScalarType::BFloat16,
                   "out must be a CUDA bfloat16 tensor");
+  STD_TORCH_CHECK(
+      output_gate_activation == "silu" || output_gate_activation == "sigmoid",
+      "output_gate_activation must be 'silu' or 'sigmoid'");
 
   STD_TORCH_CHECK(mixed_qkv.dim() == 2,
                   "mixed_qkv must have shape [L, 2 * H * 128 + HV * 128]");
@@ -547,18 +553,28 @@ void fused_gdn_decode_post_conv_mtp(
   const auto* b_ptr = static_cast<const __nv_bfloat16*>(b.data_ptr());
   const auto* output_gate_ptr =
       static_cast<const __nv_bfloat16*>(output_gate.data_ptr());
-  const auto launch = [&]<typename StateT, int ValueHeadsPerKeyHead>() {
-    launch_gdn_decode_post_conv_mtp<StateT, ValueHeadsPerKeyHead>(
+  const auto launch = [&]<typename StateT, int ValueHeadsPerKeyHead,
+                          bool SigmoidGate>() {
+    launch_gdn_decode_post_conv_mtp<StateT, ValueHeadsPerKeyHead, SigmoidGate>(
         mixed_qkv, a_log, dt_bias, state_indices, cu_seqlens,
         num_accepted_tokens, state, norm_weight, out, a_ptr, b_ptr,
         output_gate_ptr, num_key_heads, num_value_heads, scale, norm_eps,
         strides);
   };
   const auto dispatch_state_type = [&]<int ValueHeadsPerKeyHead>() {
     if (state_scalar_type == ScalarType::Float) {
-      launch.template operator()<float, ValueHeadsPerKeyHead>();
+      if (output_gate_activation == "sigmoid") {
+        launch.template operator()<float, ValueHeadsPerKeyHead, true>();
+      } else {
+        launch.template operator()<float, ValueHeadsPerKeyHead, false>();
+      }
     } else {
-      launch.template operator()<__nv_bfloat16, ValueHeadsPerKeyHead>();
+      if (output_gate_activation == "sigmoid") {
+        launch.template operator()<__nv_bfloat16, ValueHeadsPerKeyHead, true>();
+      } else {
+        launch
+            .template operator()<__nv_bfloat16, ValueHeadsPerKeyHead, false>();
+      }
     }
   };
   switch (value_heads_per_key_head) {

csrc/libtorch_stable/ops.h

@@ -402,7 +402,7 @@ void fused_gdn_decode_post_conv_mtp(
     torch::stable::Tensor const& num_accepted_tokens,
     torch::stable::Tensor& state, torch::stable::Tensor const& output_gate,
     torch::stable::Tensor const& norm_weight, torch::stable::Tensor& out,
-    double scale, double norm_eps);
+    double scale, double norm_eps, const std::string& output_gate_activation);
 
 #endif
 

csrc/libtorch_stable/torch_bindings.cpp

@@ -535,7 +535,8 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
       "Tensor mixed_qkv, Tensor a, Tensor b, Tensor A_log, Tensor dt_bias, "
       "Tensor state_indices, Tensor cu_seqlens, Tensor num_accepted_tokens, "
       "Tensor! state, Tensor output_gate, Tensor norm_weight, Tensor! out, "
-      "float scale, float norm_eps=1e-5) -> ()");
+      "float scale, float norm_eps=1e-5, "
+      "str output_gate_activation='silu') -> ()");
 #endif
 
 #ifdef VLLM_ENABLE_KIMI_K3_ATTN_RES

tests/config/test_config_utils.py

@@ -7,6 +7,7 @@
 import pytest
 
 from vllm.config.cache import CacheConfig
+from vllm.config.scheduler import SchedulerConfig
 from vllm.config.utils import get_hash_factors, hash_factors, normalize_value
 
 # Helpers
@@ -216,6 +217,24 @@ def test_cache_config_hash_ignores_kv_cache_sizing_knobs():
     assert CacheConfig(gpu_memory_utilization=0.5).compute_hash() == base_hash
 
 
+def test_scheduler_config_hash_includes_max_num_seqs():
+    """Per-request workspace sizes must invalidate compiled graphs."""
+    base_hash = SchedulerConfig(
+        max_model_len=8192,
+        is_encoder_decoder=False,
+        max_num_batched_tokens=8192,
+        max_num_seqs=128,
+    ).compute_hash()
+    larger_batch_hash = SchedulerConfig(
+        max_model_len=8192,
+        is_encoder_decoder=False,
+        max_num_batched_tokens=8192,
+        max_num_seqs=1024,
+    ).compute_hash()
+
+    assert larger_batch_hash != base_hash
+
+
 def test_cache_config_hash_ignores_prefix_cache_retention_interval():
     base_hash = CacheConfig().compute_hash()
     assert CacheConfig(prefix_cache_retention_interval=64).compute_hash() == base_hash

tests/config/test_speculative_draft_hf_overrides.py

@@ -1,6 +1,6 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""Tests for SpeculativeConfig.compose_draft_hf_overrides.
+"""Tests for draft config overrides used by SpeculativeConfig.
 
 Callable ``hf_overrides`` on the target model config (e.g. the
 ``dummy_hf_overrides`` shrink used by ``tests/models/test_initialization.py``)
@@ -12,10 +12,12 @@
 """
 
 import functools
+from unittest.mock import MagicMock, patch
 
 import pytest
 from transformers import PretrainedConfig
 
+from vllm.config.parallel import ParallelConfig
 from vllm.config.speculative import SpeculativeConfig
 
 
@@ -134,3 +136,52 @@ def test_composed_override_is_picklable():
 
     out = composed(_make_hf_config())
     assert out.num_hidden_layers == 1
+
+
+def _make_mtp_speculative_config(
+    override: bool | None,
+    checkpoint_value: bool,
+) -> SpeculativeConfig:
+    draft_hf_config = _make_hf_config(
+        architectures=["Qwen4ExpMTP"],
+        model_type="qwen4_exp_mtp",
+        n_predict=1,
+        index_share_for_mtp_iteration=checkpoint_value,
+    )
+    draft_model_config = MagicMock(
+        model="draft",
+        hf_config=draft_hf_config,
+        architectures=draft_hf_config.architectures,
+        max_model_len=128,
+    )
+    target_model_config = MagicMock(
+        model="target",
+        max_model_len=128,
+        quantization=None,
+        hf_overrides={},
+    )
+
+    with patch("vllm.config.speculative.ModelConfig", return_value=draft_model_config):
+        return SpeculativeConfig(
+            model="draft",
+            method="mtp",
+            num_speculative_tokens=1,
+            index_share_for_mtp_iteration=override,
+            target_model_config=target_model_config,
+            target_parallel_config=ParallelConfig(),
+        )
+
+
+@pytest.mark.cpu_test
+@pytest.mark.parametrize(
+    ("override", "checkpoint_value", "expected"),
+    [(None, True, True), (False, True, False), (True, False, True)],
+)
+def test_mtp_index_share_override(
+    override: bool | None, checkpoint_value: bool, expected: bool
+):
+    speculative_config = _make_mtp_speculative_config(override, checkpoint_value)
+    assert (
+        speculative_config.draft_model_config.hf_config.index_share_for_mtp_iteration
+        is expected
+    )

tests/distributed/test_custom_all_reduce.py

@@ -9,6 +9,9 @@
 import torch.distributed as dist
 
 from vllm.distributed.communication_op import tensor_model_parallel_all_reduce  # noqa
+from vllm.distributed.device_communicators.custom_all_reduce import (
+    CustomAllreduce,
+)
 from vllm.distributed.parallel_state import get_tp_group, graph_capture
 
 from ..utils import (
@@ -23,6 +26,29 @@
     test_sizes[i] -= v % 8
 
 
+@pytest.mark.parametrize(
+    ("dtype", "expected"),
+    [
+        (torch.float32, True),
+        (torch.float16, True),
+        (torch.bfloat16, True),
+        (torch.int8, False),
+        (torch.float8_e4m3fn, False),
+    ],
+)
+def test_custom_allreduce_filters_dtype(
+    dtype: torch.dtype,
+    expected: bool,
+) -> None:
+    communicator = CustomAllreduce.__new__(CustomAllreduce)
+    communicator.disabled = False
+    communicator.world_size = 2
+    communicator.max_size = 1024
+    communicator._ptr = 0
+
+    assert communicator.should_custom_ar(torch.empty(16, dtype=dtype)) is expected
+
+
 @ray.remote(num_gpus=1, max_calls=1)
 def graph_allreduce(
     monkeypatch: pytest.MonkeyPatch,

tests/evals/qwen4_exp/README.md

@@ -0,0 +1,14 @@
+# Qwen4Exp accuracy evaluation
+
+This suite starts a Qwen3.8-Flash-Next-FP8 OpenAI-compatible server once and
+uses EvalScope to evaluate GSM8K and AIME25.
+
+```bash
+# B200
+pytest -s -v tests/evals/qwen4_exp/test_accuracy.py \
+  --config-list-file=configs/models-b200.txt
+
+# H200
+pytest -s -v tests/evals/qwen4_exp/test_accuracy.py \
+  --config-list-file=configs/models-h200.txt
+```

tests/evals/qwen4_exp/__init__.py

@@ -0,0 +1,3 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+

tests/evals/qwen4_exp/configs/Qwen3.8-Flash-Next-FP8.yaml

@@ -0,0 +1,26 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+model_name: "Qwen/Qwen3.8-Flash-Next-FP8"
+datasets:
+  gsm8k:
+    metric_threshold: 0.98
+    tolerance: 0.02
+  aime25:
+    metric_threshold: 0.90
+    tolerance: 0.05
+eval_batch_size: 32
+startup_max_wait_seconds: 1800
+generation_config:
+  seed: 1236
+  do_sample: true
+  temperature: 0.7
+  top_p: 0.8
+  top_k: 20
+  repetition_penalty: 1.0
+  presence_penalty: 1.5
+  max_tokens: 32768
+  stream: true
+  extra_body:
+    chat_template_kwargs:
+      enable_thinking: true
+server_args: "--tensor-parallel-size 4"

tests/evals/qwen4_exp/configs/models-b200.txt

@@ -0,0 +1,2 @@
+# B200 configuration for Qwen3.8-Flash-Next-FP8 accuracy evaluation
+Qwen3.8-Flash-Next-FP8.yaml

tests/evals/qwen4_exp/configs/models-h200.txt

@@ -0,0 +1,2 @@
+# H200 configuration for Qwen3.8-Flash-Next-FP8 accuracy evaluation
+Qwen3.8-Flash-Next-FP8.yaml