Skip to content

Commit 4c237b4

Browse files
committed
Measure the shipped default path in the VAE conv benchmark.
The *_def columns passed DEFAULT_TILE explicitly, so they bypassed _pick_tile and no longer described what a caller without a tile argument actually gets. Drop the tile argument, record the tile that _pick_tile returned in a new tile_def column, and let the best-of-candidates start from that pick so the figure's outer bar stays an upper bound. The plotter derives its footnote ratios from the data instead of carrying the previous run's constants. Refreshed fair_baseline.json and the figure on the current tree. Weighted over the VAE call mix, the default path now sits at 1.068x of hipBLASLt on the 1024 path (10.35 ms of GEMM, was 11.51) and 1.284x on the two 1328 layers (7.12 ms, was 7.99). One regression the refresh exposes, unrelated to the tile picker: the 256x256x4x4 candidate has become 2.3-2.9x slower than 128x128x2x4 (e.g. 192->192 @512x512, 290 -> 835 us for the whole path) though it still compiles and matches torch. It used to win 7 of the 18 shapes, so the best-of-candidates upper bound drops on the 384->192 family (105.9 -> 132.1 us of GEMM on 384->192 @256x256).
1 parent 01dbd72 commit 4c237b4

4 files changed

Lines changed: 240 additions & 208 deletions

File tree

tests/perf/qwenimage_vae_conv/bench_fair.py

Lines changed: 16 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -7,6 +7,10 @@
77
Each of the 18 shapes runs in a fresh process so JIT / autotune state cannot
88
leak. Writes ``fair_baseline.json`` next to this script.
99
10+
The ``*_def`` columns are the shipped default path (no ``tile`` argument, so
11+
``_pick_tile`` decides); ``*_best`` is the fastest of that pick and the tile
12+
candidates below.
13+
1014
Usage::
1115
1216
python tests/perf/qwenimage_vae_conv/bench_fair.py
@@ -60,6 +64,7 @@
6064
("fly_def", float),
6165
("gemm_def", float),
6266
("tr_def", float),
67+
("tile_def", str),
6368
("fly_best", float),
6469
("gemm_best", float),
6570
("tr_best", float),
@@ -126,7 +131,7 @@ def run_one(sid, cin, cout, hin, stride, padding, freq, path) -> None:
126131
import torch.nn.functional as F
127132
from torch.profiler import ProfilerActivity, profile
128133

129-
from kernels.conv.conv3d_implicit import DEFAULT_TILE, conv3d_implicit
134+
from kernels.conv.conv3d_implicit import DEFAULT_TILE, _pick_tile, conv3d_implicit
130135

131136
cin, cout, hin, stride, padding, freq = map(int, (cin, cout, hin, stride, padding, freq))
132137
candidates = [
@@ -158,11 +163,13 @@ def fly(tile):
158163
def time(call):
159164
return gpu(torch, profile, ProfilerActivity, call)
160165

161-
check(fly(tuple(DEFAULT_TILE)))
162-
t_def, per_def = time(lambda: fly(tuple(DEFAULT_TILE)))
166+
# "default" is the shipped path: no tile argument, so _pick_tile decides.
167+
check(fly(None))
168+
t_def, per_def = time(lambda: fly(None))
163169
g_def, tr_def = split_fly(per_def)
164-
best = (g_def, t_def, tr_def, tuple(DEFAULT_TILE))
165-
for tile in candidates[1:]:
170+
tile_def = _pick_tile(M, N, 1, x4.device)
171+
best = (g_def, t_def, tr_def, tile_def)
172+
for tile in candidates:
166173
try:
167174
check(fly(tile))
168175
t, per = time(lambda: fly(tile))
@@ -207,7 +214,7 @@ def unfold_mm(wt=w2):
207214
t_mio, _ = time(lambda: F.conv2d(x4, w4, bias=bbf, stride=stride, padding=padding))
208215
print(
209216
f"RESULT\t{sid}\t{cin}\t{cout}\t{hin}\t{stride}\t{freq}\t{path}\t{M}\t{N}\t{K}\t{flops:.6e}\t"
210-
f"{t_def:.2f}\t{g_def:.2f}\t{tr_def:.2f}\t"
217+
f"{t_def:.2f}\t{g_def:.2f}\t{tr_def:.2f}\t{'x'.join(str(v) for v in tile_def)}\t"
211218
f"{t_best:.2f}\t{g_best:.2f}\t{tr_best:.2f}\t{'x'.join(str(v) for v in tile_best)}\t"
212219
f"{t_mm:.2f}\t{t_im2col:.2f}\t{t_unfold_mm:.2f}\t{t_mio:.2f}"
213220
)
@@ -219,7 +226,8 @@ def drive() -> None:
219226
rows = []
220227
print(
221228
f"{'shape':18s} {'x':>3s} {'M':>8s} {'K':>5s} | {'hipBLASLt':>9s} {'FlyGEMM':>8s} {'比':>5s} "
222-
f"{'BLASt T/s':>9s} {'Fly T/s':>8s} | {'im2col':>7s} {'unfold+mm':>10s} {'Fly 全':>7s} {'MIOpen':>8s} {'最优tile':>13s}"
229+
f"{'BLASt T/s':>9s} {'Fly T/s':>8s} | {'im2col':>7s} {'unfold+mm':>10s} {'Fly 全':>7s} {'MIOpen':>8s} "
230+
f"{'默认tile':>13s} {'最优tile':>13s}"
223231
)
224232
for sp in SHAPES:
225233
proc = subprocess.run(
@@ -240,7 +248,7 @@ def drive() -> None:
240248
f"{shp:18s} {r['freq']:3d} {r['M']:8d} {r['K']:5d} | {r['mm']:9.1f} {r['gemm_best']:8.1f} "
241249
f"{r['mm'] / r['gemm_best']:4.2f}x {r['flops'] / r['mm'] / 1e6:9.0f} "
242250
f"{r['flops'] / r['gemm_best'] / 1e6:8.0f} | {r['im2col']:7.1f} {r['unfold_mm']:10.1f} "
243-
f"{r['fly_best']:7.1f} {r['mio']:8.1f} {r['tile']:>13s}",
251+
f"{r['fly_best']:7.1f} {r['mio']:8.1f} {r['tile_def']:>13s} {r['tile']:>13s}",
244252
flush=True,
245253
)
246254

0 commit comments

Comments
 (0)