工程实践 知乎 - 严格鸽 2026/08/07
本文记录了一道 LeetGPU Hard 题目——FP32 滑动窗口自注意力在 T4 GPU 上的优化过程，目标是将 5000×64、window_size=16 的推理加速到极致。作者从 256 线程、每 warp 处理一个 Query 的基线出发，通过 shared memory bank conflict 消除、K/V 共享 tile、交错 query 提升 ILP、手写 PTX 倒数近似加牛顿迭代、去掉 Q 的 shared memory 缓存以提升 occupancy 等手段，将时间从约 0.28 ms 优化到 0.177 ms。文中逐一展示了各版本的关键改动和性能收益，最终方案减少了一次 shared memory staging，并通过 #pragma unroll 16 平衡了控制开销与寄存器压力，在单 T4 上取得目前第一名成绩。整体适用于固定窗口尺寸的批量自注意力加速场景，但优化策略（如 occupancy 调优、PTX 指令替换）可迁移至其他类似 GPU 内核开发中。
推荐收录，因为它不是简单的竞赛题解，而是一个完整的 GPU 内核优化案例，清晰展示了从并行设计到微架构调优的迭代过程。文中对 bank conflict 处理、warp shuffle reduction、PTX 指令精度与速度平衡、shared memory 容量与 occupancy 权衡等都有具体讨论，对学习 CUDA 优化或处理类似 reduction+attention 模式的高性能计算开发者具有直接的参考和可迁移价值。
技术文章 知乎 - 严格鸽 2026/07/28
本文以 LeetGPU 上的 Top-K 选择题目为切入点，系统介绍了在 GPU 上利用 Radix TopK 和 AIR TopK 算法高效完成最大 k 个元素选取的方法。作者首先处理了浮点数无法直接按二进制位比较的问题，通过 ordered_bits 转换将 float32 映射为保持数值顺序的无符号整数。然后详细演示了 Radix TopK 按高位到低位分桶、定位 splitter bucket 并逐步缩小候选范围的过程，给出了具体步骤和代码示例。在此基础上引入 AIR TopK（Adaptive and Iteration‑fused Radix Top‑K），说明当候选数据量下降到一定程度时，可以向专用 buffer 收集候选元素，以减少后续轮次的扫描开销。文章还提及 CUB 中的每轮 11 位处理策略，并给出性能测试结果。全文适用于了解 GPU 上的并行 Top‑K 算法实现，对大规模数据、k 较小场景有直接参考价值，但迭代融合部分的实测效果不显著，且测试环境主要基于 T4，硬件差异可能需要进一步验证。
本文从实际题目出发，深入讲解了 GPU 上 Radix TopK 和 AIR TopK 的算法原理与优化，提供了可运行的代码示例和清晰的图示，不是简单的问题解答或操作指南。对需要实现高性能 Top‑K 选择的 CUDA 开发者、并行算法研究者具有直接的参考价值，文中的 ordered_bits 转换、分桶筛选和自适应收集等思路可以迁移到其他基于 GPU 的排序与选择任务中。
工程实践 知乎 - 严格鸽 2026/07/25
文章记录了作者在 LeetGPU 平台完成 Hard 题目“GPT-2 Transformer Block”的完整优化过程。从朴素 CUDA 实现开始，逐步引入 FlashAttention 分块计算以避免完整 scores 矩阵的显存读写，但发现普通 CUDA Core 上 FlashAttention 融合并未带来理想加速。随后利用 Tensor Core 的 WMMA 指令，将数据转为 FP16 进行矩阵乘法，性能大幅提升至 22.88 ms。最终作者放弃 FlashAttention，转而直接用 WMMA 实现 QK 和 PV 计算，结合合并内存分配、手动分支消除等微调，将总耗时压至 13 ms。文中给出了核心 kernel 代码和不同方案的性能对比，揭示了在特定精度要求不高的 T4 环境下，直接使用 Tensor Core 可能比融合 kernel 更优的工程取舍，但未涉及长序列或大规模分布式推理场景，结论的环境依赖较强。
推荐收录为真实工程优化案例，因为它完整呈现了从算法理论（FlashAttention）到硬件指令（WMMA）再到性能调优的迭代链条，包含可运行的 kernel 代码和明确的性能数据，适合学习 GPU 编程和 Transformer 推理优化的开发者。文章的核心经验——“在特定硬件和精度约束下，直接使用 Tensor Core 可能优于融合内存节省的算法”——具有可迁移的权衡思维，但需注意其结论仅适用于小规模、精度要求不高的类似任务。
技术文章 知乎 - 严格鸽 2026/07/22
文章深入解析了NVIDIA cuCollections库中GPU哈希表static_map与dynamic_map的实现细节。static_map采用固定容量的开放寻址设计，默认以4个CUDA线程为一组（tile）协作插入或查找一个键，利用CAS原子操作解决并发写入冲突。文中详细说明了强类型哨兵、线性探测方案（ProbingScheme）、存储配置（Storage）等模板参数的作用，并结合代码剖析了插入和查找CUDA kernel的完整流程。dynamic_map通过维护多个static_map子表实现自动扩容，但旧数据不搬迁，导致查找需遍历所有子表，成本随扩容次数递增。整体上，文章为理解GPU上并发哈希表的协同设计与工程实现提供了清晰的代码级参考，适用于GPU加速数据库等场景，但缺乏性能评估与扩容策略的深入对比。
文章直接展示了GPU哈希表的线程协作、CAS并发控制等关键工程实现，代码解析深入，具备长期参考价值。适合CUDA开发者、数据库加速系统设计者理解GPU上数据结构的设计模式与约束。可迁移价值在于组合作、原子操作在并发容器中的实际应用，但需注意文中未给出性能基准，实际选型时还需自行测试。
工程实践 知乎 - 严格鸽 2026/07/20
文章记录了在LeetGPU平台上完成线性自注意力算子的GPU优化过程。题目要求计算给定Q、K、V矩阵的线性注意力，采用特征映射φ(x)=x+1(x>0)或exp(x)(x≤0)，并通过φ(K)ᵀV、归一化向量z和最终输出三个步骤实现。针对M=10000、d=128的规模，作者使用Split-K策略将K维度拆分为40份以增加并行度，并用atomicAdd合并局部状态矩阵S，同时将z的计算融合进同一kernel。进一步分析Shared Memory的bank conflict后，通过将q_tile的列宽从128×16改为128×17添加padding，使T4上的耗时从0.702ms降至0.656ms，提升约6.5%。该案例体现了在真实GPU受限条件下的算子优化取舍与可度量效果。
推荐收录，因为文章不是简单的调参记录，而是展示了一次完整的GPU算子优化流程：从问题定义、Split-K并行策略选择、atomicAdd合并权衡，到实测分析bank conflict并应用padding优化，每一步都有清晰的动机与性能对比。这为CUDA开发者以及需要优化Transformer类线性注意力算子的工程师提供了可迁移的实战参考。
工程实践 知乎 - 严格鸽 2026/07/18
文章记录了作者在LeetGPU平台解决三道Hard题并取得暂时排名第一的过程，覆盖了Multi-Agent Simulation、K-Means Clustering和All-Pairs Shortest Paths。对于多智能体模拟，利用随机分布的假设设计了基于空间网格的邻居查找优化；K-Means通过将聚类中心放入共享内存、展开循环并运用cooperative_groups实现跨块同步，提升并行效率；全源最短路采用分块Floyd‑Warshall算法，将矩阵划分为64×64个小块，使用共享内存和寄存器优化。每道题均给出了题意分析、核心思路、关键代码链接及性能结果，适用于测试数据范围，展示了从简单实现到充分利用GPU硬件特性的工程优化过程。
推荐收录，文章提供了三道具体GPU编程题目的完整解法与优化历程，不仅是代码片段，更展示了性能分析、网格划分、共享内存使用、cross‑block同步等实用工程技巧，对从事GPU并行算法开发和性能优化的读者有直接参考价值，相关方法可迁移至其他密集计算或聚类类问题。
技术文章 NVIDIA Technical Blog 2026/07/15
文章详细介绍了NVIDIA CUDA 13.3中新引入的进位乘法指令，这一特性填补了GPU在此之前缺乏原生进位乘法硬件的空白。文章首先回顾了进位乘法在x86 CPU上的历史及其在认证加密、纠错码和零知识证明等密码学算法中的基础作用，然后解释了新PTX指令__nv_cmul的用法和编程模型，并通过基准测试展示了其在典型密码学操作上相对于纯软件实现的显著加速。文章还指出了当前支持的GPU架构范围，并讨论了该指令在特定算法中的适用性与性能边界。
推荐收录，因为文章不是简单的版本发布公告，而是深入解释了硬件指令的原理、应用场景和性能数据，对从事GPU密码学实现或高性能计算的开发者具有直接参考价值。读者可以迁移文中介绍的指令用法和优化思路到自己的项目中，内容具备长期技术参考性。
工程实践 知乎 - 严格鸽 2026/07/12
本文以LeetGPU上的一道stream compaction题目为切入点，记录了从基础的多kernel前缀和实现到借鉴CUB的Decoupled look-back算法的完整优化过程。作者首先解释了题目要求和并行化思路，然后逐步融合算子、调整tile大小，最终引入基于状态的look-back机制，将全局扫描的读写复杂度降至2N。文章详细展示了tile状态机的设计、warp级扫描与lookback_sum的实现，以及针对写回路径的shared memory/global memory双路径优化。此外，还涵盖union复用shared memory、 sleep等待策略等工程技巧，并对比了Thrust和手写CUDA的性能差异。整体呈现了一个从简单到高性能的实战优化路径，适用于单GPU上的稳定stream compaction，但算法思想可迁移到其他并行扫描场景。
收录理由：本文不是单纯的代码片段或性能报告，而是展示了从算法选型到微观工程优化的完整决策链，包括对Decoupled look-back原理的剖析、状态机设计、warp级协同和硬件调优。对学习CUDA性能优化、并行扫描算法以及从CUB源码借鉴实践的读者具有直接参考价值，文中的look-back模式、shared memory复用和双路径写回策略均可迁移到类似的高性能GPU编程场景。
技术文章 NVIDIA Technical Blog 2026/07/10
本文深入探讨了CUDA中的内核融合技术，旨在缓解GPU计算与内存带宽之间的瓶颈。文章从GPU内存带宽不足以完全利用计算能力这一常见问题出发，详细解释了内核融合如何通过合并多个CUDA内核来减少内存传输和内核启动开销。文中介绍了多种实现融合的方法，包括手工编写融合内核、利用CUDA图动态融合以及通过编译指示自动融合等，并对比了各自的适用场景和潜在局限。文章还指出，内核融合虽然能够提升性能，但可能增加寄存器使用量和代码复杂度，需在具体情况下权衡。整体而言，该文为GPU性能优化提供了系统性的实践指导，特别适用于受内存带宽或启动延迟限制的计算密集型应用。
本文来自NVIDIA官方技术博客，从问题背景、优化原理到多种实现方式进行了系统阐述，并提供了可操作的代码示例和权衡分析。对于从事GPU编程、性能优化或HPC领域的开发者，文中关于减少内存瓶颈和启动开销的策略具有直接指导意义，可迁移至其他受类似约束的计算架构中。因此，该文具备长期参考价值，适合收录为技术深文。
工程实践 知乎 - 手抓饼熊 2026/06/23
这篇文章基于 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 调度的结论具有很强的可迁移性。
技术文章 知乎 - 手抓饼熊 2026/06/20
这篇文章系统梳理了 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 2026/04/20
文章面向大模型强化学习训练的吞吐优化问题，讨论如何在端到端流程中引入 FP8 精度，以提升训练效率并降低算力与显存开销。作者指出，RL 训练通常分为采样/生成与策略更新两类高强度阶段，二者对计算、带宽和数值稳定性的要求不同，因此单点加速往往不足。文中围绕 NVIDIA 的 GPU 与软件栈，说明 FP8 在推理、前向与反向计算中的应用方式，以及如何在保持训练可用性的前提下扩大吞吐。其核心结论是：对于支持该精度路径的硬件和框架，端到端 FP8 可以显著提升高强度 RL 训练效率，但需要配合数值验证与实现适配，不能简单地对所有模型直接替换。
收录理由直接而明确：文章围绕“端到端 FP8 精度”提升 RL 训练吞吐展开，属于可迁移的工程优化经验，而不是单纯产品宣传。适合做大模型训练、GPU 性能优化和数值精度权衡的参考；但其效果依赖支持 FP8 的硬件与软件栈，迁移时需要验证稳定性。
工程实践 Stanford Hazy Research 2026/02/19
这篇文章发布了 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 2025/09/28
这篇文章介绍了 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 2025/09/22
这篇文章系统讲解了如何在 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 2025/03/15
文章介绍了 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 2025/03/04
本文介绍 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 2024/11/27
这篇文章介绍 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 2024/10/29
这篇文章介绍了 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 2024/07/07
这篇文章讨论了高效递归语言模型在“联想回忆”（associative recall）上的缺口：虽然 Mamba、RWKV、Based 等架构在困惑度和推理效率上接近 Transformer，但在需要从长上下文中准确检索事实时仍明显落后。作者从理论上把联想回忆归约到集合不相交问题，说明仅靠固定大小的因果状态会对输入顺序高度敏感，因此需要更合适的读入顺序或非因果建模。基于这一洞察，文章提出 JRT-Prompt 通过重复上下文来帮助模型在多次扫描中决定“该记住什么”，以及 JRT-RNN 通过非因果编码器加因果解码器提升选择性记忆能力。实验显示两种方法都能在多个问答与信息抽取基准上显著提升准确率，JRT-RNN 可把强递归基线拉近到 Transformer 质量，同时保持线性时间/常数状态优势。文中也指出训练目标如何在因果与非因果模型间公平对齐仍是开放问题，结论主要适用于回忆密集型任务而非所有语言建模场景。
收录理由充分：文章不仅给出召回缺口的理论解释，还提供 JRT-Prompt/JRT-RNN 的具体方法、实验对比和 CUDA 实现结果。适合研究高效 LLM、线性注意力和长上下文检索能力的读者，迁移价值在于“重读/非因果”这类设计思路，但其收益主要集中在召回型任务。
工程实践 Stanford Hazy Research 2024/05/12
文章围绕如何让 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 2024/03/03
这篇文章围绕“检索能力（recall）—生成吞吐—显存占用”之间的权衡展开，指出许多高效架构虽然推理更快，但在长上下文回忆和 in-context learning 上会明显弱于 Transformer。作者通过可控的合成关联回忆实验和真实语言建模评估，比较了注意力、滑窗注意力、线性注意力和 Mamba 等方法，发现单一原语都难以同时兼顾局部精确对齐与全局信息传递。基于这一分析，文章提出 Based：将极小窗口的滑窗注意力与二阶 Taylor 近似的线性注意力结合，用固定大小的递归状态在 recall 与 throughput 之间移动到更优的 Pareto 前沿。实验显示它在信息抽取、阅读理解和回忆型任务上优于先前子二次架构，同时在大模型推理吞吐上显著快于 FlashAttention-2 和 Mamba；但它仍未完全追上最强 Transformer，且效果强依赖特征映射与状态设计。文章还给出 IO/数据流感知的 CUDA 实现思路，说明算法与硬件协同对最终速度至关重要。
推荐收录，因为文章同时给出了问题定义、实证曲线、架构设计和 CUDA 实现优化，直接证明了其不仅是模型概念介绍，而是可复用的研究与工程方法。适合做长上下文模型、线性注意力和高吞吐推理系统的读者参考，但需注意其结论仍受任务类型和状态规模限制，未完全超越 Transformer。
技术文章 Stanford Hazy Research 2023/07/17
这篇文章介绍了 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 2023/01/13
这篇文章介绍了 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 2022/10/13
这篇文章介绍了将 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 性能优化、生成模型推理加速和框架集成的参考，但读者也应注意其收益强依赖硬件与基线实现。