LeetCUDA

GitHub
10.3k 1.1k 较难 1 次阅读 今天GPL-3.0开发框架
AI 解读 由 AI 自动生成,仅供参考

LeetCUDA 是一套专为初学者打造的现代 CUDA 学习笔记与实战代码库,旨在通过 PyTorch 环境帮助用户轻松掌握 GPU 编程核心技能。它主要解决了开发者在学习底层高性能计算时面临的门槛高、资料分散以及缺乏高质量参考实现的痛点。

这套资源非常适合希望深入理解深度学习底层原理的 AI 研究人员、系统工程师以及想要提升模型推理速度的开发者。无论是刚接触 CUDA 的新手,还是寻求优化现有算子的资深专家,都能从中获益。

LeetCUDA 的技术亮点十分突出:它不仅收录了超过 200 个实用的 CUDA 内核示例,还深入解析了 Tensor Cores、HGEMM(半精度通用矩阵乘法)以及 FlashAttention-2 等前沿技术。其中,其手写的 HGEMM 实现能发挥出媲美 cuBLAS 库 98% 至 100% 的峰值性能,而基于纯 MMA PTX 指令优化的 FlashAttention 则展示了极致的显存访问效率。此外,项目还涵盖了 TF32、F16、BF16 等多种精度格式的实战应用,并配套了百余篇相关技术博客。通过提供从理论到落地的完整路径,LeetCUDA 成为了连接深度学习算法与硬件加速之间的桥梁。

使用场景

某 AI 初创公司的算法工程师团队正致力于将自研的大语言模型推理速度提升 30%,以满足实时对话服务的低延迟要求,他们决定对核心的矩阵乘法与注意力机制算子进行深度定制优化。

没有 LeetCUDA 时

  • 学习门槛极高:团队成员需从零摸索复杂的 Tensor Core 编程与 PTX 指令集,缺乏系统性的现代 CUDA 学习笔记,导致上手周期长达数周。
  • 性能调优困难:自行编写的 HGEMM(半精度通用矩阵乘法)内核效率低下,仅能达到 cuBLAS 库 60%-70% 的算力峰值,成为推理瓶颈。
  • 显存带宽受限:在处理大维度(Large HeadDim)的注意力机制时,无法有效利用共享内存优化,导致显存访问延迟高,难以复现 FlashAttention-2 的高性能。
  • 试错成本高昂:缺乏经过验证的 200+ 高质量 CUDA Kernel 参考代码,每次修改底层逻辑都需反复调试,严重拖慢迭代进度。

使用 LeetCUDA 后

  • 快速掌握核心:借助 LeetCUDA 提供的现代化教程与 PyTorch 绑定示例,团队在一周内便掌握了 TF32/F16/BF16 数据类型及 Tensor Core 的高效用法。
  • 逼近硬件极限:直接复用并微调 LeetCUDA 中的 HGEMM 实现,利用 WMMA 与 CuTe API 成功将矩阵乘法性能提升至 cuBLAS 的 98%-100%。
  • 突破维度限制:采用 LeetCUDA 中基于纯 MMA PTX 优化的 FlashAttention 变体,轻松支持 D=320~1024 的大维度注意力计算,显著降低 SRAM 复杂度。
  • 加速落地验证:基于仓库内 200+ 个现成的高质量 Kernel 模板进行二次开发,大幅减少底层编码时间,让团队能专注于业务逻辑的创新。

LeetCUDA 通过提供工业级性能的算子模板与现代化学法路径,帮助团队将原本需要数月完成的底层优化工作压缩至数周,成功实现了模型推理速质的飞跃。

运行环境要求

操作系统
  • 未说明
GPU

必需 NVIDIA GPU(支持 Tensor Cores),文中测试涉及 L20, RTX 4090, RTX 3080 Laptop,需安装 CUDA 工具链以编译 PTX/MMA 指令

内存

未说明

依赖
notes该项目主要是 CUDA 学习笔记和内核实现集合(包含 HGEMM, FlashAttention 等),核心代码为 C++/CUDA。虽然提到结合 PyTorch 使用(通常指通过 Python 绑定调用),但 README 未明确指定具体的 Python 版本、PyTorch 版本或操作系统要求。运行需要能够编译自定义 CUDA 内核的环境。部分高级功能依赖纯 MMA PTX 指令,对显卡架构有一定要求(如 Ampere 及以上以支持 TF32/BF16/F8 等特性)。
python未说明
PyTorch
CUDA Toolkit
Triton (部分内核)
CUTLASS (部分内核)
LeetCUDA hero image

快速开始

📚 LeetCUDA:面向初学者的现代 CUDA 学习笔记,结合 PyTorch 🐑

精选|HelloGitHub

📚 LeetCUDA:包含 Tensor/CUDA 核心、TF32/F16/BF16/F8,以及使用 PyTorch 的 📖200+ CUDA 内核🔥📖100+ LLM/CUDA🔥 博文,还有能够达到 cuBLAS 98%~100% TFLOPS 性能的 📖HGEMM⚡️,以及利用 Tensor Cores 和纯 MMA PTX 实现的 📖flash-attn⚡️。♥️ 如果可以的话,请为我点个 ⭐️ 星来支持一下吧,兄弟 ~ ♥️

©️引用🎉🎉

@misc{LeetCUDA@2025,
  title={LeetCUDA: A Modern CUDA Learn Notes with PyTorch for Beginners},
  url={https://github.com/xlite-dev/LeetCUDA.git},
  note={Open-source software available at https://github.com/xlite-dev/LeetCUDA.git},
  author={DefTruth and Many Others},
  year={2025}
}

📖 新闻 🔥🔥

  • [2026/03] Cache-DiT 🎉v1.3.0 已发布,主要更新包括:带有 batched P2PRing 注意力,USP(混合 Ring 和 Ulysses),混合 2D 和 3D 并行计算(💥USP + TP),以及 VAE-P 通信开销的降低。

arch

  • [2026/04]: 🤖ffpa-attn 发布了!这是又一个更快的 Flash 预填充注意力机制,具有 O(1)🎉SRAM 复杂度,适用于大头维度,相比 SDPA EA 提升了 1.8x~3x↑🎉:📈L20 ~1.9x↑🎉📈A30 ~1.8x↑🎉📈4090 ~2.1x↑🎉。目前,FFPA 支持自注意力、交叉注意力、分组/多查询注意力以及带因果关系的大头维度注意力(D=320~1024)。而标准的 FlashAttention-2 只支持头维度不超过 256 的情况。
image
  • [2024/12]: ⚡️HGEMM 发布了!我们从零开始使用 Tensor Cores 结合 WMMA、MMA 和 CuTe API 编写了 HGEMM,实现了峰值🎉性能。

📖 目录

📖 HGEMM 基准测试 🎉🎉

目前,在 NVIDIA L20、RTX 4090 和 RTX 3080 笔记本上,与 cuBLAS 默认的 Tensor Cores 算法相比,本仓库中的 HGEMM (WMMA/MMA/CuTe)(蓝色🔵)可以达到其(橙色🟠)性能的 98%~100%。更多详情请查看 toy-hgemm 库⚡️⚡️HGEMM⚡️⚡️ 仓库。

toy-hgemm-library

📚特性 📚特性 📚特性 📚特性
✔️CUDA/Tensor Cores ✔️循环遍历 K ✔️分块(BMxBK) ✔️线程分块(T 8x8)
✔️WMMA(m16n16k16) ✔️MMA(m16n8k16) ✔️LDST打包(128位) ✔️SMEM填充
✔️异步复制 ✔️分块 MMA ✔️分批 Warp ✔️多阶段(2~4)
✔️寄存器双缓冲 ✔️块置换 ✔️Warp置换 ✔️SMEM置换(CuTe/MMA)
✔️集体存储(Shfl) ✔️NN布局 ✔️TN布局 ✔️SGEMM FP32/TF32

📖 FA2-MMA 基准测试 🎉🎉

我还使用纯 MMA PTX 指令实现了 FlashAttention-2,它支持多阶段、分块 MMA、分块线程束、共享 KV SMEM、完全共享 QKV SMEM预取 Q s2r预取 K/V g2sQKV 精细粒度分块、集体存储等功能。更多详情请参阅 flash-attn⚡️⚡️

flash-attn-mma

📚功能 📚功能 📚功能 📚功能
✔️张量核心 ✔️循环遍历 N/D ✔️分块(Br, Bc) ✔️MMA(m16n8k16)
✔️打包 LDST(128 位) ✔️SMEM Swizzle/填充 ✔️异步复制 ✔️分块 MMAs
✔️分块线程束 ✔️多阶段(1/2) ✔️集体存储(Shfl) ✔️拆分 KV/Q
✔️共享 QKV SMEM ✔️预取 Q s2r ✔️预取 KV g2s ✔️QKV 精细粒度分块

目前,对于小规模注意力 (B<=4, H <=48, SeqLen <= 8192, D <= 64),在某些设备上它的运行速度可以超过 FA2/SDPA。例如,在 NVIDIA RTX 3080 笔记本电脑上,📚 拆分 Q + 完全共享 QKV SMEM 方法可以达到 55 TFLOPS (D=64),几乎比 FA2 快 ~1.5x 🎉。而在 NVIDIA L20 上,🤖ffpa-attn 方法可以达到 104 TFLOPS (D=512),几乎比 SDPA(高效注意力)快 ~1.8x 🎉。然而,对于大规模注意力,性能仍存在差距。敬请期待后续更新 ~(MMA 精度 F16/F32,softmax 精度 F32 与 FA2 MMA/softmax 精度 F32,👇基准测试)

算法 (B,H,N,D) RTX 3080 笔记本电脑 L20 RTX 4090
FlashAttention-2 (1,8,8192,64) 37 TFLOPS 100 TFLOPS 145 TFLOPS
share-qkv+stage2 (1,8,8192,64) 55 TFLOPS 99 TFLOPS 221 TFLOPS
FlashAttention-2 (1,48,8192,64) 37 TFLOPS 109 TFLOPS 163 TFLOPS
share-qkv+stage2 (1,48,8192,64) 48 TFLOPS 107 TFLOPS 224 TFLOPS
SDPA(高效注意力) (1,48,8192,512) 16 TFLOPS 58 TFLOPS 85 TFLOPS
🤖ffpa-attn (1,48,8192,512) 39 TFLOPS 104 TFLOPS 200 TFLOPS
精度误差 vs FA2/SDPA / 最大:< ~1e-3 最小:~0.0 平均:< ~1e-5

flash-attn⚡️⚡️ 中已经实现了 Split KVSplit Q 的实现,用于性能对比。其中,将所有 QKV 在 MMA(线程束)之间拆分的 Split KV 方法,比仅将 Q 在 MMA(线程束)之间拆分而保留所有 MMA(线程束)对 KV 的访问权限的 Split Q 方法要慢。

  • 📚 Split KV(基础版,FlashAttention-1)
// 将 QKV 在 MMA(线程束)之间拆分,采用朴素的矩阵乘法 MMA&线程束分块策略。
// 案例:8 个 MMA(2x4)的布局 [之后] kWarpTileSeqLenQxkWarpTileSeqLenK(2x2)-> 32x2,32x2=64x64:
// |  [64,64]  |    warp_KV 0    |    warp_KV 1    |    warp_KV 2    |    warp_KV 3    |
// | warp_QP 0 |-- MMA 0,MMA 0 --|-- MMA 2,MMA 2 --|-- MMA 4,MMA 4 --|-- MMA 6,MMA 6 --|
// | warp_QP 0 |-- MMA 0,MMA 0 --|-- MMA 2,MMA 2 --|-- MMA 4,MMA 4 --|-- MMA 6,MMA 6 --|
// | warp_QP 1 |-- MMA 1,MMA 1 --|-- MMA 3,MMA 2 --|-- MMA 5,MMA 5 --|-- MMA 7,MMA 7 --|
// | warp_QP 1 |-- MMA 1,MMA 1 --|-- MMA 3,MMA 2 --|-- MMA 5,MMA 5 --|-- MMA 7,MMA 7 --|
__global__ void // Q, K, V, O -> [B, H, N, D]
flash_attn_mma_stages_split_kv_kernel(half* Q, half* K, half* V, half* O, ...);
  • 📚 Split Q(更快,FlashAttention-2)
// 将 Q 在 MMA(线程束)之间拆分,并保持所有 MMA(线程束)对 KV 的访问权限,
// 以减少通过 SMEM 和线程束交换带来的通信开销。
// 案例:MMA = m16n8k16,Br=16x4=64,Bc=8x8=64,布局:4 个线程束
// |   64x64   |      warp_KV 0       |
// | warp_QP 0 | MMA 0 ... MMA 0 (x8) |
// | warp_QP 1 | MMA 1 ... MMA 1 (x8) |
// | warp_QP 2 | MMA 2 ... MMA 2 (x8) |
// | warp_QP 3 | MMA 3 ... MMA 3 (x8) |
__global__ void // Q, K, V, O -> [B, H, N, D]
flash_attn_mma_stages_split_q_kernel(half* Q, half* K, half* V, half* O, ...);
  • 📚 Split Q + 共享 KV SMEM (1/2 SRAM vs FA2)
// K 和 V 共享同一片共享内存,从而提高块占用率。
__global__ void // Q, K, V, O -> [B, H, N, D]
flash_attn_mma_stages_split_q_shared_kv_kernel(half* Q, half* K, half* V, half* O, ...);
  • 📚 Split Q + 完全共享 QKV SMEM (1/4 SRAM vs FA2)
// Q、K 和 V 完全共享同一片共享内存,并预取 Q s2r,从而提高块占用率
// 同时减少 Q 对 SMEM 的访问次数。
__global__ void // Q, K, V, O -> [B, H, N, D]
flash_attn_mma_stages_split_q_shared_qkv_kernel(half* Q, half* K, half* V, half* O, ...);
  • 📚 Split Q + QK 精细粒度分块(O(16xd) SRAM vs FA2 O(4xBrxd) SRAMHeaddim -> 1024
// 在 MMA 层面对 Q@K^T 进行精细粒度分块,使得 Q 和 K 的 SRAM 使用量恒定为
// 64 * kMmaAtomK。对于 V,SRAM 复杂度为 O(kMmaAtomK * d),因此整体 SRAM 复杂度为 O(kMmaAtomK * d)。由此,这种方法允许我们将 D(头维度)扩展到 1024。
__global__ void // Q, K, V, O -> [B, H, N, D]
flash_attn_mma_stages_split_q_tiling_qk_kernel(half* Q, half* K, half* V, half* O, ...);
  • 📚 Split Q + 完全 QKV 精细粒度分块(O(2xBrx16)~O(1) SRAM vs FA2 O(4xBrxd) SRAM
// 在 MMA 层面对所有的 Q@K^T 和 P@V 进行精细粒度分块,使得 Q、K 和 V 的 SRAM 使用量恒定为
// Br * 16 或 Bc * 16,整体 SRAM 复杂度为 O(Br * 16)。因此,这种方法让我们能够以比 SDPA 更快的速度运行,无论是否使用 MMA F32 精度。
__global__ void // Q, K, V, O -> [B, H, N, D]
flash_attn_mma_stages_split_q_tiling_qkv_kernel(half* Q, half* K, half* V, half* O, ...);

💡注意:📚 拆分 Q + 完全 QKV 精细粒度分块 已经被重构到 🤖ffpa-attn 中。

📖 200+ CUDA 核函数 🔥🔥(简单 -> 困难++)(©️返回👆🏻

此处列出的核函数将引导你逐步学习,从简单到非常具有挑战性的主题。每个主题的工作流程如下:自定义CUDA 核函数实现 -> PyTorch Python 绑定 -> 运行测试。👉提示:* = Tensor Core(WMMA、MMA、CuTe),否则为 CUDA Core;/ = 不支持;✔️ = 支持; = 待办。内容列表如下:

📚 简单📚 中等 部分涵盖了诸如 逐元素运算、矩阵转置、warp/block 归约、NMS、ReLU、GELU、Swish、层归一化、RMS 归一化、在线 Softmax、点积、嵌入 等操作,以及 FP32FP16BF16FP8 的基本用法。📚 困难📚 困难+📚 困难++ 部分则深入探讨更高级的主题,主要聚焦于 SGEMV、SGEMM、Hgemv、Hgemm 和 Flash Attention 等操作。这些部分还提供了大量使用 Tensor Core 和纯 MMA PTX 实现的核函数。

📚 简单 ⭐️ & 中等 ⭐️⭐️ (©️返回👆🏻)

📖 CUDA 内核 📖 元素数据类型 📖 累加器数据类型 📖 文档 📖 级别
✔️ elementwise_f32 f32 / link ⭐️
✔️ elementwise_f32x4 f32 / link ⭐️
✔️ elementwise_f16 f16 / link ⭐️
✔️ elementwise_f16x2 f16 / link ⭐️
✔️ elementwise_f16x8 f16 / link ⭐️
✔️ elementwise_f16x8_pack f16 / link ⭐️⭐️
✔️ histogram_i32 i32 / link ⭐️
✔️ histogram_i32x4 i32 / link ⭐️
✔️ sigmoid_f32 f32 / link ⭐️
✔️ sigmoid_f32x4 f32 / link ⭐️
✔️ sigmoid_f16 16 / link ⭐️
✔️ sigmoid_f16x2 f16 / link ⭐️
✔️ sigmoid_f16x8 f16 / link ⭐️
✔️ sigmoid_f16x8_pack f16 / link ⭐️⭐️
✔️ relu_f32 f32 / link ⭐️
✔️ relu_f32x4 f32 / link ⭐️
✔️ relu_f16 f16 / link ⭐️
✔️ relu_f16x2 f16 / link ⭐️
✔️ relu_f16x8 f16 / link ⭐️
✔️ relu_f16x8_pack f16 / link ⭐️⭐️
✔️ elu_f32 f32 / link ⭐️
✔️ elu_f32x4 f32 / link ⭐️
✔️ elu_f16 f16 / link ⭐️
✔️ elu_f16x2 f16 / link ⭐️
✔️ elu_f16x8 f16 / link ⭐️
✔️ elu_f16x8_pack f16 / link ⭐️⭐️
✔️ gelu_f32 f32 / link ⭐️
✔️ gelu_f32x4 f32 / link ⭐️
✔️ gelu_f16 f16 / link ⭐️
✔️ gelu_f16x2 f16 / link ⭐️
✔️ gelu_f16x8 f16 / link ⭐️
✔️ gelu_f16x8_pack f16 / link ⭐️⭐️
✔️ swish_f32 f32 / link ⭐️
✔️ swish_f32x4 f32 / link ⭐️
✔️ swish_f16 f16 / link ⭐️
✔️ swish_f16x2 f16 / link ⭐️
✔️ swish_f16x8 f16 / link ⭐️
✔️ swish_f16x8_pack f16 / link ⭐️⭐️
✔️ hardswish_f32 f32 / link ⭐️
✔️ hardswish_f32x4 f32 / link ⭐️
✔️ hardswish_f16 f16 / link ⭐️
✔️ hardswish_f16x2 f16 / link ⭐️
✔️ hardswish_f16x8 f16 / link ⭐️
✔️ hardswish_f16x8_pack f16 / link ⭐️⭐️
✔️ hardshrink_f32 f32 / link ⭐️
✔️ hardshrink_f32x4 f32 / link ⭐️
✔️ hardshrink_f16 f16 / link ⭐️
✔️ hardshrink_f16x2 f16 / link ⭐️
✔️ hardshrink_f16x8 f16 / link ⭐️
✔️ hardshrink_f16x8_pack f16 / link ⭐️⭐️
✔️ embedding_f32 f32 / link ⭐️
✔️ embedding_f32x4 f32 / link ⭐️
✔️ embedding_f32x4_pack f32 / link ⭐️
✔️ embedding_f16 f16 / link ⭐️
✔️ embedding_f16x2 f16 / link ⭐️
✔️ embedding_f16x8 f16 / link ⭐️
✔️ embedding_f16x8_pack f16 / link ⭐️⭐️
✔️ mat_trans_f32_col2row{2d} f32 / link ⭐️
✔️ mat_trans_f32_row2col{2d} f32 / link ⭐️
✔️ mat_trans_f32_diagonal2d f32 / link ⭐️⭐️
✔️ mat_trans_f32x4_col2row{2d} f32 / link ⭐️⭐️
✔️ mat_trans_f32x4_row2col{2d} f32 / link ⭐️⭐️
✔️ mat_trans_cute f32 / link ⭐️⭐️
✔️ warp_reduce_{all} all all link ⭐️⭐️
✔️ block_all_reduce_f32_f32 f32 f32 link ⭐️⭐️
✔️ block_all_reduce_f32x4_f32 f32 f32 link ⭐️⭐️
✔️ block_all_reduce_f16_f16 f16 f16 link ⭐️⭐️
✔️ block_all_reduce_f16_f32 f16 f32 link ⭐️⭐️
✔️ block_all_reduce_f16x2_f16 f16 f16 link ⭐️⭐️
✔️ block_all_reduce_f16x2_f32 f16 f32 link ⭐️⭐️
✔️ block_all_reduce_f16x8_pack_f16 f16 f16 link ⭐️⭐️
✔️ block_all_reduce_f16x8_pack_f32 f16 f32 link ⭐️⭐️
✔️ block_all_reduce_bf16_bf16 bf16 bf16 link ⭐️⭐️
✔️ block_all_reduce_bf16_f32 bf16 f32 link ⭐️⭐️
✔️ block_all_reduce_bf16x2_bf16 bf16 bf16 link ⭐️⭐️
✔️ block_all_reduce_bf16x2_f32 bf16 f32 link ⭐️⭐️
✔️ block_all_reduce_bf16x8_pack_bf16 bf16 bf16 link ⭐️⭐️
✔️ block_all_reduce_bf16x8_pack_f32 bf16 f32 link ⭐️⭐️
✔️ block_all_reduce_fp8_e4m3_f16 fp8_e4m3 f16 link ⭐️⭐️⭐️
✔️ block_all_reduce_fp8_e5m2_f16 fp8_e5m2 f16 link ⭐️⭐️⭐️
✔️ block_all_reduce_fp8_e4m3x16_pack_f16 fp8_e4m3 f16 link ⭐️⭐️⭐️
✔️ block_all_reduce_fp8_e5m2x16_pack_f16 fp8_e5m2 f16 link ⭐️⭐️⭐️
✔️ block_all_reduce_i8_i32 i8 i32 link ⭐️⭐️
✔️ block_all_reduce_i8x16_pack_i32 i8 i32 link ⭐️⭐️
✔️ dot_product_f32 f32 f32 link ⭐️⭐️
✔️ dot_product_f32x4 f32 f32 link ⭐️⭐️
✔️ dot_product_f16_f32 f16 f32 link ⭐️⭐️
✔️ dot_product_f16x2_f32 f16 f32 link ⭐️⭐️
✔️ dot_product_f16x8_pack_f32 f16 f32 link ⭐️⭐️
✔️ softmax_f32_per_tok f32 f32 link ⭐️⭐️
✔️ softmax_f32x4_per_tok f32 f32 link ⭐️⭐️
✔️ safe_softmax_f32_per_tok f32 f32 link ⭐️⭐️
✔️ safe_softmax_f32x4_per_tok f32 f32 link ⭐️⭐️
✔️ safe_softmax_f16_f32_per_tok f16 f32 link ⭐️⭐️
✔️ safe_softmax_f16x2_f32_per_tok f16 f32 link ⭐️⭐️
✔️ safe_softmax_f16x8_pack_f32_per_tok f16 f32 link ⭐️⭐️
✔️ online_safe_softmax_f32_per_token f32 f32 link ⭐️⭐️
✔️ online_safe_softmax_f32x4_pack_per_tok f32 f32 link ⭐️⭐️
✔️ rope_f32 f32 f32 link ⭐️⭐️
✔️ rope_f32x4_pack f32 f32 link ⭐️⭐️
✔️ layer_norm_f32 f32 f32 link ⭐️⭐️
✔️ layer_norm_f32x4 f32 f32 link ⭐️⭐️
✔️ layer_norm_f16_f16 f16 f16 link ⭐️⭐️
✔️ layer_norm_f16x2_f16 f16 f16 link ⭐️⭐️
✔️ layer_norm_f16x8_f16 f16 f16 link ⭐️⭐️
✔️ layer_norm_f16x8_pack_f16 f16 f16 link ⭐️⭐️
✔️ layer_norm_f16x8_pack_f32 f16 f32 link ⭐️⭐️
✔️ layer_norm_f16_f32 f16 f32 link ⭐️⭐️
✔️ rms_norm_f32 f32 f32 link ⭐️⭐️
✔️ rms_norm_f32x4 f32 f32 link ⭐️⭐️
✔️ rms_norm_f16_f16 f16 f16 link ⭐️⭐️
✔️ rms_norm_f16x2_f16 f16 f16 link ⭐️⭐️
✔️ rms_norm_f16x8_f16 f16 f16 link ⭐️⭐️
✔️ rms_norm_f16x8_f32 f16 f32 link ⭐️⭐️
✔️ rms_norm_f16x8_pack_f16 f16 f16 link ⭐️⭐️
✔️ rms_norm_f16x8_pack_f32 f16 f32 link ⭐️⭐️
✔️ rms_norm_f16_f32 f16 f32 link ⭐️⭐️
✔️ nms_f32 f32 / link ⭐️⭐️
✔️ merge_attn_states f16/bf16/f32 f32 link ⭐️⭐️
✔️ notes v1(deprecated) f32 f32 / ⭐️⭐️
✔️ How to use nsys/ncu(timeline/ptx/sass) / / link ⭐️⭐️

📚 困难 ⭐⭐⭐️(©️返回👆🏻

📖 CUDA 内核 📖 元素数据类型 📖 累加器数据类型 📖 文档 📖 难度
✔️ sgemv_k32_f32 f32 f32 链接 ⭐️⭐️⭐️
✔️ sgemv_k128_f32x4 f32 f32 链接 ⭐️⭐️⭐️
✔️ sgemv_k16_f32 f32 f32 链接 ⭐️⭐️⭐️
✔️ hgemv_k32_f16 f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemv_k128_f16x4 f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemv_k16_f16 f16 f16 链接 ⭐️⭐️⭐️
✔️ sgemm_naive_f32 f32 f32 链接 ⭐️⭐️
✔️ sgemm_sliced_k_f32 f32 f32 链接 ⭐️⭐️⭐️
✔️ sgemm_t_8x8_sliced_k_f32x4 f32 f32 链接 ⭐️⭐️⭐️
✔️ sgemm_t_8x8_sliced_k...bcf f32 f32 链接 ⭐️⭐️⭐️
✔️ sgemm_t_8x8_sliced_k...dbuf f32 f32 链接 ⭐️⭐️⭐️
✔️ sgemm_t_8x8_sliced_k16...dbuf f32 f32 链接 ⭐️⭐️⭐️
✔️ sgemm_t_8x8_sliced_k16...async f32 f32 链接 ⭐️⭐️⭐️
✔️ sgemm_wmma_m16n16k8...stages* tf32 f32 链接 ⭐️⭐️⭐️
✔️ sgemm_wmma_m16n16k8...swizzle* tf32 f32 链接 ⭐️⭐️⭐️
✔️ hgemm_naive_f16 f16 f16 链接 ⭐️⭐️
✔️ hgemm_sliced_k_f16 f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_t_8x8_sliced_k_f16x4 f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_t_8x8_sliced_k_f16x4_pack f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_t_8x8_sliced_k_f16x8_pack f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_t_8x8_sliced_k...dbuf f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_t_8/16x8...k16/32...dbuf f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_t_8/16x8...k16/32...async f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_wmma_m16n16k16...naive* f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_wmma_m16n16k16...mma4x2* f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_wmma_m16n16k16...mma4x4* f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_wmma_m16n16k16...dbuf* f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_wmma_m32n8k16....dbuf* f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_wmma_m16n16k16...stages* f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_wmma_m16n16k16...swizzle* f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_mma_m16n8k16...naive* f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_mma_m16n8k16...mma2x4* f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_mma_m16n8k16...stages* f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_mma_m16n8k16...swizzle* f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_mma_m16n8k16...swizzle{smem}* f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_mma_m16n8k16...swizzle{tn}{smem}* f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_mma_stages_swizzle{smem}...cute* f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_mma_cublas* f16 f16 链接 ⭐️⭐️
✔️ hgemm_wgmma_m64n128k16...tma{ws}{tn}* f16 f16 链接 ⭐️⭐️⭐️
✔️ hgemm_wgmma_m64n128k16_fp32...tma* f16 f32 链接 ⭐️⭐️⭐️

📚 更难 ⭐️⭐️⭐️⭐️ 和 非常难 ⭐️⭐️⭐️⭐️⭐️(©️返回👆🏻

  • 📚 FlashAttention-2 MMA(MMA 累加器 F32/F16,swizzle,QKV 共享片上存储,细粒度分块等🎉)
📖 CUDA 内核 📖 元素数据类型 📖 累加数据类型 📖 文档 📖 难度
✔️ flash_attn_cute(朴素) f16 f32 链接 ⭐️⭐️⭐️
✔️ 如何实现 MMA 共享内存置换* f16 f16 链接 ⭐️⭐️⭐️
✔️ flash_attn_mma 阶段分割 KV* f16 f16 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma 阶段分割 Q* f16 f16 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma 阶段…共享 KV* f16 f16 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma 阶段…共享 QKV* f16 f16 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma 阶段…分块 QK* f16 f16 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma 阶段…分块 QKV* f16 f16 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma 阶段…共享 KV{f32}* f16 f32 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma 阶段…共享 QKV{f32}* f16 f32 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma 阶段…分块 QK{f32}* f16 f32 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma 阶段…分块 QKV{f32}* f16 f32 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma…共享 KV{f32}{rr}* f16 f32 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma…共享 QKV{f32}{rr}* f16 f32 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma…共享 KV 置换{q}* f16 f16 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma…共享 KV 置换{qk}* f16 f16 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma…共享 KV 置换{qkv}* f16 f16 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma…共享 QKV 置换{q}* f16 f16 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma…共享 QKV 置换{qk}* f16 f16 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma…共享 QKV 置换{qkv}* f16 f16 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma…分块 QK 置换{q}* f16 f16 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma…分块 QK 置换{qk}* f16 f16 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma…分块 QK 置换{qkv}* f16 f16 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma…分块 QKV 置换{q}* f16 f16 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn_mma…分块 QKV 置换{qk}* f16 f16 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn…分块 QKV 置换{qkv}{f32}* f16 f32 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn…分块 QKV 置换{qk}{f32}* f16 f32 链接 ⭐️⭐️⭐️⭐️
✔️ flash_attn…分块 QKV 置换{qkv}{f32}* f16 f32 链接 ⭐️⭐️⭐️⭐️

💡注:rr 表示减少寄存器使用(适用于 d>128);f32 表示 MMA 累加时使用 FP32 数据类型,否则为 FP16。softmax 的累加数据类型始终为 FP32 以保证高精度;swizzle 目前仅支持 MMA 的共享内存置换。

  • 📚 FFPA 注意力 MMA(比 SDPA EA 快 1.8x~3x🎉,D > 256,不支持 FA2)
📖 CUDA内核 📖 元素数据类型 📖 累加器数据类型 📖 文档 📖 难度
✔️ ffpa_mma_stages_split_q_L1_F16F16F16 f16 f16 link ⭐️⭐️⭐️⭐️
✔️ ffpa_mma_stages_split_q_L1_F16F16F32 f16 f32 link ⭐️⭐️⭐️⭐️
✔️ ffpa_mma_stages_split_q_L1_mixed_acc f16 QK f32, PV f16 link ⭐️⭐️⭐️⭐️
⚠️ ffpa_mma_stages_split_q_L2_F16F16F16 f16 f16 link ⭐️⭐️⭐️⭐️
⚠️ ffpa_mma_stages_split_q_L2_F16F16F32 f16 f32 link ⭐️⭐️⭐️⭐️
⚠️ ffpa_mma_stages_split_q_L2_mixed_acc f16 QK f32, PV f16 link ⭐️⭐️⭐️⭐️
⚠️ ffpa_mma_stages_split_q_L3_F16F16F16 f16 f16 link ⭐️⭐️⭐️⭐️
⚠️ ffpa_mma_stages_split_q_L3_F16F16F32 f16 f32 link ⭐️⭐️⭐️⭐️
⚠️ ffpa_mma_stages_split_q_L3_mixed_acc f16 QK f32, PV f16 link ⭐️⭐️⭐️⭐️

💡注: 🤖ffpa-attn:📚FFPA——另一种更快的Flash Prefill注意力机制,针对head dim > 256时具有O(1)🎉SRAM复杂度,比SDPA EA快1.8倍~3倍🎉:📈L20 ~1.9倍↑🎉📈 A30 ~1.8倍↑🎉📈3080 ~2.9倍↑🎉📈4090 ~2.1倍↑🎉

📚 Triton内核(OpenAI Triton)⭐️⭐️⭐️ (©️返回👆🏻)

📖 Triton内核 📖 元素数据类型 📖 累加器数据类型 📖 文档 📖 难度
✔️ triton_vector_add_kernel all all link ⭐️⭐️
✔️ triton_fused_softmax(多阶段) f16/bf16/f32 f32 link ⭐️⭐️⭐️
✔️ triton_fused_layer_norm(前向传播) f16/bf16/f32 f32 link ⭐️⭐️⭐️
✔️ triton_fused_layer_norm(反向传播) f16/bf16/f32 f32 link ⭐️⭐️⭐️
✔️ triton_merge_attn_states_kernel(配合CUDA) f16/bf16/f32 f32 link ⭐️⭐️⭐️

📚 CUTLASS/CuTe内核 ⭐️⭐️⭐️ (©️返回👆🏻)

📖 CUTLASS/CuTe内核 📖 元素数据类型 📖 累加器数据类型 📖 文档 📖 难度
✔️ mat_transpose_cute f32 / link ⭐️⭐️
✔️ flash_attn_cute(朴素版) f16 f32 link ⭐️⭐️⭐️
✔️ hgemv_f16_cute_kernel f16 f16 link ⭐️⭐️⭐️
✔️ hgemv_f16x8_cute_kernel f16 f16 link ⭐️⭐️⭐️
✔️ hgemv_tensor_core_cute_kernel f16 f16 link ⭐️⭐️⭐️
✔️ hgemm_mma_stages_swizzle{smem}...cute* f16 f16 link ⭐️⭐️⭐️
✔️ ws_hgemm_naive_cute_kernel f16 f16 link ⭐️⭐️⭐️

📖 100+ 高性能计算与分布式-技术博客

📚 高性能计算与分布式-个人技术专栏 (©️返回👆🏻)

📖 类型-标题 📖 作者 📖 推荐
[扩散推理]📖简短的2025年总结,写在Cache-DiT v1.2.1之际 @DefTruth ⭐️⭐️
[扩散推理]📖CacheDiT支持Z-Image分布式推理和缓存加速​​ @DefTruth ⭐️⭐️
[扩散推理]📖cache-dit支持FLUX.2分布式推理和Cache @DefTruth ⭐️⭐️
[扩散推理]📖Cache加速-FoCa公式理解记录 @DefTruth ⭐️⭐️⭐
[扩散推理]📖cache-dit: BlockAdapter支持HunyuanImage-2.1 Cache加速! @DefTruth ⭐️⭐️⭐
[扩散推理]📖cache-dit + Qwen-Image-Lightning 实现 3.5 steps 推理! @DefTruth ⭐️⭐️⭐
[扩散推理]📖cache-dit: Wan2.2-MoE 2.4x 推理加速! @DefTruth ⭐️⭐️⭐
[扩散推理]📖cache-dit: Qwen-Image-Edit 2x 无损加速! @DefTruth ⭐️⭐️⭐
[扩散推理]📖cache-dit: Qwen-Image 1.5x 无损加速! @DefTruth ⭐️⭐️⭐
[扩散推理]📖Cache加速-TaylorSeer算法简析 @DefTruth ⭐️⭐️⭐
[扩散推理]📖DiT推理加速综述: Caching @DefTruth ⭐️⭐️⭐
[Triton编程][基础]📖Triton极简入门: Triton Vector Add @DefTruth ⭐️⭐️⭐
[Triton编程][基础]📖Triton Fused Softmax Kernel详解: 从Python源码到PTX @DefTruth ⭐️⭐️⭐
[Triton编程][基础]📖vLLM Triton Merge Attention States Kernel详解 @DefTruth ⭐️⭐️⭐
[Triton编程][进阶]📖vLLM Prefix Prefill Triton Kernel图解 @DefTruth ⭐️⭐️⭐️
[张量/序列并行]📖序列并行: BPT、Ring-Attention及Striped-Attention笔记 @DefTruth ⭐️⭐️⭐
[vLLM实践][算子]📖vLLM算子开发流程:”保姆级“详细记录 @DefTruth ⭐️⭐️⭐
[vLLM实践][万字]📖vLLM + DeepSeek-R1 671B 多机部署及修Bug笔记 @DefTruth ⭐️⭐️⭐
[Attention优化]📖FFPA(Split-D): FA2无限HeadDim扩展,2x↑🎉 vs SDPA EA @DefTruth ⭐️⭐️⭐️
[CUDA基础][开篇]📖LeetCUDA: v3.0 大升级-面试刷题不迷路 @DefTruth ⭐️⭐️⭐⭐️
[分布式训推][张量/序列并行]📖图解DeepSpeed-Ulysses&Megatron-LM TP/SP @DefTruth ⭐️⭐️
[VLM推理优化][InternVL系列]📖InternLM2/.../InternVL1.5系列笔记: 核心点解析 @DefTruth ⭐️⭐️
[LLM推理优化][TensorRT-LLM][5w字]📖TensorRT-LLM部署调优-指北 @DefTruth ⭐️⭐️⭐️
[LLM推理优化][KV Cache优化]📖GQA/YOCO/CLA/MLKV: 层内和层间KV Cache共享 @DefTruth ⭐️⭐️
[LLM推理优化][Prefill优化][万字]📖图解vLLM Automatic Prefix Caching: TTFT优化 @DefTruth ⭐️⭐️⭐️
[LLM推理优化][Attention优化]📖图解:从Online-Softmax到FlashAttention V1/V2/V3 @DefTruth ⭐️⭐️⭐️
[LLM推理优化][Decoding优化]📖原理&图解FlashDecoding/FlashDecoding++ @DefTruth ⭐️⭐️
[VLM推理优化][LLaVA系列]📖CLIP/LLaVA/LLaVA1.5/VILA笔记: 核心点解析 @DefTruth ⭐️⭐️
[LLM推理优化][Attention优化][万字]📖TensorRT MHA/Myelin vs FlashAttention-2 @DefTruth ⭐️⭐️⭐️
[LLM推理优化][PTX汇编]📖CUDA 12 PTX汇编: PRMT指令详解-通用模式 @DefTruth ⭐️
[LLM推理优化][PTX汇编]📖CUDA 12 PTX汇编: LOP3指令详解 @DefTruth ⭐️
[LLM推理优化][CUDA][3w字]📖高频面试题汇总-大模型手撕CUDA @DefTruth ⭐️⭐️⭐️
[LLM推理优化][Weight Only]📖WINT8/4-(00): 通俗易懂讲解-快速反量化算法 @DefTruth ⭐️⭐️
[LLM推理优化][Weight Only]📖WINT8/4-(01): PRMT指令详解及FT源码解析 @DefTruth ⭐️⭐️
[LLM推理优化][Weight Only]📖WINT8/4-(02): 快速反量化之INT8转BF16 @DefTruth ⭐️⭐️
[LLM推理优化][Weight Only]📖WINT8/4-(03): LOP3指令详解及INT4转FP16/BF16 @DefTruth ⭐️⭐️
[LLM推理优化][LLM Infra整理]📖100+篇: 大模型推理各方向新发展整理 @DefTruth ⭐️⭐️
[LLM推理优化][LLM Infra整理]📖30+篇: LLM推理论文集-500页PDF @DefTruth ⭐️⭐️
[LLM推理优化][LLM Infra整理]📖FlashDecoding++: 比FlashDecoding还要快! @DefTruth ⭐️
[LLM推理优化][LLM Infra整理]📖TensorRT-LLM开源,TensorRT 9.1也来了 @DefTruth ⭐️
[LLM推理优化][LLM Infra整理]📖20+篇: LLM推理论文集-300页PDF @DefTruth ⭐️⭐️
[LLM推理优化][LLM Infra整理]📖PagedAttention论文新鲜出炉 @DefTruth ⭐️
[推理部署][CV/NLP]📖FastDeploy三行代码搞定150+ CV、NLP模型部署 @DefTruth ⭐️
[推理部署][CV]📖如何在lite.ai.toolkit(3.6k+ stars)中增加您的模型? @DefTruth ⭐️⭐️
[推理部署][CV]📖美团 YOLOv6 ORT/MNN/TNN/NCNN C++推理部署 @DefTruth ⭐️⭐️
[推理部署][ONNX]📖ONNX推理加速技术文档-杂记 @DefTruth ⭐️
[推理部署][TensorFlow]📖Mac源码编译TensorFlow C++指北 @DefTruth ⭐️
[推理部署][CV]📖1Mb!头部姿态估计: FSANet,一个小而美的模型(C++) @DefTruth ⭐️
[推理部署][CV]📖opencv+ffmpeg编译打包全解指南 @DefTruth ⭐️⭐️
[推理部署][CV]📖RobustVideoMatting视频抠图静态ONNX模型转换 @DefTruth ⭐️
[推理部署][CV]📖190Kb!SSRNet年龄检测详细解读(含C++工程) @DefTruth ⭐️
[推理部署][CV]📖MGMatting(CVPR2021)人像抠图C++应用记录 @DefTruth ⭐️
[推理部署][CV]📖超准确人脸检测(带关键点)YOLO5Face C++工程详细记录 @DefTruth ⭐️⭐️
[推理部署][ORT]📖解决: ONNXRuntime(Python) GPU 部署配置记录 @DefTruth ⭐️
[推理部署][CV]📖记录SCRFD(CVPR2021)人脸检测C++工程化(含docker镜像) @DefTruth ⭐️⭐️
[推理部署][NCNN]📖野路子:记录一个解决onnx转ncnn时op不支持的trick @DefTruth ⭐️
[推理部署][CV]📖升级版NanoDet-Plus MNN/TNN/NCNN/ORT C++工程记录 @DefTruth ⭐️⭐️
[推理部署][CV]📖超轻量级NanoDet MNN/TNN/NCNN/ORT C++工程记录 @DefTruth ⭐️
[推理部署][CV]📖详细记录MGMatting之MNN、TNN和ORT C++移植 @DefTruth ⭐️⭐️
[推理部署][CV]📖YOLOX NCNN/MNN/TNN/ONNXRuntime C++工程简记 @DefTruth ⭐️
[推理部署][TNN]📖手动修改YoloX的tnnproto记录-TNN @DefTruth ⭐️
[推理部署][ORT]📖全网最详细 ONNXRuntime C++/Java/Python 资料! @DefTruth ⭐️
[推理部署][CV]📖RobustVideoMatting: C++工程化记录-实现篇 @DefTruth ⭐️⭐️
[推理部署][CV]📖RobustVideoMatting: C++工程化记录-应用篇 @DefTruth ⭐️⭐️
[推理部署][ORT]📖ONNXRuntime C++ CMake 工程分析及编译 @DefTruth ⭐️⭐️
[推理部署][ORT]📖如何使用ORT C++ API处理NCHW和NHWC输入? @DefTruth ⭐️
[推理部署][TNN]📖tnn-convert搭建简记-YOLOP转TNN @DefTruth ⭐️
[推理部署][CV]📖YOLOP ONNXRuntime C++工程化记录 @DefTruth ⭐️⭐️
[推理部署][NCNN]📖超有用NCNN参考资料整理 @DefTruth ⭐️
[推理部署][MNN]📖超有用MNN参考资料整理 @DefTruth ⭐️
[推理部署][TNN]📖超有用TNN参考资料整理 @DefTruth ⭐️
[推理部署][ONNX]📖超有用ONNX参考资料整理 @DefTruth ⭐️
[推理部署][ONNX]📖超有用ONNX模型结构参考资料整理 @DefTruth ⭐️
[推理部署][OpenCV-DNN]📖超有用OpenCV-DNN参考资料整理 @DefTruth ⭐️
[推理部署][Tensorflow]📖超有用Tensorflow C++工程化知识点 @DefTruth ⭐️
[推理部署][模型转换]📖深度学习模型转换资料整理 @DefTruth ⭐️
[技术随笔][C++][CMake]📖超有用CMake参考资料整理 @DefTruth ⭐️⭐️
[技术随笔][C++] [3W字]📖静态链接和静态库实践指北-原理篇 @DefTruth ⭐️⭐️⭐️
[技术随笔][C++]📖Mac下C++内存检查指北(Valgrind VS Asan) @DefTruth ⭐️
[技术随笔][CV]📖torchlm: 人脸关键点检测库 @DefTruth ⭐️⭐️
[技术随笔][ML]📖《统计学习方法-李航: 笔记-从原理到实现-基于R》 @DefTruth ⭐️⭐️
[技术随笔][Git]📖如何优雅地git clone和git submodule? @DefTruth ⭐️
[技术随笔][3D]📖人脸重建3D参考资料整理 @DefTruth ⭐️
[技术随笔][3D]📖BlendShapes参考资料整理 @DefTruth ⭐️
[技术随笔][3D]📖从源码安装Pytorch3D详细记录及学习资料 @DefTruth ⭐️
[技术随笔][ML]📖200页:《统计学习方法:李航》笔记 -从原理到实现 @DefTruth ⭐️⭐️

📚 高性能计算与分布式——技术博客推荐 (©️返回👆🏻)

💡说明: 本小节整理一些自己比较喜欢的文章。欢迎大家提PR推荐更多优秀的文章!

📖 类型-标题 📖 作者 📖 推荐
[cute系列详解][入门]📖cutlass cute 101 @朱小霖 ⭐️⭐️⭐️
[cute系列详解][入门]📖CUTLASS 2.x & CUTLASS 3.x Intro 学习笔记 @BBuf ⭐️⭐️⭐️
[cute系列详解][入门]📖写给大家看的 CuTe 教程:tiled copy @竹熙佳处 ⭐️⭐️⭐️
[cute系列详解][入门]📖写给大家看的 CuTe 教程:tiled mma @竹熙佳处 ⭐️⭐️⭐️
[cute系列详解][入门]📖写给大家看的 CuTe 教程:Layout Compose & Inverse @竹熙佳处 ⭐️⭐️⭐️
[cute系列详解][入门]📖写给大家看的 CuTe 教程: Layout Product & Divide @竹熙佳处 ⭐️⭐️⭐️
[cute系列详解][入门]📖写给大家看的 CuTe 教程:TMA Copy @竹熙佳处 ⭐️⭐️⭐️
[cute系列详解][入门]📖写给进阶开发的 CuTe 笔记:permutationMNK 参数 @竹熙佳处 ⭐️⭐️⭐️
[cute系列详解][Layout]📖cute 之 Layout @reed ⭐️⭐️⭐️
[cute系列详解][Layout]📖cute Layout 的代数和几何解释 @reed ⭐️⭐️⭐️
[cute系列详解][Tensor]📖cute 之 Tensor @reed ⭐️⭐️⭐️
[cute系列详解][MMA]📖cute 之 MMA抽象 @reed ⭐️⭐️⭐️
[cute系列详解][Copy]📖cute 之 Copy抽象 @reed ⭐️⭐️⭐️
[cute系列详解][Swizzle]📖cute 之 Swizzle @reed ⭐️⭐️⭐️
[cute系列详解][Swizzle]📖cute Swizzle细谈 @进击的Killua ⭐️⭐️⭐️
[cute系列详解][Swizzle]📖cutlass swizzle机制解析(一) @Titus ⭐️⭐️⭐️
[cute系列详解][Swizzle]📖cutlass swizzle机制解析(二) @Titus ⭐️⭐️⭐️
[cute系列详解][Swizzle]📖CUDA避免smem bank conflict的swizzle机制解析 @frankshi ⭐️⭐️⭐️
[cute系列详解][Swizzle]📖布局代数实战:Swizzle自动推导 @melonedo ⭐️⭐️⭐️
[cute系列详解][GEMM]📖cute 之 简单GEMM实现 @reed ⭐️⭐️⭐️
[cute系列详解][GEMM]📖cute 之 GEMM流水线 @reed ⭐️⭐️⭐️
[cute系列详解][GEMM]📖cute 之 高效GEMM实现 @reed ⭐️⭐️⭐️
[cute系列详解][GEMM]📖GEMM流水线: single/multi-stage、pipeline @Titus ⭐️⭐️⭐️
[cute系列详解][GEMM]📖GEMM细节分析(一): ldmatrix的选择 @Anonymous ⭐️⭐️⭐️
[cute系列详解][GEMM]📖GEMM细节分析(二): TiledCopy与cp.async @Anonymous ⭐️⭐️⭐️
[cute系列详解][GEMM]📖GEMM细节分析(三): Swizzle<B,M,S>参数取值 @Anonymous ⭐️⭐️⭐️
[cute系列详解][实践]📖Hopper Mixed GEMM的CUTLASS实现笔记 @BBuf ⭐️⭐️⭐️
[cute系列详解][实践]📖CUTLASS CuTe实战(一): 基础 @进击的Killua ⭐️⭐️⭐️
[cute系列详解][实践]📖CUTLASS CuTe实战(二): 应用 @进击的Killua ⭐️⭐️⭐️
[cute系列详解][实践]📖FlashAttention fp8实现(ada架构) @shengying.wei ⭐️⭐️⭐️
[cute系列详解][实践]📖FlashAttention 笔记: tiny-flash-attention解读 @shengying.wei ⭐️⭐️⭐️
[cute系列详解][实践]📖使用cutlass cute复现flash attention @66RING ⭐️⭐️⭐️
[cutlass教程][入门]📖cutlass 基本认知 @JoeNomad ⭐️⭐️⭐️
[cutlass教程][入门]📖cutlass 软件架构 @JoeNomad ⭐️⭐️⭐️
[cutlass教程][入门]📖CUTLASS 基础介绍 @进击的Killua ⭐️⭐️⭐️
[cutlass教程][入门]📖乱谈CUTLASS GTC2020 SLIDES @zzk again ⭐️⭐️⭐️
[cutlass教程][深入]📖cutlass block swizzle 和 tile iterator @JoeNomad ⭐️⭐️⭐️
[cutlass教程][深入]📖cutlass bank conflict free的smem layout @JoeNomad ⭐️⭐️⭐️
[cutlass教程][深入]📖cutlass 多级流水线 @JoeNomad ⭐️⭐️⭐️
[GPU指令集架构][精解]📖NVidia GPU指令集架构-前言 @reed ⭐️⭐️⭐️
[GPU指令集架构][精解]📖NVidia GPU指令集架构-寄存器 @reed ⭐️⭐️⭐️
[GPU指令集架构][精解]📖NVidia GPU指令集架构-Load和Cache @reed ⭐️⭐️⭐️
[GPU指令集架构][精解]📖NVidia GPU指令集架构-浮点运算 @reed ⭐️⭐️⭐️
[GPU指令集架构][精解]📖NVidia GPU指令集架构-整数运算 @reed ⭐️⭐️⭐️
[GPU指令集架构][精解]📖NVidia GPU指令集架构-比特和逻辑操作 @reed ⭐️⭐️⭐️
[GPU指令集架构][精解]📖NVidia GPU指令集架构-Warp级和Uniform操作 @reed ⭐️⭐️⭐️
[CUDA优化][入门]📖CUDA 入门的正确姿势:how-to-optimize-gemm @白牛 ⭐️⭐️⭐️
[CUDA优化][入门]📖CUDA(一):CUDA 编程基础 @紫气东来 ⭐️⭐️⭐️
[CUDA优化][入门]📖CUDA(二):GPU的内存体系及其优化指南 @紫气东来 ⭐️⭐️⭐️
[CUDA优化][实践]📖CUDA(三):通用矩阵乘法:从入门到熟练 @紫气东来 ⭐️⭐️⭐️
[CUDA优化][实践]📖ops(1):LayerNorm 算子的 CUDA 实现与优化 @紫气东来 ⭐️⭐️⭐️
[CUDA优化][实践]📖ops(2):SoftMax算子的 CUDA 实现 @紫气东来 ⭐️⭐️⭐️
[CUDA优化][实践]📖ops(3):Cross Entropy 的 CUDA 实现 @紫气东来 ⭐️⭐️⭐️
[CUDA优化][实践]📖ops(4):AdamW 优化器的 CUDA 实现 @紫气东来 ⭐️⭐️⭐️
[CUDA优化][实践]📖ops(5):激活函数与残差连接的 CUDA 实现 @紫气东来 ⭐️⭐️⭐️
[CUDA优化][实践]📖ops(6):embedding 层与 LM head 层的 CUDA 实现 @紫气东来 ⭐️⭐️⭐️
[CUDA优化][实践]📖ops(7):self-attention 的 CUDA 实现及优化 (上) @紫气东来 ⭐️⭐️⭐️
[CUDA优化][实践]📖ops(8):self-attention 的 CUDA 实现及优化 (下) @紫气东来 ⭐️⭐️⭐️
[CUDA优化][实践]📖CUDA(四):使用 CUDA 实现 Transformer 结构 @紫气东来 ⭐️⭐️⭐️
[CUDA优化][Copy]📖Async Copy及Memory Barrier指令的功能与实现 @Frank Wang ⭐️⭐️⭐️
[CUDA优化][GEMV]📖深入浅出GPU优化系列:gemv优化 @有了琦琦的棍子 ⭐️⭐️⭐️
[CUDA优化][实践]📖CUDA element-wise 算子详解 @懒蚂蚁呀不嘿 ⭐️⭐️⭐️
[CUDA优化][实践]📖CUDA transpose 算子详解 @懒蚂蚁呀不嘿 ⭐️⭐️⭐️
[CUDA优化][实践]📖CUDA reduce 算子详解 @懒蚂蚁呀不嘿 ⭐️⭐️⭐️
[CUDA优化][实践]📖CUDA GEMM 算子详解 @懒蚂蚁呀不嘿 ⭐️⭐️⭐️
[Tensor Cores]📖Nvidia Tensor Core初探 @木子知 ⭐️⭐️⭐️
[Tensor Cores]📖Nvidia Tensor Core-WMMA API编程入门 @木子知 ⭐️⭐️⭐️
[Tensor Cores]📖Nvidia Tensor Core-MMA PTX编程入门 @木子知 ⭐️⭐️⭐️
[Tensor Cores]📖CUDA Ampere Tensor Core HGEMM 矩阵乘法优化 @nicholaswilde ⭐️⭐️⭐️
[GPU通信架构][精解]📖NVIDIA GPGPU(四)- 通信架构 @Bruce ⭐️⭐️⭐️
[torch.compile][原理]📖Torch.compile流程解析: 介绍 @StarCap ⭐️⭐️⭐️
[torch.compile][原理]📖Torch.compile流程解析: TorchDynamo @StarCap ⭐️⭐️⭐️
[torch.compile][原理]📖Torch.compile流程解析: AOTAutograd @StarCap ⭐️⭐️⭐️
[torch.compile][原理]📖Torch.compile流程解析: TorchInductor @StarCap ⭐️⭐️⭐️
[torch.compile][原理]📖Torch.compile流程解析: 算子融合 @StarCap ⭐️⭐️⭐️
[torch.compile][实践]📖Torch.compile使用指南 @jhang ⭐️⭐️⭐️
[torch.compile][实践]📖Torch.compile详细示例解析教程 @Bbuf ⭐️⭐️⭐️
[torch.compile][原理]📖一文搞懂TorchDynamo原理 @吾乃阿尔法 ⭐️⭐️⭐️
[torch.compile][原理]📖理解torch.compile基本原理和使用方式 @俯仰 ⭐️⭐️⭐️

©️许可证 (©️返回👆🏻)

GNU 通用公共许可证 v3.0

🎉贡献 (©️返回👆🏻)

如何贡献?请给本仓库点个星,或查看 🌤🌤CONTRIBUTE🎉🎉

📖 参考文献 (©️返回👆🏻)

版本历史

v3.0.192026/04/05
v3.0.182026/03/20
v3.0.172026/02/13
v3.0.162025/10/17
v3.0.152025/09/15
v3.0.142025/08/01
v3.0.132025/06/29
v3.0.122025/06/17
v3.0.112025/06/11
v3.0.102025/06/03
v3.0.92025/05/12
v3.0.82025/05/06
v3.0.72025/04/28
v3.0.62025/04/26
v3.0.52025/04/09
v3.0.42025/03/15
v3.0.32025/03/04
v3.0.22025/02/24
v3.0.12025/02/06
v3.0.02025/01/22

常见问题

相似工具推荐

openclaw

OpenClaw 是一款专为个人打造的本地化 AI 助手,旨在让你在自己的设备上拥有完全可控的智能伙伴。它打破了传统 AI 助手局限于特定网页或应用的束缚,能够直接接入你日常使用的各类通讯渠道,包括微信、WhatsApp、Telegram、Discord、iMessage 等数十种平台。无论你在哪个聊天软件中发送消息,OpenClaw 都能即时响应,甚至支持在 macOS、iOS 和 Android 设备上进行语音交互,并提供实时的画布渲染功能供你操控。 这款工具主要解决了用户对数据隐私、响应速度以及“始终在线”体验的需求。通过将 AI 部署在本地,用户无需依赖云端服务即可享受快速、私密的智能辅助,真正实现了“你的数据,你做主”。其独特的技术亮点在于强大的网关架构,将控制平面与核心助手分离,确保跨平台通信的流畅性与扩展性。 OpenClaw 非常适合希望构建个性化工作流的技术爱好者、开发者,以及注重隐私保护且不愿被单一生态绑定的普通用户。只要具备基础的终端操作能力(支持 macOS、Linux 及 Windows WSL2),即可通过简单的命令行引导完成部署。如果你渴望拥有一个懂你

349.3k|★★★☆☆|1周前
Agent开发框架图像

stable-diffusion-webui

stable-diffusion-webui 是一个基于 Gradio 构建的网页版操作界面,旨在让用户能够轻松地在本地运行和使用强大的 Stable Diffusion 图像生成模型。它解决了原始模型依赖命令行、操作门槛高且功能分散的痛点,将复杂的 AI 绘图流程整合进一个直观易用的图形化平台。 无论是希望快速上手的普通创作者、需要精细控制画面细节的设计师,还是想要深入探索模型潜力的开发者与研究人员,都能从中获益。其核心亮点在于极高的功能丰富度:不仅支持文生图、图生图、局部重绘(Inpainting)和外绘(Outpainting)等基础模式,还独创了注意力机制调整、提示词矩阵、负向提示词以及“高清修复”等高级功能。此外,它内置了 GFPGAN 和 CodeFormer 等人脸修复工具,支持多种神经网络放大算法,并允许用户通过插件系统无限扩展能力。即使是显存有限的设备,stable-diffusion-webui 也提供了相应的优化选项,让高质量的 AI 艺术创作变得触手可及。

162.1k|★★★☆☆|1周前
开发框架图像Agent

everything-claude-code

everything-claude-code 是一套专为 AI 编程助手(如 Claude Code、Codex、Cursor 等)打造的高性能优化系统。它不仅仅是一组配置文件,而是一个经过长期实战打磨的完整框架,旨在解决 AI 代理在实际开发中面临的效率低下、记忆丢失、安全隐患及缺乏持续学习能力等核心痛点。 通过引入技能模块化、直觉增强、记忆持久化机制以及内置的安全扫描功能,everything-claude-code 能显著提升 AI 在复杂任务中的表现,帮助开发者构建更稳定、更智能的生产级 AI 代理。其独特的“研究优先”开发理念和针对 Token 消耗的优化策略,使得模型响应更快、成本更低,同时有效防御潜在的攻击向量。 这套工具特别适合软件开发者、AI 研究人员以及希望深度定制 AI 工作流的技术团队使用。无论您是在构建大型代码库,还是需要 AI 协助进行安全审计与自动化测试,everything-claude-code 都能提供强大的底层支持。作为一个曾荣获 Anthropic 黑客大奖的开源项目,它融合了多语言支持与丰富的实战钩子(hooks),让 AI 真正成长为懂上

160k|★★☆☆☆|今天
开发框架Agent语言模型

ComfyUI

ComfyUI 是一款功能强大且高度模块化的视觉 AI 引擎,专为设计和执行复杂的 Stable Diffusion 图像生成流程而打造。它摒弃了传统的代码编写模式,采用直观的节点式流程图界面,让用户通过连接不同的功能模块即可构建个性化的生成管线。 这一设计巧妙解决了高级 AI 绘图工作流配置复杂、灵活性不足的痛点。用户无需具备编程背景,也能自由组合模型、调整参数并实时预览效果,轻松实现从基础文生图到多步骤高清修复等各类复杂任务。ComfyUI 拥有极佳的兼容性,不仅支持 Windows、macOS 和 Linux 全平台,还广泛适配 NVIDIA、AMD、Intel 及苹果 Silicon 等多种硬件架构,并率先支持 SDXL、Flux、SD3 等前沿模型。 无论是希望深入探索算法潜力的研究人员和开发者,还是追求极致创作自由度的设计师与资深 AI 绘画爱好者,ComfyUI 都能提供强大的支持。其独特的模块化架构允许社区不断扩展新功能,使其成为当前最灵活、生态最丰富的开源扩散模型工具之一,帮助用户将创意高效转化为现实。

109.2k|★★☆☆☆|今天
开发框架图像Agent

gemini-cli

gemini-cli 是一款由谷歌推出的开源 AI 命令行工具,它将强大的 Gemini 大模型能力直接集成到用户的终端环境中。对于习惯在命令行工作的开发者而言,它提供了一条从输入提示词到获取模型响应的最短路径,无需切换窗口即可享受智能辅助。 这款工具主要解决了开发过程中频繁上下文切换的痛点,让用户能在熟悉的终端界面内直接完成代码理解、生成、调试以及自动化运维任务。无论是查询大型代码库、根据草图生成应用,还是执行复杂的 Git 操作,gemini-cli 都能通过自然语言指令高效处理。 它特别适合广大软件工程师、DevOps 人员及技术研究人员使用。其核心亮点包括支持高达 100 万 token 的超长上下文窗口,具备出色的逻辑推理能力;内置 Google 搜索、文件操作及 Shell 命令执行等实用工具;更独特的是,它支持 MCP(模型上下文协议),允许用户灵活扩展自定义集成,连接如图像生成等外部能力。此外,个人谷歌账号即可享受免费的额度支持,且项目基于 Apache 2.0 协议完全开源,是提升终端工作效率的理想助手。

100.8k|★★☆☆☆|1周前
插件Agent图像

markitdown

MarkItDown 是一款由微软 AutoGen 团队打造的轻量级 Python 工具,专为将各类文件高效转换为 Markdown 格式而设计。它支持 PDF、Word、Excel、PPT、图片(含 OCR)、音频(含语音转录)、HTML 乃至 YouTube 链接等多种格式的解析,能够精准提取文档中的标题、列表、表格和链接等关键结构信息。 在人工智能应用日益普及的今天,大语言模型(LLM)虽擅长处理文本,却难以直接读取复杂的二进制办公文档。MarkItDown 恰好解决了这一痛点,它将非结构化或半结构化的文件转化为模型“原生理解”且 Token 效率极高的 Markdown 格式,成为连接本地文件与 AI 分析 pipeline 的理想桥梁。此外,它还提供了 MCP(模型上下文协议)服务器,可无缝集成到 Claude Desktop 等 LLM 应用中。 这款工具特别适合开发者、数据科学家及 AI 研究人员使用,尤其是那些需要构建文档检索增强生成(RAG)系统、进行批量文本分析或希望让 AI 助手直接“阅读”本地文件的用户。虽然生成的内容也具备一定可读性,但其核心优势在于为机器

93.4k|★★☆☆☆|1周前
插件开发框架