cp分成了prefill和decode cp核心来说，就是分成几个不同的思路点： rank持有部分q和部分kv，节约容量，节约激活，得用ring attention计算 rank持有部分q和全量kv，节约不了容量，节约激活，加速（更适用于mla） rank持有全量q和部分kv，节约容量，不节约激活，加速 cp是需要每个rank都持有完整的qkv_proj权重的，所以开启cp扩充不了总容量，但是可以扩充单个请求的最长上下文 pcp pcp分成两种，一种是megatron提出的ring attention，一种是DeepSpeed Ulysses cp ring attention 每张卡上存储完整的权重 qkv_proj也是切序列，然后每张卡上存储一部分序列的kv，然后在attention计算时做ring attention通信，也就是每张卡上的部分q每轮和传来的kv做计算 这样计算n轮传递n-1轮，每个q就和历史的全量kv做完了计算，再做o_proj，就得到了输出 然后ring attention的计算本身是不平衡的，因为q对历史的长度不一样，由此引出一个zig-zag ...

gpu 连接关系拓扑 NIC : Network Interface Card 就是网卡 NVSwitch : NVLink的芯片 普通非RDMA拓扑 命令nvidia-smi topo -m 例如输出： 性能区别： 1234567891011NVLink/NVSwitch > PIX（同一 PCIe Switch） > PXB（多个 PCIe Switch） > PHB/NODE（经过 Host Bridge） > SYS（跨 CPU Socket/NUMA） > 跨机 TCP 拓扑示意图： 1234567891011121314151617============================= 【双 CPU 服务器 PCIe 拓扑示意图】 ============================= [跨 CPU 互联总线 (如 Intel UPI)] / \ ...

SGLang KV Cache Retract 机制详解:普通模式 vs PD 分离模式 一条 input 长度逼近 KV cache 容量的请求,在 decode 阶段会触发一次看起来很费时间的 “retract” 交互。本文讲清楚 retract 到底是什么、为什么会被触发、怎么实现的,以及它在普通模式和 PD(prefill-decode)分离模式下走的是两条完全不同的路。 retract 解决什么问题 decode 阶段每生成一个 token,每个跨越 page 边界的请求都要从 KV cache 池子里申请新 page。当池子见底、下一步 decode 放不下时,系统有两种选择: 直接 OOM 崩掉 scheduler; 主动撤回若干个正在 decode 的请求,释放它们的 KV cache,把腾出来的空间留给剩下的请求继续跑。 retract 就是第 2 种。它是 decode OOM 的兜底机制,用”牺牲少数请求”换”整个 batch 不崩”。 SGLang 作者自己对 retract 的评价很直白,在 PR #708 的描述里写着: Retract ...

HiSparse:把长上下文 decode 的容量瓶颈从 GPU 显存搬到 CPU 内存 长上下文推理里,稀疏注意力解决了”算不动”的问题,却没解决”装不下”的问题。一个 1M 上下文的请求,即便每步 decode 只看 top-k 个 token,它的完整 KV cache 仍然得常驻 GPU HBM,否则没法被快速访问。结果是 batch size 很快被显存撑爆,throughput 在并发还没拉起来时就见顶了。 HiSparse(Hierarchical Sparse Attention)就是 SGLang 社区针对这个容量瓶颈做的 decode 优化:把完整但冷的 KV 放进 CPU pinned memory,GPU 上每个请求只留一个固定大小的”热”buffer,decode 时按需把 top-k 的 KV 从 host swap-in 回来。实测在 GLM-5.1-FP8 的长上下文场景下吞吐最高 5×,256 并发时约 3× 基线。 这篇文档面向想搞清楚 HiSparse 到底省了什么、代价是什么、能不能扩大系统 KV 容量的工程师。代码引用基于本仓 ...

coredump使用方式 首先新建test_coredump.py脚本，脚本内容如下： 123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657import os import socket import torch from torch.utils.cpp_extension import load_inline # 提前获取 hostname 和 pid，方便打印出触发 coredump 的准确命令 hostname = socket.gethostname() pid = os.getpid() pipe_path = f"/tmp/cuda_pipe_{hostname}_{pid}" dump_path = f"/tmp/cuda_coredump_full_{hostname}_{pid}" cuda_source = """ __global__ void hang_kernel(int* dumm ...

如果把conda/miniconda换一个目录，会出现一些脚本bash bang等硬编码写死路径的问题，需要修复 可以执行以下脚本修复 12345678910111213141516171819202122232425262728293031323334353637383940 #!/bin/bash set -ex OLD_PATH="/nfs/ofs-llm-ssd/user/gogongxt/miniconda3" NEW_PATH="/nfs/gogongxt/miniconda3" # 1. 修复 base 环境的 bin 目录 if [ -d "$NEW_PATH/bin" ]; then cd "$NEW_PATH/bin" # -r 参数表示如果 grep 没有输出，xargs 不会执行后面的 sed grep -rIl "$OLD_PATH" . | xargs -r sed -i "s|$OLD_PATH|$NEW_PATH|g" fi # 2. 修复 Conda 自带的配置和环境激活脚本 if [ -d "$NEW_ ...

tbo-sbo tbo 就是 two batch overlap，最开始由deepseek提出，主要的目标是基于deepep的实现思路，优化计算和通信的overlap 在整体的计算流程中，有 attention -> alltoall -> moe -> alltoall -> ...，我们可以构造两个batch，实现一个做矩阵计算，另一个做alltoall通信，从而overlap tbo是给decode准备的，主要原因是decode是带宽瓶颈，计算耗时短，两个batch可以相互干扰影响低 不过实际蚂蚁在h20上部署deepseek时发现了tbo在高并发下反而会更慢，因为h20算力太低，高并发会延迟爆炸不满足slo 因此也是联合deepseek，sglang提出了sbo： 两点优化： 计算shared experts时进行dispatch通信 计算down的时候就边计算边发送，按照block_m粒度进行处理，一个块好了就直接发送 sbo的pr链接： SBO in SGLang: https://github.com/sgl-proje ...

通信操作 主要介绍AllGather和AllReduce，别的通信操作可以参考下面的文档 NOTE Ref: NV NCCL通信库 华为HCCL通信库 blog blog blog nv blog massively-scale-deep-learning-training-nccl nv blog fast-multi-gpu-collectives-nccl AllGather 示意图： 每张卡上都有一个矩阵的一部分，AllGather让每张卡都有完整的矩阵 ring-AllGather 算法： 所谓ring，就是环形处理，每张卡把输出传到下一张卡： 假设一个 的矩阵，每个卡有P分之一参数，每次传输每张卡都是发送了 数据量，总共P-1次传输，总通信量就是： 注意我们讲的通信量一般都是单向的，计算效率是也是除上单向的带宽 AllReduce 示意图： AllReduce操作是将通信域内所有节点的输入数据进行归约操作后（支持 ...

vllm-ascend 与 vllm 的关系 一句话总结：vllm-ascend 是 vllm 的硬件插件（out-of-tree platform plugin），通过 Python entry_points 机制注入，覆盖 GPU 相关的算子/调度/通信层，复用 vllm 的 API 服务、模型定义、请求管理等上层逻辑。 各模块的职责分工 模块 归属 说明 API Server vllm OpenAI 兼容的 HTTP 服务，vllm serve 启动，vllm-ascend 不做任何修改 请求管理/Tokenization vllm 请求接收、分词、结构化输出等 Scheduler 调度 vllm（默认） vllm 的 v1/core/sched/scheduler.py，平台无关的请求调度、KV cache 管理 Scheduler 扩展 vllm-ascend（可选） SchedulerDynamicBatch（动态batch调整SLO）、RecomputeScheduler（PD分离场景重计算），均继承自 vllm Sche ...