diff --git a/example/ck_tile/17_grouped_gemm/grouped_gemm_persistent_async.cpp b/example/ck_tile/17_grouped_gemm/grouped_gemm_persistent_async.cpp index 76cad154c0..e9aaa55ad5 100644 --- a/example/ck_tile/17_grouped_gemm/grouped_gemm_persistent_async.cpp +++ b/example/ck_tile/17_grouped_gemm/grouped_gemm_persistent_async.cpp @@ -279,17 +279,16 @@ int run_grouped_gemm_persistent_async_example(ck_tile::ArgParser& arg_parser, // Launch persistent async kernel std::cout << "\nLaunching persistent async GEMM kernel..." << std::endl; - float ave_time = - invoke_grouped_gemm_persistent, - CDataType, - ck_tile::tuple<>, - ALayout, - BLayout, - CLayout>(stream, group_count, kargs_ptr); + float ave_time = invoke_grouped_gemm_persistent, + CDataType, + ck_tile::tuple<>, + ALayout, + BLayout, + CLayout>(stream, group_count, kargs_ptr); std::size_t total_flops = 0; std::size_t total_bytes = 0; diff --git a/example/ck_tile/17_grouped_gemm/grouped_gemm_persistent_async.hpp b/example/ck_tile/17_grouped_gemm/grouped_gemm_persistent_async.hpp index 33ba25c248..325cf4bc63 100644 --- a/example/ck_tile/17_grouped_gemm/grouped_gemm_persistent_async.hpp +++ b/example/ck_tile/17_grouped_gemm/grouped_gemm_persistent_async.hpp @@ -14,8 +14,8 @@ template float invoke_grouped_gemm_persistent(const ck_tile::stream_config& s, - const ck_tile::index_t num_groups, - void* kargs_ptr) + const ck_tile::index_t num_groups, + void* kargs_ptr) { constexpr bool TransposeC = false; constexpr bool DoubleSmemBuffer = false; @@ -46,35 +46,35 @@ float invoke_grouped_gemm_persistent(const ck_tile::stream_config& s, CLayout, TransposeC>; - constexpr auto scheduler = ck_tile::GemmPipelineScheduler::Intrawave; + constexpr auto scheduler = ck_tile::GemmPipelineScheduler::Intrawave; // We create the GEMM pipeline without specifying hotloop or tailnumber. // These are automatically run inside the kernel based on the given input data. using UniversalGemmProblem = ck_tile::UniversalGemmPipelineProblem; + BDataType, + AccDataType, + GemmShape, + GemmUniversalTraits, + scheduler>; using GemmPipeline = ck_tile::GemmPipelineAgBgCrCompV3; using GemmEpilogue = ck_tile::CShuffleEpilogue< ck_tile::CShuffleEpilogueProblem>; + BDataType, + DsDataType, + AccDataType, + CDataType, + DsLayout, + CLayout, + ck_tile::element_wise::PassThrough, + TilePartitioner::MPerBlock, + TilePartitioner::NPerBlock, + GroupedGemKernelParam::M_Warp, + GroupedGemKernelParam::N_Warp, + GroupedGemKernelParam::M_Warp_Tile, + GroupedGemKernelParam::N_Warp_Tile, + GroupedGemKernelParam::K_Warp_Tile, + UniversalGemmProblem::TransposeC>>; using Kernel = ck_tile::GroupedGemmKernel; const dim3 blocks = Kernel::BlockSize(); const dim3 grids = Kernel::MaxOccupancyGridSize(s); @@ -82,16 +82,16 @@ float invoke_grouped_gemm_persistent(const ck_tile::stream_config& s, if(s.log_level_ > 0) { std::cout << "Launching kernel: " << Kernel::GetName() << " with args:" << " grid: {" - << grids.x << ", " << grids.y << ", " << grids.z << "}" << ", blocks: {" - << blocks.x << ", " << blocks.y << ", " << blocks.z << "}" << std::endl; + << grids.x << ", " << grids.y << ", " << grids.z << "}" << ", blocks: {" + << blocks.x << ", " << blocks.y << ", " << blocks.z << "}" << std::endl; } return ck_tile::launch_kernel(s, - ck_tile::make_kernel( - Kernel{}, - grids, - blocks, - 0, - ck_tile::cast_pointer_to_constant_address_space(kargs_ptr), - num_groups)); + ck_tile::make_kernel( + Kernel{}, + grids, + blocks, + 0, + ck_tile::cast_pointer_to_constant_address_space(kargs_ptr), + num_groups)); }