v1.5: two images per attention producer pass (LNC=1, <=256 tokens) + fix for <=128-token inputs
Browse filesattention_block_l1_ln packs two images into one Q/K/V projection + RoPE pass (pack2, default on for even batches
at 224 px): per instance ViT-L/16 +3.7%, ViT-H+/16 +10.6%, ViT-B/16 +7.9%, output bit-identical to v1.4.
Inputs of 128 tokens or fewer now use a single-Q-group path (previously failed to compile).
Experimental opt-ins: bf16 RoPE (VIT_ENC_ROPE_BF16=1), fused attention write-back (VIT_ENC_WB_FUSED=1).
Packing technique from the BERT-reranker fused attention block.
- CHANGELOG.md +18 -0
- README.md +5 -0
- build/torch-neuron/encoder.py +7 -2
- build/torch-neuron/nki_kernels/attention_block.py +7 -4
- build/torch-neuron/nki_kernels/attention_block_l1.py +7 -4
- build/torch-neuron/nki_kernels/attention_block_l1_ln.py +77 -36
- build/torch-neuron/nki_kernels/attention_inner.py +45 -23
CHANGELOG.md
CHANGED
|
@@ -1,5 +1,23 @@
|
|
| 1 |
# Changelog
|
| 2 |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 3 |
## v1.4 (2026-10)
|
| 4 |
|
| 5 |
* **QK^T PSUM banks:** for 1024 or fewer tokens per image (two key tiles or fewer), the QK^T matmul now rotates over
|
|
|
|
| 1 |
# Changelog
|
| 2 |
|
| 3 |
+
## v1.5 (2026-10)
|
| 4 |
+
|
| 5 |
+
* **Two images per producer pass** (LNC=1 all-group attention, `attention_block_l1_ln`, new `pack2` argument, on by
|
| 6 |
+
default from the encoder when the batch is even and two images fit one 512-token tile, i.e. up to 256 tokens such as
|
| 7 |
+
224 px). The two images' Q/K/V projections and RoPE run as one pass over 2N columns; attention stays per image.
|
| 8 |
+
Kernel device cycles, ViT-L 224 px, batch 4: -6.2% (resident weights) / -14.8% (restaged), batch 2: -8.5%.
|
| 9 |
+
End to end (real HF weights, LNC=1): per instance (8 cores, batch 8) ViT-L/16 **1,838 -> 1,906 img/s (+3.7%)**,
|
| 10 |
+
ViT-H+/16 **672 -> 743 (+10.6%)**, ViT-B/16 **4,390 -> 4,738 (+7.9%)**. Single core batch 4: ViT-L +4.6%, ViT-H+
|
| 11 |
+
+12.6%, ViT-B +4.1%. Output bit-identical to v1.4 (same rel_l2 vs fp32 on all 12 cases). `VIT_ENC_PACK2=0` disables it.
|
| 12 |
+
Technique from the BERT-reranker fused attention block.
|
| 13 |
+
* **Fix: inputs of 128 tokens or fewer** failed to compile in every attention kernel (`index 1 out of bounds`). They
|
| 14 |
+
now use a non-pipelined single-Q-group path (rel_l2 vs a float64 reference 1.9-2.5e-3). Longer inputs are
|
| 15 |
+
unchanged (bit-identical).
|
| 16 |
+
* **EXPERIMENTAL bf16 RoPE** (`VIT_ENC_ROPE_BF16=1`, off by default): RoPE tables and products in bf16. +0.6% to +3.3%
|
| 17 |
+
end to end on DINOv3; rel_l2 vs fp32 grows by up to 1.16x (V-JEPA2: 0.039 -> 0.041).
|
| 18 |
+
* **EXPERIMENTAL fused attention write-back** (`VIT_ENC_WB_FUSED=1`, off by default): one Vector multiply normalises and
|
| 19 |
+
casts O. -3.4% at 256 tokens / LNC=2, but +6.6% at 2048 tokens (Vector is the attention bottleneck), so off.
|
| 20 |
+
|
| 21 |
## v1.4 (2026-10)
|
| 22 |
|
| 23 |
* **QK^T PSUM banks:** for 1024 or fewer tokens per image (two key tiles or fewer), the QK^T matmul now rotates over
|
README.md
CHANGED
|
@@ -128,6 +128,11 @@ Notes:
|
|
| 128 |
* On a vLLM-Neuron venv (CUDA-tagged torch), `import libtorch_neuronx_lite` before `get_kernel`
|
| 129 |
so the `torch-neuron` variant is selected. On PyTorch Native this is not needed.
|
| 130 |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 131 |
**Experimental FP8 MLP** (v1.4, GELU-MLP models only): set `VIT_ENC_MLP_FP8=1` before importing to run the MLP
|
| 132 |
matmuls in FP8 e4m3 (`double_row`). That is +7-10% end to end on DINOv3 ViT-L/16, at ~7x the feature error of the
|
| 133 |
default bf16 path (CLS rel_l2 5.8e-2 vs 0.84e-2; imagenette k-NN predictions unchanged). Calibrate the per-layer
|
|
|
|
| 128 |
* On a vLLM-Neuron venv (CUDA-tagged torch), `import libtorch_neuronx_lite` before `get_kernel`
|
| 129 |
so the `torch-neuron` variant is selected. On PyTorch Native this is not needed.
|
| 130 |
|
| 131 |
+
**Environment switches** (defaults are the measured best): `VIT_ENC_PACK2=1` packs two images per attention producer
|
| 132 |
+
pass at LNC=1 when the batch is even and the images are <= 256 tokens (224 px); `VIT_ENC_ROPE_BF16=1` (experimental)
|
| 133 |
+
runs RoPE in bf16; `VIT_ENC_WB_FUSED=1` (experimental) fuses the attention normalisation; `VIT_ENC_MM1_WIDE=0`
|
| 134 |
+
restores the v1.3 QK^T PSUM layout.
|
| 135 |
+
|
| 136 |
**Experimental FP8 MLP** (v1.4, GELU-MLP models only): set `VIT_ENC_MLP_FP8=1` before importing to run the MLP
|
| 137 |
matmuls in FP8 e4m3 (`double_row`). That is +7-10% end to end on DINOv3 ViT-L/16, at ~7x the feature error of the
|
| 138 |
default bf16 path (CLS rel_l2 5.8e-2 vs 0.84e-2; imagenette k-NN predictions unchanged). Calibrate the per-layer
|
build/torch-neuron/encoder.py
CHANGED
|
@@ -28,6 +28,8 @@ from typing import Optional
|
|
| 28 |
import torch
|
| 29 |
|
| 30 |
import importlib as _il
|
|
|
|
|
|
|
| 31 |
|
| 32 |
_ab = _il.import_module(__package__ + ".nki_kernels.attention_block")
|
| 33 |
_ab1 = _il.import_module(__package__ + ".nki_kernels.attention_block_l1")
|
|
@@ -165,7 +167,10 @@ def rope_tables(cos: torch.Tensor, sin: torch.Tensor, n_prefix: int, n_pad: int,
|
|
| 165 |
S[n_prefix:n_prefix + n_rot] = sin.float().cpu()
|
| 166 |
gsz = 512 // D
|
| 167 |
nh_p = ((num_heads + gsz - 1) // gsz) * gsz
|
| 168 |
-
|
|
|
|
|
|
|
|
|
|
| 169 |
|
| 170 |
|
| 171 |
def calibrate_fp8(runner, gelu_amax, headroom: float = 2.0):
|
|
@@ -248,7 +253,7 @@ class EncoderRunner:
|
|
| 248 |
if self._all_group(N):
|
| 249 |
x, z2 = _ab1ln.attention_block_l1_ln(z1, p["wqk"], p["bqk"], tab, p["wvt"], p["bv"], p["wot"],
|
| 250 |
p["bo_full"], x, p["scale"], self.D, n_valid, p["eps2"], B,
|
| 251 |
-
self._res(_ab1ln, N))
|
| 252 |
else:
|
| 253 |
parts = self._attn_parts(z1, p, tab, n_valid, B)
|
| 254 |
x, z2 = nki_res_layernorm_fold(x, parts, p["eps2"])
|
|
|
|
| 28 |
import torch
|
| 29 |
|
| 30 |
import importlib as _il
|
| 31 |
+
import os as _osr
|
| 32 |
+
_PACK2 = _osr.environ.get("VIT_ENC_PACK2", "1") == "1" # 2 images per producer pass (N <= 256, LNC=1)
|
| 33 |
|
| 34 |
_ab = _il.import_module(__package__ + ".nki_kernels.attention_block")
|
| 35 |
_ab1 = _il.import_module(__package__ + ".nki_kernels.attention_block_l1")
|
|
|
|
| 167 |
S[n_prefix:n_prefix + n_rot] = sin.float().cpu()
|
| 168 |
gsz = 512 // D
|
| 169 |
nh_p = ((num_heads + gsz - 1) // gsz) * gsz
|
| 170 |
+
t = _ab.pack_tables(C, S, nh_p, D, rope)
|
| 171 |
+
if _osr.environ.get("VIT_ENC_ROPE_BF16", "0") == "1": # EXPERIMENTAL bf16 RoPE (tables in bf16)
|
| 172 |
+
t = t.to(torch.bfloat16)
|
| 173 |
+
return t.to(device)
|
| 174 |
|
| 175 |
|
| 176 |
def calibrate_fp8(runner, gelu_amax, headroom: float = 2.0):
|
|
|
|
| 253 |
if self._all_group(N):
|
| 254 |
x, z2 = _ab1ln.attention_block_l1_ln(z1, p["wqk"], p["bqk"], tab, p["wvt"], p["bv"], p["wot"],
|
| 255 |
p["bo_full"], x, p["scale"], self.D, n_valid, p["eps2"], B,
|
| 256 |
+
self._res(_ab1ln, N), int(_PACK2 and B % 2 == 0 and 2 * N <= 512))
|
| 257 |
else:
|
| 258 |
parts = self._attn_parts(z1, p, tab, n_valid, B)
|
| 259 |
x, z2 = nki_res_layernorm_fold(x, parts, p["eps2"])
|
build/torch-neuron/nki_kernels/attention_block.py
CHANGED
|
@@ -103,6 +103,9 @@ from torch_neuronx import nki_op
|
|
| 103 |
import nki
|
| 104 |
import nki.isa as nisa
|
| 105 |
import nki.language as nl
|
|
|
|
|
|
|
|
|
|
| 106 |
from nki.isa import engine
|
| 107 |
|
| 108 |
from nkilib.core.utils.modular_allocator import ModularAllocator
|
|
@@ -367,7 +370,7 @@ def _attention_block_kernel(h_hbm, wqk_hbm, bqk_hbm, tab_hbm, wvt_hbm, bv_hbm,
|
|
| 367 |
dge_mode=nisa.dge_mode.hwdge, engine=nisa.engine.scalar)
|
| 368 |
r_bv_bc = ra.alloc_sbuf_tensor(shape=(_V_TILE, local_heads * D), dtype=nl.float32, align_to=4)
|
| 369 |
nisa.dma_copy(dst=r_bv_bc, src=bv_hbm.ap(pattern=[[0, _V_TILE], [1, local_heads * D]], offset=vcol0))
|
| 370 |
-
r_tab = ra.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=
|
| 371 |
block_dim=[4, n_tt], num_free_tiles=[4, n_tt], align_to=4)
|
| 372 |
for s in range(n_tt):
|
| 373 |
for t_i in range(4):
|
|
@@ -504,7 +507,7 @@ def _attention_block_kernel(h_hbm, wqk_hbm, bqk_hbm, tab_hbm, wvt_hbm, bv_hbm,
|
|
| 504 |
tab = r_tab
|
| 505 |
else:
|
| 506 |
tab = allocator.alloc_sbuf_tensor(
|
| 507 |
-
shape=(_P_MAX, _PSUM_FMAX), dtype=
|
| 508 |
block_dim=[4, n_tt], num_free_tiles=[4, XT_RING], align_to=4)
|
| 509 |
|
| 510 |
if not _big:
|
|
@@ -519,9 +522,9 @@ def _attention_block_kernel(h_hbm, wqk_hbm, bqk_hbm, tab_hbm, wvt_hbm, bv_hbm,
|
|
| 519 |
proj = allocator.alloc_sbuf_tensor(
|
| 520 |
shape=(_P_MAX, _PSUM_FMAX), dtype=nl.bfloat16, block_dim=[2],
|
| 521 |
num_free_tiles=[2])
|
| 522 |
-
acc = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=
|
| 523 |
align_to=4)
|
| 524 |
-
tmp = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=
|
| 525 |
align_to=4)
|
| 526 |
roped = allocator.alloc_sbuf_tensor(
|
| 527 |
shape=(_P_MAX, _PSUM_FMAX), dtype=nl.bfloat16, block_dim=[2],
|
|
|
|
| 103 |
import nki
|
| 104 |
import nki.isa as nisa
|
| 105 |
import nki.language as nl
|
| 106 |
+
import os as _osr
|
| 107 |
+
_ROPE_BF16 = _osr.environ.get("VIT_ENC_ROPE_BF16", "0") == "1"
|
| 108 |
+
_ROPE_DT = nl.bfloat16 if _ROPE_BF16 else nl.float32
|
| 109 |
from nki.isa import engine
|
| 110 |
|
| 111 |
from nkilib.core.utils.modular_allocator import ModularAllocator
|
|
|
|
| 370 |
dge_mode=nisa.dge_mode.hwdge, engine=nisa.engine.scalar)
|
| 371 |
r_bv_bc = ra.alloc_sbuf_tensor(shape=(_V_TILE, local_heads * D), dtype=nl.float32, align_to=4)
|
| 372 |
nisa.dma_copy(dst=r_bv_bc, src=bv_hbm.ap(pattern=[[0, _V_TILE], [1, local_heads * D]], offset=vcol0))
|
| 373 |
+
r_tab = ra.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=_ROPE_DT,
|
| 374 |
block_dim=[4, n_tt], num_free_tiles=[4, n_tt], align_to=4)
|
| 375 |
for s in range(n_tt):
|
| 376 |
for t_i in range(4):
|
|
|
|
| 507 |
tab = r_tab
|
| 508 |
else:
|
| 509 |
tab = allocator.alloc_sbuf_tensor(
|
| 510 |
+
shape=(_P_MAX, _PSUM_FMAX), dtype=_ROPE_DT,
|
| 511 |
block_dim=[4, n_tt], num_free_tiles=[4, XT_RING], align_to=4)
|
| 512 |
|
| 513 |
if not _big:
|
|
|
|
| 522 |
proj = allocator.alloc_sbuf_tensor(
|
| 523 |
shape=(_P_MAX, _PSUM_FMAX), dtype=nl.bfloat16, block_dim=[2],
|
| 524 |
num_free_tiles=[2])
|
| 525 |
+
acc = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=_ROPE_DT,
|
| 526 |
align_to=4)
|
| 527 |
+
tmp = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=_ROPE_DT,
|
| 528 |
align_to=4)
|
| 529 |
roped = allocator.alloc_sbuf_tensor(
|
| 530 |
shape=(_P_MAX, _PSUM_FMAX), dtype=nl.bfloat16, block_dim=[2],
|
build/torch-neuron/nki_kernels/attention_block_l1.py
CHANGED
|
@@ -31,6 +31,9 @@ from torch_neuronx import nki_op
|
|
| 31 |
import nki
|
| 32 |
import nki.isa as nisa
|
| 33 |
import nki.language as nl
|
|
|
|
|
|
|
|
|
|
| 34 |
from nki.isa import engine
|
| 35 |
|
| 36 |
from nkilib.core.utils.modular_allocator import ModularAllocator
|
|
@@ -142,7 +145,7 @@ def _attention_block_l1_kernel(h_hbm, wqk_hbm, bqk_hbm, tab_hbm, wvt_hbm, bv_hbm
|
|
| 142 |
r_wv.append(wv_sb_)
|
| 143 |
r_b.append(b_sb_)
|
| 144 |
r_bv.append(bv_bc_)
|
| 145 |
-
r_tab = ra.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=
|
| 146 |
block_dim=[4, n_tt], num_free_tiles=[4, n_tt], align_to=4)
|
| 147 |
for s in range(n_tt):
|
| 148 |
for t_i in range(4):
|
|
@@ -187,7 +190,7 @@ def _attention_block_l1_kernel(h_hbm, wqk_hbm, bqk_hbm, tab_hbm, wvt_hbm, bv_hbm
|
|
| 187 |
if resident:
|
| 188 |
tab = r_tab
|
| 189 |
else:
|
| 190 |
-
tab = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=
|
| 191 |
block_dim=[4, n_tt], num_free_tiles=[4, n_tt], align_to=4)
|
| 192 |
for s in range(n_tt):
|
| 193 |
for t_i in range(4):
|
|
@@ -246,8 +249,8 @@ def _attention_block_l1_kernel(h_hbm, wqk_hbm, bqk_hbm, tab_hbm, wvt_hbm, bv_hbm
|
|
| 246 |
|
| 247 |
proj = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=nl.bfloat16,
|
| 248 |
block_dim=[2], num_free_tiles=[2])
|
| 249 |
-
acc = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=
|
| 250 |
-
tmp = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=
|
| 251 |
roped = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=nl.bfloat16,
|
| 252 |
block_dim=[2], num_free_tiles=[2])
|
| 253 |
vout = allocator.alloc_sbuf_tensor(shape=(_V_TILE, local_heads * D), dtype=nl.float32,
|
|
|
|
| 31 |
import nki
|
| 32 |
import nki.isa as nisa
|
| 33 |
import nki.language as nl
|
| 34 |
+
import os as _osr
|
| 35 |
+
_ROPE_BF16 = _osr.environ.get("VIT_ENC_ROPE_BF16", "0") == "1"
|
| 36 |
+
_ROPE_DT = nl.bfloat16 if _ROPE_BF16 else nl.float32
|
| 37 |
from nki.isa import engine
|
| 38 |
|
| 39 |
from nkilib.core.utils.modular_allocator import ModularAllocator
|
|
|
|
| 145 |
r_wv.append(wv_sb_)
|
| 146 |
r_b.append(b_sb_)
|
| 147 |
r_bv.append(bv_bc_)
|
| 148 |
+
r_tab = ra.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=_ROPE_DT,
|
| 149 |
block_dim=[4, n_tt], num_free_tiles=[4, n_tt], align_to=4)
|
| 150 |
for s in range(n_tt):
|
| 151 |
for t_i in range(4):
|
|
|
|
| 190 |
if resident:
|
| 191 |
tab = r_tab
|
| 192 |
else:
|
| 193 |
+
tab = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=_ROPE_DT,
|
| 194 |
block_dim=[4, n_tt], num_free_tiles=[4, n_tt], align_to=4)
|
| 195 |
for s in range(n_tt):
|
| 196 |
for t_i in range(4):
|
|
|
|
| 249 |
|
| 250 |
proj = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=nl.bfloat16,
|
| 251 |
block_dim=[2], num_free_tiles=[2])
|
| 252 |
+
acc = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=_ROPE_DT, align_to=4)
|
| 253 |
+
tmp = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=_ROPE_DT, align_to=4)
|
| 254 |
roped = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=nl.bfloat16,
|
| 255 |
block_dim=[2], num_free_tiles=[2])
|
| 256 |
vout = allocator.alloc_sbuf_tensor(shape=(_V_TILE, local_heads * D), dtype=nl.float32,
|
build/torch-neuron/nki_kernels/attention_block_l1_ln.py
CHANGED
|
@@ -40,6 +40,9 @@ from torch_neuronx import nki_op
|
|
| 40 |
import nki
|
| 41 |
import nki.isa as nisa
|
| 42 |
import nki.language as nl
|
|
|
|
|
|
|
|
|
|
| 43 |
from nki.isa import engine
|
| 44 |
|
| 45 |
from nkilib.core.utils.modular_allocator import ModularAllocator
|
|
@@ -84,7 +87,7 @@ def sbuf_estimate(H: int, Fp: int, N: int, D: int, resident: bool) -> int:
|
|
| 84 |
|
| 85 |
@nki.jit
|
| 86 |
def _attention_block_l1_ln_kernel(h_hbm, wqk_hbm, bqk_hbm, tab_hbm, wvt_hbm, bv_hbm,
|
| 87 |
-
wot_hbm, bo_hbm, x_hbm, scale, head_dim, n_valid, eps, nimg=1, resident=0):
|
| 88 |
NT, H = h_hbm.shape
|
| 89 |
N = NT // nimg
|
| 90 |
Fp = wqk_hbm.shape[2]
|
|
@@ -155,7 +158,7 @@ def _attention_block_l1_ln_kernel(h_hbm, wqk_hbm, bqk_hbm, tab_hbm, wvt_hbm, bv_
|
|
| 155 |
r_wv.append(wv_sb_)
|
| 156 |
r_b.append(b_sb_)
|
| 157 |
r_bv.append(bv_bc_)
|
| 158 |
-
r_tab = ra.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=
|
| 159 |
block_dim=[4, n_tt], num_free_tiles=[4, n_tt], align_to=4)
|
| 160 |
for s in range(n_tt):
|
| 161 |
for t_i in range(4):
|
|
@@ -170,37 +173,52 @@ def _attention_block_l1_ln_kernel(h_hbm, wqk_hbm, bqk_hbm, tab_hbm, wvt_hbm, bv_
|
|
| 170 |
nisa.dma_copy(dst=r_bo, src=bo_hbm.ap(pattern=[[0, _P_MAX], [1, H]], offset=0))
|
| 171 |
res_top = ra.get_current_address()
|
| 172 |
|
| 173 |
-
|
| 174 |
-
|
| 175 |
-
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 176 |
allocator = ModularAllocator(initial_address=res_top)
|
| 177 |
|
| 178 |
# O for ALL groups stays in SBUF until the single out-projection.
|
| 179 |
o_feat = []
|
| 180 |
for g in range(G):
|
| 181 |
-
of_ = allocator.alloc_sbuf_tensor(shape=(_P_MAX,
|
| 182 |
block_dim=[n_ochunks], num_free_tiles=[n_ochunks],
|
| 183 |
align_to=32)
|
| 184 |
o_feat.append(of_)
|
| 185 |
|
| 186 |
# x^T: shared by every group (transposed ONCE per image).
|
| 187 |
x_t = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=nl.bfloat16,
|
| 188 |
-
block_dim=[n_chunks,
|
| 189 |
align_to=32)
|
| 190 |
-
for s in range(
|
| 191 |
for c in range(n_chunks):
|
| 192 |
nisa.dma_transpose(
|
| 193 |
dst=x_t[c][s].ap([[_PSUM_FMAX, _P_MAX], [1, 1], [1, 1],
|
| 194 |
-
[1, min(_PSUM_FMAX,
|
| 195 |
-
src=h_hbm.ap([[H, min(_PSUM_FMAX,
|
| 196 |
offset=rofs + (s * _PSUM_FMAX) * H + c * _P_MAX),
|
| 197 |
)
|
| 198 |
|
| 199 |
# rope tables: shared by every group.
|
| 200 |
-
if resident:
|
| 201 |
tab = r_tab
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 202 |
else:
|
| 203 |
-
tab = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=
|
| 204 |
block_dim=[4, n_tt], num_free_tiles=[4, n_tt], align_to=4)
|
| 205 |
for s in range(n_tt):
|
| 206 |
for t_i in range(4):
|
|
@@ -218,11 +236,11 @@ def _attention_block_l1_ln_kernel(h_hbm, wqk_hbm, bqk_hbm, tab_hbm, wvt_hbm, bv_
|
|
| 218 |
v_tiles = []
|
| 219 |
for _lh in range(local_heads):
|
| 220 |
kt_ = allocator.alloc_sbuf_tensor(shape=(D, _K_TILE), dtype=nl.bfloat16,
|
| 221 |
-
block_dim=[
|
| 222 |
align_to=32)
|
| 223 |
k_tiles.append(kt_)
|
| 224 |
vt_ = allocator.alloc_sbuf_tensor(shape=(_V_TILE, D), dtype=nl.bfloat16,
|
| 225 |
-
block_dim=[
|
| 226 |
v_tiles.append(vt_)
|
| 227 |
phase2_base = allocator.get_current_address()
|
| 228 |
|
|
@@ -259,20 +277,23 @@ def _attention_block_l1_ln_kernel(h_hbm, wqk_hbm, bqk_hbm, tab_hbm, wvt_hbm, bv_
|
|
| 259 |
|
| 260 |
proj = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=nl.bfloat16,
|
| 261 |
block_dim=[2], num_free_tiles=[2])
|
| 262 |
-
acc = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=
|
| 263 |
-
tmp = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=
|
| 264 |
roped = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=nl.bfloat16,
|
| 265 |
block_dim=[2], num_free_tiles=[2])
|
| 266 |
vout = allocator.alloc_sbuf_tensor(shape=(_V_TILE, local_heads * D), dtype=nl.float32,
|
| 267 |
align_to=4)
|
| 268 |
q_priv = []
|
| 269 |
for _lh in range(local_heads):
|
| 270 |
-
|
|
|
|
|
|
|
|
|
|
| 271 |
|
| 272 |
-
# ---- producer (identical to attention_block's N<=2048 path) ----
|
| 273 |
-
for tt in range(
|
| 274 |
t0 = tt * _PSUM_FMAX
|
| 275 |
-
tw = min(_PSUM_FMAX,
|
| 276 |
for vsub in range(tw // _V_TILE):
|
| 277 |
ps_v = nl.ndarray((_V_TILE, local_heads * D), dtype=nl.float32,
|
| 278 |
buffer=nl.psum, address=(0, 7 * _PSUM_BANK))
|
|
@@ -317,11 +338,17 @@ def _attention_block_l1_ln_kernel(h_hbm, wqk_hbm, bqk_hbm, tab_hbm, wvt_hbm, bv_
|
|
| 317 |
nisa.tensor_copy(
|
| 318 |
dst=k_tiles[lh][tt][par2 * half:(par2 + 1) * half, 0:tw],
|
| 319 |
src=roped[par2][r0:r0 + half, 0:tw])
|
| 320 |
-
|
| 321 |
for par2 in range(2):
|
| 322 |
nisa.dma_copy(
|
| 323 |
-
dst=q_priv[lh][0, par2 * half:(par2 + 1) * half, t0:t0 + tw],
|
| 324 |
src=roped[par2][r0:r0 + half, 0:tw])
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 325 |
|
| 326 |
# ---- attention for this group's heads -> o_feat[g] (SBUF, feature-major) ----
|
| 327 |
allocator.set_current_address(phase2_base)
|
|
@@ -335,18 +362,31 @@ def _attention_block_l1_ln_kernel(h_hbm, wqk_hbm, bqk_hbm, tab_hbm, wvt_hbm, bv_
|
|
| 335 |
o_sb.append(grp)
|
| 336 |
attn_base = allocator.get_current_address()
|
| 337 |
for lh in range(local_heads):
|
| 338 |
-
|
| 339 |
-
|
| 340 |
-
|
| 341 |
-
|
| 342 |
-
|
| 343 |
-
|
| 344 |
-
|
| 345 |
-
|
| 346 |
-
|
| 347 |
-
|
| 348 |
-
|
| 349 |
-
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 350 |
|
| 351 |
# ---- ONE out-projection over ALL G*n_ochunks feature chunks, full bias once ----
|
| 352 |
allocator.set_current_address(img_base)
|
|
@@ -376,7 +416,7 @@ def _attention_block_l1_ln_kernel(h_hbm, wqk_hbm, bqk_hbm, tab_hbm, wvt_hbm, bv_
|
|
| 376 |
n_ob = (H + _PSUM_FMAX - 1) // _PSUM_FMAX
|
| 377 |
for tb in range(n_grp):
|
| 378 |
k = tb % _RING
|
| 379 |
-
r0 = img *
|
| 380 |
nisa.dma_copy(dst=a_t[k], src=x_hbm[r0:r0 + _P_MAX, 0:H])
|
| 381 |
for ob in range(n_ob):
|
| 382 |
ow = min(_PSUM_FMAX, H - ob * _PSUM_FMAX)
|
|
@@ -438,6 +478,7 @@ def attention_block_l1_ln(
|
|
| 438 |
eps: float,
|
| 439 |
nimg: int = 1,
|
| 440 |
resident: int = 0,
|
|
|
|
| 441 |
) -> tuple[torch.Tensor, torch.Tensor]:
|
| 442 |
"""All-group attention + residual + affine-free LN2. Returns (x_new [R,H] fp32, z2 [R,H] bf16)."""
|
| 443 |
for t in (h, wqk, bqk, tab, wvt, bv, wot, bo, x):
|
|
@@ -447,4 +488,4 @@ def attention_block_l1_ln(
|
|
| 447 |
raise NotImplementedError("attention_block_l1_ln supports N<=2048 only")
|
| 448 |
if bqk.dim() == 3 and bqk.shape[1] != 128:
|
| 449 |
bqk = _pack_qk_bias(bqk)
|
| 450 |
-
return _ATTN_L1_LN(h, wqk, bqk, tab, wvt, bv, wot, bo, x, scale, head_dim, n_valid, eps, nimg, resident)
|
|
|
|
| 40 |
import nki
|
| 41 |
import nki.isa as nisa
|
| 42 |
import nki.language as nl
|
| 43 |
+
import os as _osr
|
| 44 |
+
_ROPE_BF16 = _osr.environ.get("VIT_ENC_ROPE_BF16", "0") == "1"
|
| 45 |
+
_ROPE_DT = nl.bfloat16 if _ROPE_BF16 else nl.float32
|
| 46 |
from nki.isa import engine
|
| 47 |
|
| 48 |
from nkilib.core.utils.modular_allocator import ModularAllocator
|
|
|
|
| 87 |
|
| 88 |
@nki.jit
|
| 89 |
def _attention_block_l1_ln_kernel(h_hbm, wqk_hbm, bqk_hbm, tab_hbm, wvt_hbm, bv_hbm,
|
| 90 |
+
wot_hbm, bo_hbm, x_hbm, scale, head_dim, n_valid, eps, nimg=1, resident=0, pack2=0):
|
| 91 |
NT, H = h_hbm.shape
|
| 92 |
N = NT // nimg
|
| 93 |
Fp = wqk_hbm.shape[2]
|
|
|
|
| 158 |
r_wv.append(wv_sb_)
|
| 159 |
r_b.append(b_sb_)
|
| 160 |
r_bv.append(bv_bc_)
|
| 161 |
+
r_tab = ra.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=_ROPE_DT,
|
| 162 |
block_dim=[4, n_tt], num_free_tiles=[4, n_tt], align_to=4)
|
| 163 |
for s in range(n_tt):
|
| 164 |
for t_i in range(4):
|
|
|
|
| 173 |
nisa.dma_copy(dst=r_bo, src=bo_hbm.ap(pattern=[[0, _P_MAX], [1, H]], offset=0))
|
| 174 |
res_top = ra.get_current_address()
|
| 175 |
|
| 176 |
+
# v1.5 (pack2): two images share one producer pass (their rows are contiguous, so one 512-token x^T tile
|
| 177 |
+
# holds both when N <= 256); Q/K/V projections then stream 2N columns per matmul. Attention stays per image.
|
| 178 |
+
PK = 2 if (pack2 and nimg % 2 == 0 and 2 * N <= _PSUM_FMAX) else 1
|
| 179 |
+
NP = N * PK
|
| 180 |
+
n_ttp = (NP + _PSUM_FMAX - 1) // _PSUM_FMAX
|
| 181 |
+
n_ktp = (NP + _K_TILE - 1) // _K_TILE
|
| 182 |
+
n_vtp = NP // _V_TILE
|
| 183 |
+
n_grp = NP // _P_MAX
|
| 184 |
+
n_grp1 = N // _P_MAX
|
| 185 |
+
for img in range(nimg // PK):
|
| 186 |
+
rofs = img * NP * H
|
| 187 |
allocator = ModularAllocator(initial_address=res_top)
|
| 188 |
|
| 189 |
# O for ALL groups stays in SBUF until the single out-projection.
|
| 190 |
o_feat = []
|
| 191 |
for g in range(G):
|
| 192 |
+
of_ = allocator.alloc_sbuf_tensor(shape=(_P_MAX, NP), dtype=nl.bfloat16,
|
| 193 |
block_dim=[n_ochunks], num_free_tiles=[n_ochunks],
|
| 194 |
align_to=32)
|
| 195 |
o_feat.append(of_)
|
| 196 |
|
| 197 |
# x^T: shared by every group (transposed ONCE per image).
|
| 198 |
x_t = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=nl.bfloat16,
|
| 199 |
+
block_dim=[n_chunks, n_ttp], num_free_tiles=[n_chunks, n_ttp],
|
| 200 |
align_to=32)
|
| 201 |
+
for s in range(n_ttp):
|
| 202 |
for c in range(n_chunks):
|
| 203 |
nisa.dma_transpose(
|
| 204 |
dst=x_t[c][s].ap([[_PSUM_FMAX, _P_MAX], [1, 1], [1, 1],
|
| 205 |
+
[1, min(_PSUM_FMAX, NP - s * _PSUM_FMAX)]], offset=0),
|
| 206 |
+
src=h_hbm.ap([[H, min(_PSUM_FMAX, NP - s * _PSUM_FMAX)], [1, 1], [1, 1], [1, _P_MAX]],
|
| 207 |
offset=rofs + (s * _PSUM_FMAX) * H + c * _P_MAX),
|
| 208 |
)
|
| 209 |
|
| 210 |
# rope tables: shared by every group.
|
| 211 |
+
if resident and PK == 1:
|
| 212 |
tab = r_tab
|
| 213 |
+
elif PK == 2:
|
| 214 |
+
# pack2: the same per-image tables, side by side (one 512-wide tile)
|
| 215 |
+
tab = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=_ROPE_DT,
|
| 216 |
+
block_dim=[4, 1], num_free_tiles=[4, 1], align_to=4)
|
| 217 |
+
for t_i in range(4):
|
| 218 |
+
for sub in range(2):
|
| 219 |
+
nisa.dma_copy(dst=tab[t_i][0][:, sub * N:(sub + 1) * N], src=tab_hbm[t_i, 0:_P_MAX, 0:N])
|
| 220 |
else:
|
| 221 |
+
tab = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=_ROPE_DT,
|
| 222 |
block_dim=[4, n_tt], num_free_tiles=[4, n_tt], align_to=4)
|
| 223 |
for s in range(n_tt):
|
| 224 |
for t_i in range(4):
|
|
|
|
| 236 |
v_tiles = []
|
| 237 |
for _lh in range(local_heads):
|
| 238 |
kt_ = allocator.alloc_sbuf_tensor(shape=(D, _K_TILE), dtype=nl.bfloat16,
|
| 239 |
+
block_dim=[n_ktp], num_free_tiles=[n_ktp],
|
| 240 |
align_to=32)
|
| 241 |
k_tiles.append(kt_)
|
| 242 |
vt_ = allocator.alloc_sbuf_tensor(shape=(_V_TILE, D), dtype=nl.bfloat16,
|
| 243 |
+
block_dim=[n_vtp], num_free_tiles=[n_vtp])
|
| 244 |
v_tiles.append(vt_)
|
| 245 |
phase2_base = allocator.get_current_address()
|
| 246 |
|
|
|
|
| 277 |
|
| 278 |
proj = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=nl.bfloat16,
|
| 279 |
block_dim=[2], num_free_tiles=[2])
|
| 280 |
+
acc = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=_ROPE_DT, align_to=4)
|
| 281 |
+
tmp = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=_ROPE_DT, align_to=4)
|
| 282 |
roped = allocator.alloc_sbuf_tensor(shape=(_P_MAX, _PSUM_FMAX), dtype=nl.bfloat16,
|
| 283 |
block_dim=[2], num_free_tiles=[2])
|
| 284 |
vout = allocator.alloc_sbuf_tensor(shape=(_V_TILE, local_heads * D), dtype=nl.float32,
|
| 285 |
align_to=4)
|
| 286 |
q_priv = []
|
| 287 |
for _lh in range(local_heads):
|
| 288 |
+
qs_ = []
|
| 289 |
+
for _sub in range(PK):
|
| 290 |
+
qs_.append(nl.ndarray((1, D, N), dtype=nl.bfloat16, buffer=nl.private_hbm))
|
| 291 |
+
q_priv.append(qs_)
|
| 292 |
|
| 293 |
+
# ---- producer (identical to attention_block's N<=2048 path; NP = PK images' tokens) ----
|
| 294 |
+
for tt in range(n_ttp):
|
| 295 |
t0 = tt * _PSUM_FMAX
|
| 296 |
+
tw = min(_PSUM_FMAX, NP - t0)
|
| 297 |
for vsub in range(tw // _V_TILE):
|
| 298 |
ps_v = nl.ndarray((_V_TILE, local_heads * D), dtype=nl.float32,
|
| 299 |
buffer=nl.psum, address=(0, 7 * _PSUM_BANK))
|
|
|
|
| 338 |
nisa.tensor_copy(
|
| 339 |
dst=k_tiles[lh][tt][par2 * half:(par2 + 1) * half, 0:tw],
|
| 340 |
src=roped[par2][r0:r0 + half, 0:tw])
|
| 341 |
+
elif PK == 1:
|
| 342 |
for par2 in range(2):
|
| 343 |
nisa.dma_copy(
|
| 344 |
+
dst=q_priv[lh][0][0, par2 * half:(par2 + 1) * half, t0:t0 + tw],
|
| 345 |
src=roped[par2][r0:r0 + half, 0:tw])
|
| 346 |
+
else:
|
| 347 |
+
for par2 in range(2):
|
| 348 |
+
for sub in range(2):
|
| 349 |
+
nisa.dma_copy(
|
| 350 |
+
dst=q_priv[lh][sub][0, par2 * half:(par2 + 1) * half, 0:N],
|
| 351 |
+
src=roped[par2][r0:r0 + half, sub * N:(sub + 1) * N])
|
| 352 |
|
| 353 |
# ---- attention for this group's heads -> o_feat[g] (SBUF, feature-major) ----
|
| 354 |
allocator.set_current_address(phase2_base)
|
|
|
|
| 362 |
o_sb.append(grp)
|
| 363 |
attn_base = allocator.get_current_address()
|
| 364 |
for lh in range(local_heads):
|
| 365 |
+
for sub in range(PK):
|
| 366 |
+
if PK == 1:
|
| 367 |
+
k_l = k_tiles[lh]
|
| 368 |
+
v_l = v_tiles[lh]
|
| 369 |
+
o_l = o_sb[lh]
|
| 370 |
+
else:
|
| 371 |
+
k_l = [k_tiles[lh][0][:, sub * N:(sub + 1) * N]]
|
| 372 |
+
v_l = []
|
| 373 |
+
for _v in range(n_grp1):
|
| 374 |
+
v_l.append(v_tiles[lh][sub * n_grp1 + _v])
|
| 375 |
+
o_l = []
|
| 376 |
+
for _o in range(n_grp1):
|
| 377 |
+
o_l.append(o_sb[lh][sub * n_grp1 + _o])
|
| 378 |
+
_attention_cte(
|
| 379 |
+
q_priv[lh][sub], None, None,
|
| 380 |
+
scale=scale, causal_mask=False,
|
| 381 |
+
tp_q=False, tp_k=False, tp_out=True,
|
| 382 |
+
cache_softmax=False,
|
| 383 |
+
k_cache_sbuf=k_l,
|
| 384 |
+
v_cache_sbuf=v_l,
|
| 385 |
+
out_o_hbm=o_l,
|
| 386 |
+
init_sbuf_addr=attn_base,
|
| 387 |
+
seqlen_k_valid=n_valid,
|
| 388 |
+
)
|
| 389 |
+
allocator.set_current_address(attn_base)
|
| 390 |
|
| 391 |
# ---- ONE out-projection over ALL G*n_ochunks feature chunks, full bias once ----
|
| 392 |
allocator.set_current_address(img_base)
|
|
|
|
| 416 |
n_ob = (H + _PSUM_FMAX - 1) // _PSUM_FMAX
|
| 417 |
for tb in range(n_grp):
|
| 418 |
k = tb % _RING
|
| 419 |
+
r0 = img * NP + tb * _P_MAX
|
| 420 |
nisa.dma_copy(dst=a_t[k], src=x_hbm[r0:r0 + _P_MAX, 0:H])
|
| 421 |
for ob in range(n_ob):
|
| 422 |
ow = min(_PSUM_FMAX, H - ob * _PSUM_FMAX)
|
|
|
|
| 478 |
eps: float,
|
| 479 |
nimg: int = 1,
|
| 480 |
resident: int = 0,
|
| 481 |
+
pack2: int = 0,
|
| 482 |
) -> tuple[torch.Tensor, torch.Tensor]:
|
| 483 |
"""All-group attention + residual + affine-free LN2. Returns (x_new [R,H] fp32, z2 [R,H] bf16)."""
|
| 484 |
for t in (h, wqk, bqk, tab, wvt, bv, wot, bo, x):
|
|
|
|
| 488 |
raise NotImplementedError("attention_block_l1_ln supports N<=2048 only")
|
| 489 |
if bqk.dim() == 3 and bqk.shape[1] != 128:
|
| 490 |
bqk = _pack_qk_bias(bqk)
|
| 491 |
+
return _ATTN_L1_LN(h, wqk, bqk, tab, wvt, bv, wot, bo, x, scale, head_dim, n_valid, eps, nimg, resident, pack2)
|
build/torch-neuron/nki_kernels/attention_inner.py
CHANGED
|
@@ -110,6 +110,7 @@ _MAX_GLOBAL_CP_DEGREE = 32
|
|
| 110 |
_Q_GRP_SZ = 128
|
| 111 |
_V_TILE_SZ = 128 # V tile size for loading and MM2 operations
|
| 112 |
import os as _os
|
|
|
|
| 113 |
_MM1_WIDE = _os.environ.get("VIT_ENC_MM1_WIDE", "1") == "1" # set 0 to restore the v1.3 4-bank MM1 layout
|
| 114 |
_MM1_BANKS = (0, 1, 2, 5, 6) # MM1 rotation; bank 4 = P.V (single large tile), 3/7 = reciprocal transposes
|
| 115 |
_MM1_WIDE_MAX_KTILES = 2 # only widen when <= this many key tiles per group: at 4 key tiles (N=2048) the
|
|
@@ -480,30 +481,40 @@ def _attention_cte_impl(
|
|
| 480 |
|
| 481 |
# SBUF-KV: K and V are already resident, so the HBM K/V load block is skipped entirely.
|
| 482 |
|
| 483 |
-
# Software pipelining: num_grps=16 > 1, so the pipelined branch is the only reached one.
|
| 484 |
-
_load_q_impl(shard_seqlen_q_start, ac, atp, sp, bufs, q, batch_id, sbuf_addr)
|
| 485 |
-
_qk_and_max_impl(shard_seqlen_q_start, ac, atp, sp, bufs, batch_id)
|
| 486 |
-
_update_max_impl(shard_seqlen_q_start, ac, atp, sp, bufs, None)
|
| 487 |
-
_exp_impl(shard_seqlen_q_start, ac, atp, sp, bufs, None)
|
| 488 |
-
|
| 489 |
-
_load_q_impl(shard_seqlen_q_start + 1, ac, atp, sp, bufs, q, batch_id, sbuf_addr)
|
| 490 |
-
_qk_and_max_impl(shard_seqlen_q_start + 1, ac, atp, sp, bufs, batch_id)
|
| 491 |
-
_update_max_impl(shard_seqlen_q_start + 1, ac, atp, sp, bufs, None)
|
| 492 |
shard_seqlen_q_end = shard_seqlen_q_start + shard_seqlen_q_length
|
| 493 |
-
|
| 494 |
-
|
| 495 |
-
#
|
| 496 |
-
_load_q_impl(
|
| 497 |
-
|
| 498 |
-
|
| 499 |
-
|
| 500 |
-
|
| 501 |
-
|
| 502 |
-
|
| 503 |
-
|
| 504 |
-
|
| 505 |
-
|
| 506 |
-
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 507 |
|
| 508 |
# cache_softmax intermediate writes (out_neg_max / out_sum_recip) are DEAD (cache_softmax=False).
|
| 509 |
|
|
@@ -1081,6 +1092,17 @@ def _scale_reciprocal_write_back_impl(src_buf, grp_i, ac: AttnConfig, atp: AttnT
|
|
| 1081 |
bufs.tp_exp_sum_reciprocal_psum[grp_i].ap(pattern=[[atp.sb_p, ac.d], [1, atp.sb_p]], offset=0),
|
| 1082 |
bufs.exp_sum_reciprocal.ap(pattern=[[atp.num_grps, atp.sb_p], [0, ac.d]], offset=grp_i),
|
| 1083 |
)
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1084 |
# ---- h8 SITE 1: PSUM->SBUF evict pinned to Scalar. ----
|
| 1085 |
nisa.tensor_copy(
|
| 1086 |
bufs.tp_exp_sum_reciprocal_sb[grp_i][: ac.d, :num_f],
|
|
|
|
| 110 |
_Q_GRP_SZ = 128
|
| 111 |
_V_TILE_SZ = 128 # V tile size for loading and MM2 operations
|
| 112 |
import os as _os
|
| 113 |
+
_WB_FUSED = _os.environ.get("VIT_ENC_WB_FUSED", "0") == "1" # fused normalise+cast write-back: OFF (measured +6.6% at N=2048)
|
| 114 |
_MM1_WIDE = _os.environ.get("VIT_ENC_MM1_WIDE", "1") == "1" # set 0 to restore the v1.3 4-bank MM1 layout
|
| 115 |
_MM1_BANKS = (0, 1, 2, 5, 6) # MM1 rotation; bank 4 = P.V (single large tile), 3/7 = reciprocal transposes
|
| 116 |
_MM1_WIDE_MAX_KTILES = 2 # only widen when <= this many key tiles per group: at 4 key tiles (N=2048) the
|
|
|
|
| 481 |
|
| 482 |
# SBUF-KV: K and V are already resident, so the HBM K/V load block is skipped entirely.
|
| 483 |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 484 |
shard_seqlen_q_end = shard_seqlen_q_start + shard_seqlen_q_length
|
| 485 |
+
if shard_seqlen_q_length == 1:
|
| 486 |
+
# v1.5: one Q group (N <= 128 tokens) -- no pipelining (the pipelined prologue below assumes >= 2
|
| 487 |
+
# groups and indexed a non-existent second Q tile, so N=128 failed to compile in <= v1.4).
|
| 488 |
+
_load_q_impl(shard_seqlen_q_start, ac, atp, sp, bufs, q, batch_id, sbuf_addr)
|
| 489 |
+
_qk_and_max_impl(shard_seqlen_q_start, ac, atp, sp, bufs, batch_id)
|
| 490 |
+
_update_max_impl(shard_seqlen_q_start, ac, atp, sp, bufs, None)
|
| 491 |
+
_exp_impl(shard_seqlen_q_start, ac, atp, sp, bufs, None)
|
| 492 |
+
_pv_impl(shard_seqlen_q_start, ac, atp, sp, bufs)
|
| 493 |
+
_write_back_impl(shard_seqlen_q_start, ac, atp, sp, bufs, o, batch_id)
|
| 494 |
+
else:
|
| 495 |
+
# Software pipelining: >= 2 Q groups.
|
| 496 |
+
_load_q_impl(shard_seqlen_q_start, ac, atp, sp, bufs, q, batch_id, sbuf_addr)
|
| 497 |
+
_qk_and_max_impl(shard_seqlen_q_start, ac, atp, sp, bufs, batch_id)
|
| 498 |
+
_update_max_impl(shard_seqlen_q_start, ac, atp, sp, bufs, None)
|
| 499 |
+
_exp_impl(shard_seqlen_q_start, ac, atp, sp, bufs, None)
|
| 500 |
+
|
| 501 |
+
_load_q_impl(shard_seqlen_q_start + 1, ac, atp, sp, bufs, q, batch_id, sbuf_addr)
|
| 502 |
+
_qk_and_max_impl(shard_seqlen_q_start + 1, ac, atp, sp, bufs, batch_id)
|
| 503 |
+
_update_max_impl(shard_seqlen_q_start + 1, ac, atp, sp, bufs, None)
|
| 504 |
+
|
| 505 |
+
for grp_i in range(shard_seqlen_q_start, shard_seqlen_q_end - 2):
|
| 506 |
+
# grp_i: PV, write_back ; grp_i+1: EXP ; grp_i+2: load Q, QK+Max.
|
| 507 |
+
_load_q_impl(grp_i + 2, ac, atp, sp, bufs, q, batch_id, sbuf_addr)
|
| 508 |
+
_exp_impl(grp_i + 1, ac, atp, sp, bufs, None)
|
| 509 |
+
_fused_qkmax_and_pv_impl(grp_i, ac, atp, sp, bufs, batch_id)
|
| 510 |
+
_write_back_impl(grp_i, ac, atp, sp, bufs, o, batch_id)
|
| 511 |
+
_update_max_impl(grp_i + 2, ac, atp, sp, bufs, None)
|
| 512 |
+
|
| 513 |
+
_pv_impl(shard_seqlen_q_end - 2, ac, atp, sp, bufs)
|
| 514 |
+
_write_back_impl(shard_seqlen_q_end - 2, ac, atp, sp, bufs, o, batch_id)
|
| 515 |
+
_exp_impl(shard_seqlen_q_end - 1, ac, atp, sp, bufs, None)
|
| 516 |
+
_pv_impl(shard_seqlen_q_end - 1, ac, atp, sp, bufs)
|
| 517 |
+
_write_back_impl(shard_seqlen_q_end - 1, ac, atp, sp, bufs, o, batch_id)
|
| 518 |
|
| 519 |
# cache_softmax intermediate writes (out_neg_max / out_sum_recip) are DEAD (cache_softmax=False).
|
| 520 |
|
|
|
|
| 1092 |
bufs.tp_exp_sum_reciprocal_psum[grp_i].ap(pattern=[[atp.sb_p, ac.d], [1, atp.sb_p]], offset=0),
|
| 1093 |
bufs.exp_sum_reciprocal.ap(pattern=[[atp.num_grps, atp.sb_p], [0, ac.d]], offset=grp_i),
|
| 1094 |
)
|
| 1095 |
+
if _WB_FUSED:
|
| 1096 |
+
# v1.5: one Vector multiply reads the transposed reciprocal straight from PSUM and writes bf16 into the
|
| 1097 |
+
# caller's O slice -- replaces the Scalar PSUM evict, the fp32 multiply into mm2_final and the GpSimd
|
| 1098 |
+
# fp32->bf16 cast (same single rounding to bf16 as before: fp32 product, rounded once on write).
|
| 1099 |
+
nisa.tensor_tensor(
|
| 1100 |
+
o[grp_i][: ac.d, :num_f],
|
| 1101 |
+
src_buf[: ac.d, :num_f],
|
| 1102 |
+
bufs.tp_exp_sum_reciprocal_psum[grp_i][: ac.d, :num_f],
|
| 1103 |
+
nl.multiply,
|
| 1104 |
+
)
|
| 1105 |
+
return
|
| 1106 |
# ---- h8 SITE 1: PSUM->SBUF evict pinned to Scalar. ----
|
| 1107 |
nisa.tensor_copy(
|
| 1108 |
bufs.tp_exp_sum_reciprocal_sb[grp_i][: ac.d, :num_f],
|