如何在 CUDA 中优化算法

我还维护了一个学习深度学习框架（PyTorch和OneFlow）的仓库 https://github.com/BBuf/how-to-learn-deep-learning-framework 以及一个如何学习深度学习编译器（TVM/MLIR/LLVM）的学习仓库 https://github.com/BBuf/tvm_mlir_learn , 有需要的小伙伴可以点一点star

本工程记录如何基于 CUDA 优化一些常见的算法，同时收录了大量 GPU/大模型相关的学习笔记和博客翻译。

友情链接：https://github.com/DefTruth/CUDA-Learn-Notes

目录结构

how-to-optim-algorithm-in-cuda/
├── cuda-kernels/          # 【代码】CUDA 基础算子手写优化实现
│   ├── reduce/
│   ├── softmax/
│   ├── elementwise/
│   ├── gemv/
│   ├── fast-atomic-add/
│   ├── upsample-nearest2d/
│   ├── indexing/
│   └── linear-attention/
├── cuda-mode/             # 【笔记+代码】CUDA-MODE 课程笔记（77+ 讲）
│   ├── code/              #   实验代码（YHs_Sample、cudabmk）
│   ├── slides/            #   课程讲义 PPT/PDF
│   ├── lectures/          #   Lecture 1-77+ 笔记
│   ├── blog-translations/ #   CUDA 博客翻译
│   ├── cute-dsl/          #   CuTe DSL 笔记
│   ├── lei-mao-blogs/     #   Lei Mao CUDA 博客转载
│   ├── practice/          #   课后实战
│   └── tech-notes/        #   GPU 技术专题
├── cutlass/               # 【笔记+代码】CUTLASS / CuTe DSL 学习
│   ├── code/              #   代码（cfx-article-src、cute-examples、swizzle）
│   ├── cute/              #   CuTe Layout 笔记
│   ├── gemm/              #   GEMM 实现解析
│   ├── tma/               #   TMA 教程
│   ├── wgmma/             #   WGMMA 教程
│   ├── swizzle/           #   Swizzle 机制笔记
│   ├── instructions/      #   CUDA 指令笔记
│   └── tutorials/         #   CUTLASS 翻译教程
├── triton/                # 【笔记+代码】Triton 学习
│   ├── code/              #   Python 代码实现
│   └── meetup/            #   Triton 中国 Meetup slides
├── large-language-model/  # 【笔记】大模型推理/训练优化笔记
├── ml-engineering/        # 【笔记】ml-engineering 翻译系列
├── pytorch/               # 【笔记+代码】PyTorch 博客翻译与代码
├── papers/                # 【笔记】论文阅读
│   ├── cuda/
│   └── mlsys/
├── ptx-isa/               # 【笔记+文档】PTX ISA 学习
├── tools/                 # 工具脚本（hfd.sh 等）
└── deprecated/            # 归档：过时/低相关内容

---

0. cuda-mode

• 课程的 Slides 和 脚本：https://github.com/cuda-mode/lectures
• 课程地址：https://www.youtube.com/@CUDAMODE
• 我的课程笔记：https://github.com/BBuf/how-to-optim-algorithm-in-cuda/tree/master/cuda-mode

一直想系统看一下某个课程系统和科学的学习下 CUDA ，感觉 CUDA-MODE 这个课程能满足我的需求。这个课程是几个 PyTorch 的 Core Dev 搞的，比较系统和专业。不过由于这个课程是 Youtube 上的英语课程，所以要学习和理解这个课程还是需要花不少时间的，我这里记录一下学习这个课程的每一课的笔记，希望可以通过这个笔记帮助对这个课程以及 CUDA 感兴趣的读者更快吸收这个课程的知识。这个课程相比于以前的纯教程更加关注的是我们可以利用 CUDA 做什么事情，而不是让读者陷入到 CUDA 专业术语的细节中，那会非常痛苦。伟大无需多言，感兴趣请阅读本文件夹下的各个课程的学习笔记。

1. cuda-kernels（CUDA 基础算子）

各基础算子的 CUDA 优化实现，代码位于 cuda-kernels/ 目录。

2. reduce

这里记录学习 NIVDIA 的reduce优化官方博客 做的笔记。完整实验代码见这里 , 原理讲解请看：【BBuf的CUDA笔记】三，reduce优化入门学习笔记 。后续又添加了 PyTorch BlockReduce 模板以及在这个模板的基础上额外加了一个数据 Pack ,又获得了一些带宽的提升。详细数据如下：

性能和带宽的测试情况如下 (A100 PCIE 40G)：

3. elementwise

将 oneflow 的 elementwise 模テン抽出来方便大家使用，这个 elementwise 模テン实现了高效的性能和带宽利用率，并且用法非常灵活。完整实验代码见这里 ，原理讲解请看：【BBuf 的CUDA笔记】一，解析OneFlow Element-Wise 算子实现 。这里以逐点乘为例，性能和带宽的测试情况如下 (A100 PCIE 40G)：

优化手段	数据类型	耗时(us)	带宽利用率
naive elementwise	float	298.46us	85.88%
oneflow elementwise	float	284us	89.42%
naive elementwise	half	237.28us	52.55%
oneflow elementwise	half	140.74us	87.31%

可以看到无论是性能还是带宽，使用 oneflow 的 elementwise 模テン相比于原始实现都有较大提升。

4. FastAtomicAdd

实现的脚本是针对half数据类型做向量的内积，用到了atomicAdd，保证数据的长度以及gridsize和blocksize都是完全一致的。一共实现了3个脚本：

1. https://github.com/BBuf/how-to-optim-algorithm-in-cuda/blob/master/cuda-kernels/fast-atomic-add/atomic_add_half.cu 纯half类型的atomicAdd。
2. https://github.com/BBuf/how-to-optim-algorithm-in-cuda/blob/master/cuda-kernels/fast-atomic-add/atomic_add_half_pack2.cu half+pack，最终使用的是half2类型的atomicAdd。
3. https://github.com/BBuf/how-to-optim-algorithm-in-cuda/blob/master/cuda-kernels/fast-atomic-add/fast_atomic_add_half.cu 快速原子加，虽然没有显示的pack，但本质上也是通过对单个half补0使用上了half2的原子加。

性能和带宽的测试情况如下 (A100 PCIE 40G)：

原子加方式	性能(us)
纯half类型	422.36ms
pack half2类型	137.02ms
fastAtomicAdd	137.01ms

可以看到使用pack half的方式和直接使用half的fastAtomicAdd方式得到的性能结果一致，均比原始的half的原子加快3-4倍。

5. UpsampleNearest2D

upsample_nearest_2d.cu 展示了 oneflow 对 upsample_nearest2d 的前后向的优化 kernel 的用法，性能和带宽的测试情况如下 (A100 PCIE 40G)：

框架	数据类型	Op类型	带宽利用率	耗时
PyTorch	Float32	UpsampleNearest2D forward	28.30%	111.42us
PyTorch	Float32	UpsampleNearest2D backward	60.16%	65.12us
OneFlow	Float32	UpsampleNearest2D forward	52.18%	61.44us
OneFlow	Float32	UpsampleNearest2D backward	77.66%	50.56us
PyTorch	Float16	UpsampleNearest2D forward	16.99%	100.38us
PyTorch	Float16	UpsampleNearest2D backward	31.56%	57.38us
OneFlow	Float16	UpsampleNearest2D forward	43.26%	35.36us
OneFlow	Float16	UpsampleNearest2D backward	44.82%	40.26us

可以看到基于 oneflow upsample_nearest2d 的前后向的优化 kernel 可以获得更好的带宽利用率和性能。注意这里的 profile 使用的是 oneflow 脚本，而不是 upsample_nearest_2d.cu ，详情请看 cuda-kernels/upsample-nearest2d/README.md 。

6. indexing

在 PyTorch 中对 index_add 做了极致的优化，我这里将 PyTorch 的 index_add 实现 进行了剥离，方便大家应用于其它框架。具体请看 indexing 文件夹的 README 。其中还有和 oneflow 的 index_add 实现的各个 case 的性能比较结果。整体来说 PyTorch 在 index Tensor元素很小，但Tensor很大的情况下有较大的性能提升，其它情况和 OneFlow 基本持平。详情请看 cuda-kernels/indexing/README.md 。

7. softmax

学习了oneflow的softmax kernel实现以及Faster Transformer softmax kernel的实现，并以个人的角度分别解析了原理和代码实现，最后对性能做一个对比方便大家直观的感受到oneflow softmax kernel相比于FasterTransformer的优越性。代码位于 cuda-kernels/softmax/。

8. linear-attention

学习一些 linear attention 的 cuda 优化技巧。代码位于 cuda-kernels/linear-attention/。

9. large-language-model

收集了和大语言模型原理，训练，推理优化相关的文章和学习笔记，位于 large-language-model/。

10. papers

GPU / AI 系统论文阅读笔记，位于 papers/，分为：

• papers/cuda/：CUDA 体系结构相关论文
• papers/mlsys/：ML 系统（分布式训练、推理）相关论文

11. triton

Triton 学习过程中的代码记录和学习笔记，位于 triton/，分为：

• triton/code/：Flash Attention、LayerNorm 等 Triton / PyTorch 实现
• triton/meetup/：Triton 中国 Meetup slides 汇总

12. ptx-isa

对 CUDA PTX ISA 文档的翻译和学习，位于 ptx-isa/。

13. pytorch

对 PyTorch 团队发布的 CUDA 技术的学习笔记和博客翻译，位于 pytorch/。

14. cutlass

CUTLASS / CuTe DSL 相关的学习笔记，位于 cutlass/。

学习资源

BBuf 公众号的笔记文章列表以及 CUDA/大模型基础设施领域的优质博客资源汇总，请参阅 RESOURCES.md。

how-to-optim-algorithm-in-cuda 快速上手指南

本仓库汇集了基于 CUDA 优化常见算法的代码实现、大模型推理/训练优化笔记以及 CUDA-MODE/CUTLASS/Triton 等前沿技术的学习资料。适合希望深入理解 GPU 底层优化、算子手写及大模型 Infra 的开发者。

环境准备

在开始之前，请确保你的开发环境满足以下要求：

• 操作系统: Linux (推荐 Ubuntu 18.04/20.04/22.04)
• GPU: NVIDIA GPU (推荐 Ampere 架构及以上，如 A100, RTX 3090/4090，以体验最佳性能优化效果)
• CUDA Toolkit: 版本 >= 11.0 (部分新特性如 TMA/WGMMA 可能需要 CUDA 12.0+)
• 编译器: GCC/G++ (版本需与安装的 CUDA Toolkit 兼容)
• 构建工具: CMake (>= 3.18), Make
• 可选依赖:
  • Python 3.8+ (用于 Triton 示例及部分脚本)
  • PyTorch (用于对比测试及参考实现)
  • OneFlow (部分算子参考了其实现)

安装步骤

本项目主要包含源代码和学习笔记，无需复杂的包安装，只需克隆仓库并配置编译环境即可。

1. 克隆仓库 使用 git 克隆项目到本地：

   git clone https://github.com/BBuf/how-to-optim-algorithm-in-cuda.git
   cd how-to-optim-algorithm-in-cuda
   

   (国内用户若下载缓慢，可使用镜像加速：git clone https://ghproxy.com/https://github.com/BBuf/how-to-optim-algorithm-in-cuda.git)

2. 验证环境 检查 CUDA 编译器是否可用：

   nvcc --version
   

3. 编译示例代码 (以 cuda-kernels 为例) 进入具体的算子目录（例如 reduce），通常包含 CMakeLists.txt 或 Makefile。以 reduce 为例：

   cd cuda-kernels/reduce
   mkdir build && cd build
   cmake ..
   make
   

   注：不同子目录的编译方式可能略有差异，请参考各子目录下的 README.md。

4. 安装 Triton 依赖 (如需运行 Triton 示例)

   pip install triton
   # 国内用户推荐使用清华源加速
   # pip install triton -i https://pypi.tuna.tsinghua.edu.cn/simple
   

基本使用

本仓库的核心价值在于阅读源码和理解优化思路。以下是几个典型的使用场景：

1. 运行基础算子优化测试 (Reduce)

编译完成后，在 build 目录下运行生成的可执行文件，观察不同优化策略（如共享内存、向量化加载）对带宽利用率的影响：

./reduce_benchmark

参考文档：查看 cuda-kernels/reduce 目录下的代码及知乎专栏文章了解优化原理。

2. 学习 CUDA-MODE 课程笔记

进入 cuda-mode 目录，这里整理了 77+ 讲的课程笔记、PPT 和实验代码。

cd ../cuda-mode/lectures
# 阅读 Lecture 1 的笔记
cat Lecture-01.md
# 运行对应的实验代码
cd ../code/YHs_Sample
# (根据具体实验的 README 进行编译运行)

3. 运行 Triton 算子示例

在 triton/code 目录下，可以直接运行 Python 脚本来体验 Triton 编写的高效算子（如 Flash Attention, LayerNorm）：

cd ../triton/code
python flash_attention_triton.py

4. 查阅 CUTLASS 与 CuTe DSL

若想学习新一代 GEMM 优化技术，可前往 cutlass 目录：

cd ../cutlass/cute
# 阅读 CuTe Layout 笔记
cat layout_notes.md
# 查看 swizzle 机制示例
cd ../swizzle

提示：每个子目录（如 softmax, elementwise, linear-attention）均包含独立的 README.md 或代码注释，详细解释了该算子的优化技巧、性能对比数据（基于 A100 等显卡）及实现原理，建议结合代码对照阅读。