基准测试：当前本机快照与复验协议

证据状态（2026-08-24）：仓库没有保存 V100/A100/H100 表格对应的原始 JSON、 Nsight 报告或 Release 附件。因此第 2–5 节中的跨 GPU 数字是不可审计历史快照：不得出现在 README 首页、Release、简历、面试数字卡或性能结论中，也不得用于横向比较。
第 1.5 节保留 RTX 3060 Laptop 的本机测量上下文，供重新执行 benchmark 时做 sanity check；它同样没有成套原始 JSON，不能冒充正式跨硬件报告。下一份可引用结果必须归档硬件/软件/commit、Google Benchmark JSON、命令、原始日志和 nsys/ncu 产物。
本仓库是参考级 CUDA 实现，未接入 CUTLASS 或 cuDNN 优化管线；即使将来产出正式结果，也不代表生产级性能上限。

1. 本机快照与未来复验协议

1.1 硬件平台

GPU 型号	架构	显存	Compute Capability	驱动版本
NVIDIA V100 SXM2	Volta	32 GB HBM2	`sm_70`	535.183
NVIDIA A100-40GB PCIe	Ampere	40 GB HBM2e	`sm_80`	535.183
NVIDIA H100 SXM5	Hopper	80 GB HBM3	`sm_90`	535.183

1.2 软件栈

组件	版本
CUDA Toolkit	12.2
PyTorch (SDPA 对比基线)	2.2.0+cu122
Google Benchmark	1.8.3
GCC	11.4
CMake	3.27

1.3 测试配置

数据类型: FP16（half / __half2）
注意力类型: Causal masking enabled（下三角掩码）
head_dim: 固定为 64（本实现支持 32、64、128）
Q/K/V 布局: [batch_size, num_heads, seq_len, head_dim]
计时方式: Google Benchmark Benchmark::UseManualTime()，CUDA Event 统计 kernel-only 耗时
预热迭代: 10 次
正式采样: 30 次（取中位数）
内存统计: cudaMemGetInfo 前后差值 + nvidia-smi dmon 交叉验证

1.4 编译参数

bash

cmake --preset release \
  -DCUFASH_ATTN_BENCHMARKS=ON \
  -DCUFASH_ATTN_ARCHS="70;80;90"
cmake --build --preset release --target cuflash_attn_bench

1.5 本机测量快照（RTX 3060 Laptop，2026-08-18，非正式结果包）

本节保留当前可执行的本机测量上下文，但原始 JSON 尚未随仓库归档；它只用于 sanity check，不可替代正式结果包。

环境：

项	值
GPU	NVIDIA GeForce RTX 3060 Laptop GPU（6144 MiB，`sm_86`）
驱动	591.44
CUDA Toolkit	12.0（nvcc 12.0.140）
本仓库 commit	`6860cbc`
编译	`cmake --preset release && cmake --build --preset release -j`（默认 arch `80;86`）
配置	batch=1, heads=8, kernel-only 耗时（CUDA Event），10 预热 + 30 采样中位数

复现命令：

bash

cmake --preset release && cmake --build --preset release -j
./build/release/cuflash_attn_bench

前向（Forward，非 causal，FP16 / BF16，Tensor Core 路径），单位 ms：

seq_len	head_dim	FP16 (ms)	BF16 (ms)	FP32 (ms)
256	64	0.158	0.162	0.841
512	64	0.496	0.502	2.86
1,024	64	1.76	1.79	12.0
2,048	64	6.33	6.45	47.8
4,096	64	23.5	23.7	177
4,096	128	84.1	82.6	379

前向（Forward，Causal，FP32），单位 ms：

seq_len	head_dim	FP32 causal (ms)
1,024	64	4.55
4,096	64	56.7
4,096	128	186

反向（Backward，dQ/dK/dV，非 causal），单位 ms：

seq_len	head_dim	FP16 (ms)	BF16 (ms)	FP32 (ms)
1,024	64	29.7	29.3	28.4
4,096	64	421	415	409
4,096	128	1160	1141	994

说明：本次为 RTX 3060 Laptop（sm_86，约 220 GB/s 显存带宽）单机实测，与第 2–5 节不可审计历史快照不可直接横向比较；A100/H100 数据必须在该机型上按同一协议重跑并归档原始产物后，才可进入正式报告。

2. 不可审计历史快照：多维度 Benchmark 矩阵（禁止引用）

2.1 端到端前向传播（Forward, FP16, Causal）

单位: 毫秒（ms），数值越小越好。表格内为 kernel-only 时间。

seq_len	batch	heads	V100 (ms)	A100 (ms)	H100 (ms)
1,024	1	8	0.42	0.19	0.11
1,024	1	16	0.81	0.36	0.20
1,024	8	8	3.12	1.38	0.78
1,024	8	16	6.18	2.72	1.52
1,024	16	8	6.15	2.71	1.51
1,024	16	16	12.21	5.38	3.00
4,096	1	8	5.82	2.18	1.15
4,096	1	16	11.48	4.30	2.26
4,096	8	8	44.80	16.72	8.79
4,096	8	16	88.92	33.16	17.44
4,096	16	8	88.50	33.02	17.37
4,096	16	16	TBD	65.60	34.50
8,192	1	8	22.50	7.80	3.85
8,192	1	16	44.50	15.40	7.60
8,192	8	8	173.0	59.8	29.5
8,192	8	16	TBD	118.0	58.2
8,192	16	8	TBD	118.5	58.5
8,192	16	16	TBD	TBD	115.0
16,384	1	8	88.0	28.5	13.2
16,384	1	16	172.0	55.8	25.8
16,384	8	8	TBD	TBD	100.0
16,384	8	16	TBD	TBD	TBD
32,768	1	8	TBD	105.0	45.0
32,768	1	16	TBD	205.0	88.0

注: TBD 表示该配置因时间或硬件调度原因尚未实测。可通过下方可复现脚本补充。

2.2 反向传播（Backward, dQ/dK/dV, FP16, Causal）

反向传播通常约为前向的 2.0–2.5× 耗时（额外两次 GEMM-like 计算与重计算路径）。

seq_len	batch	heads	A100 前向 (ms)	A100 反向 (ms)	反向/前向比值
1,024	1	8	0.19	0.41	2.16×
4,096	1	8	2.18	4.75	2.18×
8,192	1	8	7.80	17.20	2.21×
16,384	1	8	28.5	63.5	2.23×

3. 不可审计历史快照：与 PyTorch SDPA 的速度对比（禁止引用）

PyTorch 2.2 的 torch.nn.functional.scaled_dot_product_attention 默认优先调用 FlashAttention-2（通过 memory_efficient 后端）。此处以 A100-40GB、FP16、causal、head_dim=64 为基准，对比端到端前向耗时。

3.1 绝对耗时对比

seq_len	batch	heads	PyTorch SDPA (ms)	CuFlash-Attn (ms)	加速比 (SDPA/ ours)
1,024	1	8	0.12	0.19	0.63×
1,024	8	16	1.82	2.72	0.67×
4,096	1	8	1.05	2.18	0.48×
4,096	8	16	15.80	33.16	0.48×
8,192	1	8	3.50	7.80	0.45×
8,192	8	16	52.00	118.0	0.44×
16,384	1	8	12.00	28.5	0.42×

3.2 结果解读

指标	观察结论
绝对性能	CuFlash-Attn 约为 PyTorch SDPA 的 42%–67%。这是预期内结果——本库为从零手写 CUDA 的教学/参考实现，未接入 cuDNN 高度调优的管线，也未使用 CUTLASS 的 warp-specialization 与自动调度。
趋势一致性	随着 `seq_len` 增大，两者的 scaling 曲线基本平行（均接近 $O (N^{2})$ 计算量），说明本实现的核心 tiling 与 softmax 重缩放逻辑是正确的。
小序列开销	在 `seq_len=1K` 时，本实现相对差距最大（0.63×），因为 kernel launch、causal mask 边界判断与 warp-level reduction 的固定开销占比更高。

面试要点: 当被问及"为什么比 PyTorch 慢"时，应区分算法正确性与工程优化深度。本库验证了 $O (1)$ HBM 内存的 FlashAttention 算法，而 PyTorch 后端调用的是 NVIDIA/cuDNN 经过数千小时调优的生产 kernel。

4. 不可审计历史快照：内存占用对比（禁止引用）

4.1 峰值显存（Forward + Backward 单步）

单位: MB。对比标准 Attention（ $O (N^{2})$ 中间激活）与 FlashAttention（ $O (N)$ 激活）。

seq_len	batch	heads	标准 Attention (MB)	CuFlash-Attn (MB)	节省比例
1,024	1	8	32.5	16.2	50.2%
4,096	1	8	512.0	65.0	87.3%
8,192	1	8	2,048.0	130.0	93.7%
16,384	1	8	8,192.0	260.0	96.8%

4.2 内存 Scaling 公式验证

标准 Attention（Materialized Softmax）: $M_{s t d} = 4 \times b a t c h \times h e a d s \times s e q_l e n^{2} \times s i z e o f (FP16)$
- 以 seq_len=8192, batch=1, heads=8 为例： $4 \times 1 \times 8 \times 8192^{2} \times 2 \approx 4096 MB = 4 GB$
- 实际测量含 PyTorch 框架开销与临时 buffer，约为 2 GB forward + 2 GB backward = 4 GB 量级。
FlashAttention（Tiled Online Softmax）: $M_{f a} = 2 \times b a t c h \times h e a d s \times s e q_l e n \times h e a d_d i m \times s i z e o f (FP16) + w o r k s p a c e$
- 以同样配置为例： $2 \times 1 \times 8 \times 8192 \times 64 \times 2 \approx 16.8 MB$ 的 Q/K/V/O 张量本身
- 加上 SRAM 大小的 softmax 统计量（m, l 向量）与 tile buffer，总计约 130 MB（含 CUDA context、cuBLAS workspace、pytorch allocator 碎片）。

核心结论: 当 seq_len ≥ 4K 时，FlashAttention 的内存优势呈指数级体现，使得在 40GB A100 上单卡可训练 seq_len=128K 的模型，而标准 Attention 在 seq_len=32K 时就会 OOM。

5. 不可审计历史快照：不同 GPU 架构的 Scaling 分析（禁止引用）

5.1 理论峰值对比

指标	V100 (sm_70)	A100 (sm_80)	H100 (sm_90)	A100/V100	H100/A100
HBM 带宽	900 GB/s	2,039 GB/s	3,350 GB/s	2.27×	1.64×
FP16 Tensor Core 算力 (dense)	31.4 TFLOPS	312 TFLOPS	989 TFLOPS	9.94×	3.17×
FP16 Tensor Core 算力 (sparse)	62.8 TFLOPS	624 TFLOPS	1,979 TFLOPS	9.94×	3.17×
带宽-算力比	34.9 FLOP/Byte	153 FLOP/Byte	295 FLOP/Byte	—	—
典型 TDP	300 W	400 W	700 W	—	—

5.2 实测 Scaling（固定配置: batch=1, heads=8, head_dim=64, causal FP16）

seq_len	V100 (TFLOPS)	A100 (TFLOPS)	H100 (TFLOPS)	V100 带宽利用率	A100 带宽利用率	H100 带宽利用率
1,024	2.1	4.5	8.2	52%	48%	46%
4,096	2.8	7.5	14.2	69%	80%	80%
8,192	3.0	8.5	17.3	74%	91%	97%
16,384	3.1	9.3	20.1	76%	>100%*	>100%*

*: 利用率超过 100% 是因为 FlashAttention 的在线 softmax 重计算减少了 HBM 流量，使得有效带宽高于原始理论值；或受限于 timer 精度与异步 overlap 的测量误差。在严谨分析中，应使用 Roofline 模型将其归一化到计算强度与实际 HBM 流量（见 Roofline 分析）。

5.3 理论 vs 实测加速比

seq_len	A100/V100 (理论)	A100/V100 (实测)	H100/A100 (理论)	H100/A100 (实测)
1,024	2.27× (带宽上限)	0.45 / 0.19 ≈ 2.37×	1.64× (带宽上限)	0.19 / 0.11 ≈ 1.73×
4,096	2.27× (带宽上限)	5.82 / 2.18 ≈ 2.67×	1.64× (带宽上限)	2.18 / 1.15 ≈ 1.90×
8,192	2.27× (带宽上限)	22.5 / 7.8 ≈ 2.88×	1.64× (带宽上限)	7.8 / 3.85 ≈ 2.03×
16,384	2.27× (带宽上限)	88.0 / 28.5 ≈ 3.09×	1.64× (带宽上限)	28.5 / 13.2 ≈ 2.16×

分析

V100 → A100: 实测加速比随 seq_len 增大而超过理论带宽比（2.27×）。这是因为：
- A100 的 Tensor Core 支持更细粒度的 FP16 MMA（ $16 \times 8 \times 16$ ），而 V100 为 $8 \times 8 \times 4$ ，tiling 效率在 A100 上更高。
- Ampere 架构引入的异步拷贝（cp.async）未被本实现显式使用，但编译器自动生成的 LDGSTS 流水仍优于 Volta。
A100 → H100: 实测加速比低于理论带宽比（1.64×）。这是因为：
- 本实现尚未针对 Hopper 的 TMA (Tensor Memory Accelerator) 与 WGMMA (Warp Group-level MMA) 进行重写。
- Hopper 的 Thread Block Cluster 与分布式共享内存特性未启用，导致未能发挥 SM90 的理论峰值。
- 预期: 若将 kernel 升级至 FlashAttention-3 风格（TMA + WGMMA），H100 上可望获得 2.5×–3.0× 于 A100 的性能。

6. 未来正式结果的复验命令

6.1 本地运行（需 CUDA GPU）

bash

# 1. 克隆仓库
git clone https://github.com/your-org/cuflash-attn.git
cd cuflash-attn

# 2. 使用 CMake Preset 构建 benchmark 目标
cmake --preset release -DCUFASH_ATTN_BENCHMARKS=ON
cmake --build --preset release --target cuflash_attn_bench

# 3. 运行全量 benchmark（约 20–30 分钟）
./build/release/bench/cuflash_attn_bench \
  --benchmark_time_unit=ms \
  --benchmark_repetitions=30 \
  --benchmark_report_aggregates_only=true \
  --benchmark_filter="BM_FlashAttention_Forward.*"

# 4. 导出 JSON 供后续分析
./build/release/bench/cuflash_attn_bench \
  --benchmark_out=results.json \
  --benchmark_out_format=json

6.2 Docker 环境（推荐，完全可复现）

dockerfile

# Dockerfile.benchmark
FROM nvidia/cuda:12.2.0-devel-ubuntu22.04

RUN apt-get update && apt-get install -y \
    cmake ninja-build git python3 python3-pip \
    libbenchmark-dev libgtest-dev \
    && rm -rf /var/lib/apt/lists/*

# 安装 Google Benchmark（若 apt 版本过旧）
RUN git clone https://github.com/google/benchmark.git /tmp/benchmark \
    && cmake -S /tmp/benchmark -B /tmp/benchmark/build \
       -DBENCHMARK_DOWNLOAD_GTEST=ON -DCMAKE_BUILD_TYPE=Release \
    && cmake --build /tmp/benchmark/build -j$(nproc) \
    && cmake --install /tmp/benchmark/build

WORKDIR /workspace
COPY . /workspace/cuflash-attn
RUN cmake --preset release -DCUFASH_ATTN_BENCHMARKS=ON \
    && cmake --build --preset release --target cuflash_attn_bench

CMD ["./build/release/bench/cuflash_attn_bench"]

bash

# 构建并运行
docker build -f Dockerfile.benchmark -t cuflash-bench .
docker run --gpus all -v $(pwd)/results:/results cuflash-bench \
  --benchmark_out=/results/bench_$(nvidia-smi --query-gpu=name --format=csv,noheader | tr ' ' '_').json

6.3 与 PyTorch SDPA 的对比脚本

python

# scripts/bench_vs_pytorch.py
import torch
import torch.nn.functional as F
import time

def benchmark_pytorch(batch, heads, seq_len, head_dim=64, causal=True, n_iters=100):
    device = torch.device("cuda")
    q = torch.randn(batch, heads, seq_len, head_dim, device=device, dtype=torch.float16)
    k = torch.randn(batch, heads, seq_len, head_dim, device=device, dtype=torch.float16)
    v = torch.randn(batch, heads, seq_len, head_dim, device=device, dtype=torch.float16)

    # 预热
    for _ in range(10):
        out = F.scaled_dot_product_attention(q, k, v, is_causal=causal)
    torch.cuda.synchronize()

    # 计时
    start = torch.cuda.Event(enable_timing=True)
    end = torch.cuda.Event(enable_timing=True)
    start.record()
    for _ in range(n_iters):
        out = F.scaled_dot_product_attention(q, k, v, is_causal=causal)
    end.record()
    torch.cuda.synchronize()
    return start.elapsed_time(end) / n_iters

if __name__ == "__main__":
    for seq_len in [1024, 4096, 8192, 16384]:
        t = benchmark_pytorch(batch=1, heads=8, seq_len=seq_len)
        print(f"seq_len={seq_len}: {t:.3f} ms/iter")

7. 正式结果发布门槛

在 L40S、A100 或其他云 GPU 上发布新的性能数字前，必须先通过数值正确性测试，并将以下产物与结果页一并提交：

metadata.json：GPU、显存、驱动、CUDA、commit、编译参数和输入矩阵；
Google Benchmark 的原始 JSON 与完整 stdout.log；
PyTorch SDPA 或官方 FlashAttention 对照的独立命令、版本与输入契约；
nsys timeline 与 ncu 指标（若云环境未开放权限，明确记录限制，不用推测替代）；
包含失败形状、OOM、数值误差和低于噪声优化的报告。

不同 GPU 的结果分目录和分表展示；禁止把不同架构、不同精度、不同实现版本的数据拼成同一加速比。

8. 版本变更记录

版本	日期	变更说明
v0.5.1	2026-08-24	跨 GPU 历史表格隔离为不可审计快照；首页撤下未经原始产物支持的数字，新增正式结果发布门槛
v0.5.0+	2026-08-18	新增 RTX 3060 Laptop 本机测量上下文（commit `6860cbc`），补充 head_dim=128 数据
v0.4.0	2026-04-15	初始 benchmark 文档；跨 GPU 数字未保留完整原始产物，现已禁止引用

基准测试：当前本机快照与复验协议 ​

1. 本机快照与未来复验协议 ​

1.1 硬件平台 ​

1.2 软件栈 ​

1.3 测试配置 ​

1.4 编译参数 ​

1.5 本机测量快照（RTX 3060 Laptop，2026-08-18，非正式结果包） ​

2. 不可审计历史快照：多维度 Benchmark 矩阵（禁止引用） ​

2.1 端到端前向传播（Forward, FP16, Causal） ​

2.2 反向传播（Backward, dQ/dK/dV, FP16, Causal） ​

3. 不可审计历史快照：与 PyTorch SDPA 的速度对比（禁止引用） ​

3.1 绝对耗时对比 ​

3.2 结果解读 ​

4. 不可审计历史快照：内存占用对比（禁止引用） ​

4.1 峰值显存（Forward + Backward 单步） ​

4.2 内存 Scaling 公式验证 ​

5. 不可审计历史快照：不同 GPU 架构的 Scaling 分析（禁止引用） ​

5.1 理论峰值对比 ​

5.2 实测 Scaling（固定配置: batch=1, heads=8, head_dim=64, causal FP16） ​

5.3 理论 vs 实测加速比 ​

分析 ​

6. 未来正式结果的复验命令 ​

6.1 本地运行（需 CUDA GPU） ​

6.2 Docker 环境（推荐，完全可复现） ​

6.3 与 PyTorch SDPA 的对比脚本 ​

7. 正式结果发布门槛 ​

8. 版本变更记录 ​

基准测试：当前本机快照与复验协议

1. 本机快照与未来复验协议

1.1 硬件平台

1.2 软件栈

1.3 测试配置

1.4 编译参数

1.5 本机测量快照（RTX 3060 Laptop，2026-08-18，非正式结果包）

2. 不可审计历史快照：多维度 Benchmark 矩阵（禁止引用）

2.1 端到端前向传播（Forward, FP16, Causal）

2.2 反向传播（Backward, dQ/dK/dV, FP16, Causal）

3. 不可审计历史快照：与 PyTorch SDPA 的速度对比（禁止引用）

3.1 绝对耗时对比

3.2 结果解读

4. 不可审计历史快照：内存占用对比（禁止引用）

4.1 峰值显存（Forward + Backward 单步）

4.2 内存 Scaling 公式验证

5. 不可审计历史快照：不同 GPU 架构的 Scaling 分析（禁止引用）

5.1 理论峰值对比

5.2 实测 Scaling（固定配置: batch=1, heads=8, head_dim=64, causal FP16）

5.3 理论 vs 实测加速比

分析

6. 未来正式结果的复验命令

6.1 本地运行（需 CUDA GPU）

6.2 Docker 环境（推荐，完全可复现）

6.3 与 PyTorch SDPA 的对比脚本

7. 正式结果发布门槛

8. 版本变更记录