Ascend

1. 部署一个vLLM服务 前一段手上拿到一批昇腾 910C，考虑到可以给大家部署一些vLLM服务来用，本来以为就是按照vLLM的官方教程一步一步走完就行了，事实上也确实是这样的，一步一步走完了流程，但是能用和好用确实是两回事情，后面为了vLLM好用踩了不少坑。 先说一下我们用到的模型的情况，经过对kimi-k3、glm-5.2、glm-5.3-flash、deepseek-v4-flash、deepseek-vision等多个模型的测试，在推理性能、占用资源数、模型能力之间取了一个平衡，最终选择部署deepseek-v4.1-flash（dsv4.1-flash），上下文选择了一个比较适中的范围：600k，同时开启了dspark加速。 2. 多实例负载均衡 最先遇到的问题是单实例服务能力不够，最直接的想法就是搞多实例然后在new-api层做负载均衡，单手后期运维起来由于new-api没有做联动，有点运维难度，且负载均衡导致缓存命中降低，所以选择在vLLM和new-api之间加了一层router来做负载均衡。 最早的router是配置文件形态：引擎列表写死在配置里，加减引擎要改配置、重载。引擎少的时候忍忍也就过去了。 后来换了新一代动态router：引擎主动注册，自带cache-aware调度（前缀亲和，同一前缀尽量打到同一引擎）、失败重试、熔断摘除。和看门狗配合之后形成两层防御： 看门狗管"引擎挂死"：真实推理探针，连续失败整组重启； Router 管"坏引擎别接流量"：探活失败摘出路由，恢复后自动回归。 3. KV缓存问题 影响vLLM的服务质量TTFT的最大因素就是KV缓存，由于我想部署多个实例，导致留给KV的空间就剩下很少了，所以就需要考虑把KV卸载到DRAM或者SSD上，最开始用到的就是MemCache来做KV缓存管理。共享缓存的第一步是 MemCache：vllm-ascend的AscendStoreConnector，在本机内存里开一块池，把前缀 KV 存出去，下一个请求来了先查池。 但是在实际上线服务的时候遇到很多奇怪的问题： 存的日志一切正常，块块都写进去了； 审计计数器显示查询也在跑，exist请求源源不断，还能查到一部分key； 但命中恒为 0，TTFT 纹丝不动。 当时给它起了个名字叫"只存不取"，一度怀疑池子根本没接上。真正的转机是把日志开到DEBUG重启，再加一个受控探针：灌入约50k的前缀token → 用流量把池驱逐干净 → 原样重发。 池侧逐组核对，exists全中：96/96、5/386……每个key都在； 紧接着一行 final hit=0； 请求全量冷预填。 所以不是"只存不取"，是查到了也算零命中。卡点在"查到之后怎么算命中长度"。 对着源码读才发现是粒度基不一致：池里存的是32-token粒度的块哈希，而vLLM core的契约写得明明白白——查命中时调用方必须传"对齐粒度 = 哈希粒度"（源码注释原话：caller passes alignment_tokens = hash_block_size）。我们部署分支的协调器传的却是引擎块大小128。128对32，取出来的哈希和池里存的根本不是同一批值，第一块就miss，后面全部作废。 先试了一行补丁：把128改成32，第一组立刻命中12288——成功了一组。但滑窗组不支持细粒度命中（直接断言），压缩器状态组的粒度又不一样，池侧还有一层按错误粒度截断的覆盖上限。修到最后确认：这个分支（vLLM 0.23）的命中路径就是个半成品，补丁救不完。 真正的解法是上游 v0.30.0rc1 把这条路径正面重写了：查询集收窄到"可前缀缓存的组"，统一调度侧和 worker 侧的命中语义。换镜像，命中日志直接出现，一行补丁都没打。顺手还赚了个惊喜：新版本把查找表挪到了主机内存，腾出来的显存全给了KV池——池容量6.88倍（98万 → 678万token），600K长文加16并发在容量上从此无障碍。 4. 共享KV池 MemCache修好之后，后续的问题也逐渐暴露出来了：池是"每节点一份"的孤立池（元服务绑在本机回环地址上），跨节点从来没有真正通过；引擎一重启，池就空。用户的文档就那几份，却要每个引擎各存一份、重启全丢、再各预填一遍。 4.1 MemCache方案 把各节点的元服务指到同一个中心点：段注册、挂载、全局rank偏移，所有前置步骤全绿。然后跨节点put，100%失败（error -1），重试的put风暴反手把引擎拖死。关掉fabric内存直传的开关复测，依旧。 4.2 Mooncake方案 换分布式KV store：master起来了、几十个段全部挂载成功、对端引擎也能连上——万事俱备。然后put计数恒为 0，设备侧内存拷贝全部超时（EE1002）。想退回普通TCP传输：vllm-ascend硬编码只认昇腾直连协议，NotImplementedError，此路不通。 4.3 问题定位 上述两个方案都是经过单实例验证过的，但是最后都卡死在KV传输上，可奇怪的是这批节点正常训练都NPU传输数据是完全没问题的，凭什么 KV 不行？ ...

1. 快速使用 HCCL（Huawei Collective Communication Library）是基于昇腾AI处理器的高性能集合通信库，其主要功能与作用与Nvidia的NCCL库相似，主要用于集合通信，CANN库种自带一套测试工具用以分析集合通信性能。 1.1 编译环境配置 前置环境配置阶段请参考Llama2部署记录，后续的所有流程均需要在具备HCCL硬件的平台上实现，需要注意的是，建议使用的Ubuntu版本大于等于20.04LTS，否则可能会遇到VSCode不支持的情况。 A. CANN CANN的配置过程参考Llama2部署记录文档的4. 安装CANN章节，并记录安装位置，以备后续使用 完成CANN后，建议在~/.bashrc文件中添加对应的环境变量，例如： # 使用nano打开~/.bashrc文件 nano ~/.bashrc 在文件的最后添加export ASCEND_HOME=/usr/local/Ascend/ascend-toolkit/latest B. OpenMPI 在OpenMPI官网找到对应版本的安装包下载，然后解压缩安装包，以mpich-4.2.0版本为例： # mpich下载链接参考：https://www.mpich.org/static/downloads wget https://www.mpich.org/static/downloads/4.2.0/mpich-4.2.0.tar.gz # 解压缩mpich压缩包 tar -zxvf mpich-4.2.0.tar.gz 配置mpi的安装路径，以/root/mpich路径为例： # 新建文件夹 mkdir ~/mpich # 进入解压缩后的mpich文件夹 cd mpich-4.2.0 # 配置安装路径 ./configure -prefix=/root/mpich --disable-fortran 使用make命令构建mpich的可执行文件，该过程可能耗时比较长 make make install 完后后将mpich添加到环境变量种，并添加对应的链接库和帮助文件 # 修改~/.bashrc文件 nano ~/.bashrc 在最后一行添加以下内容： export MPI_HOME=/root/mpich export PATH=$MPI_HOME/bin:$PATH ...

1. 创建docker 1.1 docker命令 显示全部的镜像：docker images 创建一个容器：docker run ...<查看1.2节内容> 显示全部的容器：docker ps -a 启动一个容器：docker start <id> 停止一个容器：docker stop <id> 进入一个正在运行的程序：docker exec -it <id> /bin/bash 删除一个容器：docker rm <id> 1.2 创建容器 创建一个docker容器的时候需要同时挂载本地工作目录和相关的程序驱动，因此创建docker容器的命令参数较多，以下为一个创建docker容器命令的基本格式： docker run -it --ipc=host --name <name> -v <local_path>:<docker_path> \ --workdir=<docker_path> \ --pids-limit 409600 \ --privileged --network=host \ --shm-size=128G \ --device=/dev/davinci0 \ --device=/dev/davinci1 \ --device=/dev/davinci2 \ --device=/dev/davinci3 \ --device=/dev/davinci4 \ --device=/dev/davinci5 \ --device=/dev/davinci6 \ --device=/dev/davinci7 \ --device=/dev/davinci_manager \ --device=/dev/devmm_svm \ --device=/dev/hisi_hdc \ -v /usr/local/Ascend/driver:/usr/local/Ascend/driver \ -v /usr/local/Ascend/firmware:/usr/local/Ascend/firmware \ -v /usr/local/Ascend/add-ons/:/usr/local/Ascend/add-ons/ \ -v /usr/local/bin/npu-smi:/usr/local/bin/npu-smi \ -v /usr/local/dcmi:/usr/local/dcmi \ -v /etc/ascend_install.info:/etc/ascend_install.info \ -v /etc/bashrc:/etc/bashrc \ -p 223:223 \ -p 224:6006 \ -p 225:8080 \ -p 226:8888 \ -u root \ <image_id> /bin/bash -it：已交互方式运行容器，并分配一个终端 ...

1. Nvidia GPU 参考文献： HeKun-NVIDIA. CUDA-Programming-Guide-in-Chinese \[EB/OL\]. https://github.com/HeKun-NVIDIA/CUDA-Programming-Guide-in-Chinese. Nvidia. Nvidia CUDA编程指南 \[M\]. 1.1. Nvidia, 2008. 1.1 GPU架构 软件层面架构：Nvidia GPU的架构在进行CUDA软件操作时可以抽象为：Thread->Block->Grid三层，通过指定这三个参数可以精确的控制目标运算放在目标的Thread上进行计算。 硬件层面架构：在硬件层可以讲GPU抽象为：SP(Streaming Processor)->SM(Streaming Multiprocessor)->Device三层，基本对应CUDA软件编程层面的Thread->Block->Grid，Nvidia的GPU在指令集上属于是SIMD（Single Instruction Multiple Data）架构，可以实现单条指令在多个数据的并行处理。 集群层面架构：在集群层面，一般按照Host-Device结构划分，其中一个Host可以视作一个由CPU、GPU、内存、硬盘等组成的主机，Device一般特指某一个指定的GPU设备，也就是软件层面的Grid。 图. Nvidia的GPU架构 A. SP / Thread GPU上每个SP主要由ALU（逻辑控制单元）与FPU（浮点运算单元）组成，可以进行简单的逻辑运算与浮点运算，每个SP在对应软件层面可以操作的一个Thread，也是开发者可以操作的最小单元，用户可以借助CUDA在Thread上实现SPMD（Single Program Multiple Data）操作。 B. MP / Block 多个Thread组成一个Block，同一个Block上的Thread共享同一块内存（Shared Memory），这种方法可以实现同一个Block上并行的Thread可以并行的对同一块内存中的数据进行操作。 在进行CUDA编程时，开发者可以借助一个SM控制多个Block中的Thread进行运算，SM中还包含一些必要的寄存器、共享内存、L1Cache、Scheduler、SPU、LD/ST单元等。 C. Device / Grid 多个Block一起共同组成一个Grid，也就是硬件层面与集群层面的Device，借助该概念可以实现集群中多个GPU之间的数据传输和并行计算。 1.2 GPU编程 Nvidia的GPU上有多个Thread可以并行的完成矩阵运算，例如在进行通用矩阵乘（GEMM，General matrix multiply）时，不同Thread在共享内存中取出对应位置的数据进行计算，并讲结果重新放置在共享内存中以实现GEMM的并行计算。 图. 使用共享内存并行计算矩阵乘 在代码层面的实现过程大致如下： // Host层面的矩阵乘，假定矩阵的维度是BLOCK_SIZE的整数倍 void MatMul(const Matrix A, const Matrix B, Matrix C) { // 从显存中加载A和B两个矩阵的数据 cudaMalloc(..., size); cudaMemcpy(..., ..., size, cudaMemcpyHostToDevice); // 从设备显存中申请储存乘法运算结果的矩阵C cudaMalloc(&d_C.elements, size); // 调用内核 MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C); // <<<dimGrid, dimBlock>>> // 从显存中获取计算结果C cudaMemcpy(..., ..., size, cudaMemcpyDeviceToHost); // 释放显存 cudaFree(...); } // 调用MatMul()以实现Kernel层的矩阵乘 __global__ void MatMulKernel(Matrix A, Matrix B, Matrix C) { // __global__ // 每个Thread计算矩阵C中的一个元素，并将结果累加到CValue中 float Cvalue = 0; int row = blockIdx.y * blockDim.y + threadIdx.y; int col = blockIdx.x * blockDim.x + threadIdx.x; for (int e = 0; e < A.width; ++e) Cvalue += A.elements[row * A.width + e] * B.elements[e * B.width + col]; C.elements[row * C.width + col] = Cvalue; } 2. Huawei NPU Huawei的NPU设备Ascend 910B相较于Nvidia的GPU做了一些有针对性的优化，同时由于NPU主要针对神经网络训练与推理，因此其在芯片设计架构上与Nvidia的GPU拥有完全不同的设计思路。 ...