"...composable_kernel_rocm.git" did not exist on "a65ef9030880d51dd159e4d23f1dc6093b17651c"
Commit 69d6660c authored by Adam Osewski's avatar Adam Osewski
Browse files

A/B smem pack size taken from WarpGemm attributes

parent e0d67738
...@@ -170,21 +170,17 @@ struct UniversalGemmPipelineAgBgCrPolicy ...@@ -170,21 +170,17 @@ struct UniversalGemmPipelineAgBgCrPolicy
template <typename Problem> template <typename Problem>
CK_TILE_HOST_DEVICE static constexpr auto GetSmemPackA() CK_TILE_HOST_DEVICE static constexpr auto GetSmemPackA()
{ {
using ADataType = remove_cvref_t<typename Problem::ADataType>; using BlockGemm = decltype(GetBlockGemm<Problem>());
constexpr index_t KPack = BlockGemm::Traits::KPack;
constexpr index_t MPerBlock = Problem::BlockGemmShape::kM; return KPack;
// TODO: this not alwyas has to be ture, sometimes we may want different KPack value.
return GetGlobalVectorLoadSize<Problem, ADataType, MPerBlock>();
} }
template <typename Problem> template <typename Problem>
CK_TILE_HOST_DEVICE static constexpr auto GetSmemPackB() CK_TILE_HOST_DEVICE static constexpr auto GetSmemPackB()
{ {
using BDataType = remove_cvref_t<typename Problem::BDataType>; using BlockGemm = decltype(GetBlockGemm<Problem>());
constexpr index_t KPack = BlockGemm::Traits::KPack;
constexpr index_t NPerBlock = Problem::BlockGemmShape::kN; return KPack;
// TODO: this not alwyas has to be ture, sometimes we may want different KPack value.
return GetGlobalVectorLoadSize<Problem, BDataType, NPerBlock>();
} }
template <typename Problem> template <typename Problem>
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment