ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-18 12:00:07 +00:00
Code Issues Packages Projects Releases Wiki Activity

Merge commit '80a61afb9b6fd837f5d37a5aa55e4221ab3caa4a' into develop

Browse Source
This commit is contained in:
assistant-librarian[bot]
2025-09-11 00:41:50 +00:00
parent decf932e4c
commit 70ae0149ed
10 changed files with 343 additions and 1288 deletions
Download Patch File Download Diff File Expand all files Collapse all files

26
include/ck_tile/ops/gemm_group_quant/kernel/gemm_quant_kernel.hpp
View File

@@ -70,17 +70,21 @@ struct get_bq_data_type_or<T, Default>
using type = typename T::BQDataType;
};
template <typename T, typename Default>
struct get_preshuffle_or
{
using type = Default;
template <typename T>
concept HasStaticPreshuffleQuant = requires {
{ T::PreshuffleQuant } -> std::convertible_to<decltype(T::PreshuffleQuant)>;
};
template <typename T, typename Default>
requires requires { typename T::PreshuffleQuant; }
struct get_preshuffle_or<T, Default>
template <typename T>
struct is_quantpreshuffle_enabled
{
using type = typename T::PreshuffleQuant;
static constexpr bool value = false;
};
template <HasStaticPreshuffleQuant T>
struct is_quantpreshuffle_enabled<T>
{
static constexpr auto value = T::PreshuffleQuant;
};
} // namespace detail
@@ -198,9 +202,9 @@ struct QuantGemmKernel
using BQLayout = remove_cvref_t<
typename detail::get_bq_layout_or<GemmPipeline, typename GemmPipeline::BLayout>::type>;
static constexpr index_t kBlockSize = GemmPipeline::BlockSize;
static constexpr bool PreshuffleQuant = remove_cvref_t<
typename detail::get_preshuffle_or<GemmPipeline, std::false_type>::type>::value;
static constexpr index_t kBlockSize = GemmPipeline::BlockSize;
static constexpr bool PreshuffleQuant =
detail::is_quantpreshuffle_enabled<GemmPipeline_>::value;
using ADataType = remove_cvref_t<typename GemmPipeline::ADataType>;
using BDataType = remove_cvref_t<typename GemmPipeline::BDataType>;