diff --git a/projects/composablekernel/dispatcher/bindings/ctypes/quant_bridge_common.hpp b/projects/composablekernel/dispatcher/bindings/ctypes/quant_bridge_common.hpp index 3645e81fdf41..3c6da3200d19 100644 --- a/projects/composablekernel/dispatcher/bindings/ctypes/quant_bridge_common.hpp +++ b/projects/composablekernel/dispatcher/bindings/ctypes/quant_bridge_common.hpp @@ -124,11 +124,12 @@ inline bool validate_supported_arch(const char* fn, bool allow_gfx90a = false) } const std::string arch(props.gcnArchName); const bool ok = arch.rfind("gfx950", 0) == 0 || arch.rfind("gfx942", 0) == 0 || - (allow_gfx90a && arch.rfind("gfx90a", 0) == 0); + arch.rfind("gfx1250", 0) == 0 || (allow_gfx90a && arch.rfind("gfx90a", 0) == 0); if(!ok) { std::cerr << fn << ": unsupported GPU architecture '" << arch - << "' (supported: " << (allow_gfx90a ? "gfx90a, " : "") << "gfx942, gfx950)\n"; + << "' (supported: " << (allow_gfx90a ? "gfx90a, " : "") + << "gfx942, gfx950, gfx1250)\n"; return false; } return true; diff --git a/projects/composablekernel/dispatcher/python/grouped_gemm_abquant_utils.py b/projects/composablekernel/dispatcher/python/grouped_gemm_abquant_utils.py index a55458104fac..5248d34b5ab8 100644 --- a/projects/composablekernel/dispatcher/python/grouped_gemm_abquant_utils.py +++ b/projects/composablekernel/dispatcher/python/grouped_gemm_abquant_utils.py @@ -857,3 +857,78 @@ def default_bf8_preshuffleb_config( k_block_per_cu=2, gfx_arch=gfx_arch, ) + + +# ============================================================================= +# gfx1250 (MI400) default configs +# ============================================================================= +# gfx1250 uses the standard CompV3 pipeline (NOT the gfx950-native eightwaves / +# preshuffleb paths, which require CK_GFX950_SUPPORT and TransposeC=True). +# +# Empirically (on-GPU, MI400/gfx1250, HIP_VISIBLE_DEVICES=0, fp32-dequant +# reference), the ABQuant CompV3 fp8/bf8 kernels are correct on gfx1250 only +# with warp_tile_k=128 (the FlatMM tile). The gfx9 stock CompV3 default uses +# warp_tile_k=32 (plain MFMA) which returns all-zeros on gfx1250; warp_tile_k=16 +# (gfx12 WMMA) also returns all-zeros. So the only correct gfx1250 override for +# the CompV3 pipeline is 16x16x128 with transpose_c=False. +# +# fp8: warp_tile 16x16x128 -> nonzero, max_rel_err ~3e-4 (PASS) +# bf8: warp_tile 16x16x128 -> nonzero, max_rel_err ~3e-4 (PASS) +# +# The compile path injects -DCK_USE_OCP_FP8 / -DCK_TILE_USE_OCP_FP8 for gfx12 +# archs, so fp8/bf8 use the OCP encoding on gfx1250. + +_GFX1250_ARCH = "gfx1250" + + +def default_fp8_compv3_config_gfx1250( + quant_group_k: int = 128, + bquant_group_n: int = 1, + gfx_arch: str = _GFX1250_ARCH, +) -> ABQuantKernelConfig: + """fp8 ABQuant CompV3 config for gfx1250 (FlatMM warp_tile_k=128, TransposeC=False). + + GPU-verified on MI400/gfx1250: max_rel_err ~3e-4 vs. fp32 dequant reference. + warp_tile_k=32 (gfx9 stock) and warp_tile_k=16 (gfx12 WMMA) both zero out. + """ + return ABQuantKernelConfig( + variant_key="fp8", + layout="rcr", + pipeline="compv3", + epilogue="cshuffle", + scheduler="intrawave", + tile_m=128, tile_n=128, tile_k=128, + warp_m=1, warp_n=4, warp_k=1, + warp_tile_m=16, warp_tile_n=16, warp_tile_k=128, + aquant_group_m=1, aquant_group_n=1, aquant_group_k=quant_group_k, + bquant_group_m=1, bquant_group_n=bquant_group_n, bquant_group_k=quant_group_k, + preshuffle_b=False, preshuffle_aq=False, preshuffle_bq=False, + transpose_c=False, + gfx_arch=gfx_arch, + ) + + +def default_bf8_compv3_config_gfx1250( + quant_group_k: int = 128, + bquant_group_n: int = 1, + gfx_arch: str = _GFX1250_ARCH, +) -> ABQuantKernelConfig: + """bf8 ABQuant CompV3 config for gfx1250 (FlatMM warp_tile_k=128, TransposeC=False). + + GPU-verified on MI400/gfx1250: max_rel_err ~3e-4 vs. fp32 dequant reference. + """ + return ABQuantKernelConfig( + variant_key="bf8", + layout="rcr", + pipeline="compv3", + epilogue="cshuffle", + scheduler="intrawave", + tile_m=128, tile_n=128, tile_k=128, + warp_m=1, warp_n=4, warp_k=1, + warp_tile_m=16, warp_tile_n=16, warp_tile_k=128, + aquant_group_m=1, aquant_group_n=1, aquant_group_k=quant_group_k, + bquant_group_m=1, bquant_group_n=bquant_group_n, bquant_group_k=quant_group_k, + preshuffle_b=False, preshuffle_aq=False, preshuffle_bq=False, + transpose_c=False, + gfx_arch=gfx_arch, + ) diff --git a/projects/composablekernel/dispatcher/python/grouped_gemm_aquant_utils.py b/projects/composablekernel/dispatcher/python/grouped_gemm_aquant_utils.py index 60e529446e2c..8ae843cc0c43 100644 --- a/projects/composablekernel/dispatcher/python/grouped_gemm_aquant_utils.py +++ b/projects/composablekernel/dispatcher/python/grouped_gemm_aquant_utils.py @@ -785,3 +785,81 @@ def default_bf8_preshuffleaq_config( preshuffle_aq=True, gfx_arch=gfx_arch, ) + + +# ============================================================================= +# gfx1250 (MI400) default configs +# ============================================================================= +# Empirically (on-GPU, MI400/gfx1250, HIP_VISIBLE_DEVICES=0, fp32-dequant +# reference), the AQuant fp8/bf8 kernels are correct on gfx1250 only with +# warp_tile_k=128 (the FlatMM path, same tile the gfx9 fp8/bf8 FlatMM decode +# uses). The gfx9 stock AQuant default uses warp_tile_k=32 (plain MFMA), which +# produces a WRONG result on gfx1250 (max_rel_err=1.0); warp_tile_k=16 (gfx12 +# WMMA) returns all-zeros. So the only correct gfx1250 override is 16x16x128. +# +# fp8: warp_tile 16x16x128 -> nonzero, max_rel_err ~3e-4 (PASS) +# bf8: warp_tile 16x16x128 -> nonzero, max_rel_err ~1e-4 (PASS) +# +# The compile path injects -DCK_USE_OCP_FP8 / -DCK_TILE_USE_OCP_FP8 for gfx12 +# archs, so fp8/bf8 use the OCP encoding on gfx1250. +# +# fp8i4 / bf8i4 (pk_int4 weights) do NOT compile on gfx1250 at any warp_tile_k +# (no gfx12 WMMA/FlatMM instruction for the packed-int4 quant path). Those +# variants are therefore not enabled on gfx1250: the C++ ctypes arch-gate is +# opened, but on-GPU correctness is deferred pending kernel work. No gfx1250 +# config helper is provided for i4 because it would emit a non-buildable kernel. + +_GFX1250_ARCH = "gfx1250" + + +def default_fp8_config_gfx1250( + quant_group_k: int = 128, + quant_group_m: int = 1, + gfx_arch: str = _GFX1250_ARCH, +) -> AQuantKernelConfig: + """fp8 AQuant decode config for gfx1250 (FlatMM, warp_tile_k=128). + + GPU-verified on MI400/gfx1250: max_rel_err ~3e-4 vs. fp32 dequant reference. + warp_tile_k=32 (gfx9 stock) is WRONG on gfx1250; warp_tile_k=16 zeros out. + """ + return AQuantKernelConfig( + variant_key="fp8", + layout="rcr", + pipeline="mem", + epilogue="cshuffle", + scheduler="intrawave", + tile_m=16, tile_n=64, tile_k=256, + warp_m=1, warp_n=4, warp_k=1, + warp_tile_m=16, warp_tile_n=16, warp_tile_k=128, + quant_group_m=quant_group_m, + quant_group_n=1, + quant_group_k=quant_group_k, + preshuffle_aq=False, + gfx_arch=gfx_arch, + ) + + +def default_bf8_config_gfx1250( + quant_group_k: int = 128, + quant_group_m: int = 1, + gfx_arch: str = _GFX1250_ARCH, +) -> AQuantKernelConfig: + """bf8 AQuant decode config for gfx1250 (FlatMM, warp_tile_k=128). + + GPU-verified on MI400/gfx1250: max_rel_err ~1e-4 vs. fp32 dequant reference. + """ + return AQuantKernelConfig( + variant_key="bf8", + layout="rcr", + pipeline="mem", + epilogue="cshuffle", + scheduler="intrawave", + tile_m=16, tile_n=64, tile_k=256, + warp_m=1, warp_n=4, warp_k=1, + warp_tile_m=16, warp_tile_n=16, warp_tile_k=128, + quant_group_m=quant_group_m, + quant_group_n=1, + quant_group_k=quant_group_k, + preshuffle_aq=False, + gfx_arch=gfx_arch, + ) diff --git a/projects/composablekernel/dispatcher/python/grouped_gemm_bquant_utils.py b/projects/composablekernel/dispatcher/python/grouped_gemm_bquant_utils.py index 92fb0038c8f7..6f68a5b37573 100644 --- a/projects/composablekernel/dispatcher/python/grouped_gemm_bquant_utils.py +++ b/projects/composablekernel/dispatcher/python/grouped_gemm_bquant_utils.py @@ -1214,3 +1214,62 @@ def default_mx_bf16fp4_config( quant_group_k=quant_group_k, gfx_arch=gfx_arch, ) + + +# ============================================================================= +# gfx1250 (MI400) default configs +# ============================================================================= +# Empirically (on-GPU, MI400/gfx1250, HIP_VISIBLE_DEVICES=0, via the +# test_bquant_gpu_correctness fp32-dequant reference), the BQuant fp8/bf8 kernels +# are correct on gfx1250 with the SAME config that gfx9 uses: the CompV3 FlatMM +# tile with warp_tile_k=128. warp_tile_k=16 (gfx12 WMMA) returns all-zeros on +# gfx1250. So the gfx1250 helpers simply return the stock fp8/bf8 config +# (warp_tile_k=128); no distinct gfx1250 tile is required for BQuant. +# +# fp8: warp_tile 16x16x128 -> nonzero, max_rel_err ~5e-4 (PASS) +# bf8: warp_tile 16x16x128 -> nonzero, max_rel_err ~5e-4 (PASS) +# +# The compile path injects -DCK_USE_OCP_FP8 / -DCK_TILE_USE_OCP_FP8 for gfx12 +# archs, so fp8/bf8 use the OCP encoding on gfx1250. +# +# fp8i4 / bf8i4 (pk_int4 weights) do NOT compile on gfx1250 at any warp_tile_k +# (no gfx12 instruction for the packed-int4 quant path). Those variants are not +# enabled on gfx1250: the C++ ctypes arch-gate is opened, but on-GPU correctness +# is deferred pending kernel work. No gfx1250 config helper is provided for i4 +# because it would emit a non-buildable kernel. + +_GFX1250_ARCH = "gfx1250" + + +def default_fp8_config_gfx1250( + quant_group_k: int = 128, + quant_group_n: int = 1, + gfx_arch: str = _GFX1250_ARCH, +) -> BQuantKernelConfig: + """fp8 BQuant config for gfx1250 (compv3, FlatMM warp_tile_k=128). + + Identical to the stock fp8 config; GPU-verified on MI400/gfx1250 with + max_rel_err ~5e-4 vs. fp32 dequant reference. warp_tile_k=16 zeros out. + """ + return default_fp8_config( + quant_group_k=quant_group_k, + quant_group_n=quant_group_n, + gfx_arch=gfx_arch, + ) + + +def default_bf8_config_gfx1250( + quant_group_k: int = 128, + quant_group_n: int = 1, + gfx_arch: str = _GFX1250_ARCH, +) -> BQuantKernelConfig: + """bf8 BQuant config for gfx1250 (compv3, FlatMM warp_tile_k=128). + + Identical to the stock bf8 config; GPU-verified on MI400/gfx1250 with + max_rel_err ~5e-4 vs. fp32 dequant reference. + """ + return default_bf8_config( + quant_group_k=quant_group_k, + quant_group_n=quant_group_n, + gfx_arch=gfx_arch, + ) diff --git a/projects/composablekernel/dispatcher/tests/test_grouped_gemm_abquant_utils.py b/projects/composablekernel/dispatcher/tests/test_grouped_gemm_abquant_utils.py index 1d102e28f00c..86f67b087b04 100644 --- a/projects/composablekernel/dispatcher/tests/test_grouped_gemm_abquant_utils.py +++ b/projects/composablekernel/dispatcher/tests/test_grouped_gemm_abquant_utils.py @@ -361,3 +361,55 @@ def test_transpose_c_in_name_only_when_true(self): without_tc = default_fp8_compv3_config() assert "transposec" in with_tc.name assert "transposec" not in without_tc.name + + +# ============================================================================= +# gfx1250 (MI400 / WMMA) default configs +# ============================================================================= + + +from grouped_gemm_abquant_utils import ( # noqa: E402 + default_fp8_compv3_config_gfx1250, + default_bf8_compv3_config_gfx1250, +) + + +class TestGfx1250Configs: + # gfx1250 correct config was determined empirically on MI400/gfx1250: + # ABQuant CompV3 fp8/bf8 verify only with warp_tile_k=128 (FlatMM tile); + # warp_tile_k=32 (gfx9 MFMA) and warp_tile_k=16 (WMMA) both produce all-zeros + # on gfx12. Still uses the standard CompV3 pipeline with transpose_c=False + # (NOT the gfx950-native eightwaves/preshuffleb path). + + def _all(self): + return [ + default_fp8_compv3_config_gfx1250(), + default_bf8_compv3_config_gfx1250(), + ] + + def test_gfx1250_uses_flatmm_warp_tile_k_128(self): + # GPU-verified: CompV3 fp8/bf8 are correct on gfx1250 only with + # warp_tile_k=128; warp_tile_k=32 and warp_tile_k=16 both zero out. + for cfg in self._all(): + assert cfg.warp_tile_m == 16 + assert cfg.warp_tile_n == 16 + assert cfg.warp_tile_k == 128, f"{cfg.name} must use FlatMM warp_tile_k=128" + + def test_gfx1250_uses_compv3_not_eightwaves(self): + # eightwaves / preshuffleb are gfx950-only; gfx1250 uses standard compv3. + for cfg in self._all(): + assert cfg.pipeline == "compv3" + assert cfg.transpose_c is False + assert cfg.preshuffle_b is False + + def test_gfx1250_arch_propagated(self): + for cfg in self._all(): + assert cfg.gfx_arch == "gfx1250" + + def test_gfx1250_names_16x16x128(self): + for cfg in self._all(): + assert "16x16x128" in cfg.name + + def test_gfx1250_unique_names(self): + names = [cfg.name for cfg in self._all()] + assert len(names) == len(set(names)) diff --git a/projects/composablekernel/dispatcher/tests/test_grouped_gemm_aquant_utils.py b/projects/composablekernel/dispatcher/tests/test_grouped_gemm_aquant_utils.py index be95e88c0822..9f9da689a226 100644 --- a/projects/composablekernel/dispatcher/tests/test_grouped_gemm_aquant_utils.py +++ b/projects/composablekernel/dispatcher/tests/test_grouped_gemm_aquant_utils.py @@ -408,3 +408,58 @@ def test_mem_and_compv3_pipelines_produce_different_names(self): assert mem_cfg.name != compv3_cfg.name assert "mem" in mem_cfg.name assert "compv3" in compv3_cfg.name + + +# ============================================================================= +# gfx1250 (MI400 / WMMA) default configs +# ============================================================================= + + +from grouped_gemm_aquant_utils import ( # noqa: E402 + default_fp8_config_gfx1250, + default_bf8_config_gfx1250, +) + + +class TestGfx1250Configs: + # gfx1250 correct config was determined empirically on MI400/gfx1250: + # fp8/bf8 verify only with warp_tile_k=128 (FlatMM); warp_tile_k=16 (WMMA) + # and warp_tile_k=32 (gfx9 MFMA) both produce wrong/zero output on gfx12. + # i4 variants do not compile on gfx1250, so no gfx1250 helper exists for them. + + def _all(self): + return [ + default_fp8_config_gfx1250(), + default_bf8_config_gfx1250(), + ] + + def test_gfx1250_uses_flatmm_warp_tile_k_128(self): + # GPU-verified: fp8/bf8 are correct on gfx1250 only with warp_tile_k=128. + for cfg in self._all(): + assert cfg.warp_tile_m == 16 + assert cfg.warp_tile_n == 16 + assert cfg.warp_tile_k == 128, f"{cfg.name} must use FlatMM warp_tile_k=128" + + def test_gfx1250_arch_propagated(self): + for cfg in self._all(): + assert cfg.gfx_arch == "gfx1250" + + def test_gfx1250_names_16x16x128(self): + for cfg in self._all(): + assert "16x16x128" in cfg.name + + def test_gfx1250_pipeline_is_mem(self): + for cfg in self._all(): + assert cfg.pipeline == "mem" + + def test_gfx1250_layout_is_rcr(self): + for cfg in self._all(): + assert cfg.layout == "rcr" + + def test_gfx1250_no_preshuffle(self): + for cfg in self._all(): + assert cfg.preshuffle_aq is False + + def test_gfx1250_unique_names(self): + names = [cfg.name for cfg in self._all()] + assert len(names) == len(set(names)) diff --git a/projects/composablekernel/dispatcher/tests/test_grouped_gemm_bquant_utils.py b/projects/composablekernel/dispatcher/tests/test_grouped_gemm_bquant_utils.py index 9f4ceb6face1..654b6b90d1d5 100644 --- a/projects/composablekernel/dispatcher/tests/test_grouped_gemm_bquant_utils.py +++ b/projects/composablekernel/dispatcher/tests/test_grouped_gemm_bquant_utils.py @@ -724,3 +724,55 @@ def test_name_matches_kernel_config_name(self, tmp_path): quant_group_m=1, quant_group_n=1, quant_group_k=128, ) assert configs[0].name == manual.name + + +# ============================================================================= +# gfx1250 (MI400 / WMMA) default configs +# ============================================================================= + + +from grouped_gemm_bquant_utils import ( # noqa: E402 + default_fp8_config_gfx1250, + default_bf8_config_gfx1250, +) + + +class TestGfx1250Configs: + # gfx1250 correct config was determined empirically on MI400/gfx1250: + # BQuant fp8/bf8 verify with the SAME stock config (warp_tile_k=128, FlatMM); + # warp_tile_k=16 (WMMA) produces all-zeros on gfx12. i4 variants do not + # compile on gfx1250, so no gfx1250 helper exists for them. + + def _all(self): + return [ + default_fp8_config_gfx1250(), + default_bf8_config_gfx1250(), + ] + + def test_gfx1250_uses_flatmm_warp_tile_k_128(self): + # GPU-verified: fp8/bf8 are correct on gfx1250 with warp_tile_k=128 + # (== stock). warp_tile_k=16 silently returns zeros on gfx12. + for cfg in self._all(): + assert cfg.warp_tile_m == 16 + assert cfg.warp_tile_n == 16 + assert cfg.warp_tile_k == 128, f"{cfg.name} must use FlatMM warp_tile_k=128" + + def test_gfx1250_arch_propagated(self): + for cfg in self._all(): + assert cfg.gfx_arch == "gfx1250" + + def test_gfx1250_names_16x16x128(self): + for cfg in self._all(): + assert "16x16x128" in cfg.name + + def test_gfx1250_pipeline_is_compv3(self): + for cfg in self._all(): + assert cfg.pipeline == "compv3" + + def test_gfx1250_layout_is_rcr(self): + for cfg in self._all(): + assert cfg.layout == "rcr" + + def test_gfx1250_unique_names(self): + names = [cfg.name for cfg in self._all()] + assert len(names) == len(set(names)) diff --git a/projects/composablekernel/tile_engine/ops/gemm/gemm_instance_builder.py b/projects/composablekernel/tile_engine/ops/gemm/gemm_instance_builder.py index 1026990fdf2a..90a119f90e64 100644 --- a/projects/composablekernel/tile_engine/ops/gemm/gemm_instance_builder.py +++ b/projects/composablekernel/tile_engine/ops/gemm/gemm_instance_builder.py @@ -780,20 +780,6 @@ def populate_trait_config(self, trait_combo): instance_code += """ static constexpr bool Preshuffle = false;""" - if self.kernel_name_prefix == "gemm_aquant": - instance_code += f""" - static constexpr bool APreshuffleQuant = {"true" if persistent_or_preshuffle_quant in [True, "true"] else "false"}; - static constexpr bool BPreshuffleQuant = false; - static constexpr bool PreshuffleB = false; - static constexpr ck_tile::index_t GroupSizeK = {self.group_size_k};""" - - elif self.kernel_name_prefix == "gemm_bquant": - instance_code += f""" - static constexpr bool APreshuffleQuant = false; - static constexpr bool BPreshuffleQuant = {"true" if persistent_or_preshuffle_quant in [True, "true"] else "false"}; - static constexpr bool PreshuffleB = false; - static constexpr ck_tile::index_t GroupSizeK = {self.group_size_k};""" - return instance_code def populate_initialization(self, base_pipeline_map, pipeline): @@ -1799,96 +1785,6 @@ def populate_cshuffle_gemm_aquant(self): using GemmEpilogue = ck_tile::CShuffleEpilogue;""" return instance_code - def _populate_launch_bquant( - self, - scheduler_type_map, - scheduler, - pipeline_impl_map, - pipeline, - epilogue, - k_block_per_cu, - ): - instance_code = """ - - // Launch function - static float launch(const ck_tile::QuantGemmHostArgs& args, const ck_tile::stream_config& stream) { - - // Hot loop detection - const ck_tile::index_t K_split = ck_tile::integer_least_multiple(args.K, TileShape::kK); - const ck_tile::index_t num_loop = TilePartitioner::GetLoopNum(K_split); - const bool has_hot_loop = BaseGemmPipeline::BlockHasHotloop(num_loop); - const ck_tile::TailNumber tail_num = BaseGemmPipeline::GetBlockLoopTailNum(num_loop);""" - - instance_code += f""" - - const auto Run = [&](const auto has_hot_loop_, const auto tail_number_) {{ - constexpr bool has_hot_loop_v = has_hot_loop_.value; - constexpr auto tail_number_v = tail_number_.value; - - // Full pipeline problem with hot loop and tail number - using PipelineProblem = ck_tile::GemmBQuantPipelineProblem< - ADataType, - BDataType, - BQDataType, - AccDataType, - TileShape, - Traits, - QuantGroupSize, - TransposeC, - BDataType, - {scheduler_type_map.get(scheduler)}, - has_hot_loop_v, - tail_number_v>; - - using GemmPipeline = {pipeline_impl_map.get(pipeline)};""" - - instance_code += """ - - // Epilogue""" - if epilogue == "cshuffle": - instance_code += self.populate_cshuffle_gemm_bquant() - else: - instance_code += self.populate_default_gemm_bquant() - - instance_code += f""" - - // Kernel type - using Kernel = ck_tile::QuantGemmKernel< - TilePartitioner, GemmPipeline, GemmEpilogue, - ck_tile::QuantType::BQuantGrouped>; - - // Kernel arguments - auto kargs = Kernel::MakeKernelArgs(args); - - if (!Kernel::IsSupportedArgument(kargs)) {{ - throw std::runtime_error("Unsupported kernel arguments; skipping GEMM launch."); - }} - - // Get grid and block sizes - const dim3 grids = Kernel::GridSize(args.M, args.N, args.k_batch); - const dim3 blocks = Kernel::BlockSize(); - - if(stream.log_level_ > 0) {{ - std::cout << "Launching kernel: " << Kernel::GetName() << '\\n' - << "grid: {{" << grids.x << ", " << grids.y << ", " << grids.z << "}}" - << ", blocks: {{" << blocks.x << ", " << blocks.y << ", " << blocks.z << "}}" - << std::endl; - }} - - // Launch kernel - constexpr int kBlockPerCu = {k_block_per_cu}; - float ave_time = ck_tile::launch_kernel( - stream, - ck_tile::make_kernel(Kernel{{}}, grids, blocks, 0, kargs)); - - return ave_time; - }}; - return BaseGemmPipeline::TailHandler(Run, has_hot_loop, tail_num); - }} -}}; -""" - return instance_code - def populate_default_gemm_bquant(self): instance_code = """ using EpilogueProblem = ck_tile::DefaultGemm2DEpilogueProblem<