Skip to content

Native astype #642

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 5 commits into from
Mar 11, 2021
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
11 changes: 11 additions & 0 deletions dpnp/backend/include/dpnp_iface.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -141,6 +141,17 @@ INP_DLLEXPORT void dpnp_any_c(const void* array, void* result, const size_t size
template <typename _DataType>
INP_DLLEXPORT void dpnp_arange_c(size_t start, size_t step, void* result1, size_t size);

/**
* @ingroup BACKEND_API
* @brief Copy of the array, cast to a specified type.
*
* @param [in] array Input array.
* @param [out] result Output array.
* @param [in] size Number of input elements in `array`.
*/
template <typename _DataType, typename _ResultType>
INP_DLLEXPORT void dpnp_astype_c(const void* array, void* result, const size_t size);

/**
* @ingroup BACKEND_API
* @brief Implementation of full function
Expand Down
1 change: 1 addition & 0 deletions dpnp/backend/include/dpnp_iface_fptr.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -73,6 +73,7 @@ enum class DPNPFuncName : size_t
DPNP_FN_ARGMAX, /**< Used in numpy.argmax() implementation */
DPNP_FN_ARGMIN, /**< Used in numpy.argmin() implementation */
DPNP_FN_ARGSORT, /**< Used in numpy.argsort() implementation */
DPNP_FN_ASTYPE, /**< Used in numpy.astype() implementation */
DPNP_FN_BITWISE_AND, /**< Used in numpy.bitwise_and() implementation */
DPNP_FN_BITWISE_OR, /**< Used in numpy.bitwise_or() implementation */
DPNP_FN_BITWISE_XOR, /**< Used in numpy.bitwise_xor() implementation */
Expand Down
63 changes: 63 additions & 0 deletions dpnp/backend/kernels/dpnp_krnl_common.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -35,6 +35,42 @@
namespace mkl_blas = oneapi::mkl::blas;
namespace mkl_lapack = oneapi::mkl::lapack;

template <typename _DataType, typename _ResultType>
class dpnp_astype_c_kernel;

template <typename _DataType, typename _ResultType>
void dpnp_astype_c(const void* array1_in, void* result1, const size_t size)
{
cl::sycl::event event;

const _DataType* array_in = reinterpret_cast<const _DataType*>(array1_in);
_ResultType* result = reinterpret_cast<_ResultType*>(result1);

if ((array_in == nullptr) || (result == nullptr))
{
return;
}

if (size == 0)
{
return;
}

cl::sycl::range<1> gws(size);
auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) {
size_t i = global_id[0];
result[i] = array_in[i];
};

auto kernel_func = [&](cl::sycl::handler& cgh) {
cgh.parallel_for<class dpnp_astype_c_kernel<_DataType, _ResultType>>(gws, kernel_parallel_for_func);
};

event = DPNP_QUEUE.submit(kernel_func);

event.wait();
}

template <typename _KernelNameSpecialization1, typename _KernelNameSpecialization2, typename _KernelNameSpecialization3>
class dpnp_dot_c_kernel;

Expand Down Expand Up @@ -324,6 +360,33 @@ void dpnp_matmul_c(void* array1_in, void* array2_in, void* result1, size_t size_

void func_map_init_linalg(func_map_t& fmap)
{
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_BLN][eft_BLN] = {eft_BLN, (void*)dpnp_astype_c<bool, bool>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_BLN][eft_INT] = {eft_INT, (void*)dpnp_astype_c<bool, int>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_BLN][eft_LNG] = {eft_LNG, (void*)dpnp_astype_c<bool, long>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_BLN][eft_FLT] = {eft_FLT, (void*)dpnp_astype_c<bool, float>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_BLN][eft_DBL] = {eft_DBL, (void*)dpnp_astype_c<bool, double>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_INT][eft_BLN] = {eft_BLN, (void*)dpnp_astype_c<int, bool>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_astype_c<int, int>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_INT][eft_LNG] = {eft_LNG, (void*)dpnp_astype_c<int, long>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_INT][eft_FLT] = {eft_FLT, (void*)dpnp_astype_c<int, float>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_INT][eft_DBL] = {eft_DBL, (void*)dpnp_astype_c<int, double>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_LNG][eft_BLN] = {eft_BLN, (void*)dpnp_astype_c<long, bool>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_LNG][eft_INT] = {eft_INT, (void*)dpnp_astype_c<long, int>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_astype_c<long, long>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_LNG][eft_FLT] = {eft_FLT, (void*)dpnp_astype_c<long, float>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_LNG][eft_DBL] = {eft_DBL, (void*)dpnp_astype_c<long, double>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_FLT][eft_BLN] = {eft_BLN, (void*)dpnp_astype_c<float, bool>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_FLT][eft_INT] = {eft_INT, (void*)dpnp_astype_c<float, int>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_FLT][eft_LNG] = {eft_LNG, (void*)dpnp_astype_c<float, long>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_astype_c<float, float>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_FLT][eft_DBL] = {eft_DBL, (void*)dpnp_astype_c<float, double>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_DBL][eft_BLN] = {eft_BLN, (void*)dpnp_astype_c<double, bool>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_DBL][eft_INT] = {eft_INT, (void*)dpnp_astype_c<double, int>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_DBL][eft_LNG] = {eft_LNG, (void*)dpnp_astype_c<double, long>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_DBL][eft_FLT] = {eft_FLT, (void*)dpnp_astype_c<double, float>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_astype_c<double, double>};
fmap[DPNPFuncName::DPNP_FN_ASTYPE][eft_C128][eft_C128] = {eft_C128, (void*)dpnp_astype_c<std::complex<double>, std::complex<double>>};

fmap[DPNPFuncName::DPNP_FN_DOT][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_dot_c<int, int, int>};
fmap[DPNPFuncName::DPNP_FN_DOT][eft_INT][eft_LNG] = {eft_LNG, (void*)dpnp_dot_c<int, long, long>};
fmap[DPNPFuncName::DPNP_FN_DOT][eft_INT][eft_FLT] = {eft_DBL, (void*)dpnp_dot_c<int, float, double>};
Expand Down
1 change: 1 addition & 0 deletions dpnp/dpnp_algo/dpnp_algo.pxd
Original file line number Diff line number Diff line change
Expand Up @@ -46,6 +46,7 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na
DPNP_FN_ARGMAX
DPNP_FN_ARGMIN
DPNP_FN_ARGSORT
DPNP_FN_ASTYPE
DPNP_FN_BITWISE_AND
DPNP_FN_BITWISE_OR
DPNP_FN_BITWISE_XOR
Expand Down
13 changes: 10 additions & 3 deletions dpnp/dpnp_algo/dpnp_algo.pyx
Original file line number Diff line number Diff line change
Expand Up @@ -69,6 +69,7 @@ include "dpnp_algo_trigonometric.pyx"


ctypedef void(*fptr_dpnp_arange_t)(size_t, size_t, void *, size_t)
ctypedef void(*fptr_dpnp_astype_t)(const void *, void * , const size_t)
ctypedef void(*fptr_dpnp_initval_t)(void *, void * , size_t)


Expand Down Expand Up @@ -125,10 +126,16 @@ cpdef dparray dpnp_array(obj, dtype=None):


cpdef dparray dpnp_astype(dparray array1, dtype_target):
cdef dparray result = dparray(array1.shape, dtype=dtype_target)
cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(array1.dtype)
cdef DPNPFuncType param2_type = dpnp_dtype_to_DPNPFuncType(dtype_target)
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

dtype_target might be diffrent. It can contain some object with no dtype. I think this is the cause of the error you saw


for i in range(result.size):
result[i] = array1[i]
cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_ASTYPE, param1_type, param2_type)

result_type = dpnp_DPNPFuncType_to_dtype(< size_t > kernel_data.return_type)
cdef dparray result = dparray(array1.shape, dtype=result_type)

cdef fptr_dpnp_astype_t func = <fptr_dpnp_astype_t > kernel_data.ptr
func(array1.get_data(), result.get_data(), array1.size)

return result

Expand Down
19 changes: 18 additions & 1 deletion tests/skipped_tests.tbl
Original file line number Diff line number Diff line change
@@ -1,3 +1,21 @@
tests/test_dparray.py::test_astype[[-2, -1, 0, 1, 2]-complex-float64]
tests/test_dparray.py::test_astype[[-2, -1, 0, 1, 2]-complex-float32]
tests/test_dparray.py::test_astype[[-2, -1, 0, 1, 2]-complex-int64]
tests/test_dparray.py::test_astype[[-2, -1, 0, 1, 2]-complex-int32]
tests/test_dparray.py::test_astype[[-2, -1, 0, 1, 2]-complex-bool]
tests/test_dparray.py::test_astype[[-2, -1, 0, 1, 2]-complex-bool_]
tests/test_dparray.py::test_astype[[[-2, -1], [1, 2]]-complex-float64]
tests/test_dparray.py::test_astype[[[-2, -1], [1, 2]]-complex-float32]
tests/test_dparray.py::test_astype[[[-2, -1], [1, 2]]-complex-int64]
tests/test_dparray.py::test_astype[[[-2, -1], [1, 2]]-complex-int32]
tests/test_dparray.py::test_astype[[[-2, -1], [1, 2]]-complex-bool]
tests/test_dparray.py::test_astype[[[-2, -1], [1, 2]]-complex-bool_]
tests/test_dparray.py::test_astype[[]-complex-float64]
tests/test_dparray.py::test_astype[[]-complex-float32]
tests/test_dparray.py::test_astype[[]-complex-int64]
tests/test_dparray.py::test_astype[[]-complex-int32]
tests/test_dparray.py::test_astype[[]-complex-bool]
tests/test_dparray.py::test_astype[[]-complex-bool_]
tests/test_linalg.py::test_cond[-1-[[1, 2, 3], [4, 5, 6], [7, 8, 9]]]
tests/test_linalg.py::test_cond[1-[[1, 2, 3], [4, 5, 6], [7, 8, 9]]]
tests/test_linalg.py::test_cond[-2-[[1, 0, -1], [0, 1, 0], [1, 0, 1]]]
Expand Down Expand Up @@ -154,7 +172,6 @@ tests/test_linalg.py::test_svd[(5,3)-complex128]
tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[ lambda x: x]
tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[ lambda x: dpnp.asarray(x).astype(dpnp.int8)]
tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[ lambda x: dpnp.asarray(x).astype(dpnp.complex64)]
tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[ lambda x: dpnp.asarray(x).astype(object)]
Copy link

@samir-nasibli samir-nasibli Mar 10, 2021

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

👍

tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[ lambda x: [(i, i) for i in x]]
tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[ lambda x: dpnp.vstack([x, x]).T]
tests/test_random.py::TestPermutationsTestShuffle::test_shuffle1[ lambda x: (dpnp.asarray([(i, i) for i in x], [("a", int), ("b", int)]).view(dpnp.recarray))]
Expand Down
18 changes: 18 additions & 0 deletions tests/skipped_tests_gpu.tbl
Original file line number Diff line number Diff line change
@@ -1,3 +1,21 @@
tests/test_dparray.py::test_astype[[-2, -1, 0, 1, 2]-complex-float64]
tests/test_dparray.py::test_astype[[-2, -1, 0, 1, 2]-complex-float32]
tests/test_dparray.py::test_astype[[-2, -1, 0, 1, 2]-complex-int64]
tests/test_dparray.py::test_astype[[-2, -1, 0, 1, 2]-complex-int32]
tests/test_dparray.py::test_astype[[-2, -1, 0, 1, 2]-complex-bool]
tests/test_dparray.py::test_astype[[-2, -1, 0, 1, 2]-complex-bool_]
tests/test_dparray.py::test_astype[[[-2, -1], [1, 2]]-complex-float64]
tests/test_dparray.py::test_astype[[[-2, -1], [1, 2]]-complex-float32]
tests/test_dparray.py::test_astype[[[-2, -1], [1, 2]]-complex-int64]
tests/test_dparray.py::test_astype[[[-2, -1], [1, 2]]-complex-int32]
tests/test_dparray.py::test_astype[[[-2, -1], [1, 2]]-complex-bool]
tests/test_dparray.py::test_astype[[[-2, -1], [1, 2]]-complex-bool_]
tests/test_dparray.py::test_astype[[]-complex-float64]
tests/test_dparray.py::test_astype[[]-complex-float32]
tests/test_dparray.py::test_astype[[]-complex-int64]
tests/test_dparray.py::test_astype[[]-complex-int32]
tests/test_dparray.py::test_astype[[]-complex-bool]
tests/test_dparray.py::test_astype[[]-complex-bool_]
tests/test_dot.py::test_dot_arange[float32]
tests/test_dot.py::test_dot_arange[float64]
tests/test_dot.py::test_dot_ones[float32]
Expand Down
20 changes: 20 additions & 0 deletions tests/test_dparray.py
Original file line number Diff line number Diff line change
@@ -0,0 +1,20 @@
import dpnp
import numpy
import pytest


@pytest.mark.parametrize("res_dtype",
[numpy.float64, numpy.float32, numpy.int64, numpy.int32, numpy.bool, numpy.bool_, numpy.complex],
ids=['float64', 'float32', 'int64', 'int32', 'bool', 'bool_', 'complex'])
@pytest.mark.parametrize("arr_dtype",
[numpy.float64, numpy.float32, numpy.int64, numpy.int32, numpy.bool, numpy.bool_, numpy.complex],
ids=['float64', 'float32', 'int64', 'int32', 'bool', 'bool_', 'complex'])
@pytest.mark.parametrize("arr",

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

what if array is empty?
Corner cases.

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

outdated.

[[-2, -1, 0, 1, 2], [[-2, -1], [1, 2]], []],
ids=['[-2, -1, 0, 1, 2]', '[[-2, -1], [1, 2]]', '[]'])
def test_astype(arr, arr_dtype, res_dtype):
numpy_array = numpy.array(arr, dtype=arr_dtype)
dpnp_array = dpnp.array(numpy_array)
expected = numpy_array.astype(res_dtype)
result = dpnp_array.astype(res_dtype)
numpy.testing.assert_array_equal(expected, result)
1 change: 1 addition & 0 deletions tests_external/skipped_tests_numpy.tbl
Original file line number Diff line number Diff line change
Expand Up @@ -2881,6 +2881,7 @@ tests/test_umath.py::test_rint_big_int
tests/test_umath.py::TestRoundingFunctions::test_object_direct
tests/test_umath.py::TestRoundingFunctions::test_object_indirect
tests/test_umath.py::test_signaling_nan_exceptions
tests/test_umath.py::TestSign::test_sign_dtype_nan_object

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

what is the reason?
Provide info on description, please.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This test failed with error:
DPNP error: in function dpnp_dtype_to_DPNPFuncType() type '<class 'object'>' is not supported

Copy link

@samir-nasibli samir-nasibli Mar 10, 2021

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

object is a (global) variable. By default it is bound to a built-in class which is the root of the type hierarchy.

dpnp_dtype_to_DPNPFuncType can not convert object to dpnpc function types (DPNP_FT_DOUBLE, DPNP_FT_FLOAT, DPNP_FT_LONG, ...).
See this internal func impl here.

You can fallback on numpy for case when dtype is object.

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

You can fallback on numpy for case when dtype is object.

Sorry, my fault. It is bad proposal, just leave it as is.
No need doing fallback on numpy for dparray own methods.

tests/test_umath.py::TestSign::test_sign
tests/test_umath.py::TestSign::test_sign_dtype_object
tests/test_umath.py::test_spacing
Expand Down