Compare commits
1
Commits
| Author | SHA1 | Date | |
|---|---|---|---|
|
|
3d8df4b127 |
@@ -1,50 +0,0 @@
|
||||
# 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 未配),勿加回
|
||||
@@ -1,11 +0,0 @@
|
||||
# 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
|
||||
@@ -249,33 +249,23 @@ 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, "
|
||||
"and CP=1 on both servers. DP attention is supported when "
|
||||
"both servers use the same block size and target/draft KV layout."
|
||||
"DP=1 and CP=1. Both servers must enable DSpark with the same "
|
||||
"block size and TP size."
|
||||
)
|
||||
|
||||
from sglang.srt.model_executor.cuda_graph_config import Backend, Phase, with_phase
|
||||
|
||||
prefill_graph = cfg.cuda_graph_config.prefill
|
||||
cp_breakable_prefill = (
|
||||
cfg.enable_prefill_cp
|
||||
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.
|
||||
if prefill_graph.backend != Backend.DISABLED and prefill_graph.max_seq_len is None:
|
||||
# The captured low-ratio indexer scores a static context width; 16k
|
||||
# keeps it inside the candidate window at under 1 ms per layer.
|
||||
declare_resolution(
|
||||
server_args,
|
||||
"validate_deepseek_v41_features",
|
||||
|
||||
@@ -966,14 +966,12 @@ def handle_language_model_only(server_args: Any):
|
||||
):
|
||||
if flag:
|
||||
raise ValueError(f"--language-model-only cannot be combined with {name}")
|
||||
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":
|
||||
if cfg.disaggregation_mode != "null":
|
||||
raise ValueError(
|
||||
"--language-model-only is incompatible with --disaggregation-mode "
|
||||
"prefill/decode"
|
||||
)
|
||||
architectures = hf_config.architectures
|
||||
architectures = model_config_of(server_args).hf_config.architectures
|
||||
if not any(
|
||||
a in server_args.LANGUAGE_MODEL_ONLY_ARCHITECTURES for a in architectures
|
||||
):
|
||||
|
||||
@@ -945,33 +945,9 @@ class CommonKVManager(BaseKVManager):
|
||||
"enable DSpark with the same block size and target/draft KV "
|
||||
"layout. Upgrade both servers together."
|
||||
)
|
||||
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
|
||||
):
|
||||
if info.attn_tp_size != self.attn_tp_size:
|
||||
raise RuntimeError(
|
||||
"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"
|
||||
"DeepSeek-V4.1 DSpark PD requires the same TP size on both servers"
|
||||
)
|
||||
|
||||
if self.dcp_size > 1:
|
||||
|
||||
@@ -82,115 +82,6 @@ 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:
|
||||
@@ -323,7 +214,6 @@ 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
|
||||
|
||||
@@ -337,10 +227,6 @@ 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()
|
||||
@@ -436,13 +322,6 @@ 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()
|
||||
@@ -459,24 +338,6 @@ 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):
|
||||
@@ -887,63 +748,10 @@ class MooncakeKVManager(StagingManagerMixin, CommonKVManager):
|
||||
if not transfer_blocks:
|
||||
return 0
|
||||
|
||||
if not _nvlink_intra_transport_active():
|
||||
src_addrs, dst_addrs, lengths = zip(*transfer_blocks)
|
||||
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
|
||||
src_addrs, dst_addrs, lengths = zip(*transfer_blocks)
|
||||
return self.engine.batch_transfer_sync(
|
||||
mooncake_session_id, list(src_addrs), list(dst_addrs), list(lengths)
|
||||
)
|
||||
|
||||
def _send_kvcache_generic(
|
||||
self,
|
||||
@@ -1674,10 +1482,8 @@ 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")
|
||||
or envs.SGLANG_MOONCAKE_SEND_AUX_TCP.get()
|
||||
or _nvlink_intra_transport_active()
|
||||
):
|
||||
self.enable_custom_mem_pool and self.custom_mem_pool_type == "NVLINK"
|
||||
) or envs.SGLANG_MOONCAKE_SEND_AUX_TCP.get():
|
||||
return self.send_aux_tcp(req, prefill_aux_index, dst_aux_ptrs)
|
||||
|
||||
transfer_blocks = []
|
||||
@@ -1760,71 +1566,6 @@ 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]:
|
||||
@@ -2671,8 +2412,6 @@ 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:
|
||||
@@ -2818,11 +2557,6 @@ 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])
|
||||
@@ -2851,9 +2585,6 @@ 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":
|
||||
|
||||
@@ -736,14 +736,7 @@ 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=}"
|
||||
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
|
||||
)
|
||||
dst_val.copy_(src_val)
|
||||
|
||||
# These fields are safe to replace because captured kernels only need
|
||||
# the current per-replay objects, or the field is produced inside the
|
||||
@@ -995,27 +988,6 @@ 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
|
||||
@@ -1282,11 +1254,6 @@ 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()
|
||||
@@ -1599,12 +1566,8 @@ 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()
|
||||
and get_parallel().attn_cp_size == 1
|
||||
bool(self.low_ratios) and _has_dense_fp4_indexer() and _is_sm100_or_newer()
|
||||
)
|
||||
|
||||
def can_run_prefill_cuda_graph(self, forward_batch: ForwardBatch) -> bool:
|
||||
@@ -2907,7 +2870,7 @@ class DeepseekV4AttnBackend(
|
||||
q_lora[:num_local],
|
||||
positions[:num_local].to(torch.int64),
|
||||
forward_batch,
|
||||
self._move_to_device(q_lens_cpu),
|
||||
torch.tensor(q_lens_cpu, dtype=torch.int32, device=x.device),
|
||||
q_lens_cpu,
|
||||
)
|
||||
|
||||
@@ -3273,9 +3236,7 @@ class DeepseekV4AttnBackend(
|
||||
continue
|
||||
j = torch.arange(lc, device=device)
|
||||
slot_chunks.append(
|
||||
self.req_to_token[req_pool_indices[r : r + 1], j * ratio].to(
|
||||
torch.int64
|
||||
)
|
||||
self.req_to_token[req_pool_indices[r], j * ratio].to(torch.int64)
|
||||
// ratio
|
||||
)
|
||||
start += lc
|
||||
@@ -3300,7 +3261,7 @@ class DeepseekV4AttnBackend(
|
||||
weights = indexer.head_weights(x).float()
|
||||
compress_lens = ((pos + 1) // ratio).to(torch.int32)
|
||||
ks = torch.repeat_interleave(
|
||||
self._move_to_device(starts),
|
||||
torch.tensor(starts, dtype=torch.int32, device=device),
|
||||
q_lens.to(torch.int64),
|
||||
output_size=num_tokens,
|
||||
)
|
||||
@@ -3486,44 +3447,17 @@ class DeepseekV4AttnBackend(
|
||||
self.candidate_indexer.publish_decode(inputs, page_indices, raw_indices)
|
||||
)
|
||||
return
|
||||
if isinstance(metadata.deep_gemm_metadata, list):
|
||||
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(
|
||||
(q_fp4[rows], q_sf[rows]),
|
||||
k_cache,
|
||||
weights[rows],
|
||||
metadata.compressed_seq_lens[rows],
|
||||
metadata.page_table[rows],
|
||||
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
|
||||
)
|
||||
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.
|
||||
@@ -4011,38 +3945,20 @@ class DeepseekV4AttnBackend(
|
||||
compress_ratio, core_attn_metadata, extra_page_size
|
||||
)
|
||||
n_compressed = flat_token_ids.shape[0]
|
||||
reuse_compressed = compress_ratio in (1, 2) and is_cp_active(forward_batch)
|
||||
workspace_pool = (
|
||||
self.shared_compressed_prefill_workspaces[compress_ratio]
|
||||
if reuse_compressed
|
||||
else self.sparse_prefill_workspace
|
||||
workspace = self.sparse_prefill_workspace.get(
|
||||
n_compressed + cache.swa_token_ids.shape[0]
|
||||
)
|
||||
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:
|
||||
source_key = None
|
||||
if reuse_compressed:
|
||||
source_layer = token_to_kv_pool.source_layer_of(layer_id)
|
||||
source_key = (source_layer, workspace.data_ptr())
|
||||
gather = cache.compressed[compress_ratio]
|
||||
# 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(
|
||||
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),
|
||||
)
|
||||
dequantize_k_cache_paged(
|
||||
token_to_kv_pool.get_swa_key_buffer_radix(layer_id),
|
||||
cache.swa_token_ids,
|
||||
|
||||
@@ -21,7 +21,6 @@ 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
|
||||
@@ -177,10 +176,6 @@ 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,
|
||||
@@ -236,83 +231,6 @@ 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,
|
||||
|
||||
@@ -475,40 +475,27 @@ 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,
|
||||
seq_lens,
|
||||
page_table,
|
||||
out_page_indices,
|
||||
metadata.compressed_seq_lens,
|
||||
metadata.page_table,
|
||||
page_indices,
|
||||
metadata.compressed_page_size,
|
||||
metadata.topk_metadata if topk_metadata is None else topk_metadata,
|
||||
out_raw_indices,
|
||||
metadata.topk_metadata,
|
||||
raw_indices,
|
||||
)
|
||||
else:
|
||||
topk_transform_paged(
|
||||
logits,
|
||||
seq_lens,
|
||||
page_table,
|
||||
out_page_indices,
|
||||
metadata.compressed_seq_lens,
|
||||
metadata.page_table,
|
||||
page_indices,
|
||||
metadata.compressed_page_size,
|
||||
out_raw_indices,
|
||||
raw_indices,
|
||||
)
|
||||
|
||||
|
||||
|
||||
@@ -21,7 +21,6 @@ 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__)
|
||||
@@ -307,8 +306,7 @@ class PagedIndexerMetadata:
|
||||
):
|
||||
return None
|
||||
if (
|
||||
get_is_capture_mode()
|
||||
or torch.cuda.is_current_stream_capturing()
|
||||
torch.cuda.is_current_stream_capturing()
|
||||
or is_in_breakable_cuda_graph()
|
||||
or is_in_tc_piecewise_cuda_graph()
|
||||
):
|
||||
@@ -327,27 +325,14 @@ class PagedIndexerMetadata:
|
||||
|
||||
def row_chunks(self):
|
||||
num_rows = self.compressed_seq_lens.shape[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:
|
||||
if self.row_chunk <= 0:
|
||||
return [(slice(0, num_rows), self.deep_gemm_metadata)]
|
||||
|
||||
chunks = [
|
||||
(slice(start, min(start + rows_per_chunk, num_rows)), plan)
|
||||
return [
|
||||
(slice(start, min(start + self.row_chunk, num_rows)), plan)
|
||||
for start, plan in zip(
|
||||
range(0, num_rows, rows_per_chunk), self.deep_gemm_metadata
|
||||
range(0, num_rows, self.row_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.
|
||||
|
||||
@@ -78,11 +78,10 @@ def use_dsv4_q8kv8_sparse_prefill(dsv4_prefill_backend: str = "auto") -> bool:
|
||||
class SparsePrefillWorkspace:
|
||||
"""Backend-owned scratch storage for sparse prefill KV dequantization.
|
||||
|
||||
Callers normally overwrite the entire workspace. Shared compressed-KV
|
||||
callers keep separate workspaces per ratio and track prefix validity in the
|
||||
per-forward gather cache, including the allocation address. Sparse prefill
|
||||
executes eagerly and serially on the supported paths, so the allocation can
|
||||
be replaced when a larger extent is needed.
|
||||
The workspace contents are fully overwritten before every attention call,
|
||||
so token buckets and compression ratios can safely share one buffer. Sparse
|
||||
prefill executes eagerly and serially on the supported paths, which makes it
|
||||
safe to replace the scratch allocation when a larger extent is needed.
|
||||
"""
|
||||
|
||||
def __init__(self, device: torch.device):
|
||||
@@ -276,18 +275,14 @@ 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.
|
||||
|
||||
Gather layouts depend on request/token mappings and compressed page tables.
|
||||
Shared-source dequantization keys live only for this forward; per-layer
|
||||
top-k combinations are recomputed into reused
|
||||
Fields depend on request/token mappings and compressed page tables, not
|
||||
per-layer k_cache; per-layer top-k combinations are recomputed into reused
|
||||
buffers.
|
||||
"""
|
||||
|
||||
|
||||
@@ -23,22 +23,17 @@ 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:
|
||||
@@ -55,18 +50,12 @@ 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 supports_layout
|
||||
and cfg.cp_strategy == "zigzag"
|
||||
and prefill_attention_backend == "trtllm_mha"
|
||||
)
|
||||
|
||||
|
||||
@@ -78,12 +67,8 @@ 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 configured CP strategy can run."""
|
||||
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
|
||||
)
|
||||
"""Keep only token buckets where the zigzag CP strategy can run."""
|
||||
min_num_tokens = resolved_view(server_args).attn_cp_size * 2
|
||||
filtered = [size for size in capture_num_tokens if size >= min_num_tokens]
|
||||
if not filtered:
|
||||
raise ValueError(
|
||||
@@ -111,8 +96,6 @@ 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
|
||||
|
||||
@@ -131,22 +114,12 @@ 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 the active layout."""
|
||||
"""Return the aligned CP-local rows required by a live zigzag layout."""
|
||||
strategy = get_cp_strategy()
|
||||
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):
|
||||
if not isinstance(strategy, ZigzagCPStrategy) or extend_seq_lens is None:
|
||||
return None
|
||||
|
||||
cp_segment_num = strategy.cp_size * 2
|
||||
@@ -246,7 +219,6 @@ 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
|
||||
)
|
||||
@@ -277,31 +249,12 @@ 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._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
|
||||
forward_batch.positions = positions
|
||||
self.live_local_tokens = live_local_tokens
|
||||
|
||||
|
||||
@@ -354,50 +307,10 @@ def execute_prefill_cp_bcg(
|
||||
static_forward_batch,
|
||||
torch.cuda.current_stream(),
|
||||
)
|
||||
if aux_hidden_states is not None:
|
||||
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,
|
||||
return model.logits_processor(
|
||||
forward_batch.input_ids,
|
||||
hidden_states,
|
||||
model.lm_head,
|
||||
logits_metadata,
|
||||
forward_batch,
|
||||
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
|
||||
|
||||
@@ -216,33 +216,20 @@ def _run_mega_routed(
|
||||
|
||||
if num_tokens > 0:
|
||||
router_logits = moe.gate(hidden_states, forward_batch=forward_batch)
|
||||
num_token_non_padded = (
|
||||
forward_batch.num_token_non_padded if forward_batch is not None else None
|
||||
topk_kwargs = {"input_ids": input_ids_global} if moe.is_hash else {}
|
||||
topk_output = moe.topk(
|
||||
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:
|
||||
|
||||
@@ -155,10 +155,6 @@ 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)
|
||||
|
||||
@@ -385,22 +385,14 @@ 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()(input_ids)
|
||||
input_embeds = model.get_input_embeddings()(forward_batch.input_ids)
|
||||
with cp_shard_model_inputs(
|
||||
input_embeds,
|
||||
forward_batch.positions,
|
||||
forward_batch,
|
||||
input_ids,
|
||||
forward_batch.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:
|
||||
@@ -445,7 +437,7 @@ class EagerRunner(BaseRunner):
|
||||
if aux_hidden_states is None:
|
||||
logits_kwargs["hidden_states_before_norm"] = hidden_states_before_norm
|
||||
return model.logits_processor(
|
||||
input_ids,
|
||||
forward_batch.input_ids,
|
||||
hidden_states,
|
||||
model.lm_head,
|
||||
forward_batch,
|
||||
|
||||
@@ -701,9 +701,6 @@ 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
|
||||
|
||||
@@ -785,9 +782,7 @@ 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 = getattr(
|
||||
forward_batch, "_cp_input_ids", forward_batch.input_ids
|
||||
)
|
||||
input_ids = forward_batch.input_ids
|
||||
kwargs = _build_layer_model_forward_kwargs(
|
||||
self.layer_model, forward_batch, pp_proxy_tensors
|
||||
)
|
||||
@@ -1341,9 +1336,9 @@ class PrefillCudaGraphRunner(BaseCudaGraphRunner):
|
||||
batch_max_context_len=batch_max_context_len,
|
||||
):
|
||||
return False
|
||||
if getattr(self, "enable_cp_bcg_capture", False):
|
||||
if not is_cp_active(forward_batch):
|
||||
return False
|
||||
if getattr(self, "enable_cp_bcg_capture", False) and is_cp_active(
|
||||
forward_batch
|
||||
):
|
||||
assert self.prefill_cp_bcg_input is not None
|
||||
if (
|
||||
self.prefill_cp_bcg_input.select_replay_bucket_for_batch(
|
||||
|
||||
@@ -182,7 +182,6 @@ 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,
|
||||
@@ -2040,10 +2039,7 @@ class MQALayer(MqaAttentionBase):
|
||||
if (
|
||||
forward_batch.forward_mode.is_extend()
|
||||
and is_in_breakable_cuda_graph()
|
||||
and (
|
||||
dsa_use_prefill_cp(forward_batch)
|
||||
or not getattr(attn_backend, "low_ratio_prefill_graph", False)
|
||||
)
|
||||
and not getattr(attn_backend, "low_ratio_prefill_graph", False)
|
||||
):
|
||||
bcg_deepseek_v4_low_ratio_sources(self, x, q_lora, positions)
|
||||
else:
|
||||
@@ -2654,8 +2650,7 @@ 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 not getattr(config, "language_model_only", False),
|
||||
and config.vision_n_layers > 0,
|
||||
)
|
||||
|
||||
self.input_layernorm = RMSNorm(config.hidden_size, eps=config.rms_norm_eps)
|
||||
@@ -3877,15 +3872,7 @@ 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():
|
||||
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)
|
||||
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),
|
||||
@@ -4415,18 +4402,11 @@ class DeepseekV4Model(nn.Module):
|
||||
)
|
||||
if self.engram_hasher is not None:
|
||||
if cp_extend:
|
||||
# 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.
|
||||
# n-gram hashing needs each token's predecessors: hash the whole prompt
|
||||
total = int(forward_batch.attn_cp_metadata.total_seq_lens)
|
||||
global_input_ids = forward_batch.input_ids[:total]
|
||||
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)
|
||||
hash_ids = self.engram_hasher(
|
||||
forward_batch.input_ids[:total], 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]
|
||||
@@ -4859,13 +4839,6 @@ 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
|
||||
|
||||
@@ -4891,19 +4864,14 @@ 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
|
||||
and not getattr(config, "language_model_only", False)
|
||||
):
|
||||
if config.model_type == "deepseek_v41" and config.vision_n_layers > 0:
|
||||
if (
|
||||
get_parallel().attn_cp_size != 1
|
||||
or get_pp_group().world_size != 1
|
||||
or not _v41_vision_a2a_supported()
|
||||
or get_parallel().pp_group.world_size != 1
|
||||
or not get_moe_a2a_backend().is_none()
|
||||
):
|
||||
raise ValueError(
|
||||
"V4.1 vision supports TP/EP/DP without CP or PP; "
|
||||
"MoE A2A is supported only with MegaMoE on a PD decode node"
|
||||
"V4.1 vision currently supports TP/EP/DP without CP, PP or MoE A2A"
|
||||
)
|
||||
|
||||
args = SimpleNamespace(**vars(config), dim=config.hidden_size)
|
||||
@@ -5110,19 +5078,16 @@ class DeepseekV4ForCausalLM(nn.Module):
|
||||
0 if is_shared_experts_fusion_disabled() else self.config.n_shared_experts
|
||||
)
|
||||
|
||||
def prepare_language_model_inputs(
|
||||
def forward(
|
||||
self,
|
||||
input_ids: torch.Tensor,
|
||||
positions: torch.Tensor,
|
||||
forward_batch: ForwardBatch,
|
||||
input_embeds: Optional[torch.Tensor] = None,
|
||||
) -> Tuple[torch.Tensor, Optional[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.
|
||||
"""
|
||||
pp_proxy_tensors: Optional[PPProxyTensors] = None,
|
||||
) -> torch.Tensor:
|
||||
if (
|
||||
getattr(self, "vision", None) is not None
|
||||
self.vision 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
|
||||
@@ -5131,7 +5096,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 getattr(self, "vision", None) is not None and not (
|
||||
if self.vision is not None and not (
|
||||
forward_batch.forward_mode.is_decode_or_idle()
|
||||
or forward_batch.forward_mode.is_target_verify()
|
||||
):
|
||||
@@ -5141,19 +5106,6 @@ 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
|
||||
|
||||
@@ -376,7 +376,6 @@ class ServerArgs:
|
||||
# ===== END TO BE REFACTORED ====
|
||||
|
||||
LANGUAGE_MODEL_ONLY_ARCHITECTURES = (
|
||||
"DeepseekV4ForCausalLM",
|
||||
"MuseGlimmerForConditionalGeneration",
|
||||
"Cosmos3ForConditionalGeneration",
|
||||
"Cosmos3EdgeForConditionalGeneration",
|
||||
|
||||
@@ -1,79 +0,0 @@
|
||||
"""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
|
||||
@@ -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=15, suite="base-a-test-cpu")
|
||||
register_cpu_ci(est_time=8, suite="base-a-test-cpu")
|
||||
|
||||
_GLUON_FN = "sglang.srt.layers.attention.aiter_mla_gluon._gluon_fn"
|
||||
|
||||
|
||||
@@ -13,7 +13,7 @@ from sglang.test.ci.ci_register import (
|
||||
register_xpu_ci,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=6, stage="base-b", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=7, 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")
|
||||
|
||||
|
||||
@@ -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=10, stage="base-b", runner_config="1-gpu-small")
|
||||
register_cuda_ci(est_time=9, stage="base-b", runner_config="1-gpu-small")
|
||||
register_amd_ci(est_time=10, suite="stage-b-test-1-gpu-small-amd")
|
||||
|
||||
|
||||
|
||||
@@ -24,7 +24,7 @@ from sglang.test.test_utils import (
|
||||
is_in_amd_ci,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=280, stage="base-b", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=272, 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")
|
||||
|
||||
|
||||
@@ -13,7 +13,7 @@ from sglang.test.test_utils import (
|
||||
)
|
||||
|
||||
# FlashAttention4 integration test (requires SM 100+ / Blackwell B200)
|
||||
register_cuda_ci(est_time=230, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=220, stage="base-b", runner_config="4-gpu-b200")
|
||||
|
||||
|
||||
@unittest.skipIf(get_device_sm() < 100, "Test requires CUDA SM 100 or higher")
|
||||
|
||||
@@ -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=393, stage="extra-a", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=368, stage="extra-a", runner_config="1-gpu-large")
|
||||
|
||||
|
||||
class TestHybridAttnBackendMLA(TestHybridAttnBackendBase):
|
||||
|
||||
@@ -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=25, stage="base-b", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=32, stage="base-b", runner_config="1-gpu-large")
|
||||
register_amd_ci(est_time=12, stage="stage-b", runner_config="1-gpu-large-amd")
|
||||
|
||||
|
||||
|
||||
@@ -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=11, 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=17, suite="stage-b-test-1-gpu-large-amd")
|
||||
|
||||
|
||||
|
||||
@@ -12,7 +12,7 @@ from sglang.test.test_deterministic_utils import (
|
||||
TestDeterministicBase,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=119, stage="extra-b", runner_config="4-gpu-h100")
|
||||
register_cuda_ci(est_time=135, stage="extra-b", runner_config="4-gpu-h100")
|
||||
|
||||
QWEN35 = "Qwen/Qwen3.5-35B-A3B"
|
||||
|
||||
|
||||
@@ -18,7 +18,7 @@ from sglang.test.test_utils import (
|
||||
)
|
||||
|
||||
# Torch native attention backend integration test with MMLU eval
|
||||
register_cuda_ci(est_time=310, stage="extra-a", runner_config="1-gpu-small")
|
||||
register_cuda_ci(est_time=312, stage="extra-a", runner_config="1-gpu-small")
|
||||
register_amd_ci(est_time=150, suite="stage-b-test-1-gpu-small-amd")
|
||||
|
||||
|
||||
|
||||
@@ -15,7 +15,7 @@ from sglang.test.test_utils import (
|
||||
)
|
||||
|
||||
# Sliding window attention with Triton backend (Gemma-3 model)
|
||||
register_cuda_ci(est_time=80, stage="extra-a", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=81, stage="extra-a", runner_config="1-gpu-large")
|
||||
register_amd_ci(est_time=200, suite="stage-b-test-1-gpu-small-amd")
|
||||
|
||||
|
||||
|
||||
@@ -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=10, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=8, stage="base-b", runner_config="4-gpu-b200")
|
||||
|
||||
DEVICE = "cuda"
|
||||
PAGE_SIZE = 32
|
||||
|
||||
@@ -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=16, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=12, stage="base-b", runner_config="4-gpu-b200")
|
||||
|
||||
DEVICE = "cuda"
|
||||
PAGE_SIZE = 128
|
||||
|
||||
@@ -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=12, 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=14, stage="stage-b", runner_config="1-gpu-small-amd")
|
||||
|
||||
|
||||
|
||||
@@ -27,7 +27,7 @@ from sglang.test.test_utils import (
|
||||
popen_launch_server,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=71, stage="base-b", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=90, 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"}
|
||||
|
||||
@@ -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=11, stage="base-a", runner_config="1-gpu-small")
|
||||
register_cuda_ci(est_time=10, stage="base-a", runner_config="1-gpu-small")
|
||||
|
||||
_EXTEND_CASE = DenseAttentionCase(
|
||||
name="extend_no_prefix_smoke",
|
||||
|
||||
@@ -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=16, stage="base-b", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=18, stage="base-b", runner_config="1-gpu-large")
|
||||
|
||||
|
||||
@unittest.skipIf(not torch.cuda.is_available(), "CUDA is required")
|
||||
|
||||
@@ -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=27, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=22, stage="base-b", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=23, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=15, stage="base-b", runner_config="1-gpu-large")
|
||||
|
||||
|
||||
@unittest.skipIf(not torch.cuda.is_available(), "CUDA is required")
|
||||
|
||||
@@ -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=18, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=17, stage="base-b", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=16, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=26, stage="base-b", runner_config="1-gpu-large")
|
||||
|
||||
|
||||
@unittest.skipIf(
|
||||
|
||||
@@ -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=18, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=17, stage="base-b", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=16, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=36, stage="base-b", runner_config="1-gpu-large")
|
||||
|
||||
|
||||
@unittest.skipIf(not torch.cuda.is_available(), "CUDA is required")
|
||||
|
||||
@@ -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=11, 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=10, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large")
|
||||
|
||||
|
||||
@unittest.skipIf(
|
||||
|
||||
@@ -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=11, 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="4-gpu-b200")
|
||||
register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large")
|
||||
|
||||
|
||||
@unittest.skipIf(not torch.cuda.is_available(), "CUDA is required")
|
||||
|
||||
@@ -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=12, 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=11, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=12, stage="base-b", runner_config="1-gpu-large")
|
||||
register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd")
|
||||
|
||||
|
||||
|
||||
@@ -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=18, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=19, stage="base-b", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=38, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=36, stage="base-b", runner_config="1-gpu-large")
|
||||
register_amd_ci(est_time=25, suite="stage-b-test-1-gpu-large-amd")
|
||||
|
||||
|
||||
|
||||
@@ -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=16, stage="base-b", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=15, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=18, stage="base-b", runner_config="1-gpu-large")
|
||||
|
||||
|
||||
@unittest.skipIf(
|
||||
|
||||
@@ -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=14, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=14, stage="base-b", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=22, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=26, stage="base-b", runner_config="1-gpu-large")
|
||||
|
||||
|
||||
@unittest.skipIf(not torch.cuda.is_available(), "CUDA is required")
|
||||
|
||||
@@ -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=13, 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")
|
||||
|
||||
@@ -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=15, 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=14, stage="base-b", runner_config="1-gpu-large")
|
||||
|
||||
_cuda_major = int(torch.version.cuda.split(".")[0]) if torch.version.cuda else 0
|
||||
|
||||
@@ -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=11, 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")
|
||||
|
||||
|
||||
|
||||
@@ -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=12, stage="base-b", runner_config="4-gpu-b200")
|
||||
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="1-gpu-large")
|
||||
register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd")
|
||||
|
||||
|
||||
@@ -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=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=11, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=12, stage="base-b", runner_config="1-gpu-large")
|
||||
register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd")
|
||||
|
||||
|
||||
|
||||
+1
-1
@@ -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=9, stage="base-b", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large")
|
||||
|
||||
|
||||
_KV_LORA_RANK = DEFAULT_KV_LORA_RANK
|
||||
|
||||
@@ -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=9, stage="base-b", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=8, stage="base-b", runner_config="1-gpu-large")
|
||||
|
||||
|
||||
class TestBuildFusedAcceptIndices(CustomTestCase):
|
||||
|
||||
@@ -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=15, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=20, stage="base-b", runner_config="1-gpu-large")
|
||||
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_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd")
|
||||
|
||||
|
||||
|
||||
@@ -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=12, stage="base-b", runner_config="4-gpu-b200")
|
||||
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="1-gpu-large")
|
||||
register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd")
|
||||
|
||||
|
||||
@@ -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=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="4-gpu-b200")
|
||||
register_cuda_ci(est_time=11, stage="base-b", runner_config="1-gpu-large")
|
||||
|
||||
|
||||
@unittest.skipIf(not torch.cuda.is_available(), "CUDA is required")
|
||||
|
||||
@@ -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=6, suite="base-a-test-cpu")
|
||||
register_cpu_ci(est_time=5, suite="base-a-test-cpu")
|
||||
|
||||
|
||||
class TestReplayStateIndicesValidator(unittest.TestCase):
|
||||
|
||||
@@ -32,7 +32,7 @@ MLA_SHAPE_KWARGS = dict(
|
||||
|
||||
from sglang.test.ci.ci_register import register_cuda_ci
|
||||
|
||||
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="4-gpu-b200")
|
||||
register_cuda_ci(est_time=12, stage="base-b", runner_config="1-gpu-large")
|
||||
|
||||
|
||||
|
||||
@@ -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=9, stage="base-b", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=10, stage="base-b", runner_config="1-gpu-large")
|
||||
|
||||
|
||||
@unittest.skipIf(not _SUPPORTED, _SKIP_REASON)
|
||||
|
||||
@@ -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=18, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=20, stage="base-b", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=17, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=16, stage="base-b", runner_config="1-gpu-large")
|
||||
|
||||
|
||||
@unittest.skipIf(not torch.cuda.is_available(), "CUDA is required")
|
||||
|
||||
@@ -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=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=11, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=12, stage="base-b", runner_config="1-gpu-large")
|
||||
|
||||
|
||||
@unittest.skipIf(
|
||||
|
||||
@@ -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=11, suite="base-a-test-cpu")
|
||||
register_cpu_ci(est_time=8, suite="base-a-test-cpu")
|
||||
|
||||
|
||||
class TestSWAKVPoolSetKVBuffer(CustomTestCase):
|
||||
|
||||
@@ -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=12, 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=10, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=12, stage="base-b", runner_config="1-gpu-large")
|
||||
register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd")
|
||||
|
||||
|
||||
|
||||
@@ -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=12, stage="base-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=14, stage="base-b", runner_config="1-gpu-large")
|
||||
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="1-gpu-large")
|
||||
register_amd_ci(est_time=20, suite="stage-b-test-1-gpu-large-amd")
|
||||
|
||||
|
||||
|
||||
@@ -17,7 +17,7 @@ from sglang.test.test_utils import (
|
||||
write_github_step_summary,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=243, stage="extra-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=259, stage="extra-b", runner_config="4-gpu-b200")
|
||||
|
||||
FULL_DEEPSEEK_V3_FP4_MODEL_PATH = "nvidia/DeepSeek-V3-0324-FP4"
|
||||
SERVER_LAUNCH_TIMEOUT = 1000
|
||||
|
||||
@@ -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=28, stage="base-b", runner_config="2-gpu-large")
|
||||
register_cuda_ci(est_time=26, stage="base-b", runner_config="2-gpu-large")
|
||||
|
||||
WORLD_SIZE = 2
|
||||
|
||||
|
||||
@@ -15,7 +15,7 @@ from sglang.test.test_utils import (
|
||||
popen_launch_server,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=139, stage="extra-a", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=131, stage="extra-a", runner_config="1-gpu-large")
|
||||
register_amd_ci(est_time=1100, suite="stage-b-test-1-gpu-small-amd")
|
||||
|
||||
|
||||
|
||||
@@ -13,7 +13,7 @@ from sglang.test.test_utils import (
|
||||
run_bench_serving,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=145, stage="extra-a", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=136, stage="extra-a", runner_config="1-gpu-large")
|
||||
|
||||
|
||||
class TestEagle3Latency(CustomTestCase):
|
||||
|
||||
@@ -16,7 +16,7 @@ from sglang.test.test_utils import (
|
||||
run_embeddings_benchmark_multi,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=245, 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=240, suite="stage-b-test-1-gpu-large-amd")
|
||||
|
||||
|
||||
|
||||
@@ -14,7 +14,7 @@ from sglang.test.test_utils import (
|
||||
run_bench_serving,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=490, stage="extra-a", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=247, stage="extra-a", runner_config="1-gpu-large")
|
||||
register_amd_ci(est_time=430, suite="stage-b-test-1-gpu-large-amd")
|
||||
|
||||
|
||||
|
||||
@@ -11,7 +11,7 @@ from sglang.test.test_utils import (
|
||||
run_bench_serving,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=290, stage="extra-a", runner_config="2-gpu-large")
|
||||
register_cuda_ci(est_time=271, stage="extra-a", runner_config="2-gpu-large")
|
||||
register_amd_ci(est_time=770, suite="stage-b-test-2-gpu-large-amd")
|
||||
|
||||
|
||||
|
||||
@@ -11,7 +11,7 @@ from sglang.test.test_utils import (
|
||||
run_bench_serving,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=490, stage="extra-a", runner_config="2-gpu-large")
|
||||
register_cuda_ci(est_time=325, stage="extra-a", runner_config="2-gpu-large")
|
||||
register_amd_ci(est_time=1030, suite="stage-b-test-2-gpu-large-amd")
|
||||
|
||||
|
||||
|
||||
@@ -16,7 +16,7 @@ from sglang.test.test_utils import (
|
||||
run_score_benchmark_multi,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=215, stage="extra-a", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=185, stage="extra-a", runner_config="1-gpu-large")
|
||||
register_amd_ci(est_time=210, suite="stage-b-test-1-gpu-large-amd")
|
||||
|
||||
|
||||
|
||||
@@ -10,7 +10,7 @@ from sglang.test.test_utils import (
|
||||
run_bench_serving,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=190, stage="extra-a", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=182, stage="extra-a", runner_config="1-gpu-large")
|
||||
register_amd_ci(est_time=165, suite="stage-b-test-1-gpu-large-amd")
|
||||
|
||||
|
||||
|
||||
@@ -11,7 +11,7 @@ from sglang.test.test_utils import (
|
||||
run_bench_serving,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=710, stage="extra-a", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=660, stage="extra-a", runner_config="1-gpu-large")
|
||||
register_amd_ci(est_time=810, suite="stage-b-test-1-gpu-large-amd")
|
||||
|
||||
|
||||
|
||||
@@ -10,7 +10,7 @@ from sglang.test.test_utils import (
|
||||
run_bench_offline_throughput,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=75, stage="extra-a", runner_config="2-gpu-large")
|
||||
register_cuda_ci(est_time=85, stage="extra-a", runner_config="2-gpu-large")
|
||||
register_amd_ci(est_time=280, suite="stage-b-test-2-gpu-large-amd")
|
||||
|
||||
|
||||
|
||||
@@ -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=150, stage="extra-a", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=152, stage="extra-a", runner_config="1-gpu-large")
|
||||
|
||||
|
||||
class TestVLMServingFa3(CustomTestCase):
|
||||
|
||||
@@ -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=195, stage="extra-a", runner_config="1-gpu-small")
|
||||
register_cuda_ci(est_time=191, stage="extra-a", runner_config="1-gpu-small")
|
||||
|
||||
|
||||
class TestVLMServingFlashinfer(CustomTestCase):
|
||||
|
||||
@@ -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=16, suite="base-a-test-cpu")
|
||||
register_cpu_ci(est_time=13, suite="base-a-test-cpu")
|
||||
|
||||
|
||||
def _free_port() -> int:
|
||||
|
||||
@@ -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=38, suite="base-a-test-cpu")
|
||||
register_cpu_ci(est_time=36, suite="base-a-test-cpu")
|
||||
register_cpu_ci(est_time=46, suite="stage-b-test-cpu-intel")
|
||||
|
||||
|
||||
|
||||
@@ -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=10, suite="base-a-test-cpu")
|
||||
register_cpu_ci(est_time=7, suite="base-a-test-cpu")
|
||||
|
||||
|
||||
class _StringTokenizer:
|
||||
|
||||
@@ -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=12, suite="base-a-test-cpu")
|
||||
register_cpu_ci(est_time=9, suite="base-a-test-cpu")
|
||||
|
||||
|
||||
@pytest.fixture(autouse=True)
|
||||
|
||||
@@ -10,7 +10,7 @@ from sglang.test.scripted_runtime_chunked_helpers import (
|
||||
run_until_finished,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=106, stage="extra-a", runner_config="1-gpu-small")
|
||||
register_cuda_ci(est_time=103, stage="extra-a", runner_config="1-gpu-small")
|
||||
register_amd_ci(est_time=198, stage="extra-a", runner_config="1-gpu-small-amd")
|
||||
|
||||
|
||||
|
||||
@@ -11,7 +11,7 @@ from sglang.test.scripted_runtime_chunked_helpers import (
|
||||
run_until_finished,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=55, stage="extra-b", runner_config="4-gpu-h100")
|
||||
register_cuda_ci(est_time=96, stage="extra-b", runner_config="4-gpu-h100")
|
||||
|
||||
|
||||
_CHUNK_SIZE = 64
|
||||
|
||||
@@ -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=130, stage="extra-a", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=137, stage="extra-a", runner_config="1-gpu-large")
|
||||
|
||||
|
||||
_SWA_MODEL = "openai/gpt-oss-20b"
|
||||
|
||||
@@ -17,7 +17,7 @@ from sglang.test.test_utils import (
|
||||
popen_launch_server,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=161, stage="base-b", runner_config="1-gpu-small")
|
||||
register_cuda_ci(est_time=149, stage="base-b", runner_config="1-gpu-small")
|
||||
register_amd_ci(est_time=220, suite="stage-b-test-1-gpu-small-amd")
|
||||
|
||||
|
||||
|
||||
@@ -20,7 +20,7 @@ from sglang.test.test_utils import (
|
||||
popen_launch_server,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=97, stage="base-a", runner_config="1-gpu-small")
|
||||
register_cuda_ci(est_time=112, stage="base-a", runner_config="1-gpu-small")
|
||||
register_amd_ci(est_time=160, suite="stage-a-test-1-gpu-small-amd")
|
||||
|
||||
|
||||
|
||||
@@ -20,7 +20,7 @@ from sglang.test.test_utils import (
|
||||
popen_launch_server,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=135, stage="base-a", runner_config="1-gpu-small")
|
||||
register_cuda_ci(est_time=94, stage="base-a", runner_config="1-gpu-small")
|
||||
|
||||
|
||||
class TestBasicSanityDFlash(
|
||||
|
||||
@@ -16,7 +16,7 @@ from sglang.test.test_utils import (
|
||||
popen_launch_server,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=97, stage="base-b", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=92, stage="base-b", runner_config="1-gpu-large")
|
||||
|
||||
TARGET_MODEL = "Qwen/Qwen3-14B"
|
||||
DRAFT_MODEL = "deepseek-ai/dspark_qwen3_14b_block7"
|
||||
|
||||
@@ -20,7 +20,7 @@ from sglang.test.test_utils import (
|
||||
popen_launch_server,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=167, stage="base-a", runner_config="1-gpu-small")
|
||||
register_cuda_ci(est_time=112, stage="base-a", runner_config="1-gpu-small")
|
||||
register_amd_ci(est_time=200, suite="stage-a-test-1-gpu-small-amd")
|
||||
|
||||
|
||||
|
||||
@@ -20,7 +20,7 @@ from sglang.test.test_utils import (
|
||||
CustomTestCase,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=77, stage="base-b", runner_config="1-gpu-small")
|
||||
register_cuda_ci(est_time=34, stage="base-b", runner_config="1-gpu-small")
|
||||
|
||||
|
||||
class TestEngineChildPids(CustomTestCase):
|
||||
|
||||
@@ -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=32, stage="base-b", runner_config="1-gpu-small")
|
||||
register_cuda_ci(est_time=31, 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()
|
||||
|
||||
@@ -15,7 +15,7 @@ from sglang.test.test_utils import (
|
||||
popen_launch_server,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=54, stage="base-b", runner_config="2-gpu-large")
|
||||
register_cuda_ci(est_time=53, stage="base-b", runner_config="2-gpu-large")
|
||||
|
||||
|
||||
class TestTPServerGPUProcesses(CustomTestCase):
|
||||
|
||||
@@ -17,7 +17,7 @@ from sglang.test.test_utils import (
|
||||
send_generate_requests,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=53, stage="base-b", runner_config="1-gpu-small")
|
||||
register_cuda_ci(est_time=63, stage="base-b", runner_config="1-gpu-small")
|
||||
|
||||
|
||||
class TestMaxQueuedRequests(CustomTestCase):
|
||||
|
||||
@@ -12,7 +12,7 @@ import pytest
|
||||
|
||||
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=5, 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,
|
||||
|
||||
@@ -31,7 +31,7 @@ from sglang.test.test_utils import (
|
||||
run_logprob_check,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=265, stage="base-b", runner_config="1-gpu-small")
|
||||
register_cuda_ci(est_time=250, 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"}
|
||||
|
||||
@@ -20,7 +20,7 @@ from sglang.test.test_utils import (
|
||||
CustomTestCase,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=276, stage="base-b", runner_config="1-gpu-large")
|
||||
register_cuda_ci(est_time=286, stage="base-b", runner_config="1-gpu-large")
|
||||
register_amd_ci(est_time=261, suite="stage-b-test-1-gpu-small-amd")
|
||||
|
||||
|
||||
|
||||
@@ -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=11, suite="base-a-test-cpu")
|
||||
register_cpu_ci(est_time=8, suite="base-a-test-cpu")
|
||||
|
||||
|
||||
class _ExtendMode:
|
||||
|
||||
@@ -13,7 +13,7 @@ from sglang.test.test_utils import (
|
||||
write_github_step_summary,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=275, stage="extra-b", runner_config="8-gpu-h200")
|
||||
register_cuda_ci(est_time=334, stage="extra-b", runner_config="8-gpu-h200")
|
||||
|
||||
DEEPSEEK_V3_MODEL_PATH = "deepseek-ai/DeepSeek-V3-0324"
|
||||
|
||||
|
||||
@@ -22,7 +22,7 @@ from sglang.test.test_utils import (
|
||||
try_cached_model,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=689, stage="extra-b", runner_config="4-gpu-b200")
|
||||
register_cuda_ci(est_time=654, stage="extra-b", runner_config="4-gpu-b200")
|
||||
|
||||
MODEL = "deepseek-ai/DeepSeek-V4-Flash"
|
||||
SERVER_LAUNCH_TIMEOUT = 3600
|
||||
|
||||
@@ -13,7 +13,7 @@ from sglang.test.test_utils import (
|
||||
write_github_step_summary,
|
||||
)
|
||||
|
||||
register_cuda_ci(est_time=314, stage="extra-b", runner_config="8-gpu-h200")
|
||||
register_cuda_ci(est_time=433, 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)
|
||||
|
||||
|
||||
Some files were not shown because too many files have changed in this diff Show More
Reference in New Issue
Block a user