Paper Review

July 8, 2026

Can Tensor Cores Benefit Memory-Bound Kernels? (No!)（张量核心能加速访存受限 kernel 吗？不能！）

从理论与实证两方面证明：把 GPU 张量核心（Tensor Core）用于访存受限（memory-bound）kernel 无法带来实质性能收益。核心理论：基于 Roofline 模型与机器平衡度（machine balance）推导，在完全不重叠（fully un-overlapped）极端情形下张量核心相对 CUDA 核心的加速上界为 Speedup < 2 − 2/(1+α)；对于 FP64（张量核心算力 α=2，即 V100/A100/H100）该上界仅为 1.33×，即便 α→∞ 也不超过 2×。完全重叠情形下访存主导、计算加速对总时间无影响。实证在 A100 与 GH200 上用 STREAM SCALE、SpMV（cuSPARSE vs DASP）、迭代 stencil（EBISU/Brick vs ConvStencil/LoraStencil）三类代表性访存受限 kernel 验证：张量核心实现普遍慢于同等优化的 CUDA 核心实现，主因是张量核心访存模式次优。结论：访存受限 kernel 应优先用 CUDA 核心并聚焦访存优化（cache-aware、减少访存流量、流水/重叠），而非引入张量核心。

gpu-kernel hpc memory-optimization roofline-model

CudaDMA: Optimizing GPU Memory Bandwidth via Warp Specialization（通过 Warp 特化优化 GPU 内存带宽）

SC11 经典论文，提出 CudaDMA——一个通过 warp 特化（warp specialization）优化 GPU 显存带宽的可扩展 C++ 头文件库。核心思想：把一个 CTA（线程块）内的 warp 分成两类——DMA warp 专门负责片外 DRAM 与片上 shared memory 之间的数据搬运，compute warp 专门负责计算，从而把数据传输的形状/维度与计算的形状/维度解耦。DMA warp 借助 float4 向量化访存、指针算术外提、细粒度命名屏障（PTX bar.arrive/bar.sync）生产者-消费者同步，仅用 4 个 DMA warp/SM 即可饱和内存带宽（vs 基线需 40 warp）。提供 cudaDMASequential / cudaDMAStrided / cudaDMACustom 三类实例与单缓冲/双缓冲/手动双缓冲三种缓冲策略。在 NVIDIA Fermi (Tesla C2050) 上，微基准最高 1.37×，SGEMV 最高 3.2×、3D 有限差分 stencil 1.15×，FFT 通过降低占用率省寄存器。