技术文章知乎 - 严格鸽
文章通过C++代码示例展示抽象并非零成本:直接使用 std::vector<int>& 参数时,由于LLVM IR带有nonnull、readonly、dereferenceable等属性,编译器能提取data指针并生成SIMD向量化代码;但将其封装为View类后,这些属性丢失,且&&短路求值使编译器不敢提前解引用,导致无法优化。作者通过调整&&顺序或显式添加size陷阱检查恢复了优化,说明问题在于抽象包装丢失编译期信息。文章还提到Rust存在类似现象,并已提交LLVM issue。结论指出编译器优化有边界,但通常仍优于手动优化,适用于关注性能与编译器的C++开发者。
推荐收录,因为它以具体代码和Godbolt汇编对比揭示了C++抽象导致的编译期信息丢失和向量化失败,并给出了可复现的解决方案。适合关注性能优化、编译器行为或C++抽象设计的读者,可迁移价值在于理解编译器元数据如何影响优化,以及如何通过调整代码顺序或添加断言恢复优化。
工程实践知乎 - 严格鸽
本文记录了一道 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 模式的高性能计算开发者具有直接的参考和可迁移价值。
技术文章知乎 - 严格鸽
本文以 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 的排序与选择任务中。
工程实践知乎 - 严格鸽
文章记录了作者在 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 可能优于融合内存节省的算法”——具有可迁移的权衡思维,但需注意其结论仅适用于小规模、精度要求不高的类似任务。
技术文章知乎 - 严格鸽
文章深入解析了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上数据结构的设计模式与约束。可迁移价值在于组合作、原子操作在并发容器中的实际应用,但需注意文中未给出性能基准,实际选型时还需自行测试。
工程实践知乎 - 严格鸽
文章记录了在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类线性注意力算子的工程师提供了可迁移的实战参考。
工程实践知乎 - 严格鸽
文章记录了作者在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并行算法开发和性能优化的读者有直接参考价值,相关方法可迁移至其他密集计算或聚类类问题。
工程实践知乎 - 严格鸽
本文以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编程场景。
技术文章知乎 - 严格鸽
文章围绕 C++ 中的内存别名与严格别名规则展开,用多个汇编对比例子说明编译器为何会基于类型系统做激进优化,以及这种优化为什么会让看似“合理”的强转代码在 UB 场景下产生反直觉结果。作者进一步比较了 Fortran、Java、Rust 在别名约束上的差异,并结合 `__restrict`、`char`/`uint8_t`、`vector` 内部布局等例子,说明别名问题如何直接影响性能优化、接口设计和工程实践。文章的结论是:别名分析能带来显著性能收益,但在 C++ 中必须严格遵守语言规则,否则优化会把“直觉正确”变成“语义错误”,因此高性能代码需要更谨慎的类型设计和显式约束。
推荐收录,因为它不是泛泛讲“别名会影响优化”,而是通过具体汇编和跨语言对比把编译器推理、UB 边界和性能收益讲清楚了。对写 C++、做底层性能优化或设计高性能数据结构的读者,这篇文章有很强的可迁移价值。
技术文章知乎 - 严格鸽
文章围绕“为什么主流 C++ std::vector 通常维护三个指针”展开,从严格别名、TBAA、自动向量化失败案例一路分析到 vector 的三指针布局、ABI 约束和 libc++ 的 size-based vector 提案。作者先用一个 `data + size` 结构展示当元素类型与长度字段同型时,编译器可能因别名疑虑而保守重读循环上界,再对比三指针布局为何更利于把 `size()`、`capacity()` 和循环次数提前固化,从而获得更稳定的优化结果。文章还进一步补充了 libc++ hardening、ABI break、非法 reinterpret_cast 暴露问题以及与 LLVM RFC 的关系,说明三指针并非唯一原因,但确实兼顾了优化、接口语义和历史实现路径。
推荐收录,因为文章把一个看似“vector 实现细节”的问题,拆解成了别名分析、对象布局、自动向量化、ABI 和库演进等多个层面,能长期帮助读者理解 C++ 容器实现背后的工程逻辑。它不仅给出结论,还展示了何时 `data + size` 会让编译器保守、为何三指针更稳,以及现代库为何会在性能收益和 ABI 兼容之间做取舍。
技术文章知乎 - 严格鸽
文章围绕 C++ 中 `std::function<void(int,int)>` 为什么能够接收一个返回 `bool` 的可调用对象展开,先从标准条文解释“返回类型可调用”的含义,再结合 MSVC 的实现说明 `void` 与非 `void` 的调用路径确实不同。作者进一步分析了为什么 `[[nodiscard]]` 无法阻止这种行为,并给出了基于模板约束和 `std::invoke_result_t` 的写法,用于在转换成 `std::function` 之前就限制目标可调用对象的返回类型必须为 `void`。
推荐收录,因为文章不是简单复述语法,而是把标准语义、库实现和可编译约束串起来,清楚解释了一个常见却容易误解的 C++ 行为。对于需要写接口约束、回调封装或阅读标准库实现的读者,这种分析具有很强的迁移价值。