hliyuxin的专栏

原创 昇腾与英伟达硬件架构对比

摘要：昇腾910C计算单元配置包括24个AICore、4个TensorCore和128个CudaCore，FP16稠密算力达989TFLOPS。存储方面配备80GB显存和256KB片上缓存，显存带宽3.35TB/s。高速互联采用900GB Nvlink，超节点GB300NVL72提供1.8TB双向带宽。交换芯片容量计算逻辑显示单芯片28.8Tb，整机57.6Tb，性能对标博通TH5。互联架构采用8芯片16DIE设计，Level1和Level2交换芯片带宽分别为6.272Tb/s和10.752Tb/s，性能对

原创 昇腾与英伟达编程架构对比

昇腾C++代码是显式管理数据搬运和流水线的，而CUDA代码更像是为每个线程独立编写逻辑，硬件自动完成并行。

原创 HCCL底层通信分析

摘要：本文分析了昇腾AI处理器两种通信路径——SDMA（卡间HCCS总线）和RDMA（跨机RoCE网络）的数据传输机制。SDMA路径通过HcclD2DMemcpyAsync实现HBM间直接异步传输，不经过AICore；RDMA路径则通过RoCE网卡DMA引擎直接读写注册的HBM物理地址，全程规避CPU参与。对比OpBase和Zcopy两种模式，前者需userIn到cclIn的拷贝但内存要求宽松，后者实现零拷贝但需预先注册固定大小内存。两种方案均利用专用硬件（SDMA/RoCE引擎）直接操作HBM，显著降低A

原创 Deepseek DSA 原理与训练

DSA（Deepseek Sparse Attention）技术摘要 DSA通过稀疏注意力（TopK选择）减少计算量，类似MoE思想。其核心是Lighting Indexer模块，从长上下文中筛选最相关的Token（如128K选2048），将计算复杂度从O(L²)降至O(LK)。训练分为两阶段： 密集预热阶段：冻结主模型，仅训练Indexer模仿全量注意力的分布（KL散度损失），使用2.1B Token数据。 稀疏训练阶段：联合优化Indexer和主模型，但梯度分离——Indexer通过子集KL损失

原创 NCCL 2.30版本& DeepEPv2 通信库

本文对比分析了NCCL通信模式和DeepEP V1/V2的优化策略。NCCL 2.18+提供CE（机内）、RMA（机间）和GIN（类IBGDA）三种基于对称内存的通信模式。DeepEP V2相比V1的主要改进包括：采用NCCL GIN异步操作，自主管理对称内存；将SM等待机制改为mbarrier warp级同步，SM占用从24个降至4-6个；通过TMA异步加载和GIN异步发送实现流水线，使SM始终处于计算状态；采用带宽瓶颈建模动态确定SM数量；拆分Copyepilogue为独立kernel实现PDL流水线。

转载 多模态理解模型

模态有图片、文本、视频、音频、触觉（传感器）等。每个模态可以映射到一个向量，模态之前可以做对齐（Align)

原创 PD分离推理的部署与调度

阶段决策内容依据资源池划分哪些GPU做P, 哪些GPU做D离线部署时决定Prefill候选筛选过滤不健康/过载节点实时状态Prefill打分前缀缓存命中+空闲度缓存状态+负载Decode候选筛选过滤显存不足的节点实时状态Decode打分队列长度+网络亲和性负载+拓扑联合配对选最优（P*, D*)组合综合评分执行Prefill->传输->Decode确定性流程举例：一个集群1024张GPU卡。

原创 GB300&VR200成本列表

原创 PD分离推理网络亲和性调度

一般先根据负载均衡选Decode节点（Decode节点是长期持有状态），再根据Decode节点，就近选Prefill节点（Prefill节点是无状态的，位置灵活）。更进一步，当同一个用户的请求有历史KV Cache（比如多轮对话）时，优先把新请求调度到缓存了上一轮KV Cache的Decode节点附近的Prefill节点，避免KV Cache的二次迁移。Decode节点有状态，不能随便动，一个请求一旦开始Decode，它的KV Cache就住在那个Decode节点的显存里。

原创 大模型batchsize与梯度累积

在一个step最后optimizer 执行前，W的梯度会进行num_microbatches轮的累积，每轮还需要按DP维度进行累积。micro_batch_size 设置比较大，激活值存储很容易OOM，激活值需要至少保存micro_batch_size*seq_len*hidden_size*num_layers。通常说的梯度累积是W的梯度，X的梯度会在zbv的B阶段处理完就会丢弃。micro_batch_size和num_micro_batches两个参数控制的是不同资源。

原创 CP并行通信的几种方案

输入按Seq维度汇总，按Head维度切。（切输入，非TP维度的切参数）举例： CP = 2, H = 8192, nHead = 16。，恢复按Seq维度切，按Head维度汇总。长文本CP 切分，共2次All2All。

原创 NCCL SHARP 和 TREE算法

其中reduce和broadcast由Nvswitch做，但由GPU发起控制指令驱动。其中GPU有一个UC（unicast)内存地址和MC（multicast) 内存地址, 都映射到同一块物理地址，可以由GPU和nvswitch共同访问。NVLS(nvlink sharp）算法，H100开始一代用于机内allreduce操作。

原创 NCCL 初始化COMM与AllGather通信流程

GPU和CPU之间会有一个内存映射过程，让CPU执行的RDMA指令能找到GPU的内存地址。1. runRing算法中, 会去选择Simple还是LL协议。2.GPU -> CPU Proxy 模式。

原创 NCCL Simple & LL协议

NCCL用time = latency + nByte/BW的代价模型遍历所有可用的（算法、协议）组合，选代价最小的一个。LL延迟低但带宽打五折，Simple延迟高但带宽全利用，消息越小越倾向LL，消息越大越倾向Simple, 拐点大约在几十KB到几百KB（取决于具体硬件拓扑）。总延迟=fence开销+ T1(tail传播数据包）+ B接收处理 + T2（head传播控制包）如果大包也用LL协议，带宽利用不上来，也影响整体时延。清楚可以看到，总延迟LL协议比simple协议要小。为什么大包不用LL协议？

原创 Flash Attention反向梯度优化显存

特性传统 Attention 反向Flash Attention 反向显存占用(存储整个P矩阵)(仅存储 LSE)显存访问量IO 瓶颈高 (频繁读写 HBM)低 (在 SRAM 中完成计算)计算量较低较高 (通过重计算换取空间)主要性能限制内存带宽算力 (Compute Bound)S%3DQK%5ETdQ%3DdSKO%3DP%20VO%28Nd%29。

原创 大模型显示优化之ZeRO-1/ZeRO-2/ZeRO-3

zero-1、zero-2、zero-3 是deepspeed的配置方法，对应megatron也有相应的方法，Megatron-LM 的实现方式：Distributed Optimizer（分布式优化器）。等效于 ZeRO-1，Megatron 的 Distributed Optimizer 默认行为就是将优化器状态（Optimizer States）均匀地切分并分布在数据并行（DP）组的所有 GPU 上。等效于 ZeRO-2，由于 Megatron 通常结合混合精度训练，它在计算完梯度后，会通过。

原创 Deepseek MLA和MLA absorb

主要方法是将Q K V用一个相对低秩的矩阵来存储，可以降低显存空间，以及一部分计算量， 为解决位置相关的信息丢失问题，将压缩后的QK又单独concat。MLA和MHA的区别是QKV计算的过程不同，其后的Attention计算以及linear流程基本是一致的。图中可以看到tokens输入的d维度是7168，q_down压缩之后的维度是1536+576，q_up解压后的维度是32*（128+64）=6144，整体维度都是压缩了。本质是矩阵三连乘问题，应该是先左乘还是右乘，显存占用以及计算量越少。

原创 GPU Bound问题与解决

GPU/NPU 容易出现Host bound和Rte bound问题。是一个核心性能概念，指。你可以把它理解为“CPU干活太慢，让NPU这位‘超级员工’无事可做，造成了算力浪费”

原创 AI 分享内容

标志事件：1980年， DEC（大型机小型机生产商）落地XCON系统，解决人工配置运维问题，年处理超过8万份订单，准确率高达95%-98%，为公司节省3千万美元/年。时间：2017年 Google 的几位研究员，提出Transformer架构，解决了传统神经网络处理长序列数据的难题，奠基了在自然语言处理领域的理论基础。2023年-2024年，国内大厂诞生练丹热潮，基础大模型，百度文心一言，阿里通义千问，腾讯元宝，美团龙猫等。Manus，做一个通用agent，人可以用电脑做的事，或者有手机做的事情。

原创 DeepSeek V3 训练和推理

DeepSeek-V3的训练使用了：整个训练过程（包含预训练、上下文扩展和后训练）在2048张H800 GPU上完成，总计耗时约，成本约，这在同类模型中极具性价比。根据DeepSeek官方技术报告及多方技术解读，DeepSeek-V3在训练阶段采用的。。将256个路由专家分布到不同GPU上，每张卡托管约4个专家。

原创 PP pipeline并行算法总结

DualPipeV通过创新的V形调度模式优化了传统DualPipe架构，在保持相同流水线深度和并行效率的同时，显著降低了内存占用和设备需求。其核心优势包括：1）单行道设计实现双向流效率；2）参数内存减半，解决显存冗余问题；3）设备数量减少50%。代价是通信量翻倍，但激活内存和流水线气泡保持不变。这种改进使DualPipeV在资源利用和性能之间取得了更好平衡。

原创 大模型训练CP切分（与TP、SP结合）

本文重点探讨长文本训练的CP（Context Parallelism）技术及其与TP（Tensor Parallelism）、SP（Sequence Parallelism）的结合。当处理超长上下文（如1M token）时，CP与TP正交，SP=CP。在TP=1时，CP将输入按token维度切分，通过all-gather获取全局K/V；在TP>1时，CP与TP协同工作，矩阵按列切分到各GPU，经过多次聚合操作后数据按CP维度切分，此时SP默认等于CP。

原创 FlashAttention 优化MAC/MF

FlashAttention是一种重要方法，优化显存占用MF（Memory Footprint) 和显存访问量MAC (Memory Access Cost)，目前已经应用于各大主流框架中。

原创 强化学习与思维链

摘要：强化学习领域探讨了SFT与强化学习的关系，比较了ORM和PRM两种奖励模型的优缺点，分析了基于规则和基于模型的奖励方法，并介绍了PPO和GPO算法。思维链技术部分阐述了提升模型推理能力的方法，包括增加token数和多输出生成，详细说明了思维链的工作流程（并行搜索+自我优化）以及DeepSeek-R1的训练过程。最后讨论了强化学习与蒸馏技术的结合以及推理模型的未来发展方向。全文聚焦于人工智能模型优化技术的前沿进展。

原创 推理PD分离主要流程与cache的使用

Prefill过程的result最后一行输出是第一个token的向量，这个向量经过FNN等处理后，进入Decode过程作为输入H。KV矩阵通过存储传递到Decode过程。Decode的将H通过QKV的处理，产生Q矩阵（1*h维），K，V矩阵（s+1)*h维；然后自回归迭代这个过程。

原创 大模型TP、SP、EP切分

文章摘要：本文对比了传统Dense大模型和MoE大模型在并行计算方式上的差异。传统模型采用TP切分，Attention和FNN部分都需进行前向和反向各一次AR计算。而加入SP并行后，计算模式发生变化：前向的g{}变为AG，反向的g{}变为RS。文中以FNN部分为例说明计算模式转变，指出Attention部分也可采用类似处理方式。这反映了模型并行计算策略的演进过程。

转载 大模型transformer架构训练和推理方法

因为训练可以“靠老师”，推理还得“靠自己”，这样推理时遇到的错误输出对于下次推理来说就是在训练数据分布之外（out of distribution）的异常输入，所以会导致用Teacher Forcing模式训练出来的模型在训练环节和预测环节存在行为差异。具体来说，Teacher Forcing就是每次推理给解码器输入时，不使用前次推理的输出作为下一次推理的增加输入，而是使用训练标签的真值（ground truth）作为下一次推理的增加输入。在训练时，我们可以使用与推理时相同的方法，即用自回归模式进行。

原创 大模型切分总结

矩阵运算基础 Y=XW+B， X：输入， W:参数， Y：输出。三步：正向计算Y、反向计算梯度、反向更新参数。multi-head attention, TP 为Head group。长文本输入，ring-attention或者all2all。laynorm函数, 通常与TP相等。FNN层，降低参数规模。模型按stage层切。切（all2all)

原创 集合通信算法总结

直接接口， 使用IBGDA，可以将控制面bypass cpu，原始版本支持IB，修改版本可以支持RoCE。（1）通信域大部分是单独机内或者单独机间（同号卡）， 比如TP、SP、CP机内，部分DP只有机间。借轨的思路和N卡PXN类似, RoCE发到目的同号卡， 目的同号卡再转发到机内的其它节点。有1F1B，zero-bubble流水线overlap方案，目前zero-bubble用得最多。zero-bubble， F、B、W三阶段切分。1. H卡N-1步算法，通过算法消除多打1。核心是借轨+源聚合。

原创 transformer各层的输入输出

层/组件输入形状输出形状核心功能[B, S][B, S, D]将离散token转为连续向量[B, S, D]注入位置信息[B, S, D]计算输入序列内部的上下文关系[B, S, D]非线性变换，特征加工[B, T, D]计算已生成输出序列的内部关系[B, T, D]让解码器关注编码器信息[B, T, D]非线性变换，特征加工[B, T, D][B, T, V]产生下一个词的概率分布关键记忆点维度守恒： 除了最开始的Embedding和最后的输出层，所有核心层的输入输出维度都是。信息流。

原创 美团云OVS-DPDK 2017年工作总结(SRIOV、热升级、热扩容、跨OVS-kernel热迁移）

2017年，美团云网络研发团队，对OVS-DPDK的高可用方面，做了很多自研工作，解决了很多关键难题，同时也积累了很多switch性能优化的经验。1. SRIOV方式优化local性能         KNI和TAP是常用的DPDK网络设备和内核数据交互的方案，为了提升Local port 的性能，满足一些宿主机特殊场景。美团云将SRIOV和BOND结合，优化OVS-DPDK bridge

原创 2017 DPDK summit 主讲美团云&OVS-DPDK

会议新闻链接 https://mp.weixin.qq.com/s?__biz=MzI3NDA4ODY4MA==&mid=2653334447&idx=1&sn=915787fe02c50f264a7c488dea6fb3d6&chksm=f0cb5c28c7bcd53eefef4681e3beba1fb7a9420d0c33cc09e9cd169536f91c7a420e27d939be&mps

原创 SR-IOV技术优化DPDK架构下Local port性能

转载请注明出处： http://blog.youkuaiyun.com/hliyuxin/article/details/52710749     DPDK程序中，网卡被DPDK igb_uio驱动接管，如果需要与linux系统数据交互，现在主要有KNI和TAP两种方式，两种方法都是创建虚拟设备用于收发报文。现在DPDK sample主要推荐kni方式, OVS-DPDK中bridge的local por

原创 ovs-dpdk和ovs-kernel的checksum处理

转载请注明出处： http://blog.youkuaiyun.com/hliyuxin/article/details/52710749 现在遇到问题，部分docker的机器，迁移到ovs-dpdk上之后有tcp checksum错误, ethtool关闭docker 网卡tx checksum之后，服务正常。 查看原ovs-kernel的配置，docker 网

原创 ovs 隔离vlan packet

OVS的一部分非vlan port,  会收到vlan port发出的广播包。OVS port 配置如下Bridge br_t        Port br_t            Interface br_t                type: internal        Port "client-001"            Interface "

原创 ovs-dpdk 在美团云环境中的应用与实践

OpenvSwitch 以其丰富的功能，做为多层虚拟交换机，已经广泛应用于云环境中。Open vSwitch的主要功能是为物理机上的VM提供二层网络接入，和云环境中的其它物理交换机并行工作在Layer 2层。传统host ovs工作在内核态，与guest virtio的数据传输需要多次内核态和用户态的数据切换, 带来性能瓶颈.  ovs-dpdk 结合了DPDK和vhost-user技术

原创 ovs支持dpdk-bond类型port更改

很多公司原linux的物理机网络是一块网卡的两个port做bond， ovs 需要支持bond port

原创 ovs-dpdk bridge internal port 访问延迟问题, fix

现有的云主机结构中，物理机通过bridge的brpri和linux系统通信（如dns, 云硬盘等），当主机负载比较高时，访问主机的管理口出出现延迟。 Bridge brpri Controller “tcp:127.0.0.1:56XXX” is_connected: true fail_mode: standalone

原创 vhost device still attached , ovs crash bug fix

vhost device still attached , ovs crash

原创 OVS-DPDK VM出现 virtio_net virtio4: output.0:id 30 is not a head 问题 dpdk-vhost bug fix

云环境中，部分vm运行一段时间，不对外发包，可正常收包，system log中出现：kernel: virtio_net virtio4: output.0:id 30 is not a head。      gdb ovs进程，发现rte_vhost_dequeue_burst 中，if (vq->last_used_idx == avail_idx) return ,直接判断没有数据可读。