刘芷溢
最近更新
README
CuTe FlashAttention-2 case-study benchmark This directory preserves the exact teaching kernel, benchmark harness, raw results, and figure generator used by the corresponding tom-jerr blog post.
CuTe 详解：以 FlashAttention-2 拆解 Layout、TiledCopy 与 TiledMMA
以一份可运行的 FlashAttention-2 前向 Kernel 为贯穿案例，从坐标函数和线程所有权出发，详解 CuTe 的 Layout algebra、Tensor、TiledCopy、TiledMMA、寄存器重解释与流水线，并用正确性和性能实验检查这些抽象最终生成了什么。
LLM 量化综述：数据格式、PTQ 与 QAT 的算法演进
从 INT8、FP8、INT4、MXFP4、NVFP4 的表示范围与舍入误差出发，按 8 bit、4 bit 和低于 4 bit 梳理 PTQ/QAT，解释各方法解决的问题、相对前作的创新、核心优化与数值稳定性，并对应到 llm-compressor 的工程实现。
从 EAGLE 到 DFlash 2：SGLang 与 SpecForge 的投机解码训练和推理
沿 EAGLE、EAGLE-2、EAGLE-3、DFlash、DSpark 和 DFlash 2 的演进，解释 draft 的训练目标、数据对齐、推理状态与验证预算，并对照 SpecForge 训练和 SGLang serving 源码串起完整流程。
SGLang HiCache：组件、工作流与 L1/L2/L3 读写时序
从 HiRadixCache、HiCacheController、Host KV Pool 与 Storage Backend 的职责出发，沿 Scheduler 主循环拆解 L3 query、hit、prefetch、L2 load、L2 write 和 L3 write 的准确时序。
从 Bank Conflict 到 CUTLASS Swizzle：推导 ldmatrix 的访存布局
从 shared memory 的 32 个 bank 出发，逐地址推导 ldmatrix 的行地址、寄存器分配与 bank conflict，再解释 CUTLASS 的 XOR permutation、CuTe Swizzle 参数和 producer/consumer 布局约束，附交互图与可运行验证程序。
CUDA 中的 ILP 与 MLP：从依赖链到延迟隐藏
从 warp scheduler、依赖链与 Little's Law 出发，区分 CUDA 中的指令级并行和内存级并行，并结合 GEMM、GEMV、FlashAttention、可对照的 Kernel 与 Nsight Compute 指标说明何时以及如何提升二者。
cuda basic
CUDA Basic Concept CUDA Async Copy & Pinned Memory 通常情况下，CPU 分配的内存是 Pageable（可分页） 的。操作系统为了节省物理内存，可能会把这些内存交换到磁盘上。 当你要把数据从 Pageable 内存拷贝到 GPU 时，CUDA 驱动其实做了一个幕后动作： 先在系统中申请一块临时存储区（即 Pinned Memory）。 将数据从 Pageable 内存拷贝到这块 Pinned Memory。 通过 DMA（直接内存访问） 将数据从 Pinned Memory 传给 GPU。 实际上多了一次 CPU 侧拷贝，所以性能会受到...
