From 0392b93aba9e8df3439c5b329a041d79ad33ed81 Mon Sep 17 00:00:00 2001 From: Alejandro Acosta Date: Fri, 11 Apr 2025 09:43:08 +0100 Subject: [PATCH 1/8] add gemm shapes --- benchmarks/pvc/input.in | 499 +++++++++++++++++++++++----------------- 1 file changed, 289 insertions(+), 210 deletions(-) diff --git a/benchmarks/pvc/input.in b/benchmarks/pvc/input.in index ce07ff5985..c366a4e1e7 100644 --- a/benchmarks/pvc/input.in +++ b/benchmarks/pvc/input.in @@ -1,211 +1,290 @@ -# GEMM BFloat16 benchmarks -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=4096 --n=4096 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=8192 --k=8192 --n=8192 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1 --k=5120 --n=13824 -PvcGemmBF16BF16FP32_RRR_2 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=28672 --n=8192 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=3072 --k=4096 --n=3072 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4 --k=4096 --n=12288 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=8192 --n=8192 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=8192 --n=32768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=32768 --n=8192 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=8192 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=1024 --n=8192 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=8192 --n=4096 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=4096 --n=8192 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=16384 --n=8192 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=8192 --k=16384 --n=4096 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=16384 --n=8192 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=8192 --k=16384 --n=1024 -# TODO(Codeplay): these are broken -# PvcGemmBF16BF16FP32_RRR_4 --bm_name=bf16_bf16_fp32 --l=4096 --m=8 --k=128 --n=16384 -# PvcGemmBF16BF16FP32_RRR_5 --bm_name=bf16_bf16_fp32 --l=4096 --m=8 --k=16384 --n=128 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=4 --m=32768 --k=128 --n=4096 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=4 --m=32768 --k=4096 --n=128 -PvcGemmBF16BF16FP32_RRR_3 --bm_name=bf16_bf16_fp32 --l=32 --m=4096 --k=4096 --n=128 -PvcGemmBF16BF16FP32_RCR_6 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=4096 --n=4096 -PvcGemmBF16BF16FP32_CRR_7 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=4096 --n=4096 -PvcGemmBF16BF16FP32_CCR_8 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=4096 --n=4096 +# validation error +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1 --k=768 --n=2 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2 --k=1 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=32 --k=128 --n=4 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=128 --n=6 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=512 --n=12 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=128 --n=4 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=512 --n=6 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=9 --n=3072 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=512 --n=4 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=128 --n=12 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=1 --n=3072 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=512 --n=36 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=2048 --n=12 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=1024 --n=12 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=64 --n=6 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=30522 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=2 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=2 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=30522 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=64 --n=12 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=30522 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=64 --n=4 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=128 --n=12 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=50265 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=32 --n=4 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=50265 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=128 --n=4 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=50265 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=768 --n=30522 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=1024 --n=50265 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=2 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=768 --n=2 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=128 --n=6 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=50265 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=1024 --n=50265 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=384 --k=54 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=64 --n=4 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=128 --n=30522 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1536 --n=128100 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=30522 --n=128 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=30522 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=64 --n=36 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=256 --n=2 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=768 --n=50265 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=768 --n=30522 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=512 --n=12 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=384 --n=54 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=512 --n=36 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=2 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=32005 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=128100 --n=1536 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=64 --n=12 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=512 --n=6 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=50257 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1536 --n=2 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=64 --n=6 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=512 --n=4 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=2 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=768 --n=2 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=2 --n=256 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=768 --n=32005 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=2 --n=1536 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=2 --n=4096 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1024 --n=2 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=1 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=1024 --n=12 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=768 --n=2 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=50005 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=50265 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=50265 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=4096 --n=2 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=2 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=50265 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=1024 --n=50265 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=768 --n=50265 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=64 --n=12 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2048 --k=2048 --n=12 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2048 --k=64 --n=12 +# crash +#PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=768 --n=50257 +#PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=768 --n=50005 -PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=8192 --n=8192 -PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=8192 --n=32768 -PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=32768 --n=8192 -PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=8192 --n=1024 -PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=1024 --n=8192 -PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=8192 --n=4096 -PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=4096 --n=8192 -PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=16384 --n=8192 -PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=8192 --k=16384 --n=4096 -PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=16384 --n=8192 -PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=8192 --k=16384 --n=1024 -# PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=4096 --m=8 --k=128 --n=16384 -# PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=4096 --m=8 --k=16384 --n=128 -PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=4 --m=32768 --k=128 --n=4096 -PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=4 --m=32768 --k=4096 --n=128 -PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=32 --m=4096 --k=4096 --n=128 - -PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=8192 --n=8192 -PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=8192 --n=32768 -PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=32768 --n=8192 -PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=8192 --n=1024 -PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=1024 --n=8192 -PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=8192 --n=4096 -PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=4096 --n=8192 -PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=16384 --n=8192 -PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=8192 --k=16384 --n=4096 -PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=16384 --n=8192 -PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=8192 --k=16384 --n=1024 -# PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=4096 --m=8 --k=128 --n=16384 -# PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=4096 --m=8 --k=16384 --n=128 -PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=4 --m=32768 --k=128 --n=4096 -PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=4 --m=32768 --k=4096 --n=128 -PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=32 --m=4096 --k=4096 --n=128 - -# FMHA BFloat16 benchmarks -PvcFMHABF16BF16FP32_RCR_h64_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads=32 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads=32 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=4 --num_heads=48 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=4 --num_heads=48 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=32 --num_heads=16 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=32 --num_heads=16 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=32 --num_heads=8 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=32 --num_heads=8 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=32 --num_heads=4 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=32 --num_heads=4 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=32 --num_heads=2 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=32 --num_heads=2 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=32 --num_heads=1 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=32 --num_heads=1 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=8 --num_heads=32 --head_size_vo=96 --head_size_qk=96 -PvcFMHABF16BF16FP32_RCR_h64_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=8 --num_heads=32 --head_size_vo=96 --head_size_qk=96 - -PvcFMHABF16BF16FP32_RCR_h128_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads=16 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads=16 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads=16 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads=16 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=16 --num_heads=8 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=16 --num_heads=8 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=16 --num_heads=4 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=16 --num_heads=4 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=16 --num_heads=2 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=16 --num_heads=2 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads=32 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads=32 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads=96 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads=96 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=16 --num_heads=1 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=16 --num_heads=1 --head_size_qk=128 --head_size_vo=128 - -PvcFMHABF16BF16FP32_RCR_h192_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads=128 --head_size_vo=192 --head_size_qk=192 -PvcFMHABF16BF16FP32_RCR_h192_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads=128 --head_size_vo=192 --head_size_qk=192 - -PvcFMHABF16BF16FP32_RCR_h64_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads=32 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads=32 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=4 --num_heads=48 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=4 --num_heads=48 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=32 --num_heads=16 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=32 --num_heads=16 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=32 --num_heads=8 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=32 --num_heads=8 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=32 --num_heads=4 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=32 --num_heads=4 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=32 --num_heads=2 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=32 --num_heads=2 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=32 --num_heads=1 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=32 --num_heads=1 --head_size_qk=64 --head_size_vo=64 -PvcFMHABF16BF16FP32_RCR_h64_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=8 --num_heads=32 --head_size_vo=96 --head_size_qk=96 -PvcFMHABF16BF16FP32_RCR_h64_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=8 --num_heads=32 --head_size_vo=96 --head_size_qk=96 - -PvcFMHABF16BF16FP32_RCR_h128_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads=16 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads=16 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads=16 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads=16 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=16 --num_heads=8 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=16 --num_heads=8 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=16 --num_heads=4 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=16 --num_heads=4 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=16 --num_heads=2 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=16 --num_heads=2 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads=32 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads=32 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads=96 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads=96 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=16 --num_heads=1 --head_size_qk=128 --head_size_vo=128 -PvcFMHABF16BF16FP32_RCR_h128_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=16 --num_heads=1 --head_size_qk=128 --head_size_vo=128 - -PvcFMHABF16BF16FP32_RCR_h192_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads=128 --head_size_vo=192 --head_size_qk=192 -PvcFMHABF16BF16FP32_RCR_h192_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads=128 --head_size_vo=192 --head_size_qk=192 - -# FMHA FP16 benchmarks -PvcFMHAFP16FP16FP32_RCR_h64_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads=32 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads=32 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=4 --num_heads=48 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=4 --num_heads=48 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=32 --num_heads=16 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=32 --num_heads=16 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=32 --num_heads=8 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=32 --num_heads=8 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=32 --num_heads=4 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=32 --num_heads=4 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=32 --num_heads=2 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=32 --num_heads=2 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=32 --num_heads=1 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=32 --num_heads=1 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=8 --num_heads=32 --head_size_vo=96 --head_size_qk=96 -PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=8 --num_heads=32 --head_size_vo=96 --head_size_qk=96 - -PvcFMHAFP16FP16FP32_RCR_h128_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads=16 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads=16 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads=16 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads=16 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=16 --num_heads=8 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=16 --num_heads=8 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=16 --num_heads=4 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=16 --num_heads=4 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=16 --num_heads=2 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=16 --num_heads=2 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads=32 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads=32 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads=96 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads=96 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=16 --num_heads=1 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=16 --num_heads=1 --head_size_qk=128 --head_size_vo=128 - -PvcFMHAFP16FP16FP32_RCR_h192_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads=128 --head_size_vo=192 --head_size_qk=192 -PvcFMHAFP16FP16FP32_RCR_h192_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads=128 --head_size_vo=192 --head_size_qk=192 - -PvcFMHAFP16FP16FP32_RCR_h64_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads=32 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads=32 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=4 --num_heads=48 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=4 --num_heads=48 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=32 --num_heads=16 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=32 --num_heads=16 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=32 --num_heads=8 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=32 --num_heads=8 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=32 --num_heads=4 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=32 --num_heads=4 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=32 --num_heads=2 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=32 --num_heads=2 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=32 --num_heads=1 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=32 --num_heads=1 --head_size_qk=64 --head_size_vo=64 -PvcFMHAFP16FP16FP32_RCR_h64_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=8 --num_heads=32 --head_size_vo=96 --head_size_qk=96 -PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=8 --num_heads=32 --head_size_vo=96 --head_size_qk=96 - -PvcFMHAFP16FP16FP32_RCR_h128_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads=16 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads=16 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads=16 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads=16 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=16 --num_heads=8 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=16 --num_heads=8 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=16 --num_heads=4 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=16 --num_heads=4 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=16 --num_heads=2 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=16 --num_heads=2 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads=32 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads=32 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads=96 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads=96 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=16 --num_heads=1 --head_size_qk=128 --head_size_vo=128 -PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=16 --num_heads=1 --head_size_qk=128 --head_size_vo=128 - -PvcFMHAFP16FP16FP32_RCR_h192_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads=128 --head_size_vo=192 --head_size_qk=192 -PvcFMHAFP16FP16FP32_RCR_h192_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads=128 --head_size_vo=192 --head_size_qk=192 +# working correctly +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=4096 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=512 --n=16 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1 --k=2 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1 --k=768 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2 --k=128 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2 --k=512 --n=1536 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2 --k=512 --n=4096 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2 --k=1024 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2 --k=512 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2 --k=512 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2 --k=128 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2 --k=512 --n=256 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=9 --k=64 --n=3072 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=28 --k=2560 --n=2560 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=28 --k=1024 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=32 --k=128 --n=16 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=54 --k=512 --n=384 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=512 --n=16 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=1024 --n=8 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=512 --n=64 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=512 --n=24 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=256 --n=16 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=1024 --n=16 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=128 --n=16 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=80 --k=128 --n=32 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=128 --n=32 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=256 --n=256 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=2560 --n=2560 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=256008 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=2048 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=1024 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=384 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=250112 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=250112 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=1024 --n=4096 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=2048 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=2560 --n=10240 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=32 --n=16 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=2560 --n=8008 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=256 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=768 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=384 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=80 --n=32 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=8008 --n=2560 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=2048 --n=256 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=256 --n=10000 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=4096 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=3072 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=1024 --n=256008 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=768 --n=3072 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=128 --n=16 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=128 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=1024 --n=128112 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=128112 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=10240 --n=2560 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=64 --n=16 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=128 --n=128 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=1024 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=10000 --n=256 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=128 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=256 --n=2048 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=4096 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=64 --n=16 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=4096 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=1024 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=512 --n=256 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=256 --n=16 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=512 --n=128 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=64 --n=48 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=128 --n=2048 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=1024 --n=4096 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=512 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=128 --n=256 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=768 --n=48 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=384 --k=128 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=384 --k=512 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=32000 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1024 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=512 --n=64 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=64 --n=16 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=16384 --n=4096 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=128 --n=384 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=128 --n=30000 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=128 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=256 --n=256 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1536 --n=6144 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1024 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=512 --n=24 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1024 --n=29056 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=128 --n=128 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=768 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=4096 --n=4096 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=4096 --n=16384 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=30000 --n=128 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=32000 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1024 --n=256 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=768 --n=2304 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=384 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=4096 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=128 --n=4096 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=128 --n=2048 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=4096 --n=128 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=29056 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=256 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=3072 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=768 --n=384 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=256 --n=128 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1024 --n=32000 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1024 --n=16 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1024 --n=4096 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=64 --n=64 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1536 --n=1536 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=768 --n=32000 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=2304 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=6144 --n=1536 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=128 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=128 --n=256 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=768 --n=3072 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=64 --n=24 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1024 --n=2048 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=512 --n=3072 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=1024 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=2048 --n=3072 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=512 --n=2304 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=256 --n=48 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=1024 --n=2304 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=128 --n=3072 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=512 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=1024 --n=3072 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=2048 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=128 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=2304 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=512 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=256 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=512 --n=2048 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=128 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=512 --n=256 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=768 --n=3072 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=1024 --n=4096 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=1024 --n=16 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=128 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=32128 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=3072 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=256 --n=4096 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=512 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=1024 --n=8 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=128 --n=4096 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=768 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=2048 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=768 --n=2304 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=64 --n=16 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=512 --n=4096 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=1024 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=512 --n=32128 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=64 --n=8 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1536 --k=512 --n=1536 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1536 --k=512 --n=6144 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2048 --k=768 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2048 --k=768 --n=50272 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2048 --k=3072 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2048 --k=1024 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2048 --k=768 --n=3072 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2048 --k=128 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2048 --k=50272 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2048 --k=128 --n=256 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2304 --k=512 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2560 --k=128 --n=2560 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2560 --k=128 --n=10240 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=3072 --k=2048 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=3072 --k=128 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=3072 --k=512 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=3072 --k=1024 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=512 --n=128 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=1024 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=512 --n=16384 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=256 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=512 --n=4096 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=128 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=512 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=6144 --k=512 --n=1536 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=8008 --k=128 --n=2560 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=10000 --k=128 --n=256 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=10240 --k=128 --n=2560 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=512 --n=4096 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=29056 --k=512 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=30000 --k=512 --n=128 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=30522 --k=512 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=30522 --k=512 --n=128 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=30522 --k=128 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=32000 --k=512 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=32000 --k=512 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=32005 --k=512 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=32128 --k=1024 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=50005 --k=1024 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=50257 --k=512 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=50265 --k=512 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=50265 --k=128 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=50265 --k=256 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=50265 --k=128 --n=512 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=50265 --k=1024 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=50265 --k=1024 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=50272 --k=2048 --n=768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128112 --k=128 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=250112 --k=128 --n=512 From ba5bec4e300633e1b8c8dac12104b3e7593a75d7 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Tadej=20Ciglari=C4=8D?= Date: Tue, 15 Apr 2025 10:32:32 +0100 Subject: [PATCH 2/8] hacky padding --- examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp | 60 ++++++++++++++++------- include/cutlass/gemm/kernel/xe_gemm.hpp | 2 +- tools/copy_debug/copy_debug.cpp | 8 +-- 3 files changed, 46 insertions(+), 24 deletions(-) diff --git a/examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp b/examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp index 62287fb275..dc5de2bcdc 100644 --- a/examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp +++ b/examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp @@ -156,11 +156,18 @@ struct ExampleRunner { bool verify(const ProblemShapeType& problem_size, ElementCompute alpha, ElementCompute beta) { auto [M, N, K, L] = problem_size; - - cutlass::TensorRef ref_A(block_A.get(), LayoutA::packed({M, K})); - cutlass::TensorRef ref_B(block_B.get(), LayoutB::packed({K, N})); - cutlass::TensorRef ref_C(block_C.get(), LayoutC::packed({M, N})); - cutlass::TensorRef ref_D(block_ref_D.get(), LayoutD::packed({M, N})); + + int P = 8; + //int M2 = (M+P-1)/P*P; + int N2 = (N+P-1)/P*P; + int K2 = (K+P-1)/P*P; + + //cutlass::TensorRef ref_A(block_A.get(), LayoutA::packed({M, K})); + //cutlass::TensorRef ref_B(block_B.get(), LayoutB::packed({K, N})); + cutlass::TensorRef ref_A(block_A.get(), LayoutA(K2)); + cutlass::TensorRef ref_B(block_B.get(), LayoutB(N2)); + cutlass::TensorRef ref_C(block_C.get(), LayoutC(N2)); + cutlass::TensorRef ref_D(block_ref_D.get(), LayoutD(N2)); cutlass::reference::device::GemmComplex( {M, N, K}, @@ -174,10 +181,10 @@ struct ExampleRunner { ref_D, ElementAccumulator(0), L, // batch_count - M * K, // batch_stride_A - K * N, // batch_stride_B - M * N, // batch_stride_C - M * N // batch_stride_D + M * K2, // batch_stride_A + K * N2, // batch_stride_B + M * N2, // batch_stride_C + M * N2 // batch_stride_D ); syclcompat::wait(); @@ -194,16 +201,31 @@ struct ExampleRunner { auto problem_shape_MNKL = cute::append<4>(problem_size, 1); auto [M, N, K, L] = problem_shape_MNKL; - stride_A = cutlass::make_cute_packed_stride(StrideA{}, cute::make_shape(M, K, L)); - stride_B = cutlass::make_cute_packed_stride(StrideB{}, cute::make_shape(N, K, L)); - stride_C = cutlass::make_cute_packed_stride(StrideC{}, cute::make_shape(M, N, L)); - stride_D = cutlass::make_cute_packed_stride(StrideD{}, cute::make_shape(M, N, L)); - - block_A.reset(M * K * L); - block_B.reset(K * N * L); - block_C.reset(M * N * L); - block_D.reset(M * N * L); - block_ref_D.reset(M * N * L); + /*stride_A = cutlass::make_cute_padded_stride(StrideA{}, cute::make_shape(M, K, L), 16); + stride_B = cutlass::make_cute_padded_stride(StrideB{}, cute::make_shape(N, K, L), 16); + stride_C = cutlass::make_cute_padded_stride(StrideC{}, cute::make_shape(M, N, L), 16); + stride_D = cutlass::make_cute_padded_stride(StrideD{}, cute::make_shape(M, N, L), 16);*/ + + int P = 8; + //int M2 = (M+P-1)/P*P; + int N2 = (N+P-1)/P*P; + int K2 = (K+P-1)/P*P; + + stride_A = cutlass::make_cute_packed_stride(StrideA{}, cute::make_shape(M, K2, L)); + stride_B = cutlass::make_cute_packed_stride(StrideB{}, cute::make_shape(N2, K, L)); + stride_C = cutlass::make_cute_packed_stride(StrideC{}, cute::make_shape(M, N2, L)); + stride_D = cutlass::make_cute_packed_stride(StrideD{}, cute::make_shape(M, N2, L)); + + print("stride_A "); print(stride_A); print("\n"); + print("stride_B "); print(stride_B); print("\n"); + print("stride_C "); print(stride_C); print("\n"); + print("stride_D "); print(stride_D); print("\n"); + + block_A.reset(M * K2 * L); + block_B.reset(K * N2 * L); + block_C.reset(M * N2 * L); + block_D.reset(M * N2 * L); + block_ref_D.reset(M * N2 * L); initialize_block(block_A, seed + 2023); initialize_block(block_B, seed + 2022); diff --git a/include/cutlass/gemm/kernel/xe_gemm.hpp b/include/cutlass/gemm/kernel/xe_gemm.hpp index 3f7b964bbd..4877f72190 100644 --- a/include/cutlass/gemm/kernel/xe_gemm.hpp +++ b/include/cutlass/gemm/kernel/xe_gemm.hpp @@ -166,7 +166,7 @@ class GemmUniversal< bool m_valid = m > 0; bool n_valid = n > 0 && n % 4 == 0; bool k_valid = k > 0 && k % get<2>(TileShape{}) == 0; - bool shape_implementable = (m_valid && n_valid && k_valid); + bool shape_implementable = true || (m_valid && n_valid && k_valid); bool mode_implementable = args.mode == GemmUniversalMode::kGemm || (args.mode == GemmUniversalMode::kBatched && rank(ProblemShape{}) == 4); diff --git a/tools/copy_debug/copy_debug.cpp b/tools/copy_debug/copy_debug.cpp index a62275614a..e862dbab46 100644 --- a/tools/copy_debug/copy_debug.cpp +++ b/tools/copy_debug/copy_debug.cpp @@ -113,9 +113,9 @@ void copy(int global_M, int global_N) { auto tensor_shape = make_shape(global_M, global_N, 1); int tensor_size = size(tensor_shape); - cutlass::DeviceAllocation src(tensor_size); + cutlass::DeviceAllocation src(tensor_size*32); - Tensor tensor_S = make_tensor(make_gmem_ptr(src.get()), make_layout(tensor_shape, LayoutLeft{})); + Tensor tensor_S = make_tensor(make_gmem_ptr(src.get()), make_layout(tensor_shape, make_stride(_1{},32,_0{}))); auto gridDim = syclcompat::dim3(1); auto blockDim = syclcompat::dim3(SUBGROUP_SIZE); @@ -130,7 +130,7 @@ void copy(int global_M, int global_N) { int main(){ // for 16b copies use integers as floating point types could lose precision for bigger indices // for 8b copies you have to work with overflow - copy(256, 256); - copy(256, 256); + copy(3, 5); + copy(3, 5); return 0; } From 5398c8763e0955e57e4fb7dcc706de0e8c72ad6d Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Tadej=20Ciglari=C4=8D?= Date: Tue, 22 Apr 2025 15:03:08 +0100 Subject: [PATCH 3/8] cleanup alignment and checks --- examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp | 89 +++++++++++++---------- include/cutlass/gemm/kernel/xe_gemm.hpp | 23 +++++- 2 files changed, 69 insertions(+), 43 deletions(-) diff --git a/examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp b/examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp index dc5de2bcdc..9b963191dd 100644 --- a/examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp +++ b/examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp @@ -52,6 +52,11 @@ using namespace cute; /////////////////////////////////////////////////////////////////////////////////////////////////// +// The alignment requirement in bytes on inner dimmension that will work for both PVC and BMG +constexpr int AlignmentInner = 16; +// The alignment requirement in bytes on outer dimmension that will work for both PVC and BMG +constexpr int AlignmentPtr = 64; + // Command line options parsing struct Options { @@ -127,12 +132,18 @@ struct ExampleRunner { using CollectiveEpilogue = typename Gemm::CollectiveEpilogue; using ElementC = typename Gemm::ElementC; + using ElementD = typename Gemm::ElementD; using ElementOutput = typename CollectiveEpilogue::ElementOutput; using ElementCompute = typename CollectiveEpilogue::ElementCompute; using ElementAccumulator = typename CollectiveEpilogue::ElementAccumulator; using ProblemShapeType = typename Gemm::GemmKernel::ProblemShape; + static constexpr int AlignElemA = AlignmentInner / sizeof(ElementA); + static constexpr int AlignElemB = AlignmentInner / sizeof(ElementB); + static constexpr int AlignElemC = AlignmentInner / sizeof(ElementB); + static constexpr int AlignElemD = AlignmentInner / sizeof(ElementD); + // // Data members // @@ -157,17 +168,20 @@ struct ExampleRunner { bool verify(const ProblemShapeType& problem_size, ElementCompute alpha, ElementCompute beta) { auto [M, N, K, L] = problem_size; - int P = 8; - //int M2 = (M+P-1)/P*P; - int N2 = (N+P-1)/P*P; - int K2 = (K+P-1)/P*P; - - //cutlass::TensorRef ref_A(block_A.get(), LayoutA::packed({M, K})); - //cutlass::TensorRef ref_B(block_B.get(), LayoutB::packed({K, N})); - cutlass::TensorRef ref_A(block_A.get(), LayoutA(K2)); - cutlass::TensorRef ref_B(block_B.get(), LayoutB(N2)); - cutlass::TensorRef ref_C(block_C.get(), LayoutC(N2)); - cutlass::TensorRef ref_D(block_ref_D.get(), LayoutD(N2)); + // Padded values + int N_B = cute::ceil_div(N, AlignElemB) * AlignElemB; + int N_C = cute::ceil_div(N, AlignElemC) * AlignElemC; + int N_D = cute::ceil_div(N, AlignElemD) * AlignElemD; + int K_A = cute::ceil_div(K, AlignElemA) * AlignElemA; + + int AlignmentOuter = AlignmentPtr / AlignmentInner; + int M_ACD = cute::ceil_div(M, AlignmentOuter) * AlignmentOuter; + int K_B = cute::ceil_div(K, AlignmentOuter) * AlignmentOuter; + + cutlass::TensorRef ref_A(block_A.get(), LayoutA(K_A)); + cutlass::TensorRef ref_B(block_B.get(), LayoutB(N_B)); + cutlass::TensorRef ref_C(block_C.get(), LayoutC(N_C)); + cutlass::TensorRef ref_D(block_ref_D.get(), LayoutD(N_D)); cutlass::reference::device::GemmComplex( {M, N, K}, @@ -181,10 +195,10 @@ struct ExampleRunner { ref_D, ElementAccumulator(0), L, // batch_count - M * K2, // batch_stride_A - K * N2, // batch_stride_B - M * N2, // batch_stride_C - M * N2 // batch_stride_D + M_ACD * K_A, // batch_stride_A + K_B * N_B, // batch_stride_B + M_ACD * N_C, // batch_stride_C + M_ACD * N_D // batch_stride_D ); syclcompat::wait(); @@ -201,31 +215,26 @@ struct ExampleRunner { auto problem_shape_MNKL = cute::append<4>(problem_size, 1); auto [M, N, K, L] = problem_shape_MNKL; - /*stride_A = cutlass::make_cute_padded_stride(StrideA{}, cute::make_shape(M, K, L), 16); - stride_B = cutlass::make_cute_padded_stride(StrideB{}, cute::make_shape(N, K, L), 16); - stride_C = cutlass::make_cute_padded_stride(StrideC{}, cute::make_shape(M, N, L), 16); - stride_D = cutlass::make_cute_padded_stride(StrideD{}, cute::make_shape(M, N, L), 16);*/ - - int P = 8; - //int M2 = (M+P-1)/P*P; - int N2 = (N+P-1)/P*P; - int K2 = (K+P-1)/P*P; - - stride_A = cutlass::make_cute_packed_stride(StrideA{}, cute::make_shape(M, K2, L)); - stride_B = cutlass::make_cute_packed_stride(StrideB{}, cute::make_shape(N2, K, L)); - stride_C = cutlass::make_cute_packed_stride(StrideC{}, cute::make_shape(M, N2, L)); - stride_D = cutlass::make_cute_packed_stride(StrideD{}, cute::make_shape(M, N2, L)); - - print("stride_A "); print(stride_A); print("\n"); - print("stride_B "); print(stride_B); print("\n"); - print("stride_C "); print(stride_C); print("\n"); - print("stride_D "); print(stride_D); print("\n"); - - block_A.reset(M * K2 * L); - block_B.reset(K * N2 * L); - block_C.reset(M * N2 * L); - block_D.reset(M * N2 * L); - block_ref_D.reset(M * N2 * L); + // Padded values + int N_B = cute::ceil_div(N, AlignElemB) * AlignElemB; + int N_C = cute::ceil_div(N, AlignElemC) * AlignElemC; + int N_D = cute::ceil_div(N, AlignElemD) * AlignElemD; + int K_A = cute::ceil_div(K, AlignElemA) * AlignElemA; + + int AlignmentOuter = AlignmentPtr / AlignmentInner; + int M_ACD = cute::ceil_div(M, AlignmentOuter) * AlignmentOuter; + int K_B = cute::ceil_div(K, AlignmentOuter) * AlignmentOuter; + + stride_A = cutlass::make_cute_packed_stride(StrideA{}, cute::make_shape(M_ACD, K_A, L)); + stride_B = cutlass::make_cute_packed_stride(StrideB{}, cute::make_shape(N_B, K_B, L)); + stride_C = cutlass::make_cute_packed_stride(StrideC{}, cute::make_shape(M_ACD, N_C, L)); + stride_D = cutlass::make_cute_packed_stride(StrideD{}, cute::make_shape(M_ACD, N_D, L)); + + block_A.reset(M_ACD * K_A * L); + block_B.reset(K * N_B * L); + block_C.reset(M_ACD * N_C * L); + block_D.reset(M_ACD * N_D * L); + block_ref_D.reset(M_ACD * N_D * L); initialize_block(block_A, seed + 2023); initialize_block(block_B, seed + 2022); diff --git a/include/cutlass/gemm/kernel/xe_gemm.hpp b/include/cutlass/gemm/kernel/xe_gemm.hpp index 4877f72190..a91c2b1a9a 100644 --- a/include/cutlass/gemm/kernel/xe_gemm.hpp +++ b/include/cutlass/gemm/kernel/xe_gemm.hpp @@ -162,11 +162,28 @@ class GemmUniversal< auto m = get<0>(args.problem_shape); auto n = get<1>(args.problem_shape); auto k = get<2>(args.problem_shape); + auto l = get<3>(args.problem_shape); + bool is_batch = l > 1; + auto check_stride = [is_batch](auto stride, int el_size){ + auto a = get<0>(stride); + auto b = get<1>(stride); + auto valid_is_unit = a == _1{} || b == _1{}; + auto inner = a == _1{} ? b : a; + auto valid_inner = inner % (16 / el_size) == 0; + auto valid_outer = !is_batch || get<2>(stride) % (64 / el_size) == 0; + return valid_is_unit && valid_inner && valid_outer; + }; + bool strides_valid = check_stride(args.mainloop.dA, sizeof(ElementA)) && + check_stride(args.mainloop.dB, sizeof(ElementB)) && + check_stride(args.epilogue.dC, sizeof(ElementC)) && + check_stride(args.epilogue.dD, sizeof(ElementD)); // TODO(codeplay): base *_valid on the atom shapes bool m_valid = m > 0; - bool n_valid = n > 0 && n % 4 == 0; - bool k_valid = k > 0 && k % get<2>(TileShape{}) == 0; - bool shape_implementable = true || (m_valid && n_valid && k_valid); + bool n_valid = n > 0 && n % (4 / sizeof(ElementB)) == 0 && + n % (4 / sizeof(ElementC)) == 0 && + n % (4 / sizeof(ElementD)) == 0; + bool k_valid = k > 0 && k % (4 / sizeof(ElementA)) == 0; + bool shape_implementable = m_valid && n_valid && k_valid && strides_valid; bool mode_implementable = args.mode == GemmUniversalMode::kGemm || (args.mode == GemmUniversalMode::kBatched && rank(ProblemShape{}) == 4); From 2cd39018561b619b204726978d033a804c0ece8d Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Tadej=20Ciglari=C4=8D?= Date: Tue, 22 Apr 2025 15:03:46 +0100 Subject: [PATCH 4/8] =?UTF-8?q?revert=20copy=5Fdebug=1B[H?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- tools/copy_debug/copy_debug.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/tools/copy_debug/copy_debug.cpp b/tools/copy_debug/copy_debug.cpp index e862dbab46..a62275614a 100644 --- a/tools/copy_debug/copy_debug.cpp +++ b/tools/copy_debug/copy_debug.cpp @@ -113,9 +113,9 @@ void copy(int global_M, int global_N) { auto tensor_shape = make_shape(global_M, global_N, 1); int tensor_size = size(tensor_shape); - cutlass::DeviceAllocation src(tensor_size*32); + cutlass::DeviceAllocation src(tensor_size); - Tensor tensor_S = make_tensor(make_gmem_ptr(src.get()), make_layout(tensor_shape, make_stride(_1{},32,_0{}))); + Tensor tensor_S = make_tensor(make_gmem_ptr(src.get()), make_layout(tensor_shape, LayoutLeft{})); auto gridDim = syclcompat::dim3(1); auto blockDim = syclcompat::dim3(SUBGROUP_SIZE); @@ -130,7 +130,7 @@ void copy(int global_M, int global_N) { int main(){ // for 16b copies use integers as floating point types could lose precision for bigger indices // for 8b copies you have to work with overflow - copy(3, 5); - copy(3, 5); + copy(256, 256); + copy(256, 256); return 0; } From 9663d2c2c837ce7067e4277a8f4a5d450f4ec0ce Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Tadej=20Ciglari=C4=8D?= Date: Tue, 22 Apr 2025 15:06:51 +0100 Subject: [PATCH 5/8] revert input.in --- benchmarks/pvc/input.in | 515 ++++++++++++++++++---------------------- 1 file changed, 226 insertions(+), 289 deletions(-) diff --git a/benchmarks/pvc/input.in b/benchmarks/pvc/input.in index c366a4e1e7..e48fd57e24 100644 --- a/benchmarks/pvc/input.in +++ b/benchmarks/pvc/input.in @@ -1,290 +1,227 @@ -# validation error -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1 --k=768 --n=2 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2 --k=1 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=32 --k=128 --n=4 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=128 --n=6 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=512 --n=12 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=128 --n=4 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=512 --n=6 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=9 --n=3072 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=512 --n=4 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=128 --n=12 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=1 --n=3072 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=512 --n=36 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=2048 --n=12 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=1024 --n=12 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=64 --n=6 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=30522 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=2 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=2 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=30522 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=64 --n=12 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=30522 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=64 --n=4 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=128 --n=12 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=50265 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=32 --n=4 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=50265 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=128 --n=4 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=50265 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=768 --n=30522 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=1024 --n=50265 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=2 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=768 --n=2 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=128 --n=6 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=50265 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=1024 --n=50265 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=384 --k=54 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=64 --n=4 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=128 --n=30522 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1536 --n=128100 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=30522 --n=128 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=30522 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=64 --n=36 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=256 --n=2 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=768 --n=50265 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=768 --n=30522 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=512 --n=12 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=384 --n=54 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=512 --n=36 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=2 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=32005 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=128100 --n=1536 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=64 --n=12 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=512 --n=6 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=50257 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1536 --n=2 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=64 --n=6 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=512 --n=4 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=2 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=768 --n=2 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=2 --n=256 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=768 --n=32005 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=2 --n=1536 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=2 --n=4096 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1024 --n=2 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=1 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=1024 --n=12 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=768 --n=2 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=50005 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=50265 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=50265 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=4096 --n=2 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=2 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=50265 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=1024 --n=50265 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=768 --n=50265 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=64 --n=12 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2048 --k=2048 --n=12 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2048 --k=64 --n=12 -# crash -#PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=768 --n=50257 -#PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=768 --n=50005 +# GEMM BFloat16 benchmarks +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=4096 --n=4096 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=8192 --k=8192 --n=8192 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1 --k=5120 --n=13824 +PvcGemmBF16BF16FP32_RRR_2 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=28672 --n=8192 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=3072 --k=4096 --n=3072 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4 --k=4096 --n=12288 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=8192 --n=8192 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=8192 --n=32768 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=32768 --n=8192 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=8192 --n=1024 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=1024 --n=8192 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=8192 --n=4096 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=4096 --n=8192 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=16384 --n=8192 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=8192 --k=16384 --n=4096 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=16384 --n=8192 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=8192 --k=16384 --n=1024 +# TODO(Codeplay): these are broken +# PvcGemmBF16BF16FP32_RRR_4 --bm_name=bf16_bf16_fp32 --l=4096 --m=8 --k=128 --n=16384 +# PvcGemmBF16BF16FP32_RRR_5 --bm_name=bf16_bf16_fp32 --l=4096 --m=8 --k=16384 --n=128 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=4 --m=32768 --k=128 --n=4096 +PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=4 --m=32768 --k=4096 --n=128 +PvcGemmBF16BF16FP32_RRR_3 --bm_name=bf16_bf16_fp32 --l=32 --m=4096 --k=4096 --n=128 +PvcGemmBF16BF16FP32_RCR_6 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=4096 --n=4096 +PvcGemmBF16BF16FP32_CRR_7 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=4096 --n=4096 +PvcGemmBF16BF16FP32_CCR_8 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=4096 --n=4096 -# working correctly -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=4096 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=512 --n=16 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1 --k=2 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1 --k=768 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2 --k=128 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2 --k=512 --n=1536 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2 --k=512 --n=4096 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2 --k=1024 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2 --k=512 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2 --k=512 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2 --k=128 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2 --k=512 --n=256 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=9 --k=64 --n=3072 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=28 --k=2560 --n=2560 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=28 --k=1024 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=32 --k=128 --n=16 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=54 --k=512 --n=384 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=512 --n=16 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=1024 --n=8 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=512 --n=64 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=512 --n=24 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=256 --n=16 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=1024 --n=16 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=64 --k=128 --n=16 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=80 --k=128 --n=32 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=128 --n=32 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=256 --n=256 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=2560 --n=2560 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=256008 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=2048 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=1024 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=384 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=250112 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=250112 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=1024 --n=4096 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=2048 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=2560 --n=10240 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=32 --n=16 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=2560 --n=8008 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=256 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=768 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=384 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=80 --n=32 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=8008 --n=2560 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=2048 --n=256 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=256 --n=10000 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=4096 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=3072 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=1024 --n=256008 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=768 --n=3072 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=128 --n=16 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=128 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=1024 --n=128112 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=128112 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=10240 --n=2560 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=64 --n=16 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=128 --n=128 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=1024 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=10000 --n=256 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=128 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=256 --n=2048 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128 --k=512 --n=4096 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=64 --n=16 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=4096 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=1024 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=512 --n=256 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=256 --n=16 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=512 --n=128 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=64 --n=48 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=128 --n=2048 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=1024 --n=4096 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=512 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=128 --n=256 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=256 --k=768 --n=48 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=384 --k=128 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=384 --k=512 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=32000 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1024 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=512 --n=64 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=64 --n=16 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=16384 --n=4096 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=128 --n=384 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=128 --n=30000 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=128 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=256 --n=256 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1536 --n=6144 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1024 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=512 --n=24 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1024 --n=29056 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=128 --n=128 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=768 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=4096 --n=4096 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=4096 --n=16384 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=30000 --n=128 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=32000 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1024 --n=256 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=768 --n=2304 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=384 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=4096 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=128 --n=4096 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=128 --n=2048 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=4096 --n=128 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=29056 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=256 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=3072 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=768 --n=384 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=256 --n=128 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1024 --n=32000 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1024 --n=16 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1024 --n=4096 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=64 --n=64 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1536 --n=1536 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=768 --n=32000 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=2304 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=6144 --n=1536 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=128 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=128 --n=256 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=768 --n=3072 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=64 --n=24 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=1024 --n=2048 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=512 --n=3072 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=1024 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=2048 --n=3072 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=512 --n=2304 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=256 --n=48 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=1024 --n=2304 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=128 --n=3072 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=512 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=1024 --n=3072 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=2048 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=768 --k=128 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=2304 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=512 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=256 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=512 --n=2048 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=128 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=512 --n=256 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=768 --n=3072 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=1024 --n=4096 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=1024 --n=16 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=128 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=32128 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=3072 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=256 --n=4096 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=512 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=1024 --n=8 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=128 --n=4096 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=768 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=2048 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=768 --n=2304 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=64 --n=16 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=512 --n=4096 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=1024 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=512 --n=32128 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=64 --n=8 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1536 --k=512 --n=1536 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1536 --k=512 --n=6144 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2048 --k=768 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2048 --k=768 --n=50272 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2048 --k=3072 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2048 --k=1024 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2048 --k=768 --n=3072 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2048 --k=128 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2048 --k=50272 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2048 --k=128 --n=256 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2304 --k=512 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2560 --k=128 --n=2560 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=2560 --k=128 --n=10240 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=3072 --k=2048 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=3072 --k=128 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=3072 --k=512 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=3072 --k=1024 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=512 --n=128 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=1024 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=512 --n=16384 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=256 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=512 --n=4096 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=128 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=512 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=6144 --k=512 --n=1536 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=8008 --k=128 --n=2560 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=10000 --k=128 --n=256 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=10240 --k=128 --n=2560 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=512 --n=4096 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=29056 --k=512 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=30000 --k=512 --n=128 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=30522 --k=512 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=30522 --k=512 --n=128 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=30522 --k=128 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=32000 --k=512 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=32000 --k=512 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=32005 --k=512 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=32128 --k=1024 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=50005 --k=1024 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=50257 --k=512 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=50265 --k=512 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=50265 --k=128 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=50265 --k=256 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=50265 --k=128 --n=512 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=50265 --k=1024 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=50265 --k=1024 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=50272 --k=2048 --n=768 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=128112 --k=128 --n=1024 -PvcGemmBF16BF16FP32_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=250112 --k=128 --n=512 +PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=8192 --n=8192 +PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=8192 --n=32768 +PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=32768 --n=8192 +PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=8192 --n=1024 +PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=1024 --n=8192 +PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=8192 --n=4096 +PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=4096 --n=8192 +PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=16384 --n=8192 +PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=8192 --k=16384 --n=4096 +PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=16384 --n=8192 +PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=8192 --k=16384 --n=1024 +# PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=4096 --m=8 --k=128 --n=16384 +# PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=4096 --m=8 --k=16384 --n=128 +PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=4 --m=32768 --k=128 --n=4096 +PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=4 --m=32768 --k=4096 --n=128 +PvcGemmBF16BF16FP32_StreamK_RRR_1 --bm_name=bf16_bf16_fp32 --l=32 --m=4096 --k=4096 --n=128 + +PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=8192 --n=8192 +PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=8192 --n=32768 +PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=512 --k=32768 --n=8192 +PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=8192 --n=1024 +PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=1024 --n=8192 +PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=8192 --n=4096 +PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=16384 --k=4096 --n=8192 +PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=4096 --k=16384 --n=8192 +PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=8192 --k=16384 --n=4096 +PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=1024 --k=16384 --n=8192 +PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=1 --m=8192 --k=16384 --n=1024 +# PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=4096 --m=8 --k=128 --n=16384 +# PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=4096 --m=8 --k=16384 --n=128 +PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=4 --m=32768 --k=128 --n=4096 +PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=4 --m=32768 --k=4096 --n=128 +PvcGemmBF16BF16FP32_SplitK_RRR_1 --bm_name=bf16_bf16_fp32 --l=32 --m=4096 --k=4096 --n=128 + +# FMHA BFloat16 benchmarks +PvcFMHABF16BF16FP32_RCR_h64_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads_q=32 --num_heads_kv=32 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads_q=32 --num_heads_kv=32 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=4 --num_heads_q=48 --num_heads_kv=48 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=4 --num_heads_q=48 --num_heads_kv=48 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=32 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=32 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=32 --num_heads_q=8 --num_heads_kv=8 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=32 --num_heads_q=8 --num_heads_kv=8 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=32 --num_heads_q=4 --num_heads_kv=4 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=32 --num_heads_q=4 --num_heads_kv=4 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=32 --num_heads_q=2 --num_heads_kv=2 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=32 --num_heads_q=2 --num_heads_kv=2 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=32 --num_heads_q=1 --num_heads_kv=1 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=32 --num_heads_q=1 --num_heads_kv=1 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=8 --num_heads_q=32 --num_heads_kv=32 --head_size_vo=96 --head_size_qk=96 +PvcFMHABF16BF16FP32_RCR_h64_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=8 --num_heads_q=32 --num_heads_kv=32 --head_size_vo=96 --head_size_qk=96 + +PvcFMHABF16BF16FP32_RCR_h128_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=16 --num_heads_q=8 --num_heads_kv=8 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=16 --num_heads_q=8 --num_heads_kv=8 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=16 --num_heads_q=4 --num_heads_kv=4 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=16 --num_heads_q=4 --num_heads_kv=4 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=16 --num_heads_q=2 --num_heads_kv=2 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=16 --num_heads_q=2 --num_heads_kv=2 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads_q=32 --num_heads_kv=32 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads_q=32 --num_heads_kv=32 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads_q=96--num_heads_kv=96 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads_q=96--num_heads_kv=96 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=16 --num_heads_q=1 --num_heads_kv=1 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=16 --num_heads_q=1 --num_heads_kv=1 --head_size_qk=128 --head_size_vo=128 + +PvcFMHABF16BF16FP32_RCR_h128_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=2048 --batch=1 --num_heads_q=32, --num_heads_kv=8 --head_size_vo=128 --head_size_qk=128 +PvcFMHABF16BF16FP32_RCR_h128_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=2048 --batch=1 --num_heads_q=32, --num_heads_kv=8 --head_size_vo=128 --head_size_qk=128 + +PvcFMHABF16BF16FP32_RCR_h192_Causal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads_q=128 --num_heads_kv=128 --head_size_vo=192 --head_size_qk=192 +PvcFMHABF16BF16FP32_RCR_h192_NonCausal_FixedLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads_q=128 --num_heads_kv=128 --head_size_vo=192 --head_size_qk=192 + +PvcFMHABF16BF16FP32_RCR_h64_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads_q=32 --num_heads_kv=32 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads_q=32 --num_heads_kv=32 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=4 --num_heads_q=48 --num_heads_kv=48 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=4 --num_heads_q=48 --num_heads_kv=48 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=32 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=32 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=32 --num_heads_q=8 --num_heads_kv=8 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=32 --num_heads_q=8 --num_heads_kv=8 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=32 --num_heads_q=4 --num_heads_kv=4 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=32 --num_heads_q=4 --num_heads_kv=4 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=32 --num_heads_q=2 --num_heads_kv=2 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=32 --num_heads_q=2 --num_heads_kv=2 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=32 --num_heads_q=1 --num_heads_kv=1 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=32 --num_heads_q=1 --num_heads_kv=1 --head_size_qk=64 --head_size_vo=64 +PvcFMHABF16BF16FP32_RCR_h64_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=8 --num_heads_q=32 --num_heads_kv=32 --head_size_vo=96 --head_size_qk=96 +PvcFMHABF16BF16FP32_RCR_h64_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=8 --num_heads_q=32 --num_heads_kv=32 --head_size_vo=96 --head_size_qk=96 + +PvcFMHABF16BF16FP32_RCR_h128_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=16 --num_heads_q=8 --num_heads_kv=8 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=16 --num_heads_q=8 --num_heads_kv=8 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=16 --num_heads_q=4 --num_heads_kv=4 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=16 --num_heads_q=4 --num_heads_kv=4 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=16 --num_heads_q=2 --num_heads_kv=2 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=16 --num_heads_q=2 --num_heads_kv=2 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads_q=32 --num_heads_kv=32 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads_q=32 --num_heads_kv=32 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads_q=96--num_heads_kv=96 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads_q=96--num_heads_kv=96 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=16 --num_heads_q=1 --num_heads_kv=1 --head_size_qk=128 --head_size_vo=128 +PvcFMHABF16BF16FP32_RCR_h128_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=16 --num_heads_q=1 --num_heads_kv=1 --head_size_qk=128 --head_size_vo=128 + +PvcFMHABF16BF16FP32_RCR_h128_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=2048 --batch=1 --num_heads_q=32, --num_heads_kv=8 --head_size_vo=128 --head_size_qk=128 +PvcFMHABF16BF16FP32_RCR_h128_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=2048 --batch=1 --num_heads_q=32, --num_heads_kv=8 --head_size_vo=128 --head_size_qk=128 + + + +PvcFMHABF16BF16FP32_RCR_h192_Causal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads_q=128 --num_heads_kv=128 --head_size_vo=192 --head_size_qk=192 +PvcFMHABF16BF16FP32_RCR_h192_NonCausal_VarLen --bm_name=bf16_bf16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads_q=128 --num_heads_kv=128 --head_size_vo=192 --head_size_qk=192 + +# FMHA FP16 benchmarks +PvcFMHAFP16FP16FP32_RCR_h64_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads_q=32 --num_heads_kv=32 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads_q=32 --num_heads_kv=32 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=4 --num_heads_q=48 --num_heads_kv=48 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=4 --num_heads_q=48 --num_heads_kv=48 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=32 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=32 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=32 --num_heads_q=8 --num_heads_kv=8 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=32 --num_heads_q=8 --num_heads_kv=8 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=32 --num_heads_q=4 --num_heads_kv=4 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=32 --num_heads_q=4 --num_heads_kv=4 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=32 --num_heads_q=2 --num_heads_kv=2 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=32 --num_heads_q=2 --num_heads_kv=2 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=32 --num_heads_q=1 --num_heads_kv=1 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=32 --num_heads_q=1 --num_heads_kv=1 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=8 --num_heads_q=32 --num_heads_kv=32 --head_size_vo=96 --head_size_qk=96 +PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=8 --num_heads_q=32 --num_heads_kv=32 --head_size_vo=96 --head_size_qk=96 + +PvcFMHAFP16FP16FP32_RCR_h128_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=16 --num_heads_q=8 --num_heads_kv=8 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=16 --num_heads_q=8 --num_heads_kv=8 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=16 --num_heads_q=4 --num_heads_kv=4 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=16 --num_heads_q=4 --num_heads_kv=4 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=16 --num_heads_q=2 --num_heads_kv=2 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=16 --num_heads_q=2 --num_heads_kv=2 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads_q=32 --num_heads_kv=32 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads_q=32 --num_heads_kv=32 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads_q=96--num_heads_kv=96 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads_q=96--num_heads_kv=96 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=16 --num_heads_q=1 --num_heads_kv=1 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=16 --num_heads_q=1 --num_heads_kv=1 --head_size_qk=128 --head_size_vo=128 + +PvcFMHAFP16FP16FP32_RCR_h128_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=2048 --batch=1 --num_heads_q=32, --num_heads_kv=8 --head_size_vo=128 --head_size_qk=128 +PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=2048 --batch=1 --num_heads_q=32, --num_heads_kv=8 --head_size_vo=128 --head_size_qk=128 + +PvcFMHAFP16FP16FP32_RCR_h192_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads_q=128 --num_heads_kv=128 --head_size_vo=192 --head_size_qk=192 +PvcFMHAFP16FP16FP32_RCR_h192_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads_q=128 --num_heads_kv=128 --head_size_vo=192 --head_size_qk=192 + +PvcFMHAFP16FP16FP32_RCR_h192_Causal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=2048 --batch=1 --num_heads_q=32, --num_heads_kv=8 --head_size_vo=128 --head_size_qk=128 +PvcFMHAFP16FP16FP32_RCR_h192_NonCausal_FixedLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=2048 --batch=1 --num_heads_q=32, --num_heads_kv=8 --head_size_vo=128 --head_size_qk=128 +PvcFMHAFP16FP16FP32_RCR_h64_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads_q=32 --num_heads_kv=32 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads_q=32 --num_heads_kv=32 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=4 --num_heads_q=48 --num_heads_kv=48 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=4 --num_heads_q=48 --num_heads_kv=48 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=32 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=32 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=32 --num_heads_q=8 --num_heads_kv=8 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=32 --num_heads_q=8 --num_heads_kv=8 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=32 --num_heads_q=4 --num_heads_kv=4 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=32 --num_heads_q=4 --num_heads_kv=4 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=32 --num_heads_q=2 --num_heads_kv=2 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=32 --num_heads_q=2 --num_heads_kv=2 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=32 --num_heads_q=1 --num_heads_kv=1 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=32 --num_heads_q=1 --num_heads_kv=1 --head_size_qk=64 --head_size_vo=64 +PvcFMHAFP16FP16FP32_RCR_h64_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=8 --num_heads_q=32 --num_heads_kv=32 --head_size_vo=96 --head_size_qk=96 +PvcFMHAFP16FP16FP32_RCR_h64_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=8 --num_heads_q=32 --num_heads_kv=32 --head_size_vo=96 --head_size_qk=96 + +PvcFMHAFP16FP16FP32_RCR_h128_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=2048 --batch=1 --num_heads_q=32, --num_heads_kv=8 --head_size_vo=128 --head_size_qk=128 +PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=2048 --batch=1 --num_heads_q=32, --num_heads_kv=8 --head_size_vo=128 --head_size_qk=128 + +PvcFMHAFP16FP16FP32_RCR_h128_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=512 --seq_len_kv=512 --batch=32 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads_q=16 --num_heads_kv=16 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=16 --num_heads_q=8 --num_heads_kv=8 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=16 --num_heads_q=8 --num_heads_kv=8 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=16 --num_heads_q=4 --num_heads_kv=4 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=4096 --seq_len_kv=4096 --batch=16 --num_heads_q=4 --num_heads_kv=4 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=16 --num_heads_q=2 --num_heads_kv=2 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=8192 --seq_len_kv=8192 --batch=16 --num_heads_q=2 --num_heads_kv=2 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads_q=32 --num_heads_kv=32 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads_q=32 --num_heads_kv=32 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads_q=96--num_heads_kv=96 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=2048 --seq_len_kv=2048 --batch=8 --num_heads_q=96--num_heads_kv=96 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=16 --num_heads_q=1 --num_heads_kv=1 --head_size_qk=128 --head_size_vo=128 +PvcFMHAFP16FP16FP32_RCR_h128_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=16384 --seq_len_kv=16384 --batch=16 --num_heads_q=1 --num_heads_kv=1 --head_size_qk=128 --head_size_vo=128 + +PvcFMHAFP16FP16FP32_RCR_h192_Causal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads_q=128 --num_heads_kv=128 --head_size_vo=192 --head_size_qk=192 +PvcFMHAFP16FP16FP32_RCR_h192_NonCausal_VarLen --bm_name=fp16_fp16_fp32 --seq_len_qo=1024 --seq_len_kv=1024 --batch=16 --num_heads_q=128 --num_heads_kv=128 --head_size_vo=192 --head_size_qk=192 From f28f9338a34c795da756d773a215f318a28909c3 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Tadej=20Ciglari=C4=8D?= Date: Wed, 23 Apr 2025 09:25:17 +0200 Subject: [PATCH 6/8] Update examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp Co-authored-by: Joe Todd --- examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp b/examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp index b07861b242..2ff4b9fd4f 100644 --- a/examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp +++ b/examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp @@ -199,6 +199,8 @@ struct ExampleRunner { auto [M, N, K, L] = problem_size; // Padded values + // The inner dimension is padded. Since this example is all RowMajor, + // we require the following: int N_B = cute::ceil_div(N, AlignElemB) * AlignElemB; int N_C = cute::ceil_div(N, AlignElemC) * AlignElemC; int N_D = cute::ceil_div(N, AlignElemD) * AlignElemD; From 821d949af50803fd9dabb3ede23acd74081cfb08 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Tadej=20Ciglari=C4=8D?= Date: Wed, 23 Apr 2025 09:26:12 +0200 Subject: [PATCH 7/8] Update examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp Co-authored-by: Joe Todd --- examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp b/examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp index 2ff4b9fd4f..7ff0489750 100644 --- a/examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp +++ b/examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp @@ -265,7 +265,7 @@ struct ExampleRunner { stride_D = cutlass::make_cute_packed_stride(StrideD{}, cute::make_shape(M_ACD, N_D, L)); block_A.reset(M_ACD * K_A * L); - block_B.reset(K * N_B * L); + block_B.reset(K_B * N_B * L); block_C.reset(M_ACD * N_C * L); block_D.reset(M_ACD * N_D * L); block_ref_D.reset(M_ACD * N_D * L); From 57e561471ef4d9c2c0f23bc475e536edcca3d095 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Tadej=20Ciglari=C4=8D?= Date: Wed, 23 Apr 2025 13:55:29 +0100 Subject: [PATCH 8/8] addressed review comments --- examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp | 24 +++++++++++------------ include/cutlass/gemm/kernel/xe_gemm.hpp | 17 ++++++++++------ 2 files changed, 23 insertions(+), 18 deletions(-) diff --git a/examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp b/examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp index 7ff0489750..42f47cd572 100644 --- a/examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp +++ b/examples/sycl/00_pvc_gemm/00_pvc_gemm.cpp @@ -201,14 +201,14 @@ struct ExampleRunner { // Padded values // The inner dimension is padded. Since this example is all RowMajor, // we require the following: - int N_B = cute::ceil_div(N, AlignElemB) * AlignElemB; - int N_C = cute::ceil_div(N, AlignElemC) * AlignElemC; - int N_D = cute::ceil_div(N, AlignElemD) * AlignElemD; - int K_A = cute::ceil_div(K, AlignElemA) * AlignElemA; + int N_B = cute::round_up(N, AlignElemB); + int N_C = cute::round_up(N, AlignElemC); + int N_D = cute::round_up(N, AlignElemD); + int K_A = cute::round_up(K, AlignElemA); int AlignmentOuter = AlignmentPtr / AlignmentInner; - int M_ACD = cute::ceil_div(M, AlignmentOuter) * AlignmentOuter; - int K_B = cute::ceil_div(K, AlignmentOuter) * AlignmentOuter; + int M_ACD = cute::round_up(M, AlignmentOuter); + int K_B = cute::round_up(K, AlignmentOuter); cutlass::TensorRef ref_A(block_A.get(), LayoutA(K_A)); cutlass::TensorRef ref_B(block_B.get(), LayoutB(N_B)); @@ -249,14 +249,14 @@ struct ExampleRunner { auto [M, N, K, L] = problem_shape_MNKL; // Padded values - int N_B = cute::ceil_div(N, AlignElemB) * AlignElemB; - int N_C = cute::ceil_div(N, AlignElemC) * AlignElemC; - int N_D = cute::ceil_div(N, AlignElemD) * AlignElemD; - int K_A = cute::ceil_div(K, AlignElemA) * AlignElemA; + int N_B = cute::round_up(N, AlignElemB); + int N_C = cute::round_up(N, AlignElemC); + int N_D = cute::round_up(N, AlignElemD); + int K_A = cute::round_up(K, AlignElemA); int AlignmentOuter = AlignmentPtr / AlignmentInner; - int M_ACD = cute::ceil_div(M, AlignmentOuter) * AlignmentOuter; - int K_B = cute::ceil_div(K, AlignmentOuter) * AlignmentOuter; + int M_ACD = cute::round_up(M, AlignmentOuter); + int K_B = cute::round_up(K, AlignmentOuter); // Complete the stride by combining static layout info (StrideA) with runtime size info (M,K,L) stride_A = cutlass::make_cute_packed_stride(StrideA{}, cute::make_shape(M_ACD, K_A, L)); diff --git a/include/cutlass/gemm/kernel/xe_gemm.hpp b/include/cutlass/gemm/kernel/xe_gemm.hpp index a91c2b1a9a..12827c572d 100644 --- a/include/cutlass/gemm/kernel/xe_gemm.hpp +++ b/include/cutlass/gemm/kernel/xe_gemm.hpp @@ -164,13 +164,18 @@ class GemmUniversal< auto k = get<2>(args.problem_shape); auto l = get<3>(args.problem_shape); bool is_batch = l > 1; + // all these requirements are in bytes + constexpr int inner_alignment_requirement = 16; + constexpr int outer_alignment_requirement = 64; + constexpr int width_alignment_requirement = 4; + auto check_stride = [is_batch](auto stride, int el_size){ auto a = get<0>(stride); auto b = get<1>(stride); auto valid_is_unit = a == _1{} || b == _1{}; auto inner = a == _1{} ? b : a; - auto valid_inner = inner % (16 / el_size) == 0; - auto valid_outer = !is_batch || get<2>(stride) % (64 / el_size) == 0; + auto valid_inner = inner % (inner_alignment_requirement / el_size) == 0; + auto valid_outer = !is_batch || get<2>(stride) % (outer_alignment_requirement / el_size) == 0; return valid_is_unit && valid_inner && valid_outer; }; bool strides_valid = check_stride(args.mainloop.dA, sizeof(ElementA)) && @@ -179,10 +184,10 @@ class GemmUniversal< check_stride(args.epilogue.dD, sizeof(ElementD)); // TODO(codeplay): base *_valid on the atom shapes bool m_valid = m > 0; - bool n_valid = n > 0 && n % (4 / sizeof(ElementB)) == 0 && - n % (4 / sizeof(ElementC)) == 0 && - n % (4 / sizeof(ElementD)) == 0; - bool k_valid = k > 0 && k % (4 / sizeof(ElementA)) == 0; + bool n_valid = n > 0 && n % (width_alignment_requirement / sizeof(ElementB)) == 0 && + n % (width_alignment_requirement / sizeof(ElementC)) == 0 && + n % (width_alignment_requirement / sizeof(ElementD)) == 0; + bool k_valid = k > 0 && k % (width_alignment_requirement / sizeof(ElementA)) == 0; bool shape_implementable = m_valid && n_valid && k_valid && strides_valid; bool mode_implementable = args.mode == GemmUniversalMode::kGemm ||