diff --git a/include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp b/include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp index 0e20233fab..079d3972d1 100644 --- a/include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp +++ b/include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp @@ -37,9 +37,7 @@ struct GemmHostArgs index_t K_, index_t stride_A_, index_t stride_B_, - index_t stride_E_, - bool skip_a_lds_ = false, - bool skip_b_lds_ = false) + index_t stride_E_) : a_ptr(a_ptr_), b_ptr(b_ptr_), e_ptr(e_ptr_), @@ -49,9 +47,7 @@ struct GemmHostArgs stride_A(stride_A_), stride_B(stride_B_), stride_E(stride_E_), - k_batch(k_batch_), - skip_a_lds(skip_a_lds_), - skip_b_lds(skip_b_lds_) + k_batch(k_batch_) { } @@ -76,8 +72,6 @@ struct GemmHostArgs }; index_t k_batch; - bool skip_a_lds; - bool skip_b_lds; }; template diff --git a/include/ck_tile/ops/gemm/kernel/grouped_gemm_kernel.hpp b/include/ck_tile/ops/gemm/kernel/grouped_gemm_kernel.hpp index 338517c653..921ea11720 100644 --- a/include/ck_tile/ops/gemm/kernel/grouped_gemm_kernel.hpp +++ b/include/ck_tile/ops/gemm/kernel/grouped_gemm_kernel.hpp @@ -34,9 +34,7 @@ struct GroupedGemmHostArgs index_t K_, index_t stride_A_, index_t stride_B_, - index_t stride_E_, - bool skip_a_lds_, - bool skip_b_lds_) + index_t stride_E_) : a_ptr(a_ptr_), b_ptr(b_ptr_), e_ptr(e_ptr_), @@ -46,9 +44,7 @@ struct GroupedGemmHostArgs stride_A(stride_A_), stride_B(stride_B_), stride_E(stride_E_), - k_batch(k_batch_), - skip_a_lds(skip_a_lds_), - skip_b_lds(skip_b_lds_) + k_batch(k_batch_) { } @@ -73,8 +69,6 @@ struct GroupedGemmHostArgs }; index_t k_batch; - bool skip_a_lds; - bool skip_b_lds; }; struct GemmTransKernelArg @@ -233,9 +227,7 @@ struct GroupedGemmKernel {stride_b}, {/*stride_ds*/}, stride_e, - gemm_descs[i].k_batch, - gemm_descs[i].skip_a_lds, - gemm_descs[i].skip_b_lds}; + gemm_descs[i].k_batch}; gemm_kernel_args_.emplace_back(std::move(karg), block_start, block_end); } diff --git a/include/ck_tile/ops/gemm/pipeline/tile_gemm_traits.hpp b/include/ck_tile/ops/gemm/pipeline/tile_gemm_traits.hpp index 9138d6f8db..e59a08b5cd 100644 --- a/include/ck_tile/ops/gemm/pipeline/tile_gemm_traits.hpp +++ b/include/ck_tile/ops/gemm/pipeline/tile_gemm_traits.hpp @@ -34,6 +34,12 @@ struct TileGemmTraits static constexpr index_t NumWaveGroups = NumWaveGroups_; }; +/** + * @brief Traits for a universal GEMM pipeline. + * + * @details This struct holds compile-time parameters that configure the behavior of a + * general-purpose tile-based GEMM kernel + */ template + /// registers; if false, global memory -> LDS -> registers. + static constexpr bool SkipALds = SkipALds_; + /// @brief SkipBLds_ Parameterize the data flow for matrix B: if true, global memory -> + /// registers; if false, global memory -> LDS -> registers. static constexpr bool SkipBLds = SkipBLds_; static constexpr bool UseStructuredSparsity = UseStructuredSparsity_; static constexpr bool UsePersistentKernel = UsePersistentKernel_; static constexpr index_t NumWaveGroups = NumWaveGroups_; static constexpr bool Preshuffle = Preshuffle_; - - static constexpr bool kSkipALds = SkipALds_; - static constexpr bool kSkipBLds = SkipBLds_; }; template ; - using Traits = ck_tile::TileGemmTraits; - using GemmUniversalTraits = - ck_tile::TileGemmUniversalTraits; + using Traits = ck_tile::TileGemmTraits; + using GemmUniversalTraits = ck_tile::TileGemmUniversalTraits; using GemmPipelineProblem = ck_tile::GemmPipelineProblem; @@ -444,18 +441,8 @@ class TestCkTileGroupedGemm : public ::testing::Test // TODO add support for kbatch > 1 static constexpr ck_tile::index_t k_batch = 1; - gemm_descs.push_back({p_a, - p_b, - p_c, - k_batch, - M, - N, - K, - stride_As[i], - stride_Bs[i], - stride_Cs[i], - GroupedGemKernelParam::SkipALds, - GroupedGemKernelParam::SkipBLds}); + gemm_descs.push_back( + {p_a, p_b, p_c, k_batch, M, N, K, stride_As[i], stride_Bs[i], stride_Cs[i]}); } ck_tile::DeviceMem gemm_workspace;