跳到主要内容
TileLang 系统教程

tilelang-tutorial.huandu.me

第 35 章 DeepGEMM-Ascend:把 GEMM 压到硬件极限

/code/35-deepgemm-ascend/

02-perf-ledger/README.md3614 字节
# 示例 02:把性能表换算成利用率与带宽

## 目的

正文 §5.1–§5.7 的复算脚本。它回答三件事:

1. README 里的「Hardware limit」列是怎么来的——用 `heuristics.hpp` 的硬件模型
   复算峰值,看是否吻合;
2. 各表的「Computation」列换成「占硬件极限的百分比」是多少;
3. 用**明确写出的字节口径**复算带宽,并与 README 的 Memory bandwidth 列对比,
   把比值打印出来(对不上就对不上,不掩饰)。

## 运行前提

Python 3.9+,**只用标准库**。不需要 numpy、GPU、NPU,也不 import torch/tilelang。

## 运行命令

```bash
python3 perf_ledger.py
```

## 期望输出(写作机实际输出,退出码 0)

节选(完整输出以实际运行为准):

```text
==============================================================================
1. 硬件极限的复算(heuristics.hpp:46-48,52-71)
==============================================================================
  峰值 = 8192 * 16/elem_bits * 1650 MHz * 32 核
  bf16:   8192 FLOP/cycle/核  ->    432.5 TFLOPS
   fp8:  16384 FLOP/cycle/核  ->    865.1 TFLOPS
   fp4:  32768 FLOP/cycle/核  ->   1730.2 TFLOPS
  与 README Hardware limit 列对照:bf16 432 / fp8 865 / fp4 1730

==============================================================================
2. Dense GEMM:利用率 = Computation / 本格式峰值
==============================================================================
  case                          comp     peak    util
  FP4xFP4                     1701.0   1730.2   98.3%
  FP8xFP4                      861.0    865.1   99.5%
  FP8xFP8                      861.0    865.1   99.5%
  BF16xBF16                    431.0    432.5   99.6%

==============================================================================
6. FP8 推理:带宽复算 vs README
==============================================================================
  口径:MB = A + B + SF(不含 D;对应 tests 里 c=None 的 count_bytes(a, b, c))。
       M      N      K        us   MB(A+B+SF)     模型GB/s       表格      表格/模型
     128   7168  16384      54.8        123.3       2250     2456     1.092
    4096  24576   1536     362.2         45.4        125      137     1.093
    4096  32768    512     164.3         19.5        118      129     1.089

==============================================================================
7. HC Prenorm:带宽复算 vs README(口径含 D,与 tests 一致)
==============================================================================
       M      N      K        us   MB(A+B+D)     模型GB/s       表格      表格/模型
      13     24  28672       6.7          3.5        522      519     0.994
    8192     24  28672     136.7        472.5       3457     3463     1.002
```

## 判据

* 程序正常退出(退出码 0);
* 第 1 节三行与 README 的 Hardware limit 列(432 / 865 / 1730)吻合到 0.2% 以内;
* 第 2 节四行与 README 的 Utilization 列一致(98.3 / 99.5 / 99.5 / 99.8);
  BF16 的 99.6% 与 99.8% 的差异来自 README 把极限写成 432(取整)而不是 432.5;
* 第 6 节 14 行的「表格/模型」比值全部落在 **1.089–1.095**;
* 第 7 节 5 行的比值全部落在 **0.994–1.002**。

## 已知差异(必须一起看,否则会误读)

* **占比利用率对得上**,因为它的分母就是源码里的硬件模型。
* **HC Prenorm 的带宽能复现**:脚本用 `A + B + D` 口径(与
  `tests/test_hyperconnection.py:103-107` 的 `count_bytes(a, b, d, sqr_sum)` 一致,
  `sqr_sum` 占比可忽略),比值 0.994–1.002。
* **FP8 推理的带宽复算不出来**:14 行的比值稳定在 1.089–1.095(约 +9%),
  说明表格那一列用的是另一种(快照里未给出公式的)口径。
  **本脚本只报告这个事实,不做解释**;引用时以表格列为准。
  这一项登记在 [`references.md`](../../references.md) 的「待核」小节。
* **M-Grouped / MegaMoE / MQA 不做带宽对账**:前者有 per-group 的 M、
  中者是多 rank 平均的融合 kernel、后者的列是 `(Q, KV, heads, dim)`——
  同一个 `(M, N, K)` 描述不了它们的全部流量,所以脚本只算它们的算力利用率。

## 怎么在真机上确认(本机做不到)

1. 用官方 `tests/` 在同形状上跑一次,打开 `bench_msprof`;
2. `DG_PRINT_CONFIGS=1` 记录选中的 `GemmConfig`;
3. 用 `KernelProfile` 看 `aic_mad` / `aic_mte2` / `aic_fixpipe` 哪条 pipe 高,
   与正文 §5 的瓶颈判断对照;
4. 想弄清 FP8 推理带宽列的那 9%,最直接的办法是抓一个 shape 的
   `count_bytes(a, b, c)` 实际返回值与 `KernelProfile.dur_ns`,和本脚本的模型对齐。

## 作者实际做到的验证级别

**CPU 逻辑验证**:在 macOS (Apple Silicon)、Python 3.11 上实际运行过,退出码 0。
**没有在任何加速器上测量过任何东西**;脚本里的所有数值都是复算或照抄 README。

← 浏览本章全部代码