CUDA

37 篇内容

工程实践NVIDIA Technical Blog

Speed Up Image Generation with FLUX 3 Image Klein on NVIDIA Blackwell GPUs

文章讲解如何在 NVIDIA Blackwell GPU 上用 TensorRT LLM VisualGen 加速 FLUX 3 Image Klein 的图像生成与编辑推理。它梳理了 Qwen3-VL 文本编码、扩散 Transformer 和自编码器流水线,并说明 NVFP4 量化、PrimsTS GEMM 融合内核、FP8/稀疏注意力、Cache-DiT 步缓存、多 GPU CFG/Ulysses 并行及 CUDA Graphs 的配置与取舍。在 B200 的 2048×2048 测试中,单卡组合优化使 T2I 最高提速 3.46×、I2I 4.45×;8 卡使 Base 的 T2I 中位延迟从 14.90 秒降到 3.00 秒,蒸馏模型从 2.73 秒降到 1.08 秒。文章用 LPIPS 衡量画质偏移,并给出离线推理和在线服务的 YAML 示例。结论主要适用于 Blackwell 数据中心 GPU 和 RTX PRO 6000,特定版本、分辨率和质量阈值下仍需重新验证。

推荐收录。文章不仅给出 NVFP4、融合 GEMM、稀疏/量化注意力和 Cache-DiT 的具体配置,还报告了 B200 单卡与 8 卡的端到端延迟及 LPIPS 质量代价,能帮助做扩散模型推理服务的读者理解“速度—画质—硬件”三角取舍。其风险是结论绑定 TensorRT LLM 版本、Blackwell 架构和 2048×2048 配置,迁移到其他 GPU 或模型需重新基准测试。

工程实践vLLM Blog

Migrating vLLM to the LibTorch Stable ABI

文章介绍 vLLM 从依赖 LibTorch 不稳定 ABI 迁移到 LibTorch Stable ABI 的完整工程实践。作者先解释 Stable ABI 的三层结构(C shim、C++ 包装与 headeronly 工具),说明其能让同一 vLLM wheel 在兼容 PyTorch 版本间无需重新编译。迁移按评估 API 缺口、把 AutogradFunctions 与 meta 函数移到 Python、建立独立 stable 库、用 STABLE_TORCH_LIBRARY/TORCH_BOX 注册算子、逐文件替换稳定 API,以及用 torch-abi-audit、单测和性能测试验证等步骤推进。文章还复盘 DeepGEMM 从 pybind11 转 TORCH_LIBRARY 再转 Stable ABI 的案例,并指出现有 Stable ABI 只覆盖 LibTorch 子集,部分算子需 dispatcher 调用且会带来少量 host 侧开销。最终 vLLM 0.33.0 的 CUDA wheel 达到 LibTorch ABI Stable,但兼容性仍受 Python API、CUDA 版本和第三方依赖约束。

推荐收录,因为文章给出了大型 C++/CUDA 扩展迁移到 Stable ABI 的可复用路径:分阶段改造、独立 stable 库、torch-abi-audit 审计和 CI 防回归,并诚实讨论了 API 子集与运行时开销等边界。适合维护 PyTorch 扩展、推理基础设施或跨版本二进制兼容的工程师参考,其依赖协调与逐步验证方法可迁移到类似项目。

技术文章NVIDIA Technical Blog

NVIDIA CUDA 13.5 Brings NVIDIA Vera Rubin GA Support

文章介绍 CUDA Toolkit 13.5 正式发布,重点包括 Vera Rubin 架构 GA 支持、Windows on Arm 扩展,以及 locality domain 与 MPS v3 的 GPU 调度改进。文中列出 cudaDevSmResourceSplit 等 API,说明 MPS 加权时间片、SM 分区借还、cgroup 路由和内存计费回退,并概述 CUDA Tile 多 GPU 通信与 SIMT 互操作预览。还覆盖 CUDA Python 批量拷贝、PDL、Graph 更新,CCCL 3.5 的 CUB 调优、批量 Top-K、可复现扫描与设备端计数归约,以及数学库和 Nsight 工具更新。文章属官方发布概览,部分能力为预览或限定平台,性能数据来自厂商,采用前仍需结合 release notes 与基准验证。

推荐收录:文章不是空泛发布稿,而是给出具体 API、代码示例、平台迁移限制和预览边界,能帮助 CUDA/GPU 开发者快速判断 13.5 的升级影响。适合 AI 基础设施、HPC 与性能工程读者作为版本迁移和 API 选型参考;但厂商性能数据需自行复测,且部分特性仍属预览或限定平台。

技术文章NVIDIA Technical Blog

Memory-Enforced GPU Sharing in Kubernetes with NVIDIA CUDA 13.4

文章介绍 CUDA 13.4 新增的应用透明 GPU 显存边界,以及如何通过开源 GPU Fractioning Operator 与 KAI Scheduler 在 Kubernetes 中按工作负载实施显存和计算共享。Operator 在每个 GPU 节点上管理 MPS,读取 Pod 注解为容器设置显存 request/limit 和 compute mode,并按 fraction 统计使用量;KAI Scheduler v0.18.0 的 NvFractions 模式负责选卡并写入注解。文中给出绝对显存、比例申请、可选 burst 上限、多卡分配,以及 time-slicing 与 sm-sharing 的配置示例和差异。关键边界是 sm-sharing 共享错误域、同一 GPU 不能混用两种计算模式,调度器仍需保证显存容量与模式约束;文章偏机制与配置说明,未提供性能实测。

推荐收录:文章给出了 CUDA 13.4 显存强制、GPU Fractioning Operator 与 KAI Scheduler NvFractions 的具体注解、安装方式和示例,能直接帮助平台团队设计 GPU 共享方案。适合 Kubernetes/AI 基础设施工程师;需注意它是厂商机制说明且缺少性能实测,采用前应验证 sm-sharing 错误域和调度约束。

工程实践vLLM Blog

vLLM Delivers Day-0 Support for MiniMax-M3.1-Flash

文章介绍 vLLM 对 MiniMax-M3.1-Flash 的 Day-0 支持,核心是 MSA 稀疏注意力、NVFP4 KV 缓存、NVFP4 MoE 与 DSpark 推测解码。作者分 NVIDIA Blackwell、Rubin 和 AMD MI350 三条硬件线,逐一说明 QMUL4 反量化、MegaMoE 提前释放 TMEM、indexer DCP、QKV/归一化/RoPE/KV 写入融合、FlashInfer PrimiTS 以及 AMD 查表反量化、mono kernel 等优化。AgentX 基准给出 GB300 418.9 tok/s/用户、Rubin 724.9 tok/s/用户与 191.5K tok/s/GPU、MI355X 1.333 亿 token/美元等结果。文章还展示基于 veRL、FSDP2、LoRA 的 Day-0 RL 流程和精度验证。结论依赖特定模型、GPU 与合成接受长度,属于厂商基准,独立复现前应谨慎外推。

推荐收录:文章不仅宣布 day-0 支持,还给出 NVFP4 KV 缓存、QMUL4 反量化、MegaMoE、indexer DCP、算子融合和 AMD 查表/单内核等实现细节,以及 GB300/Rubin/MI355X 的 AgentX 基准和显存收益。适合 LLM 推理服务、GPU 内核与性能优化读者,可迁移到低精度缓存、跨设备并行和 MoE 服务设计。注意其为厂商自测且强依赖特定硬件与合成接受长度,外推需独立复现。

技术文章NVIDIA Technical Blog

How to Read Data into GPU Memory in Python with NVIDIA GPUDirect Storage

文章围绕 NVIDIA KvikIO 与 GPUDirect Storage(GDS),演示如何在 Python 中把文件直接读入 GPU 内存,并解释 KvikIO 与 cuFile 两层各自独立的路径选择机制。作者给出安装 KvikIO、检查兼容模式与 GDS 阈值、用 CuPy 分配显存并用 CuFile 读取 1 GiB 文件的完整步骤。关键验证方法是设置 CUFILE_STATS_LEVEL=3 后用 gds_stats 观察 nvfs 与 posix 计数,从而区分 nvidia-fs 加速路径和 CPU 暂存兼容路径。文章还比较 KVIKIO_COMPAT_MODE 与 CUFILE_FORCE_COMPAT_MODE 的影响,并说明 cuFile 的任务拆分、并行线程和 8 MiB 阈值如何影响请求执行。其结论在特定 H100、NVMe/ext4、cuFile 1.18 环境验证,依赖 GDS 硬件与内核模块,属于厂商教程而非通用性能评测。

推荐收录:文章不仅给出 KvikIO 读取 GPU 内存的代码,还提供用 gds_stats 的 nvfs/posix 计数验证实际数据路径的可复现方法,并解释 KvikIO 与 cuFile 两层配置的边界。适合 AI 基础设施、GPU 数据加载和存储系统工程师参考;可迁移的是路径验证与参数调优思路,但需注意其结论依赖特定 GDS 硬件、驱动和文件系统配置。

工程实践vLLM Blog

Announcing Humming: Fast, Flexible Kernels for Quantized Inference

文章介绍 Ant Group 与 vLLM 社区开源的量化推理内核库 Humming,目标是用统一架构覆盖多量化格式、多 GPU 架构和 dense/MoE 负载。其核心是把量化描述与执行选择解耦,通过统一量化表示、三层配置和共享 CUDA 组件,用 JIT 为具体形状与设备生成内核。文章称其支持 Turing 起的 NVIDIA GPU 及部分 CUDA 兼容设备,并在 W4A16、W8A8、W4A8、W4A4 微基准上优于 Marlin、FlashInfer、DeepGEMM;DeepSeek-V4.1-Flash 在 8×H20 上吞吐提升约 14%–43%,TTFT 降低约 25%–32%。局限是性能多为项目方自测,收益依赖设备、形状和量化配置,autotune 仍在规划,落地前需独立验证。

推荐收录:文章具体说明了统一量化表示、三层配置、JIT 生成和共享 CUDA 组件的设计,并用多组微基准及 vLLM 端到端数据展示收益。适合推理引擎、GPU 内核和量化部署读者,其“量化描述与执行解耦”的架构可迁移到多格式、多设备内核库;需注意数据为项目方自测,应结合目标硬件复现。

工程实践NVIDIA Technical Blog

Accelerating Mamba-2 with High-Performance Forward Pass Kernels in NVIDIA CUDA Tile

文章介绍 Jane Street 与 NVIDIA 合作,用 Python DSL cuTile 为 Mamba-2 状态空间模型实现并优化前向传播的三个自定义 kernel。作者先从 Mamba-2 的分块递推与衰减项出发,推导出 log-space 累积和、状态跳跃公式和输出矩阵形式,再分别实现 chunk_cumsum、fwd_pass_state 和 compute_out,利用 ct.cumsum、ct.mma、寄存器内融合、因果掩码和 TF32 Tensor Core 加速。优化涵盖 DRAM 延迟提示、A/B 操作数交换以适配 Blackwell 的 Tensor Memory、转置内存布局以启用 TMA,以及按输入形状搜索 block size 和 occupancy 的自动调优。在单张 B300 上与 Triton/mamba_ssm 对比,前向 kernel 获得 1.14x 到 1.82x 加速,且随序列长度和 state size 增长更明显。结论仅限 kernel 级前向性能,端到端训练还依赖反向传播实现,读者需具备 CUDA/Tensor Core 和 Mamba 基础。

推荐收录,因为文章不仅给出 Mamba-2 前向的分块数学推导,还提供了三个 cuTile kernel 的完整代码、Blackwell 上的内存/Tensor Core 优化取舍和可复现 benchmark。对做 GPU 算子优化、序列模型推理/训练加速或 CUDA Tile 迁移的工程师有直接参考价值;不过结论是 kernel 级前向加速,不能直接等同于端到端训练收益。

工程实践vLLM Blog

Hyper Optimizing Kimi K3 in vLLM: Moving Fast and Responsibly

文章介绍 vLLM 团队在 NVIDIA GB300 上为 Kimi K3 构建专用解码路径:用 CuTe DSL 自定义内核替换并融合注意力、MoE 和通信算子,按批大小调优形状、重叠计算与通信、优化 MoE 调度,并在通信期间预取权重。实测 TP16、batch 1 且启用 DSpark 时达 842 tokens/s,batch 16–64 提升约 1.62–1.71×,AgentX 并发 1 的 P90 交互性提升 1.8×。文章还提出隔离改动、限定支持范围但保留服务特性、系统验证正确性和广泛扫描性能四项原则,并披露编译器重排与长上下文数值问题。适用边界为 GB300、TP8/TP16、batch≤64 及特定投机解码配置,硬件与负载针对性较强。

推荐收录:文章不仅给出 842 tokens/s、1.5–2× 吞吐和 1.8× P90 交互性等可核验数据,还把内核融合、通信重叠、MoE 调度与正确性验证方法讲清楚。适合 LLM 推理服务、GPU 性能优化和框架维护者参考,其“隔离改动+限定范围+自动验证”的协作模式可迁移到其他推理后端;主要风险是结论高度依赖 GB300/NVLink 与特定批负载。

工程实践vLLM Blog

Finish RL Training Faster: Bitwise-Consistent, High-Performance, Agent-Driven

文章讨论高性能 RL 训练中推理侧优化(低精度计算、FP8 KV 缓存、通信与 kernel 优化)改变数值,导致训练优化的策略与采样策略不一致,在 DeepSeek-V4-Flash 基线上 train-inference KL 从 2e-3 升到 29.9 且奖励回落。作者让训练前向直接调用 vLLM 推理 kernel,并改造这些高性能 kernel 使 rollout 与训练对每个采样 token 的 logprob 逐位一致(diff=0),从 kernel 确定性、批无关性、prefill/decode 阶段一致、跨 GPU 归约一致和统一数值契约五个维度展开。文章逐算子给出保留项、放弃项、提速方式与残余成本,并说明将跨 GPU 归约移出比较路径、CUDA Graph 与一致性共存的取舍。100 步训练中 diff 恒为 0,rollout 吞吐为基线 91.3%;迁移到 Mamba 混合 + NVFP4 的 Nemotron 后 SWE-bench Pro 分数保持而基线从 0.56 跌至 0.22。局限在于端到端墙钟收益尚未测量,TP>1 与解码上下文并行暂不支持,且 diff=0 不是训练质量的充分条件。

收录理由:文章以真实 100 步训练和 SWE-bench Pro 迁移实验为证据,逐算子给出数值对齐的保留、放弃、提速与残余成本,并公开 skill 与代码,属于可复现、可迁移的深度工程复盘。适合做 RL 训练、推理引擎或大模型基础设施的读者,用于理解 train-inference 数值不一致的根因、性能与一致性的取舍边界,以及如何用 Agent+skill 降低新模型的适配成本。主要风险是端到端收益与训练侧开销尚未完整量化,部分结论依赖具体硬件与并行布局。

工程实践vLLM Blog

vLLM Support for NVIDIA Vera Rubin NVL72: 7.8x Throughput over GB200 NVL72

文章介绍 vLLM 对 NVIDIA Vera Rubin NVL72 的早期支持与推理优化。Rubin 在 NVFP4 算力、HBM 带宽、NVLink 和 softmax 指数吞吐上较 GB200 显著提升;vLLM 借 Blackwell 内核兼容性实现 Day-0 支持,通过 FlashInfer 0.7.0 集成 Rubin 调优的 attention、GEMM、MoE 内核。针对 MoE decode 的内存瓶颈,作者用 CUDA 13.4 locality domain 与 split-N 权重切分,让 SM 只读本地 HBM,小 token 场景约 1.2x 加速。性能上,AgentX 中 MiniMax M3 达 GB200 每 GPU 7.84x 吞吐,MLPerf VLM 对 GB300 最高 3.7x。结果是初步的,通信未计入且优化仍在推进。

推荐收录:文章不仅给出 7.84x/3.7x 等性能数字,还公开了 Day-0 容器构建、Blackwell 内核兼容、locality domain 下的 split-N MoE 优化,以及 FlashInfer 后端配置和基准方法。对 LLM 推理性能工程师、平台/MLOps 团队有直接迁移价值;但需注意这是厂商早期预览,基准与内核配置仍可能随硬件和版本变化。

技术文章NVIDIA Technical Blog

Faster Scientific Image Analysis with NVIDIA cuPhoton

文章介绍 NVIDIA cuPhoton 开源 CUDA-X 工具包,面向天文、X 射线等高吞吐科学成像,目标是把从传感器读取到分类的完整处理链保留在 GPU 上。核心模块包括 xDataReader 直接加载 FITS、xRep 图像对齐、xPois PSF 匹配与差分、xFit 偶极拟合、xScan 候选审查以及 xRay 时域分析。文章以 Rubin Observatory LSSTCam 每 39 秒产生 3.2 吉像素曝光、每晚约 20 TB 和千万候选为背景,说明处理需在 60 到 120 秒告警窗口内完成,并给出多 GPU 下读取最高 14900 倍、信号处理最高 14550 倍的加速,但强调这是单操作而非端到端流水线数据。随后用合成天文图像演示加载、WCS 对齐、卷积核匹配差分、偶极拟合和审阅界面,并说明多 GPU/多节点 MPI 扩展方式。边界上,xDataReader 不支持 Rice 压缩和抖动浮点量化,示例依赖 CUDA 13/Linux 与特定 GPU,性能数字会随负载、实现和硬件变化。

推荐收录,因为文章给出了具体模块 API、FITS/GPUDirect Storage/nvCOMP 路径、PSF 匹配与偶极拟合公式、可运行代码和明确的不支持项,不是产品发布或泛泛宣传。它适合做 GPU 科学计算、高吞吐图像流水线和多 GPU 扩展的工程与研究人员,可迁移到需要在设备端完成加载、对齐、差分和批拟合的场景;但性能数字为单操作且依赖特定硬件,外推时需谨慎。

工程实践vLLM Blog

DeepSeek-V4.1-Flash on vLLM: 5x Agentic Throughput Since Day 0

文章介绍 vLLM 团队在 DeepSeek-V4.1-Flash 发布后三周内对其推理性能的系统优化,在 SemiAnalysis AgentX 智能体服务基准上实现低并发 1.9×、150 TPS 约束下 5.3× 的吞吐提升。优化分两类:一是 SWA 有界重放,利用因果编码器-解码器架构只缓存全局 KV,在前缀命中时重算最后 128 token 的滑动窗口 KV,并在解码端仅对每个请求末段 128 token 运行 21–39 层,配合可拆分 CUDA graphs 降低小 batch 的 kernel 启动开销;二是内核融合与并行化,集成 MegaAttention、Mega-mHC、Mega-Gate、DeepSelect 等开源 kernel,并融合 all-reduce、稀疏 MQA logits、NVFP4 压缩 KV 等。精度上 GSM8K、GPQA 未见显著差异,但 SWA 有界重放并非位精确,收益主要体现在长提示与高吞吐场景,KV offload 尚未启用。

推荐收录。文章给出可复现的性能证据:SWA 有界重放配合 trimming 与 CUDA graphs 带来约 30% TTFT 下降和 1.9×–5.3× 吞吐提升,并附 GSM8K/GPQA 精度验证,内核部分逐项列出融合与并行方案的加速比。对从事 LLM 推理服务、GPU kernel 优化和长上下文智能体部署的工程师,其中缓存与重算取舍、CUDA graph 拆分、算子融合思路可直接迁移;风险在于 SWA 有界重放非位精确,收益依赖长提示与特定部署配置。

技术文章NVIDIA Technical Blog

How DOCA GPUNetIO Unifies GPU-Initiated Networking Across the NVIDIA Software Stack

文章解析 NVIDIA DOCA GPUNetIO 与 GDA-KI 如何构成统一的 GPU 发起网络基础,使 CUDA 内核直接驱动 Ethernet、RDMA、Verbs 与 DMA,将 CPU 移出应用关键路径。文中说明 CPU 控制路径与 GPU 数据路径的分工、高/低层 API 差异,以及普通门铃、BlueFlame、CPU 协助门铃三种提交机制。它还比较开源 Verbs 版与完整 DOCA SDK 超集,指出前者可在运行时通过 dlopen 调用 SDK 功能。文章列举 NCCL GIN、NVSHMEM 3.7、NVQLink/Holoscan 等集成,并给出 Reduce-Scatter、shmem_put_bw 基准及约 2.6 微秒往返延迟数据。其边界是内容紧密绑定 NVIDIA GPU 与 ConnectX/DOCA 生态。

推荐收录,因为文章给出了统一 GDA-KI 架构的设计动机、控制/数据路径分工、门铃模式与 API 分层等可核验细节,并附 NCCL/NVSHMEM 基准数据说明 GDA-KI 对小消息和 CTA/QP 扩展的收益。适合从事分布式 GPU 通信、集合通信库或低延迟网络的工程师参考,其“共享实现点”与性能取舍思路可迁移到其他 GPU 网络栈评估,但需注意硬件与 DOCA 生态绑定。

技术文章NVIDIA Technical Blog

Control How Your GPU Shares Work with Green Contexts

文章介绍 CUDA 13.1 起通过 Runtime API 暴露的 Green Contexts 机制,允许单个进程显式切分 GPU 执行资源。核心手段是将 SM 子集和 workqueue 资源分配给指定 green context,并从中创建 stream,替代依赖线程本地隐式设备状态的传统模型。作者用延迟敏感小核与吞吐型大核并发的样例验证:在 148 SM 的 Blackwell GPU 上,专用 8 个 SM 后关键核延迟从默认高优先级下的 0.140 ms 降到 0.007 ms,比同优先级竞争的 3.727 ms 快约 500 倍。文章同时指出代价是可供 bulk 核使用的 SM 减少,且 stream 优先级无法抢占已在 SM 上执行的 block。适用于多组件并发、需要可预测性能和显式资源划分的场景,green context 为可选增量特性,现有应用可保持不变。

推荐收录,文章不仅给出 API 用法与代码骨架,还用实测数据量化了 green context 相对 stream 优先级的收益与上限(0.007 ms vs 0.140 ms),并明确说明 SM 减少的代价与适用边界。适合做 GPU 推理/训练性能优化、需要隔离延迟敏感算子的工程师参考,其显式资源划分思路可迁移到同类多工作负载共享 GPU 的场景。

工程实践NVIDIA Technical Blog

Accelerating a ROS 2 Node with an AI Agent and NVIDIA Isaac ROS

NVIDIA 技术博客讲解如何借助 AI 编码代理,把已有的 CUDA 加速 ROS 2 节点迁移到 rosidl::Buffer 抽象与 CUDA buffer backend。该后端基于 CUDA 虚拟内存管理实现零拷贝:当发布者与订阅者同主机、同 CUDA 设备、同 Linux 用户且使用受支持 RMW 实现时,可直接交换 GPU 常驻数据,否则回退 CPU 路径,并保持标准 ROS 2 消息接口不变。文章以 Depth Anything 3 TensorRT 节点为例,展示 migrate-node-to-rosidl-buffer 技能如何审计数据流、规划最小改动:订阅声明接受 cuda、输出消息分配 CUDA 缓冲、TensorRT 推理直接写入,点云与调试等可选 CPU 路径保持独立。验证用 Nsight Systems 确认 ROS 边界不再有载荷级主机-设备拷贝,并以 get_backend_type() 是否为 cuda 判断协商结果。内容依赖 ROS 2 Lyrical、Isaac ROS 5.0 与 NVIDIA 生态,未给出量化性能收益。

推荐收录:文章给出从依赖添加、订阅选项、CUDA 缓冲分配到推理直写与验证的完整迁移路径,并明确零拷贝生效前提(同主机、同 CUDA 设备、同用户、受支持 RMW)和 CPU 回退机制,可直接借鉴到 GPU 加速 ROS 2 节点与机器人数据通路优化。适合机器人中间件、边缘推理和 CUDA 数据流方向的工程师,其“审计—最小改动—验证”方法可迁移到其他节点。风险是依赖 NVIDIA 生态、缺少量化性能数据。

工程实践NVIDIA Technical Blog

Translating CUDA Tile Operations from Python to Rust Using Agentic AI

文章介绍 NVIDIA 用 Agentic AI 将 CUDA Tile 算子从 Python 翻译到 Rust 的工程实践。cuTile Python、Triton-TileIR 与 cuTile Rust 共享同一 CUDA Tile IR,因此移植不是重新优化,而是用更安全的主机语言重写同一 tile 程序,并可通过比对参考内核与生成内核的 Tile IR 做结构性验证。作者构建了一个有界多智能体流水线,覆盖分析、设备内核、主机与 FFI 代码、正确性和性能校验,每个阶段都以可机器检查的判决收尾,并配有 IR diff 分析与残余性能归因两个专职诊断子代理。团队据此移植了全部 24 个 TileGym 公开算子(约 40 个内核),在 DGX B200 上达到 cuTile Python 平均 99.5% 的性能,约三分之一算子反超。文章也指出 Rust 需显式声明特化、C ABI 之后无安全网、部分内核仍依赖非安全 API 等边界。

推荐收录。它以真实算子库为对象,给出了多智能体流水线的编排契约、判定路由、IR diff 验证与实测性能数据,并给出 softmax 的 Python/Rust 逐行对照和 C ABI 集成细节。适合做 GPU 内核、编译器前端或 AI 工程化落地的读者参考其可验证的翻译与代理编排方法,风险在于其结论依赖共享 Tile IR 与 Blackwell 工具链,迁移到其他体系需重新验证。

工程实践vLLM Blog

vLLM x Novita AI: Chord, Faster INT4 MoE for Kimi K2.x. Up to 1.3x on H200, 2.15x on Untuned B300

文章介绍 Novita AI 开源的 Chord 高性能 W4A16 MoE CUDA 算子,面向 BF16 激活、INT4 权重与 group-32 量化,针对 Kimi K2.x 推理场景。Chord 分两个内核族:基于 Humming 的 indexed 路径和源自 DeepGEMM 的 grouped SM90 路径,分别适配 prefill 与 decode 的路由形态。作者详述内核优化技术,如 WGMMA 流水线、按 expert token 数选择 block-M、stream-K 门控及 grouped 模式的 BM/BN/BK 启发式。在 H200 和 B300 上逐层测量显示,相比公共 Humming,indexed 路径取得 1.11–1.33x 加速,grouped 路径在 H200 EP8 达 1.16–1.35x,端到端吞吐提升约 4–10%。文章指出 B300 对比基于未调优的 Humming 默认配置,grouped 与 vLLM 集成仍在进行,性能数据为内核层面而非端到端保证。

推荐收录。文章不仅给出 Chord 算子的设计细节与优化手段,还提供了可复现的基准测试方法和逐层性能对比,并坦承 B300 对比的未调优前提与 grouped 集成尚在开发中。对从事 LLM 推理优化、GPU 内核开发或 MoE 量化部署的工程师而言,文中的 workload 驱动的调优策略、WGMMA 流水线设计和路由形态分析具有直接的迁移价值。

工程实践vLLM Blog

Adaptive Verification in vLLM: DSpark confidence-scheduled verification

该文介绍 vLLM 中基于 DSpark 置信度头的自适应推测解码验证机制。问题在于:固定 num_speculative_tokens(如 7)在低并发内存受限时收益高,但高并发下草稿 token 与真实 token 争抢算力,被拒 token 会浪费有效计算并拉低吞吐,且最优长度随负载变化无常量解。作者用 DSpark 置信头给出每个草稿位置的存活概率,将其转为全局 top-B 选择,B 由「每步期望产出 token / 单步耗时」最大化决定,并配合 varlen decode CUDA graph、启动期 profiled 成本表和单调化查表来降低决策开销。实测在 DeepSeek-V4-Pro-0813、8×B300 上,自适应验证使推测解码在并发 1 到 256 全程保持帕累托前沿,低并发表现为长草稿、高并发自动缩短。文末给出启用条件与限制:需 AttentionCGSupport.ALWAYS、不支持 --enforce-eager/LoRA/流水线并行,且开启后无法输出 logprobs。

推荐收录。它完整呈现了从现象(高位草稿接受率低于 10%)、成本模型、调度算法到 CUDA graph 与实测帕累托曲线的工程闭环,附可复现命令与明确限制。适合做 LLM 推理服务、吞吐优化的工程与研究者,其按负载动态分配验证预算、profiled 成本表驱动决策的思路可迁移到其他投机/批处理调度场景。

技术文章NVIDIA Technical Blog

Building Faster Cryptography with Carryless Multiplication in NVIDIA CUDA 13.3

文章详细介绍了NVIDIA CUDA 13.3中新引入的进位乘法指令,这一特性填补了GPU在此之前缺乏原生进位乘法硬件的空白。文章首先回顾了进位乘法在x86 CPU上的历史及其在认证加密、纠错码和零知识证明等密码学算法中的基础作用,然后解释了新PTX指令__nv_cmul的用法和编程模型,并通过基准测试展示了其在典型密码学操作上相对于纯软件实现的显著加速。文章还指出了当前支持的GPU架构范围,并讨论了该指令在特定算法中的适用性与性能边界。

推荐收录,因为文章不是简单的版本发布公告,而是深入解释了硬件指令的原理、应用场景和性能数据,对从事GPU密码学实现或高性能计算的开发者具有直接参考价值。读者可以迁移文中介绍的指令用法和优化思路到自己的项目中,内容具备长期技术参考性。

技术文章NVIDIA Technical Blog

Kernel Fusion in NVIDIA CUDA: Optimizing Memory Traffic and Launch Overhead

本文深入探讨了CUDA中的内核融合技术,旨在缓解GPU计算与内存带宽之间的瓶颈。文章从GPU内存带宽不足以完全利用计算能力这一常见问题出发,详细解释了内核融合如何通过合并多个CUDA内核来减少内存传输和内核启动开销。文中介绍了多种实现融合的方法,包括手工编写融合内核、利用CUDA图动态融合以及通过编译指示自动融合等,并对比了各自的适用场景和潜在局限。文章还指出,内核融合虽然能够提升性能,但可能增加寄存器使用量和代码复杂度,需在具体情况下权衡。整体而言,该文为GPU性能优化提供了系统性的实践指导,特别适用于受内存带宽或启动延迟限制的计算密集型应用。

本文来自NVIDIA官方技术博客,从问题背景、优化原理到多种实现方式进行了系统阐述,并提供了可操作的代码示例和权衡分析。对于从事GPU编程、性能优化或HPC领域的开发者,文中关于减少内存瓶颈和启动开销的策略具有直接指导意义,可迁移至其他受类似约束的计算架构中。因此,该文具备长期参考价值,适合收录为技术深文。

工程实践知乎 - 手抓饼熊

GTC 2026 技术分析:通过 CUTLASS Python 在 Blackwell GPU 上实现 GEMM 峰值 Tensor Core 性能

这篇文章基于 GTC 2026 关于 CUTLASS Python 的演讲,系统分析了如何在 Blackwell GPU 上围绕 GEMM 逐步逼近 Tensor Core 峰值性能。文章从基础 GEMM 的分块、TMA 传输、TMEM/RMEM/SMEM 流水线讲起,进一步比较了 2CTA、Warp Specialization、TMA Store、Persistent Kernel、Preferred/Fallback Cluster、Dynamic Scheduler 和 PDL 等优化手段在不同矩阵规模与瓶颈条件下的收益与边界。全文不仅给出性能数据,还强调了不同规模下瓶颈会从 DRAM 延迟、epilogue 开销转向 L2 命中率和调度效率,适合作为 Blackwell 上高性能 GEMM 编程的系统性参考。

推荐收录,因为它不是单纯介绍 CUTLASS Python,而是把 Blackwell 架构特性、内存层次、调度机制和性能结果串成了一条完整的优化路径。对做 GPU kernel、推理加速或高性能矩阵计算的读者来说,这些关于瓶颈切换、流水线隐藏和 cluster 调度的结论具有很强的可迁移性。

技术文章知乎 - 手抓饼熊

Tensor Core 编程与优化深度技术解析 — GTC 2026

这篇文章系统梳理了 NVIDIA Tensor Core 的编程模型与优化方法,覆盖 Warp Level、Warpgroup Level 到 Blackwell 第五代 Tensor Core 的演进,并把 Nsight Systems/Nsight Compute 的分析方法、TMA 流水线、CuTe/CUTLASS 编程框架串成一条完整优化路径。文章不仅解释了 MMA、tiling、寄存器/共享内存/TMEM 的数据流关系,还通过具体矩阵尺寸、SASS 形态和 profiling 指标说明如何判断瓶颈与选择合适的实现策略。适用边界主要在 NVIDIA CUDA GPU 生态,且部分 Blackwell/GTC 2026 相关特性具有较强版本依赖。

推荐收录,因为文章把 Tensor Core 的硬件代际、编程接口和性能分析工具放在同一框架下讲清楚,兼具原理、方法和实战排查路径。对于做 CUDA 算子优化、GPU 性能分析或 AI 基础设施开发的读者,它提供了可迁移的心智模型与调优抓手。

工程实践NVIDIA Technical Blog

Run High-Throughput Reinforcement Learning Training with End-to-End FP8 Precision

文章面向大模型强化学习训练的吞吐优化问题,讨论如何在端到端流程中引入 FP8 精度,以提升训练效率并降低算力与显存开销。作者指出,RL 训练通常分为采样/生成与策略更新两类高强度阶段,二者对计算、带宽和数值稳定性的要求不同,因此单点加速往往不足。文中围绕 NVIDIA 的 GPU 与软件栈,说明 FP8 在推理、前向与反向计算中的应用方式,以及如何在保持训练可用性的前提下扩大吞吐。其核心结论是:对于支持该精度路径的硬件和框架,端到端 FP8 可以显著提升高强度 RL 训练效率,但需要配合数值验证与实现适配,不能简单地对所有模型直接替换。

收录理由直接而明确:文章围绕“端到端 FP8 精度”提升 RL 训练吞吐展开,属于可迁移的工程优化经验,而不是单纯产品宣传。适合做大模型训练、GPU 性能优化和数值精度权衡的参考;但其效果依赖支持 FP8 的硬件与软件栈,迁移时需要验证稳定性。

工程实践Stanford Hazy Research

ThunderKittens 2.0: Even Faster Kernels for Your GPUs

这篇文章发布了 ThunderKittens 2.0,一个面向 GPU 的 CUDA 内嵌 DSL,并顺带给出一篇面向 Blackwell 的内核优化复盘。新版增加了 MXFP8/NVFP4 支持、调度与 tensor memory 控制能力、简化的构建结构,并把多个示例内核升级到更现代的 API。技术部分重点解释了为何原先放在 GEMM 路径中的两类 fence 实际上并非必需,作者通过 PTX 的因果顺序与 proxy 规则证明共享内存和 tensor memory 的可见性已被保证,去掉后带来约 20 TFLOP/s 的收益。文章还分析了 tcgen05.cp 与 tcgen05.mma 的隐式流水、PTX assembler 对单线程指令的保守串行化、cluster size 对占用率的影响,以及 tensor memory 触发的单 SM occupancy 限制。最后给出较完整的 GPU kernel benchmark 规范,强调输入分布、L2 冷热状态、warmup 和温度稳态都会显著影响 TFLOPs。整体内容适用于做 Blackwell/CUDA 内核、性能分析和基准设计的人,但结论主要建立在特定硬件与 PTX 行为之上,跨架构迁移时需要重新验证。

推荐收录,因为文章不是单纯发布稿,而是给出了 Blackwell GPU 内核优化的具体证据:PTX 因果/代理模型、assembler 生成差异、cluster 占用率和基准方法都配有可验证的观察。适合做 CUDA 内核、性能调优和基准设计的读者参考,但其结论强依赖 Nvidia 新架构与特定指令语义,迁移到其他 GPU 时需重新实测。

工程实践Stanford Hazy Research

How Many Llamas Can Dance in the Span of a Kernel?

这篇文章介绍了 Stanford Hazy Research 的一个吞吐优先版 Llama-70B megakernel:在 8 卡张量并行场景下,把预填充与解码、Paged KV cache、跨 GPU 通信和 CPU 侧调度尽量纳入单个内核的统一执行框架。作者沿用“解释器模板 + 指令序列”的设计,由 CPU 预先生成调度,内核内部按较粗粒度指令执行,从而减少 kernel launch、协调开销和 CPU 参与。文章重点分析了三类重叠:SM 内重叠、跨 SM 重叠和跨 GPU 重叠,并说明如何通过 warp specialization、显式同步和远程读写把这些优化放进同一套 megakernel 里。实测上,该原型在 ShareGPT 65,536 prompts 上端到端比 SGLang 快约 22%。其边界在于依赖较大的算子粒度和复杂的手工调度,当前更适合大模型高吞吐推理,而非所有模型与硬件都能直接套用。

推荐收录,因为它给出了把多 GPU 推理的调度、通信与计算统一进单核框架的直接实现证据,而不只是概念讨论。适合做大模型推理系统、CUDA 优化和高吞吐服务设计的参考,尤其对需要权衡 launch 开销、通信重叠和调度复杂度的工程场景很有迁移价值。

工程实践Stanford Hazy Research

One Kernel for All Your GPUs

这篇文章系统讲解了如何在 NVIDIA 多 GPU 平台上手写高性能通信核,重点面向 NVLink/NVSwitch 互联的单机多卡场景。作者先解释了跨进程共享 GPU 内存的三种路径:UVA、CUDA IPC 和手动 VMM,并说明为何生产环境更需要后两者以及它们的初始化开销与边界。随后文章分析了 NVSwitch 的广播/归约加速机制,以及 copy engine、TMA 和寄存器指令三种通信方式在带宽、并发和可融合性上的差异,给出在 B200 上的实测利用率。最后,作者把这些机制封装进 ThunderKittens 的 PGL 和 TKParallelTensor,展示了不到 100 行代码实现 all-reduce、all-gather、reduce-scatter 和 all-to-all,并在 8 卡 B200 上相对 NCCL 取得最高 2.6x 提升。文章的适用边界也很明确:它主要针对单机 NVLink/NVSwitch 域内的细粒度通信优化,且依赖 VMM、固定粒度显存和较强的 CUDA/编程模型理解,不直接覆盖跨节点通信。

收录价值很高,因为文章不仅给出性能结果,还把多 GPU 通信从内存映射、NVSwitch 机制到 kernel 设计完整串起来,直接提供了可复用的实现路径。适合做分布式训练、MoE、序列并行和自定义 collective 的工程参考;但读者需要接受其局限于单机 NVLink/NVSwitch 域,且实现门槛较高。

工程实践Stanford Hazy Research

ThunderKittens Now on Blackwells!

文章介绍了 ThunderKittens 针对 NVIDIA Blackwell/B200 架构的新一代 GEMM 与 Attention kernel 实现,并解释其为什么能接近或超过 cuBLAS、FA3 的性能。作者把重点放在“数据流”而非传统 CUDA 控制流上,围绕 5 代 tensor cores、tensor memory 和 CTA pairs 设计更深的流水线,通过 producer/consumer warpgroup 协作、persistent kernel、跨迭代预取 K/V、以及把输出累积器逐级写回共享内存和 HBM,尽量消除 pipeline bubble。文章还指出 B200 的 tensor core 更大,微基准上更像 128×128 systolic,因此只有当 M、N 维度足够大时才能充分吃满算力,小尺寸 GEMM 会按比例降速。它同时展示了如何把 Hopper 上的 kernel 结构迁移到 Blackwell,并说明 tensor memory 如何缓解 backward pass 的中间状态压力。整体结论是:Blackwell 的性能优化核心是提高并行数据供给深度,而 TK 的 tile 抽象恰好适配这一点,但收益依赖于特定硬件与形状假设。

推荐收录,因为文章直接给出了 Blackwell 上 GEMM/Attention kernel 的实现思路、硬件特性利用方式和明确的性能对比结果,而不是泛泛介绍新卡参数。适合做 GPU kernel、AI 加速和高性能计算的长期参考,尤其对需要把 Hopper 代码迁移到 Blackwell 的读者有可迁移的流水线设计价值;但其结论强依赖 B200 的 128×128 计算单元和特定 tile 形状。

工程实践Stanford Hazy Research

ThunderMLA: FlashMLA, Faster and Fused-er!

本文介绍 Stanford Hazy Research 的 ThunderMLA:针对 LLM 推理中变长请求和小批量 decode 的性能瓶颈,把原本分开的 attention/归约 kernel 融合成一个可由指令张量驱动的 megakernel。作者提出 ThunderKittens 的 interpreter template,在 GPU 上用虚拟指令集组织子 kernel,并通过全局 tensor 做依赖同步,从而减少 kernel launch、尾部效应和中间结果写回。文中还给出两种调度器:静态调度器与基于 makespan 反推的调度器,后者能进一步压缩执行时间约 10%。在 H100 上,ThunderMLA 相比 DeepSeek 的 FlashMLA 在多个 workload 上提升约 20–35%,但调度生成本身仍较慢,主要依赖可复用 schedule,适合推理场景而非通用低延迟单次执行。作者还强调该思路可迁移到 GQA、tensor parallel 的通信重叠以及 MoE 等数据流型 AI 工作负载。

收录价值明确:文章给出了可复现的性能证据、具体的 megakernel 设计和两类调度策略,而不是泛泛谈“更快”。适合做 LLM 推理、CUDA kernel 融合、GPU 调度与性能分析的参考,尤其对需要处理变长序列和小批量 decode 的工程场景可迁移。

工程实践Stanford Hazy Research

ThunderKittens: Bringing fp8 to theaters near you

这篇文章介绍 ThunderKittens 为 fp8 新增算子与 GEMM kernel 的实现思路,目标是在保持统一编程接口的同时支持量化数据类型。作者重点解释了 fp8 与 fp16/bf16/fp32 在寄存器布局上的差异,以及为何需要额外的线程间 shuffle 来完成数据重排,而不是简单复用原有 tile 逻辑。文章进一步分析了 ldmatrix/stmatrix 与 WGMMA 在 H100 上的使用方式,说明 fp8 在共享内存到寄存器加载时仍受 16 位指令接口限制,只能通过“先按 16 位加载再拆成两个 fp8”的方式绕过。为了降低 bank conflict,作者讨论了 32/64/128-byte swizzling 的适用条件,并把 fp8 tile 宽度下限提高到 32,以匹配更大的 core matrix 和更好的硬件利用率。整体结论是:fp8 kernel 可以复用 bf16 的整体结构,但必须围绕数据布局、bank conflict 和硬件指令约束做针对性改造,且效果高度依赖 NVIDIA H100 这类支持 WGMMA 的平台。

文中直接给出了 fp8 kernel 的布局、shuffle、swizzle 和 bank conflict 处理细节,并用 H100/WGMMA 约束解释了为何要这样设计。适合做 GPU 算子、AI 基础设施和高性能 CUDA 编程的参考,但结论明显依赖特定硬件代际。

工程实践Stanford Hazy Research

Easier, Better, Faster, Cuter

这篇文章介绍了 ThunderKittens 的第二轮升级,重点是把它从“可玩”推进到“可用”的 GPU kernel 工具箱。作者发布了多类新算子与实现,包括 fused Mamba-2、长卷积、线性注意力、RoPE、LayerNorm 和线性层,并给出在 H100 上相对现有 Triton/FlashFFTConv 实现的性能提升数据。文章还展示了 Llama3、Qwen2.5、nanoGPT、PyTorch Lightning 等集成示例,说明这些 kernel 已能用于推理和训练。除了算子本身,TK 2 还强化了构建系统、自动共享内存布局、全局 layout 描述、类型支持和大规模测试,降低了编写 kernel 的复杂度。作者强调,注意力加速的主要收益并非来自复杂算法,而是更好地利用 GPU、控制寄存器与内存流动;但当前实现仍偏向特定硬件与模型形态,且 FP8 支持尚未完善。

文中直接给出多项 kernel、注意力实现和基准对比,且附带可运行 demo 与训练集成,说明它不是概念展示而是可落地的工程经验。适合做 GPU kernel、LLM 推理/训练加速和算子设计的参考,但读者需要注意其性能结论强依赖 H100 与特定布局。

科研议题Stanford Hazy Research

Just read twice: closing the recall gap for recurrent language models

这篇文章讨论了高效递归语言模型在“联想回忆”(associative recall)上的缺口:虽然 Mamba、RWKV、Based 等架构在困惑度和推理效率上接近 Transformer,但在需要从长上下文中准确检索事实时仍明显落后。作者从理论上把联想回忆归约到集合不相交问题,说明仅靠固定大小的因果状态会对输入顺序高度敏感,因此需要更合适的读入顺序或非因果建模。基于这一洞察,文章提出 JRT-Prompt 通过重复上下文来帮助模型在多次扫描中决定“该记住什么”,以及 JRT-RNN 通过非因果编码器加因果解码器提升选择性记忆能力。实验显示两种方法都能在多个问答与信息抽取基准上显著提升准确率,JRT-RNN 可把强递归基线拉近到 Transformer 质量,同时保持线性时间/常数状态优势。文中也指出训练目标如何在因果与非因果模型间公平对齐仍是开放问题,结论主要适用于回忆密集型任务而非所有语言建模场景。

收录理由充分:文章不仅给出召回缺口的理论解释,还提供 JRT-Prompt/JRT-RNN 的具体方法、实验对比和 CUDA 实现结果。适合研究高效 LLM、线性注意力和长上下文检索能力的读者,迁移价值在于“重读/非因果”这类设计思路,但其收益主要集中在召回型任务。

工程实践Stanford Hazy Research

GPUs Go Brrr

文章围绕如何让 NVIDIA H100 的 Tensor Core 尽可能持续工作展开,作者先拆解 H100 的计算、共享内存、L2、寄存器和 TMA/WGMMA 等关键硬件资源,再用微基准说明真正的瓶颈不只是 HBM,而是共享内存延迟、地址生成开销和银行冲突。文章强调 WGMMA 与 TMA 是榨干算力的必要条件,同时指出其共享内存布局和 swizzle 规则文档混乱、易出错,需要精细控制数据布局与流水线。基于这些经验,作者发布了嵌入 CUDA 的 DSL ThunderKittens,用 tiles 抽象寄存器和共享内存中的张量操作,让复杂 kernel 代码显著简化。文中给出 FlashAttention-2 和线性注意力的实现与性能结果,说明在 H100 上可比常见实现进一步提升约 30%,但也暗示该方法高度依赖特定 GPU 架构与手工调优边界。

推荐收录,因为文章给出了 H100 上从硬件特性、布局约束到 kernel 实现的完整证据链,并以实际基准证明 ThunderKittens 能带来可观性能提升。适合做 GPU kernel、AI 加速和底层 DSL 设计的参考,但读者需注意其结论强依赖 Hopper 架构,且 swizzle/TMA 细节具有较强平台特定性。

科研议题Stanford Hazy Research

Based: Simple linear attention language models balance the recall-throughput tradeoff

这篇文章围绕“检索能力(recall)—生成吞吐—显存占用”之间的权衡展开,指出许多高效架构虽然推理更快,但在长上下文回忆和 in-context learning 上会明显弱于 Transformer。作者通过可控的合成关联回忆实验和真实语言建模评估,比较了注意力、滑窗注意力、线性注意力和 Mamba 等方法,发现单一原语都难以同时兼顾局部精确对齐与全局信息传递。基于这一分析,文章提出 Based:将极小窗口的滑窗注意力与二阶 Taylor 近似的线性注意力结合,用固定大小的递归状态在 recall 与 throughput 之间移动到更优的 Pareto 前沿。实验显示它在信息抽取、阅读理解和回忆型任务上优于先前子二次架构,同时在大模型推理吞吐上显著快于 FlashAttention-2 和 Mamba;但它仍未完全追上最强 Transformer,且效果强依赖特征映射与状态设计。文章还给出 IO/数据流感知的 CUDA 实现思路,说明算法与硬件协同对最终速度至关重要。

推荐收录,因为文章同时给出了问题定义、实证曲线、架构设计和 CUDA 实现优化,直接证明了其不仅是模型概念介绍,而是可复用的研究与工程方法。适合做长上下文模型、线性注意力和高吞吐推理系统的读者参考,但需注意其结论仍受任务类型和状态规模限制,未完全超越 Transformer。

技术文章Stanford Hazy Research

FlashAttention-2: Faster Attention with Better Parallelism and Work Partitioning

这篇文章介绍了 FlashAttention-2 的设计目标:在不做近似的前提下,继续压缩 Transformer 注意力的时间与显存开销,并把算子吞吐尽量逼近高效 GEMM。作者先回顾 FlashAttention 的分块、重算与在线 softmax 思路,再指出其瓶颈主要来自线程块与 warp 的工作划分不够理想,以及非矩阵乘法操作比例偏高。FlashAttention-2 通过减少 rescaling、边界判断和 causal mask 等非 matmul FLOPs,并将并行维度扩展到序列长度,从而在长序列、小 batch 场景下显著提升 GPU 利用率。与此同时,新版把 warp 之间的“sliced-K”改成更少通信的“sliced-Q”划分,减少 shared memory 读写与同步开销。文章还说明其支持更大的 head dimension、MQA/GQA,并在 A100/H100 上给出端到端训练与注意力基准,体现出该方法对长上下文训练和推理都有直接收益,但仍依赖具体 GPU 架构与实现细节。

文中明确给出算法改写、并行划分和 warp 通信优化的具体证据,并配有 A100/H100 与端到端训练基准,属于可复用的高质量系统/模型加速资料。适合做 GPU kernel 优化、注意力实现和长上下文训练的参考,但其收益高度依赖硬件与实现路径,迁移时需重新验证。

科研议题Stanford Hazy Research

FlashAttention: Fast Transformer Training with Long Sequences

这篇文章介绍了 FlashAttention 面向长序列训练的改进版本:在保持精确注意力、没有近似的前提下,通过 tiling、重计算和更细粒度的并行,把注意力的显存访问从二次复杂度降到线性,并进一步优化超长序列场景。作者指出,原版 FlashAttention 主要按 batch 和 head 维度并行,在长上下文但 batch 很小、head 数有限时会出现 GPU 并行度不足,因此新增了沿序列长度维度的并行。前向传播按行分块,反向传播按列分块,并借助 atomic operations 汇总梯度,从而减少 worker 间通信并提升吞吐。基准结果显示,在 8K 序列长度下,相比 PyTorch 和 Megatron-LM 实现可达 2.2-2.7 倍加速,端到端训练效率最高达 175 TFLOPs/sec/A100。实验还表明,把上下文从 2K 提升到 8K 能稳定改善困惑度和长程任务准确率,但收益主要出现在长序列、小批量训练场景。

推荐收录,因为文章给出了明确的算法改造、并行划分方式和可量化基准,不只是宣讲性能提升,而是解释了为什么长序列下原方案并行不足、如何改、改完后提升多少。适合研究注意力加速、长上下文训练和 GPU kernel 优化的读者,尤其对需要把理论收益落到真实训练吞吐的工程/研究工作很有参考价值。

工程实践Stanford Hazy Research

Fast Stable Diffusion with FlashAttention + Diffusers

这篇文章介绍了将 FlashAttention 接入 HuggingFace Diffusers,以加速 Stable Diffusion 推理的工程实践。作者先解释 FlashAttention 的核心原理:在 A100 等现代 GPU 上,注意力计算的瓶颈更多来自显存读写而非算力,因此通过融合 matmul 与 softmax、采用 tiling 和自定义 CUDA kernel 来减少内存访问。随后给出一个不到 70 行的集成方案,并证明生成结果与原版 Diffusers 一致。基准测试显示,相比未优化版本可获得 3-4 倍吞吐提升,相比 Diffusers 0.4.1 仍有约 33% 提升,A100 上最高约 1.04 images/s,T4 上收益更明显。文章还指出该方法能降低显存占用、放大 batch size,但优势主要来自注意力路径,效果依赖具体 GPU 的内存系统和原始实现是否已优化。

推荐收录,因为文中同时给出原理解释、70 行级别的集成路径和可复现的 benchmark 结果,直接证明了 FlashAttention 在扩散模型推理中的工程收益。适合做 GPU 性能优化、生成模型推理加速和框架集成的参考,但读者也应注意其收益强依赖硬件与基线实现。