Skip to content

Commit 333c368

Browse files
vtavanavtavana
authored
simplify backend implementation of dpnp.kaiser (#2472)
This PR simplifies backend implementation of `dpnp.kaiser` by getting rid of unnecessary template and address some technical debt.
1 parent b10bd15 commit 333c368

File tree

7 files changed

+45
-44
lines changed

7 files changed

+45
-44
lines changed

‎dpnp/backend/extensions/ufunc/elementwise_functions/nan_to_num.cpp‎

Lines changed: 10 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -176,16 +176,17 @@ std::pair<sycl::event, sycl::event>
176176

177177
dpctl::tensor::validation::CheckWritable::throw_if_not_writable(dst);
178178

179-
int src_nd = src.get_ndim();
179+
const int src_nd = src.get_ndim();
180180
if (src_nd != dst.get_ndim()) {
181181
throw py::value_error("Array dimensions are not the same.");
182182
}
183183

184184
const py::ssize_t *src_shape = src.get_shape_raw();
185185
const py::ssize_t *dst_shape = dst.get_shape_raw();
186186

187-
std::size_t nelems = src.get_size();
188-
bool shapes_equal = std::equal(src_shape, src_shape + src_nd, dst_shape);
187+
const std::size_t nelems = src.get_size();
188+
const bool shapes_equal =
189+
std::equal(src_shape, src_shape + src_nd, dst_shape);
189190
if (!shapes_equal) {
190191
throw py::value_error("Array shapes are not the same.");
191192
}
@@ -209,14 +210,14 @@ std::pair<sycl::event, sycl::event>
209210
char *dst_data = dst.get_data();
210211

211212
// handle contiguous inputs
212-
bool is_src_c_contig = src.is_c_contiguous();
213-
bool is_src_f_contig = src.is_f_contiguous();
213+
const bool is_src_c_contig = src.is_c_contiguous();
214+
const bool is_src_f_contig = src.is_f_contiguous();
214215

215-
bool is_dst_c_contig = dst.is_c_contiguous();
216-
bool is_dst_f_contig = dst.is_f_contiguous();
216+
const bool is_dst_c_contig = dst.is_c_contiguous();
217+
const bool is_dst_f_contig = dst.is_f_contiguous();
217218

218-
bool both_c_contig = (is_src_c_contig && is_dst_c_contig);
219-
bool both_f_contig = (is_src_f_contig && is_dst_f_contig);
219+
const bool both_c_contig = (is_src_c_contig && is_dst_c_contig);
220+
const bool both_f_contig = (is_src_f_contig && is_dst_f_contig);
220221

221222
if (both_c_contig || both_f_contig) {
222223
auto contig_fn = nan_to_num_contig_dispatch_vector[src_typeid];

‎dpnp/backend/extensions/window/bartlett.hpp‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -16,7 +16,7 @@
1616
// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
1717
// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
1818
// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
19-
// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
19+
// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, RES, OR PROFITS; OR BUSINESS
2020
// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
2121
// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
2222
// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
@@ -35,18 +35,18 @@ template <typename T>
3535
class BartlettFunctor
3636
{
3737
private:
38-
T *data = nullptr;
38+
T *res = nullptr;
3939
const std::size_t N;
4040

4141
public:
42-
BartlettFunctor(T *data, const std::size_t N) : data(data), N(N) {}
42+
BartlettFunctor(T *res, const std::size_t N) : res(res), N(N) {}
4343

4444
void operator()(sycl::id<1> id) const
4545
{
4646
const auto i = id.get(0);
4747

4848
const T alpha = (N - 1) / T(2);
49-
data[i] = T(1) - sycl::fabs(i - alpha) / alpha;
49+
res[i] = T(1) - sycl::fabs(i - alpha) / alpha;
5050
}
5151
};
5252

‎dpnp/backend/extensions/window/blackman.hpp‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -35,19 +35,19 @@ template <typename T>
3535
class BlackmanFunctor
3636
{
3737
private:
38-
T *data = nullptr;
38+
T *res = nullptr;
3939
const std::size_t N;
4040

4141
public:
42-
BlackmanFunctor(T *data, const std::size_t N) : data(data), N(N) {}
42+
BlackmanFunctor(T *res, const std::size_t N) : res(res), N(N) {}
4343

4444
void operator()(sycl::id<1> id) const
4545
{
4646
const auto i = id.get(0);
4747

4848
const T alpha = T(2) * i / (N - 1);
49-
data[i] = T(0.42) - T(0.5) * sycl::cospi(alpha) +
50-
T(0.08) * sycl::cospi(T(2) * alpha);
49+
res[i] = T(0.42) - T(0.5) * sycl::cospi(alpha) +
50+
T(0.08) * sycl::cospi(T(2) * alpha);
5151
}
5252
};
5353

‎dpnp/backend/extensions/window/common.hpp‎

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -47,16 +47,16 @@ typedef sycl::event (*window_fn_ptr_t)(sycl::queue &,
4747
const std::vector<sycl::event> &);
4848

4949
template <typename T, template <typename> class Functor>
50-
sycl::event window_impl(sycl::queue &q,
50+
sycl::event window_impl(sycl::queue &exec_q,
5151
char *result,
5252
const std::size_t nelems,
5353
const std::vector<sycl::event> &depends)
5454
{
55-
dpctl::tensor::type_utils::validate_type_for_device<T>(q);
55+
dpctl::tensor::type_utils::validate_type_for_device<T>(exec_q);
5656

5757
T *res = reinterpret_cast<T *>(result);
5858

59-
sycl::event window_ev = q.submit([&](sycl::handler &cgh) {
59+
sycl::event window_ev = exec_q.submit([&](sycl::handler &cgh) {
6060
cgh.depends_on(depends);
6161

6262
using WindowKernel = Functor<T>;
@@ -75,7 +75,7 @@ std::tuple<size_t, char *, funcPtrT>
7575
{
7676
dpctl::tensor::validation::CheckWritable::throw_if_not_writable(result);
7777

78-
int nd = result.get_ndim();
78+
const int nd = result.get_ndim();
7979
if (nd != 1) {
8080
throw py::value_error("Array should be 1d");
8181
}
@@ -87,17 +87,17 @@ std::tuple<size_t, char *, funcPtrT>
8787

8888
const bool is_result_c_contig = result.is_c_contiguous();
8989
if (!is_result_c_contig) {
90-
throw py::value_error("The result input array is not c-contiguous.");
90+
throw py::value_error("The result array is not c-contiguous.");
9191
}
9292

93-
size_t nelems = result.get_size();
93+
const std::size_t nelems = result.get_size();
9494
if (nelems == 0) {
9595
return std::make_tuple(nelems, nullptr, nullptr);
9696
}
9797

98-
int result_typenum = result.get_typenum();
98+
const int result_typenum = result.get_typenum();
9999
auto array_types = dpctl_td_ns::usm_ndarray_types();
100-
int result_type_id = array_types.typenum_to_lookup_id(result_typenum);
100+
const int result_type_id = array_types.typenum_to_lookup_id(result_typenum);
101101
funcPtrT fn = window_dispatch_vector[result_type_id];
102102

103103
if (fn == nullptr) {

‎dpnp/backend/extensions/window/hamming.hpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -35,17 +35,17 @@ template <typename T>
3535
class HammingFunctor
3636
{
3737
private:
38-
T *data = nullptr;
38+
T *res = nullptr;
3939
const std::size_t N;
4040

4141
public:
42-
HammingFunctor(T *data, const std::size_t N) : data(data), N(N) {}
42+
HammingFunctor(T *res, const std::size_t N) : res(res), N(N) {}
4343

4444
void operator()(sycl::id<1> id) const
4545
{
4646
const auto i = id.get(0);
4747

48-
data[i] = T(0.54) - T(0.46) * sycl::cospi(T(2) * i / (N - 1));
48+
res[i] = T(0.54) - T(0.46) * sycl::cospi(T(2) * i / (N - 1));
4949
}
5050
};
5151

‎dpnp/backend/extensions/window/hanning.hpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -35,17 +35,17 @@ template <typename T>
3535
class HanningFunctor
3636
{
3737
private:
38-
T *data = nullptr;
38+
T *res = nullptr;
3939
const std::size_t N;
4040

4141
public:
42-
HanningFunctor(T *data, const std::size_t N) : data(data), N(N) {}
42+
HanningFunctor(T *res, const std::size_t N) : res(res), N(N) {}
4343

4444
void operator()(sycl::id<1> id) const
4545
{
4646
const auto i = id.get(0);
4747

48-
data[i] = T(0.5) - T(0.5) * sycl::cospi(T(2) * i / (N - 1));
48+
res[i] = T(0.5) - T(0.5) * sycl::cospi(T(2) * i / (N - 1));
4949
}
5050
};
5151

‎dpnp/backend/extensions/window/kaiser.cpp‎

Lines changed: 13 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -50,13 +50,13 @@ template <typename T>
5050
class KaiserFunctor
5151
{
5252
private:
53-
T *data = nullptr;
53+
T *res = nullptr;
5454
const std::size_t N;
5555
const T beta;
5656

5757
public:
58-
KaiserFunctor(T *data, const std::size_t N, const T beta)
59-
: data(data), N(N), beta(beta)
58+
KaiserFunctor(T *res, const std::size_t N, const T beta)
59+
: res(res), N(N), beta(beta)
6060
{
6161
}
6262

@@ -67,27 +67,27 @@ class KaiserFunctor
6767
const auto i = id.get(0);
6868
const T alpha = (N - 1) / T(2);
6969
const T tmp = (i - alpha) / alpha;
70-
data[i] = cyl_bessel_i0(beta * sycl::sqrt(1 - tmp * tmp)) /
71-
cyl_bessel_i0(beta);
70+
res[i] = cyl_bessel_i0(beta * sycl::sqrt(1 - tmp * tmp)) /
71+
cyl_bessel_i0(beta);
7272
}
7373
};
7474

75-
template <typename T, template <typename> class Functor>
76-
sycl::event kaiser_impl(sycl::queue &q,
75+
template <typename T>
76+
sycl::event kaiser_impl(sycl::queue &exec_q,
7777
char *result,
7878
const std::size_t nelems,
7979
const py::object &py_beta,
8080
const std::vector<sycl::event> &depends)
8181
{
82-
dpctl::tensor::type_utils::validate_type_for_device<T>(q);
82+
dpctl::tensor::type_utils::validate_type_for_device<T>(exec_q);
8383

8484
T *res = reinterpret_cast<T *>(result);
8585
const T beta = py::cast<const T>(py_beta);
8686

87-
sycl::event kaiser_ev = q.submit([&](sycl::handler &cgh) {
87+
sycl::event kaiser_ev = exec_q.submit([&](sycl::handler &cgh) {
8888
cgh.depends_on(depends);
8989

90-
using KaiserKernel = Functor<T>;
90+
using KaiserKernel = KaiserFunctor<T>;
9191
cgh.parallel_for<KaiserKernel>(sycl::range<1>(nelems),
9292
KaiserKernel(res, nelems, beta));
9393
});
@@ -101,7 +101,7 @@ struct KaiserFactory
101101
fnT get()
102102
{
103103
if constexpr (std::is_floating_point_v<T>) {
104-
return kaiser_impl<T, KaiserFunctor>;
104+
return kaiser_impl<T>;
105105
}
106106
else {
107107
return nullptr;
@@ -115,15 +115,15 @@ std::pair<sycl::event, sycl::event>
115115
const dpctl::tensor::usm_ndarray &result,
116116
const std::vector<sycl::event> &depends)
117117
{
118-
auto [nelems, result_typeless_ptr, fn] =
118+
auto [nelems, result_typeless_ptr, kaiser_fn] =
119119
window_fn<kaiser_fn_ptr_t>(exec_q, result, kaiser_dispatch_vector);
120120

121121
if (nelems == 0) {
122122
return std::make_pair(sycl::event{}, sycl::event{});
123123
}
124124

125125
sycl::event kaiser_ev =
126-
fn(exec_q, result_typeless_ptr, nelems, py_beta, depends);
126+
kaiser_fn(exec_q, result_typeless_ptr, nelems, py_beta, depends);
127127
sycl::event args_ev =
128128
dpctl::utils::keep_args_alive(exec_q, {result}, {kaiser_ev});
129129

0 commit comments

Comments
 (0)