From d18e076d8a22598087a476dd12da3c956271bb15 Mon Sep 17 00:00:00 2001 From: Aleksander Dudek Date: Wed, 19 Nov 2025 08:10:58 -0600 Subject: [PATCH] [CK_TILE] Vector stores c col layout part4 --- .../ops/epilogue/cshuffle_epilogue.hpp | 28 +++++-------------- 1 file changed, 7 insertions(+), 21 deletions(-) diff --git a/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp b/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp index c6ff341b60..4da048de5f 100644 --- a/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp +++ b/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp @@ -191,7 +191,7 @@ struct CShuffleEpilogue return VectorSizeC; } constexpr index_t max_vector_size = 16; - if constexpr(std::is_same_v) + if constexpr(IsERowMajor) { return std::min(static_cast(NPerIteration), static_cast(max_vector_size / sizeof(ODataType))); @@ -251,7 +251,7 @@ struct CShuffleEpilogue else { constexpr index_t num_xdl_shuffles = GetVectorSizeC() / elem_per_thread; - if constexpr(std::is_same_v) + if constexpr(IsERowMajor) { static_assert((kMPerBlock % (MPerXdl * MWave) == 0) && (kMPerBlock % num_xdl_shuffles == 0), @@ -318,24 +318,10 @@ struct CShuffleEpilogue template CK_TILE_HOST_DEVICE static constexpr auto MakeLdsBlockDescriptor() { - // N is contiguous dimension - if constexpr(std::is_same_v) - { - return make_naive_tensor_descriptor( - make_tuple(number{}, number{}), - make_tuple(number{}, number<1>{})); - } - // M is contiguous dimension - else if constexpr(std::is_same_v) - { - return make_naive_tensor_descriptor( - make_tuple(number{}, number{}), - make_tuple(number{}, number<1>{})); - } - else - { - static_assert(false, "Unsupported ELayout!"); - } + return make_naive_tensor_descriptor( + make_tuple(number{}, number{}), + make_tuple(number{}, number<1>{})); + } CK_TILE_DEVICE static constexpr auto MakeLdsDistributionEncode() @@ -688,7 +674,7 @@ struct CShuffleEpilogue constexpr index_t num_access = SFC::get_num_of_access(); // TODO: Add support for Col Major Output Layout - CShuffle Epilogue - // static_assert(std::is_same_v, + // static_assert(IsERowMajor, // "Currently, the CShuffle Epilogue only supports the Row Major Output // layout"); static_assert(GetVectorSizeC() > 1, "VectorSizeC is not greater than 1!");