From c3b683fb90517afffb2172d404041bb1967b159c Mon Sep 17 00:00:00 2001 From: MadeBy561 Date: Tue, 1 Sep 2026 18:28:09 -0400 Subject: [PATCH 1/4] perf(glm5next): L2 weight prefetch for SM120 decode At decode batch sizes the dense projections of GLM-5.3-Flash are memory-bandwidth-bound streams of weights from device memory (GDDR7 on RTX PRO 6000 Blackwell): in_proj_qkvgfab is 51.5 MB per GPU at TP4 and its cuBLAS kernel runs 29.5 us at M=8 = 1.75 TB/s, the card's bandwidth; with the weights resident in L2 the same kernel runs 11 us. Meanwhile the all-reduces, mHC, routing chain and small kernels leave device memory idle for roughly half of every layer. This change issues cp.async.bulk.prefetch.L2 with an evict_last cache policy (CuTe DSL kernel, inline PTX) for the upcoming dense weights on a side stream inside those idle windows, sized so the fills finish before the routed-expert stream starts. Windows per decoder layer: A inside attention after the first projection (this layer's o_proj), B after the attention output (router weight + the next layer's first projection), C after the MoE all-reduce (the remainder). Budgets 20/35/20 MB (36 MB for the MLA window) so no fill overlaps the expert stream; the side stream is rejoined once at the end of the model forward (valid in FULL captures and eager runs). Inside breakable (PIECEWISE) captures the wrapper ends the segment at eager ops, so prefetch is skipped there. Numerics are untouched (cache hints only). The shared KDA/MLA layers only gain an optional _l2_prefetch_hook callback after their first projection; all planning lives in the GLM-5.3 model. Enabled by default on SM120, VLLM_GLM53_L2_PREFETCH=0 disables. Measured on 4x RTX PRO 6000 Blackwell Max-Q (TP4, DFlash2 K7 draft, llm_decode_bench ctx0, greedy), this exact code vs the same image without it: C1 verifier steps/s 85.06 -> 89.68 (+5.4%), Sieve coding-peak median 433 -> 449. Overlay build with identical windows: C1 85.06 -> 91.06 (+7.1%), Sieve 462; C1 trace per call: in_proj 29.9 -> 17.1 us, o_proj 11.6 -> 5.5, MLA q_b 11.7 -> 8.3. LAVD max C30/R30 29 exact / 1 near / 0 fail; Estonia max C30/R30 28/30. Co-Authored-By: Claude Fable 5.1 --- .../layers/mamba/gdn/kimi_gdn_linear_attn.py | 5 + vllm/model_executor/layers/mla.py | 5 + vllm/models/glm5next/nvidia/l2_prefetch.py | 365 ++++++++++++++++++ vllm/models/glm5next/nvidia/model.py | 88 +++++ 4 files changed, 463 insertions(+) create mode 100644 vllm/models/glm5next/nvidia/l2_prefetch.py diff --git a/vllm/model_executor/layers/mamba/gdn/kimi_gdn_linear_attn.py b/vllm/model_executor/layers/mamba/gdn/kimi_gdn_linear_attn.py index 9306834ac4bd..8891018c91dc 100644 --- a/vllm/model_executor/layers/mamba/gdn/kimi_gdn_linear_attn.py +++ b/vllm/model_executor/layers/mamba/gdn/kimi_gdn_linear_attn.py @@ -755,6 +755,11 @@ def forward( ) -> None: num_tokens = hidden_states.size(0) projected_qkvgfab = self.in_proj_qkvgfab(hidden_states)[0] + # Optional model-installed callback (e.g. GLM-5.3 L2 weight prefetch of + # o_proj while the small projections and the recurrence run). + _hook = getattr(self, "_l2_prefetch_hook", None) + if _hook is not None: + _hook(hidden_states.shape[0]) if self.use_full_rank_gate: split_sizes = [ 3 * self.local_projection_size, diff --git a/vllm/model_executor/layers/mla.py b/vllm/model_executor/layers/mla.py index f1b032858666..f64fb2b86bb9 100644 --- a/vllm/model_executor/layers/mla.py +++ b/vllm/model_executor/layers/mla.py @@ -194,6 +194,11 @@ def forward( k_pe = k_pe.unsqueeze(1) q = q_proj_layer(q_proj_input)[0] + # Optional model-installed callback (e.g. GLM-5.3 L2 weight prefetch of + # o_proj while the indexer, rope and attention core run). + _hook = getattr(self, "_l2_prefetch_hook", None) + if _hook is not None: + _hook(hidden_states.shape[0]) heads = self.num_heads if self.dcp_q_replicate: heads *= q_proj_layer.group_size diff --git a/vllm/models/glm5next/nvidia/l2_prefetch.py b/vllm/models/glm5next/nvidia/l2_prefetch.py new file mode 100644 index 000000000000..d04100288fc1 --- /dev/null +++ b/vllm/models/glm5next/nvidia/l2_prefetch.py @@ -0,0 +1,365 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""L2 weight prefetch for GLM-5.3 decode on SM120. + +At decode batch sizes (M <= 256) the dense projections are pure device-memory (GDDR7) streams +(in_proj_qkvgfab is 50 MB per GPU: 28.8 us at M=8) while the all-reduces, mHC, +routing chain and small kernels leave device memory idle for roughly half of each layer. +This module issues ``cp.async.bulk.prefetch.L2`` with an ``evict_last`` cache +policy for the *upcoming* dense weights on a side stream inside those idle +windows, sized so the fills finish before the routed-expert stream starts. +cuBLAS then reads the weights from L2 (128 MB on RTX PRO 6000 Blackwell): + + in_proj 29.9 -> 17.1 us, o_proj 11.6 -> 5.5, MLA q_b 11.7 -> 8.3 (C1 trace) + C1: 85.1 -> 91.1 verifier steps/s (+7%), Sieve coding peak 433 -> 462 tok/s + +Numerics are untouched (cache hints only). The side stream is forked per +window and rejoined once at the end of the model forward, which is valid in +FULL cudagraph captures and eager runs; inside a *breakable* (PIECEWISE) +capture the wrapper ends the capture segment at eager ops, so prefetch is +skipped there. + +Windows per decoder layer (C1, M=8): + A inside attention after the first projection: this layer's o_proj + (optionally + a head slice of the next layer's first projection) + B after the attention output: router weight + next layer's first projection + C after the MoE all-reduce: the remainder + +Environment: + VLLM_GLM53_L2_PREFETCH=0 disable (default: on for SM120) + VLLM_GLM53_L2_PREFETCH_MAX_TOKENS only prefetch for batches up to this (256) + VLLM_GLM53_L2_PREFETCH_BUDGET_{A,B,C,A_MLA}_MB per-window fill budgets + VLLM_GLM53_L2_PREFETCH_A_NEXT_MB next-layer head bytes carried in window A +""" + +from __future__ import annotations + +import os + +import torch + +from vllm.logger import init_logger + +logger = init_logger(__name__) + +_MIN_BYTES = 64 * 1024 +_CHUNK_BYTES = 4096 +_GRID = 16 +_BLOCK = 128 +_MAX_TOKENS = int(os.getenv("VLLM_GLM53_L2_PREFETCH_MAX_TOKENS", "256")) + + +def _mb(name: str, default: str) -> int: + return int(float(os.getenv(name, default)) * 1e6) + + +BUDGET_A = _mb("VLLM_GLM53_L2_PREFETCH_BUDGET_A_MB", "20") +BUDGET_B = _mb("VLLM_GLM53_L2_PREFETCH_BUDGET_B_MB", "35") +BUDGET_C = _mb("VLLM_GLM53_L2_PREFETCH_BUDGET_C_MB", "20") +BUDGET_A_MLA = _mb("VLLM_GLM53_L2_PREFETCH_BUDGET_A_MLA_MB", "36") +# Measured neutral-to-negative on C1 (fills overlap the KDA core); off by default. +A_NEXT_BYTES = _mb("VLLM_GLM53_L2_PREFETCH_A_NEXT_MB", "0") + +Segment = tuple[str, int, int] # (name, ptr, bytes) + + +def _platform_enabled() -> bool: + raw = os.getenv("VLLM_GLM53_L2_PREFETCH") + if raw is not None: + return raw != "0" + if not torch.cuda.is_available(): + return False + return torch.cuda.get_device_capability() == (12, 0) + + +ENABLED = _platform_enabled() + + +# --------------------------------------------------------------------------- +# CuTe DSL kernel: bulk L2 prefetch with an evict_last policy +# --------------------------------------------------------------------------- + +try: # CuTe DSL is optional; the feature silently disables without it. + import cutlass + import cutlass.cute as cute + from cuda.bindings.driver import CUstream + from cutlass._mlir import ir as _ir + from cutlass._mlir.dialects import llvm as _llvm + from cutlass.cutlass_dsl import dsl_user_op + + _CUTE_OK = True +except Exception: # noqa: BLE001 + _CUTE_OK = False + + +if _CUTE_OK: + + @dsl_user_op + def _createpolicy_evict_last(*, loc=None, ip=None): + i64 = _ir.IntegerType.get_signless(64) + res = _llvm.inline_asm( + i64, + [], + "createpolicy.fractional.L2::evict_last.b64 $0, 1.0;", + "=l", + has_side_effects=False, + loc=loc, + ip=ip, + ) + return cutlass.Int64(res) + + @dsl_user_op + def _bulk_prefetch_l2(addr, size, policy, *, loc=None, ip=None): + i32 = _ir.IntegerType.get_signless(32) + _llvm.inline_asm( + i32, + [ + addr.ir_value(loc=loc, ip=ip), + size.ir_value(loc=loc, ip=ip), + policy.ir_value(loc=loc, ip=ip), + ], + "cp.async.bulk.prefetch.L2.global.L2::cache_hint [$0], $1, $2; mov.u32 $3, 0;", + "l,r,l,=r", + has_side_effects=True, + loc=loc, + ip=ip, + ) + + class L2PrefetchKernel: + """grid x block threads walk the (ptr, bytes) segment table and issue + one bulk L2 prefetch per 4 KB chunk (fire-and-forget).""" + + def __init__(self, chunk_bytes: int = _CHUNK_BYTES, grid: int = _GRID, block: int = _BLOCK): + self.chunk_bytes = int(chunk_bytes) + self.grid = int(grid) + self.block = int(block) + + @cute.jit + def __call__(self, gSegs: cute.Tensor, stream: CUstream) -> None: + self.kernel(gSegs).launch( + grid=[self.grid, 1, 1], + block=[self.block, 1, 1], + stream=stream, + ) + + @cute.kernel + def kernel(self, gSegs: cute.Tensor) -> None: + tidx, _, _ = cute.arch.thread_idx() + bidx, _, _ = cute.arch.block_idx() + chunk: cutlass.Constexpr = self.chunk_bytes + stride = cutlass.Int64(self.grid * self.block) + tid = cutlass.Int64(bidx) * cutlass.Int64(self.block) + cutlass.Int64(tidx) + nseg = cutlass.Int32(cute.size(gSegs, mode=[0]) // 2) + policy = _createpolicy_evict_last() + s = cutlass.Int32(0) + while s < nseg: + base = cutlass.Int64(gSegs[2 * s]) + nbytes = cutlass.Int64(gSegs[2 * s + 1]) + nchunks = (nbytes + cutlass.Int64(chunk - 1)) // cutlass.Int64(chunk) + c = tid + while c < nchunks: + off = c * cutlass.Int64(chunk) + rem = nbytes - off + size = cutlass.Int32(chunk) + if rem < cutlass.Int64(chunk): + size = cutlass.Int32(rem) + size = (size // cutlass.Int32(16)) * cutlass.Int32(16) + if size > cutlass.Int32(0): + _bulk_prefetch_l2(base + off, size, policy) + c = c + stride + s = s + cutlass.Int32(1) + + +_compiled = None +_compile_failed = False + + +def _get_launcher(): + """Compile once (shape-dynamic segment table); returns a callable + (segs_tensor, cuda_stream_handle) -> None, or None if unavailable.""" + global _compiled, _compile_failed + if _compiled is not None: + return _compiled + if _compile_failed: + return None + if not _CUTE_OK: + _compile_failed = True + logger.warning("[l2_prefetch] disabled: CuTe DSL not available") + return None + try: + from quack.compile_utils import make_fake_tensor + + n = cute.sym_int(divisibility=2) + fake = make_fake_tensor(cutlass.Int64, (n,), divisibility=2) + stream = CUstream(torch.cuda.current_stream().cuda_stream) + compiled = cute.compile(L2PrefetchKernel(), fake, stream, options="--enable-tvm-ffi") + + def launch(segs: torch.Tensor, stream_handle: int) -> None: + compiled(segs, CUstream(stream_handle)) + + _compiled = launch + logger.info( + "[l2_prefetch] CuTe kernel ready (grid=%d block=%d chunk=%d; budgets A/B/C/A_mla=%.0f/%.0f/%.0f/%.0f MB)", + _GRID, _BLOCK, _CHUNK_BYTES, BUDGET_A / 1e6, BUDGET_B / 1e6, BUDGET_C / 1e6, BUDGET_A_MLA / 1e6, + ) + return _compiled + except Exception as exc: # noqa: BLE001 + _compile_failed = True + logger.warning("[l2_prefetch] disabled: kernel compile failed: %s", exc) + return None + + +# --------------------------------------------------------------------------- +# Planning helpers +# --------------------------------------------------------------------------- + + +def segments_of(module: torch.nn.Module, prefix: str = "", skip: tuple[str, ...] = ()) -> list[Segment]: + """Large, contiguous CUDA parameters/buffers under ``module``.""" + out: list[Segment] = [] + seen: set[int] = set() + + def add(name: str, t: torch.Tensor) -> None: + if not isinstance(t, torch.Tensor) or not t.is_cuda or not t.is_contiguous(): + return + nbytes = t.numel() * t.element_size() + ptr = t.data_ptr() + if nbytes < _MIN_BYTES or ptr % 16 != 0 or ptr in seen: + return + if any(s in name for s in skip): + return + seen.add(ptr) + out.append((name, ptr, nbytes)) + + for name, p in module.named_parameters(): + add(f"{prefix}{name}", p.data) + for m_name, m in module.named_modules(): + for attr in ("W_UK_T", "W_UV", "W_UK", "W_K", "W_V"): + t = getattr(m, attr, None) + if isinstance(t, torch.Tensor): + add(f"{prefix}{m_name}.{attr}", t) + return out + + +def take_budget(segments: list[Segment], budget: int) -> tuple[list[Segment], list[Segment]]: + """Split segments into (fits in budget, remainder); the straddling segment + is sliced so no byte is prefetched twice.""" + taken: list[Segment] = [] + rest: list[Segment] = [] + used = 0 + for name, ptr, nbytes in segments: + if used >= budget: + rest.append((name, ptr, nbytes)) + continue + room = budget - used + if nbytes <= room: + taken.append((name, ptr, nbytes)) + used += nbytes + else: + head = room - (room % _CHUNK_BYTES) + if head > 0: + taken.append((name + "[head]", ptr, head)) + used += head + rest.append((name + "[tail]", ptr + head, nbytes - head)) + return taken, rest + + +class L2PrefetchPlan: + """Device-resident [ptr, bytes] table for one prefetch launch.""" + + __slots__ = ("segs", "nseg", "total_bytes", "names") + + def __init__(self, segments: list[Segment], device: torch.device): + segments = [s for s in segments if s[2] >= 16] + self.nseg = len(segments) + self.total_bytes = sum(s[2] for s in segments) + self.names = [f"{n}:{b / 1e6:.1f}MB" for n, _, b in segments] + flat = [v for _, ptr, nbytes in segments for v in (ptr, nbytes)] + self.segs = torch.tensor(flat, dtype=torch.int64, device=device) if flat else None + + def describe(self) -> str: + return f"{self.nseg} segs {self.total_bytes / 1e6:.1f} MB: " + ", ".join(self.names[:8]) + + +def make_plan(segments: list[Segment], budget: int, device: torch.device) -> tuple[L2PrefetchPlan | None, list[Segment]]: + taken, rest = take_budget(segments, budget) + plan = L2PrefetchPlan(taken, device) + return (plan if plan.nseg else None), rest + + +# --------------------------------------------------------------------------- +# Runtime +# --------------------------------------------------------------------------- + + +def _prefetch_allowed() -> bool: + """False inside a breakable (PIECEWISE) capture, where eager breaks end + the capture segment and an un-joined side stream would be illegal.""" + try: + from vllm.compilation.breakable_cudagraph import BreakableCUDAGraphCapture + + cap = BreakableCUDAGraphCapture.current() + if cap is not None and bool(getattr(cap, "_capturing", False)): + from vllm.config import CUDAGraphMode + from vllm.forward_context import get_forward_context, is_forward_context_available + + if is_forward_context_available(): + return get_forward_context().cudagraph_runtime_mode == CUDAGraphMode.FULL + return False + except Exception: # noqa: BLE001 + return not torch.cuda.is_current_stream_capturing() + return True + + +class L2Prefetcher: + """One side stream per device: issue() forks it off the current stream and + launches the prefetch; join() (model end) rejoins.""" + + _instances: dict[int, "L2Prefetcher"] = {} + + def __init__(self, device: torch.device): + self.device = device + self.side = torch.cuda.Stream(device=device) + self.pending = False + + @classmethod + def get(cls, device: torch.device | None = None) -> "L2Prefetcher": + idx = device.index if device is not None and device.index is not None else torch.cuda.current_device() + inst = cls._instances.get(idx) + if inst is None: + inst = cls(torch.device("cuda", idx)) + cls._instances[idx] = inst + return inst + + def issue(self, plan: L2PrefetchPlan | None, num_tokens: int) -> None: + if plan is None or plan.segs is None or num_tokens > _MAX_TOKENS: + return + launch = _get_launcher() + if launch is None or not _prefetch_allowed(): + return + main = torch.cuda.current_stream(self.device) + self.side.wait_stream(main) + launch(plan.segs, self.side.cuda_stream) + self.pending = True + + def join(self) -> None: + if not self.pending: + return + torch.cuda.current_stream(self.device).wait_stream(self.side) + self.pending = False + + +def issue(plan: L2PrefetchPlan | None, num_tokens: int) -> None: + if ENABLED and plan is not None and plan.segs is not None: + L2Prefetcher.get(plan.segs.device).issue(plan, num_tokens) + + +def join_all() -> None: + if ENABLED: + L2Prefetcher.get(None).join() + + +__all__ = [ + "ENABLED", "BUDGET_A", "BUDGET_B", "BUDGET_C", "BUDGET_A_MLA", "A_NEXT_BYTES", + "L2PrefetchPlan", "L2Prefetcher", "segments_of", "take_budget", "make_plan", "issue", "join_all", +] diff --git a/vllm/models/glm5next/nvidia/model.py b/vllm/models/glm5next/nvidia/model.py index ae3fb6a73eac..8386ba3b39d4 100644 --- a/vllm/models/glm5next/nvidia/model.py +++ b/vllm/models/glm5next/nvidia/model.py @@ -87,6 +87,8 @@ from vllm.transformers_utils.configs.glm5_next import Glm5NextConfig from vllm.utils.b12x import get_b12x_mhc +from . import l2_prefetch as _l2pf + from .attention import Glm5NextMLAAttention from .kda import Glm5NextLinearAttention from .multimodal import ( @@ -415,6 +417,12 @@ def __init__( self.input_layernorm = RMSNorm(config.hidden_size, eps=config.rms_norm_eps) # Cached for the hot forward path (isinstance per layer per step). self._mlp_is_moe = isinstance(self.mlp, Glm5NextMoE) + # L2 weight prefetch (see l2_prefetch.py); plans are built lazily on + # the first forward, after weights are loaded and post-processed. + object.__setattr__(self, "_l2pf_next", None) + self._l2pf_ready = False + self._l2pf_plan_b = None + self._l2pf_plan_c = None # In SP, the attention output projection leaves a partial sum; the # decoder-layer reduce_scatter after attention completes it (DSv4 pattern). # MTP layers use the non-mHC path which has no sp_reduce_scatter, so @@ -525,6 +533,8 @@ def forward( # hc_post inputs (its ffn-pre outputs); when present, fuse that # hc_post with this layer's attn hc_pre into one kernel (inter-layer # fusion). Layer 0 has no incoming state -> standalone hc_pre. + if _l2pf.ENABLED and not self._l2pf_ready: + self._l2pf_build_plans() x = hidden_states if post is None: if self._b12x_mhc is not None: @@ -575,6 +585,10 @@ def forward( if self.is_sequence_parallel: x = sp_reduce_scatter(x) + # L2 prefetch window B: all-reduce + mHC + routing chain leave device memory idle. + if _l2pf.ENABLED: + _l2pf.issue(self._l2pf_plan_b, x.shape[0]) + # Fuse post-attn hc_post + pre-FFN hc_pre (+ RMSNorm) into one kernel. residual, post, comb, x = self.hc_fused_post_pre( x, @@ -597,6 +611,10 @@ def forward( # mHC end. The last mHC layer materializes its final hc_post (nothing # to fuse with) then contracts; every other layer defers its hc_post to # the next layer's fused pre, returning the state. + # L2 prefetch window C: MoE all-reduce + next mHC leave device memory idle. + if _l2pf.ENABLED: + _l2pf.issue(self._l2pf_plan_c, x.shape[0]) + if self.layer_idx == self.num_hidden_layers - 1: x = self.hc_post(x, residual, post, comb) x = hc_contract(x, self.n) @@ -604,6 +622,67 @@ def forward( return x, residual, post, comb + # ---- L2 weight prefetch planning (see l2_prefetch.py) ------------------- + _ATTN_SKIP = ("o_proj", "kv_b_proj", "indexer.index_kpool", "indexer.weights_proj") + + def _l2pf_build_plans(self) -> None: + self._l2pf_ready = True + try: + device = next(self.parameters()).device + segs_b: list[_l2pf.Segment] = [] + if self._mlp_is_moe: + # Router weight is read right after the post-attention mHC. + segs_b += _l2pf.segments_of(self.mlp.gate, "mlp.gate.") + # Dense-MLP layers (first 3): their 75 MB MLP weights are consumed + # right after attention with no idle window -> never prefetched. + nxt = self._l2pf_next + nxt_segs: list[_l2pf.Segment] = [] + if nxt is not None: + # The next layer's o_proj (and MLA kv_b, unused at decode) are + # prefetched inside that layer's own attention window (A). + nxt_segs = _l2pf.segments_of( + nxt.self_attn, f"L{nxt.layer_idx}.self_attn.", skip=self._ATTN_SKIP + ) + # Window A of this layer's attention: its o_proj plus a head slice + # of the next layer's first projection (it survives the expert + # stream and shortens window C so the tail is resident in time). + segs_a = _l2pf.segments_of(self.self_attn.o_proj, "o_proj.") + attn_impl = getattr(getattr(self.self_attn, "mla_attn", None), "impl", None) + for attr in ("W_UK_T", "W_UV"): + t = getattr(attn_impl, attr, None) if attn_impl is not None else None + if isinstance(t, torch.Tensor) and t.is_cuda and t.is_contiguous(): + segs_a.append((f"impl.{attr}", t.data_ptr(), t.numel() * t.element_size())) + is_mla = hasattr(self.self_attn, "mla_attn") + if nxt_segs and _l2pf.A_NEXT_BYTES > 0 and not is_mla: + head_a, rest_first = _l2pf.take_budget(nxt_segs[:1], _l2pf.A_NEXT_BYTES) + segs_a += head_a + nxt_segs = rest_first + nxt_segs[1:] + plan_a, _ = _l2pf.make_plan(segs_a, _l2pf.BUDGET_A_MLA if is_mla else _l2pf.BUDGET_A, device) + plan_b, rest = _l2pf.make_plan(segs_b + nxt_segs, _l2pf.BUDGET_B, device) + plan_c, dropped = _l2pf.make_plan(rest, _l2pf.BUDGET_C, device) + self._l2pf_plan_b = plan_b + self._l2pf_plan_c = plan_c + # Window A fires inside the attention layer, right after its first + # projection (hook in the shared KDA / MLA layers). + target = self.self_attn.mla_attn if is_mla else self.self_attn + if plan_a is not None: + object.__setattr__( + target, "_l2_prefetch_hook", lambda n, p=plan_a: _l2pf.issue(p, n) + ) + if self.layer_idx in (0, 3, 4) or self.layer_idx == self.num_hidden_layers - 1: + logger.info( + "[l2_prefetch] layer %d A: %s | B: %s | C: %s | dropped %.1f MB", + self.layer_idx, + plan_a.describe() if plan_a else "-", + plan_b.describe() if plan_b else "-", + plan_c.describe() if plan_c else "-", + sum(s[2] for s in dropped) / 1e6, + ) + except Exception as exc: # noqa: BLE001 + logger.warning("[l2_prefetch] layer %d plan failed: %s", self.layer_idx, exc) + self._l2pf_plan_b = None + self._l2pf_plan_c = None + def hc_pre( self, x: torch.Tensor, @@ -747,6 +826,12 @@ def get_layer(prefix: str): # The active slice is fixed after construction; cache it so forward # doesn't rebuild the slice (a fresh list) every step. self._active_layers = self.layers[self.start_layer : self.end_layer] + # L2 prefetch chain: each layer prefetches its successor's projections; + # the last layer prefetches the first layer's for the next step. + # (object.__setattr__ keeps the link out of the nn.Module registry.) + active = list(self._active_layers) + for i, layer in enumerate(active): + object.__setattr__(layer, "_l2pf_next", active[(i + 1) % len(active)]) if get_pp_group().is_last_rank: self.norm = RMSNorm(config.hidden_size, eps=config.rms_norm_eps) @@ -861,6 +946,9 @@ def forward( if self.is_sequence_parallel: hidden_states = sp_all_gather(hidden_states)[:full_num_tokens] + # Rejoin the L2 prefetch side stream once per forward (capture safety). + _l2pf.join_all() + hidden_states = self.norm(hidden_states) if aux_hidden_states: return hidden_states, aux_hidden_states From 439bd36eb3bba87608d3617e25b6ca23817eb213 Mon Sep 17 00:00:00 2001 From: MadeBy561 Date: Tue, 1 Sep 2026 18:39:28 -0400 Subject: [PATCH 2/4] fix(glm5next): L2 prefetch review follow-ups - ENABLED is a disable-only override: never on without CUDA. - join_all() is a no-op when nothing was issued and joins every device's side stream. - Join the prefetch side stream before the pipeline-parallel early return so a capture never ends with a forked stream on a non-last PP rank. Co-Authored-By: Claude Fable 5.1 --- vllm/models/glm5next/nvidia/l2_prefetch.py | 12 ++++++++---- vllm/models/glm5next/nvidia/model.py | 8 +++++--- 2 files changed, 13 insertions(+), 7 deletions(-) diff --git a/vllm/models/glm5next/nvidia/l2_prefetch.py b/vllm/models/glm5next/nvidia/l2_prefetch.py index d04100288fc1..c24534adb5f5 100644 --- a/vllm/models/glm5next/nvidia/l2_prefetch.py +++ b/vllm/models/glm5next/nvidia/l2_prefetch.py @@ -64,11 +64,12 @@ def _mb(name: str, default: str) -> int: def _platform_enabled() -> bool: + """Disable-only override: the feature never turns on without CUDA.""" + if not torch.cuda.is_available(): + return False raw = os.getenv("VLLM_GLM53_L2_PREFETCH") if raw is not None: return raw != "0" - if not torch.cuda.is_available(): - return False return torch.cuda.get_device_capability() == (12, 0) @@ -355,8 +356,11 @@ def issue(plan: L2PrefetchPlan | None, num_tokens: int) -> None: def join_all() -> None: - if ENABLED: - L2Prefetcher.get(None).join() + """Rejoin every pending prefetch side stream (no-op when nothing was issued).""" + if not ENABLED: + return + for inst in list(L2Prefetcher._instances.values()): + inst.join() __all__ = [ diff --git a/vllm/models/glm5next/nvidia/model.py b/vllm/models/glm5next/nvidia/model.py index 8386ba3b39d4..8c3fc6e29384 100644 --- a/vllm/models/glm5next/nvidia/model.py +++ b/vllm/models/glm5next/nvidia/model.py @@ -936,6 +936,11 @@ def forward( aux_hidden_state = sp_all_gather(aux_hidden_state)[:full_num_tokens] aux_hidden_states.append(aux_hidden_state) + # Rejoin the L2 prefetch side stream once per forward, before any + # early return, so a CUDA-graph capture never ends with a forked + # side stream (capture safety on every PP rank). + _l2pf.join_all() + if not get_pp_group().is_last_rank: # Pipeline parallelism is rejected because post/comb are the # deferred mHC state and must be propagated across rank boundaries. @@ -946,9 +951,6 @@ def forward( if self.is_sequence_parallel: hidden_states = sp_all_gather(hidden_states)[:full_num_tokens] - # Rejoin the L2 prefetch side stream once per forward (capture safety). - _l2pf.join_all() - hidden_states = self.norm(hidden_states) if aux_hidden_states: return hidden_states, aux_hidden_states From 72c39c091f50aa4ef975793279776113e324eba2 Mon Sep 17 00:00:00 2001 From: MadeBy561 Date: Tue, 1 Sep 2026 19:29:45 -0400 Subject: [PATCH 3/4] perf(glm5next): fire L2 prefetch windows before the all-reduces Windows B and C now start right before the attention-output and MoE all-reduces (optional _l2_prefetch_pre_reduce_hook on RowParallelLinear and the MoE runner, installed by the GLM-5.3 model), which adds the reduction time to each idle window; budgets become B 50 / C 15 MB. The in-forward issue points remain as the fallback when no hook target exists. C1 trace vs the post-reduce windows: MoE kernel 79 -> 71 us/layer (fills no longer spill into the expert stream), mHC finalize 7.8 -> 4.8 us; thermally matched A/B/A at C1: 89.0 / 88.3 / 89.4 verifier steps/s (+1%), Sieve coding-peak median 462 -> 477 (max 504). Co-Authored-By: Claude Fable 5.1 --- .../layers/fused_moe/runner/moe_runner.py | 5 +++++ vllm/model_executor/layers/linear.py | 5 +++++ vllm/models/glm5next/nvidia/l2_prefetch.py | 11 ++++++---- vllm/models/glm5next/nvidia/model.py | 22 +++++++++++++++---- 4 files changed, 35 insertions(+), 8 deletions(-) diff --git a/vllm/model_executor/layers/fused_moe/runner/moe_runner.py b/vllm/model_executor/layers/fused_moe/runner/moe_runner.py index 66ad630e680b..bafe4da2c8f1 100644 --- a/vllm/model_executor/layers/fused_moe/runner/moe_runner.py +++ b/vllm/model_executor/layers/fused_moe/runner/moe_runner.py @@ -487,6 +487,11 @@ def _maybe_reduce_final_output( if output_is_reduced is None: output_is_reduced = self._fused_output_is_reduced + # Optional model-installed callback fired before the final all-reduce + # (GLM-5.3 L2 weight prefetch: the reduction leaves device memory idle). + _hook = getattr(self, "_l2_prefetch_pre_reduce_hook", None) + if _hook is not None: + _hook(states.shape[0]) if ( not self.moe_config.is_sequence_parallel and not self.moe_config.skip_final_all_reduce diff --git a/vllm/model_executor/layers/linear.py b/vllm/model_executor/layers/linear.py index e723d6fee413..b195fdd550ef 100644 --- a/vllm/model_executor/layers/linear.py +++ b/vllm/model_executor/layers/linear.py @@ -1656,6 +1656,11 @@ def forward( bias_ = None if (self.tp_rank > 0 or self.skip_bias_add) else self.bias output_parallel = self.quant_method.apply(self, input_parallel, bias_) + # Optional model-installed callback fired before the all-reduce (GLM-5.3 + # L2 weight prefetch: the reduction leaves device memory idle). + _hook = getattr(self, "_l2_prefetch_pre_reduce_hook", None) + if _hook is not None: + _hook(output_parallel.shape[0]) if self.reduce_results and self.tp_size > 1: output = tensor_model_parallel_all_reduce(output_parallel) else: diff --git a/vllm/models/glm5next/nvidia/l2_prefetch.py b/vllm/models/glm5next/nvidia/l2_prefetch.py index c24534adb5f5..335cb90e6351 100644 --- a/vllm/models/glm5next/nvidia/l2_prefetch.py +++ b/vllm/models/glm5next/nvidia/l2_prefetch.py @@ -22,8 +22,9 @@ Windows per decoder layer (C1, M=8): A inside attention after the first projection: this layer's o_proj (optionally + a head slice of the next layer's first projection) - B after the attention output: router weight + next layer's first projection - C after the MoE all-reduce: the remainder + B before the attention-output all-reduce: router weight + next layer's + first projection + C before the MoE all-reduce: the remainder Environment: VLLM_GLM53_L2_PREFETCH=0 disable (default: on for SM120) @@ -54,8 +55,10 @@ def _mb(name: str, default: str) -> int: BUDGET_A = _mb("VLLM_GLM53_L2_PREFETCH_BUDGET_A_MB", "20") -BUDGET_B = _mb("VLLM_GLM53_L2_PREFETCH_BUDGET_B_MB", "35") -BUDGET_C = _mb("VLLM_GLM53_L2_PREFETCH_BUDGET_C_MB", "20") +# Windows B/C fire *before* the all-reduces (hooks in RowParallelLinear and the +# MoE runner), which adds the reduction time to each idle window. +BUDGET_B = _mb("VLLM_GLM53_L2_PREFETCH_BUDGET_B_MB", "50") +BUDGET_C = _mb("VLLM_GLM53_L2_PREFETCH_BUDGET_C_MB", "15") BUDGET_A_MLA = _mb("VLLM_GLM53_L2_PREFETCH_BUDGET_A_MLA_MB", "36") # Measured neutral-to-negative on C1 (fills overlap the KDA core); off by default. A_NEXT_BYTES = _mb("VLLM_GLM53_L2_PREFETCH_A_NEXT_MB", "0") diff --git a/vllm/models/glm5next/nvidia/model.py b/vllm/models/glm5next/nvidia/model.py index 8c3fc6e29384..10623bf72a93 100644 --- a/vllm/models/glm5next/nvidia/model.py +++ b/vllm/models/glm5next/nvidia/model.py @@ -585,8 +585,8 @@ def forward( if self.is_sequence_parallel: x = sp_reduce_scatter(x) - # L2 prefetch window B: all-reduce + mHC + routing chain leave device memory idle. - if _l2pf.ENABLED: + # L2 prefetch window B (fallback issue point when no pre-reduce hook). + if _l2pf.ENABLED and not getattr(self, "_l2pf_hooked_b", False): _l2pf.issue(self._l2pf_plan_b, x.shape[0]) # Fuse post-attn hc_post + pre-FFN hc_pre (+ RMSNorm) into one kernel. @@ -611,8 +611,8 @@ def forward( # mHC end. The last mHC layer materializes its final hc_post (nothing # to fuse with) then contracts; every other layer defers its hc_post to # the next layer's fused pre, returning the state. - # L2 prefetch window C: MoE all-reduce + next mHC leave device memory idle. - if _l2pf.ENABLED: + # L2 prefetch window C (fallback issue point when no pre-reduce hook). + if _l2pf.ENABLED and not getattr(self, "_l2pf_hooked_c", False): _l2pf.issue(self._l2pf_plan_c, x.shape[0]) if self.layer_idx == self.num_hidden_layers - 1: @@ -662,6 +662,20 @@ def _l2pf_build_plans(self) -> None: plan_c, dropped = _l2pf.make_plan(rest, _l2pf.BUDGET_C, device) self._l2pf_plan_b = plan_b self._l2pf_plan_c = plan_c + # Fire windows B/C before the all-reduces (+10 us of idle window + # each): hooks on this layer's o_proj and on the MoE runner (or the + # dense MLP's down_proj). The in-forward issue points below stay + # as the fallback when a hook target is missing. + self._l2pf_hooked_b = False + self._l2pf_hooked_c = False + o_proj = getattr(self.self_attn, "o_proj", None) + if o_proj is not None and plan_b is not None and getattr(o_proj, "reduce_results", False): + object.__setattr__(o_proj, "_l2_prefetch_pre_reduce_hook", lambda n, p=plan_b: _l2pf.issue(p, n)) + self._l2pf_hooked_b = True + target_c = self.mlp.experts if self._mlp_is_moe else getattr(self.mlp, "down_proj", None) + if target_c is not None and plan_c is not None: + object.__setattr__(target_c, "_l2_prefetch_pre_reduce_hook", lambda n, p=plan_c: _l2pf.issue(p, n)) + self._l2pf_hooked_c = True # Window A fires inside the attention layer, right after its first # projection (hook in the shared KDA / MLA layers). target = self.self_attn.mla_attn if is_mla else self.self_attn From 9ab68a1eb084d3d6e2ec50b0e61d88d6a4e4c2c8 Mon Sep 17 00:00:00 2001 From: MadeBy561 Date: Wed, 2 Sep 2026 02:32:36 -0400 Subject: [PATCH 4/4] perf(glm5next): size the persisting-L2 set-aside for weight prefetch The L2 weight prefetcher marks the upcoming dense projections with an evict_last policy, but evict_last lines are only retained against normal-priority traffic inside the CUDA persisting-L2 set-aside, whose default size is 0. On RTX PRO 6000 Blackwell about a third of a prefetched 50 MB projection was therefore evicted by the routed-expert stream before cuBLAS read it (in_proj 17.3 us instead of 11 us fully hot). Size the set-aside once per device from the prefetcher through cuCtxSetLimit(CU_LIMIT_PERSISTING_L2_CACHE_SIZE) on the device's primary context, before the first prefetch and therefore before any graph capture. VLLM_GLM53_L2_PREFETCH_PERSIST_MB selects "max" (default, the device maximum: 84 MB of the 128 MB L2), a megabyte value clamped to that maximum, or 0 to leave the driver default. Failures are logged and leave the prefetcher active. Cache-residency policy only: kernels, weights and numerics are unchanged. Co-Authored-By: Claude Fable 5.1 --- .../test_glm5next_l2_prefetch_persist.py | 109 ++++++++++++++++ vllm/models/glm5next/nvidia/l2_prefetch.py | 119 +++++++++++++++++- 2 files changed, 227 insertions(+), 1 deletion(-) create mode 100644 tests/models/test_glm5next_l2_prefetch_persist.py diff --git a/tests/models/test_glm5next_l2_prefetch_persist.py b/tests/models/test_glm5next_l2_prefetch_persist.py new file mode 100644 index 000000000000..23f4082f2735 --- /dev/null +++ b/tests/models/test_glm5next_l2_prefetch_persist.py @@ -0,0 +1,109 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Persisting-L2 set-aside sizing for the GLM-5.3 L2 weight prefetcher.""" + +import pytest +import torch + +from vllm.models.glm5next.nvidia import l2_prefetch as l2pf + +MAX = 84_000_000 + + +@pytest.mark.parametrize( + ("raw", "max_bytes", "expected"), + [ + ("max", MAX, MAX), + (" MAX ", MAX, MAX), + ("all", MAX, MAX), + ("40", MAX, 40_000_000), + ("40.5", MAX, 40_500_000), + ("200", MAX, MAX), # clamped to the device maximum + ("0", MAX, 0), + ("off", MAX, 0), + ("", MAX, 0), + (None, MAX, 0), + ("-5", MAX, 0), + ("bogus", MAX, 0), + ("max", 0, 0), # device without a persisting L2 set-aside + ], +) +def test_persisting_l2_request(raw, max_bytes, expected): + assert l2pf.persisting_l2_request(raw, max_bytes) == expected + + +def test_default_request_is_device_maximum(): + assert l2pf.PERSIST_L2 == "max" + + +def _driver(): + from cuda.bindings import driver as cu + + return cu + + +def _read_limit(cu, device: torch.device) -> int: + with torch.cuda.device(device): + torch.empty(1, device=device) # make the primary context current + err, value = cu.cuCtxGetLimit(cu.CUlimit.CU_LIMIT_PERSISTING_L2_CACHE_SIZE) + assert err == cu.CUresult.CUDA_SUCCESS + return int(value) + + +def _set_limit(cu, device: torch.device, value: int) -> None: + with torch.cuda.device(device): + torch.empty(1, device=device) + (err,) = cu.cuCtxSetLimit(cu.CUlimit.CU_LIMIT_PERSISTING_L2_CACHE_SIZE, value) + assert err == cu.CUresult.CUDA_SUCCESS + + +@pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA device required") +def test_configure_persisting_l2_applies_to_the_primary_context(): + cu = _driver() + device = torch.device("cuda", 0) + err, dev = cu.cuDeviceGet(0) + assert err == cu.CUresult.CUDA_SUCCESS + err, max_bytes = cu.cuDeviceGetAttribute( + cu.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_MAX_PERSISTING_L2_CACHE_SIZE, dev + ) + assert err == cu.CUresult.CUDA_SUCCESS + if max_bytes <= 0: + pytest.skip("device has no persisting L2 set-aside") + before = _read_limit(cu, device) + try: + assert l2pf.configure_persisting_l2(device, request="max") == max_bytes + assert _read_limit(cu, device) == max_bytes + + # A zero request never touches the driver state. + assert l2pf.configure_persisting_l2(device, request="0") == 0 + assert _read_limit(cu, device) == max_bytes + + # Over-sized requests are clamped to the device maximum. + assert l2pf.configure_persisting_l2(device, request="100000") == max_bytes + finally: + _set_limit(cu, device, before) + + +@pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA device required") +def test_prefetcher_applies_the_env_request_once(monkeypatch): + cu = _driver() + device = torch.device("cuda", 0) + before = _read_limit(cu, device) + calls: list[str | None] = [] + real = l2pf.configure_persisting_l2 + + def spy(dev, request=None): + calls.append(request) + return real(dev, request=request) + + monkeypatch.setattr(l2pf, "configure_persisting_l2", spy) + monkeypatch.setattr(l2pf.L2Prefetcher, "_instances", {}) + monkeypatch.setattr(l2pf, "_persisting_l2_applied", {}) + try: + first = l2pf.L2Prefetcher.get(device) + second = l2pf.L2Prefetcher.get(device) + assert first is second + assert calls == [None] + assert first.persisting_l2_bytes == _read_limit(cu, device) + finally: + _set_limit(cu, device, before) diff --git a/vllm/models/glm5next/nvidia/l2_prefetch.py b/vllm/models/glm5next/nvidia/l2_prefetch.py index 335cb90e6351..e76a1b209245 100644 --- a/vllm/models/glm5next/nvidia/l2_prefetch.py +++ b/vllm/models/glm5next/nvidia/l2_prefetch.py @@ -31,6 +31,15 @@ VLLM_GLM53_L2_PREFETCH_MAX_TOKENS only prefetch for batches up to this (256) VLLM_GLM53_L2_PREFETCH_BUDGET_{A,B,C,A_MLA}_MB per-window fill budgets VLLM_GLM53_L2_PREFETCH_A_NEXT_MB next-layer head bytes carried in window A + VLLM_GLM53_L2_PREFETCH_PERSIST_MB persisting-L2 set-aside per rank: "max" + (default, device maximum), megabytes, or 0 + +The ``evict_last`` policy only protects lines inside the CUDA persisting-L2 +set-aside, whose default size is 0. Without a set-aside roughly a third of a +prefetched 50 MB projection is evicted by the routed-expert stream before +cuBLAS reads it (in_proj 17.3 us instead of 11 us fully hot), so the prefetcher +sizes the set-aside once per device through ``cuCtxSetLimit``. This is a +cache-residency policy only; kernels and numerics are unchanged. """ from __future__ import annotations @@ -62,6 +71,10 @@ def _mb(name: str, default: str) -> int: BUDGET_A_MLA = _mb("VLLM_GLM53_L2_PREFETCH_BUDGET_A_MLA_MB", "36") # Measured neutral-to-negative on C1 (fills overlap the KDA core); off by default. A_NEXT_BYTES = _mb("VLLM_GLM53_L2_PREFETCH_A_NEXT_MB", "0") +# Persisting-L2 set-aside request: "max" = device maximum (84 MB of the 128 MB +# L2 on RTX PRO 6000 Blackwell), a number = megabytes clamped to that maximum, +# 0/"off" = leave the driver default (no set-aside). +PERSIST_L2 = os.getenv("VLLM_GLM53_L2_PREFETCH_PERSIST_MB", "max") Segment = tuple[str, int, int] # (name, ptr, bytes) @@ -291,6 +304,105 @@ def make_plan(segments: list[Segment], budget: int, device: torch.device) -> tup return (plan if plan.nseg else None), rest +# --------------------------------------------------------------------------- +# Persisting-L2 set-aside +# --------------------------------------------------------------------------- + + +def persisting_l2_request(raw: str | None, max_bytes: int) -> int: + """Bytes of persisting-L2 set-aside to request for the env value ``raw``. + + ``max`` maps to ``max_bytes`` (the device attribute), a number is + megabytes clamped to ``[0, max_bytes]``, and empty/0/off/invalid values + map to 0 (do not touch the driver default). + """ + if raw is None or max_bytes <= 0: + return 0 + value = raw.strip().lower() + if value in ("", "0", "off", "none", "false"): + return 0 + if value in ("max", "all"): + return int(max_bytes) + try: + want = int(float(value) * 1e6) + except ValueError: + logger.warning( + "[l2_prefetch] ignoring VLLM_GLM53_L2_PREFETCH_PERSIST_MB=%r " + "(expected megabytes, max or 0)", + raw, + ) + return 0 + return max(0, min(want, int(max_bytes))) + + +_persisting_l2_applied: dict[int, int] = {} + + +def configure_persisting_l2(device: torch.device, request: str | None = None) -> int: + """Size the persisting-L2 set-aside of ``device``'s primary context. + + Returns the set-aside in bytes read back from the driver, or 0 when the + request is 0, the device has no persisting L2, or the driver call failed + (logged, never raised). The env-driven call (``request=None``) is applied + once per device; an explicit ``request`` is always applied. + """ + idx = device.index if device.index is not None else torch.cuda.current_device() + if request is None and idx in _persisting_l2_applied: + return _persisting_l2_applied[idx] + raw = PERSIST_L2 if request is None else request + applied = 0 + try: + from cuda.bindings import driver as cu + + def check(result, what: str): + if result[0] != cu.CUresult.CUDA_SUCCESS: + raise RuntimeError(f"{what} failed: {result[0]}") + return result[1] if len(result) > 1 else None + + dev = check(cu.cuDeviceGet(idx), "cuDeviceGet") + attr = cu.CUdevice_attribute + l2_bytes = check( + cu.cuDeviceGetAttribute(attr.CU_DEVICE_ATTRIBUTE_L2_CACHE_SIZE, dev), + "cuDeviceGetAttribute(L2_CACHE_SIZE)", + ) + max_bytes = check( + cu.cuDeviceGetAttribute( + attr.CU_DEVICE_ATTRIBUTE_MAX_PERSISTING_L2_CACHE_SIZE, dev + ), + "cuDeviceGetAttribute(MAX_PERSISTING_L2_CACHE_SIZE)", + ) + want = persisting_l2_request(raw, max_bytes) + if want > 0: + # Operate on the primary context torch uses for this device; + # push/pop keeps the calling thread's current context unchanged. + limit = cu.CUlimit.CU_LIMIT_PERSISTING_L2_CACHE_SIZE + ctx = check(cu.cuDevicePrimaryCtxRetain(dev), "cuDevicePrimaryCtxRetain") + check(cu.cuCtxPushCurrent(ctx), "cuCtxPushCurrent") + try: + check(cu.cuCtxSetLimit(limit, want), "cuCtxSetLimit") + applied = int(check(cu.cuCtxGetLimit(limit), "cuCtxGetLimit")) + finally: + cu.cuCtxPopCurrent() + cu.cuDevicePrimaryCtxRelease(dev) + logger.info( + "[l2_prefetch] persisting L2 set-aside on cuda:%d: L2 %.0f MB, " + "max %.0f MB, requested %.0f MB, now %.0f MB", + idx, + l2_bytes / 1e6, + max_bytes / 1e6, + want / 1e6, + applied / 1e6, + ) + except Exception as exc: # noqa: BLE001 + logger.warning( + "[l2_prefetch] persisting L2 set-aside not applied on cuda:%d: %s", idx, exc + ) + applied = 0 + if request is None: + _persisting_l2_applied[idx] = applied + return applied + + # --------------------------------------------------------------------------- # Runtime # --------------------------------------------------------------------------- @@ -325,6 +437,9 @@ def __init__(self, device: torch.device): self.device = device self.side = torch.cuda.Stream(device=device) self.pending = False + # Before the first prefetch (and therefore before any graph capture): + # a context limit, not a captured operation. + self.persisting_l2_bytes = configure_persisting_l2(device) @classmethod def get(cls, device: torch.device | None = None) -> "L2Prefetcher": @@ -368,5 +483,7 @@ def join_all() -> None: __all__ = [ "ENABLED", "BUDGET_A", "BUDGET_B", "BUDGET_C", "BUDGET_A_MLA", "A_NEXT_BYTES", - "L2PrefetchPlan", "L2Prefetcher", "segments_of", "take_budget", "make_plan", "issue", "join_all", + "PERSIST_L2", "L2PrefetchPlan", "L2Prefetcher", "segments_of", "take_budget", + "make_plan", "issue", "join_all", "persisting_l2_request", + "configure_persisting_l2", ]