From b3165a1d3cb8594dcf8c7497596942f17664971d Mon Sep 17 00:00:00 2001 From: yy-fighting <2351884576@qq.com> Date: Tue, 8 Sep 2026 11:38:06 +0800 Subject: [PATCH] =?UTF-8?q?feat(pro6000/GLM-5.3):=20=E9=83=A8=E7=BD=B2?= =?UTF-8?q?=E6=96=B9=E6=A1=88=E5=85=A5=E5=BA=93=EF=BC=88TP8+EAGLE=20?= =?UTF-8?q?=E7=94=9F=E4=BA=A7=E6=A0=87=E5=87=86=20/=20=E5=9C=BA=E6=99=AF?= =?UTF-8?q?=E4=BA=8C=E9=AB=98=E5=B9=B6=E5=8F=91=E5=8F=98=E4=BD=93=20/=20TP?= =?UTF-8?q?4PP2+IndexCache=20/=20E7b=20CAR=20=E5=AE=9E=E9=AA=8C=E8=A1=A5?= =?UTF-8?q?=E4=B8=81=20+=20deploy=20profiles=EF=BC=89?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit - deploy_glm53_605.sh：配置 A 生产标准（60.5 在役，全 8 台 md5 fcd9109b 一致）， 支持 MEMFRAC/STEPS/TOPK/DRAFT/CHUNK/EXTRA/RESTART 调参；场景二高并发变体 仅改 mrr32 + decode 图 bs{4,8,12,16}（KV 池 16.4 驻留上限） - deploy_glm53_optimal(_s1).sh：配置 B TP4PP2+IndexCache（场景二最优 +41~79%； s1 形态唯一差异 radix-on） - deploy_glm53_607_exp.sh + car_patch/：E7b custom-AR 1stage 补丁（cc1 decode 每步 -14~-16%，实验性仅 cc1-2 验证；补丁文件与 60.7:/root/patches md5 一致） - deploy/profiles/pro6000/：两个标准 profile（sskj.deploy 可消费），关键踩坑 与场景二变体、parser 缺口均在注释中标注 --- ...nvfp4_pro6000_sglang_tp4pp2_indexcache.env | 43 ++ .../glm53_nvfp4_pro6000_sglang_tp8eagle.env | 44 ++ .../scripts/car_patch/custom_all_reduce.py | 424 ++++++++++++++ .../car_patch/custom_all_reduce_utils.py | 519 ++++++++++++++++++ .../scripts/car_patch/prep_car_patch.sh | 27 + .../scripts/deploy_glm53_605.sh | 58 ++ .../scripts/deploy_glm53_607_exp.sh | 70 +++ .../scripts/deploy_glm53_optimal.sh | 62 +++ .../scripts/deploy_glm53_optimal_s1.sh | 59 ++ 9 files changed, 1306 insertions(+) create mode 100644 deploy/profiles/pro6000/glm53_nvfp4_pro6000_sglang_tp4pp2_indexcache.env create mode 100644 deploy/profiles/pro6000/glm53_nvfp4_pro6000_sglang_tp8eagle.env create mode 100644 experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/car_patch/custom_all_reduce.py create mode 100644 experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/car_patch/custom_all_reduce_utils.py create mode 100644 experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/car_patch/prep_car_patch.sh create mode 100644 experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/deploy_glm53_605.sh create mode 100644 experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/deploy_glm53_607_exp.sh create mode 100644 experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/deploy_glm53_optimal.sh create mode 100644 experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/deploy_glm53_optimal_s1.sh diff --git a/deploy/profiles/pro6000/glm53_nvfp4_pro6000_sglang_tp4pp2_indexcache.env b/deploy/profiles/pro6000/glm53_nvfp4_pro6000_sglang_tp4pp2_indexcache.env new file mode 100644 index 0000000..23df8c5 --- /dev/null +++ b/deploy/profiles/pro6000/glm53_nvfp4_pro6000_sglang_tp4pp2_indexcache.env @@ -0,0 +1,43 @@ +# GLM-5.3-NVFP4 SGLang TP=4 PP=2 + IndexCache deployment profile (single RTX 6000D node). +# 场景二（16k 独立输入、cc8-32、prefill 主导）最优配置：吞吐 +41~79%、cc32 TTFT 减半 +# （对比 TP8+EAGLE 生产配置，2026-09-07 真实语料基线）。 +# 可执行部署脚本：experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/deploy_glm53_optimal.sh +# +# 关键点（实测踩坑，勿随意改）: +# - index_topk_freq=4 为模型原生默认（层轴索引复用省 75% indexer，无质量损失） +# - 禁投机解码：PP2 与投机框架不兼容（已实测） +# - 本文件为场景二形态（--disable-radix-cache，独立输入无前缀复用）； +# 场景一 90% 命中对比须启用 radix（唯一差异：去掉 --disable-radix-cache， +# 见 deploy_glm53_optimal_s1.sh） +# - mem 0.85：0.90 下 cuda graph 捕获余量不足会运行时 OOM（KV 池 569,600 = TP8 的 2.06 倍） +# - --disable-custom-all-reduce：TP4 over PCIe 用自定义 AR 在本栈无收益 +# - 已知缺口：未带 --tool-call-parser glm47 --reasoning-parser glm45，质量门 6/7 +# （tool call 失败纯属参数缺失，非模型问题）；上生产必须补 parser +# - 场景一（90% 命中低并发）该配置全面劣于 TP8+EAGLE（输出吞吐 −25~−65%），勿混用选型 + +PLATFORM=pro6000 +EXPERIMENT=glm53_nvfp4_pro6000_sglang_tp4pp2_indexcache +MODEL_NAME=GLM-5.3-NVFP4 +ENGINE=sglang +RUNTIME=docker +DOCKER_IMAGE=lmsysorg/sglang:nightly-dev-20260828-daf63171 +CONTAINER_NAME=glm53-nvfp4 +MODEL_PATH=/data/hf_models/GLM-5.3-NVFP4 +SERVED_MODEL_NAME=/data/hf_models/GLM-5.3-NVFP4 +PORT=30000 +HEALTH_PATH=/health +HEALTH_WAIT_S=2400 +CONTAINER_PYTHON=python3 + +TP=4 +DP=1 + +DEVICE_VARS="CUDA_VISIBLE_DEVICES=0,1,2,3,4,5,6,7" +ENGINE_ENV="PYTHONUNBUFFERED=1 HF_HUB_OFFLINE=1 TRANSFORMERS_OFFLINE=1" + +DOCKER_FLAGS="--gpus all --ipc=host --shm-size 64g --entrypoint '' -p ${PORT}:${PORT}" +VOLUMES="/data/hf_models:/data/hf_models:ro" + +BOOTSTRAP="exec python3 -m sglang.launch_server ${LAUNCH_ARGS}" + +LAUNCH_ARGS="--model-path ${MODEL_PATH} --tp-size ${TP} --pp-size 2 --mem-fraction-static 0.85 --max-running-requests 48 --disable-radix-cache --disable-shared-experts-fusion --moe-runner-backend flashinfer_cutlass --disable-flashinfer-autotune --disable-custom-all-reduce --chunked-prefill-size 16384 --json-model-override-args '{\"index_topk_freq\": 4}' --host 0.0.0.0 --port ${PORT}" diff --git a/deploy/profiles/pro6000/glm53_nvfp4_pro6000_sglang_tp8eagle.env b/deploy/profiles/pro6000/glm53_nvfp4_pro6000_sglang_tp8eagle.env new file mode 100644 index 0000000..ab548c5 --- /dev/null +++ b/deploy/profiles/pro6000/glm53_nvfp4_pro6000_sglang_tp8eagle.env @@ -0,0 +1,44 @@ +# GLM-5.3-NVFP4 SGLang TP=8 + EAGLE deployment profile (single RTX 6000D node, 8 GPUs). +# 生产标准配置（174.1.60.5 团队自用在役，2026-09-07 定稿）。 +# 可执行部署脚本：experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/deploy_glm53_605.sh +# +# 关键点（实测踩坑，勿随意改）: +# - SM120 必需三项：--disable-shared-experts-fusion --moe-runner-backend +# flashinfer_cutlass --disable-flashinfer-autotune +# - EAGLE 4/1/5 为 cc1-2 均衡点（3/1/4、5/1/6 在 ±10% 噪声内；树式 topk>1 在 DSA 上不可用； +# EAGLE3 本 nightly 不支持 GLM-5.3）；draft 模型自动从主权重加载 +# - KV fp8_e4m3 + hicache-ratio 3 是长上下文 90% 前缀命中场景的容量前提（KV 池 276,864 token） +# - chunk 8192：16384 已 A/B 证伪（MoE 工作区 OOM 风险，且 AR 带宽受限无次数红利） +# - 容器入口必须 python3 -m sglang.launch_server（镜像 entrypoint 无 shebang） +# - 重新部署前必须等显存排空（docker rm -f 异步滞留数分钟，不等会把新 KV 池压小） +# - 场景二高并发变体（16k 独立输入 cc8-32）：仅改 --max-running-requests 32 + +# --cuda-graph-max-bs-decode 16 --cuda-graph-bs-decode 4 8 12 16（KV 池 276,864 ÷ +# 16,896/req = 16.4 驻留上限，decode 批自然 ≤16，图覆盖到 bs16 即可，bs24/32 纯耗显存） +# - 上线前后过质量门 quality_gate_605.sh（PASS=7/7） + +PLATFORM=pro6000 +EXPERIMENT=glm53_nvfp4_pro6000_sglang_tp8eagle +MODEL_NAME=GLM-5.3-NVFP4 +ENGINE=sglang +RUNTIME=docker +DOCKER_IMAGE=lmsysorg/sglang:nightly-dev-20260828-daf63171 +CONTAINER_NAME=glm53-nvfp4 +MODEL_PATH=/data/hf_models/GLM-5.3-NVFP4 +SERVED_MODEL_NAME=/data/hf_models/GLM-5.3-NVFP4 +PORT=30000 +HEALTH_PATH=/health +HEALTH_WAIT_S=2400 +CONTAINER_PYTHON=python3 + +TP=8 +DP=1 + +DEVICE_VARS="CUDA_VISIBLE_DEVICES=0,1,2,3,4,5,6,7" +ENGINE_ENV="PYTHONUNBUFFERED=1 HF_HUB_OFFLINE=1 TRANSFORMERS_OFFLINE=1" + +DOCKER_FLAGS="--gpus all --ipc=host --shm-size 64g --entrypoint '' -p ${PORT}:${PORT}" +VOLUMES="/data/hf_models:/data/hf_models:ro" + +BOOTSTRAP="exec python3 -m sglang.launch_server ${LAUNCH_ARGS}" + +LAUNCH_ARGS="--model-path ${MODEL_PATH} --tp-size ${TP} --mem-fraction-static 0.90 --max-running-requests 16 --chunked-prefill-size 8192 --max-prefill-tokens 16384 --disable-shared-experts-fusion --moe-runner-backend flashinfer_cutlass --disable-flashinfer-autotune --speculative-algorithm EAGLE --speculative-num-steps 4 --speculative-eagle-topk 1 --speculative-num-draft-tokens 5 --kv-cache-dtype fp8_e4m3 --enable-hierarchical-cache --hicache-ratio 3 --cuda-graph-max-bs-decode 8 --cuda-graph-bs-decode 1 2 3 4 6 8 --cuda-graph-max-bs-prefill 8 --context-length 270336 --reasoning-parser glm45 --tool-call-parser glm47 --host 0.0.0.0 --port ${PORT}" diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/car_patch/custom_all_reduce.py b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/car_patch/custom_all_reduce.py new file mode 100644 index 0000000..4ec6995 --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/car_patch/custom_all_reduce.py @@ -0,0 +1,424 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +# Adapted from https://github.com/vllm-project/vllm/blob/v0.6.4.post1/vllm/distributed/device_communicators/custom_all_reduce.py + +import ctypes +import logging +import os +from contextlib import contextmanager +from functools import partial +from typing import Any, List, Optional, Union + +import torch +import torch.distributed as dist +from torch.distributed import ProcessGroup + +import sglang.srt.distributed.device_communicators.custom_all_reduce_ops as ops +from sglang.srt.distributed.device_communicators.cuda_wrapper import CudaRTLibrary +from sglang.srt.distributed.device_communicators.custom_all_reduce_utils import ( + can_use_custom_all_reduce_with_nvlink, + is_weak_contiguous, +) +from sglang.srt.environ import envs +from sglang.srt.model_executor.runner_backend_utils.tc_piecewise_cuda_graph import ( + is_in_tc_piecewise_cuda_graph, +) +from sglang.srt.utils import ( + get_bool_env_var, + is_cuda, + is_hip, + is_musa, + log_info_on_rank0, +) + +_is_cuda = is_cuda() +_is_hip = is_hip() +_is_musa = is_musa() + +logger = logging.getLogger(__name__) +os.environ.setdefault("SGLANG_CUSTOM_ALLREDUCE_ALGO", "1stage") # SSKJ-PATCH: C++ dispatch no-ops when full_nvlink=False; force kernel launch + + +class CustomAllreduce: + _SUPPORTED_WORLD_SIZES = [2, 4, 6, 8] + _MAX_CAR_SIZE = 8192 * 1024 + if _is_hip: + # crossover is at 16MB buffer size for ROCm + _MAX_CAR_SIZE = 2 * 8192 * 1024 + if _is_musa: + # crossover is at 128MB buffer size for MUSA + _MAX_CAR_SIZE = 16 * 8196 * 1024 + + # max_size: max supported allreduce size + def __init__( + self, + group: ProcessGroup, + device: Union[int, str, torch.device], + max_size=_MAX_CAR_SIZE, + ) -> None: + """ + Args: + group: the process group to work on. If None, it will use the + default process group. + device: the device to bind the CustomAllreduce to. If None, + it will be bind to f"cuda:{local_rank}". + It is the caller's responsibility to make sure each communicator + is bind to a unique device, and all communicators in this group + are in the same node. + """ + self._IS_CAPTURING = False + self.disabled = True # This can be modified in-place by context manager in piecewise cuda graph runner + self.original_disabled = True # To store the original state + self.use_amd_deterministic_impl = _use_amd_deterministic_impl() + + if not ops.IS_CUSTOM_AR_AVAILABLE: + # disable because of missing custom allreduce library + # e.g. in a non-cuda environment + return + + rank = dist.get_rank(group=group) + world_size = dist.get_world_size(group=group) + + if isinstance(device, int): + device = torch.device(f"cuda:{device}") + elif isinstance(device, str): + device = torch.device(device) + # now `device` is a `torch.device` object + assert isinstance(device, torch.device) + self.device = device + full_nvlink = can_use_custom_all_reduce_with_nvlink( + group=group, + device=device, + supported_world_size=self._SUPPORTED_WORLD_SIZES, + cls_name="CustomAllreduce", + ) + if full_nvlink is None: + return # fail to get nvlink status + + self.group = group + self.max_size = max_size + self.rank = rank + self.world_size = world_size + self.full_nvlink = full_nvlink + + if not _is_hip: + # Buffers memory are owned by this Python class and passed to C++. + # Meta data composes of two parts: meta data for synchronization and a + # temporary buffer for storing intermediate allreduce results. + self.meta_ptrs = self.create_shared_buffer( + ops.meta_size() + max_size, group=group + ) + # This is a pre-registered IPC buffer. In eager mode, input tensors + # are first copied into this buffer before allreduce is performed + self.buffer_ptrs = self.create_shared_buffer(max_size, group=group) + # This is a buffer for storing the tuples of pointers pointing to + # IPC buffers from all ranks. Each registered tuple has size of + # 8*world_size bytes where world_size is at most 8. Allocating 8MB + # is enough for 131072 such tuples. The largest model I've seen only + # needs less than 10000 of registered tuples. + self.rank_data = torch.empty( + max_size, dtype=torch.uint8, device=self.device + ) + self._ptr = ops.init_custom_ar( + self.meta_ptrs, self.rank_data, rank, self.full_nvlink + ) + ops.register_buffer(self._ptr, self.buffer_ptrs) + else: + # meta data buffers need to be "uncached" for signal on MI200 + self.meta = ops.allocate_meta_buffer(ops.meta_size() + max_size) + self.buffer = torch.empty(max_size, dtype=torch.uint8, device=self.device) + handle = ops.get_meta_buffer_ipc_handle(self.meta) + shard_data = ( + bytes(handle), # ipc handle to base ptr + 0, # offset of base ptr + ) + handles, offsets = self._gather_ipc_meta(shard_data) + self.rank_data = torch.empty( + max_size, dtype=torch.uint8, device=self.device + ) + self._ptr = ops.init_custom_ar( + self.meta, self.rank_data, handles, offsets, rank, self.full_nvlink + ) + self.register_buffer(self.buffer) + + self.disabled = False + self.original_disabled = False # Ensure original_disabled == disabled + logger.warning(f"SSKJ_CAR_PATCH_ACTIVE ws={self.world_size} full_nvlink={self.full_nvlink}") + self.tms_cudagraph = envs.SGLANG_MEMORY_SAVER_CUDA_GRAPH.get() + + @staticmethod + def create_shared_buffer( + size_in_bytes: int, group: Optional[ProcessGroup] = None + ) -> List[int]: + """ + Creates a shared buffer and returns a list of pointers + representing the buffer on all processes in the group. + """ + lib = CudaRTLibrary() + pointer = lib.cudaMalloc(size_in_bytes) + if _is_musa: + lib.cudaMemset(pointer, 0, size_in_bytes) + handle = lib.cudaIpcGetMemHandle(pointer) + world_size = dist.get_world_size(group=group) + rank = dist.get_rank(group=group) + handles = [None] * world_size + dist.all_gather_object(handles, handle, group=group) + + pointers: List[int] = [] + for i, h in enumerate(handles): + if i == rank: + pointers.append(pointer.value) # type: ignore + else: + pointers.append(lib.cudaIpcOpenMemHandle(h).value) # type: ignore + + return pointers + + @staticmethod + def free_shared_buffer( + pointers: List[int], group: Optional[ProcessGroup] = None + ) -> None: + rank = dist.get_rank(group=group) + lib = CudaRTLibrary() + lib.cudaFree(ctypes.c_void_p(pointers[rank])) + + @contextmanager + def capture(self): + """ + The main responsibility of this context manager is the + `register_graph_buffers` call at the end of the context. + It records all the buffer addresses used in the CUDA graph. + """ + try: + self._IS_CAPTURING = True + yield + finally: + self._IS_CAPTURING = False + if not self.disabled: + self.register_graph_buffers() + + def _get_ipc_meta(self, inp: torch.Tensor): + # _share_cuda_() doesn't accept meta buffer not allocated from + # PyTorch cache allocator, use direct HIP call to get IPC handle + handle = ops.get_meta_buffer_ipc_handle(inp) + shard_data = ( + bytes(handle), # ipc handle to base ptr + 0, # offset of base ptr + ) + return self._gather_ipc_meta(shard_data) + + def _gather_ipc_meta(self, shard_data): + # Note: don't use `[[None]] * self.world_size` here + # because it will create a list of the same reference + all_data: List[Optional[Any]] = [[None] for i in range(self.world_size)] + all_data[self.rank][0] = shard_data + + ranks = dist.get_process_group_ranks(group=self.group) + ranks.sort() + for i, rank in enumerate(ranks): + dist.broadcast_object_list( + all_data[i], src=rank, group=self.group, device="cpu" + ) + + # we cannot directly use `dist.all_gather_object` here + # because it is incompatible with `gloo` backend under inference mode. + # see https://github.com/pytorch/pytorch/issues/126032 for details. + + handles = [] + offsets = [] + for i in range(len(all_data)): + handles.append(all_data[i][0][0]) # type: ignore + offsets.append(all_data[i][0][1]) # type: ignore + return handles, offsets + + def register_buffer(self, inp: torch.Tensor): + handles, offsets = self._get_ipc_meta(inp) + ops.register_buffer(self._ptr, inp, handles, offsets) + + def register_graph_buffers(self): + if _is_hip: + handle, offset = ops.get_graph_buffer_ipc_meta(self._ptr) + handles, offsets = self._gather_ipc_meta((bytes(handle), offset)) + log_info_on_rank0(logger, f"Registering {len(offset)} cuda graph addresses") + ops.register_graph_buffers(self._ptr, handles, offsets) + else: + handle, offset = ops.get_graph_buffer_ipc_meta(self._ptr) + log_info_on_rank0(logger, f"Registering {len(offset)} cuda graph addresses") + # We cannot directly use `dist.all_gather_object` here + # because it is incompatible with `gloo` backend under inference mode. + # see https://github.com/pytorch/pytorch/issues/126032 for details. + all_data = [ + [None, None] for _ in range(dist.get_world_size(group=self.group)) + ] + all_data[self.rank] = [handle, offset] + ranks = sorted(dist.get_process_group_ranks(group=self.group)) + for i, rank in enumerate(ranks): + dist.broadcast_object_list( + all_data[i], src=rank, group=self.group, device="cpu" + ) + # Unpack list of tuples to tuple of lists. + handles = [d[0] for d in all_data] # type: ignore + offsets = [d[1] for d in all_data] # type: ignore + ops.register_graph_buffers(self._ptr, handles, offsets) + + def should_custom_ar(self, inp: torch.Tensor): + if self.disabled: + return False + inp_size = inp.numel() * inp.element_size() + # custom allreduce requires input byte size to be multiples of 16 + if inp_size % 16 != 0: + return False + if not is_weak_contiguous(inp): + return False + # for 4 or more non NVLink-capable GPUs, custom allreduce provides + # little performance improvement over NCCL. + if not _is_hip: + if True: + return inp_size <= self.max_size + return False + + if _is_hip: + if self.use_amd_deterministic_impl: + return True + if self.full_nvlink: + return inp_size <= self.max_size + return False + + return False + + def _all_reduce_impl(self, inp: torch.Tensor, registered: bool): + out = torch.empty_like(inp) + if not _is_hip: # CUDA-like + if registered: + ops.all_reduce(self._ptr, inp, out, 0, 0) + else: + ops.all_reduce( + self._ptr, inp, out, self.buffer_ptrs[self.rank], self.max_size + ) + elif self.use_amd_deterministic_impl: + inp_size = inp.numel() * inp.element_size() + if inp_size < self.max_size: + reg_buffer = self.buffer.view(inp.dtype)[: inp.numel()] + ops.deterministic_all_reduce_unreg(self._ptr, inp, reg_buffer, out) + else: + self.register_buffer(inp) + ops.deterministic_all_reduce_reg(self._ptr, inp, out) + else: # normal AMD ROCm path + if registered: + ops.all_reduce_reg(self._ptr, inp, out) + else: + ops.all_reduce_unreg(self._ptr, inp, self.buffer, out) + return out + + def custom_all_reduce(self, input: torch.Tensor) -> Optional[torch.Tensor]: + """The main allreduce API that provides support for cuda graph.""" + # When custom allreduce is disabled, this will be None. + if self.disabled or not self.should_custom_ar(input): + return None + if self._IS_CAPTURING: + if torch.cuda.is_current_stream_capturing(): + return self._all_reduce_impl(input, registered=not self.tms_cudagraph) + else: + # Could be warmup OR piecewise cuda graph split op execution. + # In piecewise cuda graph, split ops run eagerly outside the graph + # but _IS_CAPTURING is still True. We need to do real all-reduce. + if is_in_tc_piecewise_cuda_graph(): + # Split op execution - do real all-reduce + return self._all_reduce_impl(input, registered=False) + else: + # True warmup - mimic the allocation pattern since custom + # allreduce is out-of-place. + return torch.zeros_like(input) + else: + return self._all_reduce_impl(input, registered=False) + + def close(self): + if not self.disabled and self._ptr: + if ops is not None: + ops.dispose(self._ptr) + if _is_cuda: + self.free_shared_buffer(self.meta_ptrs) + self.free_shared_buffer(self.buffer_ptrs) + self._ptr = 0 + + def __del__(self): + self.close() + + +def dispatch_custom_allreduce( + group: ProcessGroup, + device: torch.device, +): + """Return the CustomAllreduce class to use (aiter on ROCm if enabled). + + On AMD with 1-stage AR enabled, use sglang's CustomAllreduce. + Otherwise use AiterCustomAllreduce if available. + + On CUDA, the JIT-compiled v2 implementation is used by default. + Set SGLANG_OPT_USE_CUSTOM_ALL_REDUCE_V2=0 to fall back to the legacy CustomAllreduce. + Multi-node v2 is admitted only for a single NVLink clique (see + can_use_custom_all_reduce_v2); other cross-node groups fall back to NCCL. + """ + if _is_cuda and envs.SGLANG_OPT_USE_CUSTOM_ALL_REDUCE_V2.get(): + from .custom_all_reduce_v2 import ( + CustomAllReduceV2, + can_use_custom_all_reduce_v2, + ) + + if can_use_custom_all_reduce_v2(group=group, device=device): + logger.debug("[AR] Using CustomAllReduceV2 (JIT-compiled)") + return CustomAllReduceV2 + + if _is_cuda or _is_musa: + return CustomAllreduce + + assert _is_hip + + if envs.SGLANG_USE_1STAGE_ALLREDUCE.is_set(): + if envs.SGLANG_USE_1STAGE_ALLREDUCE.get(): + logger.debug( + "[AR] All-reduce: 1-stage kernel (SGLANG_USE_1STAGE_ALLREDUCE=1)" + ) + else: + logger.debug("[AR] All-reduce: default (SGLANG_USE_1STAGE_ALLREDUCE=0)") + elif envs.SGLANG_ENABLE_DETERMINISTIC_INFERENCE.get(): + logger.debug( + "[AR] All-reduce: 1-stage kernel (deterministic inference enabled)" + ) + else: + logger.debug("[AR] All-reduce: default") + + # On AMD with 1-stage AR, use sglang's CustomAllreduce + # (AiterCustomAllreduce doesn't have deterministic_all_reduce method) + if _use_amd_deterministic_impl(): + return CustomAllreduce + + if get_bool_env_var("SGLANG_USE_AITER_AR", default="true"): + try: + from aiter.dist.device_communicators.custom_all_reduce import ( + CustomAllreduce as AiterCustomAllreduce, + ) + + logger.info("[AR] Using AiterCustomAllreduce (AMD default)") + tms_cudagraph = envs.SGLANG_MEMORY_SAVER_CUDA_GRAPH.get() + return partial( + AiterCustomAllreduce, + enable_register_for_capturing=not tms_cudagraph, + ) + except ImportError as e: + logger.warning( + "[AR] Aiter custom all-reduce not available; " + "falling back to sglang CustomAllreduce. Details: %s", + e, + ) + return CustomAllreduce + + return CustomAllreduce + + +def _use_amd_deterministic_impl() -> bool: + if not _is_hip: # CUDA is always deterministic + return False + if envs.SGLANG_USE_1STAGE_ALLREDUCE.is_set(): + return envs.SGLANG_USE_1STAGE_ALLREDUCE.get() + else: + return envs.SGLANG_ENABLE_DETERMINISTIC_INFERENCE.get() diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/car_patch/custom_all_reduce_utils.py b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/car_patch/custom_all_reduce_utils.py new file mode 100644 index 0000000..70a70de --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/car_patch/custom_all_reduce_utils.py @@ -0,0 +1,519 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +# Adapted from https://github.com/vllm-project/vllm/blob/v0.6.4.post1/vllm/distributed/device_communicators/custom_all_reduce_utils.py + +import ctypes +import json +import logging +import os +import pickle +import subprocess +import sys +import tempfile +from functools import wraps +from itertools import product +from typing import Callable, Dict, List, Optional, Sequence, TypeVar + +import torch +import torch.distributed as dist +import torch.multiprocessing as mp +from typing_extensions import ParamSpec + +from sglang.srt.distributed.device_communicators.cuda_wrapper import CudaRTLibrary +from sglang.srt.distributed.parallel_state import in_the_same_node_as +from sglang.srt.environ import envs as sglang_envs +from sglang.srt.utils import is_cuda, is_hip, is_musa +from sglang.srt.utils.cuda_vmm_utils import _gpu_fabric_clique + +logger = logging.getLogger(__name__) + +_is_cuda = is_cuda() +_is_hip = is_hip() +_is_musa = is_musa() + +if _is_cuda: + try: + import pynvml + except ImportError as e: + logger.warning("Failed to import pynvml with %r", e) + +if _is_musa: + try: + import pymtml as pynvml + except ImportError as e: + logger.warning("Failed to import pymtml with %r", e) + +if _is_hip: + try: + from amdsmi import ( + AmdSmiException, + amdsmi_get_processor_handles, + amdsmi_init, + amdsmi_shut_down, + amdsmi_topo_get_link_type, + ) + except ImportError as e: + logger.warning("Failed to import amdsmi with %r", e) + +_P = ParamSpec("_P") +_R = TypeVar("_R") + + +def update_environment_variables(envs: Dict[str, str]): + for k, v in envs.items(): + if k in os.environ and os.environ[k] != v: + logger.warning( + "Overwriting environment variable %s " "from '%s' to '%s'", + k, + os.environ[k], + v, + ) + os.environ[k] = v + + +def producer( + batch_src: Sequence[int], + producer_queue, + consumer_queue, + result_queue, + cuda_visible_devices: Optional[str] = None, +): + if cuda_visible_devices is not None: + update_environment_variables({"CUDA_VISIBLE_DEVICES": cuda_visible_devices}) + + lib = CudaRTLibrary() + for i in batch_src: + lib.cudaSetDevice(i) + pointer = lib.cudaMalloc(1024) + lib.cudaMemset(pointer, 1, 1024) + lib.cudaDeviceSynchronize() + handle = lib.cudaIpcGetMemHandle(pointer) + producer_queue.put(handle) + open_success = consumer_queue.get() + if open_success: + # use two queues to simulate barrier + producer_queue.put(0) + consumer_queue.get() + # check if the memory is modified + host_data = (ctypes.c_char * 1024)() + lib.cudaMemcpy(host_data, pointer, 1024) # type: ignore + for i in range(1024): + if ord(host_data[i]) != 2: + open_success = False + break + result_queue.put(open_success) + lib.cudaDeviceReset() + + +def consumer( + batch_tgt: Sequence[int], + producer_queue, + consumer_queue, + result_queue, + cuda_visible_devices: Optional[str] = None, +): + if cuda_visible_devices is not None: + update_environment_variables({"CUDA_VISIBLE_DEVICES": cuda_visible_devices}) + + lib = CudaRTLibrary() + for j in batch_tgt: + lib.cudaSetDevice(j) + handle = producer_queue.get() + open_success = False + try: + pointer = lib.cudaIpcOpenMemHandle(handle) # type: ignore + open_success = True + except RuntimeError: + # cannot error out here, because the producer process + # is still waiting for the response. + pass + consumer_queue.put(open_success) + if open_success: + # modify the memory + lib.cudaMemset(pointer, 2, 1024) + lib.cudaDeviceSynchronize() + # use two queues to simulate barrier + producer_queue.get() + consumer_queue.put(0) + # check if the memory is modified + host_data = (ctypes.c_char * 1024)() + lib.cudaMemcpy(host_data, pointer, 1024) # type: ignore + for i in range(1024): + if ord(host_data[i]) != 2: + open_success = False + break + result_queue.put(open_success) + lib.cudaDeviceReset() + + +def can_actually_p2p( + batch_src: Sequence[int], + batch_tgt: Sequence[int], +) -> Sequence[bool]: + """ + Usually, checking if P2P access is enabled can be done by + `torch.cuda.can_device_access_peer(src, tgt)`. However, sometimes + the driver might be broken, and `torch.cuda.can_device_access_peer(src, tgt)` + returns `True` even if P2P access is not actually possible. + See https://github.com/vllm-project/vllm/issues/2728 and + https://forums.developer.nvidia.com/t/direct-gpu-gpu-communication-does-not-seem-to-work-properly/283264/10 + Therefore, we have to perform a real P2P access to check if it is actually + possible. + + Note on p2p and cuda IPC: + Usually, one process uses one GPU: + GPU src --> cuda context src --> tensor src --> process src + + We need to combine p2p and cuda IPC, so that: + GPU src --> cuda context src --> tensor src --> process src + |shared| + GPU tgt --> cuda context tgt --> tensor tgt --> process tgt + That is to say, process src creates a tensor in GPU src, passes IPC handle to + process tgt, and process tgt accesses the tensor in GPU tgt. Any operation on the + tensor in process tgt will be reflected in the tensor in process src, because + they are the same memory segment. + It is important to note that process tgt accesses the tensor in GPU tgt, not + GPU src. That's why we need p2p access. + + The most time-consuming part is the process creation. To avoid creating + processes for every pair of GPUs, we use batched testing. We create two + processes for testing all pairs of GPUs in batch. The trick is to reset + the device after each test (which is not available in PyTorch). + """ # noqa + cuda_visible_devices = os.environ.get("CUDA_VISIBLE_DEVICES", None) + # pass the CUDA_VISIBLE_DEVICES to the child process + # to make sure they see the same set of GPUs + + # make sure the processes are spawned + smp = mp.get_context("spawn") + producer_queue = smp.Queue() + consumer_queue = smp.Queue() + result_queue = smp.Queue() + p_src = smp.Process( + target=producer, + args=( + batch_src, + producer_queue, + consumer_queue, + result_queue, + cuda_visible_devices, + ), + ) + p_tgt = smp.Process( + target=consumer, + args=( + batch_tgt, + producer_queue, + consumer_queue, + result_queue, + cuda_visible_devices, + ), + ) + p_src.start() + p_tgt.start() + p_src.join() + p_tgt.join() + assert p_src.exitcode == 0 and p_tgt.exitcode == 0 + result: List[bool] = [] + for src, tgt in zip(batch_src, batch_tgt): + a = result_queue.get() + b = result_queue.get() + if a != b: + logger.warning( + "Two processes do not agree on the P2P access" + " status on %d -> %d, treat as disabled.", + src, + tgt, + ) + result.append(False) + else: + result.append(a) + return result + + +# why do we need this cache? +# we are testing peer-to-peer (p2p) access between GPUs,across processes. +# if we test it every time, it will be very slow, because we need to create +# N * N * 2 processes, where N is the world size. This is very slow. +# to reduce the time, we use a cache file to store the p2p access status. +# the cache file is generated by the master process if it does not exist. +# then all the processes can read the cache file to check the p2p access status. +# Note that the cache file is suffixed by the CUDA_VISIBLE_DEVICES, so that we +# can have different cache files for different CUDA_VISIBLE_DEVICES settings, +# e.g. used by different vllm engines. The device id in the cache file is a +# **local** device id, i.e. from 0 to num_dev-1, where num_dev is the number +# of visible devices in the vllm engine. +_gpu_p2p_access_cache: Optional[Dict[str, bool]] = None + + +def gpu_p2p_access_check(src: int, tgt: int) -> bool: + """Check if GPU src can access GPU tgt.""" + + # if the cache variable is already calculated, + # read from the cache instead of checking it again + global _gpu_p2p_access_cache + if _gpu_p2p_access_cache is not None: + return _gpu_p2p_access_cache[f"{src}->{tgt}"] + + is_distributed = dist.is_initialized() + + num_dev = torch.cuda.device_count() + cuda_visible_devices = os.environ.get("CUDA_VISIBLE_DEVICES", None) + if cuda_visible_devices is None: + cuda_visible_devices = ",".join(str(i) for i in range(num_dev)) + + # VLLM_CACHE_ROOT -> SGLANG_CACHE_ROOT + # "~/.cache/vllm" -> envs.SGLANG_CACHE_DIR + SGLANG_CACHE_ROOT = os.path.expanduser(sglang_envs.SGLANG_CACHE_DIR.get()) + path = os.path.join( + SGLANG_CACHE_ROOT, f"gpu_p2p_access_cache_for_{cuda_visible_devices}.json" + ) + cache_dir = os.path.dirname(path) + try: + os.makedirs(cache_dir, exist_ok=True) + except (FileExistsError, NotADirectoryError): + if not os.path.isdir(cache_dir): + # Path exists as a file (stale cache/lock). Remove and retry. + try: + os.remove(cache_dir) + except OSError: + pass + os.makedirs(cache_dir, exist_ok=True) + from sglang.srt.distributed.parallel_state import get_world_group + + if (not is_distributed or get_world_group().local_rank == 0) and ( + not os.path.exists(path) + ): + # only the local master process (with local_rank == 0) can + # enter this block to calculate the cache + logger.info("generating GPU P2P access cache in %s", path) + cache: Dict[str, bool] = {} + ids = list(range(num_dev)) + # batch of all pairs of GPUs + batch_src, batch_tgt = zip(*list(product(ids, ids))) + # NOTE: we use `subprocess` rather than `multiprocessing` here + # because the caller might not have `if __name__ == "__main__":`, + # in that case we cannot use spawn method in multiprocessing. + # However, `can_actually_p2p` requires spawn method. + # The fix is, we use `subprocess` to call the function, + # where we have `if __name__ == "__main__":` in this file. + + # use a temporary file to store the result + # we don't use the output of the subprocess directly, + # because the subprocess might produce logging output + with tempfile.NamedTemporaryFile() as output_file: + input_bytes = pickle.dumps((batch_src, batch_tgt, output_file.name)) + returned = subprocess.run( + [sys.executable, __file__], input=input_bytes, capture_output=True + ) + # check if the subprocess is successful + try: + returned.check_returncode() + except Exception as e: + # wrap raised exception to provide more information + raise RuntimeError( + f"Error happened when batch testing " + f"peer-to-peer access from {batch_src} to {batch_tgt}:\n" + f"{returned.stderr.decode()}" + ) from e + with open(output_file.name, "rb") as f: + result = pickle.load(f) + for _i, _j, r in zip(batch_src, batch_tgt, result): + cache[f"{_i}->{_j}"] = r + with open(path, "w") as f: + json.dump(cache, f, indent=4) + if is_distributed: + get_world_group().barrier() + logger.info("reading GPU P2P access cache from %s", path) + with open(path) as f: + cache = json.load(f) + _gpu_p2p_access_cache = cache + return _gpu_p2p_access_cache[f"{src}->{tgt}"] + + +def with_nvml_context(fn: Callable[_P, _R]) -> Callable[_P, _R]: + @wraps(fn) + def wrapper(*args: _P.args, **kwargs: _P.kwargs) -> _R: + if _is_hip: + try: + amdsmi_init() + return fn(*args, **kwargs) + finally: + amdsmi_shut_down() + else: + pynvml.nvmlInit() + try: + return fn(*args, **kwargs) + finally: + pynvml.nvmlShutdown() + + return wrapper + + +@with_nvml_context +def is_full_nvlink(physical_device_ids: List[int], world_size: int) -> bool: + if _is_hip: + """ + query if the set of gpus are fully connected by xgmi (1 hop) + """ + handles = [amdsmi_get_processor_handles()[i] for i in physical_device_ids] + for i, handle in enumerate(handles): + for j, peer_handle in enumerate(handles): + if i < j: + try: + link_type = amdsmi_topo_get_link_type(handle, peer_handle) + # type is 2 for XGMI + if link_type["hops"] != 1 or link_type["type"] != 2: + return False + except AmdSmiException as error: + logger.error("AMD 1 hop XGMI detection failed.", exc_info=error) + return False + return True + else: + """ + query if the set of gpus are fully connected by nvlink (1 hop) + """ + handles = [pynvml.nvmlDeviceGetHandleByIndex(i) for i in physical_device_ids] + for i, handle in enumerate(handles): + for j, peer_handle in enumerate(handles): + if i < j: + try: + p2p_status = pynvml.nvmlDeviceGetP2PStatus( + handle, peer_handle, pynvml.NVML_P2P_CAPS_INDEX_NVLINK + ) + if p2p_status != pynvml.NVML_P2P_STATUS_OK: + return False + except pynvml.NVMLError: + logger.exception( + "NVLink detection failed. This is normal if your" + " machine has no NVLink equipped." + ) + return False + return True + + +@with_nvml_context +def is_one_nvlink_clique( + group: torch.distributed.ProcessGroup, device: torch.device +) -> bool: + """True iff every rank's GPU is in the same NVLink fabric clique (one NVL72 / + MNNVL domain). Such a clique shares a single NVLink address space even across + nodes, so custom-AR v2's symm-mem storage + fabric peer VAs are valid group-wide.""" + if _is_hip: + return False + try: + clique = _gpu_fabric_clique(device) + except Exception as e: + logger.warning( + "GPU fabric clique query failed (%r); custom-AR stays intra-node.", e + ) + clique = None + # Always all-gather (every rank calls it once) so a failed query on any rank + # resolves to a clean False rather than a collective mismatch. + world_size = dist.get_world_size(group=group) + gathered: List[object] = [None] * world_size + dist.all_gather_object(gathered, clique, group=group) + if any(c is None for c in gathered): + return False + return len(set(gathered)) == 1 + + +def is_weak_contiguous(inp: torch.Tensor): + return inp.is_contiguous() or ( + inp.storage().nbytes() - inp.storage_offset() * inp.element_size() + == inp.numel() * inp.element_size() + ) + + +def can_p2p(rank: int, world_size: int) -> bool: + # SGLANG_SKIP_P2P_CHECK can be set to False in sglang + SGLANG_SKIP_P2P_CHECK = os.getenv("SGLANG_SKIP_P2P_CHECK", "0") == "1" + for i in range(world_size): + if i == rank: + continue + if SGLANG_SKIP_P2P_CHECK: + logger.info("Skipping P2P check and trusting the driver's P2P report.") + return torch.cuda.can_device_access_peer(rank, i) + if not gpu_p2p_access_check(rank, i): + return False + return True + + +def can_use_custom_all_reduce_with_nvlink( + group: torch.distributed.ProcessGroup, + device: torch.device, + supported_world_size: List[int], + cls_name: str, +) -> Optional[bool]: # None if fail; otherwise return whether NVLink is available + assert ( + dist.get_backend(group) != dist.Backend.NCCL + ), f"{cls_name} should be attached to a non-NCCL group." + + rank = dist.get_rank(group=group) + world_size = dist.get_world_size(group=group) + + # No need to initialize custom allreduce for single GPU case. + if world_size == 1: + return + + # No need to initialize custom allreduce for multi-node case. + if not all(in_the_same_node_as(group, source_rank=0)): + logger.warning( + f"{cls_name} is disabled because this process group" " spans across nodes." + ) + return + + # For not supported world size, we disable custom allreduce. + if world_size not in supported_world_size: + logger.warning( + f"{cls_name} is disabled due to an unsupported world" + f" size: {world_size}. Supported world sizes: {supported_world_size}. " + "To silence this warning, specify disable_custom_all_reduce=True explicitly.", + ) + return + + cuda_visible_devices = os.environ.get("CUDA_VISIBLE_DEVICES", None) + if cuda_visible_devices: + device_ids = list(map(int, cuda_visible_devices.split(","))) + else: + device_ids = list(range(torch.cuda.device_count())) + physical_device_id = device_ids[device.index] + tensor = torch.tensor([physical_device_id], dtype=torch.int, device="cpu") + gather_list = [ + torch.tensor([0], dtype=torch.int, device="cpu") for _ in range(world_size) + ] + dist.all_gather(gather_list, tensor, group=group) + physical_device_ids = [int(t) for t in gather_list] + full_nvlink = is_full_nvlink(physical_device_ids, world_size) + + # test nvlink first, this will filter out most of the cases + # where custom allreduce is not supported + # this checks hardware and driver support for NVLink + if False: + logger.warning( + f"{cls_name} is disabled because it's not supported on" + " more than two PCIe-only GPUs. To silence this warning, " + "specify disable_custom_all_reduce=True explicitly." + ) + return + + # test P2P capability, this checks software/cudaruntime support + # this is expensive to compute at the first time + # then we cache the result + # On AMD GPU, p2p is always enabled between XGMI connected GPUs + if not _is_hip and not can_p2p(rank, world_size): + logger.warning( + f"{cls_name} is disabled because your platform lacks " + "GPU P2P capability or P2P test failed. To silence this " + "warning, specify disable_custom_all_reduce=True explicitly." + ) + return + + return full_nvlink + + +if __name__ == "__main__": + batch_src, batch_tgt, output_file = pickle.loads(sys.stdin.buffer.read()) + result = can_actually_p2p(batch_src, batch_tgt) + with open(output_file, "wb") as f: + f.write(pickle.dumps(result)) diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/car_patch/prep_car_patch.sh b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/car_patch/prep_car_patch.sh new file mode 100644 index 0000000..9404572 --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/car_patch/prep_car_patch.sh @@ -0,0 +1,27 @@ +#!/bin/bash +# prep_car_patch.sh — build patched custom-AR files at /root/patches/ +# Patch A: should_custom_ar allows any CUDA AR <= max_size (8MB cap keeps +# big prefill ARs on NCCL); decode ARs (<=1MB) go custom AR. +# Patch B: utils gate no longer bails on ws>2 without NVLink (P2P probe +# verified all-True on this 8-way PCIe box). +# V2 stays excluded: can_use_custom_all_reduce_v2 requires full_nvlink is True. +set -e +D=/sgl-workspace/sglang/python/sglang/srt/distributed/device_communicators +mkdir -p /root/patches +docker cp glm53-nvfp4:$D/custom_all_reduce.py /root/patches/custom_all_reduce.py +docker cp glm53-nvfp4:$D/custom_all_reduce_utils.py /root/patches/custom_all_reduce_utils.py + +ANCHOR='self.original_disabled = False # Ensure original_disabled == disabled' +n=$(grep -c "$ANCHOR" /root/patches/custom_all_reduce.py) +[ "$n" = "1" ] || { echo "ANCHOR_NOT_UNIQUE=$n"; exit 1; } + +sed -i 's/if self.world_size == 2 or self.full_nvlink:/if True:/' /root/patches/custom_all_reduce.py +sed -i 's/if world_size > 2 and not full_nvlink:/if False:/' /root/patches/custom_all_reduce_utils.py +sed -i "/${ANCHOR}/a\\ logger.warning(f\"SSKJ_CAR_PATCH_ACTIVE ws={self.world_size} full_nvlink={self.full_nvlink}\")" /root/patches/custom_all_reduce.py + +python3 -m py_compile /root/patches/custom_all_reduce.py /root/patches/custom_all_reduce_utils.py +echo "--- verify ---" +grep -n 'if True:' /root/patches/custom_all_reduce.py +grep -n 'if False:' /root/patches/custom_all_reduce_utils.py +grep -n 'SSKJ_CAR_PATCH_ACTIVE' /root/patches/custom_all_reduce.py +echo PREP_DONE diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/deploy_glm53_605.sh b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/deploy_glm53_605.sh new file mode 100644 index 0000000..fe0f108 --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/deploy_glm53_605.sh @@ -0,0 +1,58 @@ +#!/bin/bash +# deploy_glm53_605.sh — GLM-5.3-NVFP4 TP8 + EAGLE for 174.1.60.5 (team service) +# Target scenario: cc1-2, 64k/128k input 90% cache hit, single request up to 256k, +# output-throughput priority. Clean deploy: rm old container, wait VRAM drain, run. +# +# Env overrides: +# MEMFRAC (0.90) STEPS (4) TOPK (1) DRAFT (5) CTXLEN (270336) +# CHUNK (8192) MAXPRE (16384) RESTART (no|yes) EXTRA ("") +# Usage: +# bash deploy_glm53_605.sh # production 4/1/5 @ 0.90 +# STEPS=5 DRAFT=6 bash deploy_glm53_605.sh # tuning round +# CHUNK=16384 bash deploy_glm53_605.sh # prefill tuning round +# RESTART=yes bash deploy_glm53_605.sh # production finalize +set -e +MEMFRAC=${MEMFRAC:-0.90} +STEPS=${STEPS:-4} +TOPK=${TOPK:-1} +DRAFT=${DRAFT:-5} +CTXLEN=${CTXLEN:-270336} +CHUNK=${CHUNK:-8192} +MAXPRE=${MAXPRE:-16384} +EXTRA=${EXTRA:-} +if [ "$RESTART" = "yes" ]; then RP="--restart unless-stopped"; else RP="--restart no"; fi + +echo "[deploy] removing old container (if any)" +# rm -f times out on big GPU containers on this daemon; retry until really gone +for i in $(seq 1 45); do + CID=$(docker ps -a --filter name=glm53-nvfp4 -q) + [ -z "$CID" ] && break + docker rm -f glm53-nvfp4 >/dev/null 2>&1 || true + sleep 2 +done + +echo "[deploy] waiting for VRAM drain" +for i in $(seq 1 45); do + used=$(nvidia-smi --query-gpu=memory.used --format=csv,noheader,nounits | awk '{s+=$1} END {print s}') + [ "$used" -lt 2000 ] && break + sleep 2 +done +used=$(nvidia-smi --query-gpu=memory.used --format=csv,noheader,nounits | awk '{s+=$1} END {print s}') +echo "[deploy] VRAM now: ${used} MiB total" + +echo "[deploy] starting: TP8 EAGLE ${STEPS}/${TOPK}/${DRAFT} memfrac=${MEMFRAC} extra='${EXTRA}'" +docker run -d --name glm53-nvfp4 --gpus all --shm-size 64g --ipc=host $RP \ + -p 30000:30000 -v /data/hf_models:/data/hf_models \ + lmsysorg/sglang:nightly-dev-20260828-daf63171 \ + python3 -m sglang.launch_server \ + --model-path /data/hf_models/GLM-5.3-NVFP4 --tp 8 \ + --mem-fraction-static $MEMFRAC --max-running-requests 16 \ + --chunked-prefill-size $CHUNK --max-prefill-tokens $MAXPRE \ + --disable-shared-experts-fusion --moe-runner-backend flashinfer_cutlass --disable-flashinfer-autotune \ + --speculative-algorithm EAGLE --speculative-num-steps $STEPS --speculative-eagle-topk $TOPK --speculative-num-draft-tokens $DRAFT \ + --kv-cache-dtype fp8_e4m3 --enable-hierarchical-cache --hicache-ratio 3 \ + --cuda-graph-max-bs-decode 8 --cuda-graph-bs-decode 1 2 3 4 6 8 --cuda-graph-max-bs-prefill 8 \ + --context-length $CTXLEN --reasoning-parser glm45 --tool-call-parser glm47 \ + --host 0.0.0.0 --port 30000 $EXTRA + +echo "[deploy] container started; poll: docker logs -f glm53-nvfp4" diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/deploy_glm53_607_exp.sh b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/deploy_glm53_607_exp.sh new file mode 100644 index 0000000..9ede120 --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/deploy_glm53_607_exp.sh @@ -0,0 +1,70 @@ +#!/bin/bash +# deploy_glm53_607_exp.sh — GLM-5.3-NVFP4 TP8 + EAGLE experiment deploy for 174.1.60.7 +# EXACT copy of deploy_glm53_605.sh (production standard, 8-machine md5 fcd9109b) +# with only this header changed. Campaign: scenario-1 optimization (i128k/i64k, +# o512, cc1-4). Do NOT sync this file to other machines. +# +# Env overrides: +# MEMFRAC (0.90) STEPS (4) TOPK (1) DRAFT (5) CTXLEN (270336) +# CHUNK (8192) MAXPRE (16384) RESTART (no|yes) EXTRA ("") +# CAR_PATCH (1) inject custom-AR 1stage patch (E7b winner, 2026-09-07): +# small ARs <=8MB go one-shot kernel instead of NCCL RING_LL on 8-way PCIe. +# Patch files at /root/patches/ (build: bash /root/prep_car_patch.sh on a +# running stock container). Patch dies with the container; image untouched. +# Set CAR_PATCH=0 to deploy stock. +set -e +MEMFRAC=${MEMFRAC:-0.90} +STEPS=${STEPS:-4} +TOPK=${TOPK:-1} +DRAFT=${DRAFT:-5} +CTXLEN=${CTXLEN:-270336} +CHUNK=${CHUNK:-8192} +MAXPRE=${MAXPRE:-16384} +EXTRA=${EXTRA:-} +if [ "$RESTART" = "yes" ]; then RP="--restart unless-stopped"; else RP="--restart no"; fi + +echo "[deploy] removing old container (if any)" +# rm -f times out on big GPU containers on this daemon; retry until really gone +for i in $(seq 1 45); do + CID=$(docker ps -a --filter name=glm53-nvfp4 -q) + [ -z "$CID" ] && break + docker rm -f glm53-nvfp4 >/dev/null 2>&1 || true + sleep 2 +done + +echo "[deploy] waiting for VRAM drain" +for i in $(seq 1 45); do + used=$(nvidia-smi --query-gpu=memory.used --format=csv,noheader,nounits | awk '{s+=$1} END {print s}') + [ "$used" -lt 2000 ] && break + sleep 2 +done +used=$(nvidia-smi --query-gpu=memory.used --format=csv,noheader,nounits | awk '{s+=$1} END {print s}') +echo "[deploy] VRAM now: ${used} MiB total" + +echo "[deploy] starting: TP8 EAGLE ${STEPS}/${TOPK}/${DRAFT} memfrac=${MEMFRAC} chunk=${CHUNK} extra='${EXTRA}'" +docker run -d --name glm53-nvfp4 --gpus all --shm-size 64g --ipc=host $RP \ + -p 30000:30000 -v /data/hf_models:/data/hf_models \ + lmsysorg/sglang:nightly-dev-20260828-daf63171 \ + python3 -m sglang.launch_server \ + --model-path /data/hf_models/GLM-5.3-NVFP4 --tp 8 \ + --mem-fraction-static $MEMFRAC --max-running-requests 16 \ + --chunked-prefill-size $CHUNK --max-prefill-tokens $MAXPRE \ + --disable-shared-experts-fusion --moe-runner-backend flashinfer_cutlass --disable-flashinfer-autotune \ + --speculative-algorithm EAGLE --speculative-num-steps $STEPS --speculative-eagle-topk $TOPK --speculative-num-draft-tokens $DRAFT \ + --kv-cache-dtype fp8_e4m3 --enable-hierarchical-cache --hicache-ratio 3 \ + --cuda-graph-max-bs-decode 8 --cuda-graph-bs-decode 1 2 3 4 6 8 --cuda-graph-max-bs-prefill 8 \ + --context-length $CTXLEN --reasoning-parser glm45 --tool-call-parser glm47 \ + --host 0.0.0.0 --port 30000 $EXTRA + +echo "[deploy] container started; poll: docker logs -f glm53-nvfp4" + +# --- CAR 1stage patch injection (E7b winner) --- +if [ "${CAR_PATCH:-1}" != "0" ] && [ -f /root/patches/custom_all_reduce.py ]; then + echo "[deploy] CAR_PATCH: injecting custom-AR 1stage patch" + docker stop -t 20 glm53-nvfp4 >/dev/null 2>&1 || true + DPATH=/sgl-workspace/sglang/python/sglang/srt/distributed/device_communicators + docker cp /root/patches/custom_all_reduce.py glm53-nvfp4:$DPATH/custom_all_reduce.py + docker cp /root/patches/custom_all_reduce_utils.py glm53-nvfp4:$DPATH/custom_all_reduce_utils.py + docker start glm53-nvfp4 + echo "[deploy] CAR_PATCH injected, container restarted; poll health as usual" +fi diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/deploy_glm53_optimal.sh b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/deploy_glm53_optimal.sh new file mode 100644 index 0000000..7faa6f8 --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/deploy_glm53_optimal.sh @@ -0,0 +1,62 @@ +#!/bin/bash +# ============================================================ +# GLM-5.3 最优部署方案（6000D 8卡，TP4 PP2 + IndexCache freq=4） +# 2026-09-07 +# +# - 基线配置：TP4 PP2 + cps16k + mem0.85（131.6 tok/s 吞吐基线） +# - IndexCache（index_topk_freq=4）：层轴索引复用，省 75% indexer +# 16K 场景无损失；128K 长上下文并发 1.35-1.47× 提速 +# - 禁 radix cache；禁投机解码（PP2 与投机框架不兼容，已实测） +# +# 用法：bash deploy_glm53_optimal.sh +# （60.7 试用版：增加显存排空等待 + 就绪等待加长至 20 分钟） +# ============================================================ +set -uo pipefail + +CONTAINER="glm53-nvfp4" +IMAGE="lmsysorg/sglang:nightly-dev-20260828-daf63171" +MODEL="/data/hf_models/GLM-5.3-NVFP4" +PORT=30000 +TP=4; PP=2; MEM=0.85; MRR=48; CPS=16384 + +docker rm -f ${CONTAINER} 2>/dev/null || true + +# docker rm -f 后显存释放滞后数分钟，不等会把新容器 KV 池压小 +for i in $(seq 1 60); do + used=$(nvidia-smi --query-gpu=memory.used --format=csv,noheader,nounits | awk '{s+=$1} END {print s}') + if [ "$used" -lt 1500 ]; then echo "[deploy] drained: ${used} MiB"; break; fi + echo "[deploy] drain wait ${i}: ${used} MiB" + sleep 10 +done + +docker run -d --name ${CONTAINER} --gpus all --shm-size 64g --ipc=host \ + --restart unless-stopped \ + -p ${PORT}:${PORT} \ + -v /data/hf_models:/data/hf_models \ + ${IMAGE} \ + python3 -m sglang.launch_server \ + --model-path ${MODEL} \ + --tp-size ${TP} --pp-size ${PP} \ + --mem-fraction-static ${MEM} \ + --max-running-requests ${MRR} \ + --disable-radix-cache \ + --disable-shared-experts-fusion \ + --moe-runner-backend flashinfer_cutlass \ + --disable-flashinfer-autotune \ + --disable-custom-all-reduce \ + --chunked-prefill-size ${CPS} \ + --host 0.0.0.0 --port ${PORT} \ + --json-model-override-args '{"index_topk_freq": 4}' + +echo "容器已启动，等待就绪..." +for i in $(seq 1 120); do + code=$(curl -s -o /dev/null -w '%{http_code}' http://localhost:${PORT}/health 2>/dev/null) + if [ "$code" = "200" ]; then + echo "READY after ~$((i*10))s" + docker ps --filter name=${CONTAINER} --format '{{.Names}} {{.Status}}' + echo "override args: $(docker inspect ${CONTAINER} --format '{{.Config.Cmd}}' | grep -o 'index_topk_freq[^,}]*')" + exit 0 + fi + sleep 10 +done +echo "TIMEOUT"; exit 1 diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/deploy_glm53_optimal_s1.sh b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/deploy_glm53_optimal_s1.sh new file mode 100644 index 0000000..05b3461 --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_dual_scenario_bench/scripts/deploy_glm53_optimal_s1.sh @@ -0,0 +1,59 @@ +#!/bin/bash +# ============================================================ +# TP4 PP2 + IndexCache（场景一轮：启用 radix cache） +# 与 deploy_glm53_optimal.sh 唯一差异：去掉 --disable-radix-cache +# ——场景一前提是 90% 前缀命中，禁 radix 则命中率为 0 无从对比 +# 其余参数原样：TP4 PP2 / mem0.85 / mrr48 / cps16384 / index_topk_freq=4 / 无投机 +# ============================================================ +set -uo pipefail + +CONTAINER="glm53-nvfp4" +IMAGE="lmsysorg/sglang:nightly-dev-20260828-daf63171" +MODEL="/data/hf_models/GLM-5.3-NVFP4" +PORT=30000 +TP=4; PP=2; MEM=0.85; MRR=48; CPS=16384 + +docker rm -f ${CONTAINER} 2>/dev/null || true +# 补刀循环：docker rm -f 异步滞留，等容器对象彻底消失 +for i in $(seq 1 40); do + n=$(docker ps -a --filter name=${CONTAINER} -q | wc -l) + if [ "$n" = "0" ]; then echo "[deploy] container object gone"; break; fi + sleep 5 +done +# 排空显存 +for i in $(seq 1 60); do + used=$(nvidia-smi --query-gpu=memory.used --format=csv,noheader,nounits | awk '{s+=$1} END {print s}') + if [ "$used" -lt 1500 ]; then echo "[deploy] drained: ${used} MiB"; break; fi + echo "[deploy] drain wait ${i}: ${used} MiB" + sleep 10 +done + +docker run -d --name ${CONTAINER} --gpus all --shm-size 64g --ipc=host \ + --restart unless-stopped \ + -p ${PORT}:${PORT} \ + -v /data/hf_models:/data/hf_models \ + ${IMAGE} \ + python3 -m sglang.launch_server \ + --model-path ${MODEL} \ + --tp-size ${TP} --pp-size ${PP} \ + --mem-fraction-static ${MEM} \ + --max-running-requests ${MRR} \ + --disable-shared-experts-fusion \ + --moe-runner-backend flashinfer_cutlass \ + --disable-flashinfer-autotune \ + --disable-custom-all-reduce \ + --chunked-prefill-size ${CPS} \ + --host 0.0.0.0 --port ${PORT} \ + --json-model-override-args '{"index_topk_freq": 4}' + +echo "容器已启动，等待就绪..." +for i in $(seq 1 120); do + code=$(curl -s -o /dev/null -w '%{http_code}' http://localhost:${PORT}/health 2>/dev/null) + if [ "$code" = "200" ]; then + echo "READY after ~$((i*10))s" + docker ps --filter name=${CONTAINER} --format '{{.Names}} {{.Status}}' + exit 0 + fi + sleep 10 +done +echo "TIMEOUT"; exit 1