📖CUDA-Learn-Notes: 🎉CUDA/C++ 笔记 / 技术博客: fp32、fp16/bf16、fp8/int8、flash_attn、sgemm、sgemv、warp/block reduce、dot prod、elementwise、softmax、layernorm、rmsnorm、hist etc. 👉News: Most of my time now is focused on LLM/VLM/Diffusion Inference. Please check 📖Awesome-LLM-Inference , 📖Awesome-SD-Inference and 📖CUDA-Learn-Notes for more details.

0x00 📖 博客目录

📖 大模型|多模态|Diffusion|推理优化 (本人作者)

📖 类型-标题	📖 作者
[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 Prefix Prefill Triton Kernel	@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推理部署|C++|算法|技术随笔 (本人作者)

📖 类型-标题	📖 作者
[推理部署][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

📖 CUTLASS|CuTe|NCCL|CUDA|文章推荐 (其他作者)

📖 类型-标题	📖 作者
[cute系列详解][入门]📖cutlass cute 101	@朱小霖
[cute系列详解][入门]📖CUTLASS 2.x & CUTLASS 3.x Intro 学习笔记	@BBuf
[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系列详解][GEMM]📖cute 之 简单GEMM实现	@reed
[cute系列详解][GEMM]📖cute 之 GEMM流水线	@reed
[cute系列详解][GEMM]📖cute 之 高效GEMM实现	@reed
[cute系列详解][GEMM]📖GEMM流水线: single-stage、multi-stage、pipelined	@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)	@JoeNomad
[cutlass教程][深入]📖cutlass bank conflict free 的shared memory 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
[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 结构	@紫气东来
[GPU通信架构][精解]📖NVIDIA GPGPU（四）- 通信架构	@Bruce

💡说明: 大佬们写的文章实在是太棒了，学到了很多东西。欢迎大家提PR推荐更多优秀的文章！

0x01 📖 CUDA Kernel目录 (面试常考题目)

• / = not supported now.
• ✔️ = known work and already supported now.
• ❔ = in my plan, but not coming soon, maybe a few weeks later.
• workflow: custom CUDA kernel impl -> Torch python binding -> Run tests.

📖 cuda kernel	📖 elem dtype	📖 acc dtype	📖 docs
✔️ sgemm_sliced_k_f32_kernel	f32	f32	❔
✔️ sgemm_t_tile_sliced_k_f32x4_kernel	f32	f32	❔
❔ hgemm_sliced_k_f16_f32_kernel	f16	f32	❔
❔ hgemm_t_tile_sliced_k_f16x2_f32_kernel	f16	f32	❔
✔️ sgemv_k32_f32_kernel	f32	f32	❔
✔️ sgemv_k128_f32x4_kernel	f32	f32	❔
✔️ sgemv_k16_f32_kernel	f32	f32	❔
❔ hgemv_k32_f16_kernel	f16	f16	❔
❔ hgemv_k128_f16x2_kernel	f16	f16	❔
❔ hgemv_k16_f16_kernel	f16	f16	❔
✔️ warp_reduce_f32/f16/bf16_kernel	f16/bf16/f32	f16/bf16/f32	link
✔️ block_reduce_f32_kernel	f32	f32	link
✔️ block_all_reduce_sum_f32_f32_kernel	f32	f32	link
✔️ block_all_reduce_sum_f32x4_f32_kernel	f32	f32	link
✔️ block_all_reduce_sum_f16_f16_kernel	f16	f16	link
✔️ block_all_reduce_sum_f16_f32_kernel	f16	f32	link
✔️ block_all_reduce_sum_f16x2_f16_kernel	f16	f16	link
✔️ block_all_reduce_sum_f16x2_f32_kernel	f16	f32	link
✔️ block_all_reduce_sum_bf16_bf16_kernel	bf16	bf16	link
✔️ block_all_reduce_sum_bf16_f32_kernel	bf16	f32	link
✔️ block_all_reduce_sum_bf16x2_bf16_kernel	bf16	bf16	link
✔️ block_all_reduce_sum_bf16x2_f32_kernel	bf16	f32	link
✔️ block_all_reduce_sum_fp8_e4m3_f16_kernel	fp8_e4m3	f16	link
❔ block_all_reduce_sum_i8_i32_kernel	i8	i32	link
✔️ dot_product_f32_kernel	f32	f32	❔
✔️ dot_product_f32x4_kernel	f32	f32	❔
❔ dot_product_f16_f16_kernel	f16	f16	❔
❔ dot_product_f16x2_f16_kernel	f16	f16	❔
❔ dot_product_f16_f32_kernel	f16	f32	/
❔ dot_product_f16x2_f32_kernel	f16	f32	/
✔️ elementwise_f32_kernel	f32	/	/
✔️ elementwise_f32x4_kernel	f32	/	/
❔ elementwise_f16_kernel	f16	/	/
❔ elementwise_f16x2_kernel	f16	/	/
✔️ histogram_i32_kernel	i32	/	/
✔️ histogram_i32x4_kernel	i32	/	/
✔️ softmax_f32_kernel (grid level memory fence)	f32	f32	❔
✔️ softmax_f32x4_kernel (grid level memory fence)	f32	f32	❔
❔ softmax_f32x4_kernel (per token)	f32	f32	❔
❔ safe_softmax_f32x4_kernel (per token)	f32	f32	❔
✔️ sigmoid_f32_kernel	f32	/	❔
✔️ sigmoid_f32x4_kernel	f32	/	❔
✔️ relu_f32_kernel	f32	/	❔
✔️ relu_f32x4_kernel	f32	/	❔
❔ relu_f16_kernel	f16	/	❔
❔ relu_f16x2_kernel	f16	/	❔
✔️ layer_norm_f32_kernel (per token)	f32	f32	❔
✔️ layer_norm_f32x4_kernel (per token)	f32	f32	❔
❔ layer_norm_f16_kernel (per token)	f16	f16	❔
❔ layer_norm_f16x2_kernel (per token)	f16	f16	❔
✔️ rms_norm_f32_kernel (per token)	f32	f32	❔
✔️ rms_norm_f32x4_kernel (per token)	f32	f32	❔
❔ rms_norm_f16_kernel (per token)	f16	f16	❔
❔ rms_norm_f16x2_kernel (per token)	f16	f16	❔
✔️ flash_attn_1_fwd_f32_kernel	f32	f32	link
❔ flash_attn_2_fwd_f32_kernel	f32	f32	link
❔ flash_attn_2_fwd_f16_kernel	f16	f32	link
❔ flash_attn_2_fwd_bf16_kernel	bf16	f32	link
✔️ hard_nms cpp only	f32	/	❔
✔️ notes v1(deprecated)	f32	f32	/

©️License

GNU General Public License v3.0

🎉Contribute

Welcome to 🌟👆🏻star & submit a PR to this repo!

📖 References

References

• flash-attention-minimal
• tiny-flash-attention
• cute-gemm
• cutlass_flash_atten_fp8
• cuda_learning