Skip to content

Commit

Permalink
Fully functional dpnp.arange without complex number support
Browse files Browse the repository at this point in the history
  • Loading branch information
@chudur-budur
chudur-budur committed Nov 1, 2023
1 parent ac5ec10 commit 5556281
Show file tree
Hide file tree
Showing 5 changed files with 135 additions and 319 deletions.
21 changes: 10 additions & 11 deletions numba_dpex/core/runtime/kernels/tensor/include/sequences.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -19,6 +19,8 @@

#include "typeutils.hpp"

namespace dpexrt_tensor = dpex::rt::kernel::tensor;

namespace dpex
{
namespace rt
Expand Down Expand Up @@ -47,8 +49,7 @@ template <typename T> class SequenceStepFunctor
void operator()(sycl::id<1> wiid) const
{
auto i = wiid.get(0);
if constexpr (dpex::rt::kernel::tensor::typeutils::is_complex<T>::value)
{
if constexpr (dpexrt_tensor::typeutils::is_complex<T>::value) {
p[i] = T{start_v.real() + i * step_v.real(),
start_v.imag() + i * step_v.imag()};
}
Expand Down Expand Up @@ -78,8 +79,7 @@ template <typename T, typename wT> class AffineSequenceFunctor
auto i = wiid.get(0);
wT wc = wT(i) / n;
wT w = wT(n - i) / n;
if constexpr (dpex::rt::kernel::tensor::typeutils::is_complex<T>::value)
{
if constexpr (dpexrt_tensor::typeutils::is_complex<T>::value) {
using reT = typename T::value_type;
auto _w = static_cast<reT>(w);
auto _wc = static_cast<reT>(wc);
Expand All @@ -104,7 +104,7 @@ template <typename T, typename wT> class AffineSequenceFunctor
}
else {
auto affine_comb = start_v * w + end_v * wc;
p[i] = dpex::rt::kernel::tensor::typeutils::convert_impl<
p[i] = dpexrt_tensor::typeutils::convert_impl<
T, decltype(affine_comb)>(affine_comb);
}
}
Expand All @@ -119,14 +119,14 @@ sycl::event sequence_step_kernel(sycl::queue exec_q,
const std::vector<sycl::event> &depends)
{
std::cout << "sequqnce_step_kernel<"
<< dpex::rt::kernel::tensor::typeutils::demangle<T>()
<< dpexrt_tensor::typeutils::demangle<T>()
<< ">(): nelems = " << nelems << ", start_v = " << start_v
<< ", step_v = " << step_v << std::endl;

dpex::rt::kernel::tensor::typeutils::validate_type_for_device<T>(exec_q);
dpexrt_tensor::typeutils::validate_type_for_device<T>(exec_q);

std::cout << "sequqnce_step_kernel<"
<< dpex::rt::kernel::tensor::typeutils::demangle<T>()
<< dpexrt_tensor::typeutils::demangle<T>()
<< ">(): validate_type_for_device<T>(exec_q) = done" << std::endl;

sycl::event seq_step_event = exec_q.submit([&](sycl::handler &cgh) {
Expand All @@ -148,7 +148,7 @@ sycl::event affine_sequence_kernel(sycl::queue &exec_q,
char *array_data,
const std::vector<sycl::event> &depends)
{
dpex::rt::kernel::tensor::typeutils::validate_type_for_device<T>(exec_q);
dpexrt_tensor::typeutils::validate_type_for_device<T>(exec_q);
bool device_supports_doubles = exec_q.get_device().has(sycl::aspect::fp64);
sycl::event affine_seq_step_event = exec_q.submit([&](sycl::handler &cgh) {
cgh.depends_on(depends);
Expand Down Expand Up @@ -187,8 +187,7 @@ sycl::event sequence_step(sycl::queue &exec_q,
std::cerr << e.what() << std::endl;
}

std::cout << "sequqnce_step()<"
<< dpex::rt::kernel::tensor::typeutils::demangle<T>()
std::cout << "sequqnce_step()<" << dpexrt_tensor::typeutils::demangle<T>()
<< ">: nelems = " << nelems << ", *start_v = " << (*start_v)
<< ", *step_v = " << (*step_v) << std::endl;

Expand Down
112 changes: 0 additions & 112 deletions numba_dpex/core/runtime/kernels/tensor/include/typeutils.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -172,118 +172,6 @@ template <typename T> void validate_type_for_device(const sycl::queue &q)
validate_type_for_device<T>(q.get_device());
}

// template <typename Op, typename Vec, std::size_t... I>
// auto vec_cast_impl(const Vec &v, std::index_sequence<I...>)
// {
// return Op{v[I]...};
// }

// template <typename dstT,
// typename srcT,
// std::size_t N,
// typename Indices = std::make_index_sequence<N>>
// auto vec_cast(const sycl::vec<srcT, N> &s)
// {
// if constexpr (std::is_same_v<srcT, dstT>) {
// return s;
// }
// else {
// return vec_cast_impl<sycl::vec<dstT, N>, sycl::vec<srcT, N>>(s,
// Indices{});
// }
// }

// struct usm_ndarray_types
// {
// int typenum_to_lookup_id(int typenum) const
// {
// // using typenum_t = ::dpctl::tensor::type_dispatch::typenum_t;
// auto const &api = ::dpctl::detail::dpctl_capi::get();

// if (typenum == api.UAR_DOUBLE_) {
// return static_cast<int>(typenum_t::DOUBLE);
// }
// else if (typenum == api.UAR_INT64_) {
// return static_cast<int>(typenum_t::INT64);
// }
// else if (typenum == api.UAR_INT32_) {
// return static_cast<int>(typenum_t::INT32);
// }
// else if (typenum == api.UAR_BOOL_) {
// return static_cast<int>(typenum_t::BOOL);
// }
// else if (typenum == api.UAR_CDOUBLE_) {
// return static_cast<int>(typenum_t::CDOUBLE);
// }
// else if (typenum == api.UAR_FLOAT_) {
// return static_cast<int>(typenum_t::FLOAT);
// }
// else if (typenum == api.UAR_INT16_) {
// return static_cast<int>(typenum_t::INT16);
// }
// else if (typenum == api.UAR_INT8_) {
// return static_cast<int>(typenum_t::INT8);
// }
// else if (typenum == api.UAR_UINT64_) {
// return static_cast<int>(typenum_t::UINT64);
// }
// else if (typenum == api.UAR_UINT32_) {
// return static_cast<int>(typenum_t::UINT32);
// }
// else if (typenum == api.UAR_UINT16_) {
// return static_cast<int>(typenum_t::UINT16);
// }
// else if (typenum == api.UAR_UINT8_) {
// return static_cast<int>(typenum_t::UINT8);
// }
// else if (typenum == api.UAR_CFLOAT_) {
// return static_cast<int>(typenum_t::CFLOAT);
// }
// else if (typenum == api.UAR_HALF_) {
// return static_cast<int>(typenum_t::HALF);
// }
// else if (typenum == api.UAR_INT_ || typenum == api.UAR_UINT_) {
// switch (sizeof(int)) {
// case sizeof(int32_t):
// return ((typenum == api.UAR_INT_)
// ? static_cast<int>(typenum_t::INT32)
// : static_cast<int>(typenum_t::UINT32));
// case sizeof(int64_t):
// return ((typenum == api.UAR_INT_)
// ? static_cast<int>(typenum_t::INT64)
// : static_cast<int>(typenum_t::UINT64));
// default:
// throw_unrecognized_typenum_error(typenum);
// }
// }
// else if (typenum == api.UAR_LONGLONG_ || typenum ==
// api.UAR_ULONGLONG_)
// {
// switch (sizeof(long long)) {
// case sizeof(int64_t):
// return ((typenum == api.UAR_LONGLONG_)
// ? static_cast<int>(typenum_t::INT64)
// : static_cast<int>(typenum_t::UINT64));
// default:
// throw_unrecognized_typenum_error(typenum);
// }
// }
// else {
// throw_unrecognized_typenum_error(typenum);
// }
// // return code signalling error, should never be reached
// assert(false);
// return -1;
// }

// private:
// void throw_unrecognized_typenum_error(int typenum) const
// {
// throw std::runtime_error("Unrecognized typenum " +
// std::to_string(typenum) + " encountered.");
// }
// };

} // namespace typeutils
} // namespace tensor
} // namespace kernel
Expand Down
37 changes: 16 additions & 21 deletions numba_dpex/core/runtime/kernels/tensor/src/sequences.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -12,34 +12,31 @@
#include "../include/typeutils.hpp"
#include "../include/api.h"

static dpex::rt::kernel::tensor::sequence_step_ptr_t
sequence_step_dispatch_vector
[dpex::rt::kernel::tensor::typeutils::num_types];
namespace dpexrt_tensor = dpex::rt::kernel::tensor;

static dpex::rt::kernel::tensor::affine_sequence_ptr_t
affine_sequence_dispatch_vector
[dpex::rt::kernel::tensor::typeutils::num_types];
static dpexrt_tensor::sequence_step_ptr_t
sequence_step_dispatch_vector[dpexrt_tensor::typeutils::num_types];

static dpexrt_tensor::affine_sequence_ptr_t
affine_sequence_dispatch_vector[dpexrt_tensor::typeutils::num_types];

extern "C" void NUMBA_DPEX_SYCL_KERNEL_init_sequence_step_dispatch_vectors()
{
dpex::rt::kernel::tensor::dispatch::DispatchVectorBuilder<
dpex::rt::kernel::tensor::sequence_step_ptr_t,
dpex::rt::kernel::tensor::SequenceStepFactory,
dpex::rt::kernel::tensor::typeutils::num_types>
dpexrt_tensor::dispatch::DispatchVectorBuilder<
dpexrt_tensor::sequence_step_ptr_t, dpexrt_tensor::SequenceStepFactory,
dpexrt_tensor::typeutils::num_types>
dvb;
dvb.populate_dispatch_vector(sequence_step_dispatch_vector);
std::cout << "-----> init_sequence_dispatch_vectors()" << std::endl;
}

extern "C" void NUMBA_DPEX_SYCL_KERNEL_init_affine_sequence_dispatch_vectors()
{
dpex::rt::kernel::tensor::dispatch::DispatchVectorBuilder<
dpex::rt::kernel::tensor::affine_sequence_ptr_t,
dpex::rt::kernel::tensor::AffineSequenceFactory,
dpex::rt::kernel::tensor::typeutils::num_types>
dpexrt_tensor::dispatch::DispatchVectorBuilder<
dpexrt_tensor::affine_sequence_ptr_t,
dpexrt_tensor::AffineSequenceFactory,
dpexrt_tensor::typeutils::num_types>
dvb;
dvb.populate_dispatch_vector(affine_sequence_dispatch_vector);
std::cout << "-----> init_affine_sequence_dispatch_vectors()" << std::endl;
}

extern "C" uint NUMBA_DPEX_SYCL_KERNEL_populate_arystruct_sequence(
Expand All @@ -53,14 +50,12 @@ extern "C" uint NUMBA_DPEX_SYCL_KERNEL_populate_arystruct_sequence(
{
std::cout << "NUMBA_DPEX_SYCL_KERNEL_populate_arystruct_sequence:"
<< " start = "
<< dpex::rt::kernel::tensor::typeutils::caste_using_typeid(
start, dst_typeid)
<< dpexrt_tensor::typeutils::caste_using_typeid(start, dst_typeid)
<< std::endl;

std::cout << "NUMBA_DPEX_SYCL_KERNEL_populate_arystruct_sequence:"
<< " dt = "
<< dpex::rt::kernel::tensor::typeutils::caste_using_typeid(
dt, dst_typeid)
<< dpexrt_tensor::typeutils::caste_using_typeid(dt, dst_typeid)
<< std::endl;

if (ndim != 1) {
Expand Down Expand Up @@ -98,7 +93,7 @@ extern "C" uint NUMBA_DPEX_SYCL_KERNEL_populate_arystruct_sequence(
return 1;
}

// uint dpex::rt::kernel::tensor::tensor::populate_arystruct_affine_sequence(
// uint dpexrt_tensor::tensor::populate_arystruct_affine_sequence(
// void *start,
// void *end,
// arystruct_t *dst,
Expand Down
19 changes: 1 addition & 18 deletions numba_dpex/dpnp_iface/_intrinsic.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -361,25 +361,13 @@ def alloc_empty_arrayobj(context, builder, sig, queue_ref, args, is_like=False):
Returns: The LLVM IR value that stores the empty array
"""

print("alloc_empty_arrayobj: sig =", sig)
print("alloc_empty_arrayobj: args =", args)

arrtype, shape = (
_parse_empty_like_args(context, builder, sig, args)
if is_like
else _parse_empty_args(context, builder, sig, args)
)
print(
"alloc_empty_arrayobj(): arrtype =",
arrtype,
"type(arrtype) =",
type(arrtype),
)
print(
"alloc_empty_arrayobj(): shape =", shape, ", type(shape) =", type(shape)
)
ary = _empty_nd_impl(context, builder, arrtype, shape, queue_ref)
print("alloc_empty_arrayobj(): ary =", ary, ", type(ary) =", type(ary))

return ary


Expand Down Expand Up @@ -485,8 +473,6 @@ def impl_dpnp_empty(
ty_retty_ref,
)

print("--- impl_dpnp_empty()")

sycl_queue_arg_pos = -2

def codegen(context, builder, sig, args):
Expand All @@ -500,9 +486,6 @@ def codegen(context, builder, sig, args):
sycl_queue_arg=sycl_queue_arg,
)

print("impl_dpnp_empty(): sig =", sig, type(sig))
print("impl_dpnp_empty(): args =", args, type(args))

ary = alloc_empty_arrayobj(
context, builder, sig, qref_payload.queue_ref, args
)
Expand Down
Loading

0 comments on commit 5556281

Please sign in to comment.