77#include < concepts>
88#include < utility>
99#include " ck_tile/core/utility/type_traits.hpp"
10+ #include " ck_tile/core/arch/arch.hpp"
1011
1112namespace ck_tile ::builder {
1213
@@ -51,6 +52,10 @@ concept HasTupleSize = requires {
5152 { std::tuple_size<T>::value } -> std::convertible_to<size_t >;
5253};
5354
55+ // Helper for dependent static_assert
56+ template <typename >
57+ constexpr bool always_false = false ;
58+
5459// Get compile-time size of a range
5560template <typename Range>
5661constexpr size_t get_range_size ()
@@ -63,6 +68,10 @@ constexpr size_t get_range_size()
6368 {
6469 return std::tuple_size_v<Range>;
6570 }
71+ else
72+ {
73+ static_assert (always_false<Range>, " Unsupported type of range object." );
74+ }
6675}
6776
6877// Fold expression implementation for product calculation
@@ -120,8 +129,7 @@ constexpr auto get_mn_coverage()
120129template <size_t DataTypeSize>
121130constexpr auto get_data_max_vec_size ()
122131{
123- // this is arch specific - but all current gfx9 has same value
124- constexpr auto max_vec_inst_size_bytes = 16 ;
132+ constexpr auto max_vec_inst_size_bytes = get_max_mem_vec_inst_width ();
125133 static_assert (max_vec_inst_size_bytes % DataTypeSize == 0 ,
126134 " The max vec instruction size is not a multiple of given data type size." );
127135 return max_vec_inst_size_bytes / DataTypeSize;
@@ -161,11 +169,12 @@ concept ThreadsCoverCTile = requires {
161169 CBlockTransfer.scalar_per_vector ) == 0 ;
162170};
163171
172+ template <size_t Value>
173+ concept IsPowerOf2 = (Value > 0 ) && ((Value & (Value - 1 )) == 0 );
174+
164175template <size_t ScalarPerVec, size_t DataTypeSize>
165- concept IsVectorSizeValid = requires {
166- requires ((ScalarPerVec & (ScalarPerVec - 1 )) == 0 ) &&
167- ScalarPerVec <= detail::get_data_max_vec_size<DataTypeSize>();
168- };
176+ concept IsVectorSizeValid =
177+ IsPowerOf2<ScalarPerVec> && (ScalarPerVec <= detail::get_data_max_vec_size<DataTypeSize>());
169178
170179// Composite concept for input block transfer validation (A)
171180// Includes all validations: vector transfer limits, access order, cluster size,
0 commit comments