vllm.models.kimi_k3.amd.specialized_decode.mono.attention.mla_attn
¶
M1c mla_attn: Kimi-K3's MLA decode attention at TP8 + DCP1, the bf16
query (M1a qbf16) against the fp8 latent cache, split over the KV: one pass
a split, every split's partial (bf16 o, f32 LSE) for a combine (M1b's).
One launch, BLOCKS CTAs x THREADS threads. A CTA a (head group, split):
16 query rows (16 / qlen heads x qlen positions) against a chunk of the
request's KV. The split's CTAs (every head group) share an XCD, so a chunk
streams from HBM once into that XCD's L2.
block of 128 tokens (wave w its tile of 16, the waves in step), its K
rows DMA'd into one of two LDS slots a block ahead (global_load_lds, the
chunks spread over the block before: ``DMA``; this wave's chunks waited
for one at a time, so the top barrier only frees the slot)
K tile LDS -> registers (a lane 32 B a line of a token: the MFMA's A, the K
order permuted alike in Q's B), S^T = K Q^T: the fp8 K as is against
Q's exact fp8 hi + lo (``q_split``), 9 16x16x128 MFMAs a row tile
the block's row max over the waves (LDS) -> P = exp2(S - m) 2^PLOG,
fp8 hi + lo -> LDS; the running O, l rescaled (online softmax)
O^T += V^T P^T: wave w its 64 latents, the fp8 V^T straight from the
slot (``ds_read_b64_tr_b8``) against P's hi and lo, 4 16x16x128 MFMAs
a row tile and part
end: l summed over the waves -> o = O / l (bf16), LSE = m + log l
A split with no token of the request (or none before its causal bound) leaves
LSE -inf and o 0. Rows are causal: position i of the request's qlen sees its
KV up to seq_len - qlen + i.
Functions:
-
block–Block b (
stage_block) behind the top barrier, its DMA hooks. -
chunk_dma–One chunk DMA (
k_dma's):src+imm(an offset off the pool, -
cta–CTA
c["bid"]'s (head group, split): its partial stored (cache -
dma_lane–(token, 16 B cc) tail DMA lane
lanecarries (it lands at position -
dma_offs–k_dma's sources ofidx_load's indices: the tail chunk's (16 B -
fp8_pack4–Four f32 (in range) -> one dword of four E4M3, RNE (
ops' but the -
fp8_split4–Four f32 (in range) -> (hi, lo) dwords of four E4M3: x ~ hi + lo (the
-
fp8_unpack4–One dword of four E4M3 -> four f32 (exact).
-
host_args–(qb, kv, pages, indptr, scale, vscale) of the kernel (
mla_attn's -
idx_dma–Block b's page indices of this wave's tile (or tiles t0 .. t0 + n) ->
-
idx_lds–This lane's DMA tokens' page indices of block b (
idx_dma's; as -
idx_load–The page indices of this lane's DMA tokens of block b's tile of this
-
k_dma–Block b's K tile of this wave (or
tile;dma_offs) -> slot -
kpos–The 16 B position (of 64) in a tile's tail chunk of token t's 16 B cc:
-
lds_dma–A global_load_lds_dwordx4: lane i's 16 B at
src-> LDS byte -
lines_at–Token t's 16 B
colof line s: its byte in the slot (from line s's -
mfma_f8–16x16x128 E4M3 x E4M3 MFMA, unscaled (scale operands 0: LLVM's
-
mla_attn–Host launcher.
q[S, 12, 576] bf16;kv[tokens, 576] fp8 rows -
q_parts–Row tile r's B operands from
q_split's: QK step s's hi, lo -
q_split–The CTA's bf16 Q rows (slot 1:
[row][ENTRY]) -> E4M3 hi, lo of q 2^k -
run–The blocks (block 0's DMA and block 1's indices issued) and the
-
sig–The 16 B column swizzle (XOR) of token t (of 16) of a tile of parity
-
smem–The LDS of a CTA (
cta's ks / xm / xp / xi; xp P's hi | lo, then -
stage_block–One block (its K in slot b % 2) into the running (m, l, O) of each of
-
tr_read–ds_read_b64_tr_b8 at LDS byte
addr: a lane m of 16 byte m % 8 of -
v_reads–V^T of the slot's block b by transpose reads, the A of a 16x16x128 fp8
-
vmwait–s_waitcnt vmcnt(n): at most n of this wave's loads / DMAs in flight
-
vtr_lane–This lane's byte in a block's (line wave // 2) of the 8 B it gives a
-
wait_dma–This wave's DMAs (and loads) landed; the CTA's after the barrier.
-
wide_kv–The build's
widefor KV poolkv: a DMA's 32-bit offset off the -
ws_dmas–This wave's DMA hook of block bb's tile's chunks
ps(k_dma's), -
ws_l–Row
rowof row tile r's l: the QK waves' (xm) summed. -
ws_offs–This lane's DMA source of tile t of block b (
dma_offs, its index -
ws_pv–The PV waves' steps: step 0 V^T(0) and block 1's DMA, then
-
ws_pv_block–PV wave step b: O += V^T P^T of block b - 1 (its V^T
varead in -
ws_qk–The QK waves' steps (
ws_qk_block), the last block's PV's barriers, -
ws_qk_block–QK wave step b: S of block b's tiles 2 w, 2 w + 1 (
stage_block's -
ws_run–Warp-specialized blocks (
WSPEC): the QK waves', the PV waves'. -
ws_vtile–V^T latent tile i (of 8: 128 v + 16 i ..) of the slot's block b
-
xm_row–The waves' values of row
rowof row tile r in xm -
xp_at–The P dword of (row tile r, hi / lo h, tiles 4 k ..) of
lane.
block(c, b, st, masked)
¶
Block b (stage_block) behind the top barrier, its DMA hooks.
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
chunk_dma(c, src, imm, dst)
¶
One chunk DMA (k_dma's): src + imm (an offset off the pool,
or wide an address) -> LDS byte dst (uniform; lane i its 16 B at
+ 16 i).
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
cta(c)
¶
CTA c["bid"]'s (head group, split): its partial stored (cache
modifier c["cm"]), then c["done"](flag) when given (ctas).
c["args"]: qb, kv, pages, indptr, part_o, part_lse, scale, vscale.
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
1175 1176 1177 1178 1179 1180 1181 1182 1183 1184 1185 1186 1187 1188 1189 1190 1191 1192 1193 1194 1195 1196 1197 1198 1199 1200 1201 1202 1203 1204 1205 1206 1207 1208 1209 1210 1211 1212 1213 1214 1215 1216 1217 1218 1219 1220 1221 1222 1223 1224 1225 1226 1227 1228 1229 1230 1231 1232 1233 1234 1235 1236 1237 1238 1239 1240 1241 1242 1243 1244 1245 1246 1247 1248 1249 1250 1251 1252 1253 1254 1255 1256 1257 1258 1259 1260 1261 1262 1263 1264 1265 1266 1267 1268 1269 1270 1271 1272 1273 1274 1275 1276 1277 1278 1279 1280 1281 1282 1283 1284 1285 1286 | |
dma_lane(lane)
¶
(token, 16 B cc) tail DMA lane lane carries (it lands at position
lane).
dma_offs(c, idx)
¶
k_dma's sources of idx_load's indices: the tail chunk's (16 B
dcc), line chunks' tokens lane // 8 and + 8 (their 16 B column
sig-swizzled); bytes off the pool (wide: addresses).
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
fp8_pack4(a, b, c, d)
¶
Four f32 (in range) -> one dword of four E4M3, RNE (ops' but the
first pair over a poison dword: no zeroing move).
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
fp8_split4(x)
¶
Four f32 (in range) -> (hi, lo) dwords of four E4M3: x ~ hi + lo (the differences as pairs, as the unpack gives them: no shuffling moves).
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
fp8_unpack4(w)
¶
One dword of four E4M3 -> four f32 (exact).
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
host_args(*, q, kv, pages, indptr, scale, kv_scale=1.0)
¶
(qb, kv, pages, indptr, scale, vscale) of the kernel (mla_attn's
inputs but the partials; K2's fused M1c's).
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
idx_dma(c, b, t0=None, n=1)
¶
Block b's page indices of this wave's tile (or tiles t0 .. t0 + n) -> its LDS index slot b % 2 (lanes 0 .. 16 n, token lane's): a DMA, so it lands in order with the K DMAs (a load would make LLVM wait on every DMA before it to read it).
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
idx_lds(c, b)
¶
This lane's DMA tokens' page indices of block b (idx_dma's; as
idx_load).
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
idx_load(c, b)
¶
The page indices of this lane's DMA tokens of block b's tile of this wave (the tail's, lines' tokens lane // 8, + 8), clamped into the chunk (a block past the last loads its rows).
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
k_dma(c, b, offs, ps=None, tile=None)
¶
Block b's K tile of this wave (or tile; dma_offs) -> slot
b % 2: a DMA a chunk (64 lanes x 16 B: line s's 8 tokens, the tail's 16).
Inline asm: LLVM's own wait for an LDS DMA would hold this block's slot
reads for the next block's DMA too; vmwait / wait_dma are the only
waits.
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
kpos(t, cc)
¶
The 16 B position (of 64) in a tile's tail chunk of token t's 16 B cc: a lane group's K read (tokens 0..15, one cc) over the 64 banks.
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
lds_dma(src, dst)
¶
A global_load_lds_dwordx4: lane i's 16 B at src -> LDS byte
dst (uniform) + 16 i. (Every LDS DMA here an s_nop between its M0
write and the load: right after the write the load could read a stale M0,
the warp-specialized waves' K / index DMAs then landing in another slot:
nondeterministic outputs.)
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
lines_at(t, tp, col)
¶
Token t's 16 B col of line s: its byte in the slot (from line s's
chunk 2 s, a tile 1 KB: 8 tokens x 128 B), the DMA lane 8 (t % 8) +
(col ^ sig) of chunk 2 s + t // 8.
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
mfma_f8(a, b, c)
¶
16x16x128 E4M3 x E4M3 MFMA, unscaled (scale operands 0: LLVM's
v_mfma_f32_16x16x128_f8f6f4): a / b 8 dwords (QK and PV).
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
mla_attn(key, *, q, kv, pages, indptr, part_o, part_lse, scale, kv_scale=1.0)
¶
Host launcher. q [S, 12, 576] bf16; kv [tokens, 576] fp8 rows
(dense); pages int32 flat token indices, indptr [requests + 1]
int32 (a request's KV pages[indptr[r]:indptr[r + 1]]); part_o /
part_lse of partial_shapes(key), bf16 / f32. scale the softmax
scale, kv_scale the cache's dequant (folded into QK and o).
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
q_parts(c, r, at, rows)
¶
Row tile r's B operands from q_split's: QK step s's hi, lo
(s < QSTEPS: line s), 8 dwords each (a lane 32 B at 32 g of the line,
K's order); the tail's stay in LDS (qt).
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
q_split(c, rows, at)
¶
The CTA's bf16 Q rows (slot 1: [row][ENTRY]) -> E4M3 hi, lo of q 2^k
(a power of two a row: its |max| in [128, 256)): hi + lo is q exactly,
bf16's 8 bits E4M3's 4 + 4 (lo q - hi, a multiple of hi's ulp / 16; but
below 2^-9 of 2^k: 2^-17 of the row's |max|). At dword at:
[hi | lo][row][ENTRY] bytes, the tail chunk's also at qt; a row's
2^-k at xm + row. A row 16 threads, 4 latents a unit (tid % 16 +
16 i).
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
run(c)
¶
The blocks (block 0's DMA and block 1's indices issued) and the partial's stores.
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
sig(t, tp)
¶
The 16 B column swizzle (XOR) of token t (of 16) of a tile of parity tp in a line chunk: a K read (tokens 0..15, one column) and a V^T read (tokens 4 g + j, tiles 2 n, 2 n + 1) both over the 64 banks.
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
smem(qlen, rt=0, ws=-1)
¶
The LDS of a CTA (cta's ks / xm / xp / xi; xp P's hi | lo, then
the tail Q's tail hi | lo, qt).
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
stage_block(c, b, st, dma, kwait, masked)
¶
One block (its K in slot b % 2) into the running (m, l, O) of each of
the CTA's row tiles (K and V^T read once for all); dma(hook)
issues the hook's of block b + 1's K chunks (DMA), kwait(p) waits
for this wave's chunk p of block b; masked: a block with a token past
the chunk or a row's causal bound (else every score kept).
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
524 525 526 527 528 529 530 531 532 533 534 535 536 537 538 539 540 541 542 543 544 545 546 547 548 549 550 551 552 553 554 555 556 557 558 559 560 561 562 563 564 565 566 567 568 569 570 571 572 573 574 575 576 577 578 579 580 581 582 583 584 585 586 587 588 589 590 591 592 593 594 595 596 597 598 599 600 601 602 603 604 605 606 607 608 609 610 611 612 613 614 615 616 617 618 619 620 621 622 623 624 625 626 627 628 629 630 631 632 633 634 635 636 637 638 639 640 641 642 643 644 645 646 647 648 649 650 651 652 653 654 655 656 657 658 659 660 661 662 663 664 665 666 667 668 669 670 671 672 673 674 675 676 677 678 679 680 681 682 683 | |
tr_read(addr)
¶
ds_read_b64_tr_b8 at LDS byte addr: a lane m of 16 byte m % 8 of
the 8 B of lanes m // 8 + 2 e (byte e), as two dwords.
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
v_reads(c, b, dma)
¶
V^T of the slot's block b by transpose reads, the A of a 16x16x128 fp8
MFMA a latent tile i: a lane's 32 B tile w's tokens 4 g .. at 4 w (P's
order); dma the hook v<i>'s K chunks before tile i's reads.
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
vmwait(n)
¶
s_waitcnt vmcnt(n): at most n of this wave's loads / DMAs in flight (they land in order).
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
vtr_lane(lane, wave, i, lines=True)
¶
This lane's byte in a block's (line wave // 2) of the 8 B it gives a
V^T transpose read of latent tile i (tr_read; tiles 2 n .. at + 2 n
TILEW): lane k of a 16 its half k % 2 of the 16 B 4 (wave % 2) + i of
token 4 g + k // 2 % 4 of tile k // 8 -> lane m: latent 64 wave + 16 i +
m, tokens 4 g .. of tiles 2 n, 2 n + 1.
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
wait_dma()
¶
This wave's DMAs (and loads) landed; the CTA's after the barrier.
wide_kv(kv)
¶
The build's wide for KV pool kv: a DMA's 32-bit offset off the
pool's base (a token row's end at most) won't reach past 4 GiB.
ws_dmas(c, bb, ps, hooks=('all',))
¶
This wave's DMA hook of block bb's tile's chunks ps (k_dma's),
spread over hooks (the first ones one more).
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
ws_l(c, r, row)
¶
Row row of row tile r's l: the QK waves' (xm) summed.
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
ws_offs(c, b, t)
¶
This lane's DMA source of tile t of block b (dma_offs, its index
from the index slot).
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
ws_pv(c)
¶
The PV waves' steps: step 0 V^T(0) and block 1's DMA, then
ws_pv_block a block (one behind the QK waves), then o = O / l.
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
ws_pv_block(c, b, st)
¶
PV wave step b: O += V^T P^T of block b - 1 (its V^T va read in
step b - 1, P from xp, alpha from xm), B1 after WS_B1
latent tiles, V^T(b) read before it (block b's slot then free); after B1
block b + 2's chunks DMA'd (block b + 1's landed by B2).
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
ws_qk(c)
¶
The QK waves' steps (ws_qk_block), the last block's PV's barriers,
then the l over the waves and the LSE.
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
ws_qk_block(c, b, st, masked)
¶
QK wave step b: S of block b's tiles 2 w, 2 w + 1 (stage_block's
QK), the max over the QK waves (B1), block b + 2's tile DMA'd, P(b) -> xp,
alpha -> xm (B2: block b + 1's landed).
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
849 850 851 852 853 854 855 856 857 858 859 860 861 862 863 864 865 866 867 868 869 870 871 872 873 874 875 876 877 878 879 880 881 882 883 884 885 886 887 888 889 890 891 892 893 894 895 896 897 898 899 900 901 902 903 904 905 906 907 908 909 910 911 912 913 914 915 916 917 918 919 920 921 922 923 924 925 926 927 928 929 930 931 932 933 934 935 936 937 938 939 940 941 942 943 944 945 946 947 948 949 950 951 952 953 954 955 956 957 958 959 960 961 962 963 964 965 966 967 | |
ws_run(c)
¶
Warp-specialized blocks (WSPEC): the QK waves', the PV waves'.
ws_vtile(c, b, i)
¶
V^T latent tile i (of 8: 128 v + 16 i ..) of the slot's block b
(v_reads'): 8 dwords.
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
xm_row(c, r, row)
¶
The waves' values of row row of row tile r in xm
([r][row][wave]: two b128 reads).
Source code in vllm/models/kimi_k3/amd/specialized_decode/mono/attention/mla_attn.py
xp_at(r, h, k, lane)
¶
The P dword of (row tile r, hi / lo h, tiles 4 k ..) of lane.