学习让我快乐

一定要保存 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_{φ (β)} r_{φ (β)}

, \[\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

cuda pipeline 源码解析

之前在学习 cutlass 实现时, 看到其使用了 cuda pipeline 设施, 然后看了下 pipeline 的文档似懂非懂 :-( // 如下这段代码会被 thread block 中所有线程协作执行. __shared__ cuda::pipeline_shared_state<thread_scope_block, stages_count> shared_sta...

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