update .gitignore

Overseas node: original docker.io base (lmsysorg/sglang:dev-dsv41) and
default pypi. Tags follow the ymkymx convention <branch>-<sha9>-<datetime>
plus a moving <branch>-latest. Registry defaults to the Gitea instance's
container registry, override with the REGISTRY variable.

.gitea/workflows/build-image.yaml

@@ -0,0 +1,50 @@
+# Gitea Actions 自动构建 sglang 镜像（海外节点，原版源）
+# 基底 lmsysorg/sglang:dev-dsv41（docker.io），依赖走 pypi.org 默认源。
+# 触发：push 到 dsv41-pd 分支。
+#
+# 前置条件（一次性，在 Gitea 实例上配置）：
+# 1. 实例已注册 act_runner（Gitea Actions runner，标签含 ubuntu-latest）
+# 2. 仓库 Settings → Secrets 添加：
+#    REGISTRY_USERNAME / REGISTRY_PASSWORD（推送镜像的账号，如 Gitea 访问令牌）
+# 3. 可选：Settings → Actions → Variables 添加 REGISTRY（默认 git.agentwithu.com，
+#    即 Gitea 自带容器 registry；也可填 docker.io 等）
+name: build-sglang-image
+
+on:
+  push:
+    branches: [dsv41-pd]
+
+env:
+  # 直接写死 Gitea 自带 registry（vars context 在该实例上求值异常会导致回退 docker.io）
+  REGISTRY: git.agentwithu.com
+  IMAGE_NAME: minke.yu/sglang
+
+jobs:
+  build:
+    runs-on: ubuntu-latest
+    steps:
+      - uses: actions/checkout@v4
+
+      - name: Resolve tag
+        id: meta
+        run: |
+          SHA9=$(git rev-parse --short=9 HEAD)
+          echo "tag=${{ github.ref_name }}-${SHA9}-ci-$(date -u +%Y%m%d-%H%M)" >> "$GITHUB_OUTPUT"
+
+      - uses: docker/setup-buildx-action@v3
+
+      - uses: docker/login-action@v3
+        with:
+          registry: ${{ env.REGISTRY }}
+          username: ${{ secrets.REGISTRY_USERNAME }}
+          password: ${{ secrets.REGISTRY_PASSWORD }}
+
+      - uses: docker/build-push-action@v6
+        with:
+          context: .
+          file: Dockerfile.gitea
+          push: true
+          tags: |
+            ${{ env.REGISTRY }}/${{ env.IMAGE_NAME }}:${{ steps.meta.outputs.tag }}
+            ${{ env.REGISTRY }}/${{ env.IMAGE_NAME }}:${{ github.ref_name }}-latest
+          # 注意：type=gha 缓存在本实例的 runner 上会 404（cache server 未配），勿加回

Dockerfile.gitea

@@ -0,0 +1,11 @@
+# Gitea Actions 用（海外节点）：原版 docker.io 基底 + 原版 pypi
+# 与 b300 离线版（/data/ymk/build/Dockerfile）的区别：
+#   - 基底直用 docker.io 的 lmsysorg/sglang:dev-dsv41（不走 umirror/DaoCloud）
+#   - 不设 PIP_INDEX_URL，用默认 pypi.org
+#   - 源码由 CI checkout 后经 COPY 进镜像（不用 in-image clone）
+FROM lmsysorg/sglang:dev-dsv41
+
+# checkout 含 .git，setuptools-scm 可打戳（main 系分支 version 显示 dev 属正常）
+COPY . /sgl-workspace/sglang
+
+RUN pip install --no-cache-dir -e /sgl-workspace/sglang/python

python/sglang/srt/arg_groups/deepseek_v4_hook.py

@@ -249,23 +249,33 @@ def validate_deepseek_v41_features(server_args: ServerArgs) -> None:
         if (
             read_ragged_verify_mode() is not RaggedVerifyMode.STATIC
             or cfg.disaggregation_transfer_backend != "mooncake"
-            or cfg.dp_size != 1
-            or cfg.enable_dp_attention
             or cfg.attn_cp_size != 1
             or cfg.dcp_size != 1
         ):
             raise ValueError(
                 "DeepSeek-V4.1 DSpark PD requires static verify, Mooncake, "
-                "DP=1 and CP=1. Both servers must enable DSpark with the same "
+                "and CP=1 on both servers. DP attention is supported when "
-                "block size and TP size."
+                "both servers use the same block size and target/draft KV layout."
             )
 
     from sglang.srt.model_executor.cuda_graph_config import Backend, Phase, with_phase
 
     prefill_graph = cfg.cuda_graph_config.prefill
-    if prefill_graph.backend != Backend.DISABLED and prefill_graph.max_seq_len is None:
+    cp_breakable_prefill = (
-        # The captured low-ratio indexer scores a static context width; 16k
+        cfg.enable_prefill_cp
-        # keeps it inside the candidate window at under 1 ms per layer.
+        and cfg.cp_strategy == "interleave"
+        and cfg.tp_size > 1
+        and prefill_graph.backend == Backend.BREAKABLE
+    )
+    if (
+        prefill_graph.backend != Backend.DISABLED
+        and prefill_graph.max_seq_len is None
+        and not cp_breakable_prefill
+    ):
+        # The non-CP captured low-ratio indexer scores a static context width.
+        # CP BCG runs these sources eagerly with live prefix metadata, so this
+        # default would only force long-prefix CP batches back to eager.
+        # Explicit max_seq_len values still constrain both paths.
         declare_resolution(
             server_args,
             "validate_deepseek_v41_features",

python/sglang/srt/arg_groups/model_hook.py

@@ -966,12 +966,14 @@ def handle_language_model_only(server_args: Any):
     ):
         if flag:
             raise ValueError(f"--language-model-only cannot be combined with {name}")
-    if cfg.disaggregation_mode != "null":
+    hf_config = model_config_of(server_args).hf_config
+    # V4.1 text-only workers use the standard PD KV transfer path.
+    if cfg.disaggregation_mode != "null" and hf_config.model_type != "deepseek_v41":
         raise ValueError(
             "--language-model-only is incompatible with --disaggregation-mode "
             "prefill/decode"
         )
-    architectures = model_config_of(server_args).hf_config.architectures
+    architectures = hf_config.architectures
     if not any(
         a in server_args.LANGUAGE_MODEL_ONLY_ARCHITECTURES for a in architectures
     ):

python/sglang/srt/disaggregation/common/conn.py

@@ -945,9 +945,33 @@ class CommonKVManager(BaseKVManager):
                     "enable DSpark with the same block size and target/draft KV "
                     "layout. Upgrade both servers together."
                 )
-            if info.attn_tp_size != self.attn_tp_size:
+            same_tp_with_prefill_cp = (
+                info.attn_cp_size > 1
+                and (self.is_mla_backend or self.is_hybrid_mla_backend)
+                and self.attn_cp_size == 1
+                and info.attn_tp_size * info.attn_cp_size == self.attn_tp_size
+            )
+            # Combined branch (40323-series + 40177): prefill CP can also pair
+            # with a DP-attention decode server. MLA KV is replicated across
+            # prefill CP ranks, so per-rank layouts match when attn_tp matches.
+            dp_decode_with_prefill_cp = (
+                info.attn_cp_size > 1
+                and self.attn_cp_size == 1
+                and (self.is_mla_backend or self.is_hybrid_mla_backend)
+                and info.attn_tp_size == self.attn_tp_size
+            )
+            non_cp_mla_layout = info.attn_cp_size == self.attn_cp_size == 1 and (
+                self.is_mla_backend or self.is_hybrid_mla_backend
+            )
+            if info.attn_tp_size != self.attn_tp_size and not (
+                same_tp_with_prefill_cp
+                or dp_decode_with_prefill_cp
+                or non_cp_mla_layout
+            ):
                 raise RuntimeError(
-                    "DeepSeek-V4.1 DSpark PD requires the same TP size on both servers"
+                    "DeepSeek-V4.1 DSpark PD requires matching attention TP "
+                    "unless both servers use CP=1 with an MLA KV layout, "
+                    "or prefill runs CP with an MLA KV layout"
                 )
 
         if self.dcp_size > 1:

python/sglang/srt/disaggregation/mooncake/conn.py

@@ -82,6 +82,115 @@ FAILED_SESSION_RECOVERIES = Counter(
 )
 
 
+# ---------------------------------------------------------------------------
+# Intra-node NVLink transport helpers.
+#
+# Mooncake's IntraNodeNvlinkTransport can only register and reach *device*
+# memory (it IPC-opens the remote cudaMalloc segments). Host-resident regions
+# (aux buffers, some state components) cannot be registered: one host region
+# makes the whole registerLocalMemoryBatch fail, and the engine then rolls
+# back *every* region, leaving the segment descriptor empty and all KV
+# transfers failing with "Requested address ... not found". When the
+# intra-node NVLink transport is active we therefore
+#   1. register only device-memory regions, and
+#   2. route blocks whose source is host memory over the ordered zmq channel
+#      (same ordering guarantee the aux TCP path relies on) instead of the
+#      transfer engine.
+# ---------------------------------------------------------------------------
+
+
+import ctypes as _ctypes
+
+_CUDA_MEMORY_TYPE_DEVICE = 2
+
+try:
+    from cuda.bindings import runtime as _cudart
+except ImportError:  # pragma: no cover - cuda-python is always present in images
+    _cudart = None
+
+
+def _is_device_pointer(ptr: int) -> bool:
+    """Probe a *local* pointer with cudaPointerGetAttributes.
+
+    Only valid for pointers owned by this process (never probe remote
+    segment addresses). Returns False on any error so the caller falls back
+    to the safe host path.
+    """
+    if _cudart is None:
+        # Cannot tell; assume device so behavior stays unchanged.
+        return True
+    err, attr = _cudart.cudaPointerGetAttributes(int(ptr))
+    if int(err) != 0:
+        # Clear the error so subsequent CUDA calls are not poisoned.
+        _cudart.cudaGetLastError()
+        return False
+    return int(attr.type) == _CUDA_MEMORY_TYPE_DEVICE
+
+
+def _read_bytes_from_address(addr: int, length: int) -> Optional[bytes]:
+    if length <= 0:
+        return b""
+    if _is_device_pointer(addr):
+        buf = bytearray(length)
+        # cudaMemcpyDeviceToHost = 2; synchronous default-stream copy.
+        err, = _cudart.cudaMemcpy(
+            _ctypes.addressof((_ctypes.c_char * length).from_buffer(buf)),
+            int(addr),
+            length,
+            2,
+        )
+        if int(err) != 0:
+            logger.error(
+                f"cudaMemcpy D2H failed (err={err}) for addr {hex(addr)} len {length}"
+            )
+            return None
+        return bytes(buf)
+    return _ctypes.string_at(int(addr), length)
+
+
+def _write_bytes_to_address(addr: int, data: bytes) -> bool:
+    if not data:
+        return True
+    if _is_device_pointer(addr):
+        buf = _ctypes.create_string_buffer(data, len(data))
+        # cudaMemcpyHostToDevice = 1; synchronous default-stream copy.
+        err, = _cudart.cudaMemcpy(
+            int(addr), _ctypes.addressof(buf), len(data), 1
+        )
+        if int(err) != 0:
+            logger.error(
+                f"cudaMemcpy H2D failed (err={err}) for addr {hex(addr)} "
+                f"len {len(data)}"
+            )
+            return False
+        return True
+    _ctypes.memmove(int(addr), data, len(data))
+    return True
+
+
+_NVLINK_INTRA_ACTIVE = None
+
+
+def _nvlink_intra_transport_active() -> bool:
+    """Whether mooncake installed the intra-node NVLink transport.
+
+    Mirrors the env probing in mooncake's transfer_engine_impl.cpp: the
+    transport is installed iff MC_INTRANODE_NVLINK is set (any value), or an
+    equivalent protocol selection was made.
+    """
+    global _NVLINK_INTRA_ACTIVE
+    if _NVLINK_INTRA_ACTIVE is None:
+        active = bool(
+            os.environ.get("MC_INTRANODE_NVLINK")
+            or os.environ.get("MC_INTRA_NVLINK")
+        )
+        if not active:
+            proto = (os.environ.get("MOONCAKE_PROTOCOL") or "").strip().lower()
+            active = proto in ("nvlink_intra", "nvlink-intra", "intra_nvlink")
+        _NVLINK_INTRA_ACTIVE = active
+    return _NVLINK_INTRA_ACTIVE
+
+
 # decode
 @dataclasses.dataclass
 class TransferInfo:
@@ -214,6 +323,7 @@ class KVArgsRegisterInfo:
 
 class MooncakeKVManager(StagingManagerMixin, CommonKVManager):
     AUX_DATA_HEADER = b"AUX_DATA"
+    STATE_DATA_HEADER = b"STATE_DATA"
     # Implements teardown() below, so runtime PD role switching is supported.
     supports_role_switch = True
 
@@ -227,6 +337,10 @@ class MooncakeKVManager(StagingManagerMixin, CommonKVManager):
         super().__init__(args, disaggregation_mode, server_args, is_mla_backend)
         self.init_engine()
         self.register_buffer_to_engine()
+        # session_id -> (endpoint, dst_port, room), used to route host-memory
+        # transfer blocks over zmq when the intra-node NVLink transport is
+        # active (it cannot reach host memory). Populated on bootstrap.
+        self._session_endpoint_map = {}
         self.enable_staging = envs.SGLANG_DISAGG_STAGING_BUFFER.get()
         self.max_transfer_batch_indices = (
             envs.SGLANG_MOONCAKE_MAX_TRANSFER_BATCH_INDICES.get()
@@ -322,6 +436,13 @@ class MooncakeKVManager(StagingManagerMixin, CommonKVManager):
         Deduped because the unified memory pool reports one raw buffer as both
         its KV and its mamba state component, and double registration fails in
         the engine.
+
+        When the intra-node NVLink transport is active, host-memory regions
+        (aux buffers, some state components) are skipped: the transport only
+        accepts device memory, and a single host region fails the whole batch
+        and triggers a full engine-side rollback that would unregister the KV
+        pools too. Host-resident payloads are instead exchanged over the
+        ordered zmq channel (see _transfer_data / send_aux).
         """
         regions: List[Tuple[int, int]] = []
         seen: Set[Tuple[int, int]] = set()
@@ -338,6 +459,24 @@ class MooncakeKVManager(StagingManagerMixin, CommonKVManager):
             self.kv_args.state_data_ptrs, self.kv_args.state_data_lens
         ):
             add(ptrs, lens)
+
+        if _nvlink_intra_transport_active():
+            device_regions = []
+            skipped = []
+            for ptr, length in regions:
+                if _is_device_pointer(ptr):
+                    device_regions.append((ptr, length))
+                else:
+                    skipped.append((ptr, length))
+            if skipped:
+                logger.info(
+                    "Intra-node NVLink transport: skipping %d host-memory "
+                    "regions from engine registration (they will be exchanged "
+                    "over the zmq channel instead): %s",
+                    len(skipped),
+                    [(hex(p), l) for p, l in skipped[:8]],
+                )
+            regions = device_regions
         return regions
 
     def register_buffer_to_engine(self):
@@ -748,10 +887,63 @@ class MooncakeKVManager(StagingManagerMixin, CommonKVManager):
         if not transfer_blocks:
             return 0
 
-        src_addrs, dst_addrs, lengths = zip(*transfer_blocks)
+        if not _nvlink_intra_transport_active():
-        return self.engine.batch_transfer_sync(
+            src_addrs, dst_addrs, lengths = zip(*transfer_blocks)
-            mooncake_session_id, list(src_addrs), list(dst_addrs), list(lengths)
+            return self.engine.batch_transfer_sync(
-        )
+                mooncake_session_id, list(src_addrs), list(dst_addrs), list(lengths)
+            )
+
+        # Intra-node NVLink transport can only move device memory. Partition
+        # blocks by the *local source* pointer (probing a local pointer is
+        # safe; the remote dst is never probed): device-sourced blocks go
+        # through the engine as usual, host-sourced blocks are shipped over
+        # the ordered zmq channel and written into the peer's buffer by the
+        # receiver (see _handle_state_data). This mirrors the aux TCP path.
+        device_blocks = []
+        host_blocks = []
+        for src, dst, length in transfer_blocks:
+            if _is_device_pointer(src):
+                device_blocks.append((src, dst, length))
+            else:
+                host_blocks.append((src, dst, length))
+
+        rc = 0
+        if device_blocks:
+            src_addrs, dst_addrs, lengths = zip(*device_blocks)
+            rc = self.engine.batch_transfer_sync(
+                mooncake_session_id, list(src_addrs), list(dst_addrs), list(lengths)
+            )
+        if rc == 0 and host_blocks:
+            rc = self._send_host_blocks_tcp(mooncake_session_id, host_blocks)
+        return rc
+
+    def _send_host_blocks_tcp(self, mooncake_session_id, host_blocks):
+        target = self._session_endpoint_map.get(mooncake_session_id)
+        if target is None:
+            logger.error(
+                f"No zmq endpoint known for mooncake session "
+                f"{mooncake_session_id}; cannot deliver {len(host_blocks)} "
+                "host-memory transfer blocks"
+            )
+            return -1
+        endpoint, dst_port, room = target
+        na = NetworkAddress(endpoint, dst_port)
+        for src, dst, length in host_blocks:
+            data = _read_bytes_from_address(src, length)
+            if data is None:
+                return -1
+            self._send_multipart_locked(
+                na.to_tcp(),
+                [
+                    MooncakeKVManager.STATE_DATA_HEADER,
+                    str(room).encode("ascii"),
+                    str(int(dst)).encode("ascii"),
+                    struct.pack(">I", len(data)),
+                    data,
+                ],
+                is_ipv6=na.is_ipv6,
+            )
+        return 0
 
     def _send_kvcache_generic(
         self,
@@ -1482,8 +1674,10 @@ class MooncakeKVManager(StagingManagerMixin, CommonKVManager):
     ):
         # TODO(shangming): Fix me when nvlink_transport of Mooncake is bug-free
         if (
-            self.enable_custom_mem_pool and self.custom_mem_pool_type == "NVLINK"
+            (self.enable_custom_mem_pool and self.custom_mem_pool_type == "NVLINK")
-        ) or envs.SGLANG_MOONCAKE_SEND_AUX_TCP.get():
+            or envs.SGLANG_MOONCAKE_SEND_AUX_TCP.get()
+            or _nvlink_intra_transport_active()
+        ):
             return self.send_aux_tcp(req, prefill_aux_index, dst_aux_ptrs)
 
         transfer_blocks = []
@@ -1566,6 +1760,71 @@ class MooncakeKVManager(StagingManagerMixin, CommonKVManager):
             f"Received AUX_DATA for bootstrap_room {room} with length:{len(data)}"
         )
 
+    def _host_transfer_regions(self):
+        """Address ranges this process published as transfer targets.
+
+        Used to validate STATE_DATA writes. Built lazily because kv_args is
+        fully populated only after registration.
+        """
+        regions = getattr(self, "_host_transfer_regions_cache", None)
+        if regions is None:
+            regions = []
+            for ptr, length in zip(
+                self.kv_args.kv_data_ptrs or [], self.kv_args.kv_data_lens or []
+            ):
+                regions.append((int(ptr), int(ptr) + int(length)))
+            for ptr, length in zip(
+                self.kv_args.aux_data_ptrs or [], self.kv_args.aux_data_lens or []
+            ):
+                regions.append((int(ptr), int(ptr) + int(length)))
+            for ptrs, lens in zip(
+                self.kv_args.state_data_ptrs or [], self.kv_args.state_data_lens or []
+            ):
+                for ptr, length in zip(ptrs or [], lens or []):
+                    regions.append((int(ptr), int(ptr) + int(length)))
+            self._host_transfer_regions_cache = regions
+        return regions
+
+    def _handle_state_data(self, msg: List[bytes]):
+        """Handle STATE_DATA messages received by the decode thread.
+
+        Carries one host-memory transfer block that could not go through the
+        intra-node NVLink transport. Written directly into the local buffer at
+        the destination address; ordering against the final status message is
+        guaranteed by the shared per-endpoint zmq socket.
+        """
+        room = int(msg[1].decode("ascii"))
+        dst_addr = int(msg[2].decode("ascii"))
+        data_length = struct.unpack(">I", msg[3])[0]
+        data = msg[4]
+
+        if len(data) != data_length:
+            logger.error(f"STATE_DATA length mismatch for bootstrap_room {room}")
+            return
+
+        in_region = any(
+            start <= dst_addr and dst_addr + len(data) <= end
+            for start, end in self._host_transfer_regions()
+        )
+        if not in_region:
+            logger.error(
+                f"STATE_DATA for bootstrap_room {room} targets unknown region "
+                f"{hex(dst_addr)}..{hex(dst_addr + len(data))}; dropping"
+            )
+            return
+
+        if not _write_bytes_to_address(dst_addr, data):
+            logger.error(
+                f"STATE_DATA write failed for bootstrap_room {room} at "
+                f"{hex(dst_addr)} len {len(data)}"
+            )
+            return
+
+        logger.debug(
+            f"Received STATE_DATA for bootstrap_room {room} at {hex(dst_addr)} "
+            f"with length:{len(data)}"
+        )
+
     def _get_dsa_cache_transfer_skip_flags(
         self, info: Optional[KVArgsRegisterInfo]
     ) -> Tuple[bool, bool]:
@@ -2412,6 +2671,8 @@ class MooncakeKVManager(StagingManagerMixin, CommonKVManager):
                 ):
                     self._staging_outstanding.pop(kv_chunk.room, None)
                     if kv_chunk.room in self.transfer_infos:
+                        for sid in self.transfer_infos[kv_chunk.room]:
+                            self._session_endpoint_map.pop(sid, None)
                         self.transfer_infos.pop(kv_chunk.room)
                     self.req_to_decode_prefix_len.pop(kv_chunk.room, None)
                     if self.enable_staging:
@@ -2557,6 +2818,11 @@ class MooncakeKVManager(StagingManagerMixin, CommonKVManager):
                     self.transfer_infos[room][mooncake_session_id] = (
                         TransferInfo.from_zmq(waiting_req_bytes)
                     )
+                    self._session_endpoint_map[mooncake_session_id] = (
+                        self.transfer_infos[room][mooncake_session_id].endpoint,
+                        self.transfer_infos[room][mooncake_session_id].dst_port,
+                        room,
+                    )
                     # NOTE: after bootstrapping we can mark the req as waiting for input
                     if len(self.transfer_infos[room]) == required_dst_info_num:
                         self.resolve_kv_replica_factor(self.transfer_infos[room])
@@ -2585,6 +2851,9 @@ class MooncakeKVManager(StagingManagerMixin, CommonKVManager):
                 if msg[0] == MooncakeKVManager.AUX_DATA_HEADER:
                     self._handle_aux_data(msg)
                     continue
+                if msg[0] == MooncakeKVManager.STATE_DATA_HEADER:
+                    self._handle_state_data(msg)
+                    continue
 
                 # Staging: prefill notifies a chunk written to staging buffer
                 if msg[0] == b"CHUNK_READY":

python/sglang/srt/layers/attention/deepseek_v4_backend.py

@@ -736,7 +736,14 @@ class DSV4AttnMetadata:
             if src_val is None and dst_val is None:
                 continue
             assert dst_val is not None, f"{field_name=} {src_val=} {dst_val=}"
-            dst_val.copy_(src_val)
+            shape_mismatch = dst_val.shape != src_val.shape
+            assert not shape_mismatch or field_name in self._CP_GLOBAL_FIELDS, (
+                f"Only CP-global replay metadata may use a shorter live prefix, "
+                f"got {field_name=} {src_val.shape=} {dst_val.shape=}"
+            )
+            _copy_tensor_allowing_storage_alias(
+                dst_val, src_val, pad_value=0 if shape_mismatch else None
+            )
 
         # These fields are safe to replace because captured kernels only need
         # the current per-replay objects, or the field is produced inside the
@@ -988,6 +995,27 @@ def _prefill_graph_max_seq_len() -> Optional[int]:
     return get_exec().graph.cuda_graph_config.prefill.max_seq_len
 
 
+def _copy_tensor_allowing_storage_alias(
+    dst: torch.Tensor, src: torch.Tensor, *, pad_value: Optional[int] = None
+) -> None:
+    """Copy replay metadata while preserving capture-stable destination addresses."""
+    if dst is src:
+        return
+    if dst.untyped_storage().data_ptr() == src.untyped_storage().data_ptr():
+        src = src.clone()
+    if dst.shape == src.shape:
+        dst.copy_(src)
+        return
+    assert (
+        pad_value is not None
+        and dst.ndim == src.ndim
+        and dst.shape[0] >= src.shape[0]
+        and dst.shape[1:] == src.shape[1:]
+    ), f"Cannot copy replay metadata from {src.shape=} to {dst.shape=}"
+    dst.fill_(pad_value)
+    dst[: src.shape[0]].copy_(src)
+
+
 @dataclass
 class DSV4Metadata:
     core_attn_metadata: DSV4AttnMetadata
@@ -1254,6 +1282,11 @@ class DeepseekV4AttnBackend(
         ] = None
         self.online_c128_mtp = OnlineC128MTPController(self)
         self.sparse_prefill_workspace = SparsePrefillWorkspace(self.device)
+        # CP V4.1 consumers share compressed KV across layers. Separate ratio
+        # workspaces keep those prefixes intact while each layer refreshes SWA.
+        self.shared_compressed_prefill_workspaces = {
+            ratio: SparsePrefillWorkspace(self.device) for ratio in (1, 2)
+        }
         spec_alg = model_runner.spec_algorithm
         self.needs_cpu_seq_lens = not spec_alg.is_dspark() and (
             not _is_cuda or self.online_c128_mtp.enabled()
@@ -1566,8 +1599,12 @@ class DeepseekV4AttnBackend(
 
     @property
     def low_ratio_prefill_graph(self) -> bool:
+        """Whether ratio-1/2 sources use captured projections and indexer metadata."""
         return (
-            bool(self.low_ratios) and _has_dense_fp4_indexer() and _is_sm100_or_newer()
+            bool(self.low_ratios)
+            and _has_dense_fp4_indexer()
+            and _is_sm100_or_newer()
+            and get_parallel().attn_cp_size == 1
         )
 
     def can_run_prefill_cuda_graph(self, forward_batch: ForwardBatch) -> bool:
@@ -2870,7 +2907,7 @@ class DeepseekV4AttnBackend(
                 q_lora[:num_local],
                 positions[:num_local].to(torch.int64),
                 forward_batch,
-                torch.tensor(q_lens_cpu, dtype=torch.int32, device=x.device),
+                self._move_to_device(q_lens_cpu),
                 q_lens_cpu,
             )
 
@@ -3236,7 +3273,9 @@ class DeepseekV4AttnBackend(
                 continue
             j = torch.arange(lc, device=device)
             slot_chunks.append(
-                self.req_to_token[req_pool_indices[r], j * ratio].to(torch.int64)
+                self.req_to_token[req_pool_indices[r : r + 1], j * ratio].to(
+                    torch.int64
+                )
                 // ratio
             )
             start += lc
@@ -3261,7 +3300,7 @@ class DeepseekV4AttnBackend(
         weights = indexer.head_weights(x).float()
         compress_lens = ((pos + 1) // ratio).to(torch.int32)
         ks = torch.repeat_interleave(
-            torch.tensor(starts, dtype=torch.int32, device=device),
+            self._move_to_device(starts),
             q_lens.to(torch.int64),
             output_size=num_tokens,
         )
@@ -3447,17 +3486,44 @@ class DeepseekV4AttnBackend(
                 self.candidate_indexer.publish_decode(inputs, page_indices, raw_indices)
             )
             return
-        logits = deep_gemm_fp4_paged_mqa_logits(
+        if isinstance(metadata.deep_gemm_metadata, list):
-            (q_fp4, q_sf),
+            topk_plans = metadata.topk_metadata_chunks
-            k_cache,
+            assert not metadata.use_topk_v2 or topk_plans is not None
-            weights,
+            for chunk_idx, (rows, plan) in enumerate(metadata.row_chunks()):
-            metadata.compressed_seq_lens,
+                logits = deep_gemm_fp4_paged_mqa_logits(
-            metadata.page_table,
+                    (q_fp4[rows], q_sf[rows]),
-            metadata.deep_gemm_metadata,
+                    k_cache,
-            metadata.max_compressed_seq_len,
+                    weights[rows],
-        )
+                    metadata.compressed_seq_lens[rows],
-        # TODO(dark): add bf16 topk
+                    metadata.page_table[rows],
-        topk_transform_paged_from_metadata(logits, metadata, page_indices, raw_indices)
+                    plan,
+                    metadata.max_compressed_seq_len,
+                )
+                # TODO(dark): add bf16 topk
+                topk_transform_paged_from_metadata(
+                    logits,
+                    metadata,
+                    page_indices,
+                    raw_indices,
+                    rows=rows,
+                    topk_metadata=(
+                        topk_plans[chunk_idx] if topk_plans is not None else None
+                    ),
+                )
+        else:
+            logits = deep_gemm_fp4_paged_mqa_logits(
+                (q_fp4, q_sf),
+                k_cache,
+                weights,
+                metadata.compressed_seq_lens,
+                metadata.page_table,
+                metadata.deep_gemm_metadata,
+                metadata.max_compressed_seq_len,
+            )
+            # TODO(dark): add bf16 topk
+            topk_transform_paged_from_metadata(
+                logits, metadata, page_indices, raw_indices
+            )
 
     # TODO(candidate): Hopper decode still publishes / consumes masks inline (torch
     # top-k); move into the candidate indexer with the prefill paths.
@@ -3945,20 +4011,38 @@ class DeepseekV4AttnBackend(
                 compress_ratio, core_attn_metadata, extra_page_size
             )
             n_compressed = flat_token_ids.shape[0]
-            workspace = self.sparse_prefill_workspace.get(
+            reuse_compressed = compress_ratio in (1, 2) and is_cp_active(forward_batch)
-                n_compressed + cache.swa_token_ids.shape[0]
+            workspace_pool = (
+                self.shared_compressed_prefill_workspaces[compress_ratio]
+                if reuse_compressed
+                else self.sparse_prefill_workspace
             )
+            workspace = workspace_pool.get(n_compressed + cache.swa_token_ids.shape[0])
             compressed_slice = workspace[:n_compressed]
             swa_slice = workspace[n_compressed:]
 
         if compressed_slice is not None:
-            dequantize_k_cache_paged(
+            source_key = None
-                extra_k_cache,
+            if reuse_compressed:
-                flat_token_ids,
+                source_layer = token_to_kv_pool.source_layer_of(layer_id)
-                page_size=extra_page_size,
+                source_key = (source_layer, workspace.data_ptr())
-                out=compressed_slice,
+                gather = cache.compressed[compress_ratio]
-                layout=token_to_kv_pool.get_extra_key_layout(layer_id),
+            # A source layer may have just updated its cache in place. Consumer
-            )
+            # layers only reuse the compressed prefix; their top-k and SWA stay live.
+            if (
+                source_key is None
+                or layer_id == source_key[0]
+                or gather.dequantized_source != source_key
+            ):
+                dequantize_k_cache_paged(
+                    extra_k_cache,
+                    flat_token_ids,
+                    page_size=extra_page_size,
+                    out=compressed_slice,
+                    layout=token_to_kv_pool.get_extra_key_layout(layer_id),
+                )
+                if source_key is not None:
+                    gather.dequantized_source = source_key
         dequantize_k_cache_paged(
             token_to_kv_pool.get_swa_key_buffer_radix(layer_id),
             cache.swa_token_ids,

python/sglang/srt/layers/attention/dsv4/candidate_indexer_deep_gemm.py

@@ -21,6 +21,7 @@ from sglang.srt.layers.attention.dsv4.candidate_indexer import (
 )
 from sglang.srt.layers.attention.dsv4.indexer import (
     deep_gemm_fp4_paged_mqa_logits,
+    topk_transform_paged_from_metadata,
 )
 
 CANDIDATE_BLOCK_SIZE = 8  # positions per block; DeepGEMM accepts 8 or 16
@@ -176,6 +177,10 @@ class DeepGemmCandidateIndexer:
         metadata."""
         metadata = inputs.metadata
         seq_lens = metadata.compressed_seq_lens.reshape(-1)
+        if isinstance(metadata.deep_gemm_metadata, list):
+            return self._publish_decode_chunked(
+                inputs, page_indices, raw_indices, seq_lens
+            )
         logits = deep_gemm_fp4_paged_mqa_logits(
             (inputs.q_fp4, inputs.q_sf),
             inputs.k_cache,
@@ -231,6 +236,83 @@ class DeepGemmCandidateIndexer:
                 ready=ready,
             )
 
+    def _publish_decode_chunked(
+        self,
+        inputs: IndexerInputs,
+        page_indices: torch.Tensor,
+        raw_indices: Optional[torch.Tensor],
+        seq_lens: torch.Tensor,
+    ) -> SparseBlockTable:
+        """Publish an eager forward whose dense logits are bounded by row chunks.
+
+        CUDA-graph metadata always carries one tensor schedule and keeps using the
+        asynchronous fast path above.  The exceptional eager path stays on the
+        current stream so each chunk's full logits can be released before the next.
+        """
+        metadata = inputs.metadata
+        block_chunks = []
+        phys_block_chunks = []
+        valid_len_chunks = []
+        topk_plans = metadata.topk_metadata_chunks
+        assert not metadata.use_topk_v2 or topk_plans is not None
+
+        for chunk_idx, (rows, plan) in enumerate(metadata.row_chunks()):
+            logits = deep_gemm_fp4_paged_mqa_logits(
+                (inputs.q_fp4[rows], inputs.q_sf[rows]),
+                inputs.k_cache,
+                inputs.weights[rows],
+                metadata.compressed_seq_lens[rows],
+                metadata.page_table[rows],
+                plan,
+                metadata.max_compressed_seq_len,
+            )
+            topk_transform_paged_from_metadata(
+                logits,
+                metadata,
+                page_indices,
+                raw_indices,
+                rows=rows,
+                topk_metadata=(
+                    topk_plans[chunk_idx] if topk_plans is not None else None
+                ),
+            )
+
+            chunk_seq_lens = seq_lens[rows]
+            nblocks, row_valid_lens = candidate_row_lens(
+                chunk_seq_lens, self.topk_blocks
+            )
+            blocks = amax_topk_blocks(logits, chunk_seq_lens, nblocks, self.topk_blocks)
+            phys_blocks = sort_candidate_blocks(
+                blocks,
+                chunk_seq_lens,
+                metadata.page_table[rows],
+                metadata.compressed_page_size,
+            )
+            block_chunks.append(blocks)
+            phys_block_chunks.append(phys_blocks)
+            valid_len_chunks.append(row_valid_lens)
+
+        blocks = torch.cat(block_chunks)
+        phys_blocks = torch.cat(phys_block_chunks)
+        row_valid_lens = torch.cat(valid_len_chunks)
+        schedule = build_sparse_indexer_schedule(
+            blocks,
+            seq_lens,
+            metadata.page_table,
+            metadata.compressed_page_size,
+            inputs.q_fp4.dtype,
+            self._request_ids(inputs.request_ids, inputs.num_rows, blocks.device),
+        )
+        ready = torch.cuda.Event()
+        ready.record(torch.cuda.current_stream())
+        return SparseBlockTable(
+            blocks=blocks,
+            schedule=schedule,
+            phys_blocks=phys_blocks,
+            valid_lens=row_valid_lens,
+            ready=ready,
+        )
+
     def _scores(self, table: SparseBlockTable, inputs: IndexerInputs) -> torch.Tensor:
         return sparse_logits(
             inputs.q_fp4,

python/sglang/srt/layers/attention/dsv4/indexer.py

@@ -475,27 +475,40 @@ def topk_transform_paged_from_metadata(
     metadata,
     page_indices: torch.Tensor,
     raw_indices: Optional[torch.Tensor] = None,
+    *,
+    rows: Optional[slice] = None,
+    topk_metadata: Optional[torch.Tensor] = None,
 ) -> None:
     """Pool slots into ``page_indices`` (``-1`` past the valid count) and, when given,
     positions into ``raw_indices``; ``metadata`` is a ``PagedIndexerMetadata``."""
+    if rows is None:
+        seq_lens = metadata.compressed_seq_lens
+        page_table = metadata.page_table
+        out_page_indices = page_indices
+        out_raw_indices = raw_indices
+    else:
+        seq_lens = metadata.compressed_seq_lens[rows]
+        page_table = metadata.page_table[rows]
+        out_page_indices = page_indices[rows]
+        out_raw_indices = raw_indices[rows] if raw_indices is not None else None
     if metadata.use_topk_v2:
         topk_transform_paged_v2(
             logits,
-            metadata.compressed_seq_lens,
+            seq_lens,
-            metadata.page_table,
+            page_table,
-            page_indices,
+            out_page_indices,
             metadata.compressed_page_size,
-            metadata.topk_metadata,
+            metadata.topk_metadata if topk_metadata is None else topk_metadata,
-            raw_indices,
+            out_raw_indices,
         )
     else:
         topk_transform_paged(
             logits,
-            metadata.compressed_seq_lens,
+            seq_lens,
-            metadata.page_table,
+            page_table,
-            page_indices,
+            out_page_indices,
             metadata.compressed_page_size,
-            raw_indices,
+            out_raw_indices,
         )
 
 

python/sglang/srt/layers/attention/dsv4/metadata.py

@@ -21,6 +21,7 @@ from sglang.srt.model_executor.runner_backend_utils.breakable_cuda_graph.context
 from sglang.srt.model_executor.runner_backend_utils.tc_piecewise_cuda_graph import (
     is_in_tc_piecewise_cuda_graph,
 )
+from sglang.srt.model_executor.runner_utils.capture_mode import get_is_capture_mode
 from sglang.srt.utils import is_hip, is_sm120_supported, is_xpu
 
 logger = logging.getLogger(__name__)
@@ -306,7 +307,8 @@ class PagedIndexerMetadata:
         ):
             return None
         if (
-            torch.cuda.is_current_stream_capturing()
+            get_is_capture_mode()
+            or torch.cuda.is_current_stream_capturing()
             or is_in_breakable_cuda_graph()
             or is_in_tc_piecewise_cuda_graph()
         ):
@@ -325,14 +327,27 @@ class PagedIndexerMetadata:
 
     def row_chunks(self):
         num_rows = self.compressed_seq_lens.shape[0]
-        if self.row_chunk <= 0:
+        if self.row_chunk > 0:
+            rows_per_chunk = self.row_chunk
+        elif isinstance(self.deep_gemm_metadata, list):
+            assert self.rows_per_chunk is not None, (
+                "chunked DeepGEMM metadata requires rows_per_chunk"
+            )
+            rows_per_chunk = self.rows_per_chunk
+        else:
             return [(slice(0, num_rows), self.deep_gemm_metadata)]
-        return [
+
-            (slice(start, min(start + self.row_chunk, num_rows)), plan)
+        chunks = [
+            (slice(start, min(start + rows_per_chunk, num_rows)), plan)
             for start, plan in zip(
-                range(0, num_rows, self.row_chunk), self.deep_gemm_metadata
+                range(0, num_rows, rows_per_chunk), self.deep_gemm_metadata
             )
         ]
+        assert chunks and chunks[-1][0].stop == num_rows, (
+            f"chunk schedules do not cover all rows: {num_rows=} {rows_per_chunk=} "
+            f"{len(chunks)=}"
+        )
+        return chunks
 
     def copy_(self, other: PagedIndexerMetadata):
         # A chunked schedule list has no in-place copy; rebind it instead.

python/sglang/srt/layers/attention/dsv4/sparse_prefill_utils.py

@@ -78,10 +78,11 @@ def use_dsv4_q8kv8_sparse_prefill(dsv4_prefill_backend: str = "auto") -> bool:
 class SparsePrefillWorkspace:
     """Backend-owned scratch storage for sparse prefill KV dequantization.
 
-    The workspace contents are fully overwritten before every attention call,
+    Callers normally overwrite the entire workspace. Shared compressed-KV
-    so token buckets and compression ratios can safely share one buffer. Sparse
+    callers keep separate workspaces per ratio and track prefix validity in the
-    prefill executes eagerly and serially on the supported paths, which makes it
+    per-forward gather cache, including the allocation address. Sparse prefill
-    safe to replace the scratch allocation when a larger extent is needed.
+    executes eagerly and serially on the supported paths, so the allocation can
+    be replaced when a larger extent is needed.
     """
 
     def __init__(self, device: torch.device):
@@ -275,14 +276,18 @@ class CompressedGather:
     # chunk-invariant per request; subsequent layers only overwrite that prefix.
     combined_indices: Optional[torch.Tensor] = None
     combined_lens: Optional[torch.Tensor] = None
+    # Valid only for this forward's gather layout. Each ratio has its own
+    # workspace; its compressed prefix survives consumer layers' SWA writes.
+    dequantized_source: Optional[tuple[int, int]] = None
 
 
 @dataclass
 class SparsePrefillChunkCache:
     """Cache prefill-chunk metadata shared across layers.
 
-    Fields depend on request/token mappings and compressed page tables, not
+    Gather layouts depend on request/token mappings and compressed page tables.
-    per-layer k_cache; per-layer top-k combinations are recomputed into reused
+    Shared-source dequantization keys live only for this forward; per-layer
+    top-k combinations are recomputed into reused
     buffers.
     """
 

python/sglang/srt/layers/cp/bcg.py

@@ -23,17 +23,22 @@ import torch
 
 from sglang.srt.arg_groups.overrides import (
     attention_backends_of,
+    model_config_of,
     resolved_view,
     resolving_view,
 )
+from sglang.srt.configs.model_config import is_deepseek_v4
 from sglang.srt.layers.cp.base import get_cp_strategy
+from sglang.srt.layers.cp.interleave import InterleaveCPStrategy
 from sglang.srt.layers.cp.padding import get_cp_padding_align_size
 from sglang.srt.layers.cp.utils import (
     cp_gather_after_forward,
+    cp_shard_hidden_states,
     cp_split_before_forward,
     prepare_cp_forward,
 )
 from sglang.srt.layers.cp.zigzag import ZigzagCPStrategy
+from sglang.srt.layers.logits_processor import LogitsMetadata
 from sglang.srt.model_executor.forward_batch_info import PPProxyTensors
 
 if TYPE_CHECKING:
@@ -50,12 +55,18 @@ def supports_prefill_cp_bcg(server_args: ServerArgs) -> bool:
     cfg = resolving_view(server_args)
     resolved = resolved_view(server_args)
     prefill_attention_backend, _ = attention_backends_of(resolved_view(server_args))
+    supports_layout = (
+        cfg.cp_strategy == "zigzag" and prefill_attention_backend == "trtllm_mha"
+    ) or (
+        cfg.cp_strategy == "interleave"
+        and prefill_attention_backend == "dsv4"
+        and is_deepseek_v4(model_config_of(server_args).hf_config)
+    )
     return (
         cfg.enable_prefill_cp
         and cfg.pp_size == 1
         and resolved.attn_cp_size == cfg.tp_size
-        and cfg.cp_strategy == "zigzag"
+        and supports_layout
-        and prefill_attention_backend == "trtllm_mha"
     )
 
 
@@ -67,8 +78,12 @@ def enable_cp_bcg_capture(server_args: ServerArgs) -> bool:
 def filter_prefill_cp_bcg_capture_num_tokens(
     capture_num_tokens: list[int], server_args: ServerArgs
 ) -> list[int]:
-    """Keep only token buckets where the zigzag CP strategy can run."""
+    """Keep only token buckets where the configured CP strategy can run."""
-    min_num_tokens = resolved_view(server_args).attn_cp_size * 2
+    cfg = resolving_view(server_args)
+    cp_segments_per_token_block = 2 if cfg.cp_strategy == "zigzag" else 1
+    min_num_tokens = (
+        resolved_view(server_args).attn_cp_size * cp_segments_per_token_block
+    )
     filtered = [size for size in capture_num_tokens if size >= min_num_tokens]
     if not filtered:
         raise ValueError(
@@ -96,6 +111,8 @@ class PrefillCPBCGInput:
 
     input_embeds: torch.Tensor
     positions: torch.Tensor
+    input_ids: Optional[torch.Tensor] = None
+    num_token_non_padded: Optional[torch.Tensor] = None
     bucket_local_tokens: Dict[int, int] = field(default_factory=dict)
     live_local_tokens: int = 0
 
@@ -114,12 +131,22 @@ class PrefillCPBCGInput:
                     (runner.max_num_tokens,),
                     dtype=torch.int64,
                 ),
+                input_ids=torch.zeros((runner.max_num_tokens,), dtype=torch.int64),
+                num_token_non_padded=torch.zeros((), dtype=torch.int32),
             )
 
     def required_local_tokens(self, extend_seq_lens: Any) -> Optional[int]:
-        """Return the aligned CP-local rows required by a live zigzag layout."""
+        """Return the aligned CP-local rows required by the active layout."""
         strategy = get_cp_strategy()
-        if not isinstance(strategy, ZigzagCPStrategy) or extend_seq_lens is None:
+        if extend_seq_lens is None:
+            return None
+        if isinstance(strategy, InterleaveCPStrategy):
+            logical_tokens = (
+                sum(int(length) for length in extend_seq_lens) + strategy.cp_size - 1
+            ) // strategy.cp_size
+            align_size = get_cp_padding_align_size()
+            return (logical_tokens + align_size - 1) // align_size * align_size
+        if not isinstance(strategy, ZigzagCPStrategy):
             return None
 
         cp_segment_num = strategy.cp_size * 2
@@ -219,6 +246,7 @@ class PrefillCPBCGInput:
         raw_tokens = int(forward_batch.extend_num_tokens)
         global_input_ids = forward_batch.input_ids[:raw_tokens]
         global_positions = forward_batch.positions[:raw_tokens]
+        local_input_ids = cp_shard_hidden_states(global_input_ids, forward_batch)
         global_input_embeds = runner.model_runner.model.get_input_embeddings()(
             global_input_ids
         )
@@ -249,12 +277,31 @@ class PrefillCPBCGInput:
 
         input_embeds = self.input_embeds[:captured_local_tokens]
         positions = self.positions[:captured_local_tokens]
+        assert self.input_ids is not None
+        input_ids = self.input_ids[:captured_local_tokens]
         input_embeds.zero_()
         positions.zero_()
+        input_ids.zero_()
         input_embeds[:live_local_tokens].copy_(local_input_embeds)
         positions[:live_local_tokens].copy_(local_positions)
+        input_ids[:live_local_tokens].copy_(local_input_ids)
         forward_batch.input_embeds = input_embeds
-        forward_batch.positions = positions
+        forward_batch._cp_positions = positions
+        # Keep the global input_ids field intact: the runner uses its length to
+        # select the global capture bucket. The DSV4 body consumes this fixed,
+        # rank-local view for hash routing and MegaMoE.
+        forward_batch._cp_input_ids = input_ids
+        forward_batch.input_ids_global = input_ids
+        if forward_batch.num_token_non_padded is not None:
+            assert self.num_token_non_padded is not None
+            metadata = forward_batch.attn_cp_metadata
+            logical_tokens = (
+                metadata.per_rank_logical_token or metadata.per_rank_actual_token
+            )
+            strategy = get_cp_strategy()
+            assert strategy is not None
+            self.num_token_non_padded.fill_(logical_tokens[strategy.cp_rank])
+            forward_batch.num_token_non_padded = self.num_token_non_padded
         self.live_local_tokens = live_local_tokens
 
 
@@ -307,10 +354,50 @@ def execute_prefill_cp_bcg(
             static_forward_batch,
             torch.cuda.current_stream(),
         )
-        return model.logits_processor(
+        if aux_hidden_states is not None:
-            forward_batch.input_ids,
+            if torch.is_tensor(aux_hidden_states):
+                aux_hidden_states = cp_gather_after_forward(
+                    aux_hidden_states, static_forward_batch, torch.cuda.current_stream()
+                )
+            else:
+                aux_hidden_states = [
+                    cp_gather_after_forward(
+                        aux, static_forward_batch, torch.cuda.current_stream()
+                    )
+                    for aux in aux_hidden_states
+                ]
+        hidden_states_before_norm = None
+        if isinstance(hidden_states, tuple):
+            assert len(hidden_states) == 2
+            hidden_states, hidden_states_before_norm = hidden_states
+
+        input_ids = forward_batch.input_ids
+        logits_metadata = forward_batch
+        tail = None
+        language_model = getattr(model, "model", None)
+        if (
+            capture_aux_hidden_states
+            and getattr(language_model, "late_layer_start", None) is not None
+            and forward_batch.forward_mode.is_extend_without_speculative()
+        ):
+            tail_metadata = runner.model_runner.attn_backend.tail_forward_metadata
+            tail = tail_metadata.late_layer_tail
+            input_ids = tail.rows(input_ids)
+            logits_metadata = LogitsMetadata.from_forward_batch(forward_batch)
+            logits_metadata.extend_seq_lens = tail.extend_seq_lens
+            logits_metadata.extend_seq_lens_cpu = tail.extend_seq_lens_cpu
+            logits_metadata.extend_logprob_start_lens_cpu = tail.extend_seq_lens_cpu
+
+        output = model.logits_processor(
+            input_ids,
             hidden_states,
             model.lm_head,
-            forward_batch,
+            logits_metadata,
             aux_hidden_states,
+            hidden_states_before_norm=(
+                None if aux_hidden_states is not None else hidden_states_before_norm
+            ),
         )
+        if tail is not None:
+            output.hidden_states_token_indices = tail.token_indices
+        return output

python/sglang/srt/layers/moe/mega_moe.py

@@ -216,20 +216,33 @@ def _run_mega_routed(
 
     if num_tokens > 0:
         router_logits = moe.gate(hidden_states, forward_batch=forward_batch)
-        topk_kwargs = {"input_ids": input_ids_global} if moe.is_hash else {}
+        num_token_non_padded = (
-        topk_output = moe.topk(
+            forward_batch.num_token_non_padded if forward_batch is not None else None
-            hidden_states,
-            router_logits,
-            num_token_non_padded=(
-                forward_batch.num_token_non_padded
-                if forward_batch is not None
-                else None
-            ),
-            expert_location_dispatch_info=ExpertLocationDispatchInfo.init_new(
-                layer_id=moe.layer_id,
-            ),
-            **topk_kwargs,
         )
+        if isinstance(
+            getattr(moe.gate, "e_score_correction_bias_vl", None), torch.Tensor
+        ):
+            # V4.1 uses a different correction bias for image-token rows. The
+            # MegaMoE transport consumes the same routed ids/weights as TopK.
+            from sglang.srt.multimodal.dsv41.vl_routing import vision_topk
+
+            topk_output = vision_topk(
+                moe,
+                router_logits,
+                input_ids_global,
+                num_token_non_padded=num_token_non_padded,
+            )
+        else:
+            topk_kwargs = {"input_ids": input_ids_global} if moe.is_hash else {}
+            topk_output = moe.topk(
+                hidden_states,
+                router_logits,
+                num_token_non_padded=num_token_non_padded,
+                expert_location_dispatch_info=ExpertLocationDispatchInfo.init_new(
+                    layer_id=moe.layer_id,
+                ),
+                **topk_kwargs,
+            )
         topk_ids = topk_output.topk_ids
         topk_weights = topk_output.topk_weights
     else:

python/sglang/srt/mem_cache/common.py

@@ -155,6 +155,10 @@ def free_kv_row_segments(
 
 def maybe_cache_unfinished_req(req: Req, tree_cache: BasePrefixCache, **kwargs):
     if getattr(req, "skip_radix_cache_insert", False):
+        kv_indices = tree_cache.req_to_token_pool.req_to_token[
+            req.kv.req_pool_idx, : len(req.get_fill_ids())
+        ]
+        req.prefix_indices = kv_indices.to(dtype=torch.int64, copy=True)
         return
 
     tree_cache.cache_unfinished_req(req, **kwargs)

python/sglang/srt/model_executor/runner/eager_runner.py

@@ -385,14 +385,22 @@ class EagerRunner(BaseRunner):
         """
         model = self.model_runner.model
 
+        input_ids = forward_batch.input_ids
         input_embeds = kwargs.get("input_embeds")
+        # Multimodal spans must be embedded in global token order, before CP
+        # slicing. The model may also normalize image hash IDs for its router.
+        prepare_inputs = getattr(model, "prepare_language_model_inputs", None)
+        if prepare_inputs is not None:
+            input_ids, input_embeds = prepare_inputs(
+                input_ids, forward_batch, input_embeds
+            )
         if input_embeds is None:
-            input_embeds = model.get_input_embeddings()(forward_batch.input_ids)
+            input_embeds = model.get_input_embeddings()(input_ids)
         with cp_shard_model_inputs(
             input_embeds,
             forward_batch.positions,
             forward_batch,
-            forward_batch.input_ids,
+            input_ids,
         ) as (sharded_input_embeds, sharded_positions, model_input_ids):
             model_kwargs = {"input_embeds": sharded_input_embeds}
             if (pp_proxy_tensors := kwargs.get("pp_proxy_tensors")) is not None:
@@ -437,7 +445,7 @@ class EagerRunner(BaseRunner):
             if aux_hidden_states is None:
                 logits_kwargs["hidden_states_before_norm"] = hidden_states_before_norm
         return model.logits_processor(
-            forward_batch.input_ids,
+            input_ids,
             hidden_states,
             model.lm_head,
             forward_batch,

python/sglang/srt/model_executor/runner/prefill_cuda_graph_runner.py

@@ -701,6 +701,9 @@ class PrefillCudaGraphRunner(BaseCudaGraphRunner):
 
     def _get_layer_model_positions(self, forward_batch: ForwardBatch) -> torch.Tensor:
         """Mirror outer multimodal wrappers when BCG captures layer_model directly."""
+        cp_positions = getattr(forward_batch, "_cp_positions", None)
+        if cp_positions is not None:
+            return cp_positions
         if forward_batch.mrope_positions is None:
             return forward_batch.positions
 
@@ -782,7 +785,9 @@ class PrefillCudaGraphRunner(BaseCudaGraphRunner):
             if self._uses_eager_prefill_tail():
                 # BCG / Full: capture the transformer body only.
                 positions = self._get_layer_model_positions(forward_batch)
-                input_ids = forward_batch.input_ids
+                input_ids = getattr(
+                    forward_batch, "_cp_input_ids", forward_batch.input_ids
+                )
                 kwargs = _build_layer_model_forward_kwargs(
                     self.layer_model, forward_batch, pp_proxy_tensors
                 )
@@ -1336,9 +1341,9 @@ class PrefillCudaGraphRunner(BaseCudaGraphRunner):
             batch_max_context_len=batch_max_context_len,
         ):
             return False
-        if getattr(self, "enable_cp_bcg_capture", False) and is_cp_active(
+        if getattr(self, "enable_cp_bcg_capture", False):
-            forward_batch
+            if not is_cp_active(forward_batch):
-        ):
+                return False
             assert self.prefill_cp_bcg_input is not None
             if (
                 self.prefill_cp_bcg_input.select_replay_bucket_for_batch(

python/sglang/srt/models/deepseek_v4.py

@@ -182,6 +182,7 @@ from sglang.srt.multimodal.deepseek_v41_image_processing import (
 )
 from sglang.srt.runtime_context import (
     get_device,
+    get_disagg,
     get_exec,
     get_forward,
     get_parallel,
@@ -2039,7 +2040,10 @@ class MQALayer(MqaAttentionBase):
             if (
                 forward_batch.forward_mode.is_extend()
                 and is_in_breakable_cuda_graph()
-                and not getattr(attn_backend, "low_ratio_prefill_graph", False)
+                and (
+                    dsa_use_prefill_cp(forward_batch)
+                    or not getattr(attn_backend, "low_ratio_prefill_graph", False)
+                )
             ):
                 bcg_deepseek_v4_low_ratio_sources(self, x, q_lora, positions)
             else:
@@ -2650,7 +2654,8 @@ class DeepseekV4DecoderLayer(nn.Module):
             is_nextn=is_nextn,
             is_deepseek_v4=True,
             vl_correction_bias=config.model_type == "deepseek_v41"
-            and config.vision_n_layers > 0,
+            and config.vision_n_layers > 0
+            and not getattr(config, "language_model_only", False),
         )
 
         self.input_layernorm = RMSNorm(config.hidden_size, eps=config.rms_norm_eps)
@@ -3872,7 +3877,15 @@ class DeepseekV4DecoderLayer(nn.Module):
         finally:
             forward_batch.num_token_non_padded = saved_num_token_non_padded
         if _use_cp and get_moe_a2a_backend().is_none():
-            hidden_states = dsa_cp_reduce_scatter_hidden_states(hidden_states)
+            if self.config.model_type == "deepseek_v41":
+                parallel = get_parallel()
+                hidden_states = parallel.tp_group.all_reduce(hidden_states)
+                parallel = get_parallel()
+                hidden_states = hidden_states.tensor_split(parallel.attn_cp_size)[
+                    parallel.attn_cp_rank
+                ].contiguous()
+            else:
+                hidden_states = dsa_cp_reduce_scatter_hidden_states(hidden_states)
         elif _use_tp_moe_gather:
             hidden_states, global_hidden_states = (
                 get_local_dp_buffer(get_parallel().tp_group),
@@ -4402,11 +4415,18 @@ class DeepseekV4Model(nn.Module):
         )
         if self.engram_hasher is not None:
             if cp_extend:
-                # n-gram hashing needs each token's predecessors: hash the whole prompt
+                # N-gram hashing needs each token's predecessors, so hash the
+                # whole prompt before selecting this CP rank's interleaved rows.
+                # The hasher builds request-to-token indices dynamically; keep
+                # that work at an eager break during breakable graph capture.
                 total = int(forward_batch.attn_cp_metadata.total_seq_lens)
-                hash_ids = self.engram_hasher(
+                global_input_ids = forward_batch.input_ids[:total]
-                    forward_batch.input_ids[:total], forward_batch
+                if is_in_breakable_cuda_graph():
-                )
+                    hash_ids = bcg_deepseek_v4_engram_hash_ids(
+                        self.engram_hasher, global_input_ids
+                    )
+                else:
+                    hash_ids = self.engram_hasher(global_input_ids, forward_batch)
                 parallel = get_parallel()
                 hash_ids = hash_ids[parallel.attn_cp_rank :: parallel.attn_cp_size]
                 pad_rows = hidden_states.shape[0] - hash_ids.shape[0]
@@ -4839,6 +4859,13 @@ class DeepseekV4Model(nn.Module):
         return hidden_states, pre_hc_head
 
 
+def _v41_vision_a2a_supported() -> bool:
+    backend = get_moe_a2a_backend()
+    return backend.is_none() or (
+        backend.is_megamoe() and get_disagg().disaggregation_mode == "decode"
+    )
+
+
 class DeepseekV4ForCausalLM(nn.Module):
     supports_cuda_vmm_feature_transport = True
 
@@ -4864,14 +4891,19 @@ class DeepseekV4ForCausalLM(nn.Module):
         self.wo_a_fp8 = wo_a_fp8_gemm_enabled(quant_config)
         self.determine_num_fused_shared_experts()
         self.vision = None
-        if config.model_type == "deepseek_v41" and config.vision_n_layers > 0:
+        if (
+            config.model_type == "deepseek_v41"
+            and config.vision_n_layers > 0
+            and not getattr(config, "language_model_only", False)
+        ):
             if (
                 get_parallel().attn_cp_size != 1
-                or get_parallel().pp_group.world_size != 1
+                or get_pp_group().world_size != 1
-                or not get_moe_a2a_backend().is_none()
+                or not _v41_vision_a2a_supported()
             ):
                 raise ValueError(
-                    "V4.1 vision currently supports TP/EP/DP without CP, PP or MoE A2A"
+                    "V4.1 vision supports TP/EP/DP without CP or PP; "
+                    "MoE A2A is supported only with MegaMoE on a PD decode node"
                 )
 
             args = SimpleNamespace(**vars(config), dim=config.hidden_size)
@@ -5078,16 +5110,19 @@ class DeepseekV4ForCausalLM(nn.Module):
             0 if is_shared_experts_fusion_disabled() else self.config.n_shared_experts
         )
 
-    def forward(
+    def prepare_language_model_inputs(
         self,
         input_ids: torch.Tensor,
-        positions: torch.Tensor,
         forward_batch: ForwardBatch,
         input_embeds: Optional[torch.Tensor] = None,
-        pp_proxy_tensors: Optional[PPProxyTensors] = None,
+    ) -> Tuple[torch.Tensor, Optional[torch.Tensor]]:
-    ) -> torch.Tensor:
+        """Prepare full-sequence image embeddings and model IDs before CP splits.
+
+        Scheduler hash IDs stay intact for multimodal cache keys; the language
+        model uses image_token_id for Engram masking and visual MoE routing.
+        """
         if (
-            self.vision is not None
+            getattr(self, "vision", None) is not None
             and not forward_batch.forward_mode.is_decode()
             and not forward_batch.forward_mode.is_target_verify()
             and forward_batch.mm_inputs is not None
@@ -5096,7 +5131,7 @@ class DeepseekV4ForCausalLM(nn.Module):
             if input_embeds is not None:
                 raise ValueError("Cannot combine input_embeds and image inputs")
             input_embeds = self._prepare_mm_embeddings(input_ids, forward_batch)
-        if self.vision is not None and not (
+        if getattr(self, "vision", None) is not None and not (
             forward_batch.forward_mode.is_decode_or_idle()
             or forward_batch.forward_mode.is_target_verify()
         ):
@@ -5106,6 +5141,19 @@ class DeepseekV4ForCausalLM(nn.Module):
                 input_ids >= MM_PAD_SHIFT_VALUE, self.config.image_token_id
             )
 
+        return input_ids, input_embeds
+
+    def forward(
+        self,
+        input_ids: torch.Tensor,
+        positions: torch.Tensor,
+        forward_batch: ForwardBatch,
+        input_embeds: Optional[torch.Tensor] = None,
+        pp_proxy_tensors: Optional[PPProxyTensors] = None,
+    ) -> torch.Tensor:
+        input_ids, input_embeds = self.prepare_language_model_inputs(
+            input_ids, forward_batch, input_embeds
+        )
         with get_attn_tp_context().maybe_input_scattered(forward_batch):
             hidden_states = self.model.forward(
                 input_ids, positions, forward_batch, input_embeds, pp_proxy_tensors

python/sglang/srt/server_args.py

@@ -376,6 +376,7 @@ class ServerArgs:
     # ===== END TO BE REFACTORED ====
 
     LANGUAGE_MODEL_ONLY_ARCHITECTURES = (
+        "DeepseekV4ForCausalLM",
         "MuseGlimmerForConditionalGeneration",
         "Cosmos3ForConditionalGeneration",
         "Cosmos3EdgeForConditionalGeneration",

python/sglang/test/dsv41_cp_test_utils.py

@@ -0,0 +1,79 @@
+"""Small CPU tensors; production CP slicing/gather, mocked collective transport."""
+
+from contextlib import ExitStack, contextmanager, nullcontext
+from types import SimpleNamespace as NS
+from unittest.mock import patch
+
+import torch
+
+from sglang.srt.layers.cp.interleave import InterleaveCPStrategy
+from sglang.srt.layers.cp.padding import pad_logical_token_to_physical
+from sglang.srt.model_executor.forward_batch_info import ForwardMode
+
+CP = "sglang.srt.layers.cp"
+
+
+@contextmanager
+def cp_context(size, rank, lengths=(3, 6), prefix_lengths=(7, 13)):
+    """Keep real interleave indexing/padding; replace only runtime context."""
+    strategy = InterleaveCPStrategy(size)
+    parallel = NS(attn_cp_size=size, attn_cp_rank=rank, attn_cp_group=None)
+    batch = NS(
+        forward_mode=ForwardMode.EXTEND,
+        input_ids=torch.arange(1, sum(lengths) + 1),
+        positions=torch.cat(
+            [
+                torch.arange(prefix, prefix + length)
+                for prefix, length in zip(prefix_lengths, lengths)
+            ]
+        ),
+        extend_seq_lens_cpu=list(lengths),
+        extend_prefix_lens_cpu=list(prefix_lengths),
+        mm_inputs=None,
+        spec_info=None,
+    )
+    batch.attn_cp_metadata = strategy.build_metadata(
+        sum(lengths), [p + n for p, n in zip(prefix_lengths, lengths)], list(lengths)
+    )
+    with ExitStack() as stack:
+        for module in ("base", "utils", "padding", "interleave"):
+            stack.enter_context(
+                patch(CP + "." + module + ".get_parallel", return_value=parallel)
+            )
+        stack.enter_context(patch(CP + ".utils.get_cp_strategy", return_value=strategy))
+        stack.enter_context(
+            patch(CP + ".padding.get_cp_padding_align_size", return_value=size)
+        )
+        stack.enter_context(
+            patch(
+                CP + ".utils.get_moe_a2a_backend", return_value=NS(is_none=lambda: True)
+            )
+        )
+        pad_logical_token_to_physical(batch.attn_cp_metadata)
+        yield strategy, batch
+
+
+@contextmanager
+def simulated_collective(strategy, batch, global_tensor):
+    """Inject peer buffers into all-gather; retain production unpadding/reordering."""
+    physical = max(batch.attn_cp_metadata.per_rank_actual_token)
+    buffers = []
+    for rank in range(strategy.cp_size):
+        buf = global_tensor.new_zeros((physical, *global_tensor.shape[1:]))
+        local = global_tensor[rank :: strategy.cp_size]
+        buf[: len(local)] = local
+        buffers.append(buf)
+
+    def gather(output, local):
+        torch.testing.assert_close(local, buffers[strategy.cp_rank], rtol=0, atol=0)
+        output.copy_(torch.cat(buffers))
+
+    with (
+        patch(
+            CP + ".interleave.use_symmetric_memory",
+            side_effect=lambda *a, **k: nullcontext(),
+        ),
+        patch(CP + ".interleave.is_allocation_symmetric", return_value=False),
+        patch(CP + ".interleave.attn_cp_all_gather_into_tensor", side_effect=gather),
+    ):
+        yield

test/registered/attention/test_aiter_gluon_h12_fp8.py

@@ -15,7 +15,7 @@ from sglang.srt.layers.attention import aiter_mla_gluon as mod
 from sglang.test.ci.ci_register import register_cpu_ci
 from sglang.test.test_utils import CustomTestCase
 
-register_cpu_ci(est_time=8, suite="base-a-test-cpu")
+register_cpu_ci(est_time=15, suite="base-a-test-cpu")
 
 _GLUON_FN = "sglang.srt.layers.attention.aiter_mla_gluon._gluon_fn"
 

test/registered/attention/test_chunk_gated_delta_rule.py

@@ -13,7 +13,7 @@ from sglang.test.ci.ci_register import (
     register_xpu_ci,
 )
 
-register_cuda_ci(est_time=7, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=6, stage="base-b", runner_config="1-gpu-large")
 register_amd_ci(est_time=11, stage="stage-b", runner_config="1-gpu-large-amd")
 register_xpu_ci(est_time=900, suite="stage-b-test-1-gpu-xpu")
 

test/registered/attention/test_create_kvindices.py

@@ -9,7 +9,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci
 from sglang.test.test_utils import CustomTestCase
 
 # Triton kernel unit test for KV indices creation
-register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-small")
+register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-small")
 register_amd_ci(est_time=10, suite="stage-b-test-1-gpu-small-amd")
 
 

test/registered/attention/test_deterministic.py

@@ -24,7 +24,7 @@ from sglang.test.test_utils import (
     is_in_amd_ci,
 )
 
-register_cuda_ci(est_time=272, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=280, stage="base-b", runner_config="1-gpu-large")
 register_amd_ci(est_time=278, suite="stage-b-test-1-gpu-small-amd")
 register_xpu_ci(est_time=207, suite="stage-b-test-1-gpu-xpu")
 

test/registered/attention/test_flash_attention_4.py

@@ -13,7 +13,7 @@ from sglang.test.test_utils import (
 )
 
 # FlashAttention4 integration test (requires SM 100+ / Blackwell B200)
-register_cuda_ci(est_time=220, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=230, stage="base-b", runner_config="4-gpu-b200")
 
 
 @unittest.skipIf(get_device_sm() < 100, "Test requires CUDA SM 100 or higher")

test/registered/attention/test_hybrid_attn_backend.py

@@ -11,7 +11,7 @@ from sglang.test.test_utils import (
 
 # Hybrid attention backend tests (FA3 prefill + FlashInfer decode, requires SM 90+ / H100)
 # Multiple test classes: base, MLA, TorchCompile, SpecDecode variants
-register_cuda_ci(est_time=368, stage="extra-a", runner_config="1-gpu-large")
+register_cuda_ci(est_time=393, stage="extra-a", runner_config="1-gpu-large")
 
 
 class TestHybridAttnBackendMLA(TestHybridAttnBackendBase):

test/registered/attention/test_kda_kernels.py

@@ -19,7 +19,7 @@ from sglang.srt.utils.common import get_device
 from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=32, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=25, stage="base-b", runner_config="1-gpu-large")
 register_amd_ci(est_time=12, stage="stage-b", runner_config="1-gpu-large-amd")
 
 

test/registered/attention/test_normal_decode_set_metadata.py

@@ -21,7 +21,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci
 from sglang.test.test_utils import CustomTestCase
 
 # Register this test for CUDA CI in base-b (fast attention/kernel tests)
-register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large")
 register_amd_ci(est_time=17, suite="stage-b-test-1-gpu-large-amd")
 
 

test/registered/attention/test_qwen35_deterministic.py

@@ -12,7 +12,7 @@ from sglang.test.test_deterministic_utils import (
     TestDeterministicBase,
 )
 
-register_cuda_ci(est_time=135, stage="extra-b", runner_config="4-gpu-h100")
+register_cuda_ci(est_time=119, stage="extra-b", runner_config="4-gpu-h100")
 
 QWEN35 = "Qwen/Qwen3.5-35B-A3B"
 

test/registered/attention/test_torch_native_attention_backend.py

@@ -18,7 +18,7 @@ from sglang.test.test_utils import (
 )
 
 # Torch native attention backend integration test with MMLU eval
-register_cuda_ci(est_time=312, stage="extra-a", runner_config="1-gpu-small")
+register_cuda_ci(est_time=310, stage="extra-a", runner_config="1-gpu-small")
 register_amd_ci(est_time=150, suite="stage-b-test-1-gpu-small-amd")
 
 

test/registered/attention/test_triton_sliding_window.py

@@ -15,7 +15,7 @@ from sglang.test.test_utils import (
 )
 
 # Sliding window attention with Triton backend (Gemma-3 model)
-register_cuda_ci(est_time=81, stage="extra-a", runner_config="1-gpu-large")
+register_cuda_ci(est_time=80, stage="extra-a", runner_config="1-gpu-large")
 register_amd_ci(est_time=200, suite="stage-b-test-1-gpu-small-amd")
 
 

test/registered/attention/test_trtllm_mha_encoder_only.py

@@ -17,7 +17,7 @@ from sglang.test.ci.ci_register import register_cuda_ci
 from sglang.test.test_utils import CustomTestCase
 
 # trtllm_mha kernels are sm100-only; run this kernel-unit test on Blackwell.
-register_cuda_ci(est_time=8, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=10, stage="base-b", runner_config="4-gpu-b200")
 
 DEVICE = "cuda"
 PAGE_SIZE = 32

test/registered/attention/test_trtllm_mha_graph_metadata.py

@@ -23,7 +23,7 @@ from sglang.srt.model_executor.forward_batch_info import ForwardMode
 from sglang.test.ci.ci_register import register_cuda_ci
 
 # trtllm_mha kernels are sm100-only; run this kernel-unit test on Blackwell.
-register_cuda_ci(est_time=12, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=16, stage="base-b", runner_config="4-gpu-b200")
 
 DEVICE = "cuda"
 PAGE_SIZE = 128

test/registered/attention/test_trtllm_mha_page_table.py

@@ -25,7 +25,7 @@ from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci
 from sglang.test.test_utils import CustomTestCase
 
 # Triton kernel unit test for the trtllm_mha device-side page-table build.
-register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-small")
+register_cuda_ci(est_time=12, stage="base-b", runner_config="1-gpu-small")
 register_amd_ci(est_time=14, stage="stage-b", runner_config="1-gpu-small-amd")
 
 

test/registered/attention/test_unified_memory_deterministic.py

@@ -27,7 +27,7 @@ from sglang.test.test_utils import (
     popen_launch_server,
 )
 
-register_cuda_ci(est_time=90, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=71, stage="base-b", runner_config="1-gpu-large")
 
 KIMI_LINEAR_MODEL = "yujiepan/kimi-linear-tiny-random"
 SERVER_ENV = {"SGLANG_BATCH_INVARIANT_OPS_ENABLE_MM_DEEPGEMM": "0"}

test/registered/attention/unittests/dense/test_extend_init_contract.py

@@ -36,7 +36,7 @@ from sglang.test.kits.attention_unittest.attention_methods.dense_attention impor
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=10, stage="base-a", runner_config="1-gpu-small")
+register_cuda_ci(est_time=11, stage="base-a", runner_config="1-gpu-small")
 
 _EXTEND_CASE = DenseAttentionCase(
     name="extend_no_prefix_smoke",

test/registered/attention/unittests/dense/test_fa3.py

@@ -30,7 +30,7 @@ from sglang.test.kits.attention_unittest.runner_modes.split_op_runner import (
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=18, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=16, stage="base-b", runner_config="1-gpu-large")
 
 
 @unittest.skipIf(not torch.cuda.is_available(), "CUDA is required")

test/registered/attention/unittests/dense/test_fa4.py

@@ -29,8 +29,8 @@ from sglang.test.kits.attention_unittest.runner_modes.split_op_runner import (
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=23, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=27, stage="base-b", runner_config="4-gpu-b200")
-register_cuda_ci(est_time=15, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=22, stage="base-b", runner_config="1-gpu-large")
 
 
 @unittest.skipIf(not torch.cuda.is_available(), "CUDA is required")

test/registered/attention/unittests/dense/test_flashinfer.py

@@ -26,8 +26,8 @@ from sglang.test.kits.attention_unittest.runner_modes.split_op_runner import (
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=16, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=18, stage="base-b", runner_config="4-gpu-b200")
-register_cuda_ci(est_time=26, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=17, stage="base-b", runner_config="1-gpu-large")
 
 
 @unittest.skipIf(

test/registered/attention/unittests/dense/test_flex_attention.py

@@ -14,8 +14,8 @@ from sglang.test.kits.attention_unittest.runner_modes.split_op_runner import (
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=16, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=18, stage="base-b", runner_config="4-gpu-b200")
-register_cuda_ci(est_time=36, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=17, stage="base-b", runner_config="1-gpu-large")
 
 
 @unittest.skipIf(not torch.cuda.is_available(), "CUDA is required")

test/registered/attention/unittests/dense/test_hybrid_attn.py

@@ -18,8 +18,8 @@ from sglang.test.kits.attention_unittest.attention_methods.dense_attention impor
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=10, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=11, stage="base-b", runner_config="4-gpu-b200")
-register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large")
 
 
 @unittest.skipIf(

test/registered/attention/unittests/dense/test_tbo.py

@@ -21,8 +21,8 @@ from sglang.test.kits.attention_unittest.runner_modes.speculative_target_verify_
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=10, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=11, stage="base-b", runner_config="4-gpu-b200")
-register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=12, stage="base-b", runner_config="1-gpu-large")
 
 
 @unittest.skipIf(not torch.cuda.is_available(), "CUDA is required")

test/registered/attention/unittests/dense/test_torch_native.py

@@ -11,8 +11,8 @@ from sglang.test.kits.attention_unittest.attention_methods.dense_attention impor
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=11, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=12, stage="base-b", runner_config="4-gpu-b200")
-register_cuda_ci(est_time=12, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large")
 register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd")
 
 

test/registered/attention/unittests/dense/test_triton.py

@@ -30,8 +30,8 @@ from sglang.test.kits.attention_unittest.runner_modes.split_op_runner import (
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=38, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=18, stage="base-b", runner_config="4-gpu-b200")
-register_cuda_ci(est_time=36, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=19, stage="base-b", runner_config="1-gpu-large")
 register_amd_ci(est_time=25, suite="stage-b-test-1-gpu-large-amd")
 
 

test/registered/attention/unittests/dense/test_trtllm_mha.py

@@ -28,8 +28,8 @@ from sglang.test.kits.attention_unittest.runner_modes.split_op_runner import (
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=15, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=17, stage="base-b", runner_config="4-gpu-b200")
-register_cuda_ci(est_time=18, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=16, stage="base-b", runner_config="1-gpu-large")
 
 
 @unittest.skipIf(

test/registered/attention/unittests/dsa/test_dsa.py

@@ -30,8 +30,8 @@ from sglang.test.kits.attention_unittest.runner_modes.speculative_draft_runner i
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=22, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=14, stage="base-b", runner_config="4-gpu-b200")
-register_cuda_ci(est_time=26, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=14, stage="base-b", runner_config="1-gpu-large")
 
 
 @unittest.skipIf(not torch.cuda.is_available(), "CUDA is required")

test/registered/attention/unittests/dsv4/test_deepseek_v4.py

@@ -48,7 +48,7 @@ from sglang.test.kits.attention_unittest.runner_modes.speculative_target_verify_
 )
 
 register_cuda_ci(est_time=14, stage="base-b", runner_config="4-gpu-b200")
-register_cuda_ci(est_time=16, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=13, stage="base-b", runner_config="1-gpu-large")
 
 
 @unittest.skipIf(not torch.cuda.is_available(), "CUDA is required")

test/registered/attention/unittests/gdn/test_flashinfer.py

@@ -25,7 +25,7 @@ from sglang.test.kits.attention_unittest.runner_modes.split_op_runner import (
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=14, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=15, stage="base-b", runner_config="4-gpu-b200")
 register_cuda_ci(est_time=14, stage="base-b", runner_config="1-gpu-large")
 
 _cuda_major = int(torch.version.cuda.split(".")[0]) if torch.version.cuda else 0

test/registered/attention/unittests/gdn/test_linear_replayssm_decode.py

@@ -45,7 +45,7 @@ import torch
 from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large")
 register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd")
 
 

test/registered/attention/unittests/gdn/test_torch_native.py

@@ -15,7 +15,7 @@ from sglang.test.kits.attention_unittest.runner_modes.split_op_runner import (
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=10, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=12, stage="base-b", runner_config="4-gpu-b200")
 register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large")
 register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd")
 

test/registered/attention/unittests/gdn/test_triton.py

@@ -29,8 +29,8 @@ from sglang.test.kits.attention_unittest.runner_modes.split_op_runner import (
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=11, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=13, stage="base-b", runner_config="4-gpu-b200")
-register_cuda_ci(est_time=12, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large")
 register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd")
 
 

test/registered/attention/unittests/hybrid_linear/test_flashinfer_mla_chunk_metadata.py

@@ -32,7 +32,7 @@ from sglang.test.kits.attention_unittest.attention_methods.mla_attention import
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-large")
 
 
 _KV_LORA_RANK = DEFAULT_KV_LORA_RANK

test/registered/attention/unittests/hybrid_linear/test_kda_fused_accept_indices.py

@@ -24,7 +24,7 @@ from sglang.srt.layers.attention.linear.kernels.kda_flashinfer import (
 from sglang.test.ci.ci_register import register_cuda_ci
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=8, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-large")
 
 
 class TestBuildFusedAcceptIndices(CustomTestCase):

test/registered/attention/unittests/kda/test_triton.py

@@ -22,8 +22,8 @@ from sglang.test.kits.attention_unittest.runner_modes.split_op_runner import (
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=18, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=15, stage="base-b", runner_config="4-gpu-b200")
-register_cuda_ci(est_time=16, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=20, stage="base-b", runner_config="1-gpu-large")
 register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd")
 
 

test/registered/attention/unittests/lightning/test_triton.py

@@ -19,7 +19,7 @@ from sglang.test.kits.attention_unittest.runner_modes.speculative_target_verify_
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=10, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=12, stage="base-b", runner_config="4-gpu-b200")
 register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large")
 register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd")
 

test/registered/attention/unittests/mamba/test_mamba2.py

@@ -32,8 +32,8 @@ from sglang.test.kits.attention_unittest.runner_modes.speculative_target_verify_
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=11, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=13, stage="base-b", runner_config="4-gpu-b200")
-register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=12, stage="base-b", runner_config="1-gpu-large")
 
 
 @unittest.skipIf(not torch.cuda.is_available(), "CUDA is required")

test/registered/attention/unittests/mamba/test_replay_state_indices_validator.py

@@ -7,7 +7,7 @@ from sglang.srt.layers.attention.mamba.replay_state_indices_validator import (
 )
 from sglang.test.ci.ci_register import register_cpu_ci
 
-register_cpu_ci(est_time=5, suite="base-a-test-cpu")
+register_cpu_ci(est_time=6, suite="base-a-test-cpu")
 
 
 class TestReplayStateIndicesValidator(unittest.TestCase):

test/registered/attention/unittests/mla/test_flashinfer.py

@@ -32,7 +32,7 @@ MLA_SHAPE_KWARGS = dict(
 
 from sglang.test.ci.ci_register import register_cuda_ci
 
-register_cuda_ci(est_time=11, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=12, stage="base-b", runner_config="4-gpu-b200")
 register_cuda_ci(est_time=12, stage="base-b", runner_config="1-gpu-large")
 
 

test/registered/attention/unittests/mla/test_tokenspeed_mla.py

@@ -56,7 +56,7 @@ MLA_SHAPE_KWARGS = dict(
 from sglang.test.ci.ci_register import register_cuda_ci
 
 register_cuda_ci(est_time=10, stage="base-b", runner_config="4-gpu-b200")
-register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-large")
 
 
 @unittest.skipIf(not _SUPPORTED, _SKIP_REASON)

test/registered/attention/unittests/mla/test_triton.py

@@ -28,8 +28,8 @@ from sglang.test.kits.attention_unittest.runner_modes.split_op_runner import (
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=17, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=18, stage="base-b", runner_config="4-gpu-b200")
-register_cuda_ci(est_time=16, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=20, stage="base-b", runner_config="1-gpu-large")
 
 
 @unittest.skipIf(not torch.cuda.is_available(), "CUDA is required")

test/registered/attention/unittests/swa/test_flashinfer.py

@@ -24,8 +24,8 @@ from sglang.test.kits.attention_unittest.runner_modes.split_op_runner import (
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=11, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=13, stage="base-b", runner_config="4-gpu-b200")
-register_cuda_ci(est_time=12, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large")
 
 
 @unittest.skipIf(

test/registered/attention/unittests/swa/test_swa_out_cache_loc.py

@@ -24,7 +24,7 @@ from sglang.srt.model_executor.forward_batch_info import ForwardMode
 from sglang.test.ci.ci_register import register_cpu_ci
 from sglang.test.test_utils import CustomTestCase
 
-register_cpu_ci(est_time=8, suite="base-a-test-cpu")
+register_cpu_ci(est_time=11, suite="base-a-test-cpu")
 
 
 class TestSWAKVPoolSetKVBuffer(CustomTestCase):

test/registered/attention/unittests/swa/test_torch_native.py

@@ -12,8 +12,8 @@ from sglang.test.kits.attention_unittest.attention_methods.dense_attention impor
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=10, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=12, stage="base-b", runner_config="4-gpu-b200")
-register_cuda_ci(est_time=12, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large")
 register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd")
 
 

test/registered/attention/unittests/swa/test_triton.py

@@ -23,8 +23,8 @@ from sglang.test.kits.attention_unittest.runner_modes.split_op_runner import (
 )
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=11, stage="base-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=12, stage="base-b", runner_config="4-gpu-b200")
-register_cuda_ci(est_time=12, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=14, stage="base-b", runner_config="1-gpu-large")
 register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd")
 
 

test/registered/backends/test_deepseek_v3_fp4_cutedsl_moe.py

@@ -17,7 +17,7 @@ from sglang.test.test_utils import (
     write_github_step_summary,
 )
 
-register_cuda_ci(est_time=259, stage="extra-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=243, stage="extra-b", runner_config="4-gpu-b200")
 
 FULL_DEEPSEEK_V3_FP4_MODEL_PATH = "nvidia/DeepSeek-V3-0324-FP4"
 SERVER_LAUNCH_TIMEOUT = 1000

test/registered/backends/test_flashinfer_fusion_preflight.py

@@ -11,7 +11,7 @@ from sglang.srt.utils import get_cuda_driver_bindings, is_flashinfer_available
 from sglang.test.ci.ci_register import register_cuda_ci
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=26, stage="base-b", runner_config="2-gpu-large")
+register_cuda_ci(est_time=28, stage="base-b", runner_config="2-gpu-large")
 
 WORLD_SIZE = 2
 

test/registered/backends/test_torch_compile.py

@@ -15,7 +15,7 @@ from sglang.test.test_utils import (
     popen_launch_server,
 )
 
-register_cuda_ci(est_time=131, stage="extra-a", runner_config="1-gpu-large")
+register_cuda_ci(est_time=139, stage="extra-a", runner_config="1-gpu-large")
 register_amd_ci(est_time=1100, suite="stage-b-test-1-gpu-small-amd")
 
 

test/registered/basic_perf/test_eagle3_latency.py

@@ -13,7 +13,7 @@ from sglang.test.test_utils import (
     run_bench_serving,
 )
 
-register_cuda_ci(est_time=136, stage="extra-a", runner_config="1-gpu-large")
+register_cuda_ci(est_time=145, stage="extra-a", runner_config="1-gpu-large")
 
 
 class TestEagle3Latency(CustomTestCase):

test/registered/basic_perf/test_embeddings_api.py

@@ -16,7 +16,7 @@ from sglang.test.test_utils import (
     run_embeddings_benchmark_multi,
 )
 
-register_cuda_ci(est_time=215, stage="extra-a", runner_config="1-gpu-large")
+register_cuda_ci(est_time=245, stage="extra-a", runner_config="1-gpu-large")
 register_amd_ci(est_time=240, suite="stage-b-test-1-gpu-large-amd")
 
 

test/registered/basic_perf/test_lora_latency.py

@@ -14,7 +14,7 @@ from sglang.test.test_utils import (
     run_bench_serving,
 )
 
-register_cuda_ci(est_time=247, stage="extra-a", runner_config="1-gpu-large")
+register_cuda_ci(est_time=490, stage="extra-a", runner_config="1-gpu-large")
 register_amd_ci(est_time=430, suite="stage-b-test-1-gpu-large-amd")
 
 

test/registered/basic_perf/test_moe_throughput.py

@@ -11,7 +11,7 @@ from sglang.test.test_utils import (
     run_bench_serving,
 )
 
-register_cuda_ci(est_time=271, stage="extra-a", runner_config="2-gpu-large")
+register_cuda_ci(est_time=290, stage="extra-a", runner_config="2-gpu-large")
 register_amd_ci(est_time=770, suite="stage-b-test-2-gpu-large-amd")
 
 

test/registered/basic_perf/test_pp_throughput.py

@@ -11,7 +11,7 @@ from sglang.test.test_utils import (
     run_bench_serving,
 )
 
-register_cuda_ci(est_time=325, stage="extra-a", runner_config="2-gpu-large")
+register_cuda_ci(est_time=490, stage="extra-a", runner_config="2-gpu-large")
 register_amd_ci(est_time=1030, suite="stage-b-test-2-gpu-large-amd")
 
 

test/registered/basic_perf/test_score_api.py

@@ -16,7 +16,7 @@ from sglang.test.test_utils import (
     run_score_benchmark_multi,
 )
 
-register_cuda_ci(est_time=185, stage="extra-a", runner_config="1-gpu-large")
+register_cuda_ci(est_time=215, stage="extra-a", runner_config="1-gpu-large")
 register_amd_ci(est_time=210, suite="stage-b-test-1-gpu-large-amd")
 
 

test/registered/basic_perf/test_serving_latency.py

@@ -10,7 +10,7 @@ from sglang.test.test_utils import (
     run_bench_serving,
 )
 
-register_cuda_ci(est_time=182, stage="extra-a", runner_config="1-gpu-large")
+register_cuda_ci(est_time=190, stage="extra-a", runner_config="1-gpu-large")
 register_amd_ci(est_time=165, suite="stage-b-test-1-gpu-large-amd")
 
 

test/registered/basic_perf/test_serving_throughput.py

@@ -11,7 +11,7 @@ from sglang.test.test_utils import (
     run_bench_serving,
 )
 
-register_cuda_ci(est_time=660, stage="extra-a", runner_config="1-gpu-large")
+register_cuda_ci(est_time=710, stage="extra-a", runner_config="1-gpu-large")
 register_amd_ci(est_time=810, suite="stage-b-test-1-gpu-large-amd")
 
 

test/registered/basic_perf/test_torch_compile_throughput.py

@@ -10,7 +10,7 @@ from sglang.test.test_utils import (
     run_bench_offline_throughput,
 )
 
-register_cuda_ci(est_time=85, stage="extra-a", runner_config="2-gpu-large")
+register_cuda_ci(est_time=75, stage="extra-a", runner_config="2-gpu-large")
 register_amd_ci(est_time=280, suite="stage-b-test-2-gpu-large-amd")
 
 

test/registered/basic_perf/test_vlm_serving_fa3.py

@@ -6,7 +6,7 @@ from sglang.test.ci.ci_register import register_cuda_ci
 from sglang.test.kits.vlm_perf_kit import check_vlm_serving_perf
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=152, stage="extra-a", runner_config="1-gpu-large")
+register_cuda_ci(est_time=150, stage="extra-a", runner_config="1-gpu-large")
 
 
 class TestVLMServingFa3(CustomTestCase):

test/registered/basic_perf/test_vlm_serving_flashinfer.py

@@ -6,7 +6,7 @@ from sglang.test.ci.ci_register import register_cuda_ci
 from sglang.test.kits.vlm_perf_kit import check_vlm_serving_perf
 from sglang.test.test_utils import CustomTestCase
 
-register_cuda_ci(est_time=191, stage="extra-a", runner_config="1-gpu-small")
+register_cuda_ci(est_time=195, stage="extra-a", runner_config="1-gpu-small")
 
 
 class TestVLMServingFlashinfer(CustomTestCase):

test/registered/bench_fn/test_bench_serving_reasoning_stream.py

@@ -26,7 +26,7 @@ from sglang.benchmark.serving import (
 from sglang.test.ci.ci_register import register_cpu_ci
 from sglang.test.test_utils import CustomTestCase
 
-register_cpu_ci(est_time=13, suite="base-a-test-cpu")
+register_cpu_ci(est_time=16, suite="base-a-test-cpu")
 
 
 def _free_port() -> int:

test/registered/bench_fn/test_benchmark_datasets_api.py

@@ -60,7 +60,7 @@ from sglang.benchmark.serving import (
 from sglang.test.ci.ci_register import register_cpu_ci
 from sglang.test.test_utils import CustomTestCase
 
-register_cpu_ci(est_time=36, suite="base-a-test-cpu")
+register_cpu_ci(est_time=38, suite="base-a-test-cpu")
 register_cpu_ci(est_time=46, suite="stage-b-test-cpu-intel")
 
 

test/registered/bench_fn/test_steady_state_benchmark.py

@@ -13,7 +13,7 @@ from sglang.benchmark.steady_state import (
 )
 from sglang.test.ci.ci_register import register_cpu_ci
 
-register_cpu_ci(est_time=7, suite="base-a-test-cpu")
+register_cpu_ci(est_time=10, suite="base-a-test-cpu")
 
 
 class _StringTokenizer:

test/registered/chunked_prefill/test_mm_chunked_embedding_unit.py

@@ -25,7 +25,7 @@ from sglang.srt.multimodal.transport.cuda_ipc import (
 from sglang.srt.runtime_context import get_context, get_parallel
 from sglang.test.ci.ci_register import register_cpu_ci
 
-register_cpu_ci(est_time=9, suite="base-a-test-cpu")
+register_cpu_ci(est_time=12, suite="base-a-test-cpu")
 
 
 @pytest.fixture(autouse=True)

test/registered/chunked_prefill/test_scripted_core_1gpu.py

@@ -10,7 +10,7 @@ from sglang.test.scripted_runtime_chunked_helpers import (
     run_until_finished,
 )
 
-register_cuda_ci(est_time=103, stage="extra-a", runner_config="1-gpu-small")
+register_cuda_ci(est_time=106, stage="extra-a", runner_config="1-gpu-small")
 register_amd_ci(est_time=198, stage="extra-a", runner_config="1-gpu-small-amd")
 
 

test/registered/chunked_prefill/test_scripted_core_4gpu.py

@@ -11,7 +11,7 @@ from sglang.test.scripted_runtime_chunked_helpers import (
     run_until_finished,
 )
 
-register_cuda_ci(est_time=96, stage="extra-b", runner_config="4-gpu-h100")
+register_cuda_ci(est_time=55, stage="extra-b", runner_config="4-gpu-h100")
 
 
 _CHUNK_SIZE = 64

test/registered/chunked_prefill/test_scripted_swa_1gpu.py

@@ -5,7 +5,7 @@ from sglang.test.scripted_runtime.context import ScriptedContext
 from sglang.test.scripted_runtime.test_case import ScriptedTestCase
 from sglang.test.scripted_runtime_chunked_helpers import base_engine_kwargs
 
-register_cuda_ci(est_time=137, stage="extra-a", runner_config="1-gpu-large")
+register_cuda_ci(est_time=130, stage="extra-a", runner_config="1-gpu-large")
 
 
 _SWA_MODEL = "openai/gpt-oss-20b"

test/registered/constrained_decoding/test_constrained_decoding.py

@@ -17,7 +17,7 @@ from sglang.test.test_utils import (
     popen_launch_server,
 )
 
-register_cuda_ci(est_time=149, stage="base-b", runner_config="1-gpu-small")
+register_cuda_ci(est_time=161, stage="base-b", runner_config="1-gpu-small")
 register_amd_ci(est_time=220, suite="stage-b-test-1-gpu-small-amd")
 
 

test/registered/core/test_basic_sanity.py

@@ -20,7 +20,7 @@ from sglang.test.test_utils import (
     popen_launch_server,
 )
 
-register_cuda_ci(est_time=112, stage="base-a", runner_config="1-gpu-small")
+register_cuda_ci(est_time=97, stage="base-a", runner_config="1-gpu-small")
 register_amd_ci(est_time=160, suite="stage-a-test-1-gpu-small-amd")
 
 

test/registered/core/test_basic_sanity_dflash.py

@@ -20,7 +20,7 @@ from sglang.test.test_utils import (
     popen_launch_server,
 )
 
-register_cuda_ci(est_time=94, stage="base-a", runner_config="1-gpu-small")
+register_cuda_ci(est_time=135, stage="base-a", runner_config="1-gpu-small")
 
 
 class TestBasicSanityDFlash(

test/registered/core/test_basic_sanity_dspark.py

@@ -16,7 +16,7 @@ from sglang.test.test_utils import (
     popen_launch_server,
 )
 
-register_cuda_ci(est_time=92, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=97, stage="base-b", runner_config="1-gpu-large")
 
 TARGET_MODEL = "Qwen/Qwen3-14B"
 DRAFT_MODEL = "deepseek-ai/dspark_qwen3_14b_block7"

test/registered/core/test_basic_sanity_eagle3.py

@@ -20,7 +20,7 @@ from sglang.test.test_utils import (
     popen_launch_server,
 )
 
-register_cuda_ci(est_time=112, stage="base-a", runner_config="1-gpu-small")
+register_cuda_ci(est_time=167, stage="base-a", runner_config="1-gpu-small")
 register_amd_ci(est_time=200, suite="stage-a-test-1-gpu-small-amd")
 
 

test/registered/core/test_engine_child_pids.py

@@ -20,7 +20,7 @@ from sglang.test.test_utils import (
     CustomTestCase,
 )
 
-register_cuda_ci(est_time=34, stage="base-b", runner_config="1-gpu-small")
+register_cuda_ci(est_time=77, stage="base-b", runner_config="1-gpu-small")
 
 
 class TestEngineChildPids(CustomTestCase):

test/registered/core/test_hidden_states.py

@@ -8,7 +8,7 @@ from sglang.srt.utils import get_device, is_hip
 from sglang.test.ci.ci_register import register_amd_ci, register_cuda_ci
 from sglang.test.test_utils import DEFAULT_SMALL_MODEL_NAME_FOR_TEST, CustomTestCase
 
-register_cuda_ci(est_time=31, stage="base-b", runner_config="1-gpu-small")
+register_cuda_ci(est_time=32, stage="base-b", runner_config="1-gpu-small")
 register_amd_ci(est_time=55, suite="stage-b-test-1-gpu-small-amd")
 
 _is_hip = is_hip()

test/registered/core/test_no_extra_forked_cuda_context.py

@@ -15,7 +15,7 @@ from sglang.test.test_utils import (
     popen_launch_server,
 )
 
-register_cuda_ci(est_time=53, stage="base-b", runner_config="2-gpu-large")
+register_cuda_ci(est_time=54, stage="base-b", runner_config="2-gpu-large")
 
 
 class TestTPServerGPUProcesses(CustomTestCase):

test/registered/core/test_request_queue_validation.py

@@ -17,7 +17,7 @@ from sglang.test.test_utils import (
     send_generate_requests,
 )
 
-register_cuda_ci(est_time=63, stage="base-b", runner_config="1-gpu-small")
+register_cuda_ci(est_time=53, stage="base-b", runner_config="1-gpu-small")
 
 
 class TestMaxQueuedRequests(CustomTestCase):

test/registered/core/test_srt_empty_deps.py

@@ -12,7 +12,7 @@ import pytest
 
 from sglang.test.ci.ci_register import register_cpu_ci
 
-register_cpu_ci(est_time=5, suite="base-a-test-cpu")
+register_cpu_ci(est_time=7, suite="base-a-test-cpu")
 
 # Packages known to transitively depend on torch or triton.
 # If a new package is added to runtime_base and it pulls torch,

test/registered/core/test_srt_endpoint.py

@@ -31,7 +31,7 @@ from sglang.test.test_utils import (
     run_logprob_check,
 )
 
-register_cuda_ci(est_time=250, stage="base-b", runner_config="1-gpu-small")
+register_cuda_ci(est_time=265, stage="base-b", runner_config="1-gpu-small")
 register_amd_ci(est_time=260, suite="stage-b-test-1-gpu-small-amd")
 
 SERVER_ENV = {"SGLANG_USE_PICKLE_IPC": "0"}

test/registered/core/test_srt_engine.py

@@ -20,7 +20,7 @@ from sglang.test.test_utils import (
     CustomTestCase,
 )
 
-register_cuda_ci(est_time=286, stage="base-b", runner_config="1-gpu-large")
+register_cuda_ci(est_time=276, stage="base-b", runner_config="1-gpu-large")
 register_amd_ci(est_time=261, suite="stage-b-test-1-gpu-small-amd")
 
 

test/registered/cp/test_cp_strategy_unit.py

@@ -41,7 +41,7 @@ from sglang.srt.runtime_context import get_parallel
 from sglang.test.ci.ci_register import register_cpu_ci
 from sglang.test.test_utils import CustomTestCase
 
-register_cpu_ci(est_time=8, suite="base-a-test-cpu")
+register_cpu_ci(est_time=11, suite="base-a-test-cpu")
 
 
 class _ExtendMode:

test/registered/cp/test_deepseek_v3_cp_single_node.py

@@ -13,7 +13,7 @@ from sglang.test.test_utils import (
     write_github_step_summary,
 )
 
-register_cuda_ci(est_time=334, stage="extra-b", runner_config="8-gpu-h200")
+register_cuda_ci(est_time=275, stage="extra-b", runner_config="8-gpu-h200")
 
 DEEPSEEK_V3_MODEL_PATH = "deepseek-ai/DeepSeek-V3-0324"
 

test/registered/cp/test_deepseek_v4_flash_fp4_b200_cp.py

@@ -22,7 +22,7 @@ from sglang.test.test_utils import (
     try_cached_model,
 )
 
-register_cuda_ci(est_time=654, stage="extra-b", runner_config="4-gpu-b200")
+register_cuda_ci(est_time=689, stage="extra-b", runner_config="4-gpu-b200")
 
 MODEL = "deepseek-ai/DeepSeek-V4-Flash"
 SERVER_LAUNCH_TIMEOUT = 3600

test/registered/cp/test_dsa_prefill_cp.py

@@ -13,7 +13,7 @@ from sglang.test.test_utils import (
     write_github_step_summary,
 )
 
-register_cuda_ci(est_time=433, stage="extra-b", runner_config="8-gpu-h200")
+register_cuda_ci(est_time=314, stage="extra-b", runner_config="8-gpu-h200")
 GLM52_MODEL_PATH = "zai-org/GLM-5.2-FP8"
 SERVER_LAUNCH_TIMEOUT = max(DEFAULT_TIMEOUT_FOR_SERVER_LAUNCH, 1800)