Blog Tags GitHub
All tags

hpc

5 posts tagged with "hpc"

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

Optimal Software Pipelining and Warp Specialization for Tensor Core GPUs（Twill：面向 Tensor Core GPU 的最优软件流水线与 Warp 特化）

arXiv 2512.18134，来自 CudaDMA/Singe 同一血统的团队（Bauer/Aiken + Stanford/NVIDIA/Toronto）。首次把软件流水线（Software Pipelining, SWP）与 warp 特化（Warp Specialization, WS）表述为一个可交给现成约束求解器求解的联合优化问题。核心洞见：SWP 与 WS 不应分开处理——WS 恰恰是实现 Tensor Core GPU 上 SWP schedule 所必需的代码生成手段（解决多 warp 协作发射 TC、寄存器工作集、变长延迟操作、阻塞同步中断四大难题）。方法上先用模调度（modulo scheduling，ℤLP 最优求解）得到初始 initiation interval I 与 modulo schedule M，再将 M/I 展开为一段直线程序 Q，在其上叠加 SWP 约束（Uniqueness/Consistency/Completion/Dependence/Capacity）、内存感知约束（活跃变量 live 传播 + 内存容量/SSA）、warp 分配约束（Warp Uniqueness/Variable Latency/Register Limit/Cross-Warp Spills/Concurrency），交给 SMT 求解器（Yices2 QFLIA）联合求解得到 M*、I*、A*。关键实现技巧：成本归一化（cycle count 比值保持的 ℤLP，令问题可解）、变长流式操作（TMA load）零延迟建模 + 深度作为 autotuning 参数、不可满足时单调增大 I/L 重试。系统 Twill 从 Triton TTGIR 抽取依赖图，无启发式、易扩展新架构、对单层循环保证最优。评估在 H100(Hopper)/B200(Blackwell) 上对 Flash Attention 前向/后向传播自动重新发现 FA3（ping-pong scheduling）与 FA4（三组 warp 策略）专家手工方案，性能达 cuDNN/FA 手工实现的 1%~2% 以内。

gpu-kernel hpc efficient-attention warp-specialization code-generation

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 通过降低占用率省寄存器。

gpu-kernel hpc memory-optimization warp-specialization

Singe: Leveraging Warp Specialization for High Performance on GPUs（用 Warp 特化实现 GPU 高性能的 DSL 编译器）

PPoPP'14 论文，来自 CudaDMA 同一作者团队（Bauer/Treichler/Aiken, Stanford）。提出 Singe——一个面向燃烧化学（combustion chemistry）的 DSL 编译器，用 warp 特化（warp specialization）作为替代数据并行的编程模型，把一个 CTA 内的 warp 划分为不同子计算，通过硬件生产者-消费者命名屏障（named barriers, PTX bar.arrive/bar.sync）细粒度同步。核心解决燃烧 kernel 的三大难题：超大工作集（每点数百个双精度变量）、不规则计算（多阶段、部分可并行）、不规则访存。Singe 用 warp 特化把大工作集拆分到不同 warp 的片上寄存器/shared memory，避免寄存器溢出。三类 kernel 案例：Viscosity（按物种求和分块，Store 模式）、Diffusion（对 dij 矩阵按列分块，Mixed 模式）、Chemistry（反应速率存寄存器经 shared 交换，Buffer 模式 + QSSA DAG 分块 overlap）。编译器四阶段：分块→映射（FLOPS/寄存器/局部性三指标）→命名屏障放置调度（无死锁定理，16 个命名屏障=SSA 寄存器分配）→代码生成。代码生成三关键技术避免指令 cache 抖动：代码 overlay（同时遍历 AST 森林 + 位掩码/间接分支）、常量去重（Fermi shared 广播 / Kepler shuffle）、warp 索引（不规则访存）。在 Fermi C2070 与 Kepler K20c 上对 DME/Heptane 机制，相比已高度优化的数据并行 baseline 最高加速 3.75×。

gpu-kernel hpc dsl-compiler warp-specialization code-generation

Tawa: Automatic Warp Specialization for Modern GPUs with Asynchronous References（Tawa：面向现代 GPU 的异步引用自动 Warp 特化编译器）

CGO 2026 论文，Cornell + NVIDIA 合作。提出 Tawa——首个面向现代 NVIDIA GPU（Hopper+）的**全自动 warp 特化编译流程**，从未经修改、无标注的 Triton 程序出发，自动分区为 producer/consumer warp group，自动管理软件流水线。核心创新是**异步引用（asynchronous reference, aref）**——一种形式化语义保证的 IR 抽象，将 warp 间通信建模为带硬件 mbarrier 的单槽信道（put/get/consumed），封装 TMA 描述符配置、mbarrier 初始化/相位管理等底层细节。编译器三步变换：① 基于语义标记的任务感知分区（iteration stmt → producer, tile stmt → consumer，依赖闭包 + 计算复制）；② 基于 aref 的多槽环形缓冲自动构建跨 warp 数据流；③ 多级软件流水线（细粒度：CUDA 地址计算与 MMA 重叠；粗粒度：CUDA Core 变换与 TC 流水线装配线）。额外优化：协作计算 warp group（多 WG 共享 tile 寄存器预算）、持久化 kernel。在 H100 上：GEMM 达 cuBLAS 水平（FP16 +1.01×，FP8 +1.06×）；Attention 达 FA3 手工优化内核的 96%，超 Triton 1.21×。ablation 显示 +Auto WS 提升 3.78×，最终 GEMM 达 718 TFLOPs/s（baseline 的 ~7×）。

gpu-kernel hpc efficient-attention warp-specialization compiler