Skip to content

Commit 4a22407

Browse files
oleksandr-pavlykoleksandr-pavlyk
committed
Set vec_sz and n_vecs in implementations of contig_impl for each support function
1 parent 1235902 commit 4a22407

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

71 files changed

+529
-183
lines changed

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/abs.hpp‎

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -135,9 +135,13 @@ sycl::event abs_contig_impl(sycl::queue &exec_q,
135135
char *res_p,
136136
const std::vector<sycl::event> &depends = {})
137137
{
138+
using resTy = typename AbsOutputType<argTy>::value_type;
139+
constexpr auto vec_sz = VecSize_v<argTy, resTy>;
140+
constexpr unsigned int n_vec = 1u;
141+
138142
return elementwise_common::unary_contig_impl<
139-
argTy, AbsOutputType, AbsContigFunctor, abs_contig_kernel>(
140-
exec_q, nelems, arg_p, res_p, depends);
143+
argTy, AbsOutputType, AbsContigFunctor, abs_contig_kernel, vec_sz,
144+
n_vec>(exec_q, nelems, arg_p, res_p, depends);
141145
}
142146

143147
template <typename fnT, typename T> struct AbsContigFactory

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp‎

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -169,9 +169,13 @@ sycl::event acos_contig_impl(sycl::queue &exec_q,
169169
char *res_p,
170170
const std::vector<sycl::event> &depends = {})
171171
{
172+
using resTy = typename AcosOutputType<argTy>::value_type;
173+
constexpr auto vec_sz = VecSize_v<argTy, resTy>;
174+
constexpr unsigned int n_vec = 1u;
175+
172176
return elementwise_common::unary_contig_impl<
173-
argTy, AcosOutputType, AcosContigFunctor, acos_contig_kernel>(
174-
exec_q, nelems, arg_p, res_p, depends);
177+
argTy, AcosOutputType, AcosContigFunctor, acos_contig_kernel, vec_sz,
178+
n_vec>(exec_q, nelems, arg_p, res_p, depends);
175179
}
176180

177181
template <typename fnT, typename T> struct AcosContigFactory

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp‎

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -196,9 +196,13 @@ sycl::event acosh_contig_impl(sycl::queue &exec_q,
196196
char *res_p,
197197
const std::vector<sycl::event> &depends = {})
198198
{
199+
using resTy = typename AcoshOutputType<argTy>::value_type;
200+
constexpr auto vec_sz = VecSize_v<argTy, resTy>;
201+
constexpr unsigned int n_vec = 1u;
202+
199203
return elementwise_common::unary_contig_impl<
200-
argTy, AcoshOutputType, AcoshContigFunctor, acosh_contig_kernel>(
201-
exec_q, nelems, arg_p, res_p, depends);
204+
argTy, AcoshOutputType, AcoshContigFunctor, acosh_contig_kernel, vec_sz,
205+
n_vec>(exec_q, nelems, arg_p, res_p, depends);
202206
}
203207

204208
template <typename fnT, typename T> struct AcoshContigFactory

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/add.hpp‎

Lines changed: 13 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -218,10 +218,14 @@ sycl::event add_contig_impl(sycl::queue &exec_q,
218218
ssize_t res_offset,
219219
const std::vector<sycl::event> &depends = {})
220220
{
221+
using resTy = typename AddOutputType<argTy1, argTy2>::value_type;
222+
constexpr auto vec_sz = VecSize_v<argTy1, argTy2, resTy>;
223+
constexpr unsigned int n_vecs = 1;
224+
221225
return elementwise_common::binary_contig_impl<
222-
argTy1, argTy2, AddOutputType, AddContigFunctor, add_contig_kernel>(
223-
exec_q, nelems, arg1_p, arg1_offset, arg2_p, arg2_offset, res_p,
224-
res_offset, depends);
226+
argTy1, argTy2, AddOutputType, AddContigFunctor, add_contig_kernel,
227+
vec_sz, n_vecs>(exec_q, nelems, arg1_p, arg1_offset, arg2_p,
228+
arg2_offset, res_p, res_offset, depends);
225229
}
226230

227231
template <typename fnT, typename T1, typename T2> struct AddContigFactory
@@ -493,9 +497,13 @@ add_inplace_contig_impl(sycl::queue &exec_q,
493497
ssize_t res_offset,
494498
const std::vector<sycl::event> &depends = {})
495499
{
500+
constexpr auto vec_sz = VecSize_v<argTy, resTy>;
501+
constexpr unsigned int n_vecs = 1u;
502+
496503
return elementwise_common::binary_inplace_contig_impl<
497-
argTy, resTy, AddInplaceContigFunctor, add_inplace_contig_kernel>(
498-
exec_q, nelems, arg_p, arg_offset, res_p, res_offset, depends);
504+
argTy, resTy, AddInplaceContigFunctor, add_inplace_contig_kernel,
505+
vec_sz, n_vecs>(exec_q, nelems, arg_p, arg_offset, res_p, res_offset,
506+
depends);
499507
}
500508

501509
template <typename fnT, typename T1, typename T2> struct AddInplaceContigFactory

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/angle.hpp‎

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -112,9 +112,13 @@ sycl::event angle_contig_impl(sycl::queue &exec_q,
112112
char *res_p,
113113
const std::vector<sycl::event> &depends = {})
114114
{
115+
using resTy = typename AngleOutputType<argTy>::value_type;
116+
constexpr auto vec_sz = VecSize_v<argTy, resTy>;
117+
constexpr unsigned int n_vec = 1u;
118+
115119
return elementwise_common::unary_contig_impl<
116-
argTy, AngleOutputType, AngleContigFunctor, angle_contig_kernel>(
117-
exec_q, nelems, arg_p, res_p, depends);
120+
argTy, AngleOutputType, AngleContigFunctor, angle_contig_kernel, vec_sz,
121+
n_vec>(exec_q, nelems, arg_p, res_p, depends);
118122
}
119123

120124
template <typename fnT, typename T> struct AngleContigFactory

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/asin.hpp‎

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -189,9 +189,13 @@ sycl::event asin_contig_impl(sycl::queue &exec_q,
189189
char *res_p,
190190
const std::vector<sycl::event> &depends = {})
191191
{
192+
using resTy = typename AsinOutputType<argTy>::value_type;
193+
constexpr auto vec_sz = VecSize_v<argTy, resTy>;
194+
constexpr unsigned int n_vec = 1u;
195+
192196
return elementwise_common::unary_contig_impl<
193-
argTy, AsinOutputType, AsinContigFunctor, asin_contig_kernel>(
194-
exec_q, nelems, arg_p, res_p, depends);
197+
argTy, AsinOutputType, AsinContigFunctor, asin_contig_kernel, vec_sz,
198+
n_vec>(exec_q, nelems, arg_p, res_p, depends);
195199
}
196200

197201
template <typename fnT, typename T> struct AsinContigFactory

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp‎

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -172,9 +172,13 @@ sycl::event asinh_contig_impl(sycl::queue &exec_q,
172172
char *res_p,
173173
const std::vector<sycl::event> &depends = {})
174174
{
175+
using resTy = typename AsinhOutputType<argTy>::value_type;
176+
constexpr auto vec_sz = VecSize_v<argTy, resTy>;
177+
constexpr unsigned int n_vec = 1u;
178+
175179
return elementwise_common::unary_contig_impl<
176-
argTy, AsinhOutputType, AsinhContigFunctor, asinh_contig_kernel>(
177-
exec_q, nelems, arg_p, res_p, depends);
180+
argTy, AsinhOutputType, AsinhContigFunctor, asinh_contig_kernel, vec_sz,
181+
n_vec>(exec_q, nelems, arg_p, res_p, depends);
178182
}
179183

180184
template <typename fnT, typename T> struct AsinhContigFactory

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp‎

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -179,9 +179,13 @@ sycl::event atan_contig_impl(sycl::queue &exec_q,
179179
char *res_p,
180180
const std::vector<sycl::event> &depends = {})
181181
{
182+
using resTy = typename AtanOutputType<argTy>::value_type;
183+
constexpr auto vec_sz = VecSize_v<argTy, resTy>;
184+
constexpr unsigned int n_vec = 1u;
185+
182186
return elementwise_common::unary_contig_impl<
183-
argTy, AtanOutputType, AtanContigFunctor, atan_contig_kernel>(
184-
exec_q, nelems, arg_p, res_p, depends);
187+
argTy, AtanOutputType, AtanContigFunctor, atan_contig_kernel, vec_sz,
188+
n_vec>(exec_q, nelems, arg_p, res_p, depends);
185189
}
186190

187191
template <typename fnT, typename T> struct AtanContigFactory

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan2.hpp‎

Lines changed: 7 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -125,10 +125,15 @@ sycl::event atan2_contig_impl(sycl::queue &exec_q,
125125
ssize_t res_offset,
126126
const std::vector<sycl::event> &depends = {})
127127
{
128+
using resTy = typename Atan2OutputType<argTy1, argTy2>::value_type;
129+
constexpr auto vec_sz = VecSize_v<argTy1, argTy2, resTy>;
130+
constexpr unsigned int n_vecs = 1u;
131+
128132
return elementwise_common::binary_contig_impl<
129133
argTy1, argTy2, Atan2OutputType, Atan2ContigFunctor,
130-
atan2_contig_kernel>(exec_q, nelems, arg1_p, arg1_offset, arg2_p,
131-
arg2_offset, res_p, res_offset, depends);
134+
atan2_contig_kernel, vec_sz, n_vecs>(exec_q, nelems, arg1_p,
135+
arg1_offset, arg2_p, arg2_offset,
136+
res_p, res_offset, depends);
132137
}
133138

134139
template <typename fnT, typename T1, typename T2> struct Atan2ContigFactory

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp‎

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -173,9 +173,13 @@ sycl::event atanh_contig_impl(sycl::queue &exec_q,
173173
char *res_p,
174174
const std::vector<sycl::event> &depends = {})
175175
{
176+
using resTy = typename AtanhOutputType<argTy>::value_type;
177+
constexpr auto vec_sz = VecSize_v<argTy, resTy>;
178+
constexpr unsigned int n_vec = 1u;
179+
176180
return elementwise_common::unary_contig_impl<
177-
argTy, AtanhOutputType, AtanhContigFunctor, atanh_contig_kernel>(
178-
exec_q, nelems, arg_p, res_p, depends);
181+
argTy, AtanhOutputType, AtanhContigFunctor, atanh_contig_kernel, vec_sz,
182+
n_vec>(exec_q, nelems, arg_p, res_p, depends);
179183
}
180184

181185
template <typename fnT, typename T> struct AtanhContigFactory

0 commit comments

Comments
 (0)