diff --git a/include/ck_tile/ops/gemm/kernel/batched_gemm_kernel.hpp b/include/ck_tile/ops/gemm/kernel/batched_gemm_kernel.hpp index 1eed8343c5..9c1ce73eac 100644 --- a/include/ck_tile/ops/gemm/kernel/batched_gemm_kernel.hpp +++ b/include/ck_tile/ops/gemm/kernel/batched_gemm_kernel.hpp @@ -31,9 +31,7 @@ struct BatchedGemmHostArgs : public ck_tile::UniversalGemmHostArgs<> ck_tile::index_t batch_stride_A_, ck_tile::index_t batch_stride_B_, ck_tile::index_t batch_stride_C_, - ck_tile::index_t batch_count_, - bool skip_a_lds_, - bool skip_b_lds_) + ck_tile::index_t batch_count_) : UniversalGemmHostArgs<>({a_ptr_}, {b_ptr_}, {/*ds_ptr*/}, @@ -45,9 +43,7 @@ struct BatchedGemmHostArgs : public ck_tile::UniversalGemmHostArgs<> {stride_A_}, {stride_B_}, {/*stride_Ds_*/}, - stride_C_, - skip_a_lds_, - skip_b_lds_), + stride_C_), batch_stride_A(batch_stride_A_), batch_stride_B(batch_stride_B_), batch_stride_E(batch_stride_C_), @@ -144,9 +140,7 @@ struct BatchedGemmKernel hostArgs.stride_Bs, hostArgs.stride_Ds, hostArgs.stride_E, - hostArgs.k_batch, - hostArgs.skip_a_lds, - hostArgs.skip_b_lds}, + hostArgs.k_batch}, hostArgs.batch_stride_A, hostArgs.batch_stride_B, hostArgs.batch_stride_E, diff --git a/include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp b/include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp index 98a2b4278f..0e20233fab 100644 --- a/include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp +++ b/include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp @@ -158,9 +158,7 @@ struct GemmKernel {hostArgs.stride_A}, {hostArgs.stride_B}, {/*hostArgs.stride_Ds*/}, - hostArgs.stride_E, - hostArgs.skip_a_lds, - hostArgs.skip_b_lds)); + hostArgs.stride_E)); } CK_TILE_HOST static auto diff --git a/include/ck_tile/ops/gemm/kernel/universal_gemm_kernel.hpp b/include/ck_tile/ops/gemm/kernel/universal_gemm_kernel.hpp index b3cccd4b71..1d513faea3 100644 --- a/include/ck_tile/ops/gemm/kernel/universal_gemm_kernel.hpp +++ b/include/ck_tile/ops/gemm/kernel/universal_gemm_kernel.hpp @@ -41,9 +41,7 @@ struct UniversalGemmHostArgs const std::array& stride_As_, const std::array& stride_Bs_, const std::array& stride_Ds_, - index_t stride_E_, - bool skip_a_lds_, - bool skip_b_lds_) + index_t stride_E_) : as_ptr(as_ptr_), bs_ptr(bs_ptr_), ds_ptr(ds_ptr_), @@ -55,9 +53,7 @@ struct UniversalGemmHostArgs stride_Bs(stride_Bs_), stride_Ds(stride_Ds_), stride_E(stride_E_), - k_batch(k_batch_), - skip_a_lds(skip_a_lds_), - skip_b_lds(skip_b_lds_) + k_batch(k_batch_) { } @@ -82,8 +78,6 @@ struct UniversalGemmHostArgs }; index_t k_batch; - bool skip_a_lds; - bool skip_b_lds; }; /// @brief The GEMM kernel device arguments. @@ -117,10 +111,6 @@ struct UniversalGemmKernelArgs /// (in memory) of E tensor. index_t stride_E; index_t k_batch; - /// @brief Flag to skip loading A tensor tile into LDS. - bool skip_a_lds; - /// @brief Flag to skip loading B tensor tile into LDS. - bool skip_b_lds; }; /// @brief The Universal GEMM kernel template. @@ -292,9 +282,7 @@ struct UniversalGemmKernel hostArgs.stride_Bs, hostArgs.stride_Ds, hostArgs.stride_E, - hostArgs.k_batch, - hostArgs.skip_a_lds, - hostArgs.skip_b_lds}; + hostArgs.k_batch}; } CK_TILE_HOST_DEVICE static constexpr index_t GetSmemSize() 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 0046072056..9138d6f8db 100644 --- a/include/ck_tile/ops/gemm/pipeline/tile_gemm_traits.hpp +++ b/include/ck_tile/ops/gemm/pipeline/tile_gemm_traits.hpp @@ -67,6 +67,9 @@ struct TileGemmUniversalTraits 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 ; + TransposeC, + kSkipALds, + kSkipBLds>; using GemmPipelineProblem = ck_tile::GemmPipelineProblem;