跳到主要内容
TileLang 系统教程

tilelang-tutorial.huandu.me

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

/code/35-deepgemm-ascend/

03-config-search/README.md2328 字节
# 示例 03:候选与容量模型(教学版)

## 目的

正文 §4.4 说「tile 尺寸不是随便取的,四个容量约束决定它」。这个脚本把其中**可以
逐行对上的整数部分**抽出来跑一遍:

* L0C 装得下 `block_m × block_n`(`heuristics.hpp:100-102`);
* MAD 块还要再切一半给双缓冲(`:125`);
* `mad_k` 由 L0A/L0B 容量反推,并对齐到 `k_alignment`(`:131-140`);
* L1 里 A/B/SF 能放几级、event 预算怎么压它(`:164-172`)。

**它不是 `select_gemm_config` 的移植**:代价模型与那几个 refine 补丁
(`direct_store` / `resident_*` / `dual_aiv_dequant`)不在里面。这一半是「物理约束」,
另一半是「性能取舍」,后者要看带宽、频率与 event 预算——与其抄一份没法验证的副本,
不如直接读源码(正文 §4.4 就是干这个的)。

## 运行前提

Python 3.9+,**只用标准库**。

## 运行命令

```bash
python3 config_search.py
```

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

```text
L0C 容量 = 256 KiB  L1 = 512 KiB  L0A/L0B = 64 KiB

==============================================================================
bf16(block_k=128):合法候选(block_m=512/256 各列前 3 行)
==============================================================================
   blk_m blk_n  mad_m  mad_n  mad_k  l1_stg  sf_stg   L0C 占用
     512   128    256    128     64       3       0    256 KiB  ← 占满 L0C
     512   112    256    112     64       3       0    224 KiB
     512    96    256     96     64       3       0    192 KiB
     256   256    256    128     64       4       0    256 KiB  ← 占满 L0C
     256   240    128    240     64       4       0    240 KiB
     256   224    256    112     64       4       0    224 KiB
==============================================================================
结论(复算)
==============================================================================
  * block_m=256 时 block_n 最多只能是 256:256*256*4 = 262144 B = 恰好 L0C_SIZE_BYTES。
  * 想让 block_n 更大,只能把 block_m 减半——这是「tile 尺寸由片上容量决定」最直接的一例。
  * mad_k 由 L0A/L0B 容量与 num_l0_stages=2 反推,并向下对齐到 k_alignment(bf16=32,带 SF=64)。
  * 带 SF 时 num_l1_stages 被 event 预算压到 <= 4(num_pipe_events=8 减去 SF 的 2 级)。
```

## 判据

* 程序正常退出(退出码 0);
* `bf16` 那一节必须出现 `256 256 ... 256 KiB ← 占满 L0C` 这一行;
* 每种 dtype 的 `mad_k` 都能整除 `block_k`(脚本里就是靠 `while block_k % cand_k`
  达成这一点的)。

## 已知局限

* 这是**容量模型**,不是**性能模型**:它不会告诉你 `256×256×512` 与
  `128×256×512` 哪个更快——那要 `get_gemm_cost`(`heuristics.hpp:252-465`)。
* `sf_k_blocks` 在这里被简化成最小值;真实代码会在 `:174-181` 再按 L1 剩余容量
  放宽它。
* 硬件常数取的是 DeepGEMM-Ascend 自带的 Ascend 950 口径
  (`common.hpp:50-64`、`heuristics.hpp:52-71`),与第 19 章引用的 Atlas A2 规格不同。

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

**CPU 逻辑验证**:在 macOS (Apple Silicon)、Python 3.11 上实际运行过,退出码 0。
**未编译、未运行任何 kernel**。

← 浏览本章全部代码