From 3fab0cc57f59dcbf5bce13c9f9eb363bf88a2f9e Mon Sep 17 00:00:00 2001 From: Sami Aario Date: Thu, 8 Jan 2026 09:13:09 +0000 Subject: [PATCH] fixup! Add DstDataType as a template parameter to load_tile_with_elementwise, and use it for type conversion --- .../gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v3.hpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v3.hpp b/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v3.hpp index 5bff1800ef..79c4faaaa2 100644 --- a/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v3.hpp +++ b/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_v3.hpp @@ -498,7 +498,7 @@ struct GemmPipelineAgBgCrCompV3 : public BaseGemmPipelineAgBgCrCompV3 // LDS write 0 if constexpr(is_a_col_major && !is_a_load_tr_v()) { - auto a_shuffle_tmp = make_static_distributed_tensor( + auto a_shuffle_tmp = make_static_distributed_tensor( Policy::template MakeShuffledARegTileDistribution()); transpose_tile2d(a_shuffle_tmp, elementwise_As_res); Base::LocalPrefill(a_copy_lds_window, a_shuffle_tmp); @@ -509,7 +509,7 @@ struct GemmPipelineAgBgCrCompV3 : public BaseGemmPipelineAgBgCrCompV3 } if constexpr(is_b_row_major && !is_b_load_tr_v()) { - auto b_shuffle_tmp = make_static_distributed_tensor( + auto b_shuffle_tmp = make_static_distributed_tensor( Policy::template MakeShuffledBRegTileDistribution()); transpose_tile2d(b_shuffle_tmp, elementwise_Bs_res); Base::LocalPrefill(b_copy_lds_window, b_shuffle_tmp); @@ -543,7 +543,7 @@ struct GemmPipelineAgBgCrCompV3 : public BaseGemmPipelineAgBgCrCompV3 if constexpr(is_a_col_major && !is_a_load_tr_v()) { - auto a_shuffle_tmp = make_static_distributed_tensor( + auto a_shuffle_tmp = make_static_distributed_tensor( Policy::template MakeShuffledARegTileDistribution()); transpose_tile2d(a_shuffle_tmp, elementwise_As_res); Base::LocalPrefill(a_copy_lds_window, a_shuffle_tmp); @@ -554,7 +554,7 @@ struct GemmPipelineAgBgCrCompV3 : public BaseGemmPipelineAgBgCrCompV3 } if constexpr(is_b_row_major && !is_b_load_tr_v()) { - auto b_shuffle_tmp = make_static_distributed_tensor( + auto b_shuffle_tmp = make_static_distributed_tensor( Policy::template MakeShuffledBRegTileDistribution()); transpose_tile2d(b_shuffle_tmp, elementwise_Bs_res); Base::LocalPrefill(b_copy_lds_window, b_shuffle_tmp); @@ -598,7 +598,7 @@ struct GemmPipelineAgBgCrCompV3 : public BaseGemmPipelineAgBgCrCompV3 if constexpr(is_a_col_major && !is_a_load_tr_v()) { - auto a_shuffle_tmp = make_static_distributed_tensor( + auto a_shuffle_tmp = make_static_distributed_tensor( Policy::template MakeShuffledARegTileDistribution()); transpose_tile2d(a_shuffle_tmp, elementwise_As_res); Base::LocalPrefill(a_copy_lds_window, a_shuffle_tmp); @@ -609,7 +609,7 @@ struct GemmPipelineAgBgCrCompV3 : public BaseGemmPipelineAgBgCrCompV3 } if constexpr(is_b_row_major && !is_b_load_tr_v()) { - auto b_shuffle_tmp = make_static_distributed_tensor( + auto b_shuffle_tmp = make_static_distributed_tensor( Policy::template MakeShuffledBRegTileDistribution()); transpose_tile2d(b_shuffle_tmp, elementwise_Bs_res); Base::LocalPrefill(b_copy_lds_window, b_shuffle_tmp);