pid_m, pid_n → output tile
BM
2
4
8
BN
2
4
8
row 0
axis 0 ↓
col 0
axis 1 →
pid_m = tl.program_id(axis=0)
pid_n = tl.program_id(axis=1)