学习让我快乐

PD 分离中的 GDR

在这篇文章我提到过我们基于 vllm 实现了 PD 分离, 而且设计思路恰好与 Nvidia Dynamo 撞车了. 大致查询链路也与 Dynamo 相似, 简单来说是: 请求 R 从 D 节点接入, D 节点决策模块确定请求 R 应在某个 P 实例上完成, 此时 D 为 R 分配 kvcache block R.d_kv_blocks, 并会调用 P 实例 do_prefi...

Posted by w@hidva.com on April 17, 2025

一定要保存 asyncio.Task 对象啊!

很早之前, 在我刚接触到 asyncio 观其文档时, 就看到其反复提及过: Important: Save a reference to the result of this function, to avoid a task disappearing mid-execution. The event loop only keeps weak references to tasks...

Posted by w@hidva.com on March 26, 2025

NCCL AllReduce CudaGraph 时 hang 问题分析

我们同学在压测时遇到过压测进行到一段时间之后会几率性阻塞, 于是让我帮忙看一下. 于是我拉起了我的开发调试镜像, 里面存放着各种预编译好的, 带有调试信息的 pytorch, nccl 等一大坨我们会用到的组件, 然后开始启动压测复现, 复现过程不再多提, 最终我拿到了一个现场, 赶紧挂起 cuda-gdb 到一个 worker 进程, 可以看到: 0x00007f6489d08c9b i...

Posted by w@hidva.com on March 13, 2025

NCCL 源码解读(17): Primitives Simple

本文是 NCCL 源码解读系列之一, NCCL 源码解读系列使用 NCCL 版本: v2.25.1-1. 本文介绍 Primitives<ProtoSimple> 相关实现细节. 前景提要在之前关于 p2p transport 的文章中, 我们已经介绍了在发送方会在 p2pSendSetup 时分配 sizeof(ncclSendMem) 大小显存, 接收方会分配 ...

Posted by w@hidva.com on March 9, 2025

mpMath: 与微信公众号的斗智之旅(2)

这是一篇水文, 可是为了解决 mpMath 与微信公众号新版编辑器不兼容的原因, 我浪费了我美丽的周日下午, 不把这个经历水出来我实在意难平啊!意难平!! 前景提要: mpMath: 与微信公众号的斗智之旅(1) 事情好像到这里就要结束了? 但我想着趁热打铁, 总有一天我还是会面对基于 ProseMirror 的新编辑器, 这一关迟早要过的. 简单看了下 ProseMirror 文...

Posted by w@hidva.com on March 2, 2025

消失的 1 秒钟

笔者最近在参与 LLM 推理框架 PD 分离相关工作, 在完成开发压测过程中发现一个很奇怪的点: 有一些请求会在被放入 asyncio.Queue 之后要 1 秒钟左右之后才会被 Queue 消费端可见. 这让我感到有些意思, 于是准备进入深入分析一番. 首先我要有一个趁手的工具, 最起码可以展示这 1 秒内各个线程都在做什么. 当然首推的就是 py-spy 了, 但 py-spy 有个问...

Posted by w@hidva.com on March 2, 2025

mpMath: 与微信公众号的斗智之旅(1)

这是一篇水文, 可是为了解决 mpMath 与微信公众号新版编辑器不兼容的原因, 我浪费了我美丽的周日下午, 不把这个经历水出来我实在意难平啊!意难平!! mpMath 的一键公式转换一直是我做出来的东西中我第二引以为豪的! 每次点击 “公式转换” 之后 mpMath 将文档中一大坨公式一一绘制成 svg 我都感到了多巴胺爆棚. 而且 mpMath 也帮我认识好多数学科普 up 主...

Posted by w@hidva.com on March 1, 2025

asyncio 可观测性增强

如我们之前在为什么协程中提到过的: 内核中的公平调度器提供的是抢占着调度能力, 在线程调度到具体的 CPU 上执行一段时间之后, 即使这个线程对应的逻辑执行流仍有指令可以继续执行, 内核也会强行挂起当前执行线程, 转而从可运行队列中选择下一个线程来调度其到 CPU 上运行. 而位于线程中的用户态调度器提供的是协作式调度, 她仅在每个协程逻辑执行流主动让出对线程的占有时才有机会调度...

Posted by w@hidva.com on December 18, 2024

mscclpp DeviceSyncer 真的能 sync 么?

DeviceSyncer 作为一个朴素的工具类, 其实现真的很简短, 就是如下 15 行代码: MSCCLPP_DEVICE_INLINE void sync(int blockNum, int64_t maxSpinCount = 100000000) { unsigned int maxOldCnt = blockNum - 1; __syncthreads(); if (...

Posted by w@hidva.com on December 1, 2024

令人哽咽的 python asyncio 调试

笔者近来一直在参与某个不可名状的项目, 从上到下都很重视, 每天都会拉起 Scrum 会对齐进度, 虽然压力满满但也斗志昂昂. 然后最近项目出现了一个 bug, 此 bug 严重阻塞了各位同学后续进一步验证工作, 已经连续几天在 Scrum 会上被点到. 于是我开始接手这个 bug 的分析定位工作, bug 的修复到很简单, 就是普通的多线程并发问题. 时间来到了当天晚上 7 点左右, 我已...

Posted by w@hidva.com on November 23, 2024

为什么协程

老板让做一篇基础的协程分享介绍. S O W H Y C O R O U T I N E 在我们之前的文章 “为什么线程” 中, 讨论了线程的概念: 一个线程对应一个执行流. 增加一个线程意味着增加了一条逻辑执行流, 从而增强了应用的表达能力和处理能力. 当应用编写完逻辑执行流后, 负责为每条逻辑执行流创建对应的线程, 而内核则负责将这些线程调度到具体的 CPU 上执行, 以推动逻...

Posted by w@hidva.com on November 5, 2024

A note on A note on A note on the algebra of CuTe Layouts

在 A note on A note on the algebra of CuTe Layouts 末尾遗留了一个问题, 就是: In particular, I am trying to prove $M$ is left divisible by $N_{\varphi(\beta)} r_{\varphi(\beta)}$, \[\begin{align} r_{\varph...

Posted by w@hidva.com on October 21, 2024

A note on A note on the algebra of CuTe Layouts

这篇文章纪录了对 A note on the algebra of CuTe Layouts 的学习笔记, 侧重于记录我当时没有看懂的部分, 以及对书中部分知识点的一些扩展. 非常零散, 不成系统; 最好是结合 A note on the algebra of CuTe Layouts 使用. A note on the algebra of CuTe Layouts 是之前在学习 c...

Posted by w@hidva.com on October 19, 2024