From 0c7d128a41dc072c5bc33d47bd1fef967a4b20d0 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Thu, 16 Sep 2021 15:52:23 -0700 Subject: [PATCH 01/18] add bitwise_left/right_shift --- python/mxnet/ndarray/numpy/_op.py | 81 ++++++- python/mxnet/numpy/multiarray.py | 75 ++++++- .../numpy/np_elemwise_broadcast_op.cc | 18 ++ .../np_elemwise_broadcast_op_extended_thi.cc | 197 ++++++++++++++++++ .../np_elemwise_broadcast_op_extended_thi.cu | 70 +++++++ 5 files changed, 439 insertions(+), 2 deletions(-) create mode 100644 src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cc create mode 100644 src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cu diff --git a/python/mxnet/ndarray/numpy/_op.py b/python/mxnet/ndarray/numpy/_op.py index a7465865b707..b44e70a84cf3 100644 --- a/python/mxnet/ndarray/numpy/_op.py +++ b/python/mxnet/ndarray/numpy/_op.py @@ -50,7 +50,8 @@ 'true_divide', 'nonzero', 'quantile', 'percentile', 'shares_memory', 'may_share_memory', 'interp', 'diff', 'ediff1d', 'resize', 'polyval', 'nan_to_num', 'isnan', 'isinf', 'isposinf', 'isneginf', 'isfinite', 'atleast_1d', 'atleast_2d', 'atleast_3d', 'fill_diagonal', 'squeeze', - 'where', 'bincount', 'rollaxis', 'diagflat', 'repeat', 'prod', 'pad', 'cumsum', 'sum', 'diag', 'diagonal'] + 'where', 'bincount', 'rollaxis', 'diagflat', 'repeat', 'prod', 'pad', 'cumsum', 'sum', 'diag', 'diagonal', + 'bitwise_left_shift', 'bitwise_right_shift'] @set_module('mxnet.ndarray.numpy') @@ -9881,3 +9882,81 @@ def sum(a, axis=None, dtype=None, out=None, keepdims=None, initial=None, where=N raise ValueError("only where=None or where=True cases are supported for now") return _api_internal.sum(a, axis, dtype, keepdims, initial, out) # pylint:enable=redefined-outer-name, too-many-arguments + + +@set_module('mxnet.ndarray.numpy') +def bitwise_left_shift(x1, x2, out=None): + r""" + Shift the bits of and integer to the left. Bits are shifted to the left by + appending x2 0s at the right of x1. Since the internal representation of numbers + is in binary format, this operation is equivalent to ``x1 * 2**x2`` + + Parameters + ---------- + x1 : ndarray or scalar + Input values. + x2 : ndarray or scalar + Number of zeros to append to x1. Has to be non-negative. If x1.shape != x2.shape, + they must be broadcastable to a common shape (which becomes the shape of the output). + out : ndarray, optional + A location into which the result is stored. If provided, it must have a shape that the + inputs broadcast to. If not provided or None, a freshly-allocated array is returned. + + Returns + ------- + out : ndarray + Result. + + Examples + -------- + >>> np.binary_repr(5) + '101' + >>> np.left_shift(5, 2) + 20 + >>> np.binary_repr(20) + '10100' + >>> np.left_shift(5, np.array([1,2,3])) + array([10, 20, 40]) + """ + if isinstance(x1, numeric_types) and isinstance(x2, numeric_types): + return _np.left_shift(x1, x2, out=out) + return _api_internal.bitwise_left_shift(x1, x2, out) + + +@set_module('mxnet.ndarray.numpy') +def bitwise_right_shift(x1, x2, out=None): + r""" + Shift the bits of and integer to the right. Bits are shifted to the right by + x2. Because the internal representation of numbers is in binary format, + this operation is equivalent to ``x1 / 2**x2`` + + Parameters + ---------- + x1 : ndarray or scalar + Input values. + x1 : ndarray or scalar + Number of bits to remove at the right of x1. If x1.shape != x2.shape, + they must be broadcastable to a common shape (which becomes the shape of the output). + out : ndarray, optional + A location into which the result is stored. If provided, it must have a shape that the + inputs broadcast to. If not provided or None, a freshly-allocated array is returned. + + Returns + ------- + out : ndarray + Result. + + Examples + -------- + >>> np.binary_repr(10) + '1010' + >>> np.right_shift(10, 1) + 5 + >>> np.binary_repr(5) + '101' + >>> np.right_shift(10, np.array([1,2,3])) + array([5, 2, 1]) + """ + if isinstance(x1, numeric_types) and isinstance(x2, numeric_types): + return _np.right_shift(x1, x2, out=out) + return _api_internal.bitwise_right_shift(x1, x2, out) diff --git a/python/mxnet/numpy/multiarray.py b/python/mxnet/numpy/multiarray.py index 5cca1fa9225a..1754122985f7 100644 --- a/python/mxnet/numpy/multiarray.py +++ b/python/mxnet/numpy/multiarray.py @@ -80,7 +80,8 @@ 'quantile', 'percentile', 'shares_memory', 'may_share_memory', 'diff', 'ediff1d', 'resize', 'matmul', 'nan_to_num', 'isnan', 'isinf', 'isposinf', 'isneginf', 'isfinite', 'polyval', 'where', 'bincount', 'atleast_1d', 'atleast_2d', 'atleast_3d', 'fill_diagonal', 'squeeze', - 'diagflat', 'repeat', 'prod', 'pad', 'cumsum', 'sum', 'rollaxis', 'diag', 'diagonal'] + 'diagflat', 'repeat', 'prod', 'pad', 'cumsum', 'sum', 'rollaxis', 'diag', 'diagonal', + 'bitwise_left_shift', 'bitwise_right_shift'] __all__ += fallback.__all__ @@ -12234,3 +12235,75 @@ def sum(a, axis=None, dtype=None, out=None, keepdims=None, initial=None, where=N """ return _mx_nd_np.sum(a, axis=axis, dtype=dtype, out=out, keepdims=keepdims, initial=initial, where=where) # pylint: enable=redefined-outer-name, too-many-arguments + + +@set_module('mxnet.numpy') +def bitwise_left_shift(x1, x2, out=None): + r""" + Shift the bits of and integer to the left. Bits are shifted to the left by + appending x2 0s at the right of x1. Since the internal representation of numbers + is in binary format, this operation is equivalent to ``x1 * 2**x2`` + + Parameters + ---------- + x1 : ndarray or scalar + Input values. + x2 : ndarray or scalar + Number of zeros to append to x1. Has to be non-negative. If x1.shape != x2.shape, + they must be broadcastable to a common shape (which becomes the shape of the output). + out : ndarray, optional + A location into which the result is stored. If provided, it must have a shape that the + inputs broadcast to. If not provided or None, a freshly-allocated array is returned. + + Returns + ------- + out : ndarray + Result. + + Examples + -------- + >>> np.binary_repr(5) + '101' + >>> np.left_shift(5, 2) + 20 + >>> np.binary_repr(20) + '10100' + """ + return _mx_nd_np.bitwise_left_shift(x1, x2, out) + + +@set_module('mxnet.numpy') +def bitwise_right_shift(x1, x2, out=None): + r""" + Shift the bits of and integer to the right. Bits are shifted to the right by + x2. Because the internal representation of numbers is in binary format, + this operation is equivalent to ``x1 / 2**x2`` + + Parameters + ---------- + x1 : ndarray or scalar + Input values. + x1 : ndarray or scalar + Number of bits to remove at the right of x1. If x1.shape != x2.shape, + they must be broadcastable to a common shape (which becomes the shape of the output). + out : ndarray, optional + A location into which the result is stored. If provided, it must have a shape that the + inputs broadcast to. If not provided or None, a freshly-allocated array is returned. + + Returns + ------- + out : ndarray + Result. + + Examples + -------- + >>> np.binary_repr(10) + '1010' + >>> np.right_shift(10, 1) + 5 + >>> np.binary_repr(5) + '101' + >>> np.right_shift(10, np.array([1,2,3])) + array([5, 2, 1]) + """ + return _mx_nd_np.bitwise_right_shift(x1, x2, out) diff --git a/src/api/operator/numpy/np_elemwise_broadcast_op.cc b/src/api/operator/numpy/np_elemwise_broadcast_op.cc index 184a4e241eff..560223d714df 100644 --- a/src/api/operator/numpy/np_elemwise_broadcast_op.cc +++ b/src/api/operator/numpy/np_elemwise_broadcast_op.cc @@ -172,4 +172,22 @@ MXNET_REGISTER_API("_npi.ldexp").set_body([](runtime::MXNetArgs args, runtime::M UFuncHelper(args, ret, op, op_scalar, op_rscalar); }); +MXNET_REGISTER_API("_npi.bitwise_left_shift") + .set_body([](runtime::MXNetArgs args, runtime::MXNetRetValue* ret) { + using namespace runtime; + const nnvm::Op* op = Op::Get("_npi_bitwise_left_shift"); + const nnvm::Op* op_scalar = Op::Get("_npi_bitwise_left_shift_scalar"); + const nnvm::Op* op_rscalar = Op::Get("_npi_rbitwise_left_shift_scalar"); + UFuncHelper(args, ret, op, op_scalar, op_rscalar); + }); + +MXNET_REGISTER_API("_npi.bitwise_right_shift") + .set_body([](runtime::MXNetArgs args, runtime::MXNetRetValue* ret) { + using namespace runtime; + const nnvm::Op* op = Op::Get("_npi_bitwise_right_shift"); + const nnvm::Op* op_scalar = Op::Get("_npi_bitwise_right_shift_scalar"); + const nnvm::Op* op_rscalar = Op::Get("_npi_rbitwise_right_shift_scalar"); + UFuncHelper(args, ret, op, op_scalar, op_rscalar); + }); + } // namespace mxnet diff --git a/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cc b/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cc new file mode 100644 index 000000000000..23b9b878b7a3 --- /dev/null +++ b/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cc @@ -0,0 +1,197 @@ +/* + * Licensed to the Apache Software Foundation (ASF) under one + * or more contributor license agreements. See the NOTICE file + * distributed with this work for additional information + * regarding copyright ownership. The ASF licenses this file + * to you 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. + */ + +/*! + * Copyright (c) 2019 by Contributors + * \file np_elemwise_broadcast_op_extended_thi.cc + * \brief CPU Implementation of extended functions for elementwise numpy binary broadcast operator. (Third extended file) + */ + +#include "../../common/utils.h" +#include "./np_elemwise_broadcast_op.h" + +namespace mxnet { +namespace op { + +#define MXNET_OPERATOR_REGISTER_NP_BINARY_SCALAR(name) \ + NNVM_REGISTER_OP(name) \ + .set_num_inputs(1) \ + .set_num_outputs(1) \ + .set_attr_parser(ParamParser) \ + .set_attr("FInferShape", ElemwiseShape<1, 1>) \ + .set_attr("FInferType", NumpyBinaryScalarType) \ + .set_attr( \ + "FResourceRequest", \ + [](const NodeAttrs& attrs) { \ + return std::vector{ResourceRequest::kTempSpace}; \ + }) \ + .add_argument("data", "NDArray-or-Symbol", "source input") \ + .add_arguments(NumpyBinaryScalarParam::__FIELDS__()) + +NNVM_REGISTER_OP(_npi_bitwise_left_shift) + .set_num_inputs(2) + .set_num_outputs(1) + .set_attr("FListInputNames", + [](const NodeAttrs& attrs) { + return std::vector{"lhs", "rhs"}; + }) + .set_attr("FInferShape", BinaryBroadcastShape) + .set_attr("FInferType", ElemwiseIntType<2, 1>) + .set_attr("FInplaceOption", + [](const NodeAttrs& attrs) { + return std::vector >{{0, 0}, {1, 0}}; + }) + .set_attr("FCompute", BinaryBroadcastIntCompute) + .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_bitwise_left_shift"}) + .add_argument("lhs", "NDArray-or-Symbol", "First input to the function") + .add_argument("rhs", "NDArray-or-Symbol", "Second input to the function"); + +NNVM_REGISTER_OP(_npi_bitwise_left_shift_scalar) + .set_num_inputs(1) + .set_num_outputs(1) + .set_attr_parser(ParamParser) + .set_attr("FInferShape", ElemwiseShape<1, 1>) + .set_attr("FInferType", ElemwiseIntType<1, 1>) + .set_attr("FInplaceOption", + [](const NodeAttrs& attrs) { + return std::vector >{{0, 0}}; + }) + .set_attr("FCompute", BinaryScalarOp::ComputeInt) + .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_bitwise_left_shift_scalar"}) + .add_argument("data", "NDArray-or-Symbol", "source input") + .add_arguments(NumpyBinaryScalarParam::__FIELDS__()); + +NNVM_REGISTER_OP(_npi_rbitwise_left_shift_scalar) + .set_num_inputs(1) + .set_num_outputs(1) + .set_attr_parser(ParamParser) + .set_attr("FInferShape", ElemwiseShape<1, 1>) + .set_attr("FInferType", ElemwiseIntType<1, 1>) + .set_attr("FInplaceOption", + [](const NodeAttrs& attrs) { + return std::vector >{{0, 0}}; + }) + .set_attr("FCompute", BinaryScalarOp::ComputeInt) + .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_rbitwise_left_shift_scalar"}) + .add_argument("data", "NDArray-or-Symbol", "source input") + .add_arguments(NumpyBinaryScalarParam::__FIELDS__()); + +NNVM_REGISTER_OP(_backward_npi_bitwise_left_shift) + .set_num_inputs(3) + .set_num_outputs(2) + .set_attr("TIsBackward", true) + .set_attr("FInplaceOption", + [](const NodeAttrs& attrs){ + return std::vector >{{0, 1}}; + }) + .set_attr("FResourceRequest", + [](const NodeAttrs& attrs) { + return std::vector{ResourceRequest::kTempSpace}; + }) + .set_attr("FCompute", + BinaryBroadcastBackwardUseIn); + +NNVM_REGISTER_OP(_backward_npi_bitwise_left_shift_scalar) + .add_arguments(NumpyBinaryScalarParam::__FIELDS__()) + .set_attr_parser(ParamParser) + .set_attr("FCompute", BinaryScalarOp::Backward); + +MXNET_OPERATOR_REGISTER_BINARY(_backward_npi_rbitwise_left_shift_scalar) + .add_arguments(NumpyBinaryScalarParam::__FIELDS__()) + .set_attr_parser(ParamParser) + .set_attr("FCompute", BinaryScalarOp::Backward); + + +NNVM_REGISTER_OP(_npi_bitwise_right_shift) + .set_num_inputs(2) + .set_num_outputs(1) + .set_attr("FListInputNames", + [](const NodeAttrs& attrs) { + return std::vector{"lhs", "rhs"}; + }) + .set_attr("FInferShape", BinaryBroadcastShape) + .set_attr("FInferType", ElemwiseIntType<2, 1>) + .set_attr("FInplaceOption", + [](const NodeAttrs& attrs) { + return std::vector >{{0, 0}, {1, 0}}; + }) + .set_attr("FCompute", BinaryBroadcastIntCompute) + .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_bitwise_right_shift"}) + .add_argument("lhs", "NDArray-or-Symbol", "First input to the function") + .add_argument("rhs", "NDArray-or-Symbol", "Second input to the function"); + +NNVM_REGISTER_OP(_npi_bitwise_right_shift_scalar) + .set_num_inputs(1) + .set_num_outputs(1) + .set_attr_parser(ParamParser) + .set_attr("FInferShape", ElemwiseShape<1, 1>) + .set_attr("FInferType", ElemwiseIntType<1, 1>) + .set_attr("FInplaceOption", + [](const NodeAttrs& attrs) { + return std::vector >{{0, 0}}; + }) + .set_attr("FCompute", BinaryScalarOp::ComputeInt) + .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_bitwise_right_shift_scalar"}) + .add_argument("data", "NDArray-or-Symbol", "source input") + .add_arguments(NumpyBinaryScalarParam::__FIELDS__()); + +NNVM_REGISTER_OP(_npi_rbitwise_right_shift_scalar) + .set_num_inputs(1) + .set_num_outputs(1) + .set_attr_parser(ParamParser) + .set_attr("FInferShape", ElemwiseShape<1, 1>) + .set_attr("FInferType", ElemwiseIntType<1, 1>) + .set_attr("FInplaceOption", + [](const NodeAttrs& attrs) { + return std::vector >{{0, 0}}; + }) + .set_attr("FCompute", BinaryScalarOp::ComputeInt) + .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_rbitwise_right_shift_scalar"}) + .add_argument("data", "NDArray-or-Symbol", "source input") + .add_arguments(NumpyBinaryScalarParam::__FIELDS__()); + +NNVM_REGISTER_OP(_backward_npi_bitwise_right_shift) + .set_num_inputs(3) + .set_num_outputs(2) + .set_attr("TIsBackward", true) + .set_attr("FInplaceOption", + [](const NodeAttrs& attrs){ + return std::vector >{{0, 1}}; + }) + .set_attr("FResourceRequest", + [](const NodeAttrs& attrs) { + return std::vector{ResourceRequest::kTempSpace}; + }) + .set_attr("FCompute", + BinaryBroadcastBackwardUseIn); + +NNVM_REGISTER_OP(_backward_npi_bitwise_right_shift_scalar) + .add_arguments(NumpyBinaryScalarParam::__FIELDS__()) + .set_attr_parser(ParamParser) + .set_attr("FCompute", BinaryScalarOp::Backward); + +MXNET_OPERATOR_REGISTER_BINARY(_backward_npi_rbitwise_right_shift_scalar) + .add_arguments(NumpyBinaryScalarParam::__FIELDS__()) + .set_attr_parser(ParamParser) + .set_attr("FCompute", BinaryScalarOp::Backward); + +} // namespace op +} // namespace mxnet diff --git a/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cu b/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cu new file mode 100644 index 000000000000..6f7ae70af545 --- /dev/null +++ b/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cu @@ -0,0 +1,70 @@ +/* + * Licensed to the Apache Software Foundation (ASF) under one + * or more contributor license agreements. See the NOTICE file + * distributed with this work for additional information + * regarding copyright ownership. The ASF licenses this file + * to you 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. + */ + +/*! + * Copyright (c) 2019 by Contributors + * \file np_elemwise_broadcast_op_extended_thi.cu + * \brief GPU Implementation of extended functions for elementwise binary broadcast operator. (Third extended file) + */ + +#include "./np_elemwise_broadcast_op.h" + +namespace mxnet { +namespace op { + +NNVM_REGISTER_OP(_npi_bitwise_left_shift) + .set_attr("FCompute", BinaryBroadcastRTCCompute{"bitwise_left_shift"}); + +NNVM_REGISTER_OP(_npi_bitwise_left_shift_scalar) + .set_attr("FCompute", BinaryScalarRTCCompute{"bitwise_left_shift"}); + +NNVM_REGISTER_OP(_npi_rbitwise_left_shift_scalar) + .set_attr("FCompute", BinaryScalarRTCCompute{"rbitwise_left_shift"}); + +NNVM_REGISTER_OP(_backward_npi_bitwise_left_shift) + .set_attr("FCompute", BinaryBroadcastRTCBackwardUseIn{"bitwise_left_shift_grad", + "bitwise_left_shift_rgrad"}); + +NNVM_REGISTER_OP(_backward_npi_bitwise_left_shift_scalar) + .set_attr("FCompute", BinaryScalarRTCBackward{"bitwise_left_shift_grad"}); + +NNVM_REGISTER_OP(_backward_npi_rbitwise_left_shift_scalar) + .set_attr("FCompute", BinaryScalarRTCBackward{"rbitwise_left_shift_grad"}); + +NNVM_REGISTER_OP(_npi_bitwise_right_shift) + .set_attr("FCompute", BinaryBroadcastRTCCompute{"bitwise_right_shift"}); + +NNVM_REGISTER_OP(_npi_bitwise_right_shift_scalar) + .set_attr("FCompute", BinaryScalarRTCCompute{"bitwise_right_shift"}); + +NNVM_REGISTER_OP(_npi_rbitwise_right_shift_scalar) + .set_attr("FCompute", BinaryScalarRTCCompute{"rbitwise_right_shift"}); + +NNVM_REGISTER_OP(_backward_npi_bitwise_right_shift) + .set_attr("FCompute", BinaryBroadcastRTCBackwardUseIn{"bitwise_right_shift_grad", + "bitwise_right_shift_rgrad"}); + +NNVM_REGISTER_OP(_backward_npi_bitwise_right_shift_scalar) + .set_attr("FCompute", BinaryScalarRTCBackward{"bitwise_right_shift_grad"}); + +NNVM_REGISTER_OP(_backward_npi_rbitwise_right_shift_scalar) + .set_attr("FCompute", BinaryScalarRTCBackward{"rbitwise_right_shift_grad"}); + +} // namespace op +} // namespace mxnet From 15f1768909a69d47fe04a6234bbd4bf87c4d3017 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Thu, 16 Sep 2021 16:02:54 -0700 Subject: [PATCH 02/18] add more methods --- python/mxnet/numpy/multiarray.py | 30 ++++++++++++++++++++++++++++++ 1 file changed, 30 insertions(+) diff --git a/python/mxnet/numpy/multiarray.py b/python/mxnet/numpy/multiarray.py index 1754122985f7..5d52a67c101c 100644 --- a/python/mxnet/numpy/multiarray.py +++ b/python/mxnet/numpy/multiarray.py @@ -1037,6 +1037,16 @@ def __xor__(self, other): """x.__xor__(y) <=> x ^ y""" return bitwise_xor(self, other) + @wrap_mxnp_np_ufunc + def __lshift__(self, other): + """x.__lshift__(y) <=> x << y""" + return bitwise_left_shift(self, other) + + @wrap_mxnp_np_ufunc + def __rshift__(self, other): + """x.__rshift__(y) <=> x >> y""" + return bitwise_right_shift(self, other) + @wrap_mxnp_np_ufunc def __iand__(self, other): """x.__iand__(y) <=> x &= y""" @@ -1052,6 +1062,26 @@ def __ixor__(self, other): """x.__ixor__(y) <=> x ^= y""" return bitwise_xor(self, other, out=self) + @wrap_mxnp_np_ufunc + def __ilshift__(self, other): + """x.__ilshift__(y) <=> x <<= y""" + return bitwise_left_shift(self, other, out=self) + + @wrap_mxnp_np_ufunc + def __irshift__(self, other): + """x.__irshift__(y) <=> x >>= y""" + return bitwise_right_shift(self, other, out=self) + + @wrap_mxnp_np_ufunc + def __rlshift__(self, other): + """x.__rlshift__(y) <=> y << x""" + return bitwise_left_shift(other, self) + + @wrap_mxnp_np_ufunc + def __rrshift__(self, other): + """x.__rrshift__(y) <=> y >> x""" + return bitwise_right_shift(other, self) + def __round__(self, n=0): """x.__round__(n)""" return round(self, decimals=n) From 46e6e1f2f0c308718e5e91bfdbacdb0f7052022c Mon Sep 17 00:00:00 2001 From: barry-jin Date: Thu, 16 Sep 2021 16:05:13 -0700 Subject: [PATCH 03/18] add mshadow_op.h --- src/operator/mshadow_op.h | 26 ++++++++++++++++++++++++++ 1 file changed, 26 insertions(+) diff --git a/src/operator/mshadow_op.h b/src/operator/mshadow_op.h index eee928452f50..bc48c65a062f 100644 --- a/src/operator/mshadow_op.h +++ b/src/operator/mshadow_op.h @@ -690,6 +690,32 @@ MXNET_BINARY_MATH_OP(bitwise_xor, static_cast(a) ^ static_cast MXNET_BINARY_MATH_OP(bitwise_or, static_cast(a) | static_cast(b)); +/*! \brief used for generate element of bitwise_left_shift */ +MXNET_BINARY_MATH_OP(bitwise_left_shift, static_cast(a) << static_cast(b)); + +MXNET_BINARY_MATH_OP(bitwise_left_shift_grad, math::pow(2.0f, static_cast(b))); + +MXNET_BINARY_MATH_OP(bitwise_left_shift_rgrad, static_cast(a) * \ +math::pow(2.0f, static_cast(b)) * math::log(2.0f)); + +MXNET_BINARY_MATH_OP(rbitwise_left_shift, static_cast(b) << static_cast(a)); + +MXNET_BINARY_MATH_OP(rbitwise_left_shift_grad, static_cast(b) * \ +math::pow(2.0f, static_cast(a)) * math::log(2.0f)); + +/*! \brief used for generate element of bitwise_right_shift */ +MXNET_BINARY_MATH_OP(bitwise_right_shift, static_cast(a) >> static_cast(b)); + +MXNET_BINARY_MATH_OP(bitwise_right_shift_grad, math::pow(0.5f, static_cast(b))); + +MXNET_BINARY_MATH_OP(bitwise_right_shift_rgrad, static_cast(a) * \ +math::pow(0.5f, static_cast(b)) * math::log(0.5f)); + +MXNET_BINARY_MATH_OP(rbitwise_right_shift, static_cast(b) >> static_cast(a)); + +MXNET_BINARY_MATH_OP(rbitwise_right_shift_grad, -static_cast(b) * \ +math::pow(0.5f, static_cast(a)) * math::log(0.5f)); + MXNET_UNARY_MATH_OP(square_root, math::sqrt(a)); MXNET_UNARY_MATH_OP(square_root_grad, 0.5f / math::id(a)); From 27277dc6f8db09e4fb589b4cefbf8997c7167bb4 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Thu, 16 Sep 2021 16:06:20 -0700 Subject: [PATCH 04/18] fix --- src/operator/mshadow_op.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/operator/mshadow_op.h b/src/operator/mshadow_op.h index bc48c65a062f..320779a62617 100644 --- a/src/operator/mshadow_op.h +++ b/src/operator/mshadow_op.h @@ -713,7 +713,7 @@ math::pow(0.5f, static_cast(b)) * math::log(0.5f)); MXNET_BINARY_MATH_OP(rbitwise_right_shift, static_cast(b) >> static_cast(a)); -MXNET_BINARY_MATH_OP(rbitwise_right_shift_grad, -static_cast(b) * \ +MXNET_BINARY_MATH_OP(rbitwise_right_shift_grad, static_cast(b) * \ math::pow(0.5f, static_cast(a)) * math::log(0.5f)); MXNET_UNARY_MATH_OP(square_root, math::sqrt(a)); From ab41f89f92951749fd010473e330e398c1c66673 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Fri, 17 Sep 2021 17:01:20 -0700 Subject: [PATCH 05/18] fix lint & add tests --- .../np_elemwise_broadcast_op_extended_thi.cc | 42 ++++++++++----- .../np_elemwise_broadcast_op_extended_thi.cu | 10 ++-- .../unittest/test_numpy_interoperability.py | 26 ++++++++++ tests/python/unittest/test_numpy_op.py | 52 +++++++++++++++++++ 4 files changed, 112 insertions(+), 18 deletions(-) diff --git a/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cc b/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cc index 23b9b878b7a3..60ff3dcca9fa 100644 --- a/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cc +++ b/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cc @@ -57,7 +57,8 @@ NNVM_REGISTER_OP(_npi_bitwise_left_shift) [](const NodeAttrs& attrs) { return std::vector >{{0, 0}, {1, 0}}; }) - .set_attr("FCompute", BinaryBroadcastIntCompute) + .set_attr("FCompute", + BinaryBroadcastIntCompute) .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_bitwise_left_shift"}) .add_argument("lhs", "NDArray-or-Symbol", "First input to the function") .add_argument("rhs", "NDArray-or-Symbol", "Second input to the function"); @@ -72,8 +73,10 @@ NNVM_REGISTER_OP(_npi_bitwise_left_shift_scalar) [](const NodeAttrs& attrs) { return std::vector >{{0, 0}}; }) - .set_attr("FCompute", BinaryScalarOp::ComputeInt) - .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_bitwise_left_shift_scalar"}) + .set_attr("FCompute", + BinaryScalarOp::ComputeInt) + .set_attr("FGradient", + ElemwiseGradUseIn{"_backward_npi_bitwise_left_shift_scalar"}) .add_argument("data", "NDArray-or-Symbol", "source input") .add_arguments(NumpyBinaryScalarParam::__FIELDS__()); @@ -87,8 +90,10 @@ NNVM_REGISTER_OP(_npi_rbitwise_left_shift_scalar) [](const NodeAttrs& attrs) { return std::vector >{{0, 0}}; }) - .set_attr("FCompute", BinaryScalarOp::ComputeInt) - .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_rbitwise_left_shift_scalar"}) + .set_attr("FCompute", + BinaryScalarOp::ComputeInt) + .set_attr("FGradient", + ElemwiseGradUseIn{"_backward_npi_rbitwise_left_shift_scalar"}) .add_argument("data", "NDArray-or-Symbol", "source input") .add_arguments(NumpyBinaryScalarParam::__FIELDS__()); @@ -111,12 +116,14 @@ NNVM_REGISTER_OP(_backward_npi_bitwise_left_shift) NNVM_REGISTER_OP(_backward_npi_bitwise_left_shift_scalar) .add_arguments(NumpyBinaryScalarParam::__FIELDS__()) .set_attr_parser(ParamParser) - .set_attr("FCompute", BinaryScalarOp::Backward); + .set_attr("FCompute", + BinaryScalarOp::Backward); MXNET_OPERATOR_REGISTER_BINARY(_backward_npi_rbitwise_left_shift_scalar) .add_arguments(NumpyBinaryScalarParam::__FIELDS__()) .set_attr_parser(ParamParser) - .set_attr("FCompute", BinaryScalarOp::Backward); + .set_attr("FCompute", + BinaryScalarOp::Backward); NNVM_REGISTER_OP(_npi_bitwise_right_shift) @@ -132,7 +139,8 @@ NNVM_REGISTER_OP(_npi_bitwise_right_shift) [](const NodeAttrs& attrs) { return std::vector >{{0, 0}, {1, 0}}; }) - .set_attr("FCompute", BinaryBroadcastIntCompute) + .set_attr("FCompute", + BinaryBroadcastIntCompute) .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_bitwise_right_shift"}) .add_argument("lhs", "NDArray-or-Symbol", "First input to the function") .add_argument("rhs", "NDArray-or-Symbol", "Second input to the function"); @@ -147,8 +155,10 @@ NNVM_REGISTER_OP(_npi_bitwise_right_shift_scalar) [](const NodeAttrs& attrs) { return std::vector >{{0, 0}}; }) - .set_attr("FCompute", BinaryScalarOp::ComputeInt) - .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_bitwise_right_shift_scalar"}) + .set_attr("FCompute", + BinaryScalarOp::ComputeInt) + .set_attr("FGradient", + ElemwiseGradUseIn{"_backward_npi_bitwise_right_shift_scalar"}) .add_argument("data", "NDArray-or-Symbol", "source input") .add_arguments(NumpyBinaryScalarParam::__FIELDS__()); @@ -162,8 +172,10 @@ NNVM_REGISTER_OP(_npi_rbitwise_right_shift_scalar) [](const NodeAttrs& attrs) { return std::vector >{{0, 0}}; }) - .set_attr("FCompute", BinaryScalarOp::ComputeInt) - .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_rbitwise_right_shift_scalar"}) + .set_attr("FCompute", + BinaryScalarOp::ComputeInt) + .set_attr("FGradient", + ElemwiseGradUseIn{"_backward_npi_rbitwise_right_shift_scalar"}) .add_argument("data", "NDArray-or-Symbol", "source input") .add_arguments(NumpyBinaryScalarParam::__FIELDS__()); @@ -186,12 +198,14 @@ NNVM_REGISTER_OP(_backward_npi_bitwise_right_shift) NNVM_REGISTER_OP(_backward_npi_bitwise_right_shift_scalar) .add_arguments(NumpyBinaryScalarParam::__FIELDS__()) .set_attr_parser(ParamParser) - .set_attr("FCompute", BinaryScalarOp::Backward); + .set_attr("FCompute", + BinaryScalarOp::Backward); MXNET_OPERATOR_REGISTER_BINARY(_backward_npi_rbitwise_right_shift_scalar) .add_arguments(NumpyBinaryScalarParam::__FIELDS__()) .set_attr_parser(ParamParser) - .set_attr("FCompute", BinaryScalarOp::Backward); + .set_attr("FCompute", + BinaryScalarOp::Backward); } // namespace op } // namespace mxnet diff --git a/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cu b/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cu index 6f7ae70af545..00db6b4b6110 100644 --- a/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cu +++ b/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cu @@ -38,8 +38,9 @@ NNVM_REGISTER_OP(_npi_rbitwise_left_shift_scalar) .set_attr("FCompute", BinaryScalarRTCCompute{"rbitwise_left_shift"}); NNVM_REGISTER_OP(_backward_npi_bitwise_left_shift) - .set_attr("FCompute", BinaryBroadcastRTCBackwardUseIn{"bitwise_left_shift_grad", - "bitwise_left_shift_rgrad"}); + .set_attr("FCompute", + BinaryBroadcastRTCBackwardUseIn{"bitwise_left_shift_grad", + "bitwise_left_shift_rgrad"}); NNVM_REGISTER_OP(_backward_npi_bitwise_left_shift_scalar) .set_attr("FCompute", BinaryScalarRTCBackward{"bitwise_left_shift_grad"}); @@ -57,8 +58,9 @@ NNVM_REGISTER_OP(_npi_rbitwise_right_shift_scalar) .set_attr("FCompute", BinaryScalarRTCCompute{"rbitwise_right_shift"}); NNVM_REGISTER_OP(_backward_npi_bitwise_right_shift) - .set_attr("FCompute", BinaryBroadcastRTCBackwardUseIn{"bitwise_right_shift_grad", - "bitwise_right_shift_rgrad"}); + .set_attr("FCompute", + BinaryBroadcastRTCBackwardUseIn{"bitwise_right_shift_grad", + "bitwise_right_shift_rgrad"}); NNVM_REGISTER_OP(_backward_npi_bitwise_right_shift_scalar) .set_attr("FCompute", BinaryScalarRTCBackward{"bitwise_right_shift_grad"}); diff --git a/tests/python/unittest/test_numpy_interoperability.py b/tests/python/unittest/test_numpy_interoperability.py index db643d1090dc..e1b12e7c65b4 100644 --- a/tests/python/unittest/test_numpy_interoperability.py +++ b/tests/python/unittest/test_numpy_interoperability.py @@ -1530,6 +1530,30 @@ def _add_workload_bitwise_xor(): OpArgMngr.add_workload('bitwise_xor', ones, ones) +def _add_workload_bitwise_left_shift(): + for dtype in [np.int8, np.int32, np.int64]: + twenty = np.array([20], dtype=dtype) + three = np.array([3], dtype=dtype) + OpArgMngr.add_workload('bitwise_left_shift', twenty, three) + OpArgMngr.add_workload('bitwise_left_shift', twenty, three) + OpArgMngr.add_workload('bitwise_left_shift', twenty, three) + OpArgMngr.add_workload('bitwise_left_shift', twenty, three) + OpArgMngr.add_workload('bitwise_left_shift', np.array([9223372036854775807], np.int64), np.array([1], np.int64)) + OpArgMngr.add_workload('bitwise_left_shift', np.array([-9223372036854775808], np.int64), np.array([1], np.int64)) + + +def _add_workload_bitwise_right_shift(): + for dtype in [np.int8, np.int32, np.int64]: + twenty = np.array([20], dtype=dtype) + three = np.array([3], dtype=dtype) + OpArgMngr.add_workload('bitwise_right_shift', twenty, three) + OpArgMngr.add_workload('bitwise_right_shift', twenty, three) + OpArgMngr.add_workload('bitwise_right_shift', twenty, three) + OpArgMngr.add_workload('bitwise_right_shift', twenty, three) + OpArgMngr.add_workload('bitwise_right_shift', np.array([9223372036854775807], np.int64), np.array([1], np.int64)) + OpArgMngr.add_workload('bitwise_right_shift', np.array([-9223372036854775808], np.int64), np.array([1], np.int64)) + + def _add_workload_ldexp(): OpArgMngr.add_workload('ldexp', np.array(2., np.float32), np.array(3, np.int8)) OpArgMngr.add_workload('ldexp', np.array(2., np.float64), np.array(3, np.int8)) @@ -3081,6 +3105,8 @@ def _prepare_workloads(): _add_workload_bitwise_and() _add_workload_bitwise_xor() _add_workload_bitwise_or() + _add_workload_bitwise_left_shift() + _add_workload_bitwise_right_shift() _add_workload_ldexp() _add_workload_subtract(array_pool) _add_workload_multiply(array_pool) diff --git a/tests/python/unittest/test_numpy_op.py b/tests/python/unittest/test_numpy_op.py index e21e8fdc49b8..40a35b7f5288 100644 --- a/tests/python/unittest/test_numpy_op.py +++ b/tests/python/unittest/test_numpy_op.py @@ -2992,6 +2992,58 @@ def forward(self, a, *args, **kwargs): check_unary_func(func, shape, low, high) +@use_np +@pytest.mark.parametrize('ndim', [2, 3, 4]) +@pytest.mark.parametrize('func,low,high', [ + ('left_shift', -5, 5), + ('right_shift', -5, 5), +]) +def test_np_bitwise_shift(func, low, high, ndim): + def check_unary_func(func, shape, low, high): + class TestUnary(HybridBlock): + def __init__(self, func): + super(TestUnary, self).__init__() + self._func = func + + def forward(self, a, b, *args, **kwargs): + return getattr(np, self._func)(a, b) + + np_func = getattr(onp, func) + mx_func = TestUnary("bitwise_" + func) + np_test_data1 = onp.random.randint(low, high, shape).astype(onp.int64) + np_test_data2 = onp.random.randint(low + 5, high + 5, shape).astype(onp.int64) + mx_test_data1 = mx.numpy.array(np_test_data1).astype(onp.int64) + mx_test_data2 = mx.numpy.array(np_test_data2).astype(onp.int64) + for hybridize in [True, False]: + if hybridize: + mx_func.hybridize() + np_out = np_func(np_test_data1, np_test_data2) + with mx.autograd.record(): + y = mx_func(mx_test_data1, mx_test_data2) + assert y.shape == np_out.shape + assert_almost_equal(y.asnumpy(), np_out, rtol=1e-3, atol=1e-5) + if np_out.dtype == np.bool_: + assert y.dtype == np.bool_ + + np_out = getattr(onp, func)(np_test_data1, np_test_data2) + mx_out = getattr(mx.np, "bitwise_" + func)(mx_test_data1, mx_test_data2) + assert mx_out.shape == np_out.shape + assert_almost_equal(mx_out.asnumpy(), np_out, rtol=1e-3, atol=1e-5) + + assertRaises(TypeError, getattr(np, "bitwise_" + func), mx_test_data1, mx_test_data2, where=False) + assertRaises(TypeError, getattr(np, "bitwise_" + func), mx_test_data1, mx_test_data2, subok=False) + assertRaises(TypeError, getattr(np, "bitwise_" + func), mx_test_data1, mx_test_data2, dtype=onp.int8) + assertRaises(TypeError, getattr(np, "bitwise_" + func), mx_test_data1, mx_test_data2, dtype="abcdefg") + assertRaises(TypeError, getattr(np, "bitwise_" + func), mx_test_data1, mx_test_data2, casting='safe') + assertRaises(TypeError, getattr(np, "bitwise_" + func), mx_test_data1, mx_test_data2, casting='mxnet') + assertRaises(TypeError, getattr(np, "bitwise_" + func), mx_test_data1, mx_test_data2, order='C') + assertRaises(TypeError, getattr(np, "bitwise_" + func), mx_test_data1, mx_test_data2, order='mxnet') + + shape = random.choice([rand_shape_nd(ndim, dim=3), (1, 0, 2)]) + for shape in [rand_shape_nd(ndim, dim=3), (1, 0, 2)]: + check_unary_func(func, shape, low, high) + + @use_np def test_np_binary_funcs(): def check_binary_func(func, lshape, rshape, low, high, lgrads, rgrads=None, alltypes=None): From 1305a9173c31ecf68ddbef2dc65848d0371dfc34 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Fri, 17 Sep 2021 17:55:08 -0700 Subject: [PATCH 06/18] fix --- src/operator/mshadow_op.h | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/src/operator/mshadow_op.h b/src/operator/mshadow_op.h index 320779a62617..3baf845a037c 100644 --- a/src/operator/mshadow_op.h +++ b/src/operator/mshadow_op.h @@ -690,6 +690,12 @@ MXNET_BINARY_MATH_OP(bitwise_xor, static_cast(a) ^ static_cast MXNET_BINARY_MATH_OP(bitwise_or, static_cast(a) | static_cast(b)); +#pragma GCC diagnostic push +#if __GNUC__ >= 7 +#pragma GCC diagnostic ignored "-Wint-in-bool-context" +#pragma GCC diagnostic ignored "-Wbool-compare" +#endif + /*! \brief used for generate element of bitwise_left_shift */ MXNET_BINARY_MATH_OP(bitwise_left_shift, static_cast(a) << static_cast(b)); @@ -716,6 +722,8 @@ MXNET_BINARY_MATH_OP(rbitwise_right_shift, static_cast(b) >> static_cas MXNET_BINARY_MATH_OP(rbitwise_right_shift_grad, static_cast(b) * \ math::pow(0.5f, static_cast(a)) * math::log(0.5f)); +#pragma GCC diagnostic pop + MXNET_UNARY_MATH_OP(square_root, math::sqrt(a)); MXNET_UNARY_MATH_OP(square_root_grad, 0.5f / math::id(a)); From 3805946ed0b203f8081f18bba9d51a6f2477a980 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Mon, 20 Sep 2021 13:38:09 -0700 Subject: [PATCH 07/18] update operator_tune.cc --- .../numpy/np_elemwise_broadcast_op_extended_thi.cc | 12 ++++++------ src/operator/operator_tune.cc | 10 ++++++++++ 2 files changed, 16 insertions(+), 6 deletions(-) diff --git a/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cc b/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cc index 60ff3dcca9fa..b8447f547545 100644 --- a/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cc +++ b/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cc @@ -58,7 +58,7 @@ NNVM_REGISTER_OP(_npi_bitwise_left_shift) return std::vector >{{0, 0}, {1, 0}}; }) .set_attr("FCompute", - BinaryBroadcastIntCompute) + BinaryBroadcastCompute) .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_bitwise_left_shift"}) .add_argument("lhs", "NDArray-or-Symbol", "First input to the function") .add_argument("rhs", "NDArray-or-Symbol", "Second input to the function"); @@ -74,7 +74,7 @@ NNVM_REGISTER_OP(_npi_bitwise_left_shift_scalar) return std::vector >{{0, 0}}; }) .set_attr("FCompute", - BinaryScalarOp::ComputeInt) + BinaryScalarOp::Compute) .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_bitwise_left_shift_scalar"}) .add_argument("data", "NDArray-or-Symbol", "source input") @@ -91,7 +91,7 @@ NNVM_REGISTER_OP(_npi_rbitwise_left_shift_scalar) return std::vector >{{0, 0}}; }) .set_attr("FCompute", - BinaryScalarOp::ComputeInt) + BinaryScalarOp::Compute) .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_rbitwise_left_shift_scalar"}) .add_argument("data", "NDArray-or-Symbol", "source input") @@ -140,7 +140,7 @@ NNVM_REGISTER_OP(_npi_bitwise_right_shift) return std::vector >{{0, 0}, {1, 0}}; }) .set_attr("FCompute", - BinaryBroadcastIntCompute) + BinaryBroadcastCompute) .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_bitwise_right_shift"}) .add_argument("lhs", "NDArray-or-Symbol", "First input to the function") .add_argument("rhs", "NDArray-or-Symbol", "Second input to the function"); @@ -156,7 +156,7 @@ NNVM_REGISTER_OP(_npi_bitwise_right_shift_scalar) return std::vector >{{0, 0}}; }) .set_attr("FCompute", - BinaryScalarOp::ComputeInt) + BinaryScalarOp::Compute) .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_bitwise_right_shift_scalar"}) .add_argument("data", "NDArray-or-Symbol", "source input") @@ -173,7 +173,7 @@ NNVM_REGISTER_OP(_npi_rbitwise_right_shift_scalar) return std::vector >{{0, 0}}; }) .set_attr("FCompute", - BinaryScalarOp::ComputeInt) + BinaryScalarOp::Compute) .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_rbitwise_right_shift_scalar"}) .add_argument("data", "NDArray-or-Symbol", "source input") diff --git a/src/operator/operator_tune.cc b/src/operator/operator_tune.cc index f83885da1ba1..15e7887f3e01 100644 --- a/src/operator/operator_tune.cc +++ b/src/operator/operator_tune.cc @@ -449,6 +449,16 @@ IMPLEMENT_BINARY_WORKLOAD_BWD(mxnet::op::mshadow_op::logical_xor); IMPLEMENT_BINARY_WORKLOAD_FWD_WITH_BOOL(mxnet::op::mshadow_op::bitwise_and); // NOLINT() IMPLEMENT_BINARY_WORKLOAD_FWD_WITH_BOOL(mxnet::op::mshadow_op::bitwise_xor); // NOLINT() IMPLEMENT_BINARY_WORKLOAD_FWD_WITH_BOOL(mxnet::op::mshadow_op::bitwise_or); // NOLINT() +IMPLEMENT_BINARY_WORKLOAD_FWD(mxnet::op::mshadow_op::bitwise_left_shift); // NOLINT() +IMPLEMENT_BINARY_WORKLOAD_FWD(mxnet::op::mshadow_op::rbitwise_left_shift); // NOLINT() +IMPLEMENT_BINARY_WORKLOAD_BWD(mxnet::op::mshadow_op::bitwise_left_shift_grad); // NOLINT() +IMPLEMENT_BINARY_WORKLOAD_BWD(mxnet::op::mshadow_op::bitwise_left_shift_rgrad); // NOLINT() +IMPLEMENT_BINARY_WORKLOAD_BWD(mxnet::op::mshadow_op::rbitwise_left_shift_grad); // NOLINT() +IMPLEMENT_BINARY_WORKLOAD_FWD(mxnet::op::mshadow_op::bitwise_right_shift); // NOLINT() +IMPLEMENT_BINARY_WORKLOAD_FWD(mxnet::op::mshadow_op::rbitwise_right_shift); // NOLINT() +IMPLEMENT_BINARY_WORKLOAD_BWD(mxnet::op::mshadow_op::bitwise_right_shift_grad); // NOLINT() +IMPLEMENT_BINARY_WORKLOAD_BWD(mxnet::op::mshadow_op::bitwise_right_shift_rgrad); // NOLINT() +IMPLEMENT_BINARY_WORKLOAD_BWD(mxnet::op::mshadow_op::rbitwise_right_shift_grad); // NOLINT() IMPLEMENT_BINARY_WORKLOAD_FWD(mxnet::op::mshadow_op::smooth_l1_loss); // NOLINT() IMPLEMENT_BINARY_WORKLOAD_BWD(mxnet::op::mshadow_op::smooth_l1_gradient); // NOLINT() IMPLEMENT_BINARY_WORKLOAD_FWD_WITH_BOOL(mxnet::op::mshadow_op::gcd); // NOLINT() From 9e476206f9ea695d5563ca1da81fe119b4bfc8e3 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Mon, 20 Sep 2021 15:37:46 -0700 Subject: [PATCH 08/18] update amp list --- python/mxnet/amp/lists/symbol_fp16.py | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/python/mxnet/amp/lists/symbol_fp16.py b/python/mxnet/amp/lists/symbol_fp16.py index d942051c0398..ef44ca05cdf3 100644 --- a/python/mxnet/amp/lists/symbol_fp16.py +++ b/python/mxnet/amp/lists/symbol_fp16.py @@ -195,6 +195,12 @@ '_npi_bitwise_or_scalar', '_npi_bitwise_xor', '_npi_bitwise_xor_scalar', + '_npi_bitwise_left_shift', + '_npi_bitwise_left_shift_scalar', + '_npi_bitwise_right_shift', + '_npi_bitwise_right_shift_scalar', + '_npi_rbitwise_left_shift_scalar', + '_npi_rbitwise_right_shift_scalar', '_npi_blackman', '_npi_boolean_mask_assign_scalar', '_npi_boolean_mask_assign_tensor', From a9173290eeef104d09e753fca8b936b6ff1f2c7d Mon Sep 17 00:00:00 2001 From: barry-jin Date: Tue, 21 Sep 2021 10:54:29 -0700 Subject: [PATCH 09/18] add rtc functions --- src/common/cuda/rtc/backward_functions-inl.h | 42 ++++++++++++++++++++ src/common/cuda/rtc/forward_functions-inl.h | 32 +++++++++++++++ 2 files changed, 74 insertions(+) diff --git a/src/common/cuda/rtc/backward_functions-inl.h b/src/common/cuda/rtc/backward_functions-inl.h index 50d8469571d8..e2dbe6882534 100644 --- a/src/common/cuda/rtc/backward_functions-inl.h +++ b/src/common/cuda/rtc/backward_functions-inl.h @@ -426,6 +426,48 @@ copysign_grad(const DType val, return (val >= 0 && val2 >= 0) || (val < 0 && val2 < 0) ? 1 : -1; } +template +__device__ inline mixed_type +bitwise_left_shift_grad(const DType val, + const DType2 val2) { + return op::power(static_cast(2), val2); +} + +template +__device__ inline mixed_type +bitwise_left_shift_rgrad(const DType val, + const DType2 val2) { + return val * op::power(static_cast(2), val2) * op::log(static_cast(2)); +} + +template +__device__ inline mixed_type +rbitwise_left_shift_grad(const DType val, + const DType2 val2) { + return val2 * op::power(static_cast(2), val) * op::log(static_cast(2)); +} + +template +__device__ inline mixed_type +bitwise_right_shift_grad(const DType val, + const DType2 val2) { + return op::power(0.5f, val2); +} + +template +__device__ inline mixed_type +bitwise_right_shift_rgrad(const DType val, + const DType2 val2) { + return val * op::power(0.5f, val2) * op::log(0.5f); +} + +template +__device__ inline mixed_type +rbitwise_right_shift_grad(const DType val, + const DType2 val2) { + return val2 * op::power(0.5f, val) * op::log(0.5f); +} + template __device__ inline mixed_type arctan2_grad(const DType val, diff --git a/src/common/cuda/rtc/forward_functions-inl.h b/src/common/cuda/rtc/forward_functions-inl.h index 7a886a0a9aec..3bf000372e76 100644 --- a/src/common/cuda/rtc/forward_functions-inl.h +++ b/src/common/cuda/rtc/forward_functions-inl.h @@ -597,6 +597,38 @@ __device__ inline mixed_type bitwise_and(const DType a, return real_a & real_b; } +template +__device__ inline mixed_type bitwise_left_shift(const DType a, + const DType2 b) { + const mixed_type real_a = a; + const mixed_type real_b = b; + return real_a << real_b; +} + +template +__device__ inline mixed_type rbitwise_left_shift(const DType a, + const DType2 b) { + const mixed_type real_a = a; + const mixed_type real_b = b; + return real_b << real_a; +} + +template +__device__ inline mixed_type bitwise_right_shift(const DType a, + const DType2 b) { + const mixed_type real_a = a; + const mixed_type real_b = b; + return real_a >> real_b; +} + +template +__device__ inline mixed_type rbitwise_right_shift(const DType a, + const DType2 b) { + const mixed_type real_a = a; + const mixed_type real_b = b; + return real_b >> real_a; +} + DEFINE_BINARY_MATH_FUNC(arctan2, ::atan2, ::atan2f) template From 8c99fff852836ac897092020e79be429195f8eda Mon Sep 17 00:00:00 2001 From: barry-jin Date: Mon, 11 Oct 2021 20:03:23 -0700 Subject: [PATCH 10/18] fix bitwise rtc functions & numpy op gpu test overriding issue --- src/common/cuda/rtc/backward_functions-inl.h | 2 ++ tests/python/gpu/{test_numpy_op.py => test_numpy_einsum.py} | 0 2 files changed, 2 insertions(+) rename tests/python/gpu/{test_numpy_op.py => test_numpy_einsum.py} (100%) diff --git a/src/common/cuda/rtc/backward_functions-inl.h b/src/common/cuda/rtc/backward_functions-inl.h index e2dbe6882534..4b0bbc141f26 100644 --- a/src/common/cuda/rtc/backward_functions-inl.h +++ b/src/common/cuda/rtc/backward_functions-inl.h @@ -437,6 +437,7 @@ template __device__ inline mixed_type bitwise_left_shift_rgrad(const DType val, const DType2 val2) { + using type = mixed_type; return val * op::power(static_cast(2), val2) * op::log(static_cast(2)); } @@ -444,6 +445,7 @@ template __device__ inline mixed_type rbitwise_left_shift_grad(const DType val, const DType2 val2) { + using type = mixed_type; return val2 * op::power(static_cast(2), val) * op::log(static_cast(2)); } diff --git a/tests/python/gpu/test_numpy_op.py b/tests/python/gpu/test_numpy_einsum.py similarity index 100% rename from tests/python/gpu/test_numpy_op.py rename to tests/python/gpu/test_numpy_einsum.py From 519e86ff138291e51de6291a7f19e57cf0a753f4 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Mon, 11 Oct 2021 20:18:12 -0700 Subject: [PATCH 11/18] clang-format --- src/operator/mshadow_op.h | 20 +++++++++++-------- .../np_elemwise_broadcast_op_extended_thi.cc | 14 +++++++------ .../np_elemwise_broadcast_op_extended_thi.cu | 3 ++- 3 files changed, 22 insertions(+), 15 deletions(-) diff --git a/src/operator/mshadow_op.h b/src/operator/mshadow_op.h index df294bc66c1b..036ed9045466 100644 --- a/src/operator/mshadow_op.h +++ b/src/operator/mshadow_op.h @@ -700,26 +700,30 @@ MXNET_BINARY_MATH_OP(bitwise_left_shift, static_cast(a) << static_cast< MXNET_BINARY_MATH_OP(bitwise_left_shift_grad, math::pow(2.0f, static_cast(b))); -MXNET_BINARY_MATH_OP(bitwise_left_shift_rgrad, static_cast(a) * \ -math::pow(2.0f, static_cast(b)) * math::log(2.0f)); +MXNET_BINARY_MATH_OP(bitwise_left_shift_rgrad, + static_cast(a) * math::pow(2.0f, static_cast(b)) * + math::log(2.0f)); MXNET_BINARY_MATH_OP(rbitwise_left_shift, static_cast(b) << static_cast(a)); -MXNET_BINARY_MATH_OP(rbitwise_left_shift_grad, static_cast(b) * \ -math::pow(2.0f, static_cast(a)) * math::log(2.0f)); +MXNET_BINARY_MATH_OP(rbitwise_left_shift_grad, + static_cast(b) * math::pow(2.0f, static_cast(a)) * + math::log(2.0f)); /*! \brief used for generate element of bitwise_right_shift */ MXNET_BINARY_MATH_OP(bitwise_right_shift, static_cast(a) >> static_cast(b)); MXNET_BINARY_MATH_OP(bitwise_right_shift_grad, math::pow(0.5f, static_cast(b))); -MXNET_BINARY_MATH_OP(bitwise_right_shift_rgrad, static_cast(a) * \ -math::pow(0.5f, static_cast(b)) * math::log(0.5f)); +MXNET_BINARY_MATH_OP(bitwise_right_shift_rgrad, + static_cast(a) * math::pow(0.5f, static_cast(b)) * + math::log(0.5f)); MXNET_BINARY_MATH_OP(rbitwise_right_shift, static_cast(b) >> static_cast(a)); -MXNET_BINARY_MATH_OP(rbitwise_right_shift_grad, static_cast(b) * \ -math::pow(0.5f, static_cast(a)) * math::log(0.5f)); +MXNET_BINARY_MATH_OP(rbitwise_right_shift_grad, + static_cast(b) * math::pow(0.5f, static_cast(a)) * + math::log(0.5f)); #pragma GCC diagnostic pop diff --git a/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cc b/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cc index b8447f547545..2bc1e4c923e5 100644 --- a/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cc +++ b/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cc @@ -20,7 +20,8 @@ /*! * Copyright (c) 2019 by Contributors * \file np_elemwise_broadcast_op_extended_thi.cc - * \brief CPU Implementation of extended functions for elementwise numpy binary broadcast operator. (Third extended file) + * \brief CPU Implementation of extended functions for elementwise numpy binary broadcast operator. + * (Third extended file) */ #include "../../common/utils.h" @@ -102,7 +103,7 @@ NNVM_REGISTER_OP(_backward_npi_bitwise_left_shift) .set_num_outputs(2) .set_attr("TIsBackward", true) .set_attr("FInplaceOption", - [](const NodeAttrs& attrs){ + [](const NodeAttrs& attrs) { return std::vector >{{0, 1}}; }) .set_attr("FResourceRequest", @@ -110,7 +111,8 @@ NNVM_REGISTER_OP(_backward_npi_bitwise_left_shift) return std::vector{ResourceRequest::kTempSpace}; }) .set_attr("FCompute", - BinaryBroadcastBackwardUseIn); NNVM_REGISTER_OP(_backward_npi_bitwise_left_shift_scalar) @@ -125,7 +127,6 @@ MXNET_OPERATOR_REGISTER_BINARY(_backward_npi_rbitwise_left_shift_scalar) .set_attr("FCompute", BinaryScalarOp::Backward); - NNVM_REGISTER_OP(_npi_bitwise_right_shift) .set_num_inputs(2) .set_num_outputs(1) @@ -184,7 +185,7 @@ NNVM_REGISTER_OP(_backward_npi_bitwise_right_shift) .set_num_outputs(2) .set_attr("TIsBackward", true) .set_attr("FInplaceOption", - [](const NodeAttrs& attrs){ + [](const NodeAttrs& attrs) { return std::vector >{{0, 1}}; }) .set_attr("FResourceRequest", @@ -192,7 +193,8 @@ NNVM_REGISTER_OP(_backward_npi_bitwise_right_shift) return std::vector{ResourceRequest::kTempSpace}; }) .set_attr("FCompute", - BinaryBroadcastBackwardUseIn); NNVM_REGISTER_OP(_backward_npi_bitwise_right_shift_scalar) diff --git a/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cu b/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cu index 00db6b4b6110..cfdc532e8ebd 100644 --- a/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cu +++ b/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cu @@ -20,7 +20,8 @@ /*! * Copyright (c) 2019 by Contributors * \file np_elemwise_broadcast_op_extended_thi.cu - * \brief GPU Implementation of extended functions for elementwise binary broadcast operator. (Third extended file) + * \brief GPU Implementation of extended functions for elementwise binary broadcast operator. (Third + * extended file) */ #include "./np_elemwise_broadcast_op.h" From a4917fb8c13c31c8a3c0e339ff7de6882aad06ec Mon Sep 17 00:00:00 2001 From: barry-jin Date: Mon, 11 Oct 2021 22:42:48 -0700 Subject: [PATCH 12/18] fix ci --- ci/docker/Dockerfile.build.ubuntu | 1 + 1 file changed, 1 insertion(+) diff --git a/ci/docker/Dockerfile.build.ubuntu b/ci/docker/Dockerfile.build.ubuntu index 7077cac976fa..843d0f97047e 100644 --- a/ci/docker/Dockerfile.build.ubuntu +++ b/ci/docker/Dockerfile.build.ubuntu @@ -58,6 +58,7 @@ RUN export DEBIAN_FRONTEND=noninteractive && \ libprotobuf-dev \ default-jdk \ clang-6.0 \ + clang-format \ python-yaml \ clang-10 \ clang-tidy-10 \ From 76dd7341aba3ef098904019e85cbcfe860e93c4f Mon Sep 17 00:00:00 2001 From: barry-jin Date: Wed, 20 Oct 2021 10:33:39 -0700 Subject: [PATCH 13/18] add int16 support --- src/operator/tensor/elemwise_binary_broadcast_op.h | 2 +- src/operator/tensor/elemwise_binary_op.h | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/src/operator/tensor/elemwise_binary_broadcast_op.h b/src/operator/tensor/elemwise_binary_broadcast_op.h index 9bfcbc73163a..4883ab619ba8 100644 --- a/src/operator/tensor/elemwise_binary_broadcast_op.h +++ b/src/operator/tensor/elemwise_binary_broadcast_op.h @@ -216,7 +216,7 @@ void BinaryBroadcastIntCompute(const nnvm::NodeAttrs& attrs, if (outputs[0].type_flag_ == mshadow::kBool) { LOG(FATAL) << "Operator " << attrs.op->name << " does not support boolean type"; } - MXNET_INT_TYPE_SWITCH(outputs[0].type_flag_, DType, { + MXNET_INT_TYPE_SWITCH_EXT(outputs[0].type_flag_, DType, { BROADCAST_NDIM_SWITCH(ndim, NDim, { mshadow::Shape oshape = new_oshape.get(); mshadow::Shape lstride = mxnet_op::calc_stride(new_lshape.get()); diff --git a/src/operator/tensor/elemwise_binary_op.h b/src/operator/tensor/elemwise_binary_op.h index aa350b886286..2f32bbaa6fdf 100644 --- a/src/operator/tensor/elemwise_binary_op.h +++ b/src/operator/tensor/elemwise_binary_op.h @@ -449,7 +449,7 @@ class ElemwiseBinaryOp : public OpBase { CHECK_EQ(inputs.size(), 2U); CHECK_EQ(outputs.size(), 1U); MXNET_ASSIGN_REQ_SWITCH(req[0], Req, { - MXNET_INT_TYPE_SWITCH(outputs[0].type_flag_, DType, { + MXNET_INT_TYPE_SWITCH_EXT(outputs[0].type_flag_, DType, { const size_t size = (minthree(outputs[0].Size(), inputs[0].Size(), inputs[1].Size()) + DataType::kLanes - 1) / DataType::kLanes; From 49c25db2ae855a1ac594acb27ab77c879a1ae0fa Mon Sep 17 00:00:00 2001 From: barry-jin Date: Wed, 20 Oct 2021 11:16:22 -0700 Subject: [PATCH 14/18] add MXNET_INT_TYPE_SWITCH_EXT --- src/operator/mxnet_op.h | 78 +++++++++++++++++++++++++++++++++++++++++ 1 file changed, 78 insertions(+) diff --git a/src/operator/mxnet_op.h b/src/operator/mxnet_op.h index fd3b5877224e..b65861a2137a 100644 --- a/src/operator/mxnet_op.h +++ b/src/operator/mxnet_op.h @@ -419,6 +419,84 @@ struct AccType { LOG(FATAL) << "Unknown type enum " << type; \ } +#define MXNET_INT_TYPE_SWITCH_EXT(type, DType, ...) \ + switch (type) { \ + case mshadow::kFloat32: \ + { \ + LOG(FATAL) << "This operation only support " \ + "integer types, not float32"; \ + } \ + break; \ + case mshadow::kFloat64: \ + { \ + LOG(FATAL) << "This operation only support " \ + "integer types, not float64"; \ + } \ + break; \ + case mshadow::kFloat16: \ + { \ + LOG(FATAL) << "This operation only support " \ + "integer types, not float16"; \ + } \ + break; \ + case mshadow::kUint8: \ + { \ + typedef uint8_t DType; \ + {__VA_ARGS__} \ + } \ + break; \ + case mshadow::kInt8: \ + { \ + typedef int8_t DType; \ + {__VA_ARGS__} \ + } \ + break; \ + case mshadow::kInt32: \ + { \ + typedef int32_t DType; \ + {__VA_ARGS__} \ + } \ + break; \ + case mshadow::kInt64: \ + { \ + typedef int64_t DType; \ + {__VA_ARGS__} \ + } \ + break; \ + case mshadow::kInt16: \ + { \ + typedef int16_t DType; \ + {__VA_ARGS__} \ + } \ + break; \ + case mshadow::kUint16: \ + { \ + typedef uint16_t DType; \ + {__VA_ARGS__} \ + } \ + break; \ + case mshadow::kUint32: \ + { \ + typedef uint32_t DType; \ + {__VA_ARGS__} \ + } \ + break; \ + case mshadow::kUint64: \ + { \ + typedef uint64_t DType; \ + {__VA_ARGS__} \ + } \ + break; \ + case mshadow::kBool: \ + { \ + typedef bool DType; \ + {__VA_ARGS__} \ + } \ + break; \ + default: \ + LOG(FATAL) << "Unknown type enum " << type; \ + } + #define MXNET_INT32_INT64_TYPE_SWITCH(type, DType, ...) \ switch (type) { \ case mshadow::kFloat32: { \ From d3f6c07e055a170b4127d1cea09f408fcf02ed4c Mon Sep 17 00:00:00 2001 From: barry-jin Date: Sat, 30 Oct 2021 16:54:36 -0700 Subject: [PATCH 15/18] fix sanity check --- src/engine/threaded_engine.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/engine/threaded_engine.cc b/src/engine/threaded_engine.cc index 40d852b83b86..12a610c873b2 100644 --- a/src/engine/threaded_engine.cc +++ b/src/engine/threaded_engine.cc @@ -712,7 +712,7 @@ void ThreadedEngine::OnCompleteGPU(Engine* engine, void* sync_info, const dmlc:: ThreadedOpr* threaded_opr = static_cast(info->opr_block)->opr; auto* event_pool = static_cast(info->event_pool); - auto [event, event_pool_idx] = event_pool->GetNextEvent(); + auto[event, event_pool_idx] = event_pool->GetNextEvent(); auto ev = event.lock(); MSHADOW_CUDA_CALL(cudaEventRecord(*ev, worker_stream->stream_)); for (auto* read_var : threaded_opr->const_vars) { From 955f65f2f92b56674966fb1a619ca3024327293a Mon Sep 17 00:00:00 2001 From: barry-jin Date: Sat, 30 Oct 2021 17:23:10 -0700 Subject: [PATCH 16/18] fix lint --- src/engine/threaded_engine.cc | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/src/engine/threaded_engine.cc b/src/engine/threaded_engine.cc index 12a610c873b2..20d121039551 100644 --- a/src/engine/threaded_engine.cc +++ b/src/engine/threaded_engine.cc @@ -712,7 +712,8 @@ void ThreadedEngine::OnCompleteGPU(Engine* engine, void* sync_info, const dmlc:: ThreadedOpr* threaded_opr = static_cast(info->opr_block)->opr; auto* event_pool = static_cast(info->event_pool); - auto[event, event_pool_idx] = event_pool->GetNextEvent(); + // NOLINT_NEXT_LINE(whitespace/operators) + auto [event, event_pool_idx] = event_pool->GetNextEvent(); auto ev = event.lock(); MSHADOW_CUDA_CALL(cudaEventRecord(*ev, worker_stream->stream_)); for (auto* read_var : threaded_opr->const_vars) { From 2eeed8a8e9578890af3e7d026490c46f9d022cc6 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Sat, 30 Oct 2021 17:34:11 -0700 Subject: [PATCH 17/18] fix --- src/engine/threaded_engine.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/engine/threaded_engine.cc b/src/engine/threaded_engine.cc index 20d121039551..9e4c7dfc3d57 100644 --- a/src/engine/threaded_engine.cc +++ b/src/engine/threaded_engine.cc @@ -712,7 +712,7 @@ void ThreadedEngine::OnCompleteGPU(Engine* engine, void* sync_info, const dmlc:: ThreadedOpr* threaded_opr = static_cast(info->opr_block)->opr; auto* event_pool = static_cast(info->event_pool); - // NOLINT_NEXT_LINE(whitespace/operators) + // NOLINT_NEXT_LINE(whitespace/braces) auto [event, event_pool_idx] = event_pool->GetNextEvent(); auto ev = event.lock(); MSHADOW_CUDA_CALL(cudaEventRecord(*ev, worker_stream->stream_)); From c930a98917c6515719775cd2f6e341c7fafa9570 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Sat, 30 Oct 2021 17:53:42 -0700 Subject: [PATCH 18/18] fix lint --- src/engine/threaded_engine.cc | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/src/engine/threaded_engine.cc b/src/engine/threaded_engine.cc index 9e4c7dfc3d57..7639fd445987 100644 --- a/src/engine/threaded_engine.cc +++ b/src/engine/threaded_engine.cc @@ -712,8 +712,7 @@ void ThreadedEngine::OnCompleteGPU(Engine* engine, void* sync_info, const dmlc:: ThreadedOpr* threaded_opr = static_cast(info->opr_block)->opr; auto* event_pool = static_cast(info->event_pool); - // NOLINT_NEXT_LINE(whitespace/braces) - auto [event, event_pool_idx] = event_pool->GetNextEvent(); + auto [event, event_pool_idx] = event_pool->GetNextEvent(); // NOLINT(*) auto ev = event.lock(); MSHADOW_CUDA_CALL(cudaEventRecord(*ev, worker_stream->stream_)); for (auto* read_var : threaded_opr->const_vars) {