Skip to content

Refactor benchmarks for Flash Attention Prefill #447

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 6 commits into
base: sycl-develop
Choose a base branch
from
Open

Refactor benchmarks for Flash Attention Prefill #447

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
f64c7b7
Separate out type and accum type
muhammad-tanvir-1211 Jun 25, 2025
2a87b5f
Removed redundant wait from runners
muhammad-tanvir-1211 Jun 27, 2025
6a7ad57
Add benchmarks for all types with refactor
muhammad-tanvir-1211 Jun 26, 2025
3f5c19a
Merge branch 'sycl-develop' into flash_prefill_separate_out_benchmarks
muhammad-tanvir-1211 Jul 1, 2025
9503357
Update output conversion
muhammad-tanvir-1211 Jul 1, 2025
2c679d9
Hardcode head_size_vo
muhammad-tanvir-1211 Jul 7, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
10 changes: 8 additions & 2 deletions .github/workflows/intel_test.yml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -95,7 +95,8 @@ jobs:
cmake -G Ninja \
-DCUTLASS_ENABLE_SYCL=ON \
-DDPCPP_SYCL_TARGET=${{ matrix.sycl_target }} \
-DCUTLASS_SYCL_RUNNING_CI=ON
-DCUTLASS_SYCL_RUNNING_CI=ON \
-DCUTLASS_ENABLE_BENCHMARKS=OFF
cmake --build .
- name: Unit test
shell: bash
Expand All @@ -108,4 +109,9 @@ jobs:
- name: Benchmarks
shell: bash
run: |
cmake --build . --target cutlass_benchmarks
cmake -G Ninja \
-DCUTLASS_ENABLE_SYCL=ON \
-DDPCPP_SYCL_TARGET=${{ matrix.sycl_target }} \
-DCUTLASS_SYCL_RUNNING_CI=ON \
-DCUTLASS_ENABLE_BENCHMARKS=ON
cmake --build . --target cutlass_benchmarks -j 1
18 changes: 8 additions & 10 deletions applications/flash_attention_v2/collective/xe_flash_attn_prefill_epilogue.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -196,18 +196,16 @@ class FlashPrefillEpilogue<epilogue::IntelXeXMX16, MMAOperation_, TileShapeOutp
auto thread_xe_store_o = params.xe_store_o.get_thread_slice(ThreadIdxX());
Tensor tOgO = thread_xe_store_o.partition_D(gO);

Tensor final_out_reg = make_fragment_like<ElementOutput>(out_reg);
// iff ElementOutput == ElementAccumulator, then convert_type doesn't do the right conversion
// iff ElementOutput == fp8, there is no NumericConverter specialization available
// for both the above cases, we call copy() which internally performs a static_cast op on the data.
// for ElementOutput == bf16 | fp16, convert_type calls relevant NumericConverter specialization.
if constexpr (cute::is_any_of_v<ElementOutput, cute::float_e5m2_t, cute::float_e4m3_t> || cute::is_same_v<ElementOutput, ElementCompute>) {
copy(out_reg, final_out_reg);
// iff ElementOutput == ElementAccumulator, call copy directly.
// for ElementOutput == bf16 | fp16, convert_type calls relevant NumericConverter specialization.
// iff ElementOutput == fp8, there is no NumericConverter specialization available so convert_type
// performs static_cast under the hood.
if constexpr (cute::is_same_v<ElementOutput, ElementCompute>) {
copy(params.xe_store_o, out_reg, tOgO);
} else {
Tensor temp = convert_type<ElementOutput>(out_reg);
copy(temp, final_out_reg);
Tensor final_out_reg = convert_type<ElementOutput>(out_reg);
copy(params.xe_store_o, final_out_reg, tOgO);
}
copy(params.xe_store_o, final_out_reg, tOgO);
}

// SequenceLengthShapeType = Shape<int, int>
Expand Down
168 changes: 0 additions & 168 deletions benchmarks/device/pvc/input_files/input_flash_attention_prefill.in
View file
Open in desktop

This file was deleted.

72 changes: 72 additions & 0 deletions benchmarks/device/pvc/input_files/input_flash_attention_prefill_h128.in
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,72 @@
# FMHA Prefill BFloat16 benchmarks
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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

FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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

########################################################################################################################################################################################

# FMHA Prefill FP16 benchmarks

FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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

FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
16 changes: 16 additions & 0 deletions benchmarks/device/pvc/input_files/input_flash_attention_prefill_h192.in
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,16 @@
# FMHA Prefill BFloat16 benchmarks
FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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

FMHAPrefillBF16BF16FP32FP32_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
FMHAPrefillBF16BF16FP32FP32_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 Prefill FP16 benchmarks

FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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

FMHAPrefillFP16FP16FP32FP32_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
FMHAPrefillFP16FP16FP32FP32_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
Loading
Loading