diff --git a/dpctl/tensor/__init__.py b/dpctl/tensor/__init__.py index 3e470a736c..ec488cb3d4 100644 --- a/dpctl/tensor/__init__.py +++ b/dpctl/tensor/__init__.py @@ -111,6 +111,10 @@ less_equal, log, log1p, + logical_and, + logical_not, + logical_or, + logical_xor, multiply, not_equal, proj, @@ -211,6 +215,10 @@ "less", "less_equal", "log", + "logical_and", + "logical_not", + "logical_or", + "logical_xor", "log1p", "proj", "real", diff --git a/dpctl/tensor/_elementwise_funcs.py b/dpctl/tensor/_elementwise_funcs.py index dacefcc3f9..2c07ab8e6a 100644 --- a/dpctl/tensor/_elementwise_funcs.py +++ b/dpctl/tensor/_elementwise_funcs.py @@ -572,16 +572,116 @@ # FIXME: implement B15 # B16: ==== LOGICAL_AND (x1, x2) -# FIXME: implement B16 +_logical_and_docstring_ = """ +logical_and(x1, x2, out=None, order='K') + +Computes the logical AND for each element `x1_i` of the input array `x1` +with the respective element `x2_i` of the input array `x2`. + +Args: + x1 (usm_ndarray): + First input array. + x2 (usm_ndarray): + Second input array. + out ({None, usm_ndarray}, optional): + Output array to populate. + Array have the correct shape and the expected data type. + order ("C","F","A","K", optional): + Memory layout of the newly output array, if parameter `out` is `None`. + Default: "K". +Returns: + usm_narray: + An array containing the element-wise logical AND results. +""" +logical_and = BinaryElementwiseFunc( + "logical_and", + ti._logical_and_result_type, + ti._logical_and, + _logical_and_docstring_, +) # U24: ==== LOGICAL_NOT (x) -# FIXME: implement U24 +_logical_not_docstring = """ +logical_not(x, out=None, order='K') +Computes the logical NOT for each element `x_i` of input array `x`. +Args: + x (usm_ndarray): + Input array. + out (usm_ndarray): + Output array to populate. Array must have the correct + shape and the expected data type. + order ("C","F","A","K", optional): memory layout of the new + output array, if parameter `out` is `None`. + Default: "K". +Return: + usm_ndarray: + An array containing the element-wise logical NOT results. +""" + +logical_not = UnaryElementwiseFunc( + "logical_not", + ti._logical_not_result_type, + ti._logical_not, + _logical_not_docstring, +) # B17: ==== LOGICAL_OR (x1, x2) -# FIXME: implement B17 +_logical_or_docstring_ = """ +logical_or(x1, x2, out=None, order='K') + +Computes the logical OR for each element `x1_i` of the input array `x1` +with the respective element `x2_i` of the input array `x2`. + +Args: + x1 (usm_ndarray): + First input array. + x2 (usm_ndarray): + Second input array. + out ({None, usm_ndarray}, optional): + Output array to populate. + Array have the correct shape and the expected data type. + order ("C","F","A","K", optional): + Memory layout of the newly output array, if parameter `out` is `None`. + Default: "K". +Returns: + usm_narray: + An array containing the element-wise logical OR results. +""" +logical_or = BinaryElementwiseFunc( + "logical_or", + ti._logical_or_result_type, + ti._logical_or, + _logical_or_docstring_, +) # B18: ==== LOGICAL_XOR (x1, x2) -# FIXME: implement B18 +_logical_xor_docstring_ = """ +logical_xor(x1, x2, out=None, order='K') + +Computes the logical XOR for each element `x1_i` of the input array `x1` +with the respective element `x2_i` of the input array `x2`. + +Args: + x1 (usm_ndarray): + First input array. + x2 (usm_ndarray): + Second input array. + out ({None, usm_ndarray}, optional): + Output array to populate. + Array have the correct shape and the expected data type. + order ("C","F","A","K", optional): + Memory layout of the newly output array, if parameter `out` is `None`. + Default: "K". +Returns: + usm_narray: + An array containing the element-wise logical XOR results. +""" +logical_xor = BinaryElementwiseFunc( + "logical_xor", + ti._logical_xor_result_type, + ti._logical_xor, + _logical_xor_docstring_, +) # B19: ==== MULTIPLY (x1, x2) _multiply_docstring_ = """ diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_and.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_and.hpp new file mode 100644 index 0000000000..cb63b1e528 --- /dev/null +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_and.hpp @@ -0,0 +1,304 @@ +//=== logical_and.hpp - Binary function GREATER ------ +//*-C++-*--/===// +// +// Data Parallel Control (dpctl) +// +// Copyright 2020-2023 Intel Corporation +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain in1 copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// +//===---------------------------------------------------------------------===// +/// +/// \file +/// This file defines kernels for elementwise evaluation of comparison of +/// tensor elements. +//===---------------------------------------------------------------------===// + +#pragma once +#include +#include +#include +#include + +#include "utils/offset_utils.hpp" +#include "utils/type_dispatch.hpp" +#include "utils/type_utils.hpp" + +#include "kernels/elementwise_functions/common.hpp" +#include + +namespace dpctl +{ +namespace tensor +{ +namespace kernels +{ +namespace logical_and +{ + +namespace py = pybind11; +namespace td_ns = dpctl::tensor::type_dispatch; +namespace tu_ns = dpctl::tensor::type_utils; + +template +struct LogicalAndFunctor +{ + static_assert(std::is_same_v); + + using supports_sg_loadstore = std::negation< + std::disjunction, tu_ns::is_complex>>; + using supports_vec = std::conjunction< + std::is_same, + std::negation, + tu_ns::is_complex>>>; + + resT operator()(const argT1 &in1, const argT2 &in2) + { + using tu_ns::convert_impl; + + return (convert_impl(in1) && + convert_impl(in2)); + } + + template + sycl::vec operator()(const sycl::vec &in1, + const sycl::vec &in2) + { + + auto tmp = (in1 && in2); + if constexpr (std::is_same_v) { + return tmp; + } + else { + using dpctl::tensor::type_utils::vec_cast; + + return vec_cast( + tmp); + } + } +}; + +template +using LogicalAndContigFunctor = elementwise_common::BinaryContigFunctor< + argT1, + argT2, + resT, + LogicalAndFunctor, + vec_sz, + n_vecs>; + +template +using LogicalAndStridedFunctor = elementwise_common::BinaryStridedFunctor< + argT1, + argT2, + resT, + IndexerT, + LogicalAndFunctor>; + +template struct LogicalAndOutputType +{ + using value_type = typename std::disjunction< // disjunction is C++17 + // feature, supported by + // DPC++ + td_ns::BinaryTypeMapResultEntry, + td_ns:: + BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns:: + BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns:: + BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns:: + BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + T2, + std::complex, + bool>, + td_ns::BinaryTypeMapResultEntry, + T2, + std::complex, + bool>, + td_ns:: + BinaryTypeMapResultEntry, bool>, + td_ns:: + BinaryTypeMapResultEntry, T2, float, bool>, + td_ns::DefaultResultEntry>::result_type; +}; + +template +class logical_and_contig_kernel; + +template +sycl::event +logical_and_contig_impl(sycl::queue exec_q, + size_t nelems, + const char *arg1_p, + py::ssize_t arg1_offset, + const char *arg2_p, + py::ssize_t arg2_offset, + char *res_p, + py::ssize_t res_offset, + const std::vector &depends = {}) +{ + sycl::event comp_ev = exec_q.submit([&](sycl::handler &cgh) { + cgh.depends_on(depends); + + size_t lws = 64; + constexpr unsigned int vec_sz = 4; + constexpr unsigned int n_vecs = 2; + const size_t n_groups = + ((nelems + lws * n_vecs * vec_sz - 1) / (lws * n_vecs * vec_sz)); + const auto gws_range = sycl::range<1>(n_groups * lws); + const auto lws_range = sycl::range<1>(lws); + + using resTy = typename LogicalAndOutputType::value_type; + + const argTy1 *arg1_tp = + reinterpret_cast(arg1_p) + arg1_offset; + const argTy2 *arg2_tp = + reinterpret_cast(arg2_p) + arg2_offset; + resTy *res_tp = reinterpret_cast(res_p) + res_offset; + + cgh.parallel_for< + logical_and_contig_kernel>( + sycl::nd_range<1>(gws_range, lws_range), + LogicalAndContigFunctor( + arg1_tp, arg2_tp, res_tp, nelems)); + }); + return comp_ev; +} + +template struct LogicalAndContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename LogicalAndOutputType::value_type, + void>) + { + fnT fn = nullptr; + return fn; + } + else { + fnT fn = logical_and_contig_impl; + return fn; + } + } +}; + +template +struct LogicalAndTypeMapFactory +{ + /*! @brief get typeid for output type of operator()>(x, y), always bool + */ + std::enable_if_t::value, int> get() + { + using rT = typename LogicalAndOutputType::value_type; + return td_ns::GetTypeid{}.get(); + } +}; + +template +class logical_and_strided_kernel; + +template +sycl::event +logical_and_strided_impl(sycl::queue exec_q, + size_t nelems, + int nd, + const py::ssize_t *shape_and_strides, + const char *arg1_p, + py::ssize_t arg1_offset, + const char *arg2_p, + py::ssize_t arg2_offset, + char *res_p, + py::ssize_t res_offset, + const std::vector &depends, + const std::vector &additional_depends) +{ + sycl::event comp_ev = exec_q.submit([&](sycl::handler &cgh) { + cgh.depends_on(depends); + cgh.depends_on(additional_depends); + + using resTy = typename LogicalAndOutputType::value_type; + + using IndexerT = + typename dpctl::tensor::offset_utils::ThreeOffsets_StridedIndexer; + + IndexerT indexer{nd, arg1_offset, arg2_offset, res_offset, + shape_and_strides}; + + const argTy1 *arg1_tp = reinterpret_cast(arg1_p); + const argTy2 *arg2_tp = reinterpret_cast(arg2_p); + resTy *res_tp = reinterpret_cast(res_p); + + cgh.parallel_for< + logical_and_strided_kernel>( + {nelems}, LogicalAndStridedFunctor( + arg1_tp, arg2_tp, res_tp, indexer)); + }); + return comp_ev; +} + +template +struct LogicalAndStridedFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename LogicalAndOutputType::value_type, + void>) + { + fnT fn = nullptr; + return fn; + } + else { + fnT fn = logical_and_strided_impl; + return fn; + } + } +}; + +} // namespace logical_and +} // namespace kernels +} // namespace tensor +} // namespace dpctl diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_not.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_not.hpp new file mode 100644 index 0000000000..1062950461 --- /dev/null +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_not.hpp @@ -0,0 +1,162 @@ +//=== logical_not.hpp - Unary function ISNAN ------ +//*-C++-*--/===// +// +// Data Parallel Control (dpctl) +// +// Copyright 2020-2023 Intel Corporation +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// +//===---------------------------------------------------------------------===// +/// +/// \file +/// This file defines kernels for elementwise evaluation of ISNAN(x) +/// function that tests whether a tensor element is a NaN. +//===---------------------------------------------------------------------===// + +#pragma once +#include +#include +#include +#include + +#include "utils/offset_utils.hpp" +#include "utils/type_dispatch.hpp" +#include "utils/type_utils.hpp" +#include + +namespace dpctl +{ +namespace tensor +{ +namespace kernels +{ +namespace logical_not +{ + +namespace py = pybind11; +namespace td_ns = dpctl::tensor::type_dispatch; +namespace tu_ns = dpctl::tensor::type_utils; + +template struct LogicalNotFunctor +{ + static_assert(std::is_same_v); + + using is_constant = typename std::false_type; + // constexpr resT constant_value = resT{}; + using supports_vec = typename std::false_type; + using supports_sg_loadstore = typename std::negation< + std::disjunction, tu_ns::is_complex>>; + + resT operator()(const argT &in) const + { + using tu_ns::convert_impl; + return !convert_impl(in); + } +}; + +template +using LogicalNotContigFunctor = + elementwise_common::UnaryContigFunctor, + vec_sz, + n_vecs>; + +template +using LogicalNotStridedFunctor = + elementwise_common::UnaryStridedFunctor>; + +template struct LogicalNotOutputType +{ + using value_type = bool; +}; + +template +class logical_not_contig_kernel; + +template +sycl::event +logical_not_contig_impl(sycl::queue exec_q, + size_t nelems, + const char *arg_p, + char *res_p, + const std::vector &depends = {}) +{ + return elementwise_common::unary_contig_impl( + exec_q, nelems, arg_p, res_p, depends); +} + +template struct LogicalNotContigFactory +{ + fnT get() + { + fnT fn = logical_not_contig_impl; + return fn; + } +}; + +template struct LogicalNotTypeMapFactory +{ + /*! @brief get typeid for output type of sycl::logical_not(T x) */ + std::enable_if_t::value, int> get() + { + using rT = typename LogicalNotOutputType::value_type; + return td_ns::GetTypeid{}.get(); + } +}; + +template +class logical_not_strided_kernel; + +template +sycl::event +logical_not_strided_impl(sycl::queue exec_q, + size_t nelems, + int nd, + const py::ssize_t *shape_and_strides, + const char *arg_p, + py::ssize_t arg_offset, + char *res_p, + py::ssize_t res_offset, + const std::vector &depends, + const std::vector &additional_depends) +{ + return elementwise_common::unary_strided_impl( + exec_q, nelems, nd, shape_and_strides, arg_p, arg_offset, res_p, + res_offset, depends, additional_depends); +} + +template struct LogicalNotStridedFactory +{ + fnT get() + { + fnT fn = logical_not_strided_impl; + return fn; + } +}; + +} // namespace logical_not +} // namespace kernels +} // namespace tensor +} // namespace dpctl diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_or.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_or.hpp new file mode 100644 index 0000000000..ef77946ef3 --- /dev/null +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_or.hpp @@ -0,0 +1,300 @@ +//=== logical_or.hpp - Binary function GREATER ------ +//*-C++-*--/===// +// +// Data Parallel Control (dpctl) +// +// Copyright 2020-2023 Intel Corporation +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain in1 copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// +//===---------------------------------------------------------------------===// +/// +/// \file +/// This file defines kernels for elementwise evaluation of comparison of +/// tensor elements. +//===---------------------------------------------------------------------===// + +#pragma once +#include +#include +#include +#include + +#include "utils/offset_utils.hpp" +#include "utils/type_dispatch.hpp" +#include "utils/type_utils.hpp" + +#include "kernels/elementwise_functions/common.hpp" +#include + +namespace dpctl +{ +namespace tensor +{ +namespace kernels +{ +namespace logical_or +{ + +namespace py = pybind11; +namespace td_ns = dpctl::tensor::type_dispatch; +namespace tu_ns = dpctl::tensor::type_utils; + +template struct LogicalOrFunctor +{ + static_assert(std::is_same_v); + + using supports_sg_loadstore = std::negation< + std::disjunction, tu_ns::is_complex>>; + using supports_vec = std::conjunction< + std::is_same, + std::negation, + tu_ns::is_complex>>>; + + resT operator()(const argT1 &in1, const argT2 &in2) + { + using tu_ns::convert_impl; + + return (convert_impl(in1) || + convert_impl(in2)); + } + + template + sycl::vec operator()(const sycl::vec &in1, + const sycl::vec &in2) + { + + auto tmp = (in1 || in2); + if constexpr (std::is_same_v) { + return tmp; + } + else { + using dpctl::tensor::type_utils::vec_cast; + + return vec_cast( + tmp); + } + } +}; + +template +using LogicalOrContigFunctor = elementwise_common::BinaryContigFunctor< + argT1, + argT2, + resT, + LogicalOrFunctor, + vec_sz, + n_vecs>; + +template +using LogicalOrStridedFunctor = elementwise_common::BinaryStridedFunctor< + argT1, + argT2, + resT, + IndexerT, + LogicalOrFunctor>; + +template struct LogicalOrOutputType +{ + using value_type = typename std::disjunction< // disjunction is C++17 + // feature, supported by + // DPC++ + td_ns::BinaryTypeMapResultEntry, + td_ns:: + BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns:: + BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns:: + BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns:: + BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + T2, + std::complex, + bool>, + td_ns::BinaryTypeMapResultEntry, + T2, + std::complex, + bool>, + td_ns:: + BinaryTypeMapResultEntry, bool>, + td_ns:: + BinaryTypeMapResultEntry, T2, float, bool>, + td_ns::DefaultResultEntry>::result_type; +}; + +template +class logical_or_contig_kernel; + +template +sycl::event logical_or_contig_impl(sycl::queue exec_q, + size_t nelems, + const char *arg1_p, + py::ssize_t arg1_offset, + const char *arg2_p, + py::ssize_t arg2_offset, + char *res_p, + py::ssize_t res_offset, + const std::vector &depends = {}) +{ + sycl::event comp_ev = exec_q.submit([&](sycl::handler &cgh) { + cgh.depends_on(depends); + + size_t lws = 64; + constexpr unsigned int vec_sz = 4; + constexpr unsigned int n_vecs = 2; + const size_t n_groups = + ((nelems + lws * n_vecs * vec_sz - 1) / (lws * n_vecs * vec_sz)); + const auto gws_range = sycl::range<1>(n_groups * lws); + const auto lws_range = sycl::range<1>(lws); + + using resTy = typename LogicalOrOutputType::value_type; + + const argTy1 *arg1_tp = + reinterpret_cast(arg1_p) + arg1_offset; + const argTy2 *arg2_tp = + reinterpret_cast(arg2_p) + arg2_offset; + resTy *res_tp = reinterpret_cast(res_p) + res_offset; + + cgh.parallel_for< + logical_or_contig_kernel>( + sycl::nd_range<1>(gws_range, lws_range), + LogicalOrContigFunctor( + arg1_tp, arg2_tp, res_tp, nelems)); + }); + return comp_ev; +} + +template struct LogicalOrContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename LogicalOrOutputType::value_type, + void>) + { + fnT fn = nullptr; + return fn; + } + else { + fnT fn = logical_or_contig_impl; + return fn; + } + } +}; + +template struct LogicalOrTypeMapFactory +{ + /*! @brief get typeid for output type of operator()>(x, y), always bool + */ + std::enable_if_t::value, int> get() + { + using rT = typename LogicalOrOutputType::value_type; + return td_ns::GetTypeid{}.get(); + } +}; + +template +class logical_or_strided_kernel; + +template +sycl::event +logical_or_strided_impl(sycl::queue exec_q, + size_t nelems, + int nd, + const py::ssize_t *shape_and_strides, + const char *arg1_p, + py::ssize_t arg1_offset, + const char *arg2_p, + py::ssize_t arg2_offset, + char *res_p, + py::ssize_t res_offset, + const std::vector &depends, + const std::vector &additional_depends) +{ + sycl::event comp_ev = exec_q.submit([&](sycl::handler &cgh) { + cgh.depends_on(depends); + cgh.depends_on(additional_depends); + + using resTy = typename LogicalOrOutputType::value_type; + + using IndexerT = + typename dpctl::tensor::offset_utils::ThreeOffsets_StridedIndexer; + + IndexerT indexer{nd, arg1_offset, arg2_offset, res_offset, + shape_and_strides}; + + const argTy1 *arg1_tp = reinterpret_cast(arg1_p); + const argTy2 *arg2_tp = reinterpret_cast(arg2_p); + resTy *res_tp = reinterpret_cast(res_p); + + cgh.parallel_for< + logical_or_strided_kernel>( + {nelems}, LogicalOrStridedFunctor( + arg1_tp, arg2_tp, res_tp, indexer)); + }); + return comp_ev; +} + +template struct LogicalOrStridedFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename LogicalOrOutputType::value_type, + void>) + { + fnT fn = nullptr; + return fn; + } + else { + fnT fn = logical_or_strided_impl; + return fn; + } + } +}; + +} // namespace logical_or +} // namespace kernels +} // namespace tensor +} // namespace dpctl diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_xor.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_xor.hpp new file mode 100644 index 0000000000..16c925eb99 --- /dev/null +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logical_xor.hpp @@ -0,0 +1,305 @@ +//=== logical_xor.hpp - Binary function GREATER ------ +//*-C++-*--/===// +// +// Data Parallel Control (dpctl) +// +// Copyright 2020-2023 Intel Corporation +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain in1 copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// +//===---------------------------------------------------------------------===// +/// +/// \file +/// This file defines kernels for elementwise evaluation of comparison of +/// tensor elements. +//===---------------------------------------------------------------------===// + +#pragma once +#include +#include +#include +#include + +#include "utils/offset_utils.hpp" +#include "utils/type_dispatch.hpp" +#include "utils/type_utils.hpp" + +#include "kernels/elementwise_functions/common.hpp" +#include + +namespace dpctl +{ +namespace tensor +{ +namespace kernels +{ +namespace logical_xor +{ + +namespace py = pybind11; +namespace td_ns = dpctl::tensor::type_dispatch; +namespace tu_ns = dpctl::tensor::type_utils; + +template +struct LogicalXorFunctor +{ + static_assert(std::is_same_v); + + using supports_sg_loadstore = std::negation< + std::disjunction, tu_ns::is_complex>>; + using supports_vec = std::conjunction< + std::is_same, + std::negation, + tu_ns::is_complex>>>; + + resT operator()(const argT1 &in1, const argT2 &in2) + { + using tu_ns::convert_impl; + + return (convert_impl(in1) != + convert_impl(in2)); + } + + template + sycl::vec operator()(const sycl::vec &in1, + const sycl::vec &in2) + { + using tu_ns::vec_cast; + auto tmp1 = vec_cast(in1); + auto tmp2 = vec_cast(in2); + + auto tmp = (tmp1 != tmp2); + if constexpr (std::is_same_v) { + return tmp; + } + else { + return vec_cast( + tmp); + } + } +}; + +template +using LogicalXorContigFunctor = elementwise_common::BinaryContigFunctor< + argT1, + argT2, + resT, + LogicalXorFunctor, + vec_sz, + n_vecs>; + +template +using LogicalXorStridedFunctor = elementwise_common::BinaryStridedFunctor< + argT1, + argT2, + resT, + IndexerT, + LogicalXorFunctor>; + +template struct LogicalXorOutputType +{ + using value_type = typename std::disjunction< // disjunction is C++17 + // feature, supported by + // DPC++ + td_ns::BinaryTypeMapResultEntry, + td_ns:: + BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns:: + BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns:: + BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns:: + BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + td_ns::BinaryTypeMapResultEntry, + T2, + std::complex, + bool>, + td_ns::BinaryTypeMapResultEntry, + T2, + std::complex, + bool>, + td_ns:: + BinaryTypeMapResultEntry, bool>, + td_ns:: + BinaryTypeMapResultEntry, T2, float, bool>, + td_ns::DefaultResultEntry>::result_type; +}; + +template +class logical_xor_contig_kernel; + +template +sycl::event +logical_xor_contig_impl(sycl::queue exec_q, + size_t nelems, + const char *arg1_p, + py::ssize_t arg1_offset, + const char *arg2_p, + py::ssize_t arg2_offset, + char *res_p, + py::ssize_t res_offset, + const std::vector &depends = {}) +{ + sycl::event comp_ev = exec_q.submit([&](sycl::handler &cgh) { + cgh.depends_on(depends); + + size_t lws = 64; + constexpr unsigned int vec_sz = 4; + constexpr unsigned int n_vecs = 2; + const size_t n_groups = + ((nelems + lws * n_vecs * vec_sz - 1) / (lws * n_vecs * vec_sz)); + const auto gws_range = sycl::range<1>(n_groups * lws); + const auto lws_range = sycl::range<1>(lws); + + using resTy = typename LogicalXorOutputType::value_type; + + const argTy1 *arg1_tp = + reinterpret_cast(arg1_p) + arg1_offset; + const argTy2 *arg2_tp = + reinterpret_cast(arg2_p) + arg2_offset; + resTy *res_tp = reinterpret_cast(res_p) + res_offset; + + cgh.parallel_for< + logical_xor_contig_kernel>( + sycl::nd_range<1>(gws_range, lws_range), + LogicalXorContigFunctor( + arg1_tp, arg2_tp, res_tp, nelems)); + }); + return comp_ev; +} + +template struct LogicalXorContigFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename LogicalXorOutputType::value_type, + void>) + { + fnT fn = nullptr; + return fn; + } + else { + fnT fn = logical_xor_contig_impl; + return fn; + } + } +}; + +template +struct LogicalXorTypeMapFactory +{ + /*! @brief get typeid for output type of operator()>(x, y), always bool + */ + std::enable_if_t::value, int> get() + { + using rT = typename LogicalXorOutputType::value_type; + return td_ns::GetTypeid{}.get(); + } +}; + +template +class logical_xor_strided_kernel; + +template +sycl::event +logical_xor_strided_impl(sycl::queue exec_q, + size_t nelems, + int nd, + const py::ssize_t *shape_and_strides, + const char *arg1_p, + py::ssize_t arg1_offset, + const char *arg2_p, + py::ssize_t arg2_offset, + char *res_p, + py::ssize_t res_offset, + const std::vector &depends, + const std::vector &additional_depends) +{ + sycl::event comp_ev = exec_q.submit([&](sycl::handler &cgh) { + cgh.depends_on(depends); + cgh.depends_on(additional_depends); + + using resTy = typename LogicalXorOutputType::value_type; + + using IndexerT = + typename dpctl::tensor::offset_utils::ThreeOffsets_StridedIndexer; + + IndexerT indexer{nd, arg1_offset, arg2_offset, res_offset, + shape_and_strides}; + + const argTy1 *arg1_tp = reinterpret_cast(arg1_p); + const argTy2 *arg2_tp = reinterpret_cast(arg2_p); + resTy *res_tp = reinterpret_cast(res_p); + + cgh.parallel_for< + logical_xor_strided_kernel>( + {nelems}, LogicalXorStridedFunctor( + arg1_tp, arg2_tp, res_tp, indexer)); + }); + return comp_ev; +} + +template +struct LogicalXorStridedFactory +{ + fnT get() + { + if constexpr (std::is_same_v< + typename LogicalXorOutputType::value_type, + void>) + { + fnT fn = nullptr; + return fn; + } + else { + fnT fn = logical_xor_strided_impl; + return fn; + } + } +}; + +} // namespace logical_xor +} // namespace kernels +} // namespace tensor +} // namespace dpctl diff --git a/dpctl/tensor/libtensor/source/elementwise_functions.cpp b/dpctl/tensor/libtensor/source/elementwise_functions.cpp index f9c74f652d..5898f0ca7d 100644 --- a/dpctl/tensor/libtensor/source/elementwise_functions.cpp +++ b/dpctl/tensor/libtensor/source/elementwise_functions.cpp @@ -50,6 +50,10 @@ #include "kernels/elementwise_functions/less_equal.hpp" #include "kernels/elementwise_functions/log.hpp" #include "kernels/elementwise_functions/log1p.hpp" +#include "kernels/elementwise_functions/logical_and.hpp" +#include "kernels/elementwise_functions/logical_not.hpp" +#include "kernels/elementwise_functions/logical_or.hpp" +#include "kernels/elementwise_functions/logical_xor.hpp" #include "kernels/elementwise_functions/multiply.hpp" #include "kernels/elementwise_functions/not_equal.hpp" #include "kernels/elementwise_functions/proj.hpp" @@ -1024,25 +1028,149 @@ namespace impl // B16: ==== LOGICAL_AND (x1, x2) namespace impl { -// FIXME: add code for B16 +namespace logical_and_fn_ns = dpctl::tensor::kernels::logical_and; + +static binary_contig_impl_fn_ptr_t + logical_and_contig_dispatch_table[td_ns::num_types][td_ns::num_types]; +static int logical_and_output_id_table[td_ns::num_types][td_ns::num_types]; + +static binary_strided_impl_fn_ptr_t + logical_and_strided_dispatch_table[td_ns::num_types][td_ns::num_types]; + +void populate_logical_and_dispatch_tables(void) +{ + using namespace td_ns; + namespace fn_ns = logical_and_fn_ns; + + // which input types are supported, and what is the type of the result + using fn_ns::LogicalAndTypeMapFactory; + DispatchTableBuilder dtb1; + dtb1.populate_dispatch_table(logical_and_output_id_table); + + // function pointers for operation on general strided arrays + using fn_ns::LogicalAndStridedFactory; + DispatchTableBuilder + dtb2; + dtb2.populate_dispatch_table(logical_and_strided_dispatch_table); + + // function pointers for operation on contiguous inputs and output + using fn_ns::LogicalAndContigFactory; + DispatchTableBuilder + dtb3; + dtb3.populate_dispatch_table(logical_and_contig_dispatch_table); +}; } // namespace impl // U24: ==== LOGICAL_NOT (x) namespace impl { -// FIXME: add code for U24 +namespace logical_not_fn_ns = dpctl::tensor::kernels::logical_not; + +static unary_contig_impl_fn_ptr_t + logical_not_contig_dispatch_vector[td_ns::num_types]; +static int logical_not_output_typeid_vector[td_ns::num_types]; +static unary_strided_impl_fn_ptr_t + logical_not_strided_dispatch_vector[td_ns::num_types]; + +void populate_logical_not_dispatch_vectors(void) +{ + using namespace td_ns; + namespace fn_ns = logical_not_fn_ns; + + using fn_ns::LogicalNotContigFactory; + DispatchVectorBuilder + dvb1; + dvb1.populate_dispatch_vector(logical_not_contig_dispatch_vector); + + using fn_ns::LogicalNotStridedFactory; + DispatchVectorBuilder + dvb2; + dvb2.populate_dispatch_vector(logical_not_strided_dispatch_vector); + + using fn_ns::LogicalNotTypeMapFactory; + DispatchVectorBuilder dvb3; + dvb3.populate_dispatch_vector(logical_not_output_typeid_vector); +}; } // namespace impl // B17: ==== LOGICAL_OR (x1, x2) namespace impl { -// FIXME: add code for B17 +namespace logical_or_fn_ns = dpctl::tensor::kernels::logical_or; + +static binary_contig_impl_fn_ptr_t + logical_or_contig_dispatch_table[td_ns::num_types][td_ns::num_types]; +static int logical_or_output_id_table[td_ns::num_types][td_ns::num_types]; + +static binary_strided_impl_fn_ptr_t + logical_or_strided_dispatch_table[td_ns::num_types][td_ns::num_types]; + +void populate_logical_or_dispatch_tables(void) +{ + using namespace td_ns; + namespace fn_ns = logical_or_fn_ns; + + // which input types are supported, and what is the type of the result + using fn_ns::LogicalOrTypeMapFactory; + DispatchTableBuilder dtb1; + dtb1.populate_dispatch_table(logical_or_output_id_table); + + // function pointers for operation on general strided arrays + using fn_ns::LogicalOrStridedFactory; + DispatchTableBuilder + dtb2; + dtb2.populate_dispatch_table(logical_or_strided_dispatch_table); + + // function pointers for operation on contiguous inputs and output + using fn_ns::LogicalOrContigFactory; + DispatchTableBuilder + dtb3; + dtb3.populate_dispatch_table(logical_or_contig_dispatch_table); +}; } // namespace impl // B18: ==== LOGICAL_XOR (x1, x2) namespace impl { -// FIXME: add code for B18 +namespace logical_xor_fn_ns = dpctl::tensor::kernels::logical_xor; + +static binary_contig_impl_fn_ptr_t + logical_xor_contig_dispatch_table[td_ns::num_types][td_ns::num_types]; +static int logical_xor_output_id_table[td_ns::num_types][td_ns::num_types]; + +static binary_strided_impl_fn_ptr_t + logical_xor_strided_dispatch_table[td_ns::num_types][td_ns::num_types]; + +void populate_logical_xor_dispatch_tables(void) +{ + using namespace td_ns; + namespace fn_ns = logical_xor_fn_ns; + + // which input types are supported, and what is the type of the result + using fn_ns::LogicalXorTypeMapFactory; + DispatchTableBuilder dtb1; + dtb1.populate_dispatch_table(logical_xor_output_id_table); + + // function pointers for operation on general strided arrays + using fn_ns::LogicalXorStridedFactory; + DispatchTableBuilder + dtb2; + dtb2.populate_dispatch_table(logical_xor_strided_dispatch_table); + + // function pointers for operation on contiguous inputs and output + using fn_ns::LogicalXorContigFactory; + DispatchTableBuilder + dtb3; + dtb3.populate_dispatch_table(logical_xor_contig_dispatch_table); +}; } // namespace impl // B19: ==== MULTIPLY (x1, x2) @@ -2152,16 +2280,152 @@ void init_elementwise_functions(py::module_ m) // FIXME: // B16: ==== LOGICAL_AND (x1, x2) - // FIXME: + { + impl::populate_logical_and_dispatch_tables(); + using impl::logical_and_contig_dispatch_table; + using impl::logical_and_output_id_table; + using impl::logical_and_strided_dispatch_table; + + auto logical_and_pyapi = [&](dpctl::tensor::usm_ndarray src1, + dpctl::tensor::usm_ndarray src2, + dpctl::tensor::usm_ndarray dst, + sycl::queue exec_q, + const std::vector &depends = + {}) { + return py_binary_ufunc( + src1, src2, dst, exec_q, depends, logical_and_output_id_table, + // function pointers to handle operation on contiguous arrays + // (pointers may be nullptr) + logical_and_contig_dispatch_table, + // function pointers to handle operation on strided arrays (most + // general case) + logical_and_strided_dispatch_table, + // function pointers to handle operation of c-contig matrix and + // c-contig row with broadcasting (may be nullptr) + td_ns::NullPtrTable< + binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t>{}, + // function pointers to handle operation of c-contig matrix and + // c-contig row with broadcasting (may be nullptr) + td_ns::NullPtrTable< + binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t>{}); + }; + auto logical_and_result_type_pyapi = [&](py::dtype dtype1, + py::dtype dtype2) { + return py_binary_ufunc_result_type(dtype1, dtype2, + logical_and_output_id_table); + }; + m.def("_logical_and", logical_and_pyapi, "", py::arg("src1"), + py::arg("src2"), py::arg("dst"), py::arg("sycl_queue"), + py::arg("depends") = py::list()); + m.def("_logical_and_result_type", logical_and_result_type_pyapi, ""); + } // U24: ==== LOGICAL_NOT (x) - // FIXME: + { + impl::populate_logical_not_dispatch_vectors(); + using impl::logical_not_contig_dispatch_vector; + using impl::logical_not_output_typeid_vector; + using impl::logical_not_strided_dispatch_vector; + + auto logical_not_pyapi = [&](arrayT src, arrayT dst, sycl::queue exec_q, + const event_vecT &depends = {}) { + return py_unary_ufunc(src, dst, exec_q, depends, + logical_not_output_typeid_vector, + logical_not_contig_dispatch_vector, + logical_not_strided_dispatch_vector); + }; + m.def("_logical_not", logical_not_pyapi, "", py::arg("src"), + py::arg("dst"), py::arg("sycl_queue"), + py::arg("depends") = py::list()); + + auto logical_not_result_type_pyapi = [&](py::dtype dtype) { + return py_unary_ufunc_result_type(dtype, + logical_not_output_typeid_vector); + }; + m.def("_logical_not_result_type", logical_not_result_type_pyapi); + } // B17: ==== LOGICAL_OR (x1, x2) - // FIXME: + { + impl::populate_logical_or_dispatch_tables(); + using impl::logical_or_contig_dispatch_table; + using impl::logical_or_output_id_table; + using impl::logical_or_strided_dispatch_table; + + auto logical_or_pyapi = [&](dpctl::tensor::usm_ndarray src1, + dpctl::tensor::usm_ndarray src2, + dpctl::tensor::usm_ndarray dst, + sycl::queue exec_q, + const std::vector &depends = + {}) { + return py_binary_ufunc( + src1, src2, dst, exec_q, depends, logical_or_output_id_table, + // function pointers to handle operation on contiguous arrays + // (pointers may be nullptr) + logical_or_contig_dispatch_table, + // function pointers to handle operation on strided arrays (most + // general case) + logical_or_strided_dispatch_table, + // function pointers to handle operation of c-contig matrix and + // c-contig row with broadcasting (may be nullptr) + td_ns::NullPtrTable< + binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t>{}, + // function pointers to handle operation of c-contig matrix and + // c-contig row with broadcasting (may be nullptr) + td_ns::NullPtrTable< + binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t>{}); + }; + auto logical_or_result_type_pyapi = [&](py::dtype dtype1, + py::dtype dtype2) { + return py_binary_ufunc_result_type(dtype1, dtype2, + logical_or_output_id_table); + }; + m.def("_logical_or", logical_or_pyapi, "", py::arg("src1"), + py::arg("src2"), py::arg("dst"), py::arg("sycl_queue"), + py::arg("depends") = py::list()); + m.def("_logical_or_result_type", logical_or_result_type_pyapi, ""); + } // B18: ==== LOGICAL_XOR (x1, x2) - // FIXME: + { + impl::populate_logical_xor_dispatch_tables(); + using impl::logical_xor_contig_dispatch_table; + using impl::logical_xor_output_id_table; + using impl::logical_xor_strided_dispatch_table; + + auto logical_xor_pyapi = [&](dpctl::tensor::usm_ndarray src1, + dpctl::tensor::usm_ndarray src2, + dpctl::tensor::usm_ndarray dst, + sycl::queue exec_q, + const std::vector &depends = + {}) { + return py_binary_ufunc( + src1, src2, dst, exec_q, depends, logical_xor_output_id_table, + // function pointers to handle operation on contiguous arrays + // (pointers may be nullptr) + logical_xor_contig_dispatch_table, + // function pointers to handle operation on strided arrays (most + // general case) + logical_xor_strided_dispatch_table, + // function pointers to handle operation of c-contig matrix and + // c-contig row with broadcasting (may be nullptr) + td_ns::NullPtrTable< + binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t>{}, + // function pointers to handle operation of c-contig matrix and + // c-contig row with broadcasting (may be nullptr) + td_ns::NullPtrTable< + binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t>{}); + }; + auto logical_xor_result_type_pyapi = [&](py::dtype dtype1, + py::dtype dtype2) { + return py_binary_ufunc_result_type(dtype1, dtype2, + logical_xor_output_id_table); + }; + m.def("_logical_xor", logical_xor_pyapi, "", py::arg("src1"), + py::arg("src2"), py::arg("dst"), py::arg("sycl_queue"), + py::arg("depends") = py::list()); + m.def("_logical_xor_result_type", logical_xor_result_type_pyapi, ""); + } // B19: ==== MULTIPLY (x1, x2) { diff --git a/dpctl/tests/elementwise/test_logical_and.py b/dpctl/tests/elementwise/test_logical_and.py new file mode 100644 index 0000000000..12a35b06d6 --- /dev/null +++ b/dpctl/tests/elementwise/test_logical_and.py @@ -0,0 +1,294 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2023 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless_equal required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +import ctypes + +import numpy as np +import pytest + +import dpctl +import dpctl.tensor as dpt +from dpctl.tests.helper import get_queue_or_skip, skip_if_dtype_not_supported + +from .utils import _all_dtypes, _compare_dtypes, _usm_types + + +@pytest.mark.parametrize("op1_dtype", _all_dtypes) +@pytest.mark.parametrize("op2_dtype", _all_dtypes) +def test_logical_and_dtype_matrix(op1_dtype, op2_dtype): + q = get_queue_or_skip() + skip_if_dtype_not_supported(op1_dtype, q) + skip_if_dtype_not_supported(op2_dtype, q) + + sz = 127 + ar1 = dpt.asarray(np.random.randint(0, 2, sz), dtype=op1_dtype) + ar2 = dpt.asarray(np.random.randint(0, 2, sz), dtype=op2_dtype) + + r = dpt.logical_and(ar1, ar2) + assert isinstance(r, dpt.usm_ndarray) + + expected = np.logical_and(dpt.asnumpy(ar1), dpt.asnumpy(ar2)) + assert _compare_dtypes(r.dtype, expected.dtype, sycl_queue=q) + assert r.shape == ar1.shape + assert (dpt.asnumpy(r) == expected).all() + assert r.sycl_queue == ar1.sycl_queue + + r2 = dpt.empty_like(r, dtype=r.dtype) + dpt.logical_and(ar1, ar2, out=r2) + assert (dpt.asnumpy(r) == dpt.asnumpy(r2)).all() + + ar3 = dpt.zeros(sz, dtype=op1_dtype) + ar4 = dpt.ones(2 * sz, dtype=op2_dtype) + + r = dpt.logical_and(ar3[::-1], ar4[::2]) + assert isinstance(r, dpt.usm_ndarray) + expected = np.logical_and( + np.zeros(1, dtype=op1_dtype), np.ones(1, dtype=op2_dtype) + ) + assert _compare_dtypes(r.dtype, expected.dtype, sycl_queue=q) + assert r.shape == ar3.shape + assert (dpt.asnumpy(r) == expected).all() + + r2 = dpt.empty_like(r, dtype=r.dtype) + dpt.logical_and(ar3[::-1], ar4[::2], out=r2) + assert (dpt.asnumpy(r) == dpt.asnumpy(r2)).all() + + +@pytest.mark.parametrize("op_dtype", ["c8", "c16"]) +def test_logical_and_complex_matrix(op_dtype): + q = get_queue_or_skip() + skip_if_dtype_not_supported(op_dtype, q) + + sz = 127 + ar1_np_real = np.random.randint(0, 2, sz) + ar1_np_imag = np.random.randint(0, 2, sz) + ar1 = dpt.asarray(ar1_np_real + 1j * ar1_np_imag, dtype=op_dtype) + + ar2_np_real = np.random.randint(0, 2, sz) + ar2_np_imag = np.random.randint(0, 2, sz) + ar2 = dpt.asarray(ar2_np_real + 1j * ar2_np_imag, dtype=op_dtype) + + r = dpt.logical_and(ar1, ar2) + expected = np.logical_and(dpt.asnumpy(ar1), dpt.asnumpy(ar2)) + assert _compare_dtypes(r.dtype, expected.dtype, sycl_queue=q) + assert r.shape == expected.shape + assert (dpt.asnumpy(r) == expected).all() + + r1 = dpt.logical_and(ar1[::-2], ar2[::2]) + expected1 = np.logical_and(dpt.asnumpy(ar1[::-2]), dpt.asnumpy(ar2[::2])) + assert _compare_dtypes(r.dtype, expected1.dtype, sycl_queue=q) + assert r1.shape == expected1.shape + assert (dpt.asnumpy(r1) == expected1).all() + + ar3 = dpt.asarray( + [ + 2.0 + 0j, + dpt.nan, + dpt.nan * 1j, + dpt.inf, + dpt.inf * 1j, + -dpt.inf, + -dpt.inf * 1j, + ], + dtype=op_dtype, + ) + ar4 = dpt.full(ar3.shape, fill_value=1.0 + 2j, dtype=op_dtype) + r2 = dpt.logical_and(ar3, ar4) + with np.errstate(invalid="ignore"): + expected2 = np.logical_and(dpt.asnumpy(ar3), dpt.asnumpy(ar4)) + assert (dpt.asnumpy(r2) == expected2).all() + + r3 = dpt.logical_and(ar4, ar4) + with np.errstate(invalid="ignore"): + expected3 = np.logical_and(dpt.asnumpy(ar4), dpt.asnumpy(ar4)) + assert (dpt.asnumpy(r3) == expected3).all() + + +def test_logical_and_complex_float(): + get_queue_or_skip() + + ar1 = dpt.asarray([1j, 1.0 + 9j, 2.0 + 0j, 2.0 + 1j], dtype="c8") + ar2 = dpt.full(ar1.shape, 2, dtype="f4") + + r = dpt.logical_and(ar1, ar2) + expected = np.logical_and(dpt.asnumpy(ar1), dpt.asnumpy(ar2)) + assert (dpt.asnumpy(r) == expected).all() + + r1 = dpt.logical_and(ar2, ar1) + expected1 = np.logical_and(dpt.asnumpy(ar2), dpt.asnumpy(ar1)) + assert (dpt.asnumpy(r1) == expected1).all() + with np.errstate(invalid="ignore"): + for tp in [ + dpt.nan, + dpt.nan * 1j, + dpt.inf, + dpt.inf * 1j, + -dpt.inf, + -dpt.inf * 1j, + ]: + ar3 = dpt.full(ar1.shape, tp) + r2 = dpt.logical_and(ar1, ar3) + expected2 = np.logical_and(dpt.asnumpy(ar1), dpt.asnumpy(ar3)) + assert (dpt.asnumpy(r2) == expected2).all() + + r3 = dpt.logical_and(ar3, ar1) + expected3 = np.logical_and(dpt.asnumpy(ar3), dpt.asnumpy(ar1)) + assert (dpt.asnumpy(r3) == expected3).all() + + +@pytest.mark.parametrize("op1_usm_type", _usm_types) +@pytest.mark.parametrize("op2_usm_type", _usm_types) +def test_logical_and_usm_type_matrix(op1_usm_type, op2_usm_type): + get_queue_or_skip() + + sz = 128 + ar1 = dpt.asarray( + np.random.randint(0, 2, sz), dtype="i4", usm_type=op1_usm_type + ) + ar2 = dpt.asarray( + np.random.randint(0, 2, sz), dtype=ar1.dtype, usm_type=op2_usm_type + ) + + r = dpt.logical_and(ar1, ar2) + assert isinstance(r, dpt.usm_ndarray) + expected_usm_type = dpctl.utils.get_coerced_usm_type( + (op1_usm_type, op2_usm_type) + ) + assert r.usm_type == expected_usm_type + + +def test_logical_and_order(): + get_queue_or_skip() + + ar1 = dpt.ones((20, 20), dtype="i4", order="C") + ar2 = dpt.ones((20, 20), dtype="i4", order="C") + r1 = dpt.logical_and(ar1, ar2, order="C") + assert r1.flags.c_contiguous + r2 = dpt.logical_and(ar1, ar2, order="F") + assert r2.flags.f_contiguous + r3 = dpt.logical_and(ar1, ar2, order="A") + assert r3.flags.c_contiguous + r4 = dpt.logical_and(ar1, ar2, order="K") + assert r4.flags.c_contiguous + + ar1 = dpt.ones((20, 20), dtype="i4", order="F") + ar2 = dpt.ones((20, 20), dtype="i4", order="F") + r1 = dpt.logical_and(ar1, ar2, order="C") + assert r1.flags.c_contiguous + r2 = dpt.logical_and(ar1, ar2, order="F") + assert r2.flags.f_contiguous + r3 = dpt.logical_and(ar1, ar2, order="A") + assert r3.flags.f_contiguous + r4 = dpt.logical_and(ar1, ar2, order="K") + assert r4.flags.f_contiguous + + ar1 = dpt.ones((40, 40), dtype="i4", order="C")[:20, ::-2] + ar2 = dpt.ones((40, 40), dtype="i4", order="C")[:20, ::-2] + r4 = dpt.logical_and(ar1, ar2, order="K") + assert r4.strides == (20, -1) + + ar1 = dpt.ones((40, 40), dtype="i4", order="C")[:20, ::-2].mT + ar2 = dpt.ones((40, 40), dtype="i4", order="C")[:20, ::-2].mT + r4 = dpt.logical_and(ar1, ar2, order="K") + assert r4.strides == (-1, 20) + + +def test_logical_and_broadcasting(): + get_queue_or_skip() + + m = dpt.asarray(np.random.randint(0, 2, (100, 5)), dtype="i4") + v = dpt.arange(1, 6, dtype="i4") + + r = dpt.logical_and(m, v) + + expected = np.logical_and(dpt.asnumpy(m), dpt.asnumpy(v)) + assert (dpt.asnumpy(r) == expected).all() + + r2 = dpt.logical_and(v, m) + expected2 = np.logical_and(dpt.asnumpy(v), dpt.asnumpy(m)) + assert (dpt.asnumpy(r2) == expected2).all() + + r3 = dpt.empty_like(r) + dpt.logical_and(m, v, out=r3) + assert (dpt.asnumpy(r3) == expected).all() + + r4 = dpt.empty_like(r) + dpt.logical_and(v, m, out=r4) + assert (dpt.asnumpy(r4) == expected).all() + + +@pytest.mark.parametrize("arr_dt", _all_dtypes) +@pytest.mark.parametrize("scalar_val", [0, 1]) +def test_logical_and_python_scalar(arr_dt, scalar_val): + q = get_queue_or_skip() + skip_if_dtype_not_supported(arr_dt, q) + + X = dpt.asarray( + np.random.randint(0, 2, (10, 10)), dtype=arr_dt, sycl_queue=q + ) + py_ones = ( + bool(scalar_val), + int(scalar_val), + float(scalar_val), + complex(scalar_val), + np.float32(scalar_val), + ctypes.c_int(scalar_val), + ) + for sc in py_ones: + R = dpt.logical_and(X, sc) + assert isinstance(R, dpt.usm_ndarray) + E = np.logical_and(dpt.asnumpy(X), sc) + assert (dpt.asnumpy(R) == E).all() + + R = dpt.logical_and(sc, X) + assert isinstance(R, dpt.usm_ndarray) + E = np.logical_and(sc, dpt.asnumpy(X)) + assert (dpt.asnumpy(R) == E).all() + + +class MockArray: + def __init__(self, arr): + self.data_ = arr + + @property + def __sycl_usm_array_interface__(self): + return self.data_.__sycl_usm_array_interface__ + + +def test_logical_and_mock_array(): + get_queue_or_skip() + a = dpt.arange(10) + b = dpt.ones(10) + c = MockArray(b) + r = dpt.logical_and(a, c) + assert isinstance(r, dpt.usm_ndarray) + + +def test_logical_and_canary_mock_array(): + get_queue_or_skip() + a = dpt.arange(10) + + class Canary: + def __init__(self): + pass + + @property + def __sycl_usm_array_interface__(self): + return None + + c = Canary() + with pytest.raises(ValueError): + dpt.logical_and(a, c) diff --git a/dpctl/tests/elementwise/test_logical_not.py b/dpctl/tests/elementwise/test_logical_not.py new file mode 100644 index 0000000000..aec9bf31b4 --- /dev/null +++ b/dpctl/tests/elementwise/test_logical_not.py @@ -0,0 +1,176 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2023 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless_equal required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +import numpy as np +import pytest + +import dpctl.tensor as dpt +from dpctl.tests.helper import get_queue_or_skip, skip_if_dtype_not_supported + +from .utils import _all_dtypes, _compare_dtypes, _usm_types + + +@pytest.mark.parametrize("op_dtype", _all_dtypes) +def test_logical_not_dtype_matrix(op_dtype): + q = get_queue_or_skip() + skip_if_dtype_not_supported(op_dtype, q) + + sz = 7 + ar1 = dpt.asarray(np.random.randint(0, 2, sz), dtype=op_dtype) + + r = dpt.logical_not(ar1) + assert isinstance(r, dpt.usm_ndarray) + + expected = np.logical_not(dpt.asnumpy(ar1)) + assert _compare_dtypes(r.dtype, expected.dtype, sycl_queue=q) + assert r.shape == ar1.shape + assert (dpt.asnumpy(r) == expected).all() + assert r.sycl_queue == ar1.sycl_queue + + r2 = dpt.empty_like(r, dtype=r.dtype) + dpt.logical_not(ar1, out=r2) + assert (dpt.asnumpy(r) == dpt.asnumpy(r2)).all() + + ar2 = dpt.zeros(sz, dtype=op_dtype) + r = dpt.logical_not(ar2[::-1]) + assert isinstance(r, dpt.usm_ndarray) + + expected = np.logical_not(np.zeros(ar2.shape, dtype=op_dtype)) + assert _compare_dtypes(r.dtype, expected.dtype, sycl_queue=q) + assert r.shape == ar2.shape + assert (dpt.asnumpy(r) == expected).all() + + ar3 = dpt.ones(sz, dtype=op_dtype) + r2 = dpt.logical_not(ar3[::2]) + assert isinstance(r, dpt.usm_ndarray) + + expected = np.logical_not(np.ones(ar3.shape, dtype=op_dtype)[::2]) + assert _compare_dtypes(r.dtype, expected.dtype, sycl_queue=q) + assert (dpt.asnumpy(r2) == expected).all() + + r3 = dpt.empty_like(r, dtype=r.dtype) + dpt.logical_not(ar2[::-1], out=r3) + assert (dpt.asnumpy(r) == dpt.asnumpy(r3)).all() + + +@pytest.mark.parametrize("op_dtype", ["c8", "c16"]) +def test_logical_not_complex_matrix(op_dtype): + q = get_queue_or_skip() + skip_if_dtype_not_supported(op_dtype, q) + + sz = 127 + ar1_np_real = np.random.randint(0, 2, sz) + ar1_np_imag = np.random.randint(0, 2, sz) + ar1 = dpt.asarray(ar1_np_real + 1j * ar1_np_imag, dtype=op_dtype) + + r = dpt.logical_not(ar1) + expected = np.logical_not(dpt.asnumpy(ar1)) + assert _compare_dtypes(r.dtype, expected.dtype, sycl_queue=q) + assert r.shape == expected.shape + assert (dpt.asnumpy(r) == expected).all() + + r1 = dpt.logical_not(ar1[::-2]) + expected1 = np.logical_not(dpt.asnumpy(ar1[::-2])) + assert _compare_dtypes(r.dtype, expected1.dtype, sycl_queue=q) + assert r1.shape == expected1.shape + assert (dpt.asnumpy(r1) == expected1).all() + + ar2 = dpt.asarray( + [ + 2.0 + 0j, + dpt.nan, + dpt.nan * 1j, + dpt.inf, + dpt.inf * 1j, + -dpt.inf, + -dpt.inf * 1j, + ], + dtype=op_dtype, + ) + r2 = dpt.logical_not(ar2) + with np.errstate(invalid="ignore"): + expected2 = np.logical_not(dpt.asnumpy(ar2)) + assert (dpt.asnumpy(r2) == expected2).all() + + +def test_logical_not_complex_float(): + get_queue_or_skip() + + ar1 = dpt.asarray([1j, 1.0 + 9j, 2.0 + 0j, 2.0 + 1j], dtype="c8") + + r = dpt.logical_not(ar1) + expected = np.logical_not(dpt.asnumpy(ar1)) + assert (dpt.asnumpy(r) == expected).all() + + with np.errstate(invalid="ignore"): + for tp in [ + dpt.nan, + dpt.nan * 1j, + dpt.inf, + dpt.inf * 1j, + -dpt.inf, + -dpt.inf * 1j, + ]: + ar2 = dpt.full(ar1.shape, tp) + r2 = dpt.logical_not(ar2) + expected2 = np.logical_not(dpt.asnumpy(ar2)) + assert (dpt.asnumpy(r2) == expected2).all() + + +@pytest.mark.parametrize("op_usm_type", _usm_types) +def test_logical_not_usm_type_matrix(op_usm_type): + get_queue_or_skip() + + sz = 128 + ar1 = dpt.asarray( + np.random.randint(0, 2, sz), dtype="i4", usm_type=op_usm_type + ) + + r = dpt.logical_not(ar1) + assert isinstance(r, dpt.usm_ndarray) + assert r.usm_type == op_usm_type + + +def test_logical_not_order(): + get_queue_or_skip() + + ar1 = dpt.ones((20, 20), dtype="i4", order="C") + r1 = dpt.logical_not(ar1, order="C") + assert r1.flags.c_contiguous + r2 = dpt.logical_not(ar1, order="F") + assert r2.flags.f_contiguous + r3 = dpt.logical_not(ar1, order="A") + assert r3.flags.c_contiguous + r4 = dpt.logical_not(ar1, order="K") + assert r4.flags.c_contiguous + + ar1 = dpt.zeros((20, 20), dtype="i4", order="F") + r1 = dpt.logical_not(ar1, order="C") + assert r1.flags.c_contiguous + r2 = dpt.logical_not(ar1, order="F") + assert r2.flags.f_contiguous + r3 = dpt.logical_not(ar1, order="A") + assert r3.flags.f_contiguous + r4 = dpt.logical_not(ar1, order="K") + assert r4.flags.f_contiguous + + ar1 = dpt.ones((40, 40), dtype="i4", order="C")[:20, ::-2] + r4 = dpt.logical_not(ar1, order="K") + assert r4.strides == (20, -1) + + ar1 = dpt.zeros((40, 40), dtype="i4", order="C")[:20, ::-2].mT + r4 = dpt.logical_not(ar1, order="K") + assert r4.strides == (-1, 20) diff --git a/dpctl/tests/elementwise/test_logical_or.py b/dpctl/tests/elementwise/test_logical_or.py new file mode 100644 index 0000000000..f99f6758f5 --- /dev/null +++ b/dpctl/tests/elementwise/test_logical_or.py @@ -0,0 +1,294 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2023 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless_equal required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +import ctypes + +import numpy as np +import pytest + +import dpctl +import dpctl.tensor as dpt +from dpctl.tests.helper import get_queue_or_skip, skip_if_dtype_not_supported + +from .utils import _all_dtypes, _compare_dtypes, _usm_types + + +@pytest.mark.parametrize("op1_dtype", _all_dtypes) +@pytest.mark.parametrize("op2_dtype", _all_dtypes) +def test_logical_or_dtype_matrix(op1_dtype, op2_dtype): + q = get_queue_or_skip() + skip_if_dtype_not_supported(op1_dtype, q) + skip_if_dtype_not_supported(op2_dtype, q) + + sz = 127 + ar1 = dpt.asarray(np.random.randint(0, 2, sz), dtype=op1_dtype) + ar2 = dpt.asarray(np.random.randint(0, 2, sz), dtype=op2_dtype) + + r = dpt.logical_or(ar1, ar2) + assert isinstance(r, dpt.usm_ndarray) + + expected = np.logical_or(dpt.asnumpy(ar1), dpt.asnumpy(ar2)) + assert _compare_dtypes(r.dtype, expected.dtype, sycl_queue=q) + assert r.shape == ar1.shape + assert (dpt.asnumpy(r) == expected).all() + assert r.sycl_queue == ar1.sycl_queue + + r2 = dpt.empty_like(r, dtype=r.dtype) + dpt.logical_or(ar1, ar2, out=r2) + assert (dpt.asnumpy(r) == dpt.asnumpy(r2)).all() + + ar3 = dpt.zeros(sz, dtype=op1_dtype) + ar4 = dpt.ones(2 * sz, dtype=op2_dtype) + + r = dpt.logical_or(ar3[::-1], ar4[::2]) + assert isinstance(r, dpt.usm_ndarray) + expected = np.logical_or( + np.zeros(1, dtype=op1_dtype), np.ones(1, dtype=op2_dtype) + ) + assert _compare_dtypes(r.dtype, expected.dtype, sycl_queue=q) + assert r.shape == ar3.shape + assert (dpt.asnumpy(r) == expected).all() + + r2 = dpt.empty_like(r, dtype=r.dtype) + dpt.logical_or(ar3[::-1], ar4[::2], out=r2) + assert (dpt.asnumpy(r) == dpt.asnumpy(r2)).all() + + +@pytest.mark.parametrize("op_dtype", ["c8", "c16"]) +def test_logical_or_complex_matrix(op_dtype): + q = get_queue_or_skip() + skip_if_dtype_not_supported(op_dtype, q) + + sz = 127 + ar1_np_real = np.random.randint(0, 2, sz) + ar1_np_imag = np.random.randint(0, 2, sz) + ar1 = dpt.asarray(ar1_np_real + 1j * ar1_np_imag, dtype=op_dtype) + + ar2_np_real = np.random.randint(0, 2, sz) + ar2_np_imag = np.random.randint(0, 2, sz) + ar2 = dpt.asarray(ar2_np_real + 1j * ar2_np_imag, dtype=op_dtype) + + r = dpt.logical_or(ar1, ar2) + expected = np.logical_or(dpt.asnumpy(ar1), dpt.asnumpy(ar2)) + assert _compare_dtypes(r.dtype, expected.dtype, sycl_queue=q) + assert r.shape == expected.shape + assert (dpt.asnumpy(r) == expected).all() + + r1 = dpt.logical_or(ar1[::-2], ar2[::2]) + expected1 = np.logical_or(dpt.asnumpy(ar1[::-2]), dpt.asnumpy(ar2[::2])) + assert _compare_dtypes(r.dtype, expected1.dtype, sycl_queue=q) + assert r1.shape == expected1.shape + assert (dpt.asnumpy(r1) == expected1).all() + + ar3 = dpt.asarray( + [ + 2.0 + 0j, + dpt.nan, + dpt.nan * 1j, + dpt.inf, + dpt.inf * 1j, + -dpt.inf, + -dpt.inf * 1j, + ], + dtype=op_dtype, + ) + ar4 = dpt.full(ar3.shape, fill_value=1.0 + 2j, dtype=op_dtype) + r2 = dpt.logical_or(ar3, ar4) + with np.errstate(invalid="ignore"): + expected2 = np.logical_or(dpt.asnumpy(ar3), dpt.asnumpy(ar4)) + assert (dpt.asnumpy(r2) == expected2).all() + + r3 = dpt.logical_or(ar4, ar4) + with np.errstate(invalid="ignore"): + expected3 = np.logical_or(dpt.asnumpy(ar4), dpt.asnumpy(ar4)) + assert (dpt.asnumpy(r3) == expected3).all() + + +def test_logical_or_complex_float(): + get_queue_or_skip() + + ar1 = dpt.asarray([1j, 1.0 + 9j, 2.0 + 0j, 2.0 + 1j], dtype="c8") + ar2 = dpt.full(ar1.shape, 2, dtype="f4") + + r = dpt.logical_or(ar1, ar2) + expected = np.logical_or(dpt.asnumpy(ar1), dpt.asnumpy(ar2)) + assert (dpt.asnumpy(r) == expected).all() + + r1 = dpt.logical_or(ar2, ar1) + expected1 = np.logical_or(dpt.asnumpy(ar2), dpt.asnumpy(ar1)) + assert (dpt.asnumpy(r1) == expected1).all() + with np.errstate(invalid="ignore"): + for tp in [ + dpt.nan, + dpt.nan * 1j, + dpt.inf, + dpt.inf * 1j, + -dpt.inf, + -dpt.inf * 1j, + ]: + ar3 = dpt.full(ar1.shape, tp) + r2 = dpt.logical_or(ar1, ar3) + expected2 = np.logical_or(dpt.asnumpy(ar1), dpt.asnumpy(ar3)) + assert (dpt.asnumpy(r2) == expected2).all() + + r3 = dpt.logical_or(ar3, ar1) + expected3 = np.logical_or(dpt.asnumpy(ar3), dpt.asnumpy(ar1)) + assert (dpt.asnumpy(r3) == expected3).all() + + +@pytest.mark.parametrize("op1_usm_type", _usm_types) +@pytest.mark.parametrize("op2_usm_type", _usm_types) +def test_logical_or_usm_type_matrix(op1_usm_type, op2_usm_type): + get_queue_or_skip() + + sz = 128 + ar1 = dpt.asarray( + np.random.randint(0, 2, sz), dtype="i4", usm_type=op1_usm_type + ) + ar2 = dpt.asarray( + np.random.randint(0, 2, sz), dtype=ar1.dtype, usm_type=op2_usm_type + ) + + r = dpt.logical_or(ar1, ar2) + assert isinstance(r, dpt.usm_ndarray) + expected_usm_type = dpctl.utils.get_coerced_usm_type( + (op1_usm_type, op2_usm_type) + ) + assert r.usm_type == expected_usm_type + + +def test_logical_or_order(): + get_queue_or_skip() + + ar1 = dpt.ones((20, 20), dtype="i4", order="C") + ar2 = dpt.ones((20, 20), dtype="i4", order="C") + r1 = dpt.logical_or(ar1, ar2, order="C") + assert r1.flags.c_contiguous + r2 = dpt.logical_or(ar1, ar2, order="F") + assert r2.flags.f_contiguous + r3 = dpt.logical_or(ar1, ar2, order="A") + assert r3.flags.c_contiguous + r4 = dpt.logical_or(ar1, ar2, order="K") + assert r4.flags.c_contiguous + + ar1 = dpt.ones((20, 20), dtype="i4", order="F") + ar2 = dpt.ones((20, 20), dtype="i4", order="F") + r1 = dpt.logical_or(ar1, ar2, order="C") + assert r1.flags.c_contiguous + r2 = dpt.logical_or(ar1, ar2, order="F") + assert r2.flags.f_contiguous + r3 = dpt.logical_or(ar1, ar2, order="A") + assert r3.flags.f_contiguous + r4 = dpt.logical_or(ar1, ar2, order="K") + assert r4.flags.f_contiguous + + ar1 = dpt.ones((40, 40), dtype="i4", order="C")[:20, ::-2] + ar2 = dpt.ones((40, 40), dtype="i4", order="C")[:20, ::-2] + r4 = dpt.logical_or(ar1, ar2, order="K") + assert r4.strides == (20, -1) + + ar1 = dpt.ones((40, 40), dtype="i4", order="C")[:20, ::-2].mT + ar2 = dpt.ones((40, 40), dtype="i4", order="C")[:20, ::-2].mT + r4 = dpt.logical_or(ar1, ar2, order="K") + assert r4.strides == (-1, 20) + + +def test_logical_or_broadcasting(): + get_queue_or_skip() + + m = dpt.asarray(np.random.randint(0, 2, (100, 5)), dtype="i4") + v = dpt.arange(1, 6, dtype="i4") + + r = dpt.logical_or(m, v) + + expected = np.logical_or(dpt.asnumpy(m), dpt.asnumpy(v)) + assert (dpt.asnumpy(r) == expected).all() + + r2 = dpt.logical_or(v, m) + expected2 = np.logical_or(dpt.asnumpy(v), dpt.asnumpy(m)) + assert (dpt.asnumpy(r2) == expected2).all() + + r3 = dpt.empty_like(r) + dpt.logical_or(m, v, out=r3) + assert (dpt.asnumpy(r3) == expected).all() + + r4 = dpt.empty_like(r) + dpt.logical_or(v, m, out=r4) + assert (dpt.asnumpy(r4) == expected).all() + + +@pytest.mark.parametrize("arr_dt", _all_dtypes) +@pytest.mark.parametrize("scalar_val", [0, 1]) +def test_logical_or_python_scalar(arr_dt, scalar_val): + q = get_queue_or_skip() + skip_if_dtype_not_supported(arr_dt, q) + + X = dpt.asarray( + np.random.randint(0, 2, (10, 10)), dtype=arr_dt, sycl_queue=q + ) + py_ones = ( + bool(scalar_val), + int(scalar_val), + float(scalar_val), + complex(scalar_val), + np.float32(scalar_val), + ctypes.c_int(scalar_val), + ) + for sc in py_ones: + R = dpt.logical_or(X, sc) + assert isinstance(R, dpt.usm_ndarray) + E = np.logical_or(dpt.asnumpy(X), sc) + assert (dpt.asnumpy(R) == E).all() + + R = dpt.logical_or(sc, X) + assert isinstance(R, dpt.usm_ndarray) + E = np.logical_or(sc, dpt.asnumpy(X)) + assert (dpt.asnumpy(R) == E).all() + + +class MockArray: + def __init__(self, arr): + self.data_ = arr + + @property + def __sycl_usm_array_interface__(self): + return self.data_.__sycl_usm_array_interface__ + + +def test_logical_or_mock_array(): + get_queue_or_skip() + a = dpt.arange(10) + b = dpt.ones(10) + c = MockArray(b) + r = dpt.logical_or(a, c) + assert isinstance(r, dpt.usm_ndarray) + + +def test_logical_or_canary_mock_array(): + get_queue_or_skip() + a = dpt.arange(10) + + class Canary: + def __init__(self): + pass + + @property + def __sycl_usm_array_interface__(self): + return None + + c = Canary() + with pytest.raises(ValueError): + dpt.logical_or(a, c) diff --git a/dpctl/tests/elementwise/test_logical_xor.py b/dpctl/tests/elementwise/test_logical_xor.py new file mode 100644 index 0000000000..9c34e8bbb5 --- /dev/null +++ b/dpctl/tests/elementwise/test_logical_xor.py @@ -0,0 +1,294 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2023 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless_equal required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +import ctypes + +import numpy as np +import pytest + +import dpctl +import dpctl.tensor as dpt +from dpctl.tests.helper import get_queue_or_skip, skip_if_dtype_not_supported + +from .utils import _all_dtypes, _compare_dtypes, _usm_types + + +@pytest.mark.parametrize("op1_dtype", _all_dtypes) +@pytest.mark.parametrize("op2_dtype", _all_dtypes) +def test_logical_xor_dtype_matrix(op1_dtype, op2_dtype): + q = get_queue_or_skip() + skip_if_dtype_not_supported(op1_dtype, q) + skip_if_dtype_not_supported(op2_dtype, q) + + sz = 127 + ar1 = dpt.asarray(np.random.randint(0, 2, sz), dtype=op1_dtype) + ar2 = dpt.asarray(np.random.randint(0, 2, sz), dtype=op2_dtype) + + r = dpt.logical_xor(ar1, ar2) + assert isinstance(r, dpt.usm_ndarray) + + expected = np.logical_xor(dpt.asnumpy(ar1), dpt.asnumpy(ar2)) + assert _compare_dtypes(r.dtype, expected.dtype, sycl_queue=q) + assert r.shape == ar1.shape + assert (dpt.asnumpy(r) == expected).all() + assert r.sycl_queue == ar1.sycl_queue + + r2 = dpt.empty_like(r, dtype=r.dtype) + dpt.logical_xor(ar1, ar2, out=r2) + assert (dpt.asnumpy(r) == dpt.asnumpy(r2)).all() + + ar3 = dpt.zeros(sz, dtype=op1_dtype) + ar4 = dpt.ones(2 * sz, dtype=op2_dtype) + + r = dpt.logical_xor(ar3[::-1], ar4[::2]) + assert isinstance(r, dpt.usm_ndarray) + expected = np.logical_xor( + np.zeros(1, dtype=op1_dtype), np.ones(1, dtype=op2_dtype) + ) + assert _compare_dtypes(r.dtype, expected.dtype, sycl_queue=q) + assert r.shape == ar3.shape + assert (dpt.asnumpy(r) == expected).all() + + r2 = dpt.empty_like(r, dtype=r.dtype) + dpt.logical_xor(ar3[::-1], ar4[::2], out=r2) + assert (dpt.asnumpy(r) == dpt.asnumpy(r2)).all() + + +@pytest.mark.parametrize("op_dtype", ["c8", "c16"]) +def test_logical_xor_complex_matrix(op_dtype): + q = get_queue_or_skip() + skip_if_dtype_not_supported(op_dtype, q) + + sz = 127 + ar1_np_real = np.random.randint(0, 2, sz) + ar1_np_imag = np.random.randint(0, 2, sz) + ar1 = dpt.asarray(ar1_np_real + 1j * ar1_np_imag, dtype=op_dtype) + + ar2_np_real = np.random.randint(0, 2, sz) + ar2_np_imag = np.random.randint(0, 2, sz) + ar2 = dpt.asarray(ar2_np_real + 1j * ar2_np_imag, dtype=op_dtype) + + r = dpt.logical_xor(ar1, ar2) + expected = np.logical_xor(dpt.asnumpy(ar1), dpt.asnumpy(ar2)) + assert _compare_dtypes(r.dtype, expected.dtype, sycl_queue=q) + assert r.shape == expected.shape + assert (dpt.asnumpy(r) == expected).all() + + r1 = dpt.logical_xor(ar1[::-2], ar2[::2]) + expected1 = np.logical_xor(dpt.asnumpy(ar1[::-2]), dpt.asnumpy(ar2[::2])) + assert _compare_dtypes(r.dtype, expected1.dtype, sycl_queue=q) + assert r1.shape == expected1.shape + assert (dpt.asnumpy(r1) == expected1).all() + + ar3 = dpt.asarray( + [ + 2.0 + 0j, + dpt.nan, + dpt.nan * 1j, + dpt.inf, + dpt.inf * 1j, + -dpt.inf, + -dpt.inf * 1j, + ], + dtype=op_dtype, + ) + ar4 = dpt.full(ar3.shape, fill_value=1.0 + 2j, dtype=op_dtype) + r2 = dpt.logical_xor(ar3, ar4) + with np.errstate(invalid="ignore"): + expected2 = np.logical_xor(dpt.asnumpy(ar3), dpt.asnumpy(ar4)) + assert (dpt.asnumpy(r2) == expected2).all() + + r3 = dpt.logical_xor(ar4, ar4) + with np.errstate(invalid="ignore"): + expected3 = np.logical_xor(dpt.asnumpy(ar4), dpt.asnumpy(ar4)) + assert (dpt.asnumpy(r3) == expected3).all() + + +def test_logical_xor_complex_float(): + get_queue_or_skip() + + ar1 = dpt.asarray([1j, 1.0 + 9j, 2.0 + 0j, 2.0 + 1j], dtype="c8") + ar2 = dpt.full(ar1.shape, 2, dtype="f4") + + r = dpt.logical_xor(ar1, ar2) + expected = np.logical_xor(dpt.asnumpy(ar1), dpt.asnumpy(ar2)) + assert (dpt.asnumpy(r) == expected).all() + + r1 = dpt.logical_xor(ar2, ar1) + expected1 = np.logical_xor(dpt.asnumpy(ar2), dpt.asnumpy(ar1)) + assert (dpt.asnumpy(r1) == expected1).all() + with np.errstate(invalid="ignore"): + for tp in [ + dpt.nan, + dpt.nan * 1j, + dpt.inf, + dpt.inf * 1j, + -dpt.inf, + -dpt.inf * 1j, + ]: + ar3 = dpt.full(ar1.shape, tp) + r2 = dpt.logical_xor(ar1, ar3) + expected2 = np.logical_xor(dpt.asnumpy(ar1), dpt.asnumpy(ar3)) + assert (dpt.asnumpy(r2) == expected2).all() + + r3 = dpt.logical_xor(ar3, ar1) + expected3 = np.logical_xor(dpt.asnumpy(ar3), dpt.asnumpy(ar1)) + assert (dpt.asnumpy(r3) == expected3).all() + + +@pytest.mark.parametrize("op1_usm_type", _usm_types) +@pytest.mark.parametrize("op2_usm_type", _usm_types) +def test_logical_xor_usm_type_matrix(op1_usm_type, op2_usm_type): + get_queue_or_skip() + + sz = 128 + ar1 = dpt.asarray( + np.random.randint(0, 2, sz), dtype="i4", usm_type=op1_usm_type + ) + ar2 = dpt.asarray( + np.random.randint(0, 2, sz), dtype=ar1.dtype, usm_type=op2_usm_type + ) + + r = dpt.logical_xor(ar1, ar2) + assert isinstance(r, dpt.usm_ndarray) + expected_usm_type = dpctl.utils.get_coerced_usm_type( + (op1_usm_type, op2_usm_type) + ) + assert r.usm_type == expected_usm_type + + +def test_logical_xor_order(): + get_queue_or_skip() + + ar1 = dpt.ones((20, 20), dtype="i4", order="C") + ar2 = dpt.ones((20, 20), dtype="i4", order="C") + r1 = dpt.logical_xor(ar1, ar2, order="C") + assert r1.flags.c_contiguous + r2 = dpt.logical_xor(ar1, ar2, order="F") + assert r2.flags.f_contiguous + r3 = dpt.logical_xor(ar1, ar2, order="A") + assert r3.flags.c_contiguous + r4 = dpt.logical_xor(ar1, ar2, order="K") + assert r4.flags.c_contiguous + + ar1 = dpt.ones((20, 20), dtype="i4", order="F") + ar2 = dpt.ones((20, 20), dtype="i4", order="F") + r1 = dpt.logical_xor(ar1, ar2, order="C") + assert r1.flags.c_contiguous + r2 = dpt.logical_xor(ar1, ar2, order="F") + assert r2.flags.f_contiguous + r3 = dpt.logical_xor(ar1, ar2, order="A") + assert r3.flags.f_contiguous + r4 = dpt.logical_xor(ar1, ar2, order="K") + assert r4.flags.f_contiguous + + ar1 = dpt.ones((40, 40), dtype="i4", order="C")[:20, ::-2] + ar2 = dpt.ones((40, 40), dtype="i4", order="C")[:20, ::-2] + r4 = dpt.logical_xor(ar1, ar2, order="K") + assert r4.strides == (20, -1) + + ar1 = dpt.ones((40, 40), dtype="i4", order="C")[:20, ::-2].mT + ar2 = dpt.ones((40, 40), dtype="i4", order="C")[:20, ::-2].mT + r4 = dpt.logical_xor(ar1, ar2, order="K") + assert r4.strides == (-1, 20) + + +def test_logical_xor_broadcasting(): + get_queue_or_skip() + + m = dpt.asarray(np.random.randint(0, 2, (100, 5)), dtype="i4") + v = dpt.arange(1, 6, dtype="i4") + + r = dpt.logical_xor(m, v) + + expected = np.logical_xor(dpt.asnumpy(m), dpt.asnumpy(v)) + assert (dpt.asnumpy(r) == expected).all() + + r2 = dpt.logical_xor(v, m) + expected2 = np.logical_xor(dpt.asnumpy(v), dpt.asnumpy(m)) + assert (dpt.asnumpy(r2) == expected2).all() + + r3 = dpt.empty_like(r) + dpt.logical_xor(m, v, out=r3) + assert (dpt.asnumpy(r3) == expected).all() + + r4 = dpt.empty_like(r) + dpt.logical_xor(v, m, out=r4) + assert (dpt.asnumpy(r4) == expected).all() + + +@pytest.mark.parametrize("arr_dt", _all_dtypes) +@pytest.mark.parametrize("scalar_val", [0, 1]) +def test_logical_xor_python_scalar(arr_dt, scalar_val): + q = get_queue_or_skip() + skip_if_dtype_not_supported(arr_dt, q) + + X = dpt.asarray( + np.random.randint(0, 2, (10, 10)), dtype=arr_dt, sycl_queue=q + ) + py_ones = ( + bool(scalar_val), + int(scalar_val), + float(scalar_val), + complex(scalar_val), + np.float32(scalar_val), + ctypes.c_int(scalar_val), + ) + for sc in py_ones: + R = dpt.logical_xor(X, sc) + assert isinstance(R, dpt.usm_ndarray) + E = np.logical_xor(dpt.asnumpy(X), sc) + assert (dpt.asnumpy(R) == E).all() + + R = dpt.logical_xor(sc, X) + assert isinstance(R, dpt.usm_ndarray) + E = np.logical_xor(sc, dpt.asnumpy(X)) + assert (dpt.asnumpy(R) == E).all() + + +class MockArray: + def __init__(self, arr): + self.data_ = arr + + @property + def __sycl_usm_array_interface__(self): + return self.data_.__sycl_usm_array_interface__ + + +def test_logical_xor_mock_array(): + get_queue_or_skip() + a = dpt.arange(10) + b = dpt.ones(10) + c = MockArray(b) + r = dpt.logical_xor(a, c) + assert isinstance(r, dpt.usm_ndarray) + + +def test_logical_xor_canary_mock_array(): + get_queue_or_skip() + a = dpt.arange(10) + + class Canary: + def __init__(self): + pass + + @property + def __sycl_usm_array_interface__(self): + return None + + c = Canary() + with pytest.raises(ValueError): + dpt.logical_xor(a, c)