CHANGELOG and comment wording
Browse files
CHANGELOG.md
CHANGED
|
@@ -28,7 +28,7 @@
|
|
| 28 |
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%)**,
|
| 29 |
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+
|
| 30 |
+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.
|
| 31 |
-
|
| 32 |
* **Fix: inputs of 128 tokens or fewer** failed to compile in every attention kernel (`index 1 out of bounds`). They
|
| 33 |
now use a non-pipelined single-Q-group path (rel_l2 vs a float64 reference 1.9-2.5e-3). Longer inputs are
|
| 34 |
unchanged (bit-identical).
|
|
|
|
| 28 |
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%)**,
|
| 29 |
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+
|
| 30 |
+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.
|
| 31 |
+
A fused-attention-block technique (two images per producer pass).
|
| 32 |
* **Fix: inputs of 128 tokens or fewer** failed to compile in every attention kernel (`index 1 out of bounds`). They
|
| 33 |
now use a non-pipelined single-Q-group path (rel_l2 vs a float64 reference 1.9-2.5e-3). Longer inputs are
|
| 34 |
unchanged (bit-identical).
|
build/torch-neuron/nki_kernels/attention_block.py
CHANGED
|
@@ -626,7 +626,7 @@ def _attention_block_kernel(h_hbm, wqk_hbm, bqk_hbm, tab_hbm, wvt_hbm, bv_hbm,
|
|
| 626 |
# Scalar's 612 ns fee is faster than the 744 ns Vector cell, so the
|
| 627 |
# zero-lead rope chain consuming proj cannot lengthen. ========
|
| 628 |
# NKI 0.7: a Scalar-engine pin here fails ISA opcode_engine_check on newer compilers
|
| 629 |
-
#
|
| 630 |
nisa.tensor_scalar(dst=proj[par][:, 0:tw], data=ps[:, 0:tw],
|
| 631 |
op0=nl.add, operand0=b_sb[qk][:, g:g + 1])
|
| 632 |
|
|
|
|
| 626 |
# Scalar's 612 ns fee is faster than the 744 ns Vector cell, so the
|
| 627 |
# zero-lead rope chain consuming proj cannot lengthen. ========
|
| 628 |
# NKI 0.7: a Scalar-engine pin here fails ISA opcode_engine_check on newer compilers
|
| 629 |
+
# newer compilers -> unpinned (compiler picks the engine).
|
| 630 |
nisa.tensor_scalar(dst=proj[par][:, 0:tw], data=ps[:, 0:tw],
|
| 631 |
op0=nl.add, operand0=b_sb[qk][:, g:g + 1])
|
| 632 |
|