刘芷溢

最近更新

cuda

12 篇文章

  • FlashAttention v1-v4：从 IO 感知到异步流水与硬件协同

    沿着 FlashAttention v1-v4 的论文、作者博客与源码，拆解 Ampere、Hopper、Blackwell 如何改变 attention 的数据路径、线程分工和瓶颈，并用调度图解释 copy、MMA、softmax 与输出校正如何真正重叠。

  • 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、寄存器重解释与流水线，并用正确性和性能实验检查这些抽象最终生成了什么。

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

    • cuda performance checklist

      CUDA Performance Checklist DRAM vs SRAM DRAM由1个晶体管和1个电容器构成；SRAM: 由6个晶体管构成 SRAM 比 DRAM 更快，但也更贵；SRAM 占用更多空间且发热更多； 实际上SRAM就对应了GPU的Shared Memory，而DRAM对应的则是Shared Memory。 Performance Checklist 合并全局内存访问（Coalesced Global Memory Access） 最大化占用率（Maximize occupancy） 理解是内存受限还是计算受限（Understand if memory or comput...

    • quantization cuda

      Quantization CUDA 动态量化流程 (Dynamic Quantization Flow): 权重和激活都从浮点开始 权重在预处理阶段量化 激活在运行时量化 使用Int8进行乘法运算 使用Int32进行累加 最后rescale回浮点数 未量化 (Not Quantized): 权重和激活都保持浮点格式 所有运算（乘法和累加）都使用浮点数进行 仅权重量化 (Weight Only Quantization): 权重在预处理阶段量化 随后立即反量化回浮点数 激活保持浮点格式 乘法和累加都使用浮点数进行 最后有一个rescale步骤 Dynamic Quantization 数学表达：...

    • 从 GEMM 实践 CUDA 优化

      本文将从最 naive 的 GEMM 实现开始，使用 nsight compute 工具进行性能分析寻找瓶颈并一步步进行优化。通过这种方式来实践 CUDA 中的各种优化技巧，包括 Tiling、Free Bank Conflict、Double Buffer、wmma 指令优化等。

    • GPU 内存系统演进：最大化带宽利用与延迟隐藏的技术路径

      本文探讨了 GPU 内存系统的演进路径，重点介绍了如何通过提升带宽利用率和隐藏延迟来优化 GPU 性能。涵盖了 Little's Law、ILP/DLP 并行扩展、异步内存加载以及启动延迟优化等关键技术。

    • CUDA Optimization for LLM Inference

      在大规模语言模型（LLM）推理中，优化 CUDA 代码对于提升性能和效率至关重要。本文档介绍了一些关键的 CUDA 优化技术，帮助开发者更好地利用 GPU 资源进行 LLM 推理。

    • 一步步实现 CUDA Vector Add 优化

      本文将通过一个简单的 CUDA Vector Add 例子，介绍如何使用 Nsight Compute 工具进行性能分析，并一步步进行优化。