diff --git a/include/ck_tile/host/mx_processing.hpp b/include/ck_tile/host/mx_processing.hpp index a11fa18481..c15df6f509 100644 --- a/include/ck_tile/host/mx_processing.hpp +++ b/include/ck_tile/host/mx_processing.hpp @@ -68,7 +68,11 @@ auto preShuffleScale(ck_tile::HostTensor& src, const bool kLast) HostTensor shuffled(HostTensorDescriptor({static_cast(MNPadded * K)}, {static_cast(1)})); - assert(K % (KXdlPack * XdlKThread) == 0); + if(K % (KXdlPack * XdlKThread) != 0) + { + throw std::runtime_error("wrong! K must be a multiple of (KXdlPack * XdlKThread)"); + } + const index_t K0 = K / KXdlPack / XdlKThread; for(index_t n = 0; n < MNPadded; ++n) @@ -120,7 +124,10 @@ auto preShuffleScalePermuteN(const HostTensor& src, const bool kLast) HostTensor shuffled(HostTensorDescriptor({static_cast(MNPadded * K)}, {static_cast(1)})); - assert(K % (KXdlPack * XdlKThread) == 0); + if(K % (KXdlPack * XdlKThread) != 0) + { + throw std::runtime_error("wrong! K must be a multiple of (KXdlPack * XdlKThread)"); + } const index_t K0 = K / KXdlPack / XdlKThread; for(index_t n = 0; n < MNPadded; ++n) diff --git a/include/ck_tile/host/tensor_shuffle_utils.hpp b/include/ck_tile/host/tensor_shuffle_utils.hpp index e597753cc4..63e74cdb23 100644 --- a/include/ck_tile/host/tensor_shuffle_utils.hpp +++ b/include/ck_tile/host/tensor_shuffle_utils.hpp @@ -164,9 +164,10 @@ auto shuffle_b_permuteN(const ck_tile::HostTensor& t, number) { assert(t.get_lengths().size() == 2); - int n_ = t.get_lengths()[1]; - int k_ = t.get_lengths()[0]; - int NRepeat = gemmConfig.N_Tile / gemmConfig.N_Warp_Tile / gemmConfig.N_Warp; + int n_ = t.get_lengths()[1]; + int k_ = t.get_lengths()[0]; + constexpr ck_tile::index_t NRepeat = + GemmConfig::N_Tile / GemmConfig::N_Warp_Tile / GemmConfig::N_Warp; if(ck_tile::is_gfx12_supported()) { constexpr int divisor = 2; @@ -185,7 +186,8 @@ auto shuffle_b_permuteN(const ck_tile::HostTensor& t, } else { - assert(NRepeat % BlockedXDLNPerWarp == 0); + static_assert(NRepeat % BlockedXDLNPerWarp == 0, + "wrong! NRepeat must be a multiple of BlockedXDLNPerWarp"); constexpr int KLane = ck_tile::get_warp_size() / GemmConfig::N_Warp_Tile; constexpr int ItemsPerAccess = std::min(detail::b_contiguous_items_per_access::value,