GEMM, MoE & Quantization¶
20 ops, 251 workloads — Linear Algebra (GEMM) 18 · Quantization 2.
One table per op, one row per workload. Ratio is the baseline device time divided by ours in the same measurement regime, so green is faster than it, plain is level with it, red is slower. Times are in ms. How these numbers are taken.
Linear Algebra (GEMM)¶
GemmFp8Fwd¶
dtype=fp8e4m3
- W1m=128n=7168k=2048scale_mode=block128
ds-v3-decode-down-block128 - W2m=128n=7168k=2048scale_mode=per_tensor
ds-v3-decode-down-per-tensor - W3m=128n=2112k=7168scale_mode=block128
ds-v3-decode-gate-up-block128 - W5m=128n=2112k=7168scale_mode=per_tensor
ds-v3-decode-gate-up-per-tensor - W6m=4096n=4096k=7168scale_mode=block128
ds-v3-prefill-attn-proj-block128 - W7m=4096n=7168k=2048scale_mode=block128
ds-v3-prefill-down-block128 - W8m=4096n=7168k=2048scale_mode=per_tensor
ds-v3-prefill-down-per-tensor - W9m=4096n=2112k=7168scale_mode=block128
ds-v3-prefill-gate-up-block128 - W10m=4096n=2112k=7168scale_mode=per_tensor
ds-v3-prefill-gate-up-per-tensor - W11m=1n=7168k=2048scale_mode=block128
gemv-down-m1-block128 - W12m=1n=7168k=2048scale_mode=per_tensor
gemv-down-m1-per-tensor - W13m=4096n=7168k=16384scale_mode=block128
k-dominant-7168x16384-block128 - W14m=8n=7168k=2048scale_mode=per_tensor
small-batch-down-m8-per-tensor - W15m=4096n=24576k=1536scale_mode=block128
wide-n-24576-block128
bias: [N], bf16
dtype=fp8e4m3N=2112m=128n=2112k=7168scale_mode=per_tensor
- W4
ds-v3-decode-gate-up-per-tensor-bias
| Workload | Ratio | Device time | Alternatives | Throughput | SOL | Bound | |
|---|---|---|---|---|---|---|---|
| alt / ours | ms | name · tier | ms | TFLOP/s | of ceiling | by | |
| W1 | 3.22× | 0.2166 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library ⓘ | 0.6985 1.0308 | · | · | · |
| W2 | 5.11× | 0.1100 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library ⓘ | 0.5644 0.6743 | · | · | · |
| W3 | 3.06× | 0.2300 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library ⓘ | 0.7031 1.0401 | · | · | · |
| W4 | 5.86× | 0.1308 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis ⓘ | 0.7050 0.7047 0.7046 | · | · | · |
| W5 | 5.30× | 0.1085 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library ⓘ | 0.5749 0.7051 | · | · | · |
| W6 | 4.53× | 1.3802 | torch_compile_aclgraph torch_compile basistorch_compile_ge torch_compiletorch_npu eager vendor library ⓘ | 6.2960 6.2969 6.2966 | · | · | · |
| W7 | 1.80× | 0.7394 | torch_compile_aclgraph torch_compile basistorch_compile_ge torch_compiletorch_npu eager vendor library ⓘ | 1.3319 1.7572 1.7561 | · | · | · |
| W8 | 1.92× | 0.5847 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library ⓘ | 1.1205 1.2816 | · | · | · |
| W9 | 3.39× | 0.9463 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library ⓘ | 3.2060 4.0818 | · | · | · |
| W10 | 4.57× | 0.6191 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library ⓘ | 2.8299 3.2950 | · | · | · |
| W11 | 4.63× | 0.2034 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis ⓘ | 0.9085 0.9089 0.9083 | · | · | · |
| W12 | 6.85× | 0.0968 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis ⓘ | 0.6019 0.6010 0.6009 | · | · | · |
| W13 | 3.34× | 6.8591 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis ⓘ | 22.8992 22.9017 22.8991 | · | · | · |
| W14 | 7.14× | 0.0963 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis ⓘ | 0.6411 0.6416 0.6410 | · | · | · |
| W15 | 2.34× | 2.2824 | torch_compile_aclgraph torch_compile basistorch_compile_ge torch_compiletorch_npu eager vendor library ⓘ | 5.3496 5.3502 5.3505 | · | · | · |
AllGatherGemmFwd¶
- W1dtype=bf16k=2048m=128n=7168trans_b=true
ds-v3-decode-down@rank0 - W2dtype=bf16k=2048m=128n=7168trans_b=true
ds-v3-decode-down@rank1 - W3dtype=bf16k=7168m=128n=2112trans_b=true
ds-v3-decode-gate-up@rank0 - W4dtype=bf16k=7168m=128n=2112trans_b=true
ds-v3-decode-gate-up@rank1 - W5dtype=bf16k=7168m=4096n=4096trans_b=true
ds-v3-prefill-attn-proj@rank0 - W6dtype=f16k=7168m=4096n=4096trans_b=true
ds-v3-prefill-attn-proj@rank0 - W7dtype=bf16k=7168m=4096n=4096trans_b=true
ds-v3-prefill-attn-proj@rank1 - W8dtype=f16k=7168m=4096n=4096trans_b=true
ds-v3-prefill-attn-proj@rank1 - W9dtype=bf16k=2048m=4096n=7168trans_b=true
ds-v3-prefill-down@rank0 - W10dtype=bf16k=2048m=4096n=7168trans_b=true
ds-v3-prefill-down@rank1 - W11dtype=bf16k=7168m=4096n=2112trans_b=true
ds-v3-prefill-gate-up@rank0 - W12dtype=bf16k=7168m=4096n=2112trans_b=true
ds-v3-prefill-gate-up@rank1 - W13dtype=bf16k=16384m=4096n=7168trans_b=true
k-dominant-7168x16384@rank0 - W14dtype=bf16k=16384m=4096n=7168trans_b=true
k-dominant-7168x16384@rank1 - W15dtype=bf16k=7168m=16n=4096trans_b=true
mid-m16-attn@rank0 - W16dtype=bf16k=7168m=16n=4096trans_b=true
mid-m16-attn@rank1 - W17dtype=bf16k=7168m=32n=4096trans_b=true
mid-m32-attn@rank0 - W18dtype=bf16k=7168m=32n=4096trans_b=true
mid-m32-attn@rank1 - W19dtype=bf16k=2048m=64n=7168trans_b=true
mid-m64-down@rank0 - W20dtype=bf16k=2048m=64n=7168trans_b=true
mid-m64-down@rank1 - W21dtype=bf16k=7168m=96n=2112trans_b=true
mid-m96-gate-up@rank0 - W22dtype=bf16k=7168m=96n=2112trans_b=true
mid-m96-gate-up@rank1 - W23dtype=bf16k=1024m=1024n=1024
square-1k-nn@rank0 - W24dtype=f16k=1024m=1024n=1024
square-1k-nn@rank0 - W25dtype=bf16k=1024m=1024n=1024
square-1k-nn@rank1 - W26dtype=f16k=1024m=1024n=1024
square-1k-nn@rank1 - W27dtype=bf16k=1536m=4096n=24576trans_b=true
wide-n-24576@rank0 - W28dtype=bf16k=1536m=4096n=24576trans_b=true
wide-n-24576@rank1
| Workload | Ratio | Device time | Alternatives | Throughput | SOL | Bound | |
|---|---|---|---|---|---|---|---|
| alt / ours | ms | name | ms | TFLOP/s | of ceiling | by | |
| W1 | 1.93× | 0.0975 | Torch-NPU semantic implementation | 0.1886 | · | · | · |
| W2 | 1.92× | 0.1021 | Torch-NPU semantic implementation | 0.1963 | · | · | · |
| W3 | 1.57× | 0.1379 | Torch-NPU semantic implementation | 0.2162 | · | · | · |
| W4 | 1.56× | 0.1420 | Torch-NPU semantic implementation | 0.2220 | · | · | · |
| W5 | 1.94× | 2.4007 | Torch-NPU semantic implementation | 4.6677 | · | · | · |
| W6 | 1.95× | 2.3924 | Torch-NPU semantic implementation | 4.6666 | · | · | · |
| W7 | 1.96× | 2.4015 | Torch-NPU semantic implementation | 4.7040 | · | · | · |
| W8 | 1.96× | 2.3937 | Torch-NPU semantic implementation | 4.6899 | · | · | · |
| W9 | 2.07× | 0.9513 | Torch-NPU semantic implementation | 1.9730 | · | · | · |
| W10 | 2.09× | 0.9426 | Torch-NPU semantic implementation | 1.9680 | · | · | · |
| W11 | 1.62× | 2.0751 | Torch-NPU semantic implementation | 3.3533 | · | · | · |
| W12 | 1.61× | 2.0659 | Torch-NPU semantic implementation | 3.3201 | · | · | · |
| W13 | 2.03× | 8.2814 | Torch-NPU semantic implementation | 16.7741 | · | · | · |
| W14 | 2.01× | 8.3184 | Torch-NPU semantic implementation | 16.7129 | · | · | · |
| W15 | 2.07× | 0.1389 | Torch-NPU semantic implementation | 0.2869 | · | · | · |
| W16 | 2.18× | 0.1323 | Torch-NPU semantic implementation | 0.2886 | · | · | · |
| W17 | 1.73× | 0.1512 | Torch-NPU semantic implementation | 0.2619 | · | · | · |
| W18 | 1.70× | 0.1560 | Torch-NPU semantic implementation | 0.2657 | · | · | · |
| W19 | 1.82× | 0.0902 | Torch-NPU semantic implementation | 0.1640 | · | · | · |
| W20 | 1.78× | 0.0895 | Torch-NPU semantic implementation | 0.1595 | · | · | · |
| W21 | 1.60× | 0.1393 | Torch-NPU semantic implementation | 0.2226 | · | · | · |
| W22 | 1.58× | 0.1368 | Torch-NPU semantic implementation | 0.2166 | · | · | · |
| W23 | 1.49× | 0.1123 | Torch-NPU semantic implementation | 0.1678 | · | · | · |
| W24 | 1.45× | 0.1251 | Torch-NPU semantic implementation | 0.1814 | · | · | · |
| W25 | 1.55× | 0.1027 | Torch-NPU semantic implementation | 0.1593 | · | · | · |
| W26 | 1.58× | 0.1032 | Torch-NPU semantic implementation | 0.1631 | · | · | · |
| W27 | 2.87× | 1.6191 | Torch-NPU semantic implementation | 4.6493 | · | · | · |
| W28 | 2.87× | 1.6142 | Torch-NPU semantic implementation | 4.6391 | · | · | · |
GemmBiasReluFwd¶
- W1dtype=bf16m=128n=7168k=2048
ds-v3-decode-down - W2dtype=bf16m=128n=2112k=7168
ds-v3-decode-gate-up - W3dtype=bf16m=4096n=4096k=7168
ds-v3-prefill-attn-proj - W4dtype=f16m=4096n=4096k=7168
ds-v3-prefill-attn-proj - W5dtype=bf16m=4096n=7168k=2048
ds-v3-prefill-down - W6dtype=bf16m=4096n=2112k=7168
ds-v3-prefill-gate-up - W7dtype=bf16m=4096n=7168k=16384
k-dominant-7168x16384 - W8dtype=bf16m=16n=4096k=7168
mid-m16-attn - W9dtype=bf16m=32n=4096k=7168
mid-m32-attn - W10dtype=bf16m=64n=7168k=2048
mid-m64-down - W11dtype=bf16m=96n=2112k=7168
mid-m96-gate-up - W12dtype=bf16m=1024n=1024k=1024trans_b=false
square-1k-nn - W13dtype=f16m=1024n=1024k=1024trans_b=false
square-1k-nn - W14dtype=bf16m=4096n=24576k=1536
wide-n-24576
| Workload | Ratio | Device time | Alternatives | Throughput | SOL | Bound | |
|---|---|---|---|---|---|---|---|
| alt / ours | ms | name · tier | ms | TFLOP/s | of ceiling | by | |
| W1 | 1.13× | 0.0757 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.0855 0.0664 0.0643 | · | · | · |
| W2 | 0.96× | 0.0813 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.0783 0.0693 0.0660 | · | · | · |
| W3 | 1.02× | 0.8664 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.8825 0.8825 0.8821 | · | · | · |
| W4 | 1.00× | 0.8652 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.8552 0.8524 0.8545 | · | · | · |
| W5 | 0.88× | 0.5421 | torch_compile_aclgraph torch_compile basistorch_compile_ge torch_compiletorch_npu eager vendor library | 0.4699 0.4703 0.4703 | · | · | · |
| W6 | 1.06× | 0.5220 | torch_compile_aclgraph torch_compile basistorch_compile_ge torch_compiletorch_npu eager vendor library | 0.5449 0.5461 0.5451 | · | · | · |
| W7 | 0.75× | 4.7685 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 3.5644 3.5620 3.5609 | · | · | · |
| W8 | 0.99× | 0.1098 | torch_compile_aclgraph torch_compile basistorch_compile_ge torch_compiletorch_npu eager vendor library | 0.1047 0.1046 0.1052 | · | · | · |
| W9 | 1.02× | 0.1125 | torch_compile_aclgraph torch_compile basistorch_compile_ge torch_compiletorch_npu eager vendor library | 0.1081 0.1085 0.1082 | · | · | · |
| W10 | 1.11× | 0.0722 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.0600 0.0599 0.0599 | · | · | · |
| W11 | 1.05× | 0.0803 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.0641 0.0640 0.0641 | · | · | · |
| W12 | 1.04× | 0.0375 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.0395 0.0272 0.0266 | · | · | · |
| W13 | 0.89× | 0.0408 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.0367 0.0248 0.0253 | · | · | · |
| W14 | 0.94× | 1.6584 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 1.5746 1.5749 1.5733 | · | · | · |
GemmBiasGeluFwd¶
- W1dtype=bf16m=128n=7168k=2048
ds-v3-decode-down - W2dtype=bf16m=128n=2112k=7168
ds-v3-decode-gate-up - W3dtype=bf16m=4096n=4096k=7168
ds-v3-prefill-attn-proj - W4dtype=f16m=4096n=4096k=7168
ds-v3-prefill-attn-proj - W5dtype=bf16m=4096n=7168k=2048
ds-v3-prefill-down - W6dtype=bf16m=4096n=2112k=7168
ds-v3-prefill-gate-up - W7dtype=bf16m=4096n=7168k=16384
k-dominant-7168x16384 - W8dtype=bf16m=16n=4096k=7168
mid-m16-attn - W9dtype=bf16m=32n=4096k=7168
mid-m32-attn - W10dtype=bf16m=64n=7168k=2048
mid-m64-down - W11dtype=bf16m=96n=2112k=7168
mid-m96-gate-up - W12dtype=bf16m=1024n=1024k=1024trans_b=false
square-1k-nn - W13dtype=f16m=1024n=1024k=1024trans_b=false
square-1k-nn - W14dtype=bf16m=4096n=24576k=1536
wide-n-24576
| Workload | Ratio | Device time | Alternatives | Throughput | SOL | Bound | |
|---|---|---|---|---|---|---|---|
| alt / ours | ms | name · tier | ms | TFLOP/s | of ceiling | by | |
| W1 | 0.99× | 0.0860 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.0859 0.0660 0.0661 | · | · | · |
| W2 | 0.97× | 0.0818 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.0795 0.0685 0.0680 | · | · | · |
| W3 | 0.98× | 0.9435 | torch_compile_aclgraph torch_compile basistorch_compile_ge torch_compiletorch_npu eager vendor library | 0.9164 0.9164 0.9165 | · | · | · |
| W4 | 0.96× | 0.9353 | torch_compile_aclgraph torch_compile basistorch_compile_ge torch_compiletorch_npu eager vendor library | 0.8981 0.8942 0.8982 | · | · | · |
| W5 | 0.81× | 0.6883 | torch_compile_aclgraph torch_compile basistorch_compile_ge torch_compiletorch_npu eager vendor library | 0.5459 0.5459 0.5463 | · | · | · |
| W6 | 1.01× | 0.5693 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.5640 0.5639 0.5640 | · | · | · |
| W7 | 0.74× | 4.9324 | torch_compile_aclgraph torch_compile basistorch_compile_ge torch_compiletorch_npu eager vendor library | 3.6449 3.6480 3.6484 | · | · | · |
| W8 | 1.02× | 0.1134 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.1055 0.1057 0.1054 | · | · | · |
| W9 | 1.01× | 0.1126 | torch_compile_aclgraph torch_compile basistorch_compile_ge torch_compiletorch_npu eager vendor library | 0.1098 0.1092 0.1099 | · | · | · |
| W10 | 0.94× | 0.0798 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.0595 0.0594 0.0595 | · | · | · |
| W11 | 0.93× | 0.0843 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.0654 0.0650 0.0650 | · | · | · |
| W12 | 0.86× | 0.0484 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.0418 0.0302 0.0304 | · | · | · |
| W13 | 0.84× | 0.0465 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.0398 0.0300 0.0288 | · | · | · |
| W14 | 0.77× | 2.0700 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 1.6100 1.6104 1.6074 | · | · | · |
GemmBiasFwd¶
- W1dtype=bf16m=128n=7168k=2048
ds-v3-decode-down - W2dtype=bf16m=128n=2112k=7168
ds-v3-decode-gate-up - W3dtype=bf16m=4096n=4096k=7168
ds-v3-prefill-attn-proj - W4dtype=f16m=4096n=4096k=7168
ds-v3-prefill-attn-proj - W5dtype=bf16m=4096n=7168k=2048
ds-v3-prefill-down - W6dtype=bf16m=4096n=2112k=7168
ds-v3-prefill-gate-up - W7dtype=bf16m=4096n=7168k=16384
k-dominant-7168x16384 - W8dtype=bf16m=16n=4096k=7168
mid-m16-attn - W9dtype=bf16m=32n=4096k=7168
mid-m32-attn - W10dtype=bf16m=64n=7168k=2048
mid-m64-down - W11dtype=bf16m=96n=2112k=7168
mid-m96-gate-up - W12dtype=bf16m=1024n=1024k=1024trans_b=false
square-1k-nn - W13dtype=f16m=1024n=1024k=1024trans_b=false
square-1k-nn - W14dtype=bf16m=4096n=24576k=1536
wide-n-24576
| Workload | Ratio | Device time | Alternatives | Throughput | SOL | Bound | |
|---|---|---|---|---|---|---|---|
| alt / ours | ms | name · tier | ms | TFLOP/s | of ceiling | by | |
| W1 | 1.02× | 0.0789 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.0810 0.0592 0.0633 | · | · | · |
| W2 | 0.89× | 0.0824 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.0709 0.0625 0.0656 | · | · | · |
| W3 | 1.01× | 0.8590 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.8554 0.8548 0.8542 | · | · | · |
| W4 | 1.00× | 0.8508 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.8489 0.8466 0.8485 | · | · | · |
| W5 | 0.84× | 0.5411 | torch_compile_aclgraph torch_compile basistorch_compile_ge torch_compiletorch_npu eager vendor library | 0.4534 0.4540 0.4540 | · | · | · |
| W6 | 1.04× | 0.5171 | torch_compile_aclgraph torch_compile basistorch_compile_ge torch_compiletorch_npu eager vendor library | 0.5327 0.5335 0.5331 | · | · | · |
| W7 | 0.70× | 5.0534 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 3.5349 3.5321 3.5321 | · | · | · |
| W8 | 0.99× | 0.1065 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.0995 0.0995 0.0994 | · | · | · |
| W9 | 1.03× | 0.1095 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.1074 0.1075 0.1074 | · | · | · |
| W10 | 0.96× | 0.0712 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.0568 0.0568 0.0568 | · | · | · |
| W11 | 0.81× | 0.0865 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.0628 0.0630 0.0625 | · | · | · |
| W12 | 0.78× | 0.0381 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.0290 0.0185 0.0236 | · | · | · |
| W13 | 0.74× | 0.0386 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 0.0290 0.0184 0.0232 | · | · | · |
| W14 | 0.73× | 1.6672 | torch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis | 1.2080 1.2071 1.2064 | · | · | · |
DualGemmFwd¶
- W1dtype=bf16m=128n=7168k=2048trans_b=true
ds-v3-decode-down - W2dtype=bf16m=128n=2112k=7168trans_b=true
ds-v3-decode-gate-up - W3dtype=bf16m=4096n=4096k=7168trans_b=true
ds-v3-prefill-attn-proj - W4dtype=f16m=4096n=4096k=7168trans_b=true
ds-v3-prefill-attn-proj - W5dtype=bf16m=4096n=7168k=2048trans_b=true
ds-v3-prefill-down - W6dtype=bf16m=4096n=2112k=7168trans_b=true
ds-v3-prefill-gate-up - W7dtype=bf16m=4096n=7168k=16384trans_b=true
k-dominant-7168x16384 - W8dtype=bf16m=16n=4096k=7168trans_b=true
mid-m16-attn - W9dtype=bf16m=32n=4096k=7168trans_b=true
mid-m32-attn - W10dtype=bf16m=64n=7168k=2048trans_b=true
mid-m64-down - W11dtype=bf16m=96n=2112k=7168trans_b=true
mid-m96-gate-up - W12dtype=bf16m=1024n=1024k=1024
square-1k-nn - W13dtype=f16m=1024n=1024k=1024
square-1k-nn - W14dtype=bf16m=4096n=24576k=1536trans_b=true
wide-n-24576
| Workload | Ratio | Device time | Alternatives | Throughput | SOL | Bound | |
|---|---|---|---|---|---|---|---|
| alt / ours | ms | name · tier | ms | TFLOP/s | of ceiling | by | |
| W1 | 0.87× | 0.1237 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.1074 0.1046 | · | · | · |
| W2 | 0.81× | 0.1401 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.1130 0.1110 | · | · | · |
| W3 | 1.01× | 1.6120 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 1.6186 1.6170 | · | · | · |
| W4 | 1.01× | 1.6231 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 1.6374 1.6320 | · | · | · |
| W5 | 0.87× | 0.9234 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.8027 0.7984 | · | · | · |
| W6 | 1.02× | 0.9414 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 0.9561 0.9766 | · | · | · |
| W7 | 0.73× | 9.3761 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 6.8450 6.8978 | · | · | · |
| W8 | 0.92× | 0.2335 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 0.2132 0.2140 | · | · | · |
| W9 | 0.83× | 0.2385 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.1945 0.1943 | · | · | · |
| W10 | 0.88× | 0.1128 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0985 0.0985 | · | · | · |
| W11 | 0.76× | 0.1497 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 0.1120 0.1121 | · | · | · |
| W12 | 0.73× | 0.0462 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0340 0.0315 | · | · | · |
| W13 | 0.71× | 0.0473 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0338 0.0312 | · | · | · |
| W14 | 0.78× | 2.5046 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 1.9526 1.9511 | · | · | · |
StridedBatchedGemmFwd¶
- W1dtype=bf16b=64m=128n=128k=2048
mha-decode-b64-pv - W2dtype=f16b=64m=128n=128k=2048
mha-decode-b64-pv - W3dtype=bf16b=64m=128n=2048k=128
mha-decode-b64-qk - W4dtype=f16b=64m=128n=2048k=128
mha-decode-b64-qk - W5dtype=bf16b=128m=512n=512k=2048
moe-prefill-b128 - W6dtype=bf16b=8m=128n=128k=128
small-b8-128 - W7dtype=f16b=8m=128n=128k=128
small-b8-128 - W8dtype=bf16b=16m=512n=512k=512
square-b16-512 - W9dtype=f16b=16m=512n=512k=512
square-b16-512 - W10dtype=bf16b=32m=256n=256k=256
square-b32-256 - W11dtype=f16b=32m=256n=256k=256
square-b32-256 - W12dtype=bf16b=4m=4096n=4096k=4096
square-b4-4k - W13dtype=bf16b=8m=1024n=1024k=1024
square-b8-1k - W14dtype=f16b=8m=1024n=1024k=1024
square-b8-1k - W15dtype=bf16b=8m=2048n=2048k=2048
square-b8-2k - W16dtype=f16b=8m=2048n=2048k=2048
square-b8-2k
| Workload | Ratio | Device time | Alternatives | Throughput | SOL | Bound | |
|---|---|---|---|---|---|---|---|
| alt / ours | ms | name · tier | ms | TFLOP/s | of ceiling | by | |
| W1 | 0.73× | 0.1310 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 0.0950 0.0953 | · | · | · |
| W2 | 0.75× | 0.1285 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0949 0.0948 | · | · | · |
| W3 | 0.70× | 0.0885 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0605 0.0605 | · | · | · |
| W4 | 0.69× | 0.0895 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0610 0.0610 | · | · | · |
| W5 | 0.97× | 0.8046 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 0.7837 0.7840 | · | · | · |
| W6 | 1.02× | 0.0100 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0103 0.00575 | · | · | · |
| W7 | 1.08× | 0.00975 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0105 0.00575 | · | · | · |
| W8 | 0.75× | 0.0498 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0370 0.0355 | · | · | · |
| W9 | 0.77× | 0.0493 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0377 0.0360 | · | · | · |
| W10 | 0.65× | 0.0275 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0177 0.0173 | · | · | · |
| W11 | 0.65× | 0.0283 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0192 0.0168 | · | · | · |
| W12 | 1.00× | 1.8698 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 1.8770 1.8784 | · | · | · |
| W13 | 0.84× | 0.1096 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0915 0.0899 | · | · | · |
| W14 | 0.83× | 0.1095 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0912 0.0907 | · | · | · |
| W15 | 0.91× | 0.5674 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.5180 0.5174 | · | · | · |
| W16 | 1.00× | 0.5647 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 0.5646 0.5653 | · | · | · |
GemmSplitKFwd¶
- W1dtype=bf16m=128n=7168k=2048
ds-v3-decode-down - W2dtype=bf16m=128n=2112k=7168
ds-v3-decode-gate-up - W3dtype=bf16m=4096n=4096k=7168
ds-v3-prefill-attn-proj - W4dtype=f16m=4096n=4096k=7168
ds-v3-prefill-attn-proj - W5dtype=bf16m=4096n=7168k=2048
ds-v3-prefill-down - W6dtype=bf16m=4096n=2112k=7168
ds-v3-prefill-gate-up - W7dtype=bf16m=4096n=7168k=16384
k-dominant-7168x16384 - W8dtype=bf16m=16n=4096k=7168
mid-m16-attn - W9dtype=bf16m=32n=4096k=7168
mid-m32-attn - W10dtype=bf16m=64n=7168k=2048
mid-m64-down - W11dtype=bf16m=96n=2112k=7168
mid-m96-gate-up - W12dtype=bf16m=1024n=1024k=1024trans_b=false
square-1k-nn - W13dtype=f16m=1024n=1024k=1024trans_b=false
square-1k-nn - W14dtype=bf16m=4096n=24576k=1536
wide-n-24576
| Workload | Ratio | Device time | Alternatives | Throughput | SOL | Bound | |
|---|---|---|---|---|---|---|---|
| alt / ours | ms | name · tier | ms | TFLOP/s | of ceiling | by | |
| W1 | 0.81× | 0.0661 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0540 0.0525 | · | · | · |
| W2 | 0.80× | 0.0775 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0620 0.0595 | · | · | · |
| W3 | 1.01× | 0.8244 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.8304 0.8249 | · | · | · |
| W4 | 1.00× | 0.8383 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.8454 0.8387 | · | · | · |
| W5 | 0.85× | 0.4841 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.4113 0.4091 | · | · | · |
| W6 | 1.05× | 0.4938 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.5186 0.5138 | · | · | · |
| W7 | 0.72× | 4.8048 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 3.4614 3.4606 | · | · | · |
| W8 | 0.85× | 0.1128 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 0.0953 0.0956 | · | · | · |
| W9 | 0.90× | 0.1139 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.1006 0.1005 | · | · | · |
| W10 | 0.81× | 0.0648 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0515 0.0515 | · | · | · |
| W11 | 0.76× | 0.0772 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0565 0.0565 | · | · | · |
| W12 | 0.62× | 0.0295 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0185 0.0168 | · | · | · |
| W13 | 0.65× | 0.0280 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0182 0.0170 | · | · | · |
| W14 | 0.78× | 1.2536 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 0.9785 0.9786 | · | · | · |
GemmFwd¶
- W1dtype=bf16m=128n=7168k=2048
ds-v3-decode-down - W2dtype=bf16m=128n=2112k=7168
ds-v3-decode-gate-up - W3dtype=bf16m=4096n=4096k=7168
ds-v3-prefill-attn-proj - W4dtype=f16m=4096n=4096k=7168
ds-v3-prefill-attn-proj - W5dtype=bf16m=4096n=7168k=2048
ds-v3-prefill-down - W6dtype=bf16m=4096n=2112k=7168
ds-v3-prefill-gate-up - W7dtype=bf16m=4096n=7168k=16384
k-dominant-7168x16384 - W8dtype=bf16m=16n=4096k=7168
mid-m16-attn - W9dtype=bf16m=32n=4096k=7168
mid-m32-attn - W10dtype=bf16m=64n=7168k=2048
mid-m64-down - W11dtype=bf16m=96n=2112k=7168
mid-m96-gate-up - W12dtype=bf16m=1024n=1024k=1024trans_b=false
square-1k-nn - W13dtype=f16m=1024n=1024k=1024trans_b=false
square-1k-nn - W14dtype=bf16m=4096n=24576k=1536
wide-n-24576
| Workload | Ratio | Device time | Alternatives | Throughput | SOL | Bound | |
|---|---|---|---|---|---|---|---|
| alt / ours | ms | name · tier | ms | TFLOP/s | of ceiling | by | |
| W1 | 0.82× | 0.0670 | catlass-r252:r252_gemm open-source ⓘops-nn-gemm:aclnnMm vendor library basistilelang-ascend tilelang-ascend ⓘtorch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library ⓘ | 0.0565 0.0535 0.1031 0.0558 0.0545 0.0535 | · | · | · |
| W2 | 0.83× | 0.0721 | catlass-r252:r252_gemm open-source ⓘops-nn-gemm:aclnnMm vendor librarytilelang-ascend tilelang-ascend ⓘtorch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis ⓘ | 0.0720 0.0575 0.1194 0.0600 0.0576 0.0574 | · | · | · |
| W3 | 1.02× | 0.8204 | catlass-r252:r252_gemm open-source ⓘops-nn-gemm:aclnnMm open-source ⓘtilelang-ascend tilelang-ascend ⓘtorch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis ⓘ | 0.8364 0.8311 3.1139 0.8310 0.8294 0.8296 | · | · | · |
| W4 | 1.00× | 0.8433 | catlass-r252:r252_gemm open-source ⓘops-nn-gemm:aclnnMm open-source ⓘtilelang-ascend tilelang-ascend ⓘtorch_compile_aclgraph torch_compile basistorch_compile_ge torch_compiletorch_npu eager vendor library ⓘ | 0.8515 0.8383 3.1155 0.8377 0.8351 0.8379 | · | · | · |
| W5 | 0.85× | 0.4728 | catlass-r252:r252_gemm open-source ⓘ basisops-nn-gemm:aclnnMm open-source ⓘtilelang-ascend tilelang-ascend ⓘtorch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library ⓘ | 0.3954 0.4088 1.5580 0.4089 0.4088 0.4088 | · | · | · |
| W6 | 0.92× | 0.4936 | catlass-r252:r252_gemm open-source ⓘ basisops-nn-gemm:aclnnMm open-source ⓘtilelang-ascend tilelang-ascend ⓘtorch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library ⓘ | 0.4481 0.5140 1.9768 0.5122 0.5121 0.5132 | · | · | · |
| W7 | 0.75× | 4.6385 | catlass-r252:r252_gemm open-source ⓘops-nn-gemm:aclnnMm open-source ⓘ basistilelang-ascend tilelang-ascend ⓘtorch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library ⓘ | 4.8831 3.4577 13.6544 3.4600 3.4590 3.4599 | · | · | · |
| W8 | 0.92× | 0.1035 | catlass-r252:r252_gemm open-source ⓘops-nn-gemm:aclnnMm vendor librarytilelang-ascend tilelang-ascend ⓘtorch_compile_aclgraph torch_compile basistorch_compile_ge torch_compiletorch_npu eager vendor library ⓘ | 0.1008 0.0943 0.1630 0.0940 0.0941 0.0943 | · | · | · |
| W9 | 0.92× | 0.1128 | catlass-r252:r252_gemm open-source ⓘops-nn-gemm:aclnnMm vendor librarytilelang-ascend tilelang-ascend ⓘtorch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis ⓘ | 0.1106 0.1015 0.1772 0.1013 0.1015 0.1011 | · | · | · |
| W10 | 0.80× | 0.0622 | catlass-r252:r252_gemm open-source ⓘops-nn-gemm:aclnnMm vendor library basistilelang-ascend tilelang-ascend ⓘtorch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library ⓘ | 0.0515 0.0489 0.0823 0.0490 0.0489 0.0490 | · | · | · |
| W11 | 0.76× | 0.0747 | catlass-r252:r252_gemm open-source ⓘops-nn-gemm:aclnnMm vendor librarytilelang-ascend tilelang-ascend ⓘtorch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library basis ⓘ | 0.0705 0.0548 0.1030 0.0548 0.0545 0.0545 | · | · | · |
| W12 | 0.63× | 0.0285 | catlass-r252:r252_gemm open-source ⓘops-nn-gemm:aclnnMm vendor library basistilelang-ascend tilelang-ascend ⓘtorch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor librarytilelang-mlir-ascend:example_gemm tilelang-mlir-ascend ⓘ eager · n=5 ⓘ | 0.0187 0.0160 0.0173 0.0177 0.0168 0.0160 0.0449 | · | · | · |
| W13 | 0.62× | 0.0293 | catlass:catlass_gemm_fp16 open-source ⓘcatlass-r252:r252_gemm open-source ⓘops-nn-gemm:aclnnMm vendor library basistilelang-ascend tilelang-ascend ⓘtorch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor librarytilelang-mlir-ascend:example_gemm tilelang-mlir-ascend ⓘ eager · n=5 ⓘ | 0.0192 0.0187 0.0164 0.0175 0.0182 0.0165 0.0165 0.0468 | · | · | · |
| W14 | 0.78× | 1.2525 | catlass-r252:r252_gemm open-source ⓘops-nn-gemm:aclnnMm open-source ⓘ basistilelang-ascend tilelang-ascend ⓘtorch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library ⓘ | 1.0828 0.9759 4.0015 0.9764 0.9772 0.9774 | · | · | · |
GemmReduceScatterFwd¶
- W1dtype=bf16k=2048m=128n=7168trans_b=true
ds-v3-decode-down@rank0 - W2dtype=bf16k=2048m=128n=7168trans_b=true
ds-v3-decode-down@rank1 - W3dtype=bf16k=7168m=128n=2112trans_b=true
ds-v3-decode-gate-up@rank0 - W4dtype=bf16k=7168m=128n=2112trans_b=true
ds-v3-decode-gate-up@rank1 - W5dtype=bf16k=7168m=4096n=4096trans_b=true
ds-v3-prefill-attn-proj@rank0 - W6dtype=f16k=7168m=4096n=4096trans_b=true
ds-v3-prefill-attn-proj@rank0 - W7dtype=bf16k=7168m=4096n=4096trans_b=true
ds-v3-prefill-attn-proj@rank1 - W8dtype=f16k=7168m=4096n=4096trans_b=true
ds-v3-prefill-attn-proj@rank1 - W9dtype=bf16k=2048m=4096n=7168trans_b=true
ds-v3-prefill-down@rank0 - W10dtype=bf16k=2048m=4096n=7168trans_b=true
ds-v3-prefill-down@rank1 - W11dtype=bf16k=7168m=4096n=2112trans_b=true
ds-v3-prefill-gate-up@rank0 - W12dtype=bf16k=7168m=4096n=2112trans_b=true
ds-v3-prefill-gate-up@rank1 - W13dtype=bf16k=16384m=4096n=7168trans_b=true
k-dominant-7168x16384@rank0 - W14dtype=bf16k=16384m=4096n=7168trans_b=true
k-dominant-7168x16384@rank1 - W15dtype=bf16k=7168m=16n=4096trans_b=true
mid-m16-attn@rank0 - W16dtype=bf16k=7168m=16n=4096trans_b=true
mid-m16-attn@rank1 - W17dtype=bf16k=7168m=32n=4096trans_b=true
mid-m32-attn@rank0 - W18dtype=bf16k=7168m=32n=4096trans_b=true
mid-m32-attn@rank1 - W19dtype=bf16k=2048m=64n=7168trans_b=true
mid-m64-down@rank0 - W20dtype=bf16k=2048m=64n=7168trans_b=true
mid-m64-down@rank1 - W21dtype=bf16k=7168m=96n=2112trans_b=true
mid-m96-gate-up@rank0 - W22dtype=bf16k=7168m=96n=2112trans_b=true
mid-m96-gate-up@rank1 - W23dtype=bf16k=1024m=1024n=1024
square-1k-nn@rank0 - W24dtype=f16k=1024m=1024n=1024
square-1k-nn@rank0 - W25dtype=bf16k=1024m=1024n=1024
square-1k-nn@rank1 - W26dtype=f16k=1024m=1024n=1024
square-1k-nn@rank1 - W27dtype=bf16k=1536m=4096n=24576trans_b=true
wide-n-24576@rank0 - W28dtype=bf16k=1536m=4096n=24576trans_b=true
wide-n-24576@rank1
| Workload | Ratio | Device time | Alternatives | Throughput | SOL | Bound | |
|---|---|---|---|---|---|---|---|
| alt / ours | ms | name | ms | TFLOP/s | of ceiling | by | |
| W1 | 0.87× | 0.4290 | Torch-NPU semantic implementation | 0.3713 | · | · | · |
| W2 | 0.87× | 0.4259 | Torch-NPU semantic implementation | 0.3686 | · | · | · |
| W3 | 0.60× | 0.3752 | Torch-NPU semantic implementation | 0.2235 | · | · | · |
| W4 | 0.59× | 0.3794 | Torch-NPU semantic implementation | 0.2254 | · | · | · |
| W5 | 0.38× | 17.6604 | Torch-NPU semantic implementation | 6.6771 | · | · | · |
| W6 | 0.38× | 17.6441 | Torch-NPU semantic implementation | 6.6766 | · | · | · |
| W7 | 0.38× | 17.6571 | Torch-NPU semantic implementation | 6.6735 | · | · | · |
| W8 | 0.38× | 17.6532 | Torch-NPU semantic implementation | 6.6738 | · | · | · |
| W9 | 0.69× | 11.3592 | Torch-NPU semantic implementation | 7.8670 | · | · | · |
| W10 | 0.69× | 11.3441 | Torch-NPU semantic implementation | 7.8548 | · | · | · |
| W11 | 0.40× | 9.0945 | Torch-NPU semantic implementation | 3.6114 | · | · | · |
| W12 | 0.40× | 9.0930 | Torch-NPU semantic implementation | 3.6069 | · | · | · |
| W13 | 0.26× | 73.3225 | Torch-NPU semantic implementation | 19.0761 | · | · | · |
| W14 | 0.26× | 73.2879 | Torch-NPU semantic implementation | 19.0469 | · | · | · |
| W15 | 1.35× | 0.2186 | Torch-NPU semantic implementation | 0.2949 | · | · | · |
| W16 | 1.35× | 0.2171 | Torch-NPU semantic implementation | 0.2930 | · | · | · |
| W17 | 1.19× | 0.2590 | Torch-NPU semantic implementation | 0.3079 | · | · | · |
| W18 | 1.20× | 0.2308 | Torch-NPU semantic implementation | 0.2770 | · | · | · |
| W19 | 1.01× | 0.2522 | Torch-NPU semantic implementation | 0.2542 | · | · | · |
| W20 | 1.01× | 0.2525 | Torch-NPU semantic implementation | 0.2549 | · | · | · |
| W21 | 0.67× | 0.3191 | Torch-NPU semantic implementation | 0.2137 | · | · | · |
| W22 | 0.67× | 0.3197 | Torch-NPU semantic implementation | 0.2127 | · | · | · |
| W23 | 0.96× | 0.3449 | Torch-NPU semantic implementation | 0.3295 | · | · | · |
| W24 | 0.96× | 0.3307 | Torch-NPU semantic implementation | 0.3171 | · | · | · |
| W25 | 0.96× | 0.3270 | Torch-NPU semantic implementation | 0.3137 | · | · | · |
| W26 | 0.95× | 0.3291 | Torch-NPU semantic implementation | 0.3135 | · | · | · |
| W27 | 0.80× | 31.9567 | Torch-NPU semantic implementation | 25.6279 | · | · | · |
| W28 | 0.80× | 31.9584 | Torch-NPU semantic implementation | 25.6348 | · | · | · |
GroupedGemmFwd¶
batch_sum=4096batch_count=16n=4096k=4096
- W1dtype=f16transpose_b=false
nn-batch16-m4096-n4096-k4096 - W2dtype=bf16
nt-batch16-m4096-n4096-k4096 - W3dtype=f16
nt-batch16-m4096-n4096-k4096 - W4dtype=f16transpose_a=truetranspose_b=false
tn-batch16-m4096-n4096-k4096
| Workload | Ratio | Device time | Alternatives | Throughput | SOL | Bound | |
|---|---|---|---|---|---|---|---|
| alt / ours | ms | name · tier | ms | TFLOP/s | of ceiling | by | |
| W1 | 0.63× | 1.1623 | catlass-r252:r252_gmm open-source ⓘ basistorch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library ⓘ | 0.7312 1.0529 1.4347 1.0926 | · | · | · |
| W2 | 0.63× | 1.1694 | catlass-r252:r252_gmm open-source ⓘ basistorch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library ⓘ | 0.7260 0.9878 1.3709 1.0328 | · | · | · |
| W3 | 0.63× | 1.1579 | catlass-r252:r252_gmm open-source ⓘ basistorch_compile_aclgraph torch_compiletorch_compile_ge torch_compiletorch_npu eager vendor library ⓘ | 0.7200 0.9752 1.3673 1.0191 | · | · | · |
| W4 | 0.68× | 2.2741 | torch_compile_aclgraph torch_compile basistorch_compile_ge torch_compiletorch_npu eager vendor library ⓘ | 1.5355 0.7739 1.5376 | · | · | · |
GemmSparse2to4Fwd¶
- W1dtype=bf16m=128n=7168k=2048trans_b=true
ds-v3-decode-down - W2dtype=bf16m=128n=2112k=7168trans_b=true
ds-v3-decode-gate-up - W3dtype=bf16m=4096n=4096k=7168trans_b=true
ds-v3-prefill-attn-proj - W4dtype=f16m=4096n=4096k=7168trans_b=true
ds-v3-prefill-attn-proj - W5dtype=bf16m=4096n=7168k=2048trans_b=true
ds-v3-prefill-down - W6dtype=bf16m=4096n=2112k=7168trans_b=true
ds-v3-prefill-gate-up - W7dtype=bf16m=4096n=7168k=16384trans_b=true
k-dominant-7168x16384 - W8dtype=bf16m=16n=4096k=7168trans_b=true
mid-m16-attn - W9dtype=bf16m=32n=4096k=7168trans_b=true
mid-m32-attn - W10dtype=bf16m=64n=7168k=2048trans_b=true
mid-m64-down - W11dtype=bf16m=96n=2112k=7168trans_b=true
mid-m96-gate-up - W12dtype=bf16m=1024n=1024k=1024
square-1k-nn - W13dtype=f16m=1024n=1024k=1024
square-1k-nn - W14dtype=bf16m=4096n=24576k=1536trans_b=true
wide-n-24576
| Workload | Ratio | Device time | Alternatives | Throughput | SOL | Bound | |
|---|---|---|---|---|---|---|---|
| alt / ours | ms | name · tier | ms | TFLOP/s | of ceiling | by | |
| W1 | 0.63× | 0.3563 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.2205 0.1933 | · | · | · |
| W2 | 0.28× | 1.1927 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.3330 0.2950 | · | · | · |
| W3 | 0.23× | 36.0448 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 8.2523 8.2260 | · | · | · |
| W4 | 0.23× | 36.0423 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 8.2621 8.2255 | · | · | · |
| W5 | 0.32× | 9.3550 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 2.9786 2.9536 | · | · | · |
| W6 | 0.19× | 35.7067 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 6.9085 6.8755 | · | · | · |
| W7 | 0.29× | 85.3295 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 24.6160 24.6036 | · | · | · |
| W8 | 1.10× | 0.2732 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.2747 0.2744 | · | · | · |
| W9 | 0.68× | 0.4065 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 0.2626 0.2627 | · | · | · |
| W10 | 0.85× | 0.2097 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.1509 0.1507 | · | · | · |
| W11 | 0.30× | 0.9061 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 0.2606 0.2607 | · | · | · |
| W12 | 0.24× | 1.1515 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.2747 0.2502 | · | · | · |
| W13 | 0.24× | 1.1559 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.2767 0.2505 | · | · | · |
| W14 | 0.69× | 7.7070 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 5.3418 5.3405 | · | · | · |
GemmInt8Fwd¶
- W1dtype=bf16m=128n=7168k=2048trans_b=true
ds-v3-decode-down - W2dtype=bf16m=128n=2112k=7168trans_b=true
ds-v3-decode-gate-up - W3dtype=bf16m=4096n=4096k=7168trans_b=true
ds-v3-prefill-attn-proj - W4dtype=f16m=4096n=4096k=7168trans_b=true
ds-v3-prefill-attn-proj - W5dtype=bf16m=4096n=7168k=2048trans_b=true
ds-v3-prefill-down - W6dtype=bf16m=4096n=2112k=7168trans_b=true
ds-v3-prefill-gate-up - W7dtype=bf16m=4096n=7168k=16384trans_b=true
k-dominant-7168x16384 - W8dtype=bf16m=16n=4096k=7168trans_b=true
mid-m16-attn - W9dtype=bf16m=32n=4096k=7168trans_b=true
mid-m32-attn - W10dtype=bf16m=64n=7168k=2048trans_b=true
mid-m64-down - W11dtype=bf16m=96n=2112k=7168trans_b=true
mid-m96-gate-up - W12dtype=bf16m=1024n=1024k=1024
square-1k-nn - W13dtype=f16m=1024n=1024k=1024
square-1k-nn - W14dtype=bf16m=4096n=24576k=1536trans_b=true
wide-n-24576
| Workload | Ratio | Device time | Alternatives | Throughput | SOL | Bound | |
|---|---|---|---|---|---|---|---|
| alt / ours | ms | name · tier | ms | TFLOP/s | of ceiling | by | |
| W1 | 0.48× | 0.3344 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.1593 0.1214 | · | · | · |
| W2 | 0.42× | 0.3334 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.1396 0.1123 | · | · | · |
| W3 | 0.19× | 15.8272 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 3.0452 3.0318 | · | · | · |
| W4 | 0.19× | 15.8409 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 3.0438 3.0275 | · | · | · |
| W5 | 0.19× | 8.6520 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 1.6146 1.6729 | · | · | · |
| W6 | 0.21× | 8.1717 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 1.7127 1.7018 | · | · | · |
| W7 | 0.21× | 61.1816 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 12.8299 12.7981 | · | · | · |
| W8 | 1.19× | 0.1844 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 0.2042 0.2045 | · | · | · |
| W9 | 0.93× | 0.1983 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 0.1700 0.1706 | · | · | · |
| W10 | 0.69× | 0.1903 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.1036 0.1036 | · | · | · |
| W11 | 0.51× | 0.2854 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.1284 0.1281 | · | · | · |
| W12 | 0.33× | 0.2530 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0828 0.0646 | · | · | · |
| W13 | 0.32× | 0.2500 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0805 0.0648 | · | · | · |
| W14 | 0.24× | 22.5699 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 5.3871 5.3886 | · | · | · |
GemmA16W8Fwd¶
- W1dtype=bf16m=128n=7168k=2048trans_b=true
ds-v3-decode-down - W2dtype=bf16m=128n=2112k=7168trans_b=true
ds-v3-decode-gate-up - W3dtype=bf16m=4096n=4096k=7168trans_b=true
ds-v3-prefill-attn-proj - W4dtype=f16m=4096n=4096k=7168trans_b=true
ds-v3-prefill-attn-proj - W5dtype=bf16m=4096n=7168k=2048trans_b=true
ds-v3-prefill-down - W6dtype=bf16m=4096n=2112k=7168trans_b=true
ds-v3-prefill-gate-up - W7dtype=bf16m=4096n=7168k=16384trans_b=true
k-dominant-7168x16384 - W8dtype=bf16m=16n=4096k=7168trans_b=true
mid-m16-attn - W9dtype=bf16m=32n=4096k=7168trans_b=true
mid-m32-attn - W10dtype=bf16m=64n=7168k=2048trans_b=true
mid-m64-down - W11dtype=bf16m=96n=2112k=7168trans_b=true
mid-m96-gate-up - W12dtype=bf16m=1024n=1024k=1024
square-1k-nn - W13dtype=f16m=1024n=1024k=1024
square-1k-nn - W14dtype=bf16m=4096n=24576k=1536trans_b=true
wide-n-24576
| Workload | Ratio | Device time | Alternatives | Throughput | SOL | Bound | |
|---|---|---|---|---|---|---|---|
| alt / ours | ms | name · tier | ms | TFLOP/s | of ceiling | by | |
| W1 | 0.22× | 0.6624 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.1521 0.1268 | · | · | · |
| W2 | 0.24× | 0.5459 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.1315 0.1159 | · | · | · |
| W3 | 0.18× | 16.7480 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 3.0372 3.0380 | · | · | · |
| W4 | 0.19× | 16.3590 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 3.0372 3.0412 | · | · | · |
| W5 | 0.17× | 8.9730 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 1.5276 1.7535 | · | · | · |
| W6 | 0.20× | 8.4763 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 1.7203 1.7145 | · | · | · |
| W7 | 0.17× | 75.0704 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 12.8368 12.7846 | · | · | · |
| W8 | 0.26× | 0.8434 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.2011 0.2010 | · | · | · |
| W9 | 0.21× | 0.8475 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 0.1680 0.1681 | · | · | · |
| W10 | 0.24× | 0.5192 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.1019 0.1017 | · | · | · |
| W11 | 0.25× | 0.4860 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 0.1172 0.1174 | · | · | · |
| W12 | 0.28× | 0.2704 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0725 0.0649 | · | · | · |
| W13 | 0.31× | 0.2389 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0730 0.0654 | · | · | · |
| W14 | 0.21× | 23.1819 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 4.7480 4.7484 | · | · | · |
GemmEpilogueSwiGLUFwd¶
- W1dtype=bf16m=128n=7168k=2048trans_b=true
ds-v3-decode-down - W2dtype=bf16m=128n=2112k=7168trans_b=true
ds-v3-decode-gate-up - W3dtype=bf16m=4096n=4096k=7168trans_b=true
ds-v3-prefill-attn-proj - W4dtype=f16m=4096n=4096k=7168trans_b=true
ds-v3-prefill-attn-proj - W5dtype=bf16m=4096n=7168k=2048trans_b=true
ds-v3-prefill-down - W6dtype=bf16m=4096n=2112k=7168trans_b=true
ds-v3-prefill-gate-up - W7dtype=bf16m=4096n=7168k=16384trans_b=true
k-dominant-7168x16384 - W8dtype=bf16m=16n=4096k=7168trans_b=true
mid-m16-attn - W9dtype=bf16m=32n=4096k=7168trans_b=true
mid-m32-attn - W10dtype=bf16m=64n=7168k=2048trans_b=true
mid-m64-down - W11dtype=bf16m=96n=2112k=7168trans_b=true
mid-m96-gate-up - W12dtype=bf16m=1024n=1024k=1024
square-1k-nn - W13dtype=f16m=1024n=1024k=1024
square-1k-nn - W14dtype=bf16m=4096n=24576k=1536trans_b=true
wide-n-24576
| Workload | Ratio | Device time | Alternatives | Throughput | SOL | Bound | |
|---|---|---|---|---|---|---|---|
| alt / ours | ms | name · tier | ms | TFLOP/s | of ceiling | by | |
| W1 | 0.18× | 1.5390 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.2765 0.2531 | · | · | · |
| W2 | 0.19× | 1.5446 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.2980 0.2884 | · | · | · |
| W3 | 0.10× | 63.5631 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 6.0729 6.1739 | · | · | · |
| W4 | 0.10× | 63.5561 | torch_npu eager vendor library basis | 6.1647 | · | · | · |
| W5 | 0.11× | 33.2796 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 3.7622 3.8369 | · | · | · |
| W6 | 0.12× | 26.1990 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 3.1896 3.2142 | · | · | · |
| W7 | 0.10× | 263.9803 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 26.0779 26.0617 | · | · | · |
| W8 | 0.56× | 1.0939 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 0.5913 0.5920 | · | · | · |
| W9 | 0.49× | 1.2266 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 0.6012 0.6020 | · | · | · |
| W10 | 0.28× | 0.8626 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.2235 0.2234 | · | · | · |
| W11 | 0.19× | 1.3798 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 0.2585 0.2587 | · | · | · |
| W12 | 0.22× | 0.5600 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.1236 0.1163 | · | · | · |
| W13 | 0.22× | 0.5586 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.1222 0.1135 | · | · | · |
| W14 | 0.13× | 79.9905 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 10.8063 10.8057 | · | · | · |
GemmBlockScaledFwd¶
- W1dtype=bf16m=128n=7168k=2048trans_b=true
ds-v3-decode-down - W2dtype=bf16m=128n=2112k=7168trans_b=true
ds-v3-decode-gate-up - W3dtype=bf16m=4096n=4096k=7168trans_b=true
ds-v3-prefill-attn-proj - W4dtype=f16m=4096n=4096k=7168trans_b=true
ds-v3-prefill-attn-proj - W5dtype=bf16m=4096n=7168k=2048trans_b=true
ds-v3-prefill-down - W6dtype=bf16m=4096n=2112k=7168trans_b=true
ds-v3-prefill-gate-up - W7dtype=bf16m=4096n=7168k=16384trans_b=true
k-dominant-7168x16384 - W8dtype=bf16m=16n=4096k=7168trans_b=true
mid-m16-attn - W9dtype=bf16m=32n=4096k=7168trans_b=true
mid-m32-attn - W10dtype=bf16m=64n=7168k=2048trans_b=true
mid-m64-down - W11dtype=bf16m=96n=2112k=7168trans_b=true
mid-m96-gate-up - W12dtype=bf16m=1024n=1024k=1024
square-1k-nn - W13dtype=f16m=1024n=1024k=1024
square-1k-nn - W14dtype=bf16m=4096n=24576k=1536trans_b=true
wide-n-24576
| Workload | Ratio | Device time | Alternatives | Throughput | SOL | Bound | |
|---|---|---|---|---|---|---|---|
| alt / ours | ms | name · tier | ms | TFLOP/s | of ceiling | by | |
| W1 | 0.15× | 102.2184 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 15.2870 15.5379 | · | · | · |
| W2 | 0.14× | 114.7999 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 16.4260 16.6836 | · | · | · |
| W3 | 0.01× | 5356.4300 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 62.4492 63.5335 | · | · | · |
| W4 | 0.01× | 5356.4076 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 62.3709 63.4989 | · | · | · |
| W5 | 0.01× | 2714.5082 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 24.6761 24.9730 | · | · | · |
| W6 | 0.02× | 3046.2231 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 46.3914 47.2873 | · | · | · |
| W7 | 0.01× | 21732.5278 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 203.1527 203.2515 | · | · | · |
| W8 | 0.51× | 60.2694 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 30.8980 30.8871 | · | · | · |
| W9 | 0.39× | 79.6118 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 31.0940 31.0756 | · | · | · |
| W10 | 0.25× | 60.7450 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 15.3833 15.3793 | · | · | · |
| W11 | 0.18× | 91.0631 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 16.4288 16.4565 | · | · | · |
| W12 | 0.05× | 50.3330 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 2.6625 2.7111 | · | · | · |
| W13 | 0.05× | 50.3204 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 2.6656 2.6984 | · | · | · |
| W14 | 0.01× | 6907.4176 | torch_compile_aclgraph torch_compile basistorch_npu eager vendor library | 49.5547 49.6379 | · | · | · |
GemmW4A16Fwd¶
| Workload | Ratio | Device time | Alternatives | Throughput | SOL | Bound | |
|---|---|---|---|---|---|---|---|
| alt / ours | ms | name | ms | TFLOP/s | of ceiling | by | |
MoeGroupedGemmFwd¶
| Workload | Ratio | Device time | Alternatives | Throughput | SOL | Bound | |
|---|---|---|---|---|---|---|---|
| alt / ours | ms | name | ms | TFLOP/s | of ceiling | by | |
Quantization¶
FP8QuantFwd¶
input_tensor: [batch, seq_len_kv, kv_group, index_dim]
batch=1kv_group=1
- W1dtype=f32seq_len_kv=4096index_dim=128
kv-index-4k-d128 - W2dtype=bf16seq_len_kv=8192index_dim=64
kv-index-8k-d64 - W3dtype=f16seq_len_kv=8192index_dim=64
kv-index-8k-d64
| Workload | Ratio | Device time | Alternatives | Throughput | SOL | Bound | |
|---|---|---|---|---|---|---|---|
| alt / ours | ms | name · tier | ms | TFLOP/s | of ceiling | by | |
| W1 | 12.05× | 0.0220 | torch_npu eager vendor library basis ⓘ | 0.2650 | · | · | · |
| W2 | 11.60× | 0.0248 | torch_npu eager vendor library basis ⓘ | 0.2870 | · | · | · |
| W3 | 12.60× | 0.0222 | torch_npu eager vendor library basis ⓘ | 0.2802 | · | · | · |
DequantizeFwd¶
input_tensor: [batch, seq_len_kv, kv_group, index_dim]
batch=1kv_group=1
- W1dtype=f32seq_len_kv=4096index_dim=128
kv-index-4k-d128 - W2dtype=bf16seq_len_kv=8192index_dim=64
kv-index-8k-d64 - W3dtype=f16seq_len_kv=8192index_dim=64
kv-index-8k-d64
| Workload | Ratio | Device time | Alternatives | Throughput | SOL | Bound | |
|---|---|---|---|---|---|---|---|
| alt / ours | ms | name · tier | ms | TFLOP/s | of ceiling | by | |
| W1 | 0.99× | 0.0175 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0203 0.0152 | · | · | · |
| W2 | 1.55× | 0.0163 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0253 0.0173 | · | · | · |
| W3 | 1.73× | 0.0168 | torch_compile_aclgraph torch_compiletorch_npu eager vendor library basis | 0.0261 0.0173 | · | · | · |