Fix benchmark

Author: leonling-lly

benchmarks/flash_attention/flash_attention_decode/benchmark_runner.hpp

@@ -630,7 +630,7 @@ template <class FMHADecodeConfiguration> struct BenchmarkRunnerFMHADecode {
       sycl::ext::oneapi::experimental::sub_group_size<FMHADecodeKernel::DispatchPolicy::SubgroupSize>
     };
     cutlasscompat::experimental::launch_policy policy{sycl_grid, sycl_block, launch_props, kernel_props};
-    auto event = cutlasscompat::experimental::launch<cutlass::device_kernel<FMHADecodeKernel>>(policy, params);
+    auto event = cutlasscompat::experimental::launch<cutlass::device_kernel<FMHADecodeKernel>, FMHADecodeKernel>(policy, params);
 #endif
 
     EventManager::getInstance().addEvent(event);

benchmarks/flash_attention/flash_attention_prefill/benchmark_runner.hpp

@@ -490,7 +490,7 @@ template <class FMHAPrefillConfiguration> struct BenchmarkRunnerFMHA {
       sycl::ext::oneapi::experimental::sub_group_size<GemmKernel::DispatchPolicy::SubgroupSize>
     };
     cutlasscompat::experimental::launch_policy policy{sycl_grid, sycl_block, launch_props, kernel_props};
-    auto event = cutlasscompat::experimental::launch<cutlass::device_kernel<GemmKernel>>(policy, params);
+    auto event = cutlasscompat::experimental::launch<cutlass::device_kernel<GemmKernel>, GemmKernel>(policy, params);
 #endif
 
     EventManager::getInstance().addEvent(event);

benchmarks/flash_attention/flash_attention_prefill_cachedKV/benchmark_runner.hpp

@@ -583,7 +583,7 @@ template <class FMHAPrefillConfiguration> struct BenchmarkRunnerFMHA {
       sycl::ext::oneapi::experimental::sub_group_size<GemmKernel::DispatchPolicy::SubgroupSize>
     };
     cutlasscompat::experimental::launch_policy policy{sycl_grid, sycl_block, launch_props, kernel_props};
-    auto event = cutlasscompat::experimental::launch<cutlass::device_kernel<GemmKernel>>(policy, params);
+    auto event = cutlasscompat::experimental::launch<cutlass::device_kernel<GemmKernel>, GemmKernel>(policy, params);
 #endif
 
     EventManager::getInstance().addEvent(event);