蠕虫式驻留:劫持整个 GPU 的调度权
作者梳理 GPU 内核边界的演变:传统模型中 host 逐个派发 kernel,kernel 结束后再协调下一步;Hexagon-MLIR 的融合思路把多个算子合并到一个可推理区域,减少中间值的搬运与读回;Hazy Research 的 megakernel 进一步把解释器、依赖调度和资源管理搬进一个常驻 GPU 的程序,让 GPU 内部直接消费任务,不必每次回到 host 做原生 dispatch。 文章的核心论点:去掉 kernel 边界不是免费的,同步、竞争、死锁、共享内存/register/SM 时间限制都要在常驻程序里自己管;不同层次(编译器、host runtime、设备侧软件、硬件)各自掌握不同信息,把调度决策下移到更靠近硬件的一层只有在该层拥有 host 用不上的实时信息时才划算,否则只是把复杂度搬了个位置。 还提到 AMD RX 7900 XTX 上常驻 tinygrad UOps 的实验、Flapping Airplanes 的 Worm 方案,以及 ThunderKittens、HipKittens、TileLang 等对 tile/worker/pipeline 控制权的不同取舍。 作者强调 capacity ≠ coordination,机器可以并行度极高但调度差就会浪费;常驻 kernel 一旦能持续接受新任务、解释执行、协调资源,就开始像 runtime 本身的一部分,离"劫持整个 GPU"只差一步。 













