Skip to content

Commit 1105d05

Browse files
steffenlarsensteffenlarsen
committed
Geometric marray functions
Signed-off-by: Larsen, Steffen <[email protected]>
1 parent 7fb0f90 commit 1105d05

File tree

3 files changed

+136
-292
lines changed

3 files changed

+136
-292
lines changed

‎sycl/include/sycl/builtins_marray.hpp

Lines changed: 3 additions & 256 deletions
Original file line numberDiff line numberDiff line change
@@ -34,198 +34,9 @@ namespace __sycl_std = __host_std;
3434
/* ----------------- 4.13.3 Math functions. ---------------------------------*/
3535

3636
template <typename T, size_t N>
37-
inline __SYCL_ALWAYS_INLINE
38-
std::enable_if_t<detail::is_sgenfloat<T>::value, marray<int, N>>
39-
ilogb(marray<T, N> x) __NOEXC {
40-
marray<int, N> res;
41-
for (size_t i = 0; i < N / 2; i++) {
42-
vec<int, 2> partial_res =
43-
__sycl_std::__invoke_ilogb<vec<int, 2>>(detail::to_vec2(x, i * 2));
44-
std::memcpy(&res[i * 2], &partial_res, sizeof(vec<int, 2>));
45-
}
46-
if (N % 2) {
47-
res[N - 1] = __sycl_std::__invoke_ilogb<int>(x[N - 1]);
48-
}
49-
return res;
50-
}
51-
52-
#define __SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD(NAME) \
53-
template <typename T, size_t N> \
54-
inline __SYCL_ALWAYS_INLINE \
55-
std::enable_if_t<detail::is_sgenfloat<T>::value, marray<T, N>> \
56-
NAME(marray<T, N> x, T y) __NOEXC { \
57-
marray<T, N> res; \
58-
sycl::vec<T, 2> y_vec{y, y}; \
59-
for (size_t i = 0; i < N / 2; i++) { \
60-
auto partial_res = __sycl_std::__invoke_##NAME<vec<T, 2>>( \
61-
detail::to_vec2(x, i * 2), y_vec); \
62-
std::memcpy(&res[i * 2], &partial_res, sizeof(vec<T, 2>)); \
63-
} \
64-
if (N % 2) { \
65-
res[N - 1] = __sycl_std::__invoke_##NAME<T>(x[N - 1], y_vec[0]); \
66-
} \
67-
return res; \
68-
}
69-
70-
__SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD(fmax)
71-
// clang-format off
72-
__SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD(fmin)
73-
74-
#undef __SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD
75-
76-
template <typename T, size_t N>
77-
inline __SYCL_ALWAYS_INLINE
78-
std::enable_if_t<detail::is_sgenfloat<T>::value, marray<T, N>>
79-
ldexp(marray<T, N> x, marray<int, N> k) __NOEXC {
80-
// clang-format on
81-
marray<T, N> res;
82-
for (size_t i = 0; i < N; i++) {
83-
res[i] = __sycl_std::__invoke_ldexp<T>(x[i], k[i]);
84-
}
85-
return res;
86-
}
87-
88-
template <typename T, size_t N>
89-
inline __SYCL_ALWAYS_INLINE
90-
std::enable_if_t<detail::is_sgenfloat<T>::value, marray<T, N>>
91-
ldexp(marray<T, N> x, int k) __NOEXC {
92-
marray<T, N> res;
93-
for (size_t i = 0; i < N; i++) {
94-
res[i] = __sycl_std::__invoke_ldexp<T>(x[i], k);
95-
}
96-
return res;
97-
}
98-
99-
#define __SYCL_MATH_FUNCTION_2_GENINT_Y_OVERLOAD_IMPL(NAME) \
100-
marray<T, N> res; \
101-
for (size_t i = 0; i < N; i++) { \
102-
res[i] = __sycl_std::__invoke_##NAME<T>(x[i], y[i]); \
103-
} \
104-
return res;
105-
106-
template <typename T, size_t N>
107-
inline __SYCL_ALWAYS_INLINE
108-
std::enable_if_t<detail::is_sgenfloat<T>::value, marray<T, N>>
109-
pown(marray<T, N> x, marray<int, N> y) __NOEXC {
110-
__SYCL_MATH_FUNCTION_2_GENINT_Y_OVERLOAD_IMPL(pown)
111-
}
112-
113-
template <typename T, size_t N>
114-
inline __SYCL_ALWAYS_INLINE
115-
std::enable_if_t<detail::is_sgenfloat<T>::value, marray<T, N>>
116-
rootn(marray<T, N> x, marray<int, N> y) __NOEXC {
117-
__SYCL_MATH_FUNCTION_2_GENINT_Y_OVERLOAD_IMPL(rootn)
118-
}
119-
120-
#undef __SYCL_MATH_FUNCTION_2_GENINT_Y_OVERLOAD_IMPL
121-
122-
#define __SYCL_MATH_FUNCTION_2_INT_Y_OVERLOAD_IMPL(NAME) \
123-
marray<T, N> res; \
124-
for (size_t i = 0; i < N; i++) { \
125-
res[i] = __sycl_std::__invoke_##NAME<T>(x[i], y); \
126-
} \
127-
return res;
128-
129-
template <typename T, size_t N>
130-
inline __SYCL_ALWAYS_INLINE
131-
std::enable_if_t<detail::is_sgenfloat<T>::value, marray<T, N>>
132-
pown(marray<T, N> x, int y) __NOEXC {
133-
__SYCL_MATH_FUNCTION_2_INT_Y_OVERLOAD_IMPL(pown)
134-
}
135-
136-
template <typename T, size_t N>
137-
inline __SYCL_ALWAYS_INLINE
138-
std::enable_if_t<detail::is_sgenfloat<T>::value, marray<T, N>>
139-
rootn(marray<T, N> x,
140-
int y) __NOEXC{__SYCL_MATH_FUNCTION_2_INT_Y_OVERLOAD_IMPL(rootn)}
141-
142-
#undef __SYCL_MATH_FUNCTION_2_INT_Y_OVERLOAD_IMPL
143-
144-
// other marray math functions
145-
146-
// TODO: can be optimized in the way marray math functions above are optimized
147-
// (usage of vec<T, 2>)
148-
#define __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, ARGPTR, \
149-
...) \
150-
marray<T, N> res; \
151-
for (int j = 0; j < N; j++) { \
152-
res[j] = \
153-
NAME(__VA_ARGS__, \
154-
address_space_cast<AddressSpace, IsDecorated, \
155-
detail::marray_element_t<T2>>(&(*ARGPTR)[j])); \
156-
} \
157-
return res;
158-
159-
#define __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENFLOATPTR_OVERLOAD( \
160-
NAME, ARG1, ARG2, ...) \
161-
template <typename T, size_t N, typename T2, \
162-
access::address_space AddressSpace, access::decorated IsDecorated> \
163-
std::enable_if_t< \
164-
detail::is_svgenfloat<T>::value && \
165-
detail::is_genfloatptr_marray<T2, AddressSpace, IsDecorated>::value, \
166-
marray<T, N>> \
167-
NAME(marray<T, N> ARG1, multi_ptr<T2, AddressSpace, IsDecorated> ARG2) \
168-
__NOEXC { \
169-
__SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, ARG2, \
170-
__VA_ARGS__) \
171-
}
172-
173-
__SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENFLOATPTR_OVERLOAD(fract, x,
174-
iptr, x[j])
175-
__SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENFLOATPTR_OVERLOAD(modf, x,
176-
iptr,
177-
x[j])
178-
__SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENFLOATPTR_OVERLOAD(
179-
sincos, x, cosval, x[j])
180-
181-
#undef __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_GENFLOATPTR_OVERLOAD
182-
183-
#define __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENINTPTR_OVERLOAD( \
184-
NAME, ARG1, ARG2, ...) \
185-
template <typename T, size_t N, typename T2, \
186-
access::address_space AddressSpace, access::decorated IsDecorated> \
187-
std::enable_if_t< \
188-
detail::is_svgenfloat<T>::value && \
189-
detail::is_genintptr_marray<T2, AddressSpace, IsDecorated>::value, \
190-
marray<T, N>> \
191-
NAME(marray<T, N> ARG1, multi_ptr<T2, AddressSpace, IsDecorated> ARG2) \
192-
__NOEXC { \
193-
__SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, ARG2, \
194-
__VA_ARGS__) \
195-
}
196-
197-
__SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENINTPTR_OVERLOAD(
198-
frexp, x, exp, x[j])
199-
__SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENINTPTR_OVERLOAD(
200-
lgamma_r, x, signp, x[j])
201-
202-
#undef __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_GENINTPTR_OVERLOAD
203-
204-
#define __SYCL_MARRAY_MATH_FUNCTION_REMQUO_OVERLOAD(NAME, ...) \
205-
template <typename T, size_t N, typename T2, \
206-
access::address_space AddressSpace, access::decorated IsDecorated> \
207-
std::enable_if_t< \
208-
detail::is_svgenfloat<T>::value && \
209-
detail::is_genintptr_marray<T2, AddressSpace, IsDecorated>::value, \
210-
marray<T, N>> \
211-
NAME(marray<T, N> x, marray<T, N> y, \
212-
multi_ptr<T2, AddressSpace, IsDecorated> quo) __NOEXC { \
213-
__SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, quo, \
214-
__VA_ARGS__) \
215-
}
216-
217-
__SYCL_MARRAY_MATH_FUNCTION_REMQUO_OVERLOAD(remquo,
218-
x[j], y[j])
219-
220-
#undef __SYCL_MARRAY_MATH_FUNCTION_REMQUO_OVERLOAD
221-
222-
#undef __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL
223-
224-
template <typename T, size_t N>
225-
std::enable_if_t<
226-
detail::is_nan_type<T>::value,
227-
marray<detail::nan_return_t<T>,
228-
N>> nan(marray<T, N> nancode) __NOEXC {
37+
std::enable_if_t<detail::is_nan_type<T>::value,
38+
marray<detail::nan_return_t<T>, N>>
39+
nan(marray<T, N> nancode) __NOEXC {
22940
marray<detail::nan_return_t<T>, N> res;
23041
for (int j = 0; j < N; j++) {
23142
res[j] = nan(nancode[j]);
@@ -545,70 +356,6 @@ __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_IU_OVERLOAD(upsample, 32bit)
545356
#undef __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_UU_OVERLOAD
546357
#undef __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_OVERLOAD_IMPL
547358

548-
// marray geometric functions
549-
550-
#define __SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL(NAME, ...) \
551-
vec<detail::marray_element_t<T>, T::size()> result_v; \
552-
result_v = NAME(__VA_ARGS__); \
553-
return detail::to_marray(result_v);
554-
555-
template <typename T>
556-
std::enable_if_t<detail::is_gencrossmarray<T>::value, T> cross(T p0,
557-
T p1) __NOEXC {
558-
__SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL(cross, detail::to_vec(p0),
559-
detail::to_vec(p1))
560-
}
561-
562-
template <typename T>
563-
std::enable_if_t<detail::is_gengeomarray<T>::value, T> normalize(T p) __NOEXC {
564-
__SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL(normalize, detail::to_vec(p))
565-
}
566-
567-
template <typename T>
568-
std::enable_if_t<detail::is_gengeomarrayfloat<T>::value, T>
569-
fast_normalize(T p) __NOEXC {
570-
__SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL(fast_normalize,
571-
detail::to_vec(p))
572-
}
573-
574-
#undef __SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL
575-
576-
#define __SYCL_MARRAY_GEOMETRIC_FUNCTION_IS_GENGEOMARRAY_BINOP_OVERLOAD(NAME) \
577-
template <typename T> \
578-
std::enable_if_t<detail::is_gengeomarray<T>::value, \
579-
detail::marray_element_t<T>> \
580-
NAME(T p0, T p1) __NOEXC { \
581-
return NAME(detail::to_vec(p0), detail::to_vec(p1)); \
582-
}
583-
584-
// clang-format off
585-
__SYCL_MARRAY_GEOMETRIC_FUNCTION_IS_GENGEOMARRAY_BINOP_OVERLOAD(dot)
586-
__SYCL_MARRAY_GEOMETRIC_FUNCTION_IS_GENGEOMARRAY_BINOP_OVERLOAD(distance)
587-
// clang-format on
588-
589-
#undef __SYCL_MARRAY_GEOMETRIC_FUNCTION_IS_GENGEOMARRAY_BINOP_OVERLOAD
590-
591-
template <typename T>
592-
std::enable_if_t<detail::is_gengeomarray<T>::value, detail::marray_element_t<T>>
593-
length(T p) __NOEXC {
594-
return __sycl_std::__invoke_length<detail::marray_element_t<T>>(
595-
detail::to_vec(p));
596-
}
597-
598-
template <typename T>
599-
std::enable_if_t<detail::is_gengeomarrayfloat<T>::value,
600-
detail::marray_element_t<T>>
601-
fast_distance(T p0, T p1) __NOEXC {
602-
return fast_distance(detail::to_vec(p0), detail::to_vec(p1));
603-
}
604-
605-
template <typename T>
606-
std::enable_if_t<detail::is_gengeomarrayfloat<T>::value,
607-
detail::marray_element_t<T>>
608-
fast_length(T p) __NOEXC {
609-
return fast_length(detail::to_vec(p));
610-
}
611-
612359
// marray relational functions
613360

614361
#define __SYCL_MARRAY_RELATIONAL_FUNCTION_BINOP_OVERLOAD(NAME) \

‎sycl/include/sycl/builtins_utils.hpp

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -207,6 +207,19 @@ template <typename T, int N> struct same_size_float<vec<T, N>> {
207207
template <typename T>
208208
using same_size_float_t = typename same_size_float<T>::type;
209209

210+
template <typename T> struct int_elements {
211+
using type = int;
212+
};
213+
template <typename T, size_t N> struct int_elements<marray<T, N>> {
214+
using type = marray<typename int_elements<T>::type, N>;
215+
};
216+
template <typename T, int N> struct int_elements<vec<T, N>> {
217+
using type = vec<typename int_elements<T>::type, N>;
218+
};
219+
// TODO: Swizzle variant of this?
220+
221+
template <typename T> using int_elements_t = typename int_elements<T>::type;
222+
210223
// For upsampling we look for an integer of double the size of the specified
211224
// type.
212225
template <typename T> struct upsampled_int {

0 commit comments

Comments
 (0)