Stabilize CP HiCache residency under L1/L2 pressure
CP shared KV now keeps explicit L1 and host free-room targets so pressure is handled by planned eviction instead of repeated capacity-edge retries. The host allocator gains contiguous-preferred page reservation, L1 owner-lane allocation prefers contiguous physical pages, and CP HiCache metadata preserves pending backup safety for page-granular radix updates. Mooncake transfer stats and allocator microbenchmarks are included to make the remaining transfer bottlenecks measurable rather than inferred. Constraint: CP shared KV uses decode CP size 1 with all prefill CP ranks participating in transfer, so L1/L2 cache residency must remain page-granular and avoid extra collectives.\nConstraint: Production HiCache can be hundreds of GB, so allocator metadata overhead must be visible before enabling aggressive contiguous allocation broadly.\nRejected: Evict only the exact deficit | this keeps the cache at the cliff and causes repeated evict/allocate pressure.\nRejected: Rely on allocator scans alone for contiguity | remote microbenchmarks show fragmented 220GB-equivalent host metadata can make contiguous-preferred scans multi-ms.\nConfidence: medium\nScope-risk: moderate\nDirective: Do not increase L1/L2 free-room defaults or add new CP collectives without ETE evidence and transfer/allocator measurements.\nTested: python -m py_compile on touched runtime/test/benchmark files.\nTested: PYTHONPATH=. python -m pytest -q test/registered/unit/benchmark/test_cp_hicache_allocator_bench.py => 4 passed, 1 warning.\nTested: Remote g0034 log /mnt/beegfs/cjy/log/sglang_cp_hicache_20260601_233723.log shows active prefill process with L1/L2 free-room args, 702 HTTP 200 chat completions, 6272 prefill batches, and no fatal scheduler traceback in latest scan.\nTested: User-reported L1/L2 cache ETE validation passed on remote run.\nNot-tested: Full local pytest suite; local environment is missing several runtime dependencies.\nNot-tested: CUDA allocator microbenchmark during active production prefill process.\nNot-tested: Mooncake straggler fix; stats show transfer tail latency remains a separate bottleneck.
This commit is contained in:
@@ -40,3 +40,36 @@ def group_concurrent_contiguous(
|
||||
dst_groups = [g.tolist() for g in dst_groups]
|
||||
|
||||
return src_groups, dst_groups
|
||||
|
||||
|
||||
def contiguous_group_stats(
|
||||
src_indices: npt.NDArray[np.int32],
|
||||
dst_indices: npt.NDArray[np.int32],
|
||||
src_groups: List[npt.NDArray[np.int32]],
|
||||
dst_groups: List[npt.NDArray[np.int32]],
|
||||
) -> dict[str, object]:
|
||||
"""Summarize contiguous-group coalescing without exposing full index arrays."""
|
||||
|
||||
del dst_groups # Same grouping shape as src_groups; keep arg for call-site clarity.
|
||||
|
||||
src_indices = np.asarray(src_indices).reshape(-1)
|
||||
dst_indices = np.asarray(dst_indices).reshape(-1)
|
||||
group_lens = [len(group) for group in src_groups]
|
||||
page_count = int(src_indices.size)
|
||||
group_count = int(len(group_lens))
|
||||
|
||||
def _diff_head(indices: npt.NDArray[np.int32], limit: int = 16) -> list[int]:
|
||||
if indices.size <= 1:
|
||||
return []
|
||||
sample = indices[: min(indices.size, limit + 1)]
|
||||
return np.diff(sample).astype(np.int64).tolist()
|
||||
|
||||
return {
|
||||
"pages": page_count,
|
||||
"groups": group_count,
|
||||
"min_group_pages": int(min(group_lens)) if group_lens else 0,
|
||||
"max_group_pages": int(max(group_lens)) if group_lens else 0,
|
||||
"avg_group_pages": float(page_count / group_count) if group_count else 0.0,
|
||||
"src_diff_head": _diff_head(src_indices),
|
||||
"dst_diff_head": _diff_head(dst_indices),
|
||||
}
|
||||
|
||||
@@ -23,6 +23,7 @@ from sglang.srt.disaggregation.common.conn import (
|
||||
)
|
||||
from sglang.srt.disaggregation.common.utils import (
|
||||
FastQueue,
|
||||
contiguous_group_stats,
|
||||
group_concurrent_contiguous,
|
||||
)
|
||||
from sglang.srt.disaggregation.mooncake.utils import (
|
||||
@@ -62,6 +63,24 @@ def _cp_draft_shared_kv_debug(message: str, *args, limit: int = 64) -> None:
|
||||
logger.info("[CP_DRAFT_SHARED_KV] " + message, *args)
|
||||
|
||||
|
||||
def _mooncake_transfer_stats_enabled() -> bool:
|
||||
return envs.SGLANG_DISAGGREGATION_TRANSFER_STATS.get()
|
||||
|
||||
|
||||
def _mooncake_transfer_stats_log(message: str, *args) -> None:
|
||||
if not _mooncake_transfer_stats_enabled():
|
||||
return
|
||||
limit = envs.SGLANG_DISAGGREGATION_TRANSFER_STATS_LIMIT.get()
|
||||
if limit is not None and limit <= 0:
|
||||
return
|
||||
key = "mooncake_transfer_stats"
|
||||
count = _CP_SHARED_DEBUG_COUNTS.get(key, 0)
|
||||
if limit is not None and count >= limit:
|
||||
return
|
||||
_CP_SHARED_DEBUG_COUNTS[key] = count + 1
|
||||
logger.info("[Mooncake-transfer-stats] " + message, *args)
|
||||
|
||||
|
||||
def _np_summary(arr) -> str:
|
||||
if arr is None:
|
||||
return "None"
|
||||
@@ -310,6 +329,7 @@ class MooncakeKVManager(CommonKVManager):
|
||||
prefill_data_indices: npt.NDArray[np.int32],
|
||||
dst_data_indices: npt.NDArray[np.int32],
|
||||
executor: concurrent.futures.ThreadPoolExecutor,
|
||||
debug_room: Optional[int] = None,
|
||||
) -> int:
|
||||
"""
|
||||
Generic KV cache transfer supporting both MHA and MLA architectures.
|
||||
@@ -319,6 +339,15 @@ class MooncakeKVManager(CommonKVManager):
|
||||
prefill_kv_blocks, dst_kv_blocks = group_concurrent_contiguous(
|
||||
prefill_data_indices, dst_data_indices
|
||||
)
|
||||
transfer_stats_enabled = _mooncake_transfer_stats_enabled()
|
||||
grouping_stats = None
|
||||
if transfer_stats_enabled:
|
||||
grouping_stats = contiguous_group_stats(
|
||||
prefill_data_indices,
|
||||
dst_data_indices,
|
||||
prefill_kv_blocks,
|
||||
dst_kv_blocks,
|
||||
)
|
||||
|
||||
layers_params = None
|
||||
|
||||
@@ -387,6 +416,7 @@ class MooncakeKVManager(CommonKVManager):
|
||||
transfer_blocks.extend(set_transfer_blocks(src_ptr, dst_ptr, item_len))
|
||||
return self._transfer_data(mooncake_session_id, transfer_blocks)
|
||||
|
||||
start_time = time.perf_counter() if transfer_stats_enabled else 0.0
|
||||
if self.enable_custom_mem_pool:
|
||||
futures = [
|
||||
executor.submit(
|
||||
@@ -402,12 +432,84 @@ class MooncakeKVManager(CommonKVManager):
|
||||
if status != 0:
|
||||
for f in futures:
|
||||
f.cancel()
|
||||
if transfer_stats_enabled:
|
||||
self._log_kvcache_transfer_stats(
|
||||
mooncake_session_id=mooncake_session_id,
|
||||
debug_room=debug_room,
|
||||
grouping_stats=grouping_stats,
|
||||
layers_params=layers_params,
|
||||
elapsed_ms=(time.perf_counter() - start_time) * 1000,
|
||||
status=status,
|
||||
custom_mem_pool=True,
|
||||
)
|
||||
return status
|
||||
return 0
|
||||
status = 0
|
||||
else:
|
||||
# Combining all layers' params in one batch transfer is more efficient
|
||||
# compared to using multiple threads
|
||||
return process_layers(layers_params)
|
||||
status = process_layers(layers_params)
|
||||
|
||||
if transfer_stats_enabled:
|
||||
self._log_kvcache_transfer_stats(
|
||||
mooncake_session_id=mooncake_session_id,
|
||||
debug_room=debug_room,
|
||||
grouping_stats=grouping_stats,
|
||||
layers_params=layers_params,
|
||||
elapsed_ms=(time.perf_counter() - start_time) * 1000,
|
||||
status=status,
|
||||
custom_mem_pool=self.enable_custom_mem_pool,
|
||||
)
|
||||
return status
|
||||
|
||||
def _log_kvcache_transfer_stats(
|
||||
self,
|
||||
mooncake_session_id: str,
|
||||
debug_room: Optional[int],
|
||||
grouping_stats: Optional[dict[str, object]],
|
||||
layers_params: List[Tuple[int, int, int]],
|
||||
elapsed_ms: float,
|
||||
status: int,
|
||||
custom_mem_pool: bool,
|
||||
) -> None:
|
||||
if grouping_stats is None:
|
||||
return
|
||||
page_count = int(grouping_stats["pages"])
|
||||
group_count = int(grouping_stats["groups"])
|
||||
layer_count = len(layers_params)
|
||||
item_bytes_per_page = sum(int(item_len) for _, _, item_len in layers_params)
|
||||
total_bytes = page_count * item_bytes_per_page
|
||||
transfer_blocks = group_count * layer_count
|
||||
avg_block_bytes = (
|
||||
float(total_bytes / transfer_blocks) if transfer_blocks else 0.0
|
||||
)
|
||||
bandwidth_gbps = (
|
||||
float(total_bytes / elapsed_ms / 1e6) if elapsed_ms > 0 else 0.0
|
||||
)
|
||||
|
||||
_mooncake_transfer_stats_log(
|
||||
"cp_rank=%s room=%s session=%s status=%s custom_mem_pool=%s "
|
||||
"pages=%s groups=%s avg_group_pages=%.2f max_group_pages=%s "
|
||||
"layers=%s transfer_blocks=%s total_bytes=%.3fGiB "
|
||||
"avg_block_bytes=%.1fKiB elapsed_ms=%.3f bandwidth=%.2fGB/s "
|
||||
"src_diff_head=%s dst_diff_head=%s",
|
||||
self.attn_cp_rank,
|
||||
debug_room,
|
||||
mooncake_session_id,
|
||||
status,
|
||||
custom_mem_pool,
|
||||
page_count,
|
||||
group_count,
|
||||
float(grouping_stats["avg_group_pages"]),
|
||||
grouping_stats["max_group_pages"],
|
||||
layer_count,
|
||||
transfer_blocks,
|
||||
total_bytes / (1024**3),
|
||||
avg_block_bytes / 1024,
|
||||
elapsed_ms,
|
||||
bandwidth_gbps,
|
||||
grouping_stats["src_diff_head"],
|
||||
grouping_stats["dst_diff_head"],
|
||||
)
|
||||
|
||||
def send_kvcache(
|
||||
self,
|
||||
@@ -416,6 +518,7 @@ class MooncakeKVManager(CommonKVManager):
|
||||
dst_kv_ptrs: list[int],
|
||||
dst_kv_indices: npt.NDArray[np.int32],
|
||||
executor: concurrent.futures.ThreadPoolExecutor,
|
||||
debug_room: Optional[int] = None,
|
||||
):
|
||||
return self._send_kvcache_generic(
|
||||
mooncake_session_id=mooncake_session_id,
|
||||
@@ -425,6 +528,7 @@ class MooncakeKVManager(CommonKVManager):
|
||||
prefill_data_indices=prefill_kv_indices,
|
||||
dst_data_indices=dst_kv_indices,
|
||||
executor=executor,
|
||||
debug_room=debug_room,
|
||||
)
|
||||
|
||||
def send_kvcache_slice(
|
||||
@@ -973,6 +1077,7 @@ class MooncakeKVManager(CommonKVManager):
|
||||
target_rank_registration_info.dst_kv_ptrs,
|
||||
chunked_dst_kv_indice,
|
||||
executor,
|
||||
debug_room=kv_chunk.room,
|
||||
)
|
||||
else:
|
||||
ret = self.send_kvcache_slice(
|
||||
|
||||
@@ -261,6 +261,8 @@ class Envs:
|
||||
SGLANG_DISAGGREGATION_WAITING_TIMEOUT = EnvInt(300)
|
||||
SGLANG_DISAGGREGATION_NIXL_BACKEND = EnvStr("UCX")
|
||||
SGLANG_DISAGGREGATION_ALL_CP_RANKS_TRANSFER = EnvBool(False)
|
||||
SGLANG_DISAGGREGATION_TRANSFER_STATS = EnvBool(False)
|
||||
SGLANG_DISAGGREGATION_TRANSFER_STATS_LIMIT = EnvInt(128)
|
||||
|
||||
# Scheduler: others:
|
||||
SGLANG_EMPTY_CACHE_INTERVAL = EnvFloat(-1) # in seconds. Set if you observe high memory accumulation over a long serving period.
|
||||
|
||||
@@ -980,7 +980,12 @@ class HiCacheController:
|
||||
|
||||
owned_logical_indices = padded_device_indices[owned_mask]
|
||||
physical_device_indices = layout.logical_locs_to_physical(owned_logical_indices)
|
||||
host_indices = self.mem_pool_host.alloc(len(physical_device_indices))
|
||||
host_alloc = getattr(
|
||||
self.mem_pool_host,
|
||||
"alloc_contiguous_preferred",
|
||||
self.mem_pool_host.alloc,
|
||||
)
|
||||
host_indices = host_alloc(len(physical_device_indices))
|
||||
if host_indices is None:
|
||||
logger.info(
|
||||
"[CacheCtrl-write] reserve_write_cp FAILED (host full): node_id=%d logical_len=%d owned=%d",
|
||||
@@ -992,9 +997,12 @@ class HiCacheController:
|
||||
|
||||
draft_host_indices = None
|
||||
if self.has_draft_hicache:
|
||||
draft_host_indices = self.draft_mem_pool_host.alloc(
|
||||
len(physical_device_indices)
|
||||
draft_host_alloc = getattr(
|
||||
self.draft_mem_pool_host,
|
||||
"alloc_contiguous_preferred",
|
||||
self.draft_mem_pool_host.alloc,
|
||||
)
|
||||
draft_host_indices = draft_host_alloc(len(physical_device_indices))
|
||||
if draft_host_indices is None:
|
||||
self.mem_pool_host.free(host_indices)
|
||||
logger.info(
|
||||
|
||||
@@ -20,6 +20,7 @@ Page-aligned memory pool.
|
||||
"""
|
||||
|
||||
import abc
|
||||
import math
|
||||
from typing import TYPE_CHECKING, List, Optional
|
||||
|
||||
import torch
|
||||
@@ -36,6 +37,33 @@ if TYPE_CHECKING:
|
||||
from sglang.srt.mem_cache.memory_pool import KVCache
|
||||
|
||||
|
||||
def compute_owner_lane_free_room_deficits(
|
||||
*,
|
||||
required: List[int],
|
||||
available: List[int],
|
||||
capacities: List[int],
|
||||
target_ratio: float,
|
||||
trigger_ratio: float,
|
||||
) -> List[int]:
|
||||
deficits: List[int] = []
|
||||
for req, avail, capacity in zip(required, available, capacities):
|
||||
target_room = (
|
||||
int(math.ceil(float(capacity) * float(target_ratio)))
|
||||
if capacity > 0 and target_ratio > 0
|
||||
else 0
|
||||
)
|
||||
trigger_room = (
|
||||
int(math.ceil(float(capacity) * float(trigger_ratio)))
|
||||
if capacity > 0 and trigger_ratio > 0
|
||||
else 0
|
||||
)
|
||||
if int(avail) >= int(req) + trigger_room:
|
||||
deficits.append(0)
|
||||
else:
|
||||
deficits.append(max(0, int(req) + target_room - int(avail)))
|
||||
return deficits
|
||||
|
||||
|
||||
def _debug_sort_nvtx_enabled() -> bool:
|
||||
return _SORT_NVTX_ENABLED
|
||||
|
||||
@@ -632,6 +660,73 @@ class CPSharedPagedTokenToKVPoolAllocator(PagedTokenToKVPoolAllocator):
|
||||
]
|
||||
return required, available, deficits
|
||||
|
||||
def compute_owner_lane_capacity_pages(self) -> List[int]:
|
||||
capacity_pages = int(self.physical_size // self.page_size)
|
||||
return [capacity_pages for _ in range(self.cp_size)]
|
||||
|
||||
def compute_owner_lane_free_room_stats(
|
||||
self,
|
||||
page_compute_owners: List[int],
|
||||
*,
|
||||
target_ratio: float,
|
||||
trigger_ratio: float,
|
||||
) -> tuple[List[int], List[int], List[int]]:
|
||||
"""Return owner-lane deficits with reactive trigger/target free room.
|
||||
|
||||
Units are pages, matching ``page_compute_owners`` and
|
||||
``owner_lane_deficits``. With both ratios at zero this is identical to
|
||||
``compute_owner_lane_stats``.
|
||||
"""
|
||||
|
||||
required, available, _exact_deficits = self.compute_owner_lane_stats(
|
||||
page_compute_owners
|
||||
)
|
||||
deficits = compute_owner_lane_free_room_deficits(
|
||||
required=required,
|
||||
available=available,
|
||||
capacities=self.compute_owner_lane_capacity_pages(),
|
||||
target_ratio=target_ratio,
|
||||
trigger_ratio=trigger_ratio,
|
||||
)
|
||||
return required, available, deficits
|
||||
|
||||
def _select_owner_free_pages_prefer_contiguous(
|
||||
self,
|
||||
owner_mask: torch.Tensor,
|
||||
required_count: int,
|
||||
) -> torch.Tensor:
|
||||
selected_prefix_mask = owner_mask & (
|
||||
torch.cumsum(owner_mask.to(torch.int64), dim=0) <= required_count
|
||||
)
|
||||
if required_count <= 1:
|
||||
return selected_prefix_mask
|
||||
|
||||
owner_positions = owner_mask.nonzero(as_tuple=True)[0]
|
||||
if owner_positions.numel() < required_count:
|
||||
return selected_prefix_mask
|
||||
|
||||
owner_pages = self.free_pages[owner_positions]
|
||||
run_edges = owner_pages[1:] - owner_pages[:-1] == self.cp_size
|
||||
if run_edges.numel() < required_count - 1:
|
||||
return selected_prefix_mask
|
||||
|
||||
eligible_starts = run_edges.unfold(0, required_count - 1, 1).all(dim=1)
|
||||
if eligible_starts.numel() == 0:
|
||||
return selected_prefix_mask
|
||||
|
||||
chosen_start = torch.argmax(eligible_starts.to(torch.int64))
|
||||
contiguous_positions = owner_positions[
|
||||
chosen_start
|
||||
+ torch.arange(
|
||||
required_count, dtype=torch.int64, device=owner_positions.device
|
||||
)
|
||||
]
|
||||
selected_contiguous_mask = torch.zeros_like(owner_mask)
|
||||
selected_contiguous_mask[contiguous_positions] = True
|
||||
return torch.where(
|
||||
eligible_starts.any(), selected_contiguous_mask, selected_prefix_mask
|
||||
)
|
||||
|
||||
def _select_compute_owner_pages(
|
||||
self,
|
||||
page_compute_owners: List[int],
|
||||
@@ -659,8 +754,10 @@ class CPSharedPagedTokenToKVPoolAllocator(PagedTokenToKVPoolAllocator):
|
||||
continue
|
||||
|
||||
owner_mask = torch.remainder(self.free_pages - 1, self.cp_size) == owner
|
||||
selected_owner_free_mask = owner_mask & (
|
||||
torch.cumsum(owner_mask.to(torch.int64), dim=0) <= required_count
|
||||
selected_owner_free_mask = (
|
||||
self._select_owner_free_pages_prefer_contiguous(
|
||||
owner_mask, required_count
|
||||
)
|
||||
)
|
||||
selected_owner_pages = self.free_pages[selected_owner_free_mask]
|
||||
|
||||
|
||||
@@ -12,6 +12,7 @@ from sglang.srt.mem_cache.base_prefix_cache import (
|
||||
EvictParams,
|
||||
EvictResult,
|
||||
)
|
||||
from sglang.srt.mem_cache.allocator import compute_owner_lane_free_room_deficits
|
||||
from sglang.srt.mem_cache.cp_shared_kv_compute_owner import (
|
||||
build_in_seq_page_compute_owners,
|
||||
get_in_seq_page_compute_owner_unavailable_reason,
|
||||
@@ -65,6 +66,42 @@ def _log_cp_shared_kv_alloc_fallback(
|
||||
)
|
||||
|
||||
|
||||
def _compute_owner_lane_stats_for_eviction(
|
||||
*,
|
||||
tree_cache: BasePrefixCache,
|
||||
allocator,
|
||||
page_compute_owners: list[int],
|
||||
) -> tuple[list[int], list[int], list[int]]:
|
||||
target_ratio = float(getattr(tree_cache, "hicache_l1_free_room_ratio", 0.0) or 0.0)
|
||||
trigger_ratio = float(
|
||||
getattr(tree_cache, "hicache_l1_free_room_trigger_ratio", 0.0) or 0.0
|
||||
)
|
||||
free_room_stats = getattr(allocator, "compute_owner_lane_free_room_stats", None)
|
||||
if free_room_stats is not None:
|
||||
return free_room_stats(
|
||||
page_compute_owners,
|
||||
target_ratio=target_ratio,
|
||||
trigger_ratio=trigger_ratio,
|
||||
)
|
||||
required, available, exact_deficits = allocator.compute_owner_lane_stats(
|
||||
page_compute_owners
|
||||
)
|
||||
lane_capacity_pages = getattr(allocator, "compute_owner_lane_capacity_pages", None)
|
||||
if lane_capacity_pages is None:
|
||||
return required, available, exact_deficits
|
||||
return (
|
||||
required,
|
||||
available,
|
||||
compute_owner_lane_free_room_deficits(
|
||||
required=required,
|
||||
available=available,
|
||||
capacities=lane_capacity_pages(),
|
||||
target_ratio=target_ratio,
|
||||
trigger_ratio=trigger_ratio,
|
||||
),
|
||||
)
|
||||
|
||||
|
||||
@triton.jit
|
||||
def write_req_to_token_pool_triton(
|
||||
req_to_token_ptr, # [max_batch, max_context_len]
|
||||
@@ -331,7 +368,11 @@ def _evict_for_compute_owner_lanes(
|
||||
|
||||
max_attempts = max(2, min(8, int(getattr(allocator, "cp_size", 1))))
|
||||
for attempt in range(max_attempts):
|
||||
_required, _available, deficits = compute_owner_lane_stats(page_compute_owners)
|
||||
_required, _available, deficits = _compute_owner_lane_stats_for_eviction(
|
||||
tree_cache=tree_cache,
|
||||
allocator=allocator,
|
||||
page_compute_owners=page_compute_owners,
|
||||
)
|
||||
deficit_pages = sum(deficits)
|
||||
if deficit_pages <= 0:
|
||||
return
|
||||
|
||||
@@ -4,6 +4,7 @@ import atexit
|
||||
import heapq
|
||||
import json
|
||||
import logging
|
||||
import math
|
||||
import os
|
||||
import threading
|
||||
import time
|
||||
@@ -20,7 +21,10 @@ from sglang.srt.managers.cache_controller import (
|
||||
HiCacheWriteReservation,
|
||||
PrefetchOperation,
|
||||
)
|
||||
from sglang.srt.mem_cache.allocator import CPSharedPagedTokenToKVPoolAllocator
|
||||
from sglang.srt.mem_cache.allocator import (
|
||||
CPSharedPagedTokenToKVPoolAllocator,
|
||||
compute_owner_lane_free_room_deficits,
|
||||
)
|
||||
from sglang.srt.mem_cache.base_prefix_cache import (
|
||||
DecLockRefParams,
|
||||
DecLockRefResult,
|
||||
@@ -126,6 +130,31 @@ def _compute_shared_hicache_token_capacities(
|
||||
return target_tokens, draft_tokens
|
||||
|
||||
|
||||
def _page_aligned_room_from_ratio(capacity: int, ratio: float, page_size: int) -> int:
|
||||
if page_size <= 0:
|
||||
raise ValueError(f"page_size must be positive, got {page_size}")
|
||||
if capacity <= 0 or ratio <= 0:
|
||||
return 0
|
||||
raw_room = int(math.ceil(float(capacity) * float(ratio)))
|
||||
return ((raw_room + page_size - 1) // page_size) * page_size
|
||||
|
||||
|
||||
def _free_room_deficit(
|
||||
*,
|
||||
required: int,
|
||||
available: int,
|
||||
capacity: int,
|
||||
page_size: int,
|
||||
target_ratio: float,
|
||||
trigger_ratio: float,
|
||||
) -> int:
|
||||
target_room = _page_aligned_room_from_ratio(capacity, target_ratio, page_size)
|
||||
trigger_room = _page_aligned_room_from_ratio(capacity, trigger_ratio, page_size)
|
||||
if int(available) >= int(required) + trigger_room:
|
||||
return 0
|
||||
return max(0, int(required) + target_room - int(available))
|
||||
|
||||
|
||||
@dataclass
|
||||
class CpHiCacheNodeMetadata:
|
||||
# Legacy name kept for existing call sites. Under the page-aligned cache
|
||||
@@ -570,6 +599,14 @@ class HiRadixCache(RadixCache):
|
||||
1 if server_args.hicache_write_policy == "write_through" else 2
|
||||
)
|
||||
self.load_back_threshold = 10
|
||||
self.hicache_host_free_room_ratio = server_args.hicache_host_free_room_ratio
|
||||
self.hicache_host_free_room_trigger_ratio = (
|
||||
server_args.hicache_host_free_room_trigger_ratio
|
||||
)
|
||||
self.hicache_l1_free_room_ratio = server_args.hicache_l1_free_room_ratio
|
||||
self.hicache_l1_free_room_trigger_ratio = (
|
||||
server_args.hicache_l1_free_room_trigger_ratio
|
||||
)
|
||||
|
||||
# Detach storage backend automatically on process shutdown
|
||||
atexit.register(self.shutdown)
|
||||
@@ -973,7 +1010,7 @@ class HiRadixCache(RadixCache):
|
||||
f"page_size={self.page_size} page_owners={len(page_owners)}"
|
||||
)
|
||||
|
||||
required, available, deficits = compute_owner_lane_stats(page_owners)
|
||||
required, available, deficits = self._cp_load_back_owner_lane_stats(page_owners)
|
||||
return CpLoadBackPlan(
|
||||
page_owners=list(page_owners),
|
||||
required_by_owner=[int(v) for v in required],
|
||||
@@ -982,9 +1019,42 @@ class HiRadixCache(RadixCache):
|
||||
host_hit_len=host_hit_len,
|
||||
)
|
||||
|
||||
def _cp_load_back_owner_lane_stats(
|
||||
self, page_owners: List[int]
|
||||
) -> Tuple[List[int], List[int], List[int]]:
|
||||
allocator = self.token_to_kv_pool_allocator
|
||||
target_ratio = float(getattr(self, "hicache_l1_free_room_ratio", 0.0) or 0.0)
|
||||
trigger_ratio = float(
|
||||
getattr(self, "hicache_l1_free_room_trigger_ratio", 0.0) or 0.0
|
||||
)
|
||||
free_room_stats = getattr(allocator, "compute_owner_lane_free_room_stats", None)
|
||||
if free_room_stats is not None:
|
||||
return free_room_stats(
|
||||
page_owners,
|
||||
target_ratio=target_ratio,
|
||||
trigger_ratio=trigger_ratio,
|
||||
)
|
||||
required, available, exact_deficits = allocator.compute_owner_lane_stats(
|
||||
page_owners
|
||||
)
|
||||
lane_capacity_pages = getattr(allocator, "compute_owner_lane_capacity_pages", None)
|
||||
if lane_capacity_pages is None:
|
||||
return required, available, exact_deficits
|
||||
return (
|
||||
required,
|
||||
available,
|
||||
compute_owner_lane_free_room_deficits(
|
||||
required=required,
|
||||
available=available,
|
||||
capacities=lane_capacity_pages(),
|
||||
target_ratio=target_ratio,
|
||||
trigger_ratio=trigger_ratio,
|
||||
),
|
||||
)
|
||||
|
||||
def _refresh_cp_load_back_plan(self, plan: CpLoadBackPlan) -> CpLoadBackPlan:
|
||||
required, available, deficits = (
|
||||
self.token_to_kv_pool_allocator.compute_owner_lane_stats(plan.page_owners)
|
||||
required, available, deficits = self._cp_load_back_owner_lane_stats(
|
||||
plan.page_owners
|
||||
)
|
||||
return CpLoadBackPlan(
|
||||
page_owners=plan.page_owners,
|
||||
@@ -1344,11 +1414,42 @@ class HiRadixCache(RadixCache):
|
||||
snapshot = self._cp_host_capacity_snapshot()
|
||||
target_available = self._cp_host_available_tokens_by_rank(snapshot)
|
||||
draft_available = self._cp_host_available_tokens_by_rank(snapshot, draft=True)
|
||||
target_deficit = tuple(
|
||||
max(0, req - avail) for req, avail in zip(required, target_available)
|
||||
target_ratio = float(getattr(self, "hicache_host_free_room_ratio", 0.0) or 0.0)
|
||||
trigger_ratio = float(
|
||||
getattr(self, "hicache_host_free_room_trigger_ratio", 0.0) or 0.0
|
||||
)
|
||||
target_deficit = tuple(
|
||||
_free_room_deficit(
|
||||
required=req,
|
||||
available=avail,
|
||||
capacity=capacity,
|
||||
page_size=self.page_size,
|
||||
target_ratio=target_ratio,
|
||||
trigger_ratio=trigger_ratio,
|
||||
)
|
||||
for req, avail, capacity in zip(
|
||||
required, target_available, snapshot.target_capacity
|
||||
)
|
||||
)
|
||||
draft_capacity = snapshot.draft_capacity
|
||||
draft_deficit = tuple(
|
||||
max(0, req - avail) for req, avail in zip(required, draft_available)
|
||||
(
|
||||
_free_room_deficit(
|
||||
required=req,
|
||||
available=avail,
|
||||
capacity=capacity,
|
||||
page_size=self.page_size,
|
||||
target_ratio=target_ratio,
|
||||
trigger_ratio=trigger_ratio,
|
||||
)
|
||||
if draft_capacity is not None
|
||||
else 0
|
||||
)
|
||||
for req, avail, capacity in zip(
|
||||
required,
|
||||
draft_available,
|
||||
draft_capacity if draft_capacity is not None else snapshot.target_capacity,
|
||||
)
|
||||
)
|
||||
deficit = tuple(max(a, b) for a, b in zip(target_deficit, draft_deficit))
|
||||
eviction_plan = self._plan_cp_host_evictions(deficit)
|
||||
|
||||
@@ -338,6 +338,88 @@ class HostKVCache(abc.ABC):
|
||||
|
||||
return select_index
|
||||
|
||||
@synchronized
|
||||
def alloc_contiguous_preferred(self, need_size: int) -> Optional[torch.Tensor]:
|
||||
"""Allocate page-shaped host slots, preferring physical page runs.
|
||||
|
||||
``free_slots`` is still the source of truth and the fallback remains FIFO.
|
||||
The CP HiCache fast paths only require best-effort contiguity: if a
|
||||
contiguous physical run is readily available, use it to reduce transfer
|
||||
descriptor fragmentation; otherwise preserve the historical allocator
|
||||
behavior by returning ``alloc(need_size)``.
|
||||
"""
|
||||
|
||||
assert (
|
||||
need_size % self.page_size == 0
|
||||
), "The requested size should be a multiple of the page size."
|
||||
if need_size > self.available_size():
|
||||
return None
|
||||
if need_size == 0:
|
||||
return self.alloc(need_size)
|
||||
|
||||
fifo_prefix = self.free_slots[:need_size]
|
||||
expected_prefix = fifo_prefix[:1] + torch.arange(
|
||||
need_size, dtype=fifo_prefix.dtype, device=fifo_prefix.device
|
||||
)
|
||||
if torch.equal(fifo_prefix, expected_prefix):
|
||||
return self.alloc(need_size)
|
||||
|
||||
page_size = int(self.page_size)
|
||||
need_pages = need_size // page_size
|
||||
page_count = int(self.free_slots.numel()) // page_size
|
||||
if page_count < need_pages:
|
||||
return self.alloc(need_size)
|
||||
|
||||
page_slots = self.free_slots[: page_count * page_size].view(
|
||||
page_count, page_size
|
||||
)
|
||||
page_offsets = torch.arange(
|
||||
page_size, dtype=page_slots.dtype, device=page_slots.device
|
||||
)
|
||||
page_is_contiguous = torch.all(
|
||||
page_slots == (page_slots[:, :1] + page_offsets), dim=1
|
||||
)
|
||||
if not bool(page_is_contiguous.any()):
|
||||
return self.alloc(need_size)
|
||||
|
||||
chunk_indices = torch.arange(page_count, dtype=torch.int64)[
|
||||
page_is_contiguous.cpu()
|
||||
]
|
||||
page_starts = page_slots[:, 0][page_is_contiguous]
|
||||
sorted_starts, order = torch.sort(page_starts)
|
||||
sorted_chunks = chunk_indices[order.cpu()]
|
||||
|
||||
run_start = -1
|
||||
if need_pages == 1:
|
||||
run_start = 0
|
||||
else:
|
||||
current_run_start = 0
|
||||
current_run_len = 1
|
||||
diffs = sorted_starts[1:] - sorted_starts[:-1]
|
||||
for offset, diff in enumerate(diffs.tolist(), start=1):
|
||||
if diff == page_size:
|
||||
current_run_len += 1
|
||||
if current_run_len >= need_pages:
|
||||
run_start = current_run_start
|
||||
break
|
||||
else:
|
||||
current_run_start = offset
|
||||
current_run_len = 1
|
||||
|
||||
if run_start < 0:
|
||||
return self.alloc(need_size)
|
||||
|
||||
selected_chunks = sorted_chunks[run_start : run_start + need_pages]
|
||||
token_offsets = (
|
||||
selected_chunks[:, None] * page_size
|
||||
+ torch.arange(page_size, dtype=torch.int64)[:, None].T
|
||||
).reshape(-1)
|
||||
select_index = self.free_slots[token_offsets]
|
||||
keep_mask = torch.ones(self.free_slots.numel(), dtype=torch.bool)
|
||||
keep_mask[token_offsets] = False
|
||||
self.free_slots = self.free_slots[keep_mask]
|
||||
return select_index
|
||||
|
||||
@synchronized
|
||||
def free(self, indices: torch.Tensor) -> int:
|
||||
self.free_slots = torch.cat([self.free_slots, indices.cpu()])
|
||||
|
||||
@@ -558,6 +558,10 @@ class ServerArgs:
|
||||
hicache_write_policy: str = "write_through"
|
||||
hicache_io_backend: str = "kernel"
|
||||
hicache_mem_layout: str = "layer_first"
|
||||
hicache_host_free_room_ratio: float = 0.0
|
||||
hicache_host_free_room_trigger_ratio: float = 0.0
|
||||
hicache_l1_free_room_ratio: float = 0.0
|
||||
hicache_l1_free_room_trigger_ratio: float = 0.0
|
||||
disable_hicache_numa_detect: bool = False
|
||||
hicache_storage_backend: Optional[str] = None
|
||||
hicache_storage_prefetch_policy: str = "best_effort"
|
||||
@@ -2898,6 +2902,7 @@ class ServerArgs:
|
||||
or self.disaggregation_decode_enable_offload_kvcache
|
||||
):
|
||||
return
|
||||
self._validate_hicache_free_room_args()
|
||||
|
||||
# Step 1: Initial layout-io compatibility normalization.
|
||||
self._resolve_layout_io_compatibility()
|
||||
@@ -2912,6 +2917,34 @@ class ServerArgs:
|
||||
if io_changed:
|
||||
self._resolve_layout_io_compatibility()
|
||||
|
||||
def _validate_hicache_free_room_args(self):
|
||||
def check_ratio(name: str, value: float):
|
||||
if value < 0 or value >= 1:
|
||||
raise ValueError(f"{name} must be in [0, 1), got {value}")
|
||||
|
||||
check_ratio(
|
||||
"hicache_host_free_room_ratio", self.hicache_host_free_room_ratio
|
||||
)
|
||||
check_ratio(
|
||||
"hicache_host_free_room_trigger_ratio",
|
||||
self.hicache_host_free_room_trigger_ratio,
|
||||
)
|
||||
check_ratio("hicache_l1_free_room_ratio", self.hicache_l1_free_room_ratio)
|
||||
check_ratio(
|
||||
"hicache_l1_free_room_trigger_ratio",
|
||||
self.hicache_l1_free_room_trigger_ratio,
|
||||
)
|
||||
if self.hicache_host_free_room_trigger_ratio > self.hicache_host_free_room_ratio:
|
||||
raise ValueError(
|
||||
"hicache_host_free_room_trigger_ratio must be <= "
|
||||
"hicache_host_free_room_ratio"
|
||||
)
|
||||
if self.hicache_l1_free_room_trigger_ratio > self.hicache_l1_free_room_ratio:
|
||||
raise ValueError(
|
||||
"hicache_l1_free_room_trigger_ratio must be <= "
|
||||
"hicache_l1_free_room_ratio"
|
||||
)
|
||||
|
||||
def _resolve_layout_io_compatibility(self):
|
||||
if (
|
||||
self.hicache_mem_layout == "page_first_direct"
|
||||
@@ -5136,6 +5169,42 @@ class ServerArgs:
|
||||
default=ServerArgs.hicache_mem_layout,
|
||||
help="The layout of host memory pool for hierarchical cache.",
|
||||
)
|
||||
parser.add_argument(
|
||||
"--hicache-host-free-room-ratio",
|
||||
type=float,
|
||||
default=ServerArgs.hicache_host_free_room_ratio,
|
||||
help=(
|
||||
"L2/host free-room target ratio after a triggered CP HiCache "
|
||||
"host eviction. 0 preserves exact-deficit eviction."
|
||||
),
|
||||
)
|
||||
parser.add_argument(
|
||||
"--hicache-host-free-room-trigger-ratio",
|
||||
type=float,
|
||||
default=ServerArgs.hicache_host_free_room_trigger_ratio,
|
||||
help=(
|
||||
"L2/host free-room trigger ratio. Eviction starts only when "
|
||||
"available host slots are below required plus this room."
|
||||
),
|
||||
)
|
||||
parser.add_argument(
|
||||
"--hicache-l1-free-room-ratio",
|
||||
type=float,
|
||||
default=ServerArgs.hicache_l1_free_room_ratio,
|
||||
help=(
|
||||
"L1/device owner-lane free-room target ratio after a triggered "
|
||||
"CP HiCache owner-lane eviction."
|
||||
),
|
||||
)
|
||||
parser.add_argument(
|
||||
"--hicache-l1-free-room-trigger-ratio",
|
||||
type=float,
|
||||
default=ServerArgs.hicache_l1_free_room_trigger_ratio,
|
||||
help=(
|
||||
"L1/device owner-lane free-room trigger ratio. Eviction starts "
|
||||
"only when owner-lane pages are below required plus this room."
|
||||
),
|
||||
)
|
||||
parser.add_argument(
|
||||
"--disable-hicache-numa-detect",
|
||||
action="store_true",
|
||||
|
||||
Reference in New Issue
Block a user