- 博客(13)
- 收藏
- 关注
RSS订阅原创 [CUDA 优化实战] 纯手搓 flash decoding sm120 : 超越 flashinfer.single_decode_with_kv_cache
文章摘要 本文介绍了针对 LLM decoding 场景优化的 flash decoding 实现，对比了 PyTorch native 和 flashinfer 两个基线。作者详细阐述了 kernel 设计思路，包括： 采用 BN=64 的 kv tile 分块策略 配置 128 线程的 block 大小 动态计算 chunk_size 以充分利用 SM 使用 TMA 和 float4 向量化优化访存 实现双缓冲 Ks 缓存提升性能 实验在 RTX 5060（26 SM）上进行，完整代码已开源。文章深入探
2026-06-03 22:19:50
551
原创 [CUDA 优化实战] fmha sm120 : 超越 torch.sdpa(flash-attention-2)
本文适用于有一定 CUDA 编程基础，熟悉 GEMM/multi-head-attention 优化，对进阶 tensor core / 嵌入 PTX 指令 性能调优感兴趣的读者阅读实际上阅读本文前最好先阅读本人先发布的 hgemm sm120 、safe online softmax 以及 gemm 系列文章。因为可能有重合的知识点，本人就一笔带过了，而朋友们还云里雾里。
2026-05-23 17:04:40
530
原创 [CUDA 优化实战] hgemm sm120 - 100KB SMEM 中的“微雕”战争：Tensor-core、TMA、ldmatrix、mma
本文探讨了在RTX 5060笔记本GPU上利用TMA（Tensor Memory Accelerator）和mma指令优化hgemm（半精度矩阵乘法）的实现方法。作者通过对比cublas基准和自研的4种kernel实现（包括传统ldmatrix+mma、三级流水线优化、TMA加载优化等版本），发现尽管进行了深度优化，性能提升仍较为有限。文章详细介绍了各版本kernel的设计思路，特别是如何通过TMA硬件加速数据搬运、双缓冲寄存器等技术挑战性能极限，最终实现了一个近乎完美的kernel版本，解决了共享内存访问
2026-05-15 10:51:05
592
原创 [CUDA 优化实战] hgemm - 超越 cuBLAS：Tensor-core、cp.async、ldmatrix、mma
本文适用于有一定 CUDA 编程基础，熟悉 GEMM 优化，对进阶 tensor core / 嵌入 PTX 指令 性能调优感兴趣的读者阅读完整 kernel 和测试代码可以点击查看没错，这是超越 cuBLAS 系列之三（NV 还不针对移动端显卡调优的话，我们就还是超越。不过这估计是 GEMM 系列最后一篇了，也可能会写 fp8 的 kernel 但应该不会是纯血 fp8 gemm 了，感觉有点同质化。而且 fp8 本身就要搭配各种量化姿势使用）
2026-05-11 21:26:06
513
原创 [CUDA 优化实战] sgemm tf32 - 超越 cuBLAS：Tensor-core、cp.async、ldmatrix、mma
摘要：Tensor Core优化实战——超越cuBLAS的SGEMM TF32实现 本文详细介绍了如何通过CUDA编程实现超越NVIDIA cuBLAS库性能的SGEMM TF32计算。作者采用Tensor Core、cp.async、ldmatrix和mma等PTX指令进行优化，在RTX 5060移动显卡上成功超越了cuBLAS性能。文章包含5个逐步优化的kernel实现：从基础的cp.async+双ldmatrix+mma组合开始，通过swizzle解决bank conflict、grid swizzl
2026-05-05 20:16:22
566
原创 [CUDA 优化实战] RoPE - 手写算子的作用之 kernel fusion：减少访存次数、减少启动开销的优化技巧
本文探讨了手写CUDA算子在大模型RoPE位置编码中的性能优势。通过分析PyTorch原生实现的访存瓶颈，展示了手写算子如何通过内核融合和向量化技术大幅提升性能：将理论读写总量从约11倍压缩到2倍，削减80%无效访存。测试结果显示，手写向量化实现(317.36 GB/s)比PyTorch原生(57.13 GB/s)提升5.5倍性能。文章强调在AI编译器时代，手写算子通过减少显存带宽消耗、消除中间变量等优化仍具有不可替代的价值，特别针对内存密集型算子。完整代码已开源供参考。
2026-04-24 17:02:23
619
原创 [CUDA 优化实战] sgemm - 超越 cuBLAS：带你学会极致优化的矩阵乘法 cuda c++ 实现
本文详细介绍了如何通过CUDA优化实现超越cuBLAS的矩阵乘法(SGEMM)性能。作者从基础实现出发，逐步引入6种优化技术：向量化读写、block tiling共享内存、矩阵转置、swizzle解决bank conflict、写回事务合并和双buffer流水线。在RTX 5060显卡上，最终将4096x4096矩阵乘法耗时从16.59ms优化到13.85ms，超越cuBLAS的14.33ms。文章深入探讨了数据分块策略、线程分配、访存计算比等关键优化思路，展示了如何通过纯C++ CUDA代码充分发挥硬件性
2026-04-18 18:47:07
918
原创 白嫖 github action + Gemini 自动化每日精选论文
本文介绍了一个基于GitHub Action和Google AI Studio的AI论文精选工具，可自动从arXiv抓取论文并使用大模型进行评分分类。该工具支持自定义研究领域关键词和arXiv分类，利用Gemini等模型每天处理最多500篇论文，精选各领域最高分论文并生成摘要。项目开源在GitHub，用户可自由配置和扩展功能。
2026-04-13 11:35:31
676
原创 [CUDA 入门] 认识 CUDA “不存在的存储层级” - local memory
本文摘要：探讨了CUDA中local memory的使用场景与优化技巧。文章首先指出local memory虽逻辑上属于线程私有，但物理上仍位于global memory空间，访问性能远低于寄存器。重点分析了三种编译器会将变量分配到local memory的情况：非常量索引数组、大型结构体/数组以及寄存器溢出。特别强调了取地址操作(&)会导致变量强制分配到local memory这一隐蔽陷阱。通过具体代码示例对比了错误写法与两种正确优化方法：直接使用原生向量类型或union强绑定。最后警告开发者不要
2026-04-10 22:49:32
553
原创 WSL2 shell 使用优化 - 交互响应速度，网络配置
在 WSL2 中，所有涉及 cd 或 prompt 触发的逻辑，必须尽可能避免 Fork 外部进程。能用 Rust/C++ 二进制解决的，绝不用 Shell 脚本。能减少 I/O 往返的，绝不读写磁盘。
2026-04-09 21:04:52
232
原创 CUDA 开发者应该熟悉的数
本文总结了CUDA编程中影响性能的关键硬件参数和延迟数据。主要内容包括：Warp Size=32是基本执行单位，影响分支发散和内存合并；内存层级延迟从寄存器到全局内存呈数量级增长；共享内存存在32路Bank Conflict问题；PCIe/NVLink数据传输速度差异显著；Kernel启动开销约2-5微秒。文章强调，掌握这些"黄金数字"对实现极致性能至关重要，并提供了优化清单：确保Warp满载、内存访问合并、避免Bank Conflict、减少PCIe传输小数据、保证Kernel足够大。
2026-04-04 19:58:02
469
原创 [CUDA 优化实战] 矩阵转置 - 从 Padding 到 Swizzle：基础共享内存优化的艺术
矩阵转置（Transpose）是深度学习和高性能计算中极其基础的操作。看似简单的坐标交换 B[y][x]=A[x][y]，在 CPU 上可能只是两层循环，但在 GPU 这种吞吐导向的架构上，访存模式（Memory Access Pattern） 往往比计算逻辑更能决定性能的生死。如何写出一个高效的 Transpose Kernel 往往是考察 CUDA 工程师功底的试金石。
2026-04-03 13:21:38
442
原创 深度解读 DeviceQuery：理解你的 GPU 硬件属性
CUDA高性能计算开发/入门必读，深度解读 DeviceQuery：理解你的 GPU 硬件属性，懂 GPU 硬件是 kernel 优化的前提。
2026-04-03 11:17:12
638
空空如也
空空如也
TA创建的收藏夹 TA关注的收藏夹
TA关注的人