刘芷溢

最近更新

  • 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 侧拷贝，所以性能会受到...

    code

    1 篇文章

    • 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.