@@ -18,26 +18,24 @@ typedef uint64_t v4u64_t __attribute__((ext_vector_type(4)));
1818// Check if a given type is a vector type or not. Used in submitAndCheck to
1919// branch the check: we need element-wise comparison for vector types. Default
2020// case: T is not a vector type.
21- template <typename T>
22- struct is_vector : std::false_type {};
21+ template <typename T> struct is_vector : std::false_type {};
2322// Specialization for vector types. If T has
2423// __attribute__((ext_vector_type(N))), then it's a vector type.
2524template <typename T, std::size_t N>
2625struct is_vector <T __attribute__ ((ext_vector_type(N)))> : std::true_type {};
27- template <typename T>
28- inline constexpr bool is_vector_v = is_vector<T>::value;
26+ template <typename T> inline constexpr bool is_vector_v = is_vector<T>::value;
2927
3028// Get the length of a vector type. Used in submitAndCheck to iterate over the
3129// elements of the vector type. Default case: length is 1.
3230template <typename T>
3331struct vector_length {
34- static constexpr std::size_t value = 1 ;
32+ static constexpr std::size_t value = 1 ;
3533};
3634// Specialization for vector types. If T has
3735// __attribute__((ext_vector_type(N))), then the length is N.
3836template <typename T, std::size_t N>
3937struct vector_length <T __attribute__ ((ext_vector_type(N)))> {
40- static constexpr std::size_t value = N;
38+ static constexpr std::size_t value = N;
4139};
4240template <typename T>
4341inline constexpr std::size_t vector_length_v = vector_length<T>::value;
@@ -50,7 +48,7 @@ struct element_type;
5048// Specialization for vector types. If T has __attribute__((ext_vector_type(N))), return T.
5149template <typename T, int N>
5250struct element_type <T __attribute__ ((ext_vector_type(N)))> {
53- using type = T;
51+ using type = T;
5452};
5553// Helper alias template.
5654template <typename T>
@@ -64,41 +62,41 @@ struct TypeList {};
6462// Function to trigger llvm.scmp/ucmp.
6563template <typename RetTy, typename ArgTy>
6664void compare (RetTy &res, ArgTy x, ArgTy y) {
67- auto lessOrEq = (x <= y);
68- auto lessThan = (x < y);
69- res = lessOrEq ? (lessThan ? RetTy (-1 ) : RetTy (0 )) : RetTy (1 );
65+ auto lessOrEq = (x <= y);
66+ auto lessThan = (x < y);
67+ res = lessOrEq ? (lessThan ? RetTy (-1 ) : RetTy (0 )) : RetTy (1 );
7068}
7169
7270// Function to submit kernel and check device result with host result.
7371template <typename RetTy, typename ArgTy>
7472void submitAndCheck (sycl::queue &q, ArgTy x, ArgTy y) {
75- RetTy res;
76- {
77- sycl::buffer<RetTy, 1 > res_b{&res, 1 };
78- q.submit ([&](sycl::handler &cgh) {
79- sycl::accessor acc{res_b, cgh, sycl::write_only};
80- cgh.single_task <>([=] {
81- RetTy tmp;
82- compare<RetTy, ArgTy>(tmp, x, y);
83- acc[0 ] = tmp;
84- });
85- });
86- }
87- RetTy expectedRes;
88- compare<RetTy, ArgTy>(expectedRes, x, y);
89- if constexpr (is_vector_v<RetTy>) {
90- for (int i = 0 ; i < vector_length_v<RetTy>; ++i) {
91- assert (res[i] == expectedRes[i]);
92- }
93- } else {
94- assert (res == expectedRes);
73+ RetTy res;
74+ {
75+ sycl::buffer<RetTy, 1 > res_b{&res, 1 };
76+ q.submit ([&](sycl::handler &cgh) {
77+ sycl::accessor acc{res_b, cgh, sycl::write_only};
78+ cgh.single_task <>([=] {
79+ RetTy tmp;
80+ compare<RetTy, ArgTy>(tmp, x, y);
81+ acc[0 ] = tmp;
82+ });
83+ });
84+ }
85+ RetTy expectedRes;
86+ compare<RetTy, ArgTy>(expectedRes, x, y);
87+ if constexpr (is_vector_v<RetTy>) {
88+ for (int i = 0 ; i < vector_length_v<RetTy>; ++i) {
89+ assert (res[i] == expectedRes[i]);
9590 }
91+ } else {
92+ assert (res == expectedRes);
93+ }
9694}
9795
9896// Helper to call submitAndCheck for each combination.
9997template <typename RetTypes, typename ArgTypes>
10098void submitAndCheckCombination (sycl::queue &q, int x, int y) {
101- submitAndCheck<RetTypes, ArgTypes>(q, x, y);
99+ submitAndCheck<RetTypes, ArgTypes>(q, x, y);
102100}
103101
104102// Function to generate all the combinations possible with the two type lists.
@@ -110,8 +108,8 @@ void submitAndCheckCombination(sycl::queue &q, int x, int y) {
110108// Recursive case to generate combinations.
111109template <typename RetType, typename ... RetTypes, typename ... ArgTypes>
112110void submitCombinations (sycl::queue &q, int x, int y, TypeList<RetType, RetTypes...>, TypeList<ArgTypes...>) {
113- (submitAndCheckCombination<RetType, ArgTypes>(q, x, y), ...);
114- submitCombinations (q, x, y, TypeList<RetTypes...>{}, TypeList<ArgTypes...>{});
111+ (submitAndCheckCombination<RetType, ArgTypes>(q, x, y), ...);
112+ submitCombinations (q, x, y, TypeList<RetTypes...>{}, TypeList<ArgTypes...>{});
115113}
116114// Base case to stop recursion.
117115template <typename ... ArgTypes>
@@ -125,26 +123,29 @@ void submitCombinations(sycl::queue &, int, int, TypeList<>, TypeList<ArgTypes..
125123// Recursive case to generate combinations.
126124template <typename ArgType, typename ... ArgTypes>
127125void submitVecCombinations (sycl::queue &q, int x, int y, TypeList<ArgType, ArgTypes...>) {
128- // Use signed types for return type, as it may return -1.
129- using ElemType = std::make_signed_t <element_type_t <ArgType>>;
130- using RetType = ElemType __attribute__ ((ext_vector_type (vector_length_v<ArgType>)));
131- submitAndCheckCombination<RetType, ArgType>(q, x, y);
132- submitVecCombinations (q, x, y, TypeList<ArgTypes...>{});
126+ // Use signed types for return type, as it may return -1.
127+ using ElemType = std::make_signed_t <element_type_t <ArgType>>;
128+ using RetType =
129+ ElemType __attribute__ ((ext_vector_type (vector_length_v<ArgType>)));
130+ submitAndCheckCombination<RetType, ArgType>(q, x, y);
131+ submitVecCombinations (q, x, y, TypeList<ArgTypes...>{});
133132}
134133// Base case to stop recursion.
135134void submitVecCombinations (sycl::queue &, int , int , TypeList<>) {}
136135
137136int main (int argc, char **argv) {
138- sycl::queue q;
139- // RetTypes includes only signed types because it may return -1.
140- using RetTypes = TypeList<int8_t , int16_t , int32_t , int64_t >;
141- using ArgTypes = TypeList<int8_t , int16_t , int32_t , int64_t , uint8_t , uint16_t , uint32_t , uint64_t >;
142- submitCombinations (q, 50 , 49 , RetTypes{}, ArgTypes{});
143- submitCombinations (q, 50 , 50 , RetTypes{}, ArgTypes{});
144- submitCombinations (q, 50 , 51 , RetTypes{}, ArgTypes{});
145- using VecTypes = TypeList<v4i8_t , v4i16_t , v4i32_t , v4i64_t , v4u8_t , v4u16_t , v4u32_t , v4u64_t >;
146- submitVecCombinations (q, 50 , 49 , VecTypes{});
147- submitVecCombinations (q, 50 , 50 , VecTypes{});
148- submitVecCombinations (q, 50 , 51 , VecTypes{});
149- return 0 ;
137+ sycl::queue q;
138+ // RetTypes includes only signed types because it may return -1.
139+ using RetTypes = TypeList<int8_t , int16_t , int32_t , int64_t >;
140+ using ArgTypes = TypeList<int8_t , int16_t , int32_t , int64_t , uint8_t ,
141+ uint16_t , uint32_t , uint64_t >;
142+ submitCombinations (q, 50 , 49 , RetTypes{}, ArgTypes{});
143+ submitCombinations (q, 50 , 50 , RetTypes{}, ArgTypes{});
144+ submitCombinations (q, 50 , 51 , RetTypes{}, ArgTypes{});
145+ using VecTypes = TypeList<v4i8_t , v4i16_t , v4i32_t , v4i64_t , v4u8_t , v4u16_t ,
146+ v4u32_t , v4u64_t >;
147+ submitVecCombinations (q, 50 , 49 , VecTypes{});
148+ submitVecCombinations (q, 50 , 50 , VecTypes{});
149+ submitVecCombinations (q, 50 , 51 , VecTypes{});
150+ return 0 ;
150151}
0 commit comments