diff --git a/tests/kernels/quantization/golden/hybrid_w4a16_gfx1151.json b/tests/kernels/quantization/golden/hybrid_w4a16_gfx1151.json index 9a0bef551c28..3cd23c4eb558 100644 --- a/tests/kernels/quantization/golden/hybrid_w4a16_gfx1151.json +++ b/tests/kernels/quantization/golden/hybrid_w4a16_gfx1151.json @@ -13,67 +13,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 1.4128 + "tflops": 1.1559 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.808 + "tflops": 0.8269 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 1.7654 + "tflops": 2.5282 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.4264 + "tflops": 0.8062 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 6.7844 + "tflops": 1.5831 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 12.8477 + "tflops": 3.0903 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 12.6978 + "tflops": 5.9227 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 18.4082 + "tflops": 10.0094 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 15.6968 + "tflops": 15.6352 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 16.3975 + "tflops": 17.9323 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 16.6954 + "tflops": 19.6268 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 21.4047 + "tflops": 26.6381 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 21.6875 + "tflops": 28.5911 } ] }, @@ -83,67 +83,1795 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 1.2419 + "tflops": 1.0299 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7243 + "tflops": 1.9072 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 2.567 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.3091 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 2.4677 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 4.6223 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 8.6363 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 16.937 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 16.872 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 18.4228 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 20.1996 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 27.0369 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 29.0645 + } + ] + }, + { + "provider": "hybrid-w4a16-bf16", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 1.1168 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 2.0441 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 2.586 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 0.892 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.7454 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 3.3339 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 6.3861 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 12.754 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 16.3416 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 17.4334 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 17.5329 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 20.8247 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 22.2488 + } + ] + }, + { + "provider": "hybrid-w4a16-zp-bf16", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 1.0381 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 1.9298 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 2.3942 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 0.8668 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.6879 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 3.2901 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 6.2698 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 12.1989 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 12.8709 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 16.2518 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 15.6363 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 19.0931 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 20.1319 + } + ] + } + ] + }, + { + "in_features": 2048, + "out_features": 2560, + "group_size": 128, + "comment": "gemma-2b qkv_proj", + "providers": [ + { + "provider": "hybrid-w4a16", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 1.2674 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 2.1182 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 2.9598 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 0.8344 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.6471 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 3.1964 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 6.0487 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 11.6793 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 15.8469 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 16.5678 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 21.4805 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 27.1865 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 28.9609 + } + ] + }, + { + "provider": "hybrid-w4a16-zp", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 1.1961 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 1.9056 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 2.8585 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.5451 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 2.878 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 5.3435 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 9.4857 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 18.8774 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 16.5786 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 16.9639 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 22.1413 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 27.1424 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 28.9564 + } + ] + }, + { + "provider": "hybrid-w4a16-bf16", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 1.2847 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 1.9547 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 2.8409 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.0433 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.9714 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 3.7096 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 7.3696 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 14.5038 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 14.9605 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 16.5258 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 19.0271 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 21.7741 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 22.2774 + } + ] + }, + { + "provider": "hybrid-w4a16-zp-bf16", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 1.2086 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 1.9406 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 2.6503 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.025 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.9926 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 3.8434 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 7.3482 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 14.5068 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 13.4321 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 14.4775 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 17.1488 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 19.8503 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 19.8949 + } + ] + } + ] + }, + { + "in_features": 2048, + "out_features": 4096, + "group_size": 128, + "comment": "Qwen3-1.7B qkv_proj", + "providers": [ + { + "provider": "hybrid-w4a16", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 1.0602 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 1.7898 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 3.2288 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.3152 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 2.5115 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 4.8448 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 9.0274 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 13.1714 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 14.663 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 17.9598 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 22.917 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 27.7708 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 28.203 + } + ] + }, + { + "provider": "hybrid-w4a16-zp", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 1.0037 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 1.6677 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 2.8777 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 2.3602 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 4.5198 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 8.3918 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 14.9622 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 14.9472 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 15.9053 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 18.6496 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 23.1686 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 27.6254 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 27.7915 + } + ] + }, + { + "provider": "hybrid-w4a16-bf16", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 1.0629 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 1.7551 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 3.1428 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.6645 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 3.2262 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 6.257 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 12.1619 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 13.9512 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 15.9841 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 16.443 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 19.328 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 21.1399 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 21.7073 + } + ] + }, + { + "provider": "hybrid-w4a16-zp-bf16", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 1.0194 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 1.7098 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 3.0184 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.6122 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 3.1753 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 6.0621 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 11.493 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 10.6608 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 15.3041 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 14.558 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 17.5505 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 19.9005 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 20.1614 + } + ] + } + ] + }, + { + "in_features": 2048, + "out_features": 12288, + "group_size": 128, + "comment": "Qwen3-1.7B gate_up_proj", + "providers": [ + { + "provider": "hybrid-w4a16", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 0.982 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 1.8416 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 3.3548 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.213 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 2.4009 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 4.8468 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 9.3221 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 11.0839 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 14.2381 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 23.205 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 26.5631 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 27.9003 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 28.3698 + } + ] + }, + { + "provider": "hybrid-w4a16-zp", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 0.9442 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 1.7612 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 3.0483 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.1697 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 2.3037 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 4.6095 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 8.9134 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 10.9909 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 14.2394 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 23.1619 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 26.029 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 28.2538 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 28.7351 + } + ] + }, + { + "provider": "hybrid-w4a16-bf16", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 0.9807 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 1.822 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 3.3335 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.4233 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 2.7281 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 5.188 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 9.7842 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 9.5725 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 14.5524 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 19.8716 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 20.924 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 22.0217 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 21.8262 + } + ] + }, + { + "provider": "hybrid-w4a16-zp-bf16", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 0.9521 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 1.7693 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 3.1597 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.0719 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 2.1211 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 4.1848 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 8.1774 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 8.4795 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 13.2832 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 17.6926 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 19.0574 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 20.2834 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 20.2207 + } + ] + } + ] + }, + { + "in_features": 2048, + "out_features": 32768, + "group_size": 128, + "comment": "gemma-2b gate_up_proj", + "providers": [ + { + "provider": "hybrid-w4a16", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 0.9246 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 1.7985 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 3.4215 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.3394 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 2.678 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 5.2058 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 9.6103 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 16.3538 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 21.5046 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 28.049 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 29.9442 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 30.5017 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 29.5685 + } + ] + }, + { + "provider": "hybrid-w4a16-zp", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 0.8924 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 1.7004 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 3.2662 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.2378 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 2.4308 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 4.7817 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 9.0649 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 15.3796 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 21.7906 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 28.9106 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 30.4191 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 30.3716 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 30.4503 + } + ] + }, + { + "provider": "hybrid-w4a16-bf16", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 0.9213 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 1.7896 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 3.4045 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 0.9777 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.9358 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 3.7733 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 7.165 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 11.9042 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 17.3186 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 20.4312 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 20.5412 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 20.8689 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 21.0126 + } + ] + }, + { + "provider": "hybrid-w4a16-zp-bf16", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 0.9023 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 1.7237 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 3.3535 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 0.9296 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.8282 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 3.5961 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 6.9498 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 12.0869 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 16.9059 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 19.6894 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 19.9446 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 19.7747 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 19.8709 + } + ] + } + ] + }, + { + "in_features": 2560, + "out_features": 2560, + "group_size": 128, + "comment": "Qwen3-4B o_proj", + "providers": [ + { + "provider": "hybrid-w4a16", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 1.0688 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 2.0041 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 3.1493 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.3138 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 2.5424 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 4.8358 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 8.9581 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 18.6402 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 25.926 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 28.9225 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 27.7678 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 29.4615 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 29.8245 + } + ] + }, + { + "provider": "hybrid-w4a16-zp", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 0.9766 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 1.8533 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 2.9915 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.681 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 2.9461 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 5.6372 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 10.5582 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 21.7627 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 27.1269 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 29.262 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 27.5948 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 29.4727 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 29.6714 + } + ] + }, + { + "provider": "hybrid-w4a16-bf16", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 1.0598 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 1.9762 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 3.0839 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.1388 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 2.213 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 4.2882 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 8.198 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 16.5417 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 20.829 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 21.7958 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 21.9211 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 22.077 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 22.4176 + } + ] + }, + { + "provider": "hybrid-w4a16-zp-bf16", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 1.0001 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 1.8728 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 3.0412 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.098 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 2.1583 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 4.097 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 7.5094 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 15.287 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 19.242 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 20.8961 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 20.9543 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 20.8801 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 20.8178 + } + ] + } + ] + }, + { + "in_features": 2560, + "out_features": 3840, + "group_size": 128, + "comment": "Qwen3-4B qkv_proj", + "providers": [ + { + "provider": "hybrid-w4a16", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 0.9948 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 1.8961 + }, + { + "batch_size": 4, + "kernel": "wvsplitk_int4", + "tflops": 3.8036 + }, + { + "batch_size": 8, + "kernel": "hybrid_triton_w4a16", + "tflops": 1.9931 + }, + { + "batch_size": 16, + "kernel": "hybrid_triton_w4a16", + "tflops": 3.9479 + }, + { + "batch_size": 32, + "kernel": "hybrid_triton_w4a16", + "tflops": 7.7181 + }, + { + "batch_size": 64, + "kernel": "hybrid_triton_w4a16", + "tflops": 14.6163 + }, + { + "batch_size": 128, + "kernel": "hybrid_triton_w4a16", + "tflops": 21.6956 + }, + { + "batch_size": 256, + "kernel": "hybrid_triton_w4a16", + "tflops": 28.7607 + }, + { + "batch_size": 512, + "kernel": "hybrid_triton_w4a16", + "tflops": 28.9012 + }, + { + "batch_size": 1024, + "kernel": "hybrid_triton_w4a16", + "tflops": 28.6021 + }, + { + "batch_size": 2048, + "kernel": "hybrid_triton_w4a16", + "tflops": 30.4936 + }, + { + "batch_size": 4096, + "kernel": "hybrid_triton_w4a16", + "tflops": 30.9068 + } + ] + }, + { + "provider": "hybrid-w4a16-zp", + "baselines": [ + { + "batch_size": 1, + "kernel": "wvsplitk_int4", + "tflops": 0.9202 + }, + { + "batch_size": 2, + "kernel": "wvsplitk_int4", + "tflops": 1.7588 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 1.7317 + "tflops": 3.4879 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.4744 + "tflops": 2.4667 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 6.8001 + "tflops": 4.4544 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 12.9224 + "tflops": 8.6625 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 12.1375 + "tflops": 16.7537 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 19.0644 + "tflops": 23.1952 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 15.5782 + "tflops": 29.099 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 16.2083 + "tflops": 29.5405 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 16.3604 + "tflops": 28.7058 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 20.9934 + "tflops": 30.4157 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 21.2141 + "tflops": 31.1741 } ] }, @@ -153,67 +1881,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 1.3524 + "tflops": 0.9865 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7727 + "tflops": 1.8774 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 1.7153 + "tflops": 3.7692 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.3279 + "tflops": 1.741 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 6.537 + "tflops": 3.4319 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 12.6312 + "tflops": 6.7199 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 12.9988 + "tflops": 13.0163 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 19.7027 + "tflops": 16.6828 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 18.2978 + "tflops": 21.3457 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 19.0515 + "tflops": 21.7726 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 18.7332 + "tflops": 22.5085 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 23.7608 + "tflops": 22.6894 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 23.8789 + "tflops": 22.5829 } ] }, @@ -223,77 +1951,77 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 1.2658 + "tflops": 0.9366 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7165 + "tflops": 1.7981 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 1.6581 + "tflops": 3.6229 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.0787 + "tflops": 1.6258 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 6.1021 + "tflops": 3.2067 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 11.821 + "tflops": 6.2312 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 12.8133 + "tflops": 11.8048 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 18.5013 + "tflops": 16.4615 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 17.0859 + "tflops": 20.2181 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 18.1149 + "tflops": 19.3784 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 18.0408 + "tflops": 21.2318 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 23.729 + "tflops": 21.1794 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 23.7454 + "tflops": 20.4921 } ] } ] }, { - "in_features": 2048, - "out_features": 2560, + "in_features": 2560, + "out_features": 4096, "group_size": 128, - "comment": "gemma-2b qkv_proj", + "comment": "Gemma3-4B qkv_proj", "providers": [ { "provider": "hybrid-w4a16", @@ -301,67 +2029,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 1.2864 + "tflops": 0.9657 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7909 + "tflops": 1.822 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 1.8736 + "tflops": 3.7163 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.8731 + "tflops": 2.1201 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 7.5097 + "tflops": 4.1406 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 14.3792 + "tflops": 7.9402 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 13.0927 + "tflops": 14.5963 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 21.705 + "tflops": 22.0096 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 16.4913 + "tflops": 26.6971 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 16.6033 + "tflops": 27.2121 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 17.6707 + "tflops": 27.9626 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 23.0623 + "tflops": 29.2287 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 23.4599 + "tflops": 28.8006 } ] }, @@ -371,67 +2099,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 1.1639 + "tflops": 0.8949 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7231 + "tflops": 1.6724 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 1.8702 + "tflops": 3.3103 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.6878 + "tflops": 2.6151 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 7.222 + "tflops": 4.7196 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 13.4316 + "tflops": 8.8714 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 12.8495 + "tflops": 16.6964 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 18.4533 + "tflops": 24.0235 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 16.1602 + "tflops": 26.3426 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 16.5573 + "tflops": 26.6898 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 17.2919 + "tflops": 27.8513 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 22.668 + "tflops": 29.317 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 23.3596 + "tflops": 29.1983 } ] }, @@ -441,67 +2169,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 1.2682 + "tflops": 0.96 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7146 + "tflops": 1.7871 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 1.8464 + "tflops": 3.6257 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.5959 + "tflops": 1.8182 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 6.9594 + "tflops": 3.5398 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 13.3677 + "tflops": 6.847 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 14.0524 + "tflops": 13.1463 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 22.9062 + "tflops": 17.8187 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 18.3427 + "tflops": 19.6771 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 18.1713 + "tflops": 20.7722 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 19.3614 + "tflops": 20.8275 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 24.3312 + "tflops": 22.0856 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 24.5709 + "tflops": 22.3205 } ] }, @@ -511,77 +2239,77 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 1.1972 + "tflops": 0.9122 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7096 + "tflops": 1.7162 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 1.8188 + "tflops": 3.4726 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.4401 + "tflops": 1.7116 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 6.7777 + "tflops": 3.387 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 13.0765 + "tflops": 6.519 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 13.7838 + "tflops": 11.9366 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 20.377 + "tflops": 17.1079 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 17.4448 + "tflops": 18.6979 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 17.6678 + "tflops": 20.0784 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 18.9946 + "tflops": 20.2248 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 23.791 + "tflops": 20.161 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 24.435 + "tflops": 20.2285 } ] } ] }, { - "in_features": 2048, - "out_features": 32768, + "in_features": 2560, + "out_features": 19456, "group_size": 128, - "comment": "gemma-2b gate_up_proj", + "comment": "Qwen3-4B gate_up_proj", "providers": [ { "provider": "hybrid-w4a16", @@ -589,67 +2317,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9247 + "tflops": 0.8908 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7686 + "tflops": 1.7186 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.4351 + "tflops": 3.3222 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.3068 + "tflops": 2.1262 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 8.2607 + "tflops": 4.1906 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 16.0032 + "tflops": 8.0963 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 9.257 + "tflops": 15.2998 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 20.6918 + "tflops": 27.5738 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 24.2849 + "tflops": 29.6011 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 24.5337 + "tflops": 30.5525 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 23.9014 + "tflops": 31.3004 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 21.6818 + "tflops": 31.5412 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 19.6505 + "tflops": 31.789 } ] }, @@ -659,67 +2387,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8991 + "tflops": 0.8518 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.6974 + "tflops": 1.6112 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.2795 + "tflops": 3.1533 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.0476 + "tflops": 1.5867 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 7.9376 + "tflops": 3.1625 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 14.9119 + "tflops": 6.1146 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 7.4613 + "tflops": 12.1102 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 20.293 + "tflops": 26.2346 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 24.0634 + "tflops": 29.1859 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 24.3204 + "tflops": 30.5001 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 23.6918 + "tflops": 31.4039 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 21.3473 + "tflops": 31.8688 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 19.499 + "tflops": 32.2731 } ] }, @@ -729,67 +2457,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9152 + "tflops": 0.8871 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.758 + "tflops": 1.71 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.3974 + "tflops": 3.3187 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.0196 + "tflops": 1.4679 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 7.8085 + "tflops": 2.8666 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 15.1246 + "tflops": 5.5965 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 9.0186 + "tflops": 10.7474 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 22.442 + "tflops": 20.3358 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 25.9514 + "tflops": 20.9652 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 25.9164 + "tflops": 20.9976 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 24.9292 + "tflops": 21.0775 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 22.6886 + "tflops": 21.7142 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 21.0021 + "tflops": 22.0322 } ] }, @@ -799,67 +2527,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8995 + "tflops": 0.8591 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.6913 + "tflops": 1.6495 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.1925 + "tflops": 3.2264 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.8311 + "tflops": 1.4249 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 7.4835 + "tflops": 2.8336 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 14.5089 + "tflops": 5.5024 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 7.34 + "tflops": 10.5204 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 20.5148 + "tflops": 19.8473 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 25.7319 + "tflops": 20.6832 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 25.7786 + "tflops": 20.4263 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 24.832 + "tflops": 20.962 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 22.711 + "tflops": 21.5173 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 20.923 + "tflops": 21.7468 } ] } @@ -867,9 +2595,9 @@ }, { "in_features": 2560, - "out_features": 2560, + "out_features": 20480, "group_size": 128, - "comment": "Qwen3-4B o_proj", + "comment": "Gemma3-4B gate_up_proj", "providers": [ { "provider": "hybrid-w4a16", @@ -877,67 +2605,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 1.0641 + "tflops": 0.8946 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.8888 + "tflops": 1.7328 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 1.964 + "tflops": 3.3671 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.2454 + "tflops": 2.1915 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 8.3269 + "tflops": 4.327 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 15.3406 + "tflops": 8.3624 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 15.9707 + "tflops": 15.6054 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 18.9499 + "tflops": 28.2873 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 20.3059 + "tflops": 30.418 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 22.1459 + "tflops": 30.7887 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 22.8095 + "tflops": 31.3324 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 23.7538 + "tflops": 31.533 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 23.9946 + "tflops": 31.0749 } ] }, @@ -947,67 +2675,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.968 + "tflops": 0.856 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7862 + "tflops": 1.6257 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 1.9541 + "tflops": 3.2029 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.1948 + "tflops": 1.6474 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 8.0128 + "tflops": 3.2787 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 15.5936 + "tflops": 6.4326 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 15.3383 + "tflops": 12.6516 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 18.4911 + "tflops": 26.4926 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 20.1208 + "tflops": 29.4912 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 22.4179 + "tflops": 30.846 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 22.8093 + "tflops": 31.4098 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 23.6467 + "tflops": 31.5553 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 23.7181 + "tflops": 31.5647 } ] }, @@ -1017,67 +2745,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 1.0738 + "tflops": 0.8917 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.8448 + "tflops": 1.7176 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 1.9545 + "tflops": 3.3428 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.9924 + "tflops": 1.5102 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 7.9266 + "tflops": 2.9912 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 15.2644 + "tflops": 5.817 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 17.4998 + "tflops": 11.1144 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 20.0723 + "tflops": 20.844 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 23.1021 + "tflops": 21.001 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 24.022 + "tflops": 21.8362 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 24.2114 + "tflops": 21.7072 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 25.2344 + "tflops": 21.8859 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 25.3184 + "tflops": 21.8004 } ] }, @@ -1087,77 +2815,77 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9984 + "tflops": 0.8628 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7989 + "tflops": 1.6575 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 1.9533 + "tflops": 3.2835 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.8566 + "tflops": 1.4749 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 7.6189 + "tflops": 2.9016 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 14.7116 + "tflops": 5.622 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 17.693 + "tflops": 10.746 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 19.2561 + "tflops": 19.9005 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 21.4957 + "tflops": 21.1398 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 22.6414 + "tflops": 21.1544 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 22.7318 + "tflops": 21.1007 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 25.1294 + "tflops": 21.2922 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 25.3521 + "tflops": 20.5278 } ] } ] }, { - "in_features": 2560, - "out_features": 3840, + "in_features": 3584, + "out_features": 3584, "group_size": 128, - "comment": "Qwen3-4B qkv_proj", + "comment": "Qwen2.5-7B o_proj", "providers": [ { "provider": "hybrid-w4a16", @@ -1165,67 +2893,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9951 + "tflops": 0.9827 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.8274 + "tflops": 1.8462 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.5253 + "tflops": 3.7054 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.3449 + "tflops": 1.2996 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 8.5221 + "tflops": 2.5657 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 16.2684 + "tflops": 4.9792 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 17.1006 + "tflops": 9.5315 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 22.8312 + "tflops": 17.4976 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 23.5846 + "tflops": 25.2939 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 20.577 + "tflops": 31.7169 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 23.3841 + "tflops": 31.1318 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 24.6554 + "tflops": 31.2303 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 24.9148 + "tflops": 31.1835 } ] }, @@ -1235,67 +2963,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9157 + "tflops": 0.9205 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7331 + "tflops": 1.7106 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.3217 + "tflops": 3.3912 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.235 + "tflops": 1.4295 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 8.2574 + "tflops": 2.805 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 15.7746 + "tflops": 5.4748 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 17.1391 + "tflops": 10.4818 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 23.2947 + "tflops": 19.4789 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 23.3903 + "tflops": 25.6961 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 20.7392 + "tflops": 30.8347 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 23.7488 + "tflops": 31.1464 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 24.5053 + "tflops": 31.4989 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 24.716 + "tflops": 31.9244 } ] }, @@ -1305,67 +3033,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9856 + "tflops": 0.9806 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.827 + "tflops": 1.8334 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.4782 + "tflops": 3.6956 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.1246 + "tflops": 0.8782 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 8.1326 + "tflops": 1.7326 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 15.6296 + "tflops": 3.367 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 19.2748 + "tflops": 6.5339 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 25.0966 + "tflops": 12.4797 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 24.4635 + "tflops": 17.4311 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 20.0961 + "tflops": 17.4813 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 24.5463 + "tflops": 21.2096 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 25.7778 + "tflops": 21.1003 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 25.8978 + "tflops": 21.4669 } ] }, @@ -1375,77 +3103,77 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9305 + "tflops": 0.9375 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7363 + "tflops": 1.7612 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.284 + "tflops": 3.5538 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.0933 + "tflops": 0.8855 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 8.1012 + "tflops": 1.7441 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 15.6989 + "tflops": 3.3787 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 19.8573 + "tflops": 6.5585 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 23.2563 + "tflops": 12.53 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 23.0665 + "tflops": 16.8721 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 18.8356 + "tflops": 17.1796 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 23.08 + "tflops": 20.9009 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 25.6339 + "tflops": 20.8974 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 25.7817 + "tflops": 20.7896 } ] } ] }, { - "in_features": 2560, - "out_features": 19456, + "in_features": 3584, + "out_features": 4608, "group_size": 128, - "comment": "Qwen3-4B gate_up_proj", + "comment": "Qwen2.5-7B qkv_proj", "providers": [ { "provider": "hybrid-w4a16", @@ -1453,67 +3181,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8928 + "tflops": 0.9599 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7141 + "tflops": 1.8092 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.2879 + "tflops": 3.6468 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.359 + "tflops": 1.6439 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 8.489 + "tflops": 3.2469 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 16.3508 + "tflops": 6.2961 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 21.7559 + "tflops": 11.8297 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 23.7674 + "tflops": 21.8399 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 24.238 + "tflops": 29.6812 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 24.5302 + "tflops": 29.0643 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 24.3167 + "tflops": 30.1562 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 24.2346 + "tflops": 31.663 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 24.0488 + "tflops": 31.6427 } ] }, @@ -1523,67 +3251,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.854 + "tflops": 0.9003 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.6112 + "tflops": 1.6655 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.1849 + "tflops": 3.3366 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.1882 + "tflops": 1.7787 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 8.2449 + "tflops": 3.5706 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 15.686 + "tflops": 6.8764 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 15.3656 + "tflops": 12.8607 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 22.8964 + "tflops": 23.8148 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 23.6275 + "tflops": 30.0127 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 23.9546 + "tflops": 30.3585 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 24.1914 + "tflops": 30.9038 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 24.057 + "tflops": 32.0528 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 23.812 + "tflops": 32.4043 } ] }, @@ -1593,67 +3321,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8868 + "tflops": 0.9515 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.6983 + "tflops": 1.7941 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.3113 + "tflops": 3.6226 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.2113 + "tflops": 1.1208 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 8.2385 + "tflops": 2.212 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 16.1452 + "tflops": 4.2907 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 21.0316 + "tflops": 8.2158 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 25.6797 + "tflops": 15.5746 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 25.1757 + "tflops": 19.9207 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 25.594 + "tflops": 20.3574 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 25.6878 + "tflops": 20.6498 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 25.7688 + "tflops": 20.7403 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 25.5703 + "tflops": 21.6341 } ] }, @@ -1663,67 +3391,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8546 + "tflops": 0.9133 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.638 + "tflops": 1.7193 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.1231 + "tflops": 3.5064 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.132 + "tflops": 1.1218 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 8.023 + "tflops": 2.2065 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 15.8907 + "tflops": 4.2488 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 14.5537 + "tflops": 8.169 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 23.7896 + "tflops": 15.5241 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 25.3109 + "tflops": 19.7499 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 25.5795 + "tflops": 20.0352 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 25.4976 + "tflops": 20.3374 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 25.4903 + "tflops": 20.3859 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 25.5424 + "tflops": 20.9096 } ] } @@ -1731,9 +3459,9 @@ }, { "in_features": 3584, - "out_features": 3584, + "out_features": 37888, "group_size": 128, - "comment": "Qwen2.5-7B o_proj", + "comment": "Qwen2.5-7B gate_up_proj", "providers": [ { "provider": "hybrid-w4a16", @@ -1741,67 +3469,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9835 + "tflops": 0.9474 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.8347 + "tflops": 1.792 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.5694 + "tflops": 3.5143 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.1539 + "tflops": 2.0716 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 7.8149 + "tflops": 4.0453 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 15.8862 + "tflops": 7.853 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 16.8782 + "tflops": 14.7785 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 18.0556 + "tflops": 27.3903 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 22.0844 + "tflops": 31.564 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 20.633 + "tflops": 32.0038 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 23.0266 + "tflops": 32.0317 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 24.5028 + "tflops": 32.066 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 24.4662 + "tflops": 32.3823 } ] }, @@ -1811,67 +3539,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9292 + "tflops": 0.9414 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7102 + "tflops": 1.7192 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.3607 + "tflops": 3.3846 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.1325 + "tflops": 1.7197 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 7.9536 + "tflops": 3.4423 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 15.5861 + "tflops": 6.8306 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 16.4198 + "tflops": 13.3952 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 17.5615 + "tflops": 25.771 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 22.0996 + "tflops": 31.2743 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 20.7903 + "tflops": 32.3797 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 22.6602 + "tflops": 32.5806 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 24.2616 + "tflops": 33.0038 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 24.5029 + "tflops": 32.7214 } ] }, @@ -1881,67 +3609,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9774 + "tflops": 0.9473 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.802 + "tflops": 1.7757 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.5337 + "tflops": 3.4757 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.0088 + "tflops": 1.4147 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 7.9151 + "tflops": 2.7907 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 15.4545 + "tflops": 5.4895 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 18.795 + "tflops": 10.6477 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 18.476 + "tflops": 20.1035 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 23.5884 + "tflops": 21.0172 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 19.2127 + "tflops": 22.0944 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 23.6789 + "tflops": 21.9033 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 25.5239 + "tflops": 22.5496 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 25.7562 + "tflops": 22.2133 } ] }, @@ -1951,77 +3679,77 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9331 + "tflops": 0.9446 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7483 + "tflops": 1.7525 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.325 + "tflops": 3.4817 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.9838 + "tflops": 1.3255 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 7.8087 + "tflops": 2.6733 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 15.214 + "tflops": 5.1914 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 18.6939 + "tflops": 9.9831 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 17.975 + "tflops": 19.5256 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 22.3571 + "tflops": 20.7682 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 18.1997 + "tflops": 21.4711 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 22.0545 + "tflops": 21.7887 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 25.4933 + "tflops": 21.9907 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 25.5425 + "tflops": 22.038 } ] } ] }, { - "in_features": 3584, - "out_features": 4608, + "in_features": 4096, + "out_features": 4096, "group_size": 128, - "comment": "Qwen2.5-7B qkv_proj", + "comment": "Qwen3-8B o_proj", "providers": [ { "provider": "hybrid-w4a16", @@ -2029,67 +3757,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9585 + "tflops": 0.891 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7914 + "tflops": 1.6342 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.5332 + "tflops": 2.9701 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.1761 + "tflops": 0.6892 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 7.4575 + "tflops": 1.38 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 16.2209 + "tflops": 2.7171 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 18.5579 + "tflops": 5.2821 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 22.0935 + "tflops": 8.0793 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 21.567 + "tflops": 7.8258 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 22.2707 + "tflops": 8.3051 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 22.8004 + "tflops": 17.5122 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 24.4321 + "tflops": 26.2673 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 24.7165 + "tflops": 27.3317 } ] }, @@ -2099,217 +3827,217 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9067 + "tflops": 0.8416 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.669 + "tflops": 1.5287 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.355 + "tflops": 2.6861 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.8338 + "tflops": 1.9476 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 7.2186 + "tflops": 3.4268 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 14.2532 + "tflops": 5.5264 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 17.164 + "tflops": 7.9997 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 21.2933 + "tflops": 8.0331 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 20.9253 + "tflops": 8.3338 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 21.5604 + "tflops": 8.9581 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 22.6155 + "tflops": 16.7889 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 24.303 + "tflops": 25.6035 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 24.6696 + "tflops": 27.5938 } ] }, { - "provider": "hybrid-w4a16-bf16", + "provider": "hybrid-w4a16-zp-bf16", "baselines": [ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9427 + "tflops": 0.8471 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7606 + "tflops": 1.5567 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.4726 + "tflops": 2.7851 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.6749 + "tflops": 1.3646 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 7.3571 + "tflops": 2.6915 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 14.5447 + "tflops": 4.5839 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 18.9263 + "tflops": 6.641 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 23.929 + "tflops": 7.1821 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 23.2479 + "tflops": 7.1563 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 22.8656 + "tflops": 9.0672 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 23.1451 + "tflops": 15.9096 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 25.7038 + "tflops": 19.4167 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 25.6065 + "tflops": 19.9405 } ] }, { - "provider": "hybrid-w4a16-zp-bf16", + "provider": "hybrid-w4a16-bf16", "baselines": [ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9113 + "tflops": 0.8887 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7064 + "tflops": 1.6254 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.307 + "tflops": 2.9424 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.6973 + "tflops": 1.7243 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 7.3129 + "tflops": 3.4983 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 14.4331 + "tflops": 6.7653 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 19.4637 + "tflops": 12.4652 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 22.2201 + "tflops": 10.2476 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 21.6125 + "tflops": 7.7119 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 21.4651 + "tflops": 10.2066 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 21.627 + "tflops": 17.7252 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 25.4997 + "tflops": 21.7368 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 25.0739 + "tflops": 22.5135 } ] } ] }, { - "in_features": 3584, - "out_features": 37888, + "in_features": 4096, + "out_features": 6144, "group_size": 128, - "comment": "Qwen2.5-7B gate_up_proj", + "comment": "Qwen3-8B qkv_proj", "providers": [ { "provider": "hybrid-w4a16", @@ -2317,67 +4045,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9485 + "tflops": 0.9058 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7864 + "tflops": 1.6866 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.4789 + "tflops": 2.7386 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.495 + "tflops": 0.7778 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 8.715 + "tflops": 1.5534 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 16.5926 + "tflops": 3.2561 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 22.2507 + "tflops": 6.0536 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 23.5897 + "tflops": 7.1498 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 24.9993 + "tflops": 7.4862 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 25.0296 + "tflops": 10.7365 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 25.158 + "tflops": 20.7286 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 25.1095 + "tflops": 26.8236 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 24.6831 + "tflops": 27.5498 } ] }, @@ -2387,67 +4115,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9442 + "tflops": 0.8701 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7205 + "tflops": 1.6223 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.4637 + "tflops": 2.5724 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.3038 + "tflops": 1.6043 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 8.2927 + "tflops": 2.5839 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 16.0686 + "tflops": 4.5336 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 14.764 + "tflops": 7.6643 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 23.2801 + "tflops": 7.1183 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 24.3827 + "tflops": 7.1238 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 24.6459 + "tflops": 9.5693 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 24.7479 + "tflops": 19.4538 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 24.7876 + "tflops": 27.4105 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 24.2924 + "tflops": 28.1118 } ] }, @@ -2457,67 +4185,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9278 + "tflops": 0.9053 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7707 + "tflops": 1.6817 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.4841 + "tflops": 2.7306 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.2226 + "tflops": 1.6987 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 8.2607 + "tflops": 3.3307 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 15.9568 + "tflops": 6.5596 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 21.5344 + "tflops": 11.6812 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 26.0977 + "tflops": 7.0207 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 25.8018 + "tflops": 7.4687 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 25.8168 + "tflops": 12.2557 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 25.9784 + "tflops": 19.7106 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 26.19 + "tflops": 21.721 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 25.7026 + "tflops": 22.9786 } ] }, @@ -2527,67 +4255,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.922 + "tflops": 0.8674 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7209 + "tflops": 1.6294 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.375 + "tflops": 2.6357 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.9634 + "tflops": 1.312 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 7.8804 + "tflops": 2.2099 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 15.5574 + "tflops": 3.9258 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 14.6318 + "tflops": 6.7379 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 23.7604 + "tflops": 6.793 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 25.6745 + "tflops": 6.5367 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 25.4901 + "tflops": 11.0368 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 25.5034 + "tflops": 17.518 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 23.9144 + "tflops": 19.9604 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 25.5529 + "tflops": 20.1446 } ] } @@ -2595,9 +4323,9 @@ }, { "in_features": 4096, - "out_features": 4096, + "out_features": 24576, "group_size": 128, - "comment": "Qwen3-8B o_proj", + "comment": "Qwen3-8B gate_up_proj", "providers": [ { "provider": "hybrid-w4a16", @@ -2605,67 +4333,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8801 + "tflops": 0.9069 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.5812 + "tflops": 1.7219 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 2.9409 + "tflops": 3.3879 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 2.5406 + "tflops": 0.6201 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 5.0901 + "tflops": 1.359 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 9.9565 + "tflops": 2.3977 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 5.3127 + "tflops": 4.5619 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 18.489 + "tflops": 8.0076 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 7.6342 + "tflops": 14.1235 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 10.0903 + "tflops": 23.5719 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 16.7122 + "tflops": 29.2092 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 24.6697 + "tflops": 30.923 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 18.1801 + "tflops": 29.6173 } ] }, @@ -2675,207 +4403,207 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8389 + "tflops": 0.9056 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.5446 + "tflops": 1.6626 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 2.6576 + "tflops": 3.252 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 2.4447 + "tflops": 0.5237 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 5.3615 + "tflops": 1.0348 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 9.263 + "tflops": 1.9542 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 4.8619 + "tflops": 3.727 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 17.686 + "tflops": 6.995 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 6.6291 + "tflops": 14.1578 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 9.03 + "tflops": 22.0651 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 15.7897 + "tflops": 30.3687 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 24.4615 + "tflops": 32.2342 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 18.2463 + "tflops": 29.4532 } ] }, { - "provider": "hybrid-w4a16-zp-bf16", + "provider": "hybrid-w4a16-bf16", "baselines": [ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8441 + "tflops": 0.9096 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.5661 + "tflops": 1.7191 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 2.6097 + "tflops": 3.3668 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 2.2277 + "tflops": 0.6024 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 4.4826 + "tflops": 1.1609 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 8.8083 + "tflops": 2.2757 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 4.9037 + "tflops": 4.443 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 17.8947 + "tflops": 8.1429 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 6.5683 + "tflops": 13.0335 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 9.0265 + "tflops": 18.7773 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 15.7467 + "tflops": 20.4773 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 25.01 + "tflops": 21.2594 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 18.4844 + "tflops": 20.724 } ] }, { - "provider": "hybrid-w4a16-bf16", + "provider": "hybrid-w4a16-zp-bf16", "baselines": [ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8726 + "tflops": 0.9098 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.62 + "tflops": 1.6769 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 2.9255 + "tflops": 3.3439 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 2.4974 + "tflops": 0.5032 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 4.9684 + "tflops": 1.0236 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 9.5407 + "tflops": 2.0517 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 5.3852 + "tflops": 3.9999 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 17.9886 + "tflops": 7.8184 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 7.5107 + "tflops": 12.2582 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 10.2528 + "tflops": 17.8501 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 17.0682 + "tflops": 19.8511 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 25.1099 + "tflops": 20.3894 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 18.4845 + "tflops": 19.2773 } ] } @@ -2883,9 +4611,9 @@ }, { "in_features": 4096, - "out_features": 6144, + "out_features": 2048, "group_size": 128, - "comment": "Qwen3-8B qkv_proj", + "comment": "Qwen3.5-35B-A3B GDN out_proj", "providers": [ { "provider": "hybrid-w4a16", @@ -2893,67 +4621,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.876 + "tflops": 0.9142 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.6678 + "tflops": 1.5583 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 2.7312 + "tflops": 3.0526 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 2.1969 + "tflops": 0.7447 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 4.3383 + "tflops": 1.4758 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 8.4491 + "tflops": 2.9006 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 4.7209 + "tflops": 5.6179 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 20.1226 + "tflops": 9.0536 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 5.972 + "tflops": 14.4081 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 11.0402 + "tflops": 10.9412 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 19.3295 + "tflops": 12.8371 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 24.4893 + "tflops": 24.4081 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 17.3335 + "tflops": 25.5683 } ] }, @@ -2963,67 +4691,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8636 + "tflops": 0.8409 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.5756 + "tflops": 1.4476 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 2.6067 + "tflops": 2.8446 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 2.2277 + "tflops": 1.0974 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 4.341 + "tflops": 2.1132 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 8.1686 + "tflops": 3.8893 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 4.5626 + "tflops": 7.0552 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 19.2137 + "tflops": 12.0781 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 5.2924 + "tflops": 13.6385 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 9.8484 + "tflops": 10.1117 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 18.0353 + "tflops": 11.9815 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 24.6201 + "tflops": 23.4816 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 17.3055 + "tflops": 26.5053 } ] }, @@ -3033,67 +4761,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8996 + "tflops": 0.8939 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.6717 + "tflops": 1.5493 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 2.7599 + "tflops": 2.9442 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 2.2078 + "tflops": 0.8556 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 4.4118 + "tflops": 1.7243 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 8.2191 + "tflops": 3.4399 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 4.7325 + "tflops": 6.7338 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 19.5135 + "tflops": 13.0712 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 5.9658 + "tflops": 14.8378 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 11.2653 + "tflops": 10.3963 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 19.4455 + "tflops": 15.1154 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 24.9241 + "tflops": 20.406 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 17.6045 + "tflops": 21.6576 } ] }, @@ -3103,77 +4831,77 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8511 + "tflops": 0.839 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.6016 + "tflops": 1.488 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 2.567 + "tflops": 2.558 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 2.0103 + "tflops": 0.7746 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 4.0027 + "tflops": 1.4999 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 7.7614 + "tflops": 2.774 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 4.5797 + "tflops": 5.2506 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 19.4189 + "tflops": 9.4703 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 5.1948 + "tflops": 12.1404 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 10.0063 + "tflops": 11.7443 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 17.957 + "tflops": 13.7116 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 24.6564 + "tflops": 18.8508 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 17.3753 + "tflops": 19.3658 } ] } ] }, { - "in_features": 4096, - "out_features": 24576, + "in_features": 6144, + "out_features": 2048, "group_size": 128, - "comment": "Qwen3-8B gate_up_proj", + "comment": "Qwen3-1.7B down_proj", "providers": [ { "provider": "hybrid-w4a16", @@ -3181,67 +4909,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9091 + "tflops": 0.8833 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7382 + "tflops": 1.6901 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.3076 + "tflops": 3.2731 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 2.6973 + "tflops": 0.8417 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 4.954 + "tflops": 1.6725 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 9.139 + "tflops": 3.3041 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 4.0481 + "tflops": 6.4718 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 19.8157 + "tflops": 11.4941 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 23.7086 + "tflops": 17.3671 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 25.3975 + "tflops": 20.9193 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 25.2058 + "tflops": 18.5265 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 25.2144 + "tflops": 25.7565 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 20.1374 + "tflops": 20.2314 } ] }, @@ -3251,67 +4979,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8986 + "tflops": 0.8251 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.6747 + "tflops": 1.5559 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.2698 + "tflops": 3.0386 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 2.2405 + "tflops": 1.182 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 4.3141 + "tflops": 2.3199 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 8.6709 + "tflops": 4.4475 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 3.2115 + "tflops": 8.2273 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 18.8777 + "tflops": 14.9229 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 23.1642 + "tflops": 17.3573 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 24.7568 + "tflops": 19.3475 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 25.1365 + "tflops": 17.5902 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 25.0005 + "tflops": 25.6497 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 20.0364 + "tflops": 20.5957 } ] }, @@ -3321,67 +5049,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9054 + "tflops": 0.8724 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.738 + "tflops": 1.6681 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.3016 + "tflops": 3.1439 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 2.3695 + "tflops": 0.938 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 4.5947 + "tflops": 1.8408 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 8.6985 + "tflops": 3.5858 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 4.328 + "tflops": 6.944 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 19.5395 + "tflops": 13.6228 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 23.745 + "tflops": 16.7535 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 25.4509 + "tflops": 18.759 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 25.593 + "tflops": 17.8364 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 25.5 + "tflops": 20.8599 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 20.4605 + "tflops": 19.7268 } ] }, @@ -3391,77 +5119,77 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8882 + "tflops": 0.8139 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.6885 + "tflops": 1.596 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 3.1548 + "tflops": 2.8025 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 2.2772 + "tflops": 0.6596 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 4.5175 + "tflops": 1.3097 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 8.5687 + "tflops": 2.5729 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 3.2218 + "tflops": 5.0456 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 19.7153 + "tflops": 8.9995 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 23.1478 + "tflops": 13.5697 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 25.1433 + "tflops": 17.2579 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 25.2589 + "tflops": 16.0322 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 25.4194 + "tflops": 19.1452 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 20.3847 + "tflops": 16.7398 } ] } ] }, { - "in_features": 4096, - "out_features": 2048, + "in_features": 8192, + "out_features": 512, "group_size": 128, - "comment": "Qwen3.5-35B-A3B GDN out_proj", + "comment": "L2 2MiB at", "providers": [ { "provider": "hybrid-w4a16", @@ -3469,67 +5197,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9383 + "tflops": 0.9918 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.6086 + "tflops": 1.7155 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 1.9359 + "tflops": 2.5448 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.0483 + "tflops": 0.2725 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 5.989 + "tflops": 0.5327 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 11.9035 + "tflops": 1.0472 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 11.1658 + "tflops": 2.0586 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 17.5029 + "tflops": 3.8608 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 20.1903 + "tflops": 7.4167 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 18.6465 + "tflops": 13.9223 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 20.4598 + "tflops": 20.3801 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 19.8067 + "tflops": 21.4062 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 18.49 + "tflops": 14.6262 } ] }, @@ -3539,67 +5267,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8659 + "tflops": 0.8573 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.5013 + "tflops": 1.4576 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 1.9112 + "tflops": 2.4248 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 2.6511 + "tflops": 0.3361 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 5.2954 + "tflops": 0.6302 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 10.1059 + "tflops": 1.2325 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 10.0962 + "tflops": 2.4044 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 16.4971 + "tflops": 4.5934 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 18.023 + "tflops": 9.2078 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 18.0744 + "tflops": 17.6097 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 20.0834 + "tflops": 19.5936 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 19.0754 + "tflops": 22.9366 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 17.7431 + "tflops": 16.138 } ] }, @@ -3609,67 +5337,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8974 + "tflops": 0.9757 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.5892 + "tflops": 1.6621 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 1.9001 + "tflops": 2.487 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 2.8712 + "tflops": 0.2421 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 5.6464 + "tflops": 0.4791 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 10.8051 + "tflops": 0.9388 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 11.2481 + "tflops": 1.8058 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 17.6082 + "tflops": 3.4986 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 20.6487 + "tflops": 7.096 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 23.003 + "tflops": 13.8743 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 24.9117 + "tflops": 17.9302 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 21.4438 + "tflops": 19.2357 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 18.7492 + "tflops": 16.0936 } ] }, @@ -3679,77 +5407,77 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8679 + "tflops": 0.8633 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.5197 + "tflops": 1.5072 }, { "batch_size": 4, "kernel": "wvsplitk_int4", - "tflops": 1.8586 + "tflops": 2.2311 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 2.6793 + "tflops": 0.2071 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 5.2758 + "tflops": 0.4069 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 10.1368 + "tflops": 0.8075 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 10.4729 + "tflops": 1.5918 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 16.4199 + "tflops": 2.8465 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 18.4962 + "tflops": 5.5724 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 20.469 + "tflops": 10.3583 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 23.6211 + "tflops": 14.7445 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 20.9558 + "tflops": 17.0884 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 18.393 + "tflops": 13.6706 } ] } ] }, { - "in_features": 8192, + "in_features": 8320, "out_features": 512, "group_size": 128, - "comment": "L2 2MiB at", + "comment": "L2 2MiB above", "providers": [ { "provider": "hybrid-w4a16", @@ -3757,67 +5485,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 1.1493 + "tflops": 0.9682 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.583 + "tflops": 1.8149 }, { "batch_size": 4, - "kernel": "wvsplitk_int4", - "tflops": 1.7119 + "kernel": "hybrid_triton_w4a16", + "tflops": 0.0977 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 1.67 + "tflops": 0.1951 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 3.1469 + "tflops": 0.3887 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 5.4256 + "tflops": 0.7726 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 3.5442 + "tflops": 1.5347 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 9.6996 + "tflops": 3.0304 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 16.0883 + "tflops": 6.0345 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 18.2808 + "tflops": 11.7608 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 21.8062 + "tflops": 22.4916 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 16.5493 + "tflops": 28.4092 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 14.2201 + "tflops": 27.2318 } ] }, @@ -3827,67 +5555,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9742 + "tflops": 0.8567 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.4265 + "tflops": 1.6121 }, { "batch_size": 4, - "kernel": "wvsplitk_int4", - "tflops": 1.679 + "kernel": "hybrid_triton_w4a16", + "tflops": 0.1065 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 1.5253 + "tflops": 0.2127 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 2.8178 + "tflops": 0.4238 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 5.0579 + "tflops": 0.8403 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 3.3119 + "tflops": 1.6588 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 9.3804 + "tflops": 3.2698 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 15.5742 + "tflops": 6.5458 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 18.3559 + "tflops": 12.8844 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 21.6903 + "tflops": 24.125 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 15.8783 + "tflops": 28.9633 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 12.0947 + "tflops": 27.8209 } ] }, @@ -3897,67 +5625,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 1.1474 + "tflops": 0.9513 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.5826 + "tflops": 1.7603 }, { "batch_size": 4, - "kernel": "wvsplitk_int4", - "tflops": 1.6762 + "kernel": "hybrid_triton_w4a16", + "tflops": 0.0654 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 1.4824 + "tflops": 0.1306 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 2.8497 + "tflops": 0.2604 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 5.0031 + "tflops": 0.519 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 3.6669 + "tflops": 1.0325 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 9.6239 + "tflops": 2.0403 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 18.3148 + "tflops": 4.0647 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 19.8355 + "tflops": 8.0148 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 24.0519 + "tflops": 15.5757 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 16.7656 + "tflops": 19.5266 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 13.5411 + "tflops": 18.8673 } ] }, @@ -3967,77 +5695,77 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9991 + "tflops": 0.8826 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.4644 + "tflops": 1.6999 }, { "batch_size": 4, - "kernel": "wvsplitk_int4", - "tflops": 1.6134 + "kernel": "hybrid_triton_w4a16", + "tflops": 0.0648 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 1.5056 + "tflops": 0.1295 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 2.9368 + "tflops": 0.2577 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 5.0178 + "tflops": 0.5131 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 3.4393 + "tflops": 1.0236 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 8.8621 + "tflops": 2.0325 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 16.4891 + "tflops": 3.9919 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 18.7655 + "tflops": 7.7657 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 22.1401 + "tflops": 14.8809 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 15.8708 + "tflops": 19.256 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 11.9382 + "tflops": 18.1177 } ] } ] }, { - "in_features": 8320, - "out_features": 512, + "in_features": 9728, + "out_features": 2560, "group_size": 128, - "comment": "L2 2MiB above", + "comment": "Qwen3-4B down_proj", "providers": [ { "provider": "hybrid-w4a16", @@ -4045,67 +5773,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 1.1777 + "tflops": 0.9364 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.6729 + "tflops": 1.8359 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 0.8472 + "tflops": 0.4614 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 1.685 + "tflops": 0.9223 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 3.3672 + "tflops": 1.8322 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 6.647 + "tflops": 3.614 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 4.5277 + "tflops": 7.0769 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 10.7157 + "tflops": 13.5563 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 16.0478 + "tflops": 25.7784 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 18.5271 + "tflops": 31.947 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 23.0259 + "tflops": 31.1157 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 19.5149 + "tflops": 30.357 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 19.8677 + "tflops": 30.9075 } ] }, @@ -4115,67 +5843,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 1.0332 + "tflops": 0.9039 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.5737 + "tflops": 1.7022 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 0.7516 + "tflops": 0.5257 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 1.4951 + "tflops": 1.0417 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 2.9602 + "tflops": 2.0626 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 5.8492 + "tflops": 4.0694 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 4.2091 + "tflops": 7.9322 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 10.8307 + "tflops": 15.4123 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 15.6779 + "tflops": 28.1909 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 18.411 + "tflops": 33.7798 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 22.7803 + "tflops": 32.864 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 19.3718 + "tflops": 31.6789 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 20.2113 + "tflops": 31.8298 } ] }, @@ -4185,67 +5913,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 1.1746 + "tflops": 0.9334 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.6628 + "tflops": 1.8264 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 0.758 + "tflops": 0.3199 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 1.5047 + "tflops": 0.6367 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 2.9989 + "tflops": 1.2627 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 5.9566 + "tflops": 2.4912 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 4.5895 + "tflops": 4.904 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 10.441 + "tflops": 9.4703 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 18.8975 + "tflops": 18.1389 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 20.5872 + "tflops": 22.9398 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 24.6901 + "tflops": 22.9036 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 20.2411 + "tflops": 22.3608 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 20.9146 + "tflops": 21.8316 } ] }, @@ -4255,77 +5983,77 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 1.0683 + "tflops": 0.8953 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.5969 + "tflops": 1.7645 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 0.7481 + "tflops": 0.3006 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 1.4903 + "tflops": 0.6002 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 2.9715 + "tflops": 1.1995 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 5.8624 + "tflops": 2.3797 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 4.4103 + "tflops": 4.7054 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 9.9836 + "tflops": 9.2171 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 16.5573 + "tflops": 17.5776 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 18.7919 + "tflops": 22.6747 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 23.0312 + "tflops": 22.285 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 19.9999 + "tflops": 21.4776 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 20.902 + "tflops": 18.0842 } ] } ] }, { - "in_features": 9728, + "in_features": 10240, "out_features": 2560, "group_size": 128, - "comment": "Qwen3-4B down_proj", + "comment": "Gemma3-4B down_proj", "providers": [ { "provider": "hybrid-w4a16", @@ -4333,67 +6061,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9361 + "tflops": 0.9474 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.8381 + "tflops": 1.7571 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 2.1756 + "tflops": 0.4814 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.1222 + "tflops": 0.9583 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 8.3762 + "tflops": 1.9021 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 15.3302 + "tflops": 3.7433 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 15.7077 + "tflops": 7.251 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 21.8582 + "tflops": 12.9912 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 24.4545 + "tflops": 17.1377 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 24.053 + "tflops": 16.4297 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 24.4355 + "tflops": 19.6575 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 23.6322 + "tflops": 24.0721 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 23.6437 + "tflops": 19.9253 } ] }, @@ -4403,67 +6131,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9072 + "tflops": 0.9052 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7243 + "tflops": 1.6662 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 2.1943 + "tflops": 0.7336 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.2395 + "tflops": 1.4512 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 8.208 + "tflops": 2.8262 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 16.0093 + "tflops": 5.2582 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 15.4266 + "tflops": 9.4762 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 23.9891 + "tflops": 17.7795 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 23.7865 + "tflops": 15.9361 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 23.2798 + "tflops": 15.0453 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 24.1159 + "tflops": 18.2305 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 22.8202 + "tflops": 24.1354 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 22.0967 + "tflops": 20.7125 } ] }, @@ -4473,67 +6201,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.9103 + "tflops": 0.9483 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.8069 + "tflops": 1.738 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 2.1057 + "tflops": 0.5704 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.1229 + "tflops": 1.1335 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 8.0928 + "tflops": 2.2456 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 15.8411 + "tflops": 4.3923 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 16.668 + "tflops": 8.0964 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 23.4208 + "tflops": 15.4211 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 26.2979 + "tflops": 16.2339 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 26.768 + "tflops": 16.6561 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 26.3123 + "tflops": 19.4106 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 23.9842 + "tflops": 20.4185 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 24.0119 + "tflops": 19.2072 } ] }, @@ -4543,67 +6271,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8914 + "tflops": 0.8973 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.7599 + "tflops": 1.7129 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 1.9605 + "tflops": 0.3455 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.7542 + "tflops": 0.689 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 7.7122 + "tflops": 1.37 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 15.0475 + "tflops": 2.7153 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 16.5715 + "tflops": 5.3378 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 22.5244 + "tflops": 10.0232 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 23.6942 + "tflops": 13.2986 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 24.1609 + "tflops": 14.0095 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 23.9898 + "tflops": 17.7409 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 23.1156 + "tflops": 17.8051 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 22.9039 + "tflops": 16.4209 } ] } @@ -4621,67 +6349,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8492 + "tflops": 0.8502 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.6561 + "tflops": 1.6487 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 1.7992 + "tflops": 0.441 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.6172 + "tflops": 0.8873 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 7.2098 + "tflops": 1.7754 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 13.7937 + "tflops": 3.4795 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 9.163 + "tflops": 6.8175 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 18.7062 + "tflops": 8.8599 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 17.657 + "tflops": 12.437 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 22.0426 + "tflops": 10.9013 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 24.3151 + "tflops": 18.0574 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 16.9969 + "tflops": 23.4308 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 16.6227 + "tflops": 20.4357 } ] }, @@ -4691,67 +6419,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8005 + "tflops": 0.8066 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.5564 + "tflops": 1.5468 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 1.6645 + "tflops": 0.9584 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.2999 + "tflops": 1.9449 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 6.5801 + "tflops": 3.7708 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 12.6058 + "tflops": 6.9938 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 6.9698 + "tflops": 11.5851 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 17.945 + "tflops": 11.244 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 15.6539 + "tflops": 11.7804 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 21.0149 + "tflops": 11.0301 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 23.8531 + "tflops": 17.2307 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 15.6563 + "tflops": 23.8923 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 15.1108 + "tflops": 20.8318 } ] }, @@ -4761,67 +6489,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8386 + "tflops": 0.8533 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.656 + "tflops": 1.6559 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 1.7196 + "tflops": 0.7821 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.4173 + "tflops": 1.5125 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 6.7982 + "tflops": 2.9433 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 13.1231 + "tflops": 5.5077 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 9.8564 + "tflops": 9.5952 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 17.8348 + "tflops": 9.244 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 17.4486 + "tflops": 10.8992 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 22.8572 + "tflops": 11.7715 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 25.0252 + "tflops": 18.0982 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 16.1874 + "tflops": 20.3325 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 16.1473 + "tflops": 19.4472 } ] }, @@ -4831,67 +6559,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.7746 + "tflops": 0.8175 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.5219 + "tflops": 1.5811 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 1.497 + "tflops": 0.3104 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.1086 + "tflops": 0.6194 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 6.2707 + "tflops": 1.4297 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 12.0282 + "tflops": 2.862 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 7.1846 + "tflops": 4.8809 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 18.2821 + "tflops": 6.5418 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 15.5474 + "tflops": 8.5851 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 20.3352 + "tflops": 9.6767 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 22.9198 + "tflops": 15.2068 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 14.949 + "tflops": 16.8002 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 14.6586 + "tflops": 15.8888 } ] } @@ -4909,67 +6637,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.6628 + "tflops": 0.6626 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.2668 + "tflops": 1.2552 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 1.7605 + "tflops": 0.3631 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.3027 + "tflops": 0.7252 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 6.4611 + "tflops": 1.4298 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 12.7037 + "tflops": 2.9187 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 12.2372 + "tflops": 5.4955 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 18.419 + "tflops": 9.6645 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 21.427 + "tflops": 14.5118 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 19.3701 + "tflops": 14.5374 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 20.9742 + "tflops": 11.7759 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 14.4975 + "tflops": 18.8955 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 15.57 + "tflops": 19.7764 } ] }, @@ -4979,67 +6707,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.6178 + "tflops": 0.6324 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.1842 + "tflops": 1.1878 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 1.6867 + "tflops": 0.5385 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.1813 + "tflops": 1.0745 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 6.2884 + "tflops": 2.1385 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 12.054 + "tflops": 4.1494 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 10.3469 + "tflops": 8.164 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 16.7178 + "tflops": 14.4897 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 20.6519 + "tflops": 14.6282 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 18.8697 + "tflops": 13.0479 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 20.7922 + "tflops": 13.0198 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 13.2217 + "tflops": 18.9139 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 14.1922 + "tflops": 20.2203 } ] }, @@ -5049,67 +6777,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.664 + "tflops": 0.6732 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.2698 + "tflops": 1.2697 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 1.6986 + "tflops": 0.4511 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.3525 + "tflops": 0.8977 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 6.5969 + "tflops": 1.7666 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 12.6248 + "tflops": 3.3677 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 12.468 + "tflops": 6.4155 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 18.9792 + "tflops": 11.9891 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 23.004 + "tflops": 14.5136 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 21.1417 + "tflops": 13.1772 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 22.0724 + "tflops": 13.8007 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 14.2558 + "tflops": 18.0767 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 14.977 + "tflops": 19.4068 } ] }, @@ -5119,67 +6847,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.6384 + "tflops": 0.6353 }, { "batch_size": 2, "kernel": "wvsplitk_int4", - "tflops": 1.2171 + "tflops": 1.2122 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 1.5234 + "tflops": 0.2488 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.042 + "tflops": 0.4927 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 6.0245 + "tflops": 0.9923 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 11.6611 + "tflops": 1.9496 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 10.7429 + "tflops": 3.8487 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 16.9008 + "tflops": 7.0428 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 20.7589 + "tflops": 11.0587 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 18.9066 + "tflops": 12.228 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 20.8939 + "tflops": 11.9532 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 13.0903 + "tflops": 15.0609 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 13.9524 + "tflops": 16.0397 } ] } @@ -5197,67 +6925,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.885 + "tflops": 0.8834 }, { "batch_size": 2, "kernel": "hybrid_triton_w4a16", - "tflops": 1.1035 + "tflops": 0.3112 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 2.1331 + "tflops": 0.6217 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.1602 + "tflops": 1.2358 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 8.2714 + "tflops": 2.3913 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 15.8243 + "tflops": 4.6372 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 17.1115 + "tflops": 8.9877 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 17.2223 + "tflops": 16.7528 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 21.7413 + "tflops": 25.7123 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 23.4659 + "tflops": 32.0589 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 23.6745 + "tflops": 31.6928 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 22.2756 + "tflops": 31.9897 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 22.8942 + "tflops": 31.9872 } ] }, @@ -5267,67 +6995,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8619 + "tflops": 0.8613 }, { "batch_size": 2, "kernel": "hybrid_triton_w4a16", - "tflops": 1.0915 + "tflops": 0.34 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 2.069 + "tflops": 0.6773 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.943 + "tflops": 1.3479 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 7.8982 + "tflops": 2.5977 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 15.3026 + "tflops": 4.9653 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 15.3981 + "tflops": 9.4957 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 16.8937 + "tflops": 18.089 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 21.6004 + "tflops": 26.6892 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 22.5615 + "tflops": 33.1583 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 23.6987 + "tflops": 32.4939 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 21.6381 + "tflops": 31.7021 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 21.5786 + "tflops": 32.0314 } ] }, @@ -5337,67 +7065,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8809 + "tflops": 0.8817 }, { "batch_size": 2, "kernel": "hybrid_triton_w4a16", - "tflops": 1.0253 + "tflops": 0.2102 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 2.0285 + "tflops": 0.4184 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 4.0189 + "tflops": 0.8331 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 7.9354 + "tflops": 1.6359 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 15.4095 + "tflops": 3.1951 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 18.0048 + "tflops": 6.2688 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 17.3915 + "tflops": 12.0409 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 22.9484 + "tflops": 17.6806 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 25.4824 + "tflops": 17.344 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 25.1519 + "tflops": 21.4057 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 22.6301 + "tflops": 21.3762 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 23.1348 + "tflops": 21.1222 } ] }, @@ -5407,67 +7135,67 @@ { "batch_size": 1, "kernel": "wvsplitk_int4", - "tflops": 0.8576 + "tflops": 0.8712 }, { "batch_size": 2, "kernel": "hybrid_triton_w4a16", - "tflops": 0.9376 + "tflops": 0.2052 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 1.8925 + "tflops": 0.4104 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.7333 + "tflops": 0.818 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 7.3705 + "tflops": 1.6027 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 14.2827 + "tflops": 3.083 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 16.4865 + "tflops": 6.0017 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 16.557 + "tflops": 11.5772 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 21.0001 + "tflops": 17.2713 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 22.8224 + "tflops": 17.0321 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 23.1666 + "tflops": 20.9435 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 22.1723 + "tflops": 21.0122 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 22.0013 + "tflops": 17.2395 } ] } @@ -5485,67 +7213,67 @@ { "batch_size": 1, "kernel": "hybrid_triton_w4a16", - "tflops": 0.424 + "tflops": 0.0993 }, { "batch_size": 2, "kernel": "hybrid_triton_w4a16", - "tflops": 0.8228 + "tflops": 0.198 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 1.5955 + "tflops": 0.3962 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 3.1065 + "tflops": 0.7841 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 5.4339 + "tflops": 1.5226 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 9.1693 + "tflops": 2.9346 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 6.7792 + "tflops": 5.2995 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 17.9705 + "tflops": 9.3127 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 22.3488 + "tflops": 16.8916 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 20.9475 + "tflops": 18.1418 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 19.9299 + "tflops": 15.3485 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 15.6152 + "tflops": 21.2461 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 16.0634 + "tflops": 18.7303 } ] }, @@ -5555,67 +7283,67 @@ { "batch_size": 1, "kernel": "hybrid_triton_w4a16", - "tflops": 0.3508 + "tflops": 0.1251 }, { "batch_size": 2, "kernel": "hybrid_triton_w4a16", - "tflops": 0.6919 + "tflops": 0.2476 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 1.3579 + "tflops": 0.4891 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 2.568 + "tflops": 0.9438 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 4.5435 + "tflops": 1.7556 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 7.6452 + "tflops": 3.3161 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 6.5037 + "tflops": 5.984 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 17.3892 + "tflops": 11.1626 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 22.0346 + "tflops": 17.0845 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 20.5223 + "tflops": 17.331 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 19.4163 + "tflops": 14.682 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 13.7689 + "tflops": 19.4676 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 14.7043 + "tflops": 19.1259 } ] }, @@ -5625,67 +7353,67 @@ { "batch_size": 1, "kernel": "hybrid_triton_w4a16", - "tflops": 0.3789 + "tflops": 0.0948 }, { "batch_size": 2, "kernel": "hybrid_triton_w4a16", - "tflops": 0.7366 + "tflops": 0.1894 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 1.4298 + "tflops": 0.3739 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 2.691 + "tflops": 0.7294 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 4.9467 + "tflops": 1.3775 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 8.4458 + "tflops": 2.6444 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 6.9205 + "tflops": 4.9281 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 18.3407 + "tflops": 9.0306 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 24.098 + "tflops": 15.6083 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 22.5525 + "tflops": 17.2335 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 19.3027 + "tflops": 17.2162 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 15.0436 + "tflops": 18.595 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 15.5594 + "tflops": 18.0164 } ] }, @@ -5695,67 +7423,67 @@ { "batch_size": 1, "kernel": "hybrid_triton_w4a16", - "tflops": 0.3533 + "tflops": 0.0772 }, { "batch_size": 2, "kernel": "hybrid_triton_w4a16", - "tflops": 0.7015 + "tflops": 0.1547 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 1.4 + "tflops": 0.307 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 2.7728 + "tflops": 0.6088 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 5.4103 + "tflops": 1.2005 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 9.9436 + "tflops": 2.3358 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 6.8438 + "tflops": 4.261 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 17.0035 + "tflops": 7.7459 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 22.0357 + "tflops": 12.7208 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 20.4016 + "tflops": 14.7178 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 19.316 + "tflops": 14.4769 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 13.6848 + "tflops": 15.7824 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 14.0292 + "tflops": 14.7408 } ] } @@ -5773,67 +7501,67 @@ { "batch_size": 1, "kernel": "hybrid_triton_w4a16", - "tflops": 0.2254 + "tflops": 0.0417 }, { "batch_size": 2, "kernel": "hybrid_triton_w4a16", - "tflops": 0.4478 + "tflops": 0.0833 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 0.8837 + "tflops": 0.1667 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 1.7114 + "tflops": 0.3308 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 3.3617 + "tflops": 0.6525 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 6.0897 + "tflops": 1.2675 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 5.2481 + "tflops": 2.4392 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 13.251 + "tflops": 4.6133 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 17.8496 + "tflops": 10.2592 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 17.0149 + "tflops": 13.4064 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 19.0741 + "tflops": 12.2536 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 13.2297 + "tflops": 17.7417 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 14.7638 + "tflops": 17.7476 } ] }, @@ -5843,67 +7571,67 @@ { "batch_size": 1, "kernel": "hybrid_triton_w4a16", - "tflops": 0.1845 + "tflops": 0.0561 }, { "batch_size": 2, "kernel": "hybrid_triton_w4a16", - "tflops": 0.3661 + "tflops": 0.115 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 0.7269 + "tflops": 0.2272 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 1.3789 + "tflops": 0.4525 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 2.6411 + "tflops": 0.8805 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 4.6507 + "tflops": 1.6714 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 4.5358 + "tflops": 3.0679 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 12.7032 + "tflops": 6.451 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 16.6265 + "tflops": 9.9335 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 15.7705 + "tflops": 13.3028 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 17.5845 + "tflops": 13.2313 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 12.3266 + "tflops": 17.8574 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 13.3552 + "tflops": 17.4941 } ] }, @@ -5913,67 +7641,67 @@ { "batch_size": 1, "kernel": "hybrid_triton_w4a16", - "tflops": 0.2229 + "tflops": 0.0499 }, { "batch_size": 2, "kernel": "hybrid_triton_w4a16", - "tflops": 0.4398 + "tflops": 0.0991 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 0.8398 + "tflops": 0.1965 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 1.5984 + "tflops": 0.3906 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 3.0584 + "tflops": 0.7668 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 5.2429 + "tflops": 1.4463 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 5.2641 + "tflops": 2.6992 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 12.7121 + "tflops": 5.0535 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 18.687 + "tflops": 10.0347 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 18.4856 + "tflops": 12.9205 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 18.8484 + "tflops": 13.4907 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 12.7829 + "tflops": 17.7347 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 13.8779 + "tflops": 16.9656 } ] }, @@ -5983,67 +7711,67 @@ { "batch_size": 1, "kernel": "hybrid_triton_w4a16", - "tflops": 0.2218 + "tflops": 0.034 }, { "batch_size": 2, "kernel": "hybrid_triton_w4a16", - "tflops": 0.4401 + "tflops": 0.0678 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 0.8749 + "tflops": 0.1325 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 1.727 + "tflops": 0.2665 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 3.3501 + "tflops": 0.5328 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 6.1865 + "tflops": 1.0522 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 4.6114 + "tflops": 2.1253 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 12.5741 + "tflops": 4.2888 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 16.5046 + "tflops": 8.5998 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 15.7934 + "tflops": 11.9984 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 17.7758 + "tflops": 13.0135 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 11.586 + "tflops": 14.2636 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 12.6513 + "tflops": 14.0604 } ] } @@ -6061,67 +7789,67 @@ { "batch_size": 1, "kernel": "hybrid_triton_w4a16", - "tflops": 0.262 + "tflops": 0.0446 }, { "batch_size": 2, "kernel": "hybrid_triton_w4a16", - "tflops": 0.5001 + "tflops": 0.0893 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 0.969 + "tflops": 0.1786 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 2.0626 + "tflops": 0.3532 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 4.0089 + "tflops": 0.7256 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 7.4095 + "tflops": 1.4251 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 6.0759 + "tflops": 2.8447 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 9.6046 + "tflops": 5.7896 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 11.9319 + "tflops": 11.4974 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 18.156 + "tflops": 11.4417 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 21.0747 + "tflops": 15.2352 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 14.6117 + "tflops": 17.344 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 14.811 + "tflops": 17.4632 } ] }, @@ -6131,67 +7859,67 @@ { "batch_size": 1, "kernel": "hybrid_triton_w4a16", - "tflops": 0.1896 + "tflops": 0.0894 }, { "batch_size": 2, "kernel": "hybrid_triton_w4a16", - "tflops": 0.372 + "tflops": 0.1689 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 0.7141 + "tflops": 0.3262 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 1.4814 + "tflops": 0.6347 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 3.0375 + "tflops": 1.1713 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 5.8263 + "tflops": 2.2273 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 5.305 + "tflops": 4.1185 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 10.1078 + "tflops": 6.2779 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 11.497 + "tflops": 10.33 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 15.3731 + "tflops": 10.2453 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 19.4851 + "tflops": 15.0855 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 13.4057 + "tflops": 18.019 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 12.9403 + "tflops": 17.1495 } ] }, @@ -6201,67 +7929,67 @@ { "batch_size": 1, "kernel": "hybrid_triton_w4a16", - "tflops": 0.246 + "tflops": 0.0747 }, { "batch_size": 2, "kernel": "hybrid_triton_w4a16", - "tflops": 0.4896 + "tflops": 0.1472 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 0.9348 + "tflops": 0.2831 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 1.9293 + "tflops": 0.5603 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 3.877 + "tflops": 1.0748 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 7.2312 + "tflops": 2.0612 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 5.8312 + "tflops": 3.8849 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 14.6063 + "tflops": 5.629 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 12.3716 + "tflops": 10.4749 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 19.6078 + "tflops": 11.7025 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 19.7853 + "tflops": 15.2197 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 14.0654 + "tflops": 17.2658 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 14.574 + "tflops": 16.8417 } ] }, @@ -6271,67 +7999,67 @@ { "batch_size": 1, "kernel": "hybrid_triton_w4a16", - "tflops": 0.2389 + "tflops": 0.0502 }, { "batch_size": 2, "kernel": "hybrid_triton_w4a16", - "tflops": 0.464 + "tflops": 0.1023 }, { "batch_size": 4, "kernel": "hybrid_triton_w4a16", - "tflops": 0.9057 + "tflops": 0.2006 }, { "batch_size": 8, "kernel": "hybrid_triton_w4a16", - "tflops": 1.8822 + "tflops": 0.4278 }, { "batch_size": 16, "kernel": "hybrid_triton_w4a16", - "tflops": 3.6763 + "tflops": 0.7971 }, { "batch_size": 32, "kernel": "hybrid_triton_w4a16", - "tflops": 6.7128 + "tflops": 1.4958 }, { "batch_size": 64, "kernel": "hybrid_triton_w4a16", - "tflops": 5.0969 + "tflops": 2.7336 }, { "batch_size": 128, "kernel": "hybrid_triton_w4a16", - "tflops": 13.7095 + "tflops": 5.1791 }, { "batch_size": 256, "kernel": "hybrid_triton_w4a16", - "tflops": 10.2957 + "tflops": 9.0843 }, { "batch_size": 512, "kernel": "hybrid_triton_w4a16", - "tflops": 14.9646 + "tflops": 9.1883 }, { "batch_size": 1024, "kernel": "hybrid_triton_w4a16", - "tflops": 19.1935 + "tflops": 13.5532 }, { "batch_size": 2048, "kernel": "hybrid_triton_w4a16", - "tflops": 12.6953 + "tflops": 14.2143 }, { "batch_size": 4096, "kernel": "hybrid_triton_w4a16", - "tflops": 12.7153 + "tflops": 13.9764 } ] } diff --git a/tests/kernels/quantization/test_hybrid_w4a16_perf.py b/tests/kernels/quantization/test_hybrid_w4a16_perf.py index ee3138572fc7..302a31590e17 100644 --- a/tests/kernels/quantization/test_hybrid_w4a16_perf.py +++ b/tests/kernels/quantization/test_hybrid_w4a16_perf.py @@ -217,6 +217,43 @@ def _log_temp(config: Any, label: str) -> float: "group_size": 128, "comment": "L2 2MiB above", }, + { + "in_features": 2048, + "out_features": 4096, + "group_size": 128, + "comment": "Qwen3-1.7B qkv_proj", + }, + { + "in_features": 2048, + "out_features": 12288, + "group_size": 128, + "comment": "Qwen3-1.7B gate_up_proj", + }, + { + "in_features": 6144, + "out_features": 2048, + "group_size": 128, + "comment": "Qwen3-1.7B down_proj", + }, + # RedHatAI/gemma-3-4b-it-quantized.w4a16 (g=128) + { + "in_features": 2560, + "out_features": 4096, + "group_size": 128, + "comment": "Gemma3-4B qkv_proj", + }, + { + "in_features": 2560, + "out_features": 20480, + "group_size": 128, + "comment": "Gemma3-4B gate_up_proj", + }, + { + "in_features": 10240, + "out_features": 2560, + "group_size": 128, + "comment": "Gemma3-4B down_proj", + }, ] # Provider naming convention: "[-zp][-bf16]". Suffix -zp selects the @@ -356,6 +393,34 @@ def prepare_hybrid_weights( } +def _compute_packed_scale_zp( + w_s: torch.Tensor, + w_zp: torch.Tensor | None, + dtype: torch.dtype, +) -> torch.Tensor | None: + """Pack the per-group packed scale/zp carrier into one fp32, matching the load-time + carrier built in ``HybridW4A16LinearKernel.process_weights_after_loading``. + + Built ONLY for asymmetric layers (``w_zp`` given); returns None for symmetric + (the kernel uses the constant -8 offset there, no carrier). Low 16 bits = + scale; high 16 bits are dtype-specific: + fp16: bias_eff = -(8*scale + (zp-8)*scale) — magic-constant fp16 FMA dequant. + bf16: zp_int = raw zp 0..15 — int-domain subtract (bit-identical to the + separate scale + zp loads). + """ + if w_zp is None or dtype not in (torch.float16, torch.bfloat16): + return None + scale_u16 = w_s.view(torch.uint16).to(torch.int32) & 0xFFFF + if dtype == torch.float16: + w_s_f32 = w_s.to(torch.float32) + scaled_zp_f32 = (w_zp.to(torch.float32) - 8.0) * w_s_f32 + bias_eff = (-(8.0 * w_s_f32 + scaled_zp_f32)).to(dtype) + hi_u16 = bias_eff.contiguous().view(torch.uint16).to(torch.int32) & 0xFFFF + else: + hi_u16 = w_zp.to(torch.int32) & 0xFFFF + return ((hi_u16 << 16) | scale_u16).view(torch.float32).contiguous() + + # --------------------------------------------------------------------------- # Core measurement # --------------------------------------------------------------------------- @@ -400,6 +465,10 @@ def measure_tflops( cu_count = num_compute_units() use_zp = _provider_use_zp(provider) + # packing zero point and scales for faster access + packed_scale_zp = _compute_packed_scale_zp( + weights["w_s_skinny"], weights["w_zp"] if use_zp else None, dtype + ) def run(): return _hybrid_w4a16_apply_impl( @@ -411,6 +480,7 @@ def run(): None, # bias cu_count, group_size, + packed_scale_zp, ) ms = triton.testing.do_bench_cudagraph(run, quantiles=[0.5]) diff --git a/tests/kernels/quantization/test_hybrid_w4a16_triton.py b/tests/kernels/quantization/test_hybrid_w4a16_triton.py index ffef412318c8..d6cbaa629c00 100644 --- a/tests/kernels/quantization/test_hybrid_w4a16_triton.py +++ b/tests/kernels/quantization/test_hybrid_w4a16_triton.py @@ -105,7 +105,6 @@ def test_triton_w4a16_skinny_fmt_gemm_matches_reference( b_q=b_packed, scales=scales, group_size=G, - zp_bias=8, ) ref = _w4a16_skinny_reference( a, @@ -146,6 +145,24 @@ def _w4a16_skinny_reference_asymmetric( return out.to(a_mk.dtype) +def _pack_scale_zp( + scales_nkg: torch.Tensor, zp_nkg: torch.Tensor, dtype: torch.dtype +) -> torch.Tensor: + """Build the asymmetric PackedSb carrier [N, K//G] fp32 that + triton_w4a16_skinny_fmt_gemm consumes (low 16 bits = scale; high 16 bits = + fp16 bias_eff = -(8 + (zp-8))*scale, or bf16 integer zp). Mirrors + HybridW4A16LinearKernel.process_weights_after_loading. + """ + scale_u16 = scales_nkg.contiguous().view(torch.uint16).to(torch.int32) & 0xFFFF + if dtype == torch.float16: + s32 = scales_nkg.to(torch.float32) + bias_eff = (-(8.0 * s32 + (zp_nkg.to(torch.float32) - 8.0) * s32)).to(dtype) + hi_u16 = bias_eff.contiguous().view(torch.uint16).to(torch.int32) & 0xFFFF + else: + hi_u16 = zp_nkg.to(torch.int32) & 0xFFFF + return ((hi_u16 << 16) | scale_u16).view(torch.float32).contiguous() + + @pytest.mark.skipif(not current_platform.is_rocm(), reason="ROCm only") @pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16]) @pytest.mark.parametrize( @@ -185,7 +202,7 @@ def test_triton_w4a16_skinny_fmt_gemm_asymmetric(dtype, M, K, N, G, random_seed: b_q=b_packed, scales=scales, group_size=G, - zp=zp, + packed_scale_zp=_pack_scale_zp(scales, zp, dtype), ) ref = _w4a16_skinny_reference_asymmetric( a, diff --git a/vllm/model_executor/kernels/linear/mixed_precision/hybrid_w4a16.py b/vllm/model_executor/kernels/linear/mixed_precision/hybrid_w4a16.py index 96219d4ebc52..322426677e42 100644 --- a/vllm/model_executor/kernels/linear/mixed_precision/hybrid_w4a16.py +++ b/vllm/model_executor/kernels/linear/mixed_precision/hybrid_w4a16.py @@ -13,6 +13,7 @@ """ from contextlib import nullcontext +import os import torch @@ -44,13 +45,46 @@ # --------------------------------------------------------------------------- +@tl.target_info.constexpr_function +def _target_is_gfx1x() -> bool: + """Compile-time True on RDNA gfx11/gfx12 (where the v_and_or_b32 packed + dequant is validated/tuned).""" + target = tl.target_info.current_target() + if target is None or target.backend != "hip": + return False + arch = str(target.arch) + return arch.startswith("gfx11") or arch.startswith("gfx12") + + +@triton.jit +def _int4_pair_to_fp16x2(x): + """Unpack two packed int4 nibbles into a uint32 holding two fp16 lanes, + each equal to 1024 + nibble, with one ``v_and_or_b32`` + (``(x & 0x000F000F) | 0x64006400``). + + OR-ing a 4-bit nibble into the low mantissa of fp16 1024.0 (0x6400) + bitcasts to exactly 1024+n (CK's i4_to_half trick). Doing it on a full + 32-bit lane dequants two nibbles per instruction, vs the scalar + v_and_b16 + v_or_b16 pair Triton emits from the elementwise form. + """ + mask = tl.full(x.shape, 0x000F000F, tl.int32) + return tl.inline_asm_elementwise( + asm="v_and_or_b32 $0, $1, $2, 0x64006400", + constraints="=v,v,v", + args=[x, mask], + dtype=tl.uint32, + is_pure=True, + pack=1, + ) + + @triton.jit def _triton_w4a16_skinny_fmt_kernel( # Pointers a_ptr, # [M, K] fp16/bf16 activations b_ptr, # [N, K//8] int32 packed (ExLlama shuffle, K is packed dim) - scales_ptr, # [N, K//G] fp16/bf16 scales (skinny layout) - zp_ptr, # [N, K//G] fp16/bf16 raw zero-points (when HAS_ZP=True) + scales_ptr, # [N, K//G] fp16/bf16 scales (sym path, HAS_ZP=False) + packed_scale_zp_ptr, # [N, K//G] int32 scale/zp carrier (asym, HAS_ZP) c_ptr, # [M, N] fp16/bf16 output # Dimensions M, @@ -58,27 +92,32 @@ def _triton_w4a16_skinny_fmt_kernel( K, K8, # K // 8 num_groups, # K // group_size - # Quantization parameters group_size, - ZP_BIAS: tl.constexpr, - HAS_ZP: tl.constexpr, + HAS_ZP: tl.constexpr, # asym: read the scale/zp carrier; sym: scales + (-8) # Block sizes BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr, ): """ - Fused W4A16 GEMM reading weights from skinny format [N, K//8]. + Fused W4A16 GEMM reading skinny weights [N, K//8]. B is stored as [N, K//8] int32 using ExLlama shuffle packing: each int32 packs 8 K-values with interleave [0,2,4,6,1,3,5,7]: packed = val[0] | (val[2]<<4) | (val[4]<<8) | (val[6]<<12) | (val[1]<<16) | (val[3]<<20) | (val[5]<<24) | (val[7]<<28) - Scales are [N, K//G] (skinny layout, NOT transposed). - When HAS_ZP=True, raw zero-points zp_raw are loaded from zp_ptr [N, K//G] - and subtracted directly: (nibble - zp_raw) * scale. - When HAS_ZP=False, only the constant ZP_BIAS is subtracted (symmetric). + Two dequant paths, chosen at the layer's sym/asym nature: + - HAS_ZP=True (asymmetric): read the carrier ``packed_scale_zp_ptr`` + [N, K//G] (one fp32 per (n, group)) — it folds the per-group scale AND + the zero-point offset into a single load, replacing the separate scale + + zp loads. Layout: fp16 = scale | bias_eff (= -8*scale - scaled_zp), + dequant (nibble-1024)*scale + bias_eff via the magic-const fp16 unpack; + bf16 = scale | zp_int, dequant (nibble - zp_int)*scale. + - HAS_ZP=False (symmetric): the -8 offset is a constant, so there is + no second load to fold — read ``scales_ptr`` directly and subtract the + constant 8. fp16: (nibble - 1032)*scale via the magic unpack; bf16: + (nibble - 8)*scale. (No carrier overhead for the sym fast path.) """ pid_m = tl.program_id(0) pid_n = tl.program_id(1) @@ -114,43 +153,85 @@ def _triton_w4a16_skinny_fmt_kernel( mask_b = (offs_n[:, None] < N) & (offs_k8[None, :] < K8) b_packed = tl.load(b_ptrs, mask=mask_b, other=0) - # ---- Unpack int4 weights with ExLlama unshuffle ---- - b = tl.interleave(b_packed, b_packed) - b = tl.interleave(b, b) - b = tl.interleave(b, b) - b = (b >> shifts_full) & 0xF # [BLOCK_N, BLOCK_K] + # ---- Unpack int4 weights ---- + # The packed v_and_or_b32 / v_pk_fma dequant is fp16-only (the 1024+n + # magic trick needs fp16's mantissa) and only validated/tuned on RDNA + # gfx11/gfx12. Decided here at compile time (dtype + target arch) so + # callers pass no flag; everything else uses the scalar unpack. The + # condition is written inline (not via a local) so Triton constexpr- + # eliminates the dead arm — the per-dtype vars below are only defined + # on the taken path. + if (a.dtype == tl.float16) and _target_is_gfx1x(): + # Packed dequant (fp16). The ExLlama int32 holds the + # paired nibbles val[2p] @ bits[4p:4p+4] and val[2p+1] @ + # bits[16+4p:20+4p], so for pre-shift 4p (p=0..3), + # (x >> 4p) & 0x000F000F | 0x64006400 + # is one v_and_or_b32 producing a half2 = (1024+val[2p], + # 1024+val[2p+1]) in K order (signed shift is fine: the sign fill + # lands above bit 20, masked out). This dequants TWO nibbles per + # instruction; the elementwise form lowers to scalar v_and_b16 + + # v_or_b16 (1 nibble each). The interleave(lo, hi) lays b_raw out as + # half2 so the downstream affine also packs into v_pk_fma_f16. The + # dequant inner loop is VALU-issue-bound on gfx11, so this ~halves + # the dequant instruction count per WMMA and matches CK. + shifts4 = (tl.arange(0, 4) * 4)[None, None, :] + bp_shift = tl.reshape( + b_packed[:, :, None] >> shifts4, (BLOCK_N, BLOCK_K // 2) + ) + packed_hl = _int4_pair_to_fp16x2(bp_shift) # u32 half2: 1024+nibble pair + lo = (packed_hl & 0xFFFF).to(tl.uint16).to(tl.float16, bitcast=True) + hi = (packed_hl >> 16).to(tl.uint16).to(tl.float16, bitcast=True) + b_raw = tl.interleave(lo, hi) # [BLOCK_N, BLOCK_K] fp16 = 1024+nibble + else: + # ExLlama unshuffle: replicate each int32 8x then per-lane shift+mask. + b = tl.interleave(b_packed, b_packed) + b = tl.interleave(b, b) + b = tl.interleave(b, b) + b = (b >> shifts_full) & 0xF # [BLOCK_N, BLOCK_K] - # ---- Load scales from [N, K//G] layout ---- + # ---- Per-group quant params from [N, K//G] layout ---- g_idx = (k_start * BLOCK_K) // group_size - scale_ptrs = scales_ptr + offs_n * num_groups + g_idx scale_mask = offs_n < N - scales = tl.load(scale_ptrs, mask=scale_mask, other=1.0) # ---- Dequantize ---- if HAS_ZP: - zp_ptrs = zp_ptr + offs_n * num_groups + g_idx - zp_raw = tl.load(zp_ptrs, mask=scale_mask, other=0.0) - if scales.dtype == tl.bfloat16: - # bf16: subtract zp in INT (zp values are 0..15, exact - # roundtrip), then cast once. This mirrors the symmetric - # path and avoids the per-tile int->bf16 cast of the full - # [BLOCK_N, BLOCK_K] nibble block, which is the bottleneck - # for asymmetric w4a16 bf16 prefill on RDNA3.5 (Strix Halo, - # gfx1151). Casting zp_raw (only BLOCK_N elements) to int - # is cheap. Recovers ~14% TFLOPS across Qwen3-8B w4a16 - # prefill projections without changing fp16 behavior. - zp_int = zp_raw.to(b.dtype) - b_fp = (b - zp_int[:, None]).to(scales.dtype) * scales[:, None] + # Asymmetric: packed scale/zp carrier (one fp32/group folds scale + zp). + psz = tl.load( + packed_scale_zp_ptr + offs_n * num_groups + g_idx, + mask=scale_mask, + other=0, + ) + psz_u = psz.to(tl.uint32, bitcast=True) + if a.dtype == tl.float16: + # fp16: low16 = scale, high16 = bias_eff (= -8*scale - scaled_zp). + # ONE fp16 FMA per group via the magic-constant i4->fp16 unpack. + scale = (psz_u & 0xFFFF).to(tl.uint16).to(tl.float16, bitcast=True) + bias_eff = (psz_u >> 16).to(tl.uint16).to(tl.float16, bitcast=True) + if not _target_is_gfx1x(): + b_raw = (b | 0x6400).to(tl.uint16).to(tl.float16, bitcast=True) + c1024 = tl.full((), 1024.0, tl.float16) + b_fp = (b_raw - c1024) * scale[:, None] + bias_eff[:, None] else: - # fp16: original asymmetric path. The int->fp16 cast on - # RDNA3.5 has a direct ISA path and fuses well with the - # subsequent subtraction, so keeping the cast-first order - # avoids a small fp16 regression observed when switching - # both dtypes to the int-subtract-first form. - b_fp = (b.to(scales.dtype) - zp_raw[:, None]) * scales[:, None] + # bf16: low16 = scale, high16 = zp_int. Cheap int-domain subtract + # before the single bf16 multiply (RDNA3 has no v_pk_fma_bf16). + scale = (psz_u & 0xFFFF).to(tl.uint16).to(tl.bfloat16, bitcast=True) + zp_int = ((psz_u >> 16) & 0xFFFF).to(b.dtype) + b_fp = (b - zp_int[:, None]).to(scale.dtype) * scale[:, None] else: - # Symmetric: (w - 8) * scale - b_fp = (b - ZP_BIAS).to(scales.dtype) * scales[:, None] + # Symmetric: the -8 offset is constant (no zp to fold), so read the + # scale directly — no carrier overhead. + scales = tl.load( + scales_ptr + offs_n * num_groups + g_idx, mask=scale_mask, other=1.0 + ) + if a.dtype == tl.float16: + # (nibble - 8) * scale == (b_raw - (1024+8)) * scale, via magic. + if not _target_is_gfx1x(): + b_raw = (b | 0x6400).to(tl.uint16).to(tl.float16, bitcast=True) + c = tl.full((), float(1024 + 8), tl.float16) + b_fp = (b_raw - c) * scales[:, None] + else: + # bf16: (nibble - 8) * scale, int subtract before the cast. + b_fp = (b - 8).to(scales.dtype) * scales[:, None] # ---- Transpose to [BLOCK_K, BLOCK_N] for matmul ---- b_fp_t = tl.trans(b_fp) @@ -165,26 +246,83 @@ def _triton_w4a16_skinny_fmt_kernel( tl.store(c_ptrs, c, mask=mask_c) +# Explicit gfx11 prefill tile selection, distilled from a per-shape +# autotune sweep over the F.2 GEMM catalog (Qwen3 / Gemma / Llama prefill +# shapes, M in {2048, 3968}). The autotuned winner is overwhelmingly +# BLOCK_N=256 / num_warps=8 / BLOCK_K=32 — a wide-N tile keeps the workgroup +# grid large enough to saturate the 40 CUs at M~2-4k while the wide N-tile makes +# the tl.trans -> ds_read_b128 LDS readback efficient. BLOCK_K is fixed at 32 +# (<= every AWQ group size, so no group-boundary scale aliasing; BK>32 was +# measured to regress). This replaces the legacy heuristic on gfx11x; other +# arches keep it. +def _select_skinny_gfx11_config(M: int, N: int, K: int) -> tuple[int, int, int]: + """Return (BLOCK_M, BLOCK_N, num_warps) for the gfx11 prefill skinny GEMM. + + BLOCK_M=128 vs 64, re-tuned on the packed-dequant kernel (the kernel is no + longer ALU-bound, so this differs from the earlier scalar-kernel rule) via a + low-noise interleaved A/B over the full F.2 catalog (<2% CoV, 31 shapes, + zero mispicks): + * deep, non-cliff K (K>=3072 and K%2048!=0) -> BLOCK_M=128: the long + K-loop amortizes the wider M-tile (Qwen7B down -12.5%, gate_up -9.1%). + * very wide N (N>=16384) -> BLOCK_M=128: enough N-tiles that the halved + M-tile count still saturates the 40 CUs (lm_head, wide gate_up). + * otherwise BLOCK_M=64. K%2048==0 is the gfx1151 power-of-2 global-load + cliff where BLOCK_M=64's smaller A footprint is much faster unless N is + very wide (Gemma2B q/o, LLaMA8B q/o, Qwen3.5 GDN would lose 17-31% at + BLOCK_M=128); shallow K=2560 at small N also prefers 64. + """ + forced_prefill_tile = os.getenv("VLLM_W4A16_PREFILL_TILE") or os.getenv( + "VLLM_W4A16_671_TILE" + ) + if forced_prefill_tile: + parts = [p.strip() for p in forced_prefill_tile.split(",")] + if len(parts) == 3 and all(p.isdigit() for p in parts): + return int(parts[0]), int(parts[1]), int(parts[2]) + + # Profile-guided default for Qwen3-VL-like single-request multimodal prefill + # in the common token-length band around the current workload. + # Keep this limited to recurring Qwen3 projection shapes to reduce risk for + # unrelated models while preserving the measured TTFT win. + qwen3_prefill_shapes = { + (19456, 2560), # gate_up_proj-like + (2560, 9728), # down_proj-like + (6144, 2560), # qkv_proj-like + (2560, 4096), # o_proj-like + } + if 576 <= M <= 832 and (N, K) in qwen3_prefill_shapes: + return 64, 256, 8 + + block_n, num_warps = 256, 8 + block_m = 128 if ((K >= 3072 and K % 2048 != 0) or N >= 16384) else 64 + return block_m, block_n, num_warps + + def triton_w4a16_skinny_fmt_gemm( a: torch.Tensor, # [M, K] fp16/bf16 b_q: torch.Tensor, # [N, K//8] int32 (ExLlama shuffle packed) - scales: torch.Tensor, # [N, K//G] fp16/bf16 + scales: torch.Tensor, # [N, K//G] fp16/bf16 (used for the symmetric path) group_size: int, - zp_bias: int = 8, - zp: torch.Tensor | None = None, # [N, K//G] per-group zero-points + out: torch.Tensor | None = None, # [M, N] optional pre-allocated output + packed_scale_zp: torch.Tensor | None = None, # [N, K//G] fp32 carrier (asym only) ) -> torch.Tensor: """ - Fused W4A16 GEMM reading from skinny weight format [N, K//8]. + Fused W4A16 GEMM reading skinny weights [N, K//8]. + + Asymmetric layers pass ``packed_scale_zp`` (the carrier that folds scale + + zero-point into one load); symmetric layers leave it None and the kernel + reads ``scales`` directly with a constant -8 offset (no carrier overhead — + sym has no second load to fold). Args: a: Activation matrix [M, K], float16 or bfloat16. b_q: Packed weight matrix [N, K//8], int32 (ExLlama shuffle). - scales: Per-group scales [N, K//G], same dtype as a. + scales: Per-group scales [N, K//G], same dtype as a (symmetric path). group_size: Quantization group size (resolved from -1 to K by caller). - zp_bias: Constant zero bias (default 8 for unsigned int4). - zp: Raw per-group zero-points [N, K//G] (asymmetric), - stored as zp_raw in activation dtype. When provided, - dequant is (nibble - zp_raw) * scale. + out: Optional pre-allocated [M, N] output. + packed_scale_zp: Optional packed scale/zp carrier [N, K//G] fp32 for asymmetric + layers; layout is dtype-specific (fp16: scale|bias_eff; bf16: + scale|zp_int) — see the kernel docstring. When None, the + symmetric path is used. Returns: Output matrix [M, N], same dtype as a. @@ -202,14 +340,56 @@ def triton_w4a16_skinny_fmt_gemm( assert scales.shape == (N, num_groups), ( f"scales shape mismatch: {scales.shape} vs ({N}, {num_groups})" ) - if zp is not None: - assert zp.is_contiguous(), "Zero-points must be contiguous" - assert zp.shape == (N, num_groups), ( - f"zp shape mismatch: {zp.shape} vs ({N}, {num_groups})" + has_zp = packed_scale_zp is not None + if packed_scale_zp is not None: + assert packed_scale_zp.is_contiguous(), "packed_scale_zp must be contiguous" + assert packed_scale_zp.shape == (N, num_groups), ( + f"packed_scale_zp shape mismatch: {packed_scale_zp.shape} " + f"vs ({N}, {num_groups})" ) - has_zp = zp is not None + packed_scale_zp_i32 = packed_scale_zp.view(torch.int32) + else: + packed_scale_zp_i32 = scales # dummy pointer (unused when HAS_ZP=False) - c = torch.empty((M, N), dtype=a.dtype, device=a.device) + if out is None: + c = torch.empty((M, N), dtype=a.dtype, device=a.device) + else: + assert out.shape == (M, N), f"out shape mismatch: {out.shape} vs ({M}, {N})" + assert out.dtype == a.dtype, f"out dtype mismatch: {out.dtype} vs {a.dtype}" + assert out.device == a.device, ( + f"out device mismatch: {out.device} vs {a.device}" + ) + assert out.is_contiguous(), "out must be contiguous" + c = out + + # On gfx11x, select the tile config per (M, N, K) from the + # distilled table (see _select_skinny_gfx11_config). BLOCK_K is fixed at 32, + # which is <= every supported AWQ group size (no group-boundary scale + # aliasing). + if on_gfx1x(): + block_m, block_n, num_warps = _select_skinny_gfx11_config(M, N, K) + grid = (triton.cdiv(M, block_m), triton.cdiv(N, block_n)) + # The kernel picks the packed (fp16/gfx1x) vs scalar unpack itself. + _triton_w4a16_skinny_fmt_kernel[grid]( + a, + b_q, + scales, + packed_scale_zp_i32, + c, + M, + N, + K, + K8, + num_groups, + group_size=group_size, + HAS_ZP=has_zp, + BLOCK_M=block_m, + BLOCK_N=block_n, + BLOCK_K=32, + num_warps=num_warps, + num_stages=1, # >1 regresses badly for this kernel (no SW pipeline) + ) + return c # AMD-specific scheduling hint; only consumed by the HIP backend below # (see compiler.py amdgpu-waves-per-eu attribute). Set to 0 by default @@ -250,55 +430,6 @@ def triton_w4a16_skinny_fmt_gemm( BLOCK_M, BLOCK_N, BLOCK_K, num_warps = 256, 64, 64, 8 else: BLOCK_M, BLOCK_N, BLOCK_K, num_warps = 128, 128, 32, 8 - elif on_gfx1x(): - # Tuned on gfx1151 (Strix Halo, 40 CUs, 32-wide wavefronts) - # using Qwen3-4B weight shapes with group_size=128. - # waves_per_eu=0 means no constraint; specific values pin LLVM - # to a target VGPR budget per occupancy.md (gfx1151 has 1536 - # VGPRs/SIMD; waves_per_eu=N sets max VGPRs to ~1536/N). - if M <= 32: - BLOCK_M, BLOCK_N, BLOCK_K, num_warps = 32, 32, 128, 4 - elif M <= 64: - BLOCK_M, BLOCK_N, BLOCK_K, num_warps = 64, 64, 32, 4 - elif M <= 128: - # For K >= 4096 AND N >= 4096, a single config (BN=32, BK=128, - # NW=4) wins on every projection shape across Qwen3-8B and - # Llama-3.1-8B (down/qkv/gate_up/o_proj all gain +23%..+35% vs - # prior shape-specific configs). The wider K-tile escapes WMMA - # latency-bound regime (wmma.md: >= 2 waves/SIMD), and BN=32 - # keeps the workgroup grid large enough to saturate 40 CUs even - # at N up to ~28k. - # - # Small-N or small-K shapes (Qwen3-VL-4B / Qwen3-4B) need the - # legacy shape-specific configs — at N=2560 the BN=32 grid drops - # below the saturation point. - if K >= 4096 and N >= 4096: - BLOCK_M, BLOCK_N, BLOCK_K, num_warps = 64, 32, 128, 4 - # waves_per_eu=6 matches the natural VGPR-bound occupancy - # but explicitly pinning the target gives LLVM a single - # register count to optimize against (compiler.py: "forces - # LLVM to focus on a single register count, simplifies some - # heuristics and may improve scheduling"). +5-8% across all - # 4 K=N=4096 projection shapes. - waves_per_eu = 6 - elif K >= 2 * N: # tall K, small-N down (e.g. Qwen3-VL-4B down) - BLOCK_M, BLOCK_N, BLOCK_K, num_warps = 64, 16, 64, 1 - elif N > K: # wide N, small K (e.g. Qwen3-VL-4B qkv/gate_up) - BLOCK_M, BLOCK_N, BLOCK_K, num_warps = 64, 64, 64, 4 - else: # N ~= K, small K (e.g. Qwen3-VL-4B o_proj) - BLOCK_M, BLOCK_N, BLOCK_K, num_warps = 64, 32, 64, 4 - elif M <= 1024: - if K >= 2 * N: # tall K (e.g. down_proj) - BLOCK_M, BLOCK_N, BLOCK_K, num_warps = 64, 64, 64, 4 - elif N >= 4 * K: # very wide N (e.g. gate_up_proj) - BLOCK_M, BLOCK_N, BLOCK_K, num_warps = 128, 64, 64, 8 - else: - BLOCK_M, BLOCK_N, BLOCK_K, num_warps = 64, 128, 32, 4 - else: - if K >= 2 * N: # tall K (e.g. down_proj) - BLOCK_M, BLOCK_N, BLOCK_K, num_warps = 128, 512, 32, 16 - else: - BLOCK_M, BLOCK_N, BLOCK_K, num_warps = 128, 64, 64, 8 else: num_warps = 4 if M <= 32: @@ -319,7 +450,7 @@ def triton_w4a16_skinny_fmt_gemm( a, b_q, scales, - zp if has_zp else scales, # dummy pointer when no zp (unused) + packed_scale_zp_i32, c, M, N, @@ -327,7 +458,6 @@ def triton_w4a16_skinny_fmt_gemm( K8, num_groups, group_size=group_size, - ZP_BIAS=zp_bias, HAS_ZP=has_zp, BLOCK_M=BLOCK_M, BLOCK_N=BLOCK_N, @@ -377,6 +507,7 @@ def _hybrid_w4a16_apply_impl( bias: torch.Tensor | None, cu_count: int, group_size: int, + packed_scale_zp: torch.Tensor | None = None, ) -> torch.Tensor: """Dispatch between skinny GEMM and Triton based on batch size M. @@ -387,6 +518,8 @@ def _hybrid_w4a16_apply_impl( w_zp: [N, K//G] raw zero-points (zp_raw) in act dtype, or None for symmetric. Both HIP skinny and Triton use this single format: dequant = (nibble - zp_raw) * scale. + packed_scale_zp: [N, K//G] fp32 carrier packing scale + zero-point per + group (Triton prefill, asymmetric only), or None for symmetric. Registered as a custom op so torch.compile treats it as opaque. """ @@ -413,12 +546,15 @@ def _hybrid_w4a16_apply_impl( else torch.profiler.record_function(f"hybrid_triton_w4a16 {M}x{N}x{K}") ) with ctx: + # Asymmetric layers carry the packed scale/zp carrier (scale + zero-point folded + # into one load); symmetric layers pass packed_scale_zp=None and the kernel + # reads scales directly with a constant -8 offset (no carrier overhead). output = triton_w4a16_skinny_fmt_gemm( - a=x_2d, - b_q=w_q_i32, - scales=w_s, - group_size=group_size, - zp=w_zp, + x_2d, + w_q_i32, + w_s, + group_size, + packed_scale_zp=packed_scale_zp, ) if bias is not None: output.add_(bias) @@ -434,6 +570,7 @@ def _hybrid_w4a16_apply_fake( bias: torch.Tensor | None, cu_count: int, group_size: int, + packed_scale_zp: torch.Tensor | None = None, ) -> torch.Tensor: M = x_2d.size(0) N = w_q.size(0) @@ -564,6 +701,35 @@ def process_weights_after_loading(self, layer: torch.nn.Module) -> None: torch.nn.Parameter(w_q_skinny_i32, requires_grad=False), ) + # Packed scale/zp carrier for the Triton prefill path — built ONLY + # for asymmetric layers, where it folds the two per-group loads (scale + + # zp) into one. Symmetric layers skip it: the -8 offset is a constant, so + # there is no second load to fold and the carrier would be pure overhead + # (measured ~+8% on fp16 sym); sym reads scales directly instead. + # Layout (matches the kernel's HAS_ZP dequant): + # fp16: low16 = scale, high16 = bias_eff (= -8*scale - (zp-8)*scale). + # Consumed via one fp16 FMA with the magic-constant i4->fp16 unpack. + # bf16: low16 = scale (bf16 bits), high16 = zp_int (raw zp 0..15), as a + # plain integer. Consumed by the int-domain subtract (RDNA3 has no + # v_pk_fma_bf16). Bit-identical to the separate scale+zp loads. + if c.zero_points and c.act_type in (torch.float16, torch.bfloat16): + scale_u16 = w_s_skinny.view(torch.uint16).to(torch.int32) & 0xFFFF + if c.act_type == torch.float16: + w_s_f32 = w_s_skinny.to(torch.float32) + scaled_zp_f32 = (w_zp.to(torch.float32) - 8.0) * w_s_f32 + bias_eff = (-(8.0 * w_s_f32 + scaled_zp_f32)).to(c.act_type) + bias_u16 = bias_eff.contiguous().view(torch.uint16) + hi_u16 = bias_u16.to(torch.int32) & 0xFFFF + else: + hi_u16 = w_zp.to(torch.int32) & 0xFFFF # raw zp 0..15 + packed_scale_zp = ( + ((hi_u16 << 16) | scale_u16).view(torch.float32).contiguous() + ) + layer.register_parameter( + "_hybrid_w_packed_scale_zp", + torch.nn.Parameter(packed_scale_zp, requires_grad=False), + ) + def apply_weights( self, layer: torch.nn.Module, @@ -575,6 +741,8 @@ def apply_weights( c = self.config w_q, w_s, w_zp, _ = self._get_weight_params(layer) w_q_i32 = layer._hybrid_w_q_i32 + # Packed scale/zp carrier (asymmetric layers only; None for sym). + packed_scale_zp = getattr(layer, "_hybrid_w_packed_scale_zp", None) x_2d = x.reshape(-1, x.shape[-1]) N = w_q.shape[0] @@ -590,5 +758,6 @@ def apply_weights( bias, cu_count, c.group_size, + packed_scale_zp, ) return output.reshape(out_shape) diff --git a/vllm/v1/attention/ops/triton_unified_attention.py b/vllm/v1/attention/ops/triton_unified_attention.py index df40f5b65053..76f221b4c7e2 100644 --- a/vllm/v1/attention/ops/triton_unified_attention.py +++ b/vllm/v1/attention/ops/triton_unified_attention.py @@ -35,6 +35,19 @@ float8_info = torch.finfo(current_platform.fp8_dtype()) +def _env_int(name: str) -> int | None: + raw = envs.environment_variables.get(name) if hasattr(envs, "environment_variables") else None + if raw is None: + import os + raw = os.getenv(name) + if not raw: + return None + raw = raw.strip() + if not raw.isdigit(): + return None + return int(raw) + + @triton.jit def _cast_kv_tile(data, Q, tensor_scale, KV_QUANT_MODE: tl.constexpr): """Cast a loaded KV tile to Q's dtype, dequantizing if needed. @@ -492,6 +505,25 @@ def _is_gemma3_attention(head_size: int, sliding_window: int) -> bool: return sliding_window == 1024 and head_size in (128, 256) +def _is_qwen3_vl_4b_prefill_signature( + *, + head_size: int, + num_query_heads: int, + num_kv_heads: int, + max_seqlen_q: int, +) -> bool: + """Detect Qwen3-VL-4B text prefill attention shape. + + This signature is intentionally strict to avoid affecting unrelated models. + """ + return ( + head_size == 128 + and num_query_heads == 32 + and num_kv_heads == 8 + and max_seqlen_q >= 256 + ) + + def _get_tile_size( head_size: int, sliding_window: int, @@ -782,15 +814,57 @@ def unified_attention( grid: tuple[Any, ...] config = {} + tile_size_prefill = TILE_SIZE_PREFILL use_swapped_grid = not use_3d and current_platform.is_gfx1151() + # Optional tuning overrides for targeted profiling/experiments. + # These are intentionally no-op unless explicitly set. + if not use_3d: + # Qwen3-VL-4B prefill default on Navi/gfx11: + # BM64/T32/W4/S1/EU4 was the best end-to-end TTFT setting in local sweeps. + if current_platform.is_navi() and _is_qwen3_vl_4b_prefill_signature( + head_size=head_size, + num_query_heads=num_query_heads, + num_kv_heads=num_kv_heads, + max_seqlen_q=max_seqlen_q, + ): + BLOCK_M = 64 + BLOCK_Q = BLOCK_M // num_queries_per_kv + total_num_q_blocks = q.shape[0] // BLOCK_Q + num_seqs + tile_size_prefill = 32 + num_warps = 4 + num_stages = 1 + waves_per_eu = 4 + + forced_block_m = _env_int("VLLM_UA_PREFILL_BLOCK_M") + if forced_block_m is not None and forced_block_m > 0: + BLOCK_M = max(forced_block_m, triton.next_power_of_2(num_queries_per_kv)) + BLOCK_Q = BLOCK_M // num_queries_per_kv + total_num_q_blocks = q.shape[0] // BLOCK_Q + num_seqs + + forced_tile = _env_int("VLLM_UA_PREFILL_TILE_SIZE") + if forced_tile is not None and forced_tile > 0: + tile_size_prefill = forced_tile + + forced_num_warps = _env_int("VLLM_UA_PREFILL_NUM_WARPS") + if forced_num_warps is not None and forced_num_warps > 0: + num_warps = forced_num_warps + + forced_num_stages = _env_int("VLLM_UA_PREFILL_NUM_STAGES") + if forced_num_stages is not None and forced_num_stages > 0: + num_stages = forced_num_stages + + forced_waves = _env_int("VLLM_UA_PREFILL_WAVES_PER_EU") + if forced_waves is not None and forced_waves > 0: + waves_per_eu = forced_waves + if not use_3d: if use_swapped_grid: grid = (num_kv_heads, total_num_q_blocks) else: grid = (total_num_q_blocks, num_kv_heads) - tile_size = TILE_SIZE_PREFILL + tile_size = tile_size_prefill else: tile_size = TILE_SIZE_DECODE