From ab24c2433ea838088657a0dfeab21070ceff28a7 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Wed, 3 Nov 2021 20:42:12 -0700 Subject: [PATCH 01/39] [API] Standardize and add more array api tests --- ci/docker/runtime_functions.sh | 20 + python/mxnet/ndarray/numpy/_op.py | 6 + python/mxnet/numpy/multiarray.py | 2 + src/common/utils.h | 115 +++-- src/operator/contrib/boolean_mask.cc | 2 +- src/operator/mshadow_op.h | 24 +- src/operator/numpy/np_elemwise_broadcast_op.h | 392 +++++++++++++++++- .../np_elemwise_broadcast_op_extended.cc | 88 +--- .../np_elemwise_broadcast_op_extended_thi.cc | 42 +- .../numpy/np_elemwise_broadcast_op_lae.cc | 24 +- src/operator/numpy/np_true_divide-inl.h | 70 +++- src/operator/numpy/np_true_divide.cc | 2 +- src/operator/tensor/elemwise_unary_op.h | 8 +- 13 files changed, 643 insertions(+), 152 deletions(-) diff --git a/ci/docker/runtime_functions.sh b/ci/docker/runtime_functions.sh index 0f7913557935..0e64775ff986 100755 --- a/ci/docker/runtime_functions.sh +++ b/ci/docker/runtime_functions.sh @@ -874,6 +874,26 @@ unittest_array_api_standardization() { export DMLC_LOG_STACK_TRACE_DEPTH=100 python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ array_api_tests/test_type_promotion.py::test_elementwise_function_two_arg_bool_type_promotion + python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ + array_api_tests/test_type_promotion.py::test_elementwise_function_two_arg_promoted_type_promotion + python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ + array_api_tests/test_type_promotion.py::test_elementwise_function_one_arg_bool + python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ + array_api_tests/test_type_promotion.py::test_elementwise_function_one_arg_type_promotion + python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ + array_api_tests/test_type_promotion.py::test_operator_one_arg_type_promotion + python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ + array_api_tests/test_type_promotion.py::test_operator_two_arg_bool_promotion + python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ + array_api_tests/test_type_promotion.py::test_operator_two_arg_promoted_promotion + python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ + array_api_tests/test_type_promotion.py::test_operator_inplace_two_arg_promoted_promotion + python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ + array_api_tests/test_elementwise_functions.py + python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ + array_api_tests/test_indexing.py + python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ + array_api_tests/test_constants.py popd } diff --git a/python/mxnet/ndarray/numpy/_op.py b/python/mxnet/ndarray/numpy/_op.py index 7378bd6e7a8a..ddd7b9edbe88 100644 --- a/python/mxnet/ndarray/numpy/_op.py +++ b/python/mxnet/ndarray/numpy/_op.py @@ -3731,6 +3731,8 @@ def ceil(x, out=None, **kwargs): >>> a array(4.) """ + if _np.issubdtype(x.dtype, _np.integer): + return x return _pure_unary_func_helper(x, _api_internal.ceil, _np.ceil, out=out, **kwargs) @@ -3770,6 +3772,8 @@ def floor(x, out=None, **kwargs): >>> a array(3.) """ + if _np.issubdtype(x.dtype, _np.integer): + return x return _pure_unary_func_helper(x, _api_internal.floor, _np.floor, out=out, **kwargs) @@ -3915,6 +3919,8 @@ def trunc(x, out=None, **kwargs): >>> np.trunc(a) array([-1., -1., -0., 0., 1., 1., 2.]) """ + if _np.issubdtype(x.dtype, _np.integer): + return x return _pure_unary_func_helper(x, _api_internal.trunc, _np.trunc, out=out, **kwargs) diff --git a/python/mxnet/numpy/multiarray.py b/python/mxnet/numpy/multiarray.py index 148b1298842c..404db69ddfb1 100644 --- a/python/mxnet/numpy/multiarray.py +++ b/python/mxnet/numpy/multiarray.py @@ -581,6 +581,8 @@ def _get_np_boolean_indexing(self, key, ndim, shape): remaining_dims = shape[key_ndim:] data = _reshape_view(self, -1, *remaining_dims) key = _reshape_view(key, -1) + if data.size == 0 and key.size == 0: + return data return _reshape_view(_npi.boolean_mask(data, key), -1, *remaining_dims) def _set_np_boolean_indexing(self, key, value): diff --git a/src/common/utils.h b/src/common/utils.h index 15e676c816c9..825b9fcec3a3 100644 --- a/src/common/utils.h +++ b/src/common/utils.h @@ -901,13 +901,44 @@ inline bool is_float(const int dtype) { } inline bool is_int(const int dtype) { - return dtype == mshadow::kUint8 || dtype == mshadow::kInt8 || dtype == mshadow::kInt32 || - dtype == mshadow::kInt64; + return dtype == mshadow::kUint8 || dtype == mshadow::kInt8 || + dtype == mshadow::kUint16 || dtype == mshadow::kInt16 || + dtype == mshadow::kUint32 || dtype == mshadow::kInt32 || + dtype == mshadow::kUint64 || dtype == mshadow::kInt64; } -inline int get_more_precise_type(const int type1, const int type2) { - if (type1 == type2) - return type1; +inline bool is_signed_int(const int dtype) { + return dtype == mshadow::kInt8 || dtype == mshadow::kInt16 || + dtype == mshadow::kInt32 || dtype == mshadow::kInt64; +} + +inline bool is_unsigned_int(const int dtype) { + return dtype == mshadow::kUint8 || dtype == mshadow::kUint16 || + dtype == mshadow::kUint32 || dtype == mshadow::kUint64; +} + +static int bits_of(const int type_flag) { + switch (type_flag) { + case mshadow::kFloat32: return sizeof(float) * CHAR_BIT; + case mshadow::kFloat64: return sizeof(double) * CHAR_BIT; + case mshadow::kUint8: return sizeof(uint8_t) * CHAR_BIT; + case mshadow::kInt32: return sizeof(int32_t) * CHAR_BIT; + case mshadow::kInt8: return sizeof(int8_t) * CHAR_BIT; + case mshadow::kInt64: return sizeof(int64_t) * CHAR_BIT; + case mshadow::kBool: return sizeof(bool) * CHAR_BIT; + case mshadow::kInt16: return sizeof(int16_t) * CHAR_BIT; + case mshadow::kUint16: return sizeof(uint16_t) * CHAR_BIT; + case mshadow::kUint32: return sizeof(uint32_t) * CHAR_BIT; + case mshadow::kUint64: return sizeof(uint64_t) * CHAR_BIT; + default: { + LOG(FATAL) << "Unknown type_flag=" << type_flag; + return -1; + } + } +} + +inline int type_promotion(const int type1, const int type2) { + if (type1 == type2) return type1; if (is_float(type1) && is_float(type2)) { if (type1 == mshadow::kFloat64 || type2 == mshadow::kFloat64) { return mshadow::kFloat64; @@ -919,27 +950,63 @@ inline int get_more_precise_type(const int type1, const int type2) { } else if (is_float(type1) || is_float(type2)) { return is_float(type1) ? type1 : type2; } - if (type1 == mshadow::kInt64 || type2 == mshadow::kInt64) { - return mshadow::kInt64; - } - if (type1 == mshadow::kInt32 || type2 == mshadow::kInt32) { - return mshadow::kInt32; - } - CHECK(!((type1 == mshadow::kUint8 && type2 == mshadow::kInt8) || - (type1 == mshadow::kInt8 && type2 == mshadow::kUint8))) - << "1 is UInt8 and 1 is Int8 should not get here"; - if (type1 == mshadow::kUint8 || type2 == mshadow::kUint8) { + if (is_signed_int(type1) && is_signed_int(type2)) { + if (type1 == mshadow::kInt64 || type2 == mshadow::kInt64) { + return mshadow::kInt64; + } + if (type1 == mshadow::kInt32 || type2 == mshadow::kInt32) { + return mshadow::kInt32; + } + if (type1 == mshadow::kInt16 || type2 == mshadow::kInt16) { + return mshadow::kInt16; + } + return mshadow::kInt8; + } else if (is_unsigned_int(type1) && is_unsigned_int(type2)) { + if (type1 == mshadow::kUint64 || type2 == mshadow::kUint64) { + return mshadow::kUint64; + } + if (type1 == mshadow::kUint32 || type2 == mshadow::kUint32) { + return mshadow::kUint32; + } + if (type1 == mshadow::kUint16 || type2 == mshadow::kUint16) { + return mshadow::kUint16; + } return mshadow::kUint8; + } else if (is_unsigned_int(type1) || is_unsigned_int(type2)) { + if (bits_of(type1) < bits_of(type2)) { + if (type1 == mshadow::kInt8 && type2 == mshadow::kUint16) { + return mshadow::kInt32; + } else if (type1 == mshadow::kInt8 && type2 == mshadow::kUint32) { + return mshadow::kInt64; + } else if (type1 == mshadow::kInt16 && type2 == mshadow::kUint32) { + return mshadow::kInt64; + } else { + return type2; + } + } else if (bits_of(type2) < bits_of(type1)) { + if (type2 == mshadow::kInt8 && type1 == mshadow::kUint16) { + return mshadow::kInt32; + } else if (type2 == mshadow::kInt8 && type1 == mshadow::kUint32) { + return mshadow::kInt64; + } else if (type2 == mshadow::kInt16 && type1 == mshadow::kUint32) { + return mshadow::kInt64; + } else { + return type1; + } + } else { + if (type1 == mshadow::kUint8 || type2 == mshadow::kUint8) { + return mshadow::kInt16; + } + if (type1 == mshadow::kUint16 || type2 == mshadow::kUint16) { + return mshadow::kInt32; + } + if (type1 == mshadow::kUint32 || type2 == mshadow::kUint32) { + return mshadow::kInt64; + } + } } - return mshadow::kInt8; -} - -inline int np_binary_out_infer_type(const int type1, const int type2) { - if ((type1 == mshadow::kUint8 && type2 == mshadow::kInt8) || - (type1 == mshadow::kInt8 && type2 == mshadow::kUint8)) { - return mshadow::kInt32; - } - return get_more_precise_type(type1, type2); + LOG(FATAL) << "should not reach here "; + return -1; } inline const std::string NodeAttrsGetProfilerScope(const nnvm::NodeAttrs& attrs) { diff --git a/src/operator/contrib/boolean_mask.cc b/src/operator/contrib/boolean_mask.cc index f3ba7f9f638f..b9307ea7d1dd 100644 --- a/src/operator/contrib/boolean_mask.cc +++ b/src/operator/contrib/boolean_mask.cc @@ -133,7 +133,7 @@ inline void BooleanMaskForward(const nnvm::NodeAttrs& attrs, const_cast(out).Init(s); // do the copy - MSHADOW_TYPE_SWITCH_WITH_BOOL(data.dtype(), DType, { + MSHADOW_TYPE_SWITCH_EXT_WITH_BOOL(data.dtype(), DType, { size_t input_size = data.shape().Size(); size_t col_size = input_size / idx_size; mshadow::Stream* stream = ctx.get_stream(); diff --git a/src/operator/mshadow_op.h b/src/operator/mshadow_op.h index 9a14794a47da..5743667001a0 100644 --- a/src/operator/mshadow_op.h +++ b/src/operator/mshadow_op.h @@ -819,7 +819,15 @@ MXNET_BINARY_MATH_OP(bitwise_or, static_cast(a) | static_cast( #endif /*! \brief used for generate element of bitwise_left_shift */ -MXNET_BINARY_MATH_OP(bitwise_left_shift, static_cast(a) << static_cast(b)); +struct bitwise_left_shift : public mxnet_op::tunable { + template + MSHADOW_XINLINE static DType Map(DType a, DType b) { + if (b >= DType(sizeof(DType) * CHAR_BIT)) { + return DType(0); + } + return static_cast(a) << static_cast(b); + } +}; MXNET_BINARY_MATH_OP(bitwise_left_shift_grad, math::pow(2.0f, static_cast(b))); @@ -834,7 +842,19 @@ MXNET_BINARY_MATH_OP(rbitwise_left_shift_grad, 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)); +struct bitwise_right_shift : public mxnet_op::tunable { + template + MSHADOW_XINLINE static DType Map(DType a, DType b) { + if (b >= DType(sizeof(DType) * CHAR_BIT)) { + if (a < 0) { + return DType(-1); + } else { + return DType(0); + } + } + return static_cast(a) >> static_cast(b); + } +}; MXNET_BINARY_MATH_OP(bitwise_right_shift_grad, math::pow(0.5f, static_cast(b))); diff --git a/src/operator/numpy/np_elemwise_broadcast_op.h b/src/operator/numpy/np_elemwise_broadcast_op.h index da40fe4044e7..569358d3f326 100644 --- a/src/operator/numpy/np_elemwise_broadcast_op.h +++ b/src/operator/numpy/np_elemwise_broadcast_op.h @@ -116,7 +116,7 @@ void MixedIntRealBinaryElemwiseCompute(const OpContext& ctx, if (size == 0) return; - MXNET_INT_TYPE_SWITCH(rhs.type_flag_, IType, { + MXNET_INT_TYPE_SWITCH_EXT(rhs.type_flag_, IType, { MXNET_ASSIGN_REQ_SWITCH(req, Req, { Kernel, xpu>::Launch( s, size, out.dptr(), rhs.dptr(), lhs.dptr()); @@ -125,7 +125,81 @@ void MixedIntRealBinaryElemwiseCompute(const OpContext& ctx, }); } -template +template +void MixedIntBinaryElemwiseCompute(const nnvm::NodeAttrs& attrs, + const OpContext& ctx, + const TBlob& lhs, + const TBlob& rhs, + const TBlob& out, + const OpReqType req) { + using namespace mshadow; + using namespace mxnet_op; + + Stream *s = ctx.get_stream(); + TBlob temp_tblob; + if (lhs.type_flag_ == out.type_flag_) { + MXNET_INT_TYPE_SWITCH_EXT(lhs.type_flag_, LType, { + Tensor temp_tensor = + ctx.requested[0].get_space_typed(Shape1(rhs.Size()), s); + temp_tblob = TBlob(temp_tensor); + }); + CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob}); + MXNET_ASSIGN_REQ_SWITCH(req, Req, { + MXNET_INT_TYPE_SWITCH_EXT(out.type_flag_, DType, { + const size_t size = (ElemwiseBinaryOp::minthree(out.Size(), lhs.Size(), temp_tblob.Size()) + + DataType::kLanes - 1) / DataType::kLanes; + if (size != 0) { + Kernel, xpu>::Launch(s, size, + out.dptr(), + lhs.dptr(), temp_tblob.dptr()); + } + }); + }); + } else if (rhs.type_flag_ == out.type_flag_) { + MXNET_INT_TYPE_SWITCH_EXT(rhs.type_flag_, RType, { + Tensor temp_tensor = + ctx.requested[0].get_space_typed(Shape1(lhs.Size()), s); + temp_tblob = TBlob(temp_tensor); + }); + CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob}); + MXNET_ASSIGN_REQ_SWITCH(req, Req, { + MXNET_INT_TYPE_SWITCH_EXT(out.type_flag_, DType, { + const size_t size = (ElemwiseBinaryOp::minthree(out.Size(), temp_tblob.Size(), rhs.Size()) + + DataType::kLanes - 1) / DataType::kLanes; + if (size != 0) { + Kernel, xpu>::Launch(s, size, + out.dptr(), + temp_tblob.dptr(), rhs.dptr()); + } + }); + }); + } else { + TBlob temp_tblob_l; + TBlob temp_tblob_r; + MXNET_INT_TYPE_SWITCH_EXT(out.type_flag_, OType, { + Tensor workspace = + ctx.requested[0].get_space_typed(Shape1(lhs.Size() + rhs.Size()), s); + TBlob temp_tblob = TBlob(workspace); + temp_tblob_l = TBlob(reinterpret_cast(temp_tblob.dptr_), lhs.shape_, temp_tblob.dev_mask(), temp_tblob.dev_id()); + temp_tblob_r = TBlob(reinterpret_cast(temp_tblob.dptr_) + lhs.Size() + 1, rhs.shape_, temp_tblob.dev_mask(), temp_tblob.dev_id()); + }); + CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob_l}); + CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob_r}); + MXNET_ASSIGN_REQ_SWITCH(req, Req, { + MXNET_INT_TYPE_SWITCH_EXT(out.type_flag_, DType, { + const size_t size = (ElemwiseBinaryOp::minthree(out.Size(), temp_tblob_l.Size(), temp_tblob_r.Size()) + + DataType::kLanes - 1) / DataType::kLanes; + if (size != 0) { + Kernel, xpu>::Launch(s, size, + out.dptr(), + temp_tblob_l.dptr(), temp_tblob_r.dptr()); + } + }); + }); + } +} + +template void MixedBinaryElemwiseCompute(const nnvm::NodeAttrs& attrs, const OpContext& ctx, const std::vector& inputs, @@ -152,7 +226,7 @@ void MixedBinaryElemwiseCompute(const nnvm::NodeAttrs& attrs, MixedIntRealBinaryElemwiseCompute(ctx, rhs, lhs, out, req[0]); } } else { - PrintErrorMessage(attrs.op->name, lhs.type_flag_, rhs.type_flag_); + MixedIntBinaryElemwiseCompute(attrs, ctx, lhs, rhs, out, req[0]); } } @@ -250,7 +324,7 @@ void MixedBinaryBroadcastCompute(const nnvm::NodeAttrs& attrs, int ndim = BinaryBroadcastShapeCompact( lhs.shape_, rhs.shape_, out.shape_, &new_lshape, &new_rshape, &new_oshape); if (!ndim) { - MixedBinaryElemwiseCompute(attrs, ctx, inputs, req, outputs); + MixedBinaryElemwiseCompute(attrs, ctx, inputs, req, outputs); } else { mshadow::Stream* s = ctx.get_stream(); if (common::is_float(lhs.type_flag_) && common::is_float(rhs.type_flag_)) { @@ -270,7 +344,7 @@ void MixedBinaryBroadcastCompute(const nnvm::NodeAttrs& attrs, mshadow::Shape rstride = mxnet_op::calc_stride(new_rshape.get()); if (lhs.type_flag_ == out.type_flag_) { MSHADOW_REAL_TYPE_SWITCH(out.type_flag_, LType, { - MXNET_INT_TYPE_SWITCH(rhs.type_flag_, RType, { + MXNET_INT_TYPE_SWITCH_EXT(rhs.type_flag_, RType, { mxnet_op::Kernel, xpu>::template LaunchEx(s, new_oshape.Size(), @@ -285,7 +359,7 @@ void MixedBinaryBroadcastCompute(const nnvm::NodeAttrs& attrs, }); } else { MSHADOW_REAL_TYPE_SWITCH(out.type_flag_, RType, { - MXNET_INT_TYPE_SWITCH(lhs.type_flag_, LType, { + MXNET_INT_TYPE_SWITCH_EXT(lhs.type_flag_, LType, { mxnet_op::Kernel, xpu>::template LaunchEx(s, new_oshape.Size(), @@ -303,7 +377,7 @@ void MixedBinaryBroadcastCompute(const nnvm::NodeAttrs& attrs, } else if (!common::is_float(lhs.type_flag_) && !common::is_float(rhs.type_flag_)) { TBlob temp_tblob; if (lhs.type_flag_ == out.type_flag_) { - MXNET_INT_TYPE_SWITCH(lhs.type_flag_, LType, { + MXNET_INT_TYPE_SWITCH_EXT(lhs.type_flag_, LType, { Tensor temp_tensor = ctx.requested[0].get_space_typed(Shape1(rhs.Size()), s); temp_tblob = TBlob(temp_tensor); @@ -311,7 +385,7 @@ void MixedBinaryBroadcastCompute(const nnvm::NodeAttrs& attrs, CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob}); BinaryBroadcastCompute( attrs, ctx, {lhs, temp_tblob.reshape(rhs.shape_)}, req, outputs); - } else { + } else if (rhs.type_flag_ == out.type_flag_) { MXNET_INT_TYPE_SWITCH(rhs.type_flag_, RType, { Tensor temp_tensor = ctx.requested[0].get_space_typed(Shape1(lhs.Size()), s); @@ -320,6 +394,20 @@ void MixedBinaryBroadcastCompute(const nnvm::NodeAttrs& attrs, CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob}); BinaryBroadcastCompute( attrs, ctx, {temp_tblob.reshape(lhs.shape_), rhs}, req, outputs); + } else { + TBlob temp_tblob_l; + TBlob temp_tblob_r; + MXNET_INT_TYPE_SWITCH_EXT(out.type_flag_, OType, { + Tensor workspace = + ctx.requested[0].get_space_typed(Shape1(lhs.Size() + rhs.Size()), s); + TBlob temp_tblob = TBlob(workspace); + temp_tblob_l = TBlob(reinterpret_cast(temp_tblob.dptr_), lhs.shape_, temp_tblob.dev_mask(), temp_tblob.dev_id()); + temp_tblob_r = TBlob(reinterpret_cast(temp_tblob.dptr_) + lhs.Size() + 1, rhs.shape_, temp_tblob.dev_mask(), temp_tblob.dev_id()); + }); + CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob_l}); + CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob_r}); + BinaryBroadcastCompute( + attrs, ctx, {temp_tblob_l, temp_tblob_r}, req, outputs); } } else { PrintErrorMessage(attrs.op->name, lhs.type_flag_, rhs.type_flag_); @@ -379,7 +467,7 @@ void NumpyBinaryBroadcastComputeWithBool(const nnvm::NodeAttrs& attrs, Stream* s = ctx.get_stream(); TBlob temp_tblob; if (lhs.type_flag_ == out.type_flag_) { - MXNET_INT_TYPE_SWITCH(lhs.type_flag_, LType, { + MXNET_INT_TYPE_SWITCH_EXT(lhs.type_flag_, LType, { Tensor temp_tensor = ctx.requested[0].get_space_typed(Shape1(rhs.Size()), s); temp_tblob = TBlob(temp_tensor); @@ -387,8 +475,8 @@ void NumpyBinaryBroadcastComputeWithBool(const nnvm::NodeAttrs& attrs, CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob}); BinaryBroadcastCompute( attrs, ctx, {lhs, temp_tblob.reshape(rhs.shape_)}, req, outputs); - } else { - MXNET_INT_TYPE_SWITCH(rhs.type_flag_, RType, { + } else if (rhs.type_flag_ == out.type_flag_) { + MXNET_INT_TYPE_SWITCH_EXT(rhs.type_flag_, RType, { Tensor temp_tensor = ctx.requested[0].get_space_typed(Shape1(lhs.Size()), s); temp_tblob = TBlob(temp_tensor); @@ -396,12 +484,213 @@ void NumpyBinaryBroadcastComputeWithBool(const nnvm::NodeAttrs& attrs, CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob}); BinaryBroadcastCompute( attrs, ctx, {temp_tblob.reshape(lhs.shape_), rhs}, req, outputs); + } else { + TBlob temp_tblob_l; + TBlob temp_tblob_r; + MXNET_INT_TYPE_SWITCH_EXT(out.type_flag_, OType, { + Tensor workspace = + ctx.requested[0].get_space_typed(Shape1(lhs.Size() + rhs.Size()), s); + TBlob temp_tblob = TBlob(workspace); + temp_tblob_l = TBlob(reinterpret_cast(temp_tblob.dptr_), lhs.shape_, temp_tblob.dev_mask(), temp_tblob.dev_id()); + temp_tblob_r = TBlob(reinterpret_cast(temp_tblob.dptr_) + lhs.Size() + 1, rhs.shape_, temp_tblob.dev_mask(), temp_tblob.dev_id()); + }); + CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob_l}); + CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob_r}); + BinaryBroadcastCompute( + attrs, ctx, {temp_tblob_l, temp_tblob_r}, req, outputs); } return; } MixedBinaryBroadcastCompute(attrs, ctx, inputs, req, outputs); } + +template +void NumpyBinaryBroadcastIntComputeWithBool(const nnvm::NodeAttrs& attrs, + const OpContext& ctx, + const std::vector& inputs, + const std::vector& req, + const std::vector& outputs) { + using namespace mshadow; + using namespace mxnet_op; + CHECK_EQ(inputs.size(), 2U); + CHECK_EQ(outputs.size(), 1U); + + const TBlob& lhs = inputs[0]; + const TBlob& rhs = inputs[1]; + const TBlob& out = outputs[0]; + + if ((out.shape_.Size() == 0U) || (req[0] == kNullOp)) return; + + if (lhs.type_flag_ == rhs.type_flag_) { + BinaryBroadcastIntCompute(attrs, ctx, inputs, req, outputs); + return; + } + Stream *s = ctx.get_stream(); + TBlob temp_tblob; + if (lhs.type_flag_ == out.type_flag_) { + MXNET_INT_TYPE_SWITCH_EXT(lhs.type_flag_, LType, { + Tensor temp_tensor = + ctx.requested[0].get_space_typed(Shape1(rhs.Size()), s); + temp_tblob = TBlob(temp_tensor); + }); + CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob}); + BinaryBroadcastIntCompute( + attrs, ctx, {lhs, temp_tblob.reshape(rhs.shape_)}, req, outputs); + } else if (rhs.type_flag_ == out.type_flag_) { + MXNET_INT_TYPE_SWITCH_EXT(rhs.type_flag_, RType, { + Tensor temp_tensor = + ctx.requested[0].get_space_typed(Shape1(lhs.Size()), s); + temp_tblob = TBlob(temp_tensor); + }); + CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob}); + BinaryBroadcastIntCompute( + attrs, ctx, {temp_tblob.reshape(lhs.shape_), rhs}, req, outputs); + } else { + TBlob temp_tblob_l; + TBlob temp_tblob_r; + MXNET_INT_TYPE_SWITCH_EXT(out.type_flag_, OType, { + Tensor workspace = + ctx.requested[0].get_space_typed(Shape1(lhs.Size() + rhs.Size()), s); + TBlob temp_tblob = TBlob(workspace); + temp_tblob_l = TBlob(reinterpret_cast(temp_tblob.dptr_), lhs.shape_, temp_tblob.dev_mask(), temp_tblob.dev_id()); + temp_tblob_r = TBlob(reinterpret_cast(temp_tblob.dptr_) + lhs.Size() + 1, rhs.shape_, temp_tblob.dev_mask(), temp_tblob.dev_id()); + }); + CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob_l}); + CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob_r}); + BinaryBroadcastIntCompute( + attrs, ctx, {temp_tblob_l, temp_tblob_r}, req, outputs); + } + return; +} + + +template +void NumpyBinaryBroadcastIntCompute(const nnvm::NodeAttrs& attrs, + const OpContext& ctx, + const std::vector& inputs, + const std::vector& req, + const std::vector& outputs) { + using namespace mshadow; + using namespace mxnet_op; + CHECK_EQ(inputs.size(), 2U); + CHECK_EQ(outputs.size(), 1U); + + const TBlob& lhs = inputs[0]; + const TBlob& rhs = inputs[1]; + const TBlob& out = outputs[0]; + + if ((out.shape_.Size() == 0U) || (req[0] == kNullOp)) return; + + if (lhs.type_flag_ == rhs.type_flag_) { + BinaryBroadcastIntCompute(attrs, ctx, inputs, req, outputs); + return; + } + Stream *s = ctx.get_stream(); + TBlob temp_tblob; + if (lhs.type_flag_ == out.type_flag_) { + MXNET_INT_TYPE_SWITCH_EXT(lhs.type_flag_, LType, { + Tensor temp_tensor = + ctx.requested[0].get_space_typed(Shape1(rhs.Size()), s); + temp_tblob = TBlob(temp_tensor); + }); + CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob}); + BinaryBroadcastIntCompute( + attrs, ctx, {lhs, temp_tblob.reshape(rhs.shape_)}, req, outputs); + } else if (rhs.type_flag_ == out.type_flag_) { + MXNET_INT_TYPE_SWITCH_EXT(rhs.type_flag_, RType, { + Tensor temp_tensor = + ctx.requested[0].get_space_typed(Shape1(lhs.Size()), s); + temp_tblob = TBlob(temp_tensor); + }); + CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob}); + BinaryBroadcastIntCompute( + attrs, ctx, {temp_tblob.reshape(lhs.shape_), rhs}, req, outputs); + } else { + TBlob temp_tblob_l; + TBlob temp_tblob_r; + MXNET_INT_TYPE_SWITCH_EXT(out.type_flag_, OType, { + Tensor workspace = + ctx.requested[0].get_space_typed(Shape1(lhs.Size() + rhs.Size()), s); + TBlob temp_tblob = TBlob(workspace); + temp_tblob_l = TBlob(reinterpret_cast(temp_tblob.dptr_), lhs.shape_, temp_tblob.dev_mask(), temp_tblob.dev_id()); + temp_tblob_r = TBlob(reinterpret_cast(temp_tblob.dptr_) + lhs.Size() + 1, rhs.shape_, temp_tblob.dev_mask(), temp_tblob.dev_id()); + }); + CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob_l}); + CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob_r}); + BinaryBroadcastIntCompute( + attrs, ctx, {temp_tblob_l, temp_tblob_r}, req, outputs); + } + return; +} + +inline bool NumpyBinaryMixedFloatingType(const nnvm::NodeAttrs& attrs, + std::vector* in_attrs, + std::vector* out_attrs) { + CHECK_EQ(in_attrs->size(), 2U); + CHECK_EQ(out_attrs->size(), 1U); + const int ltype = in_attrs->at(0); + const int rtype = in_attrs->at(1); + + if (ltype != -1 && rtype != -1 && (ltype != rtype)) { + // Only when both input types are known and not the same, we enter the mixed-precision mode + TYPE_ASSIGN_CHECK(*out_attrs, 0, common::type_promotion(ltype, rtype)); + } else { + TYPE_ASSIGN_CHECK(*out_attrs, 0, in_attrs->at(0)); + TYPE_ASSIGN_CHECK(*out_attrs, 0, in_attrs->at(1)); + TYPE_ASSIGN_CHECK(*in_attrs, 0, out_attrs->at(0)); + TYPE_ASSIGN_CHECK(*in_attrs, 1, out_attrs->at(0)); + } + // check if it is float16, float32 or float64. If not, raise error. + CHECK(common::is_float(in_attrs->at(0))) << "Do not support `int` as input.\n"; + return out_attrs->at(0) != -1; +} + +template +void NumpyBinaryMixedFloatingCompute(const nnvm::NodeAttrs& attrs, + const OpContext& ctx, + const std::vector& inputs, + const std::vector& req, + const std::vector& outputs) { + using namespace mshadow; + using namespace mxnet_op; + CHECK_EQ(inputs.size(), 2U); + CHECK_EQ(outputs.size(), 1U); + + const TBlob& lhs = inputs[0]; + const TBlob& rhs = inputs[1]; + const TBlob& out = outputs[0]; + + if ((out.shape_.Size() == 0U) || (req[0] == kNullOp)) return; + + if (lhs.type_flag_ == rhs.type_flag_) { + BinaryBroadcastCompute(attrs, ctx, inputs, req, outputs); + return; + } + Stream *s = ctx.get_stream(); + TBlob temp_tblob; + if (lhs.type_flag_ == out.type_flag_) { + MSHADOW_REAL_TYPE_SWITCH(lhs.type_flag_, LType, { + Tensor temp_tensor = + ctx.requested[0].get_space_typed(Shape1(rhs.Size()), s); + temp_tblob = TBlob(temp_tensor); + }); + CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob}); + BinaryBroadcastCompute( + attrs, ctx, {lhs, temp_tblob.reshape(rhs.shape_)}, req, outputs); + } else { + MSHADOW_REAL_TYPE_SWITCH(rhs.type_flag_, RType, { + Tensor temp_tensor = + ctx.requested[0].get_space_typed(Shape1(lhs.Size()), s); + temp_tblob = TBlob(temp_tensor); + }); + CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob}); + BinaryBroadcastCompute( + attrs, ctx, {temp_tblob.reshape(lhs.shape_), rhs}, req, outputs); + } + return; +} + template void NumpyBinaryBackwardUseIn(const nnvm::NodeAttrs& attrs, const OpContext& ctx, @@ -557,7 +846,7 @@ inline bool NumpyBinaryMixedPrecisionType(const nnvm::NodeAttrs& attrs, const int rtype = in_attrs->at(1); if (ltype != -1 && rtype != -1 && (ltype != rtype)) { // Only when both input types are known and not the same, we enter the mixed-precision mode - TYPE_ASSIGN_CHECK(*out_attrs, 0, common::np_binary_out_infer_type(ltype, rtype)); + TYPE_ASSIGN_CHECK(*out_attrs, 0, common::type_promotion(ltype, rtype)); } else { return ElemwiseType<2, 1>(attrs, in_attrs, out_attrs); } @@ -586,6 +875,85 @@ inline bool NumpyBinaryMixedPrecisionType(const nnvm::NodeAttrs& attrs, .add_argument("lhs", "NDArray-or-Symbol", "First input to the function") \ .add_argument("rhs", "NDArray-or-Symbol", "Second input to the function") +inline bool NumpyBinaryMixedIntPrecisionTypeWithBool(const nnvm::NodeAttrs& attrs, + std::vector* in_attrs, + std::vector* out_attrs) { + CHECK_EQ(in_attrs->size(), 2U); + CHECK_EQ(out_attrs->size(), 1U); + const int ltype = in_attrs->at(0); + const int rtype = in_attrs->at(1); + CHECK(common::is_int(ltype) || ltype == mshadow::kBool) << "1st input only supports integer types or bool types."; + CHECK(common::is_int(rtype) || rtype == mshadow::kBool) << "2nd input only supports integer types or bool types."; + if (ltype != -1 && rtype != -1 && (ltype != rtype)) { + // Only when both input types are known and not the same, we enter the mixed-precision mode + TYPE_ASSIGN_CHECK(*out_attrs, 0, common::type_promotion(ltype, rtype)); + } else { + return ElemwiseType<2, 1>(attrs, in_attrs, out_attrs); + } + return true; +} + +#define MXNET_OPERATOR_REGISTER_NP_BINARY_MIXED_INT_PRECISION_WITH_BOOL(name) \ + NNVM_REGISTER_OP(name) \ + .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", NumpyBinaryMixedIntPrecisionTypeWithBool) \ + .set_attr("FInplaceOption", \ + [](const NodeAttrs& attrs){ \ + return std::vector >{{0, 0}, {1, 0}}; \ + }) \ + .set_attr("FResourceRequest", \ + [](const NodeAttrs& attrs) { \ + return std::vector{ResourceRequest::kTempSpace}; \ + }) \ + .add_argument("lhs", "NDArray-or-Symbol", "First input to the function") \ + .add_argument("rhs", "NDArray-or-Symbol", "Second input to the function") + + +inline bool NumpyBinaryMixedIntPrecisionType(const nnvm::NodeAttrs& attrs, + std::vector* in_attrs, + std::vector* out_attrs) { + CHECK_EQ(in_attrs->size(), 2U); + CHECK_EQ(out_attrs->size(), 1U); + const int ltype = in_attrs->at(0); + const int rtype = in_attrs->at(1); + CHECK(common::is_int(ltype)) << "1st input only supports integer types."; + CHECK(common::is_int(rtype)) << "2nd input only supports integer types."; + if (ltype != -1 && rtype != -1 && (ltype != rtype)) { + // Only when both input types are known and not the same, we enter the mixed-precision mode + TYPE_ASSIGN_CHECK(*out_attrs, 0, common::type_promotion(ltype, rtype)); + } else { + return ElemwiseType<2, 1>(attrs, in_attrs, out_attrs); + } + return true; +} + +#define MXNET_OPERATOR_REGISTER_NP_BINARY_MIXED_INT_PRECISION(name) \ + NNVM_REGISTER_OP(name) \ + .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", NumpyBinaryMixedIntPrecisionType) \ + .set_attr("FInplaceOption", \ + [](const NodeAttrs& attrs){ \ + return std::vector >{{0, 0}, {1, 0}}; \ + }) \ + .set_attr("FResourceRequest", \ + [](const NodeAttrs& attrs) { \ + return std::vector{ResourceRequest::kTempSpace}; \ + }) \ + .add_argument("lhs", "NDArray-or-Symbol", "First input to the function") \ + .add_argument("rhs", "NDArray-or-Symbol", "Second input to the function") + } // namespace op } // namespace mxnet #endif // MXNET_OPERATOR_NUMPY_NP_ELEMWISE_BROADCAST_OP_H_ diff --git a/src/operator/numpy/np_elemwise_broadcast_op_extended.cc b/src/operator/numpy/np_elemwise_broadcast_op_extended.cc index 98a4688002ce..fff0f7c66184 100644 --- a/src/operator/numpy/np_elemwise_broadcast_op_extended.cc +++ b/src/operator/numpy/np_elemwise_broadcast_op_extended.cc @@ -130,23 +130,11 @@ NNVM_REGISTER_OP(_npi_lcm_scalar) .add_arguments(NumpyBinaryScalarParam::__FIELDS__()) .set_attr("FCompute", BinaryScalarOp::ComputeInt); -NNVM_REGISTER_OP(_npi_bitwise_and) - .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("FGradient", MakeZeroGradNodes) - .set_attr("FCompute", BinaryBroadcastIntCompute) - .add_argument("lhs", "NDArray-or-Symbol", "First input to the function") - .add_argument("rhs", "NDArray-or-Symbol", "Second input to the function"); +MXNET_OPERATOR_REGISTER_NP_BINARY_MIXED_INT_PRECISION_WITH_BOOL(_npi_bitwise_and) +.set_attr( + "FCompute", + NumpyBinaryBroadcastIntComputeWithBool) +.set_attr("FGradient", MakeZeroGradNodes); NNVM_REGISTER_OP(_npi_bitwise_and_scalar) .set_num_inputs(1) @@ -163,41 +151,17 @@ NNVM_REGISTER_OP(_npi_bitwise_and_scalar) .add_arguments(NumpyBinaryScalarParam::__FIELDS__()) .set_attr("FCompute", BinaryScalarOp::ComputeInt); -NNVM_REGISTER_OP(_npi_bitwise_xor) - .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("FGradient", MakeZeroGradNodes) - .set_attr("FCompute", BinaryBroadcastIntCompute) - .add_argument("lhs", "NDArray-or-Symbol", "First input to the function") - .add_argument("rhs", "NDArray-or-Symbol", "Second input to the function"); +MXNET_OPERATOR_REGISTER_NP_BINARY_MIXED_INT_PRECISION_WITH_BOOL(_npi_bitwise_xor) +.set_attr( + "FCompute", + NumpyBinaryBroadcastIntComputeWithBool) +.set_attr("FGradient", MakeZeroGradNodes); -NNVM_REGISTER_OP(_npi_bitwise_or) - .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("FGradient", MakeZeroGradNodes) - .set_attr("FCompute", BinaryBroadcastIntCompute) - .add_argument("lhs", "NDArray-or-Symbol", "First input to the function") - .add_argument("rhs", "NDArray-or-Symbol", "Second input to the function"); +MXNET_OPERATOR_REGISTER_NP_BINARY_MIXED_INT_PRECISION_WITH_BOOL(_npi_bitwise_or) +.set_attr( + "FCompute", + NumpyBinaryBroadcastIntComputeWithBool) +.set_attr("FGradient", MakeZeroGradNodes); NNVM_REGISTER_OP(_npi_bitwise_xor_scalar) .set_num_inputs(1) @@ -240,20 +204,6 @@ MXNET_OPERATOR_REGISTER_NP_BINARY_SCALAR(_npi_rcopysign_scalar) MXNET_OPERATOR_REGISTER_NP_BINARY_SCALAR(_backward_npi_copysign_scalar) .set_attr("FCompute", BinaryScalarOp::Backward); -inline bool Arctan2OpType(const nnvm::NodeAttrs& attrs, - std::vector* in_attrs, - std::vector* out_attrs) { - CHECK_EQ(in_attrs->size(), 2U); - CHECK_EQ(out_attrs->size(), 1U); - - TYPE_ASSIGN_CHECK(*out_attrs, 0, in_attrs->at(0)); - TYPE_ASSIGN_CHECK(*out_attrs, 0, in_attrs->at(1)); - TYPE_ASSIGN_CHECK(*in_attrs, 0, out_attrs->at(0)); - TYPE_ASSIGN_CHECK(*in_attrs, 1, out_attrs->at(0)); - // check if it is float16, float32 or float64. If not, raise error. - CHECK(common::is_float(in_attrs->at(0))) << "Do not support `int` as input.\n"; - return out_attrs->at(0) != -1; -} NNVM_REGISTER_OP(_npi_arctan2) .set_num_inputs(2) @@ -263,13 +213,17 @@ NNVM_REGISTER_OP(_npi_arctan2) return std::vector{"x1", "x2"}; }) .set_attr("FInferShape", BinaryBroadcastShape) - .set_attr("FInferType", Arctan2OpType) - .set_attr("FCompute", BinaryBroadcastCompute) + .set_attr("FInferType", NumpyBinaryMixedFloatingType) + .set_attr("FCompute", NumpyBinaryMixedFloatingCompute) .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_arctan2"}) .set_attr("FInplaceOption", [](const NodeAttrs& attrs) { return std::vector >{{0, 0}}; }) + .set_attr("FResourceRequest", + [](const NodeAttrs& attrs) { + return std::vector{ResourceRequest::kTempSpace}; + }) .add_argument("x1", "NDArray-or-Symbol", "The input array") .add_argument("x2", "NDArray-or-Symbol", "The input array"); 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 90ecd6e2387a..7f13b79f0e2d 100644 --- a/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cc +++ b/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cc @@ -44,24 +44,9 @@ namespace op { .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", - 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"); +MXNET_OPERATOR_REGISTER_NP_BINARY_MIXED_INT_PRECISION(_npi_bitwise_left_shift) + .set_attr("FCompute", NumpyBinaryBroadcastIntCompute) + .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_bitwise_left_shift"}); NNVM_REGISTER_OP(_npi_bitwise_left_shift_scalar) .set_num_inputs(1) @@ -126,24 +111,9 @@ 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) - .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", - 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"); +MXNET_OPERATOR_REGISTER_NP_BINARY_MIXED_INT_PRECISION(_npi_bitwise_right_shift) + .set_attr("FCompute", NumpyBinaryBroadcastIntCompute) + .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_bitwise_right_shift"}); NNVM_REGISTER_OP(_npi_bitwise_right_shift_scalar) .set_num_inputs(1) diff --git a/src/operator/numpy/np_elemwise_broadcast_op_lae.cc b/src/operator/numpy/np_elemwise_broadcast_op_lae.cc index 05d83d819dc9..e53df7e463dd 100644 --- a/src/operator/numpy/np_elemwise_broadcast_op_lae.cc +++ b/src/operator/numpy/np_elemwise_broadcast_op_lae.cc @@ -27,9 +27,27 @@ namespace mxnet { namespace op { -MXNET_OPERATOR_REGISTER_BINARY_BROADCAST(_npi_logaddexp) - .set_attr("FCompute", BinaryBroadcastCompute) - .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_logaddexp"}); +NNVM_REGISTER_OP(_npi_logaddexp) +.set_num_inputs(2) +.set_num_outputs(1) +.set_attr("FListInputNames", + [](const NodeAttrs& attrs) { + return std::vector{"x1", "x2"}; + }) +.set_attr("FInferShape", BinaryBroadcastShape) +.set_attr("FInferType", NumpyBinaryMixedFloatingType) +.set_attr("FCompute", NumpyBinaryMixedFloatingCompute) +.set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_logaddexp"}) +.set_attr("FResourceRequest", + [](const NodeAttrs& attrs) { + return std::vector{ResourceRequest::kTempSpace}; + }) +.set_attr("FInplaceOption", + [](const NodeAttrs& attrs) { + return std::vector >{{0, 0}}; + }) +.add_argument("x1", "NDArray-or-Symbol", "The input array") +.add_argument("x2", "NDArray-or-Symbol", "The input array"); MXNET_OPERATOR_REGISTER_NP_BINARY_SCALAR(_npi_logaddexp_scalar) .set_attr("FCompute", BinaryScalarOp::Compute) diff --git a/src/operator/numpy/np_true_divide-inl.h b/src/operator/numpy/np_true_divide-inl.h index 6424e22ad209..acaf4cd25434 100644 --- a/src/operator/numpy/np_true_divide-inl.h +++ b/src/operator/numpy/np_true_divide-inl.h @@ -117,7 +117,34 @@ void TrueDivideElemwiseCompute(const nnvm::NodeAttrs& attrs, // Case when types of the 2 input tensors are different if (common::is_float(lhs.type_flag_) && common::is_float(rhs.type_flag_)) { // both lhs and rhs are float types, output type is the more precise one - LOG(FATAL) << "not implemented yet..."; + TBlob temp_tblob; + if (lhs.type_flag_ == out.type_flag_) { + MSHADOW_REAL_TYPE_SWITCH(lhs.type_flag_, LType, { + Tensor temp_tensor = + ctx.requested[0].get_space_typed(Shape1(rhs.Size()), s); + temp_tblob = TBlob(temp_tensor); + }); + CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob}); + MXNET_ASSIGN_REQ_SWITCH(req[0], Req, { + MSHADOW_REAL_TYPE_SWITCH(out.type_flag_, DType, { + Kernel, xpu>::Launch( + s, out.Size(), out.dptr(), lhs.dptr(), temp_tblob.dptr()); + }); + }); + } else { + MSHADOW_REAL_TYPE_SWITCH(rhs.type_flag_, RType, { + Tensor temp_tensor = + ctx.requested[0].get_space_typed(Shape1(lhs.Size()), s); + temp_tblob = TBlob(temp_tensor); + }); + CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob}); + MXNET_ASSIGN_REQ_SWITCH(req[0], Req, { + MSHADOW_REAL_TYPE_SWITCH(out.type_flag_, DType, { + Kernel, xpu>::Launch( + s, out.Size(), out.dptr(), temp_tblob.dptr(), rhs.dptr()); + }); + }); + } } else if (common::is_float(lhs.type_flag_) || common::is_float(rhs.type_flag_)) { // one is float type, the other is integer type, the output type should be the same as float CHECK_EQ(out.type_flag_, common::is_float(lhs.type_flag_) ? lhs.type_flag_ : rhs.type_flag_) @@ -213,7 +240,46 @@ void TrueDivideBroadcastCompute(const nnvm::NodeAttrs& attrs, } else { if (common::is_float(lhs.type_flag_) && common::is_float(rhs.type_flag_)) { // lhs and rhs have different float types, the output is the more precise one - LOG(FATAL) << "not implemented yet..."; + TBlob temp_tblob; + if (lhs.type_flag_ == out.type_flag_) { + MSHADOW_REAL_TYPE_SWITCH(lhs.type_flag_, LType, { + Tensor temp_tensor = + ctx.requested[0].get_space_typed(Shape1(rhs.Size()), s); + temp_tblob = TBlob(temp_tensor); + }); + CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob}); + MSHADOW_REAL_TYPE_SWITCH(out.type_flag_, DType, { + Kernel, xpu>::template LaunchEx( + s, + new_oshape.Size(), + req[0], + lstride, + rstride, + oshape, + lhs.dptr(), + temp_tblob.dptr(), + out.dptr()); + }); + } else { + MSHADOW_REAL_TYPE_SWITCH(rhs.type_flag_, RType, { + Tensor temp_tensor = + ctx.requested[0].get_space_typed(Shape1(lhs.Size()), s); + temp_tblob = TBlob(temp_tensor); + }); + CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob}); + MSHADOW_REAL_TYPE_SWITCH(out.type_flag_, DType, { + Kernel, xpu>::template LaunchEx( + s, + new_oshape.Size(), + req[0], + lstride, + rstride, + oshape, + temp_tblob.dptr(), + rhs.dptr(), + out.dptr()); + }); + } } else if (common::is_float(lhs.type_flag_) || common::is_float(rhs.type_flag_)) { // one of lhs and rhs is float, the output is the same type as the float one if (common::is_float(lhs.type_flag_)) { diff --git a/src/operator/numpy/np_true_divide.cc b/src/operator/numpy/np_true_divide.cc index 13fb72ca970a..4b0b98131fa1 100644 --- a/src/operator/numpy/np_true_divide.cc +++ b/src/operator/numpy/np_true_divide.cc @@ -30,7 +30,7 @@ namespace op { int TrueDivideOutType(int ltype, int rtype) { if (common::is_float(ltype) && common::is_float(rtype)) { // If both inputs are float, return the one with the higher precision - return common::get_more_precise_type(ltype, rtype); + return common::type_promotion(ltype, rtype); } else if (common::is_float(ltype) || common::is_float(rtype)) { // If only one of the inputs is float, return that float type return (common::is_float(ltype)) ? ltype : rtype; diff --git a/src/operator/tensor/elemwise_unary_op.h b/src/operator/tensor/elemwise_unary_op.h index 38949f1769ed..46855cb5d937 100644 --- a/src/operator/tensor/elemwise_unary_op.h +++ b/src/operator/tensor/elemwise_unary_op.h @@ -275,7 +275,7 @@ class UnaryOp : public OpBase { UnaryOp::Compute(attrs, ctx, inputs, req, outputs); } else { MSHADOW_REAL_TYPE_SWITCH(outputs[0].type_flag_, DType, { - MXNET_INT_TYPE_SWITCH(inputs[0].type_flag_, IType, { + MXNET_INT_TYPE_SWITCH_EXT(inputs[0].type_flag_, IType, { MXNET_ASSIGN_REQ_SWITCH(req[0], Req, { if (inputs[0].Size() != 0) { mxnet_op::Kernel, xpu>::Launch( @@ -294,7 +294,7 @@ class UnaryOp : public OpBase { const std::vector& req, const std::vector& outputs) { mshadow::Stream* s = ctx.get_stream(); - MXNET_INT_TYPE_SWITCH(outputs[0].type_flag_, DType, { + MXNET_INT_TYPE_SWITCH_EXT(outputs[0].type_flag_, DType, { MXNET_ASSIGN_REQ_SWITCH(req[0], Req, { if (inputs[0].Size() != 0) { mxnet_op::Kernel, xpu>::Launch( @@ -311,7 +311,7 @@ class UnaryOp : public OpBase { const std::vector& req, const std::vector& outputs) { mshadow::Stream* s = ctx.get_stream(); - MSHADOW_TYPE_SWITCH_WITH_BOOL(inputs[0].type_flag_, DType, { + MSHADOW_TYPE_SWITCH_EXT_WITH_BOOL(inputs[0].type_flag_, DType, { MXNET_ASSIGN_REQ_SWITCH(req[0], Req, { if (inputs[0].Size() != 0) { mxnet_op::Kernel, xpu>::Launch( @@ -700,7 +700,7 @@ void AroundOpForward(const nnvm::NodeAttrs& attrs, s, out_data.Size(), out_data.dptr(), in_data.dptr()); }); } else { - MSHADOW_TYPE_SWITCH(out_data.type_flag_, DType, { + MSHADOW_TYPE_SWITCH_EXT(out_data.type_flag_, DType, { MXNET_ASSIGN_REQ_SWITCH(req[0], req_type, { Kernel, xpu>::Launch( s, out_data.Size(), out_data.dptr(), in_data.dptr(), param.decimals); From 27dbc99414786676174457ecd85b00f4c8734903 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Wed, 3 Nov 2021 21:35:51 -0700 Subject: [PATCH 02/39] fix lint --- .github/workflows/os_x_staticbuild.yml | 33 +++ src/common/utils.h | 49 ++-- src/operator/numpy/np_elemwise_broadcast_op.h | 247 ++++++++++-------- .../np_elemwise_broadcast_op_extended.cc | 28 +- .../np_elemwise_broadcast_op_extended_thi.cc | 6 +- .../numpy/np_elemwise_broadcast_op_lae.cc | 41 +-- src/operator/numpy/np_true_divide-inl.h | 48 ++-- 7 files changed, 263 insertions(+), 189 deletions(-) diff --git a/.github/workflows/os_x_staticbuild.yml b/.github/workflows/os_x_staticbuild.yml index 019069ac32e6..2588d1e246cb 100644 --- a/.github/workflows/os_x_staticbuild.yml +++ b/.github/workflows/os_x_staticbuild.yml @@ -54,3 +54,36 @@ jobs: python3 -m pytest -n 4 --durations=50 --verbose tests/python/unittest/ -k 'not test_operator and not (test_subgraph or test_custom_op or test_external_op or test_recordimage_dataset_with_data_loader_multiworker or test_multi_worker or test_multi_worker_shape or test_multi_worker_forked_data_loader or test_multi_worker_dataloader_release_pool)' -m 'not serial' MXNET_ENGINE_TYPE=NaiveEngine python3 -m pytest -n 4 --durations=50 --verbose tests/python/unittest/ -k 'test_operator and not (test_subgraph or test_custom_op or test_external_op or test_recordimage_dataset_with_data_loader_multiworker or test_multi_worker or test_multi_worker_shape or test_multi_worker_forked_data_loader or test_multi_worker_dataloader_release_pool)' -m 'not serial' python3 -m pytest --durations=50 --verbose tests/python/unittest/ -k 'not (test_subgraph or test_custom_op or test_external_op or test_recordimage_dataset_with_data_loader_multiworker or test_multi_worker or test_multi_worker_shape or test_multi_worker_forked_data_loader or test_multi_worker_dataloader_release_pool)' -m 'serial' + + - name: Test Array API + env: + MXNET_ENFORCE_CYTHON: 0 + run: | + cd .. + git clone https://github.com/data-apis/array-api-tests.git + cd array-api-tests + git checkout c1dba80a196a03f880d2e0a998a272fb3867b720 + export ARRAY_API_TESTS_MODULE=mxnet.numpy pytest + export DMLC_LOG_STACK_TRACE_DEPTH=100 + python3 -m pytest --durations=50 --verbose \ + array_api_tests/test_type_promotion.py::test_elementwise_function_two_arg_bool_type_promotion + python3 -m pytest --durations=50 --verbose \ + array_api_tests/test_type_promotion.py::test_elementwise_function_two_arg_promoted_type_promotion + python3 -m pytest --durations=50 --verbose \ + array_api_tests/test_type_promotion.py::test_elementwise_function_one_arg_bool + python3 -m pytest --durations=50 --verbose \ + array_api_tests/test_type_promotion.py::test_elementwise_function_one_arg_type_promotion + python3 -m pytest --durations=50 --verbose \ + array_api_tests/test_type_promotion.py::test_operator_one_arg_type_promotion + python3 -m pytest --durations=50 --verbose \ + array_api_tests/test_type_promotion.py::test_operator_two_arg_bool_promotion + python3 -m pytest --durations=50 --verbose \ + array_api_tests/test_type_promotion.py::test_operator_two_arg_promoted_promotion + python3 -m pytest --durations=50 --verbose \ + array_api_tests/test_type_promotion.py::test_operator_inplace_two_arg_promoted_promotion + python3 -m pytest --durations=50 --verbose \ + array_api_tests/test_elementwise_functions.py + python3 -m pytest --durations=50 --verbose \ + array_api_tests/test_indexing.py + python3 -m pytest --durations=50 --verbose \ + array_api_tests/test_constants.py diff --git a/src/common/utils.h b/src/common/utils.h index 825b9fcec3a3..5f65eb44be20 100644 --- a/src/common/utils.h +++ b/src/common/utils.h @@ -901,35 +901,45 @@ inline bool is_float(const int dtype) { } inline bool is_int(const int dtype) { - return dtype == mshadow::kUint8 || dtype == mshadow::kInt8 || - dtype == mshadow::kUint16 || dtype == mshadow::kInt16 || - dtype == mshadow::kUint32 || dtype == mshadow::kInt32 || + return dtype == mshadow::kUint8 || dtype == mshadow::kInt8 || dtype == mshadow::kUint16 || + dtype == mshadow::kInt16 || dtype == mshadow::kUint32 || dtype == mshadow::kInt32 || dtype == mshadow::kUint64 || dtype == mshadow::kInt64; } inline bool is_signed_int(const int dtype) { - return dtype == mshadow::kInt8 || dtype == mshadow::kInt16 || - dtype == mshadow::kInt32 || dtype == mshadow::kInt64; + return dtype == mshadow::kInt8 || dtype == mshadow::kInt16 || dtype == mshadow::kInt32 || + dtype == mshadow::kInt64; } inline bool is_unsigned_int(const int dtype) { - return dtype == mshadow::kUint8 || dtype == mshadow::kUint16 || - dtype == mshadow::kUint32 || dtype == mshadow::kUint64; + return dtype == mshadow::kUint8 || dtype == mshadow::kUint16 || dtype == mshadow::kUint32 || + dtype == mshadow::kUint64; } static int bits_of(const int type_flag) { switch (type_flag) { - case mshadow::kFloat32: return sizeof(float) * CHAR_BIT; - case mshadow::kFloat64: return sizeof(double) * CHAR_BIT; - case mshadow::kUint8: return sizeof(uint8_t) * CHAR_BIT; - case mshadow::kInt32: return sizeof(int32_t) * CHAR_BIT; - case mshadow::kInt8: return sizeof(int8_t) * CHAR_BIT; - case mshadow::kInt64: return sizeof(int64_t) * CHAR_BIT; - case mshadow::kBool: return sizeof(bool) * CHAR_BIT; - case mshadow::kInt16: return sizeof(int16_t) * CHAR_BIT; - case mshadow::kUint16: return sizeof(uint16_t) * CHAR_BIT; - case mshadow::kUint32: return sizeof(uint32_t) * CHAR_BIT; - case mshadow::kUint64: return sizeof(uint64_t) * CHAR_BIT; + case mshadow::kFloat32: + return sizeof(float) * CHAR_BIT; + case mshadow::kFloat64: + return sizeof(double) * CHAR_BIT; + case mshadow::kUint8: + return sizeof(uint8_t) * CHAR_BIT; + case mshadow::kInt32: + return sizeof(int32_t) * CHAR_BIT; + case mshadow::kInt8: + return sizeof(int8_t) * CHAR_BIT; + case mshadow::kInt64: + return sizeof(int64_t) * CHAR_BIT; + case mshadow::kBool: + return sizeof(bool) * CHAR_BIT; + case mshadow::kInt16: + return sizeof(int16_t) * CHAR_BIT; + case mshadow::kUint16: + return sizeof(uint16_t) * CHAR_BIT; + case mshadow::kUint32: + return sizeof(uint32_t) * CHAR_BIT; + case mshadow::kUint64: + return sizeof(uint64_t) * CHAR_BIT; default: { LOG(FATAL) << "Unknown type_flag=" << type_flag; return -1; @@ -938,7 +948,8 @@ static int bits_of(const int type_flag) { } inline int type_promotion(const int type1, const int type2) { - if (type1 == type2) return type1; + if (type1 == type2) + return type1; if (is_float(type1) && is_float(type2)) { if (type1 == mshadow::kFloat64 || type2 == mshadow::kFloat64) { return mshadow::kFloat64; diff --git a/src/operator/numpy/np_elemwise_broadcast_op.h b/src/operator/numpy/np_elemwise_broadcast_op.h index 569358d3f326..1c02802e1ef9 100644 --- a/src/operator/numpy/np_elemwise_broadcast_op.h +++ b/src/operator/numpy/np_elemwise_broadcast_op.h @@ -125,7 +125,7 @@ void MixedIntRealBinaryElemwiseCompute(const OpContext& ctx, }); } -template +template void MixedIntBinaryElemwiseCompute(const nnvm::NodeAttrs& attrs, const OpContext& ctx, const TBlob& lhs, @@ -135,41 +135,41 @@ void MixedIntBinaryElemwiseCompute(const nnvm::NodeAttrs& attrs, using namespace mshadow; using namespace mxnet_op; - Stream *s = ctx.get_stream(); + Stream* s = ctx.get_stream(); TBlob temp_tblob; if (lhs.type_flag_ == out.type_flag_) { MXNET_INT_TYPE_SWITCH_EXT(lhs.type_flag_, LType, { Tensor temp_tensor = - ctx.requested[0].get_space_typed(Shape1(rhs.Size()), s); + ctx.requested[0].get_space_typed(Shape1(rhs.Size()), s); temp_tblob = TBlob(temp_tensor); }); CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob}); MXNET_ASSIGN_REQ_SWITCH(req, Req, { MXNET_INT_TYPE_SWITCH_EXT(out.type_flag_, DType, { - const size_t size = (ElemwiseBinaryOp::minthree(out.Size(), lhs.Size(), temp_tblob.Size()) - + DataType::kLanes - 1) / DataType::kLanes; + const size_t size = (ElemwiseBinaryOp::minthree(out.Size(), lhs.Size(), temp_tblob.Size()) + + DataType::kLanes - 1) / + DataType::kLanes; if (size != 0) { - Kernel, xpu>::Launch(s, size, - out.dptr(), - lhs.dptr(), temp_tblob.dptr()); + Kernel, xpu>::Launch( + s, size, out.dptr(), lhs.dptr(), temp_tblob.dptr()); } }); }); } else if (rhs.type_flag_ == out.type_flag_) { MXNET_INT_TYPE_SWITCH_EXT(rhs.type_flag_, RType, { Tensor temp_tensor = - ctx.requested[0].get_space_typed(Shape1(lhs.Size()), s); + ctx.requested[0].get_space_typed(Shape1(lhs.Size()), s); temp_tblob = TBlob(temp_tensor); }); CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob}); MXNET_ASSIGN_REQ_SWITCH(req, Req, { MXNET_INT_TYPE_SWITCH_EXT(out.type_flag_, DType, { - const size_t size = (ElemwiseBinaryOp::minthree(out.Size(), temp_tblob.Size(), rhs.Size()) - + DataType::kLanes - 1) / DataType::kLanes; + const size_t size = (ElemwiseBinaryOp::minthree(out.Size(), temp_tblob.Size(), rhs.Size()) + + DataType::kLanes - 1) / + DataType::kLanes; if (size != 0) { - Kernel, xpu>::Launch(s, size, - out.dptr(), - temp_tblob.dptr(), rhs.dptr()); + Kernel, xpu>::Launch( + s, size, out.dptr(), temp_tblob.dptr(), rhs.dptr()); } }); }); @@ -178,21 +178,28 @@ void MixedIntBinaryElemwiseCompute(const nnvm::NodeAttrs& attrs, TBlob temp_tblob_r; MXNET_INT_TYPE_SWITCH_EXT(out.type_flag_, OType, { Tensor workspace = - ctx.requested[0].get_space_typed(Shape1(lhs.Size() + rhs.Size()), s); + ctx.requested[0].get_space_typed(Shape1(lhs.Size() + rhs.Size()), s); TBlob temp_tblob = TBlob(workspace); - temp_tblob_l = TBlob(reinterpret_cast(temp_tblob.dptr_), lhs.shape_, temp_tblob.dev_mask(), temp_tblob.dev_id()); - temp_tblob_r = TBlob(reinterpret_cast(temp_tblob.dptr_) + lhs.Size() + 1, rhs.shape_, temp_tblob.dev_mask(), temp_tblob.dev_id()); + temp_tblob_l = TBlob(reinterpret_cast(temp_tblob.dptr_), + lhs.shape_, + temp_tblob.dev_mask(), + temp_tblob.dev_id()); + temp_tblob_r = TBlob(reinterpret_cast(temp_tblob.dptr_) + lhs.Size() + 1, + rhs.shape_, + temp_tblob.dev_mask(), + temp_tblob.dev_id()); }); CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob_l}); CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob_r}); MXNET_ASSIGN_REQ_SWITCH(req, Req, { MXNET_INT_TYPE_SWITCH_EXT(out.type_flag_, DType, { - const size_t size = (ElemwiseBinaryOp::minthree(out.Size(), temp_tblob_l.Size(), temp_tblob_r.Size()) - + DataType::kLanes - 1) / DataType::kLanes; + const size_t size = + (ElemwiseBinaryOp::minthree(out.Size(), temp_tblob_l.Size(), temp_tblob_r.Size()) + + DataType::kLanes - 1) / + DataType::kLanes; if (size != 0) { - Kernel, xpu>::Launch(s, size, - out.dptr(), - temp_tblob_l.dptr(), temp_tblob_r.dptr()); + Kernel, xpu>::Launch( + s, size, out.dptr(), temp_tblob_l.dptr(), temp_tblob_r.dptr()); } }); }); @@ -399,15 +406,20 @@ void MixedBinaryBroadcastCompute(const nnvm::NodeAttrs& attrs, TBlob temp_tblob_r; MXNET_INT_TYPE_SWITCH_EXT(out.type_flag_, OType, { Tensor workspace = - ctx.requested[0].get_space_typed(Shape1(lhs.Size() + rhs.Size()), s); + ctx.requested[0].get_space_typed(Shape1(lhs.Size() + rhs.Size()), s); TBlob temp_tblob = TBlob(workspace); - temp_tblob_l = TBlob(reinterpret_cast(temp_tblob.dptr_), lhs.shape_, temp_tblob.dev_mask(), temp_tblob.dev_id()); - temp_tblob_r = TBlob(reinterpret_cast(temp_tblob.dptr_) + lhs.Size() + 1, rhs.shape_, temp_tblob.dev_mask(), temp_tblob.dev_id()); + temp_tblob_l = TBlob(reinterpret_cast(temp_tblob.dptr_), + lhs.shape_, + temp_tblob.dev_mask(), + temp_tblob.dev_id()); + temp_tblob_r = TBlob(reinterpret_cast(temp_tblob.dptr_) + lhs.Size() + 1, + rhs.shape_, + temp_tblob.dev_mask(), + temp_tblob.dev_id()); }); CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob_l}); CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob_r}); - BinaryBroadcastCompute( - attrs, ctx, {temp_tblob_l, temp_tblob_r}, req, outputs); + BinaryBroadcastCompute(attrs, ctx, {temp_tblob_l, temp_tblob_r}, req, outputs); } } else { PrintErrorMessage(attrs.op->name, lhs.type_flag_, rhs.type_flag_); @@ -489,23 +501,27 @@ void NumpyBinaryBroadcastComputeWithBool(const nnvm::NodeAttrs& attrs, TBlob temp_tblob_r; MXNET_INT_TYPE_SWITCH_EXT(out.type_flag_, OType, { Tensor workspace = - ctx.requested[0].get_space_typed(Shape1(lhs.Size() + rhs.Size()), s); + ctx.requested[0].get_space_typed(Shape1(lhs.Size() + rhs.Size()), s); TBlob temp_tblob = TBlob(workspace); - temp_tblob_l = TBlob(reinterpret_cast(temp_tblob.dptr_), lhs.shape_, temp_tblob.dev_mask(), temp_tblob.dev_id()); - temp_tblob_r = TBlob(reinterpret_cast(temp_tblob.dptr_) + lhs.Size() + 1, rhs.shape_, temp_tblob.dev_mask(), temp_tblob.dev_id()); + temp_tblob_l = TBlob(reinterpret_cast(temp_tblob.dptr_), + lhs.shape_, + temp_tblob.dev_mask(), + temp_tblob.dev_id()); + temp_tblob_r = TBlob(reinterpret_cast(temp_tblob.dptr_) + lhs.Size() + 1, + rhs.shape_, + temp_tblob.dev_mask(), + temp_tblob.dev_id()); }); CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob_l}); CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob_r}); - BinaryBroadcastCompute( - attrs, ctx, {temp_tblob_l, temp_tblob_r}, req, outputs); + BinaryBroadcastCompute(attrs, ctx, {temp_tblob_l, temp_tblob_r}, req, outputs); } return; } MixedBinaryBroadcastCompute(attrs, ctx, inputs, req, outputs); } - -template +template void NumpyBinaryBroadcastIntComputeWithBool(const nnvm::NodeAttrs& attrs, const OpContext& ctx, const std::vector& inputs, @@ -520,52 +536,57 @@ void NumpyBinaryBroadcastIntComputeWithBool(const nnvm::NodeAttrs& attrs, const TBlob& rhs = inputs[1]; const TBlob& out = outputs[0]; - if ((out.shape_.Size() == 0U) || (req[0] == kNullOp)) return; + if ((out.shape_.Size() == 0U) || (req[0] == kNullOp)) + return; if (lhs.type_flag_ == rhs.type_flag_) { BinaryBroadcastIntCompute(attrs, ctx, inputs, req, outputs); return; } - Stream *s = ctx.get_stream(); + Stream* s = ctx.get_stream(); TBlob temp_tblob; if (lhs.type_flag_ == out.type_flag_) { MXNET_INT_TYPE_SWITCH_EXT(lhs.type_flag_, LType, { Tensor temp_tensor = - ctx.requested[0].get_space_typed(Shape1(rhs.Size()), s); + ctx.requested[0].get_space_typed(Shape1(rhs.Size()), s); temp_tblob = TBlob(temp_tensor); }); CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob}); BinaryBroadcastIntCompute( - attrs, ctx, {lhs, temp_tblob.reshape(rhs.shape_)}, req, outputs); + attrs, ctx, {lhs, temp_tblob.reshape(rhs.shape_)}, req, outputs); } else if (rhs.type_flag_ == out.type_flag_) { MXNET_INT_TYPE_SWITCH_EXT(rhs.type_flag_, RType, { Tensor temp_tensor = - ctx.requested[0].get_space_typed(Shape1(lhs.Size()), s); + ctx.requested[0].get_space_typed(Shape1(lhs.Size()), s); temp_tblob = TBlob(temp_tensor); }); CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob}); BinaryBroadcastIntCompute( - attrs, ctx, {temp_tblob.reshape(lhs.shape_), rhs}, req, outputs); + attrs, ctx, {temp_tblob.reshape(lhs.shape_), rhs}, req, outputs); } else { TBlob temp_tblob_l; TBlob temp_tblob_r; MXNET_INT_TYPE_SWITCH_EXT(out.type_flag_, OType, { Tensor workspace = - ctx.requested[0].get_space_typed(Shape1(lhs.Size() + rhs.Size()), s); + ctx.requested[0].get_space_typed(Shape1(lhs.Size() + rhs.Size()), s); TBlob temp_tblob = TBlob(workspace); - temp_tblob_l = TBlob(reinterpret_cast(temp_tblob.dptr_), lhs.shape_, temp_tblob.dev_mask(), temp_tblob.dev_id()); - temp_tblob_r = TBlob(reinterpret_cast(temp_tblob.dptr_) + lhs.Size() + 1, rhs.shape_, temp_tblob.dev_mask(), temp_tblob.dev_id()); + temp_tblob_l = TBlob(reinterpret_cast(temp_tblob.dptr_), + lhs.shape_, + temp_tblob.dev_mask(), + temp_tblob.dev_id()); + temp_tblob_r = TBlob(reinterpret_cast(temp_tblob.dptr_) + lhs.Size() + 1, + rhs.shape_, + temp_tblob.dev_mask(), + temp_tblob.dev_id()); }); CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob_l}); CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob_r}); - BinaryBroadcastIntCompute( - attrs, ctx, {temp_tblob_l, temp_tblob_r}, req, outputs); + BinaryBroadcastIntCompute(attrs, ctx, {temp_tblob_l, temp_tblob_r}, req, outputs); } return; } - -template +template void NumpyBinaryBroadcastIntCompute(const nnvm::NodeAttrs& attrs, const OpContext& ctx, const std::vector& inputs, @@ -580,46 +601,52 @@ void NumpyBinaryBroadcastIntCompute(const nnvm::NodeAttrs& attrs, const TBlob& rhs = inputs[1]; const TBlob& out = outputs[0]; - if ((out.shape_.Size() == 0U) || (req[0] == kNullOp)) return; + if ((out.shape_.Size() == 0U) || (req[0] == kNullOp)) + return; if (lhs.type_flag_ == rhs.type_flag_) { BinaryBroadcastIntCompute(attrs, ctx, inputs, req, outputs); return; } - Stream *s = ctx.get_stream(); + Stream* s = ctx.get_stream(); TBlob temp_tblob; if (lhs.type_flag_ == out.type_flag_) { MXNET_INT_TYPE_SWITCH_EXT(lhs.type_flag_, LType, { Tensor temp_tensor = - ctx.requested[0].get_space_typed(Shape1(rhs.Size()), s); + ctx.requested[0].get_space_typed(Shape1(rhs.Size()), s); temp_tblob = TBlob(temp_tensor); }); CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob}); BinaryBroadcastIntCompute( - attrs, ctx, {lhs, temp_tblob.reshape(rhs.shape_)}, req, outputs); + attrs, ctx, {lhs, temp_tblob.reshape(rhs.shape_)}, req, outputs); } else if (rhs.type_flag_ == out.type_flag_) { MXNET_INT_TYPE_SWITCH_EXT(rhs.type_flag_, RType, { Tensor temp_tensor = - ctx.requested[0].get_space_typed(Shape1(lhs.Size()), s); + ctx.requested[0].get_space_typed(Shape1(lhs.Size()), s); temp_tblob = TBlob(temp_tensor); }); CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob}); BinaryBroadcastIntCompute( - attrs, ctx, {temp_tblob.reshape(lhs.shape_), rhs}, req, outputs); + attrs, ctx, {temp_tblob.reshape(lhs.shape_), rhs}, req, outputs); } else { TBlob temp_tblob_l; TBlob temp_tblob_r; MXNET_INT_TYPE_SWITCH_EXT(out.type_flag_, OType, { Tensor workspace = - ctx.requested[0].get_space_typed(Shape1(lhs.Size() + rhs.Size()), s); + ctx.requested[0].get_space_typed(Shape1(lhs.Size() + rhs.Size()), s); TBlob temp_tblob = TBlob(workspace); - temp_tblob_l = TBlob(reinterpret_cast(temp_tblob.dptr_), lhs.shape_, temp_tblob.dev_mask(), temp_tblob.dev_id()); - temp_tblob_r = TBlob(reinterpret_cast(temp_tblob.dptr_) + lhs.Size() + 1, rhs.shape_, temp_tblob.dev_mask(), temp_tblob.dev_id()); + temp_tblob_l = TBlob(reinterpret_cast(temp_tblob.dptr_), + lhs.shape_, + temp_tblob.dev_mask(), + temp_tblob.dev_id()); + temp_tblob_r = TBlob(reinterpret_cast(temp_tblob.dptr_) + lhs.Size() + 1, + rhs.shape_, + temp_tblob.dev_mask(), + temp_tblob.dev_id()); }); CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob_l}); CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob_r}); - BinaryBroadcastIntCompute( - attrs, ctx, {temp_tblob_l, temp_tblob_r}, req, outputs); + BinaryBroadcastIntCompute(attrs, ctx, {temp_tblob_l, temp_tblob_r}, req, outputs); } return; } @@ -646,7 +673,7 @@ inline bool NumpyBinaryMixedFloatingType(const nnvm::NodeAttrs& attrs, return out_attrs->at(0) != -1; } -template +template void NumpyBinaryMixedFloatingCompute(const nnvm::NodeAttrs& attrs, const OpContext& ctx, const std::vector& inputs, @@ -661,32 +688,33 @@ void NumpyBinaryMixedFloatingCompute(const nnvm::NodeAttrs& attrs, const TBlob& rhs = inputs[1]; const TBlob& out = outputs[0]; - if ((out.shape_.Size() == 0U) || (req[0] == kNullOp)) return; + if ((out.shape_.Size() == 0U) || (req[0] == kNullOp)) + return; if (lhs.type_flag_ == rhs.type_flag_) { BinaryBroadcastCompute(attrs, ctx, inputs, req, outputs); return; } - Stream *s = ctx.get_stream(); + Stream* s = ctx.get_stream(); TBlob temp_tblob; if (lhs.type_flag_ == out.type_flag_) { MSHADOW_REAL_TYPE_SWITCH(lhs.type_flag_, LType, { Tensor temp_tensor = - ctx.requested[0].get_space_typed(Shape1(rhs.Size()), s); + ctx.requested[0].get_space_typed(Shape1(rhs.Size()), s); temp_tblob = TBlob(temp_tensor); }); CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob}); BinaryBroadcastCompute( - attrs, ctx, {lhs, temp_tblob.reshape(rhs.shape_)}, req, outputs); + attrs, ctx, {lhs, temp_tblob.reshape(rhs.shape_)}, req, outputs); } else { MSHADOW_REAL_TYPE_SWITCH(rhs.type_flag_, RType, { Tensor temp_tensor = - ctx.requested[0].get_space_typed(Shape1(lhs.Size()), s); + ctx.requested[0].get_space_typed(Shape1(lhs.Size()), s); temp_tblob = TBlob(temp_tensor); }); CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob}); BinaryBroadcastCompute( - attrs, ctx, {temp_tblob.reshape(lhs.shape_), rhs}, req, outputs); + attrs, ctx, {temp_tblob.reshape(lhs.shape_), rhs}, req, outputs); } return; } @@ -882,8 +910,10 @@ inline bool NumpyBinaryMixedIntPrecisionTypeWithBool(const nnvm::NodeAttrs& attr CHECK_EQ(out_attrs->size(), 1U); const int ltype = in_attrs->at(0); const int rtype = in_attrs->at(1); - CHECK(common::is_int(ltype) || ltype == mshadow::kBool) << "1st input only supports integer types or bool types."; - CHECK(common::is_int(rtype) || rtype == mshadow::kBool) << "2nd input only supports integer types or bool types."; + CHECK(common::is_int(ltype) || ltype == mshadow::kBool) + << "1st input only supports integer types or bool types."; + CHECK(common::is_int(rtype) || rtype == mshadow::kBool) + << "2nd input only supports integer types or bool types."; if (ltype != -1 && rtype != -1 && (ltype != rtype)) { // Only when both input types are known and not the same, we enter the mixed-precision mode TYPE_ASSIGN_CHECK(*out_attrs, 0, common::type_promotion(ltype, rtype)); @@ -893,27 +923,27 @@ inline bool NumpyBinaryMixedIntPrecisionTypeWithBool(const nnvm::NodeAttrs& attr return true; } -#define MXNET_OPERATOR_REGISTER_NP_BINARY_MIXED_INT_PRECISION_WITH_BOOL(name) \ - NNVM_REGISTER_OP(name) \ - .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", NumpyBinaryMixedIntPrecisionTypeWithBool) \ - .set_attr("FInplaceOption", \ - [](const NodeAttrs& attrs){ \ - return std::vector >{{0, 0}, {1, 0}}; \ - }) \ - .set_attr("FResourceRequest", \ - [](const NodeAttrs& attrs) { \ - return std::vector{ResourceRequest::kTempSpace}; \ - }) \ - .add_argument("lhs", "NDArray-or-Symbol", "First input to the function") \ - .add_argument("rhs", "NDArray-or-Symbol", "Second input to the function") - +#define MXNET_OPERATOR_REGISTER_NP_BINARY_MIXED_INT_PRECISION_WITH_BOOL(name) \ + NNVM_REGISTER_OP(name) \ + .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", NumpyBinaryMixedIntPrecisionTypeWithBool) \ + .set_attr("FInplaceOption", \ + [](const NodeAttrs& attrs) { \ + return std::vector >{{0, 0}, {1, 0}}; \ + }) \ + .set_attr( \ + "FResourceRequest", \ + [](const NodeAttrs& attrs) { \ + return std::vector{ResourceRequest::kTempSpace}; \ + }) \ + .add_argument("lhs", "NDArray-or-Symbol", "First input to the function") \ + .add_argument("rhs", "NDArray-or-Symbol", "Second input to the function") inline bool NumpyBinaryMixedIntPrecisionType(const nnvm::NodeAttrs& attrs, std::vector* in_attrs, @@ -933,26 +963,27 @@ inline bool NumpyBinaryMixedIntPrecisionType(const nnvm::NodeAttrs& attrs, return true; } -#define MXNET_OPERATOR_REGISTER_NP_BINARY_MIXED_INT_PRECISION(name) \ - NNVM_REGISTER_OP(name) \ - .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", NumpyBinaryMixedIntPrecisionType) \ - .set_attr("FInplaceOption", \ - [](const NodeAttrs& attrs){ \ - return std::vector >{{0, 0}, {1, 0}}; \ - }) \ - .set_attr("FResourceRequest", \ - [](const NodeAttrs& attrs) { \ - return std::vector{ResourceRequest::kTempSpace}; \ - }) \ - .add_argument("lhs", "NDArray-or-Symbol", "First input to the function") \ - .add_argument("rhs", "NDArray-or-Symbol", "Second input to the function") +#define MXNET_OPERATOR_REGISTER_NP_BINARY_MIXED_INT_PRECISION(name) \ + NNVM_REGISTER_OP(name) \ + .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", NumpyBinaryMixedIntPrecisionType) \ + .set_attr("FInplaceOption", \ + [](const NodeAttrs& attrs) { \ + return std::vector >{{0, 0}, {1, 0}}; \ + }) \ + .set_attr( \ + "FResourceRequest", \ + [](const NodeAttrs& attrs) { \ + return std::vector{ResourceRequest::kTempSpace}; \ + }) \ + .add_argument("lhs", "NDArray-or-Symbol", "First input to the function") \ + .add_argument("rhs", "NDArray-or-Symbol", "Second input to the function") } // namespace op } // namespace mxnet diff --git a/src/operator/numpy/np_elemwise_broadcast_op_extended.cc b/src/operator/numpy/np_elemwise_broadcast_op_extended.cc index fff0f7c66184..33fddab3149b 100644 --- a/src/operator/numpy/np_elemwise_broadcast_op_extended.cc +++ b/src/operator/numpy/np_elemwise_broadcast_op_extended.cc @@ -131,10 +131,9 @@ NNVM_REGISTER_OP(_npi_lcm_scalar) .set_attr("FCompute", BinaryScalarOp::ComputeInt); MXNET_OPERATOR_REGISTER_NP_BINARY_MIXED_INT_PRECISION_WITH_BOOL(_npi_bitwise_and) -.set_attr( - "FCompute", - NumpyBinaryBroadcastIntComputeWithBool) -.set_attr("FGradient", MakeZeroGradNodes); + .set_attr("FCompute", + NumpyBinaryBroadcastIntComputeWithBool) + .set_attr("FGradient", MakeZeroGradNodes); NNVM_REGISTER_OP(_npi_bitwise_and_scalar) .set_num_inputs(1) @@ -152,16 +151,14 @@ NNVM_REGISTER_OP(_npi_bitwise_and_scalar) .set_attr("FCompute", BinaryScalarOp::ComputeInt); MXNET_OPERATOR_REGISTER_NP_BINARY_MIXED_INT_PRECISION_WITH_BOOL(_npi_bitwise_xor) -.set_attr( - "FCompute", - NumpyBinaryBroadcastIntComputeWithBool) -.set_attr("FGradient", MakeZeroGradNodes); + .set_attr("FCompute", + NumpyBinaryBroadcastIntComputeWithBool) + .set_attr("FGradient", MakeZeroGradNodes); MXNET_OPERATOR_REGISTER_NP_BINARY_MIXED_INT_PRECISION_WITH_BOOL(_npi_bitwise_or) -.set_attr( - "FCompute", - NumpyBinaryBroadcastIntComputeWithBool) -.set_attr("FGradient", MakeZeroGradNodes); + .set_attr("FCompute", + NumpyBinaryBroadcastIntComputeWithBool) + .set_attr("FGradient", MakeZeroGradNodes); NNVM_REGISTER_OP(_npi_bitwise_xor_scalar) .set_num_inputs(1) @@ -204,7 +201,6 @@ MXNET_OPERATOR_REGISTER_NP_BINARY_SCALAR(_npi_rcopysign_scalar) MXNET_OPERATOR_REGISTER_NP_BINARY_SCALAR(_backward_npi_copysign_scalar) .set_attr("FCompute", BinaryScalarOp::Backward); - NNVM_REGISTER_OP(_npi_arctan2) .set_num_inputs(2) .set_num_outputs(1) @@ -221,9 +217,9 @@ NNVM_REGISTER_OP(_npi_arctan2) return std::vector >{{0, 0}}; }) .set_attr("FResourceRequest", - [](const NodeAttrs& attrs) { - return std::vector{ResourceRequest::kTempSpace}; - }) + [](const NodeAttrs& attrs) { + return std::vector{ResourceRequest::kTempSpace}; + }) .add_argument("x1", "NDArray-or-Symbol", "The input array") .add_argument("x2", "NDArray-or-Symbol", "The input array"); 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 7f13b79f0e2d..7fc8d9a9635f 100644 --- a/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cc +++ b/src/operator/numpy/np_elemwise_broadcast_op_extended_thi.cc @@ -45,7 +45,8 @@ namespace op { .add_arguments(NumpyBinaryScalarParam::__FIELDS__()) MXNET_OPERATOR_REGISTER_NP_BINARY_MIXED_INT_PRECISION(_npi_bitwise_left_shift) - .set_attr("FCompute", NumpyBinaryBroadcastIntCompute) + .set_attr("FCompute", + NumpyBinaryBroadcastIntCompute) .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_bitwise_left_shift"}); NNVM_REGISTER_OP(_npi_bitwise_left_shift_scalar) @@ -112,7 +113,8 @@ MXNET_OPERATOR_REGISTER_BINARY(_backward_npi_rbitwise_left_shift_scalar) BinaryScalarOp::Backward); MXNET_OPERATOR_REGISTER_NP_BINARY_MIXED_INT_PRECISION(_npi_bitwise_right_shift) - .set_attr("FCompute", NumpyBinaryBroadcastIntCompute) + .set_attr("FCompute", + NumpyBinaryBroadcastIntCompute) .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_bitwise_right_shift"}); NNVM_REGISTER_OP(_npi_bitwise_right_shift_scalar) diff --git a/src/operator/numpy/np_elemwise_broadcast_op_lae.cc b/src/operator/numpy/np_elemwise_broadcast_op_lae.cc index e53df7e463dd..0298d1dbd3ff 100644 --- a/src/operator/numpy/np_elemwise_broadcast_op_lae.cc +++ b/src/operator/numpy/np_elemwise_broadcast_op_lae.cc @@ -28,26 +28,27 @@ namespace mxnet { namespace op { NNVM_REGISTER_OP(_npi_logaddexp) -.set_num_inputs(2) -.set_num_outputs(1) -.set_attr("FListInputNames", - [](const NodeAttrs& attrs) { - return std::vector{"x1", "x2"}; - }) -.set_attr("FInferShape", BinaryBroadcastShape) -.set_attr("FInferType", NumpyBinaryMixedFloatingType) -.set_attr("FCompute", NumpyBinaryMixedFloatingCompute) -.set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_logaddexp"}) -.set_attr("FResourceRequest", - [](const NodeAttrs& attrs) { - return std::vector{ResourceRequest::kTempSpace}; - }) -.set_attr("FInplaceOption", - [](const NodeAttrs& attrs) { - return std::vector >{{0, 0}}; - }) -.add_argument("x1", "NDArray-or-Symbol", "The input array") -.add_argument("x2", "NDArray-or-Symbol", "The input array"); + .set_num_inputs(2) + .set_num_outputs(1) + .set_attr("FListInputNames", + [](const NodeAttrs& attrs) { + return std::vector{"x1", "x2"}; + }) + .set_attr("FInferShape", BinaryBroadcastShape) + .set_attr("FInferType", NumpyBinaryMixedFloatingType) + .set_attr("FCompute", + NumpyBinaryMixedFloatingCompute) + .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_logaddexp"}) + .set_attr("FResourceRequest", + [](const NodeAttrs& attrs) { + return std::vector{ResourceRequest::kTempSpace}; + }) + .set_attr("FInplaceOption", + [](const NodeAttrs& attrs) { + return std::vector >{{0, 0}}; + }) + .add_argument("x1", "NDArray-or-Symbol", "The input array") + .add_argument("x2", "NDArray-or-Symbol", "The input array"); MXNET_OPERATOR_REGISTER_NP_BINARY_SCALAR(_npi_logaddexp_scalar) .set_attr("FCompute", BinaryScalarOp::Compute) diff --git a/src/operator/numpy/np_true_divide-inl.h b/src/operator/numpy/np_true_divide-inl.h index acaf4cd25434..047489f648cc 100644 --- a/src/operator/numpy/np_true_divide-inl.h +++ b/src/operator/numpy/np_true_divide-inl.h @@ -121,7 +121,7 @@ void TrueDivideElemwiseCompute(const nnvm::NodeAttrs& attrs, if (lhs.type_flag_ == out.type_flag_) { MSHADOW_REAL_TYPE_SWITCH(lhs.type_flag_, LType, { Tensor temp_tensor = - ctx.requested[0].get_space_typed(Shape1(rhs.Size()), s); + ctx.requested[0].get_space_typed(Shape1(rhs.Size()), s); temp_tblob = TBlob(temp_tensor); }); CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob}); @@ -134,7 +134,7 @@ void TrueDivideElemwiseCompute(const nnvm::NodeAttrs& attrs, } else { MSHADOW_REAL_TYPE_SWITCH(rhs.type_flag_, RType, { Tensor temp_tensor = - ctx.requested[0].get_space_typed(Shape1(lhs.Size()), s); + ctx.requested[0].get_space_typed(Shape1(lhs.Size()), s); temp_tblob = TBlob(temp_tensor); }); CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob}); @@ -244,40 +244,40 @@ void TrueDivideBroadcastCompute(const nnvm::NodeAttrs& attrs, if (lhs.type_flag_ == out.type_flag_) { MSHADOW_REAL_TYPE_SWITCH(lhs.type_flag_, LType, { Tensor temp_tensor = - ctx.requested[0].get_space_typed(Shape1(rhs.Size()), s); + ctx.requested[0].get_space_typed(Shape1(rhs.Size()), s); temp_tblob = TBlob(temp_tensor); }); CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob}); MSHADOW_REAL_TYPE_SWITCH(out.type_flag_, DType, { - Kernel, xpu>::template LaunchEx( - s, - new_oshape.Size(), - req[0], - lstride, - rstride, - oshape, - lhs.dptr(), - temp_tblob.dptr(), - out.dptr()); + Kernel, + xpu>::template LaunchEx(s, + new_oshape.Size(), + req[0], + lstride, + rstride, + oshape, + lhs.dptr(), + temp_tblob.dptr(), + out.dptr()); }); } else { MSHADOW_REAL_TYPE_SWITCH(rhs.type_flag_, RType, { Tensor temp_tensor = - ctx.requested[0].get_space_typed(Shape1(lhs.Size()), s); + ctx.requested[0].get_space_typed(Shape1(lhs.Size()), s); temp_tblob = TBlob(temp_tensor); }); CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob}); MSHADOW_REAL_TYPE_SWITCH(out.type_flag_, DType, { - Kernel, xpu>::template LaunchEx( - s, - new_oshape.Size(), - req[0], - lstride, - rstride, - oshape, - temp_tblob.dptr(), - rhs.dptr(), - out.dptr()); + Kernel, + xpu>::template LaunchEx(s, + new_oshape.Size(), + req[0], + lstride, + rstride, + oshape, + temp_tblob.dptr(), + rhs.dptr(), + out.dptr()); }); } } else if (common::is_float(lhs.type_flag_) || common::is_float(rhs.type_flag_)) { From 9e80789a922a118327f8fe577d8ecf307becb8fa Mon Sep 17 00:00:00 2001 From: barry-jin Date: Thu, 4 Nov 2021 10:04:45 -0700 Subject: [PATCH 03/39] fix lint --- src/operator/mxnet_op.h | 104 ++++++++++++++++++++-------------------- 1 file changed, 52 insertions(+), 52 deletions(-) diff --git a/src/operator/mxnet_op.h b/src/operator/mxnet_op.h index c805a2e858da..c8a00fbfeefa 100644 --- a/src/operator/mxnet_op.h +++ b/src/operator/mxnet_op.h @@ -473,58 +473,58 @@ 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: { \ - LOG(FATAL) << "This operation only support " \ - "integer types, not bool type"; \ - } break; \ - default: \ - 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: { \ + LOG(FATAL) << "This operation only support " \ + "integer types, not bool type"; \ + } break; \ + default: \ + LOG(FATAL) << "Unknown type enum " << type; \ } #define MXNET_INT32_INT64_TYPE_SWITCH(type, DType, ...) \ From 0de420111492dc289d12b16edd4b311d4a1e77a5 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Thu, 4 Nov 2021 10:27:32 -0700 Subject: [PATCH 04/39] fix --- src/operator/numpy/np_elemwise_broadcast_op.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/operator/numpy/np_elemwise_broadcast_op.h b/src/operator/numpy/np_elemwise_broadcast_op.h index a09fc8e4881e..9a7a4981c45d 100644 --- a/src/operator/numpy/np_elemwise_broadcast_op.h +++ b/src/operator/numpy/np_elemwise_broadcast_op.h @@ -176,7 +176,7 @@ void MixedIntBinaryElemwiseCompute(const nnvm::NodeAttrs& attrs, } else { TBlob temp_tblob_l; TBlob temp_tblob_r; - MXNET_INT_TYPE_SWITCH_EXT_WITH(out.type_flag_, OType, { + MXNET_INT_TYPE_SWITCH_EXT(out.type_flag_, OType, { Tensor workspace = ctx.requested[0].get_space_typed(Shape1(lhs.Size() + rhs.Size()), s); TBlob temp_tblob = TBlob(workspace); From 4e20e9c6013115176072227c5b53fafbc59a6216 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Thu, 4 Nov 2021 11:02:37 -0700 Subject: [PATCH 05/39] fix build --- src/operator/numpy/np_elemwise_broadcast_op.h | 8 ++-- .../tensor/elemwise_binary_broadcast_op.h | 40 +++++++++++++++++++ src/operator/tensor/elemwise_binary_op.h | 25 ++++++++++++ 3 files changed, 69 insertions(+), 4 deletions(-) diff --git a/src/operator/numpy/np_elemwise_broadcast_op.h b/src/operator/numpy/np_elemwise_broadcast_op.h index 9a7a4981c45d..bf2094167d61 100644 --- a/src/operator/numpy/np_elemwise_broadcast_op.h +++ b/src/operator/numpy/np_elemwise_broadcast_op.h @@ -540,7 +540,7 @@ void NumpyBinaryBroadcastIntComputeWithBool(const nnvm::NodeAttrs& attrs, return; if (lhs.type_flag_ == rhs.type_flag_) { - BinaryBroadcastIntCompute(attrs, ctx, inputs, req, outputs); + BinaryBroadcastIntComputeWithBool(attrs, ctx, inputs, req, outputs); return; } Stream* s = ctx.get_stream(); @@ -552,7 +552,7 @@ void NumpyBinaryBroadcastIntComputeWithBool(const nnvm::NodeAttrs& attrs, temp_tblob = TBlob(temp_tensor); }); CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob}); - BinaryBroadcastIntCompute( + BinaryBroadcastIntComputeWithBool( attrs, ctx, {lhs, temp_tblob.reshape(rhs.shape_)}, req, outputs); } else if (rhs.type_flag_ == out.type_flag_) { MXNET_INT_TYPE_SWITCH_EXT_WITH_BOOL(rhs.type_flag_, RType, { @@ -561,7 +561,7 @@ void NumpyBinaryBroadcastIntComputeWithBool(const nnvm::NodeAttrs& attrs, temp_tblob = TBlob(temp_tensor); }); CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob}); - BinaryBroadcastIntCompute( + BinaryBroadcastIntComputeWithBool( attrs, ctx, {temp_tblob.reshape(lhs.shape_), rhs}, req, outputs); } else { TBlob temp_tblob_l; @@ -581,7 +581,7 @@ void NumpyBinaryBroadcastIntComputeWithBool(const nnvm::NodeAttrs& attrs, }); CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob_l}); CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob_r}); - BinaryBroadcastIntCompute(attrs, ctx, {temp_tblob_l, temp_tblob_r}, req, outputs); + BinaryBroadcastIntComputeWithBool(attrs, ctx, {temp_tblob_l, temp_tblob_r}, req, outputs); } return; } diff --git a/src/operator/tensor/elemwise_binary_broadcast_op.h b/src/operator/tensor/elemwise_binary_broadcast_op.h index e793f1014c3e..3485607f37e0 100644 --- a/src/operator/tensor/elemwise_binary_broadcast_op.h +++ b/src/operator/tensor/elemwise_binary_broadcast_op.h @@ -209,6 +209,46 @@ void BinaryBroadcastIntCompute(const nnvm::NodeAttrs& attrs, inputs[0].shape_, inputs[1].shape_, outputs[0].shape_, &new_lshape, &new_rshape, &new_oshape); if (!ndim) { ElemwiseBinaryOp::ComputeInt(attrs, ctx, inputs, req, outputs); + } else { + if (req[0] == kNullOp) + return; + mshadow::Stream* s = ctx.get_stream(); + if (outputs[0].type_flag_ == mshadow::kBool) { + LOG(FATAL) << "Operator " << attrs.op->name << " does not support boolean type"; + } + 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()); + mshadow::Shape rstride = mxnet_op::calc_stride(new_rshape.get()); + mxnet_op::Kernel, xpu>::template LaunchEx( + s, + new_oshape.Size(), + req[0], + lstride, + rstride, + oshape, + inputs[0].dptr(), + inputs[1].dptr(), + outputs[0].dptr()); + }); + }); + } +} + +template +void BinaryBroadcastIntComputeWithBool(const nnvm::NodeAttrs& attrs, + const OpContext& ctx, + const std::vector& inputs, + const std::vector& req, + const std::vector& outputs) { + if (outputs[0].shape_.Size() == 0U) + return; + mxnet::TShape new_lshape, new_rshape, new_oshape; + int ndim = BinaryBroadcastShapeCompact( + inputs[0].shape_, inputs[1].shape_, outputs[0].shape_, &new_lshape, &new_rshape, &new_oshape); + if (!ndim) { + ElemwiseBinaryOp::ComputeIntWithBool(attrs, ctx, inputs, req, outputs); } else { if (req[0] == kNullOp) return; diff --git a/src/operator/tensor/elemwise_binary_op.h b/src/operator/tensor/elemwise_binary_op.h index 58267e558cb1..1904bf9ef4f4 100644 --- a/src/operator/tensor/elemwise_binary_op.h +++ b/src/operator/tensor/elemwise_binary_op.h @@ -448,6 +448,31 @@ class ElemwiseBinaryOp : public OpBase { Stream* s = ctx.get_stream(); CHECK_EQ(inputs.size(), 2U); CHECK_EQ(outputs.size(), 1U); + MXNET_ASSIGN_REQ_SWITCH(req[0], Req, { + 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; + if (size != 0) { + Kernel, xpu>::Launch( + s, size, outputs[0].dptr(), inputs[0].dptr(), inputs[1].dptr()); + } + }); + }); + } + + template + static void ComputeIntWithBool(const nnvm::NodeAttrs& attrs, + const OpContext& ctx, + const std::vector& inputs, + const std::vector& req, + const std::vector& outputs) { + using namespace mxnet_op; + if (req[0] == kNullOp) + return; + Stream* s = ctx.get_stream(); + CHECK_EQ(inputs.size(), 2U); + CHECK_EQ(outputs.size(), 1U); MXNET_ASSIGN_REQ_SWITCH(req[0], Req, { MXNET_INT_TYPE_SWITCH_EXT_WITH_BOOL(outputs[0].type_flag_, DType, { const size_t size = (minthree(outputs[0].Size(), inputs[0].Size(), inputs[1].Size()) + From 19d501e071231f6b906f72d7ca89de54787cab86 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Thu, 4 Nov 2021 11:05:27 -0700 Subject: [PATCH 06/39] fix lint --- src/operator/numpy/np_elemwise_broadcast_op.h | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/src/operator/numpy/np_elemwise_broadcast_op.h b/src/operator/numpy/np_elemwise_broadcast_op.h index bf2094167d61..fb79dc70080b 100644 --- a/src/operator/numpy/np_elemwise_broadcast_op.h +++ b/src/operator/numpy/np_elemwise_broadcast_op.h @@ -581,7 +581,8 @@ void NumpyBinaryBroadcastIntComputeWithBool(const nnvm::NodeAttrs& attrs, }); CastCompute(attrs, ctx, {lhs}, {kWriteTo}, {temp_tblob_l}); CastCompute(attrs, ctx, {rhs}, {kWriteTo}, {temp_tblob_r}); - BinaryBroadcastIntComputeWithBool(attrs, ctx, {temp_tblob_l, temp_tblob_r}, req, outputs); + BinaryBroadcastIntComputeWithBool( + attrs, ctx, {temp_tblob_l, temp_tblob_r}, req, outputs); } return; } From f6bc4d13e26d89e8c12dc1487a321eba28d43cd1 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Thu, 4 Nov 2021 13:10:55 -0700 Subject: [PATCH 07/39] update --- ci/docker/runtime_functions.sh | 2 +- src/common/utils.h | 4 ++++ src/operator/mshadow_op.h | 4 ++-- src/operator/tensor/elemwise_binary_broadcast_op.h | 3 --- src/operator/tensor/elemwise_binary_op.h | 2 +- 5 files changed, 8 insertions(+), 7 deletions(-) diff --git a/ci/docker/runtime_functions.sh b/ci/docker/runtime_functions.sh index 59d06e7c1416..eaff0f4b2c18 100755 --- a/ci/docker/runtime_functions.sh +++ b/ci/docker/runtime_functions.sh @@ -874,7 +874,7 @@ unittest_array_api_standardization() { export DMLC_LOG_STACK_TRACE_DEPTH=100 python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_creation_functions.py python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_indexing.py - python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_elementwise_functions.py + python3 -m pytest -s --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_elementwise_functions.py python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_constants.py python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ array_api_tests/test_type_promotion.py::test_elementwise_function_two_arg_bool_type_promotion diff --git a/src/common/utils.h b/src/common/utils.h index 5f65eb44be20..3965f9fcf35a 100644 --- a/src/common/utils.h +++ b/src/common/utils.h @@ -1015,6 +1015,10 @@ inline int type_promotion(const int type1, const int type2) { return mshadow::kInt64; } } + } else if (type1 == mshadow::kBool) { + return type2; + } else if (type2 == mshadow::kBool) { + return type1; } LOG(FATAL) << "should not reach here "; return -1; diff --git a/src/operator/mshadow_op.h b/src/operator/mshadow_op.h index 2404abee6317..65796476cf9e 100644 --- a/src/operator/mshadow_op.h +++ b/src/operator/mshadow_op.h @@ -822,7 +822,7 @@ MXNET_BINARY_MATH_OP(bitwise_or, static_cast(a) | static_cast( struct bitwise_left_shift : public mxnet_op::tunable { template MSHADOW_XINLINE static DType Map(DType a, DType b) { - if (b >= static_cast(sizeof(DType) * CHAR_BIT)) { + if (static_cast(b) >= (sizeof(DType) * CHAR_BIT)) { return DType(0); } return static_cast(a) << static_cast(b); @@ -845,7 +845,7 @@ MXNET_BINARY_MATH_OP(rbitwise_left_shift_grad, struct bitwise_right_shift : public mxnet_op::tunable { template MSHADOW_XINLINE static DType Map(DType a, DType b) { - if (b >= static_cast(sizeof(DType) * CHAR_BIT)) { + if (static_cast(b) >= (sizeof(DType) * CHAR_BIT)) { if (a < 0) { return DType(-1); } else { diff --git a/src/operator/tensor/elemwise_binary_broadcast_op.h b/src/operator/tensor/elemwise_binary_broadcast_op.h index 3485607f37e0..713e24932f8b 100644 --- a/src/operator/tensor/elemwise_binary_broadcast_op.h +++ b/src/operator/tensor/elemwise_binary_broadcast_op.h @@ -253,9 +253,6 @@ void BinaryBroadcastIntComputeWithBool(const nnvm::NodeAttrs& attrs, if (req[0] == kNullOp) return; mshadow::Stream* s = ctx.get_stream(); - if (outputs[0].type_flag_ == mshadow::kBool) { - LOG(FATAL) << "Operator " << attrs.op->name << " does not support boolean type"; - } MXNET_INT_TYPE_SWITCH_EXT_WITH_BOOL(outputs[0].type_flag_, DType, { BROADCAST_NDIM_SWITCH(ndim, NDim, { mshadow::Shape oshape = new_oshape.get(); diff --git a/src/operator/tensor/elemwise_binary_op.h b/src/operator/tensor/elemwise_binary_op.h index 1904bf9ef4f4..4f36b8acd404 100644 --- a/src/operator/tensor/elemwise_binary_op.h +++ b/src/operator/tensor/elemwise_binary_op.h @@ -502,7 +502,7 @@ class ElemwiseBinaryOp : public OpBase { LOG(FATAL) << "Operator " << attrs.op->name << " does not support boolean type"; } MXNET_ASSIGN_REQ_SWITCH(req[0], Req, { - MSHADOW_TYPE_SWITCH(outputs[0].type_flag_, DType, { + MSHADOW_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 9dfe283585df549af5e7568f0b7c937f23f3f958 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Thu, 4 Nov 2021 14:58:04 -0700 Subject: [PATCH 08/39] fix --- ci/docker/runtime_functions.sh | 2 +- src/operator/numpy/np_true_divide.cc | 4 ++++ src/operator/tensor/elemwise_unary_op.h | 2 +- 3 files changed, 6 insertions(+), 2 deletions(-) diff --git a/ci/docker/runtime_functions.sh b/ci/docker/runtime_functions.sh index eaff0f4b2c18..59d06e7c1416 100755 --- a/ci/docker/runtime_functions.sh +++ b/ci/docker/runtime_functions.sh @@ -874,7 +874,7 @@ unittest_array_api_standardization() { export DMLC_LOG_STACK_TRACE_DEPTH=100 python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_creation_functions.py python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_indexing.py - python3 -m pytest -s --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_elementwise_functions.py + python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_elementwise_functions.py python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_constants.py python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ array_api_tests/test_type_promotion.py::test_elementwise_function_two_arg_bool_type_promotion diff --git a/src/operator/numpy/np_true_divide.cc b/src/operator/numpy/np_true_divide.cc index 4b0b98131fa1..28758756a34c 100644 --- a/src/operator/numpy/np_true_divide.cc +++ b/src/operator/numpy/np_true_divide.cc @@ -74,6 +74,10 @@ NNVM_REGISTER_OP(_npi_true_divide) [](const NodeAttrs& attrs) { return std::vector >{{0, 0}, {1, 0}}; }) + .set_attr("FResourceRequest", + [](const NodeAttrs& attrs) { + return std::vector{ResourceRequest::kTempSpace}; + }) .set_attr("FCompute", TrueDivideBroadcastCompute) .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_broadcast_div"}) .add_argument("lhs", "NDArray-or-Symbol", "Dividend array") diff --git a/src/operator/tensor/elemwise_unary_op.h b/src/operator/tensor/elemwise_unary_op.h index 4fd9e4008967..f516a7858c62 100644 --- a/src/operator/tensor/elemwise_unary_op.h +++ b/src/operator/tensor/elemwise_unary_op.h @@ -243,7 +243,7 @@ class UnaryOp : public OpBase { const std::vector& inputs, const std::vector& req, const std::vector& outputs) { - MSHADOW_TYPE_SWITCH(outputs[0].type_flag_, DType, { + MSHADOW_TYPE_SWITCH_EXT(outputs[0].type_flag_, DType, { MXNET_ASSIGN_REQ_SWITCH(req[0], Req, { if (inputs[0].Size() != 0) { mxnet_op::Kernel, cpu>::Launch( From c02603e785e6de0c09992ab4f5667c26f7d59b05 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Thu, 4 Nov 2021 15:00:46 -0700 Subject: [PATCH 09/39] fix lint --- src/operator/numpy/np_true_divide.cc | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/src/operator/numpy/np_true_divide.cc b/src/operator/numpy/np_true_divide.cc index 28758756a34c..99cd4c4718bc 100644 --- a/src/operator/numpy/np_true_divide.cc +++ b/src/operator/numpy/np_true_divide.cc @@ -75,9 +75,9 @@ NNVM_REGISTER_OP(_npi_true_divide) return std::vector >{{0, 0}, {1, 0}}; }) .set_attr("FResourceRequest", - [](const NodeAttrs& attrs) { - return std::vector{ResourceRequest::kTempSpace}; - }) + [](const NodeAttrs& attrs) { + return std::vector{ResourceRequest::kTempSpace}; + }) .set_attr("FCompute", TrueDivideBroadcastCompute) .set_attr("FGradient", ElemwiseGradUseIn{"_backward_npi_broadcast_div"}) .add_argument("lhs", "NDArray-or-Symbol", "Dividend array") From 46d696b81dd1f8d8f4d4b6f5be4df0a36770dfa1 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Thu, 4 Nov 2021 18:44:09 -0700 Subject: [PATCH 10/39] fix --- python/mxnet/numpy/utils.py | 26 +++++++++++------------ src/operator/tensor/broadcast_reduce_op.h | 4 ++-- 2 files changed, 15 insertions(+), 15 deletions(-) diff --git a/python/mxnet/numpy/utils.py b/python/mxnet/numpy/utils.py index 15b83c7f2b73..35ea80a8e392 100644 --- a/python/mxnet/numpy/utils.py +++ b/python/mxnet/numpy/utils.py @@ -29,19 +29,19 @@ py_bool = bool -float16 = onp.float16 -float32 = onp.float32 -float64 = onp.float64 -uint8 = onp.uint8 -int32 = onp.int32 -int8 = onp.int8 -int64 = onp.int64 -bool_ = onp.bool_ -bool = onp.bool -int16 = onp.int16 -uint16 = onp.uint16 -uint32 = onp.uint32 -uint64 = onp.uint64 +float16 = onp.dtype(onp.float16) +float32 = onp.dtype(onp.float32) +float64 = onp.dtype(onp.float64) +uint8 = onp.dtype(onp.uint8) +int32 = onp.dtype(onp.int32) +int8 = onp.dtype(onp.int8) +int64 = onp.dtype(onp.int64) +bool_ = onp.dtype(onp.bool_) +bool = onp.dtype(onp.bool) +int16 = onp.dtype(onp.int16) +uint16 = onp.dtype(onp.uint16) +uint32 = onp.dtype(onp.uint32) +uint64 = onp.dtype(onp.uint64) pi = onp.pi inf = onp.inf diff --git a/src/operator/tensor/broadcast_reduce_op.h b/src/operator/tensor/broadcast_reduce_op.h index d5ba8c2f60c0..b8f2902444fa 100644 --- a/src/operator/tensor/broadcast_reduce_op.h +++ b/src/operator/tensor/broadcast_reduce_op.h @@ -1359,8 +1359,8 @@ inline void BroadcastComputeImpl(const nnvm::NodeAttrs& attrs, BroadcastReduceShapeCompact(outputs[0].shape_, small, &dst_shape, &src_shape); Stream* s = ctx.get_stream(); bool isCPU = std::is_same::value; - MSHADOW_TYPE_SWITCH_WITH_BOOL(inputs[0].type_flag_, IType, { - MSHADOW_TYPE_SWITCH_WITH_BOOL(outputs[0].type_flag_, OType, { + MSHADOW_TYPE_SWITCH_EXT_WITH_BOOL(inputs[0].type_flag_, IType, { + MSHADOW_TYPE_SWITCH_EXT_WITH_BOOL(outputs[0].type_flag_, OType, { mshadow::Shape in_shape; mshadow::Shape out_shape; for (int i = 0; i < MXNET_SPECIAL_MAX_NDIM; ++i) { From 7a89f563e43f234ab96b8f990fdcf60b2f223931 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Thu, 4 Nov 2021 20:15:03 -0700 Subject: [PATCH 11/39] update remainder --- src/operator/mshadow_op.h | 45 ++++++++++++++++++++++++++++++++------- 1 file changed, 37 insertions(+), 8 deletions(-) diff --git a/src/operator/mshadow_op.h b/src/operator/mshadow_op.h index 65796476cf9e..51b8eb59ad90 100644 --- a/src/operator/mshadow_op.h +++ b/src/operator/mshadow_op.h @@ -1011,10 +1011,15 @@ struct mod : public mxnet_op::tunable { MSHADOW_XINLINE static typename enable_if::value, DType>::type Map(DType a, DType b) { if (b == DType(0)) { + if (a < DType(0)) { + return -DType(0); + } return DType(0); } else if (b < DType(0)) { if (a < DType(0)) { return DType(-::fmod(-static_cast(a), -static_cast(b))); + } else if (a == DType(0)){ + return -DType(0); } else { return DType( ::fmod(static_cast(a), -static_cast(b)) + @@ -1022,8 +1027,8 @@ struct mod : public mxnet_op::tunable { } } else { if (a < DType(0)) { - return DType( - -::fmod(-static_cast(a), static_cast(b)) + + return -DType( + ::fmod(-static_cast(a), static_cast(b)) + (::fmod(-static_cast(a), static_cast(b)) != DType(0) ? b : DType(0))); } else { return DType(::fmod(static_cast(a), static_cast(b))); @@ -1044,7 +1049,11 @@ struct mod : public mxnet_op::tunable { struct mixed_mod { template ::value, int>::type = 0> MSHADOW_XINLINE static mshadow::half::half_t Map(DType a, mshadow::half::half_t b) { - return mod::Map(static_cast(a), b); + mshadow::half::half_t ret = mod::Map(static_cast(a), b); + if ((ret == 0) && ((a < 0) != (b < 0))) { + return -ret; + } + return ret; } template ::value, int>::type = 0> MSHADOW_XINLINE static float Map(DType a, float b) { - return mod::Map(static_cast(a), b); + float ret = mod::Map(static_cast(a), b); + if ((ret == 0) && ((a < 0) != (b < 0))) { + return -ret; + } + return ret; } template ::value, int>::type = 0> MSHADOW_XINLINE static double Map(DType a, double b) { - return mod::Map(static_cast(a), b); + double ret = mod::Map(static_cast(a), b); + if ((ret == 0) && ((a < 0) != (b < 0))) { + return -ret; + } + return ret; } }; struct mixed_rmod { template ::value, int>::type = 0> MSHADOW_XINLINE static mshadow::half::half_t Map(DType a, mshadow::half::half_t b) { - return mod::Map(b, static_cast(a)); + mshadow::half::half_t ret = mod::Map(b, static_cast(a)); + if ((ret == 0) && ((a < 0) != (b < 0))) { + return -ret; + } + return ret; } template ::value, int>::type = 0> MSHADOW_XINLINE static float Map(DType a, float b) { - return mod::Map(b, static_cast(a)); + float ret = mod::Map(b, static_cast(a)); + if ((ret == 0) && ((a < 0) != (b < 0))) { + return -ret; + } + return ret; } template ::value, int>::type = 0> MSHADOW_XINLINE static double Map(DType a, double b) { - return mod::Map(b, static_cast(a)); + double ret = mod::Map(b, static_cast(a)); + if ((ret == 0) && ((a < 0) != (b < 0))) { + return -ret; + } + return ret; } }; From 270e60d428ca76e4c7d060ef8b3a20716b996bc6 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Thu, 4 Nov 2021 20:22:41 -0700 Subject: [PATCH 12/39] fix lint --- 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 51b8eb59ad90..6aa70b410647 100644 --- a/src/operator/mshadow_op.h +++ b/src/operator/mshadow_op.h @@ -1018,7 +1018,7 @@ struct mod : public mxnet_op::tunable { } else if (b < DType(0)) { if (a < DType(0)) { return DType(-::fmod(-static_cast(a), -static_cast(b))); - } else if (a == DType(0)){ + } else if (a == DType(0)) { return -DType(0); } else { return DType( From fe9813b7db68a5eef846f5257b187c07cd550c66 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Fri, 5 Nov 2021 13:02:40 -0700 Subject: [PATCH 13/39] switch to no tvmop --- ci/jenkins/Jenkinsfile_unix_cpu | 3 +- src/operator/mshadow_op.h | 57 ++---- .../unittest/test_numpy_interoperability.py | 168 +++++++++--------- tests/python/unittest/test_numpy_ndarray.py | 12 +- 4 files changed, 110 insertions(+), 130 deletions(-) diff --git a/ci/jenkins/Jenkinsfile_unix_cpu b/ci/jenkins/Jenkinsfile_unix_cpu index 9681270d8905..22fc536592c2 100644 --- a/ci/jenkins/Jenkinsfile_unix_cpu +++ b/ci/jenkins/Jenkinsfile_unix_cpu @@ -46,7 +46,8 @@ core_logic: { utils.parallel_stage('Tests', [ custom_steps.test_unix_python3_cpu('cpu'), custom_steps.test_unix_python3_onnx_cpu('cpu'), - custom_steps.test_unix_python3_array_api('cpu'), + // TVMOP has issue with NAN, see https://github.com/apache/incubator-mxnet/issues/20729 + custom_steps.test_unix_python3_array_api('cpu_openblas_no_tvm_op'), custom_steps.test_unix_python3_mkl_cpu('cpu_mkl'), custom_steps.test_unix_python3_onednn_cpu('onednn_cpu'), custom_steps.test_unix_python3_onednn_mkl_cpu('onednn_mkl_cpu'), diff --git a/src/operator/mshadow_op.h b/src/operator/mshadow_op.h index 6aa70b410647..985975b3dcf6 100644 --- a/src/operator/mshadow_op.h +++ b/src/operator/mshadow_op.h @@ -1011,24 +1011,27 @@ struct mod : public mxnet_op::tunable { MSHADOW_XINLINE static typename enable_if::value, DType>::type Map(DType a, DType b) { if (b == DType(0)) { - if (a < DType(0)) { - return -DType(0); - } return DType(0); } else if (b < DType(0)) { if (a < DType(0)) { - return DType(-::fmod(-static_cast(a), -static_cast(b))); - } else if (a == DType(0)) { - return -DType(0); + DType ret = DType(-::fmod(-static_cast(a), -static_cast(b))); + if (ret == 0) { + return -ret; + } + return ret; } else { - return DType( + DType ret = DType( ::fmod(static_cast(a), -static_cast(b)) + (::fmod(static_cast(a), -static_cast(b)) != DType(0) ? b : DType(0))); + if (ret == 0) { + return -ret; + } + return ret; } } else { if (a < DType(0)) { - return -DType( - ::fmod(-static_cast(a), static_cast(b)) + + return DType( + -::fmod(-static_cast(a), static_cast(b)) + (::fmod(-static_cast(a), static_cast(b)) != DType(0) ? b : DType(0))); } else { return DType(::fmod(static_cast(a), static_cast(b))); @@ -1049,11 +1052,7 @@ struct mod : public mxnet_op::tunable { struct mixed_mod { template ::value, int>::type = 0> MSHADOW_XINLINE static mshadow::half::half_t Map(DType a, mshadow::half::half_t b) { - mshadow::half::half_t ret = mod::Map(static_cast(a), b); - if ((ret == 0) && ((a < 0) != (b < 0))) { - return -ret; - } - return ret; + return mod::Map(static_cast(a), b); } template ::value, int>::type = 0> MSHADOW_XINLINE static float Map(DType a, float b) { - float ret = mod::Map(static_cast(a), b); - if ((ret == 0) && ((a < 0) != (b < 0))) { - return -ret; - } - return ret; + return mod::Map(static_cast(a), b); } template ::value, int>::type = 0> MSHADOW_XINLINE static double Map(DType a, double b) { - double ret = mod::Map(static_cast(a), b); - if ((ret == 0) && ((a < 0) != (b < 0))) { - return -ret; - } - return ret; + return mod::Map(static_cast(a), b); } }; struct mixed_rmod { template ::value, int>::type = 0> MSHADOW_XINLINE static mshadow::half::half_t Map(DType a, mshadow::half::half_t b) { - mshadow::half::half_t ret = mod::Map(b, static_cast(a)); - if ((ret == 0) && ((a < 0) != (b < 0))) { - return -ret; - } - return ret; + return mod::Map(b, static_cast(a)); } template ::value, int>::type = 0> MSHADOW_XINLINE static float Map(DType a, float b) { - float ret = mod::Map(b, static_cast(a)); - if ((ret == 0) && ((a < 0) != (b < 0))) { - return -ret; - } - return ret; + return mod::Map(b, static_cast(a)); } template ::value, int>::type = 0> MSHADOW_XINLINE static double Map(DType a, double b) { - double ret = mod::Map(b, static_cast(a)); - if ((ret == 0) && ((a < 0) != (b < 0))) { - return -ret; - } - return ret; + return mod::Map(b, static_cast(a)); } }; diff --git a/tests/python/unittest/test_numpy_interoperability.py b/tests/python/unittest/test_numpy_interoperability.py index 6060f32a9587..b4dcf0b4f485 100644 --- a/tests/python/unittest/test_numpy_interoperability.py +++ b/tests/python/unittest/test_numpy_interoperability.py @@ -263,7 +263,7 @@ def _add_workload_percentile(): q3 = np.array([25, 50, 100]) q4 = 65 x4 = np.arange(11 * 2).reshape(11, 1, 2, 1) - x5 = np.array([0, np.nan]) + x5 = np.array([0, _np.nan]) OpArgMngr.add_workload('percentile', x1, q1, None, None, None) OpArgMngr.add_workload('percentile', x1, q1, None, None, None, 'linear') @@ -760,9 +760,9 @@ def _add_workload_tril(): [[1, 1], [0, 0]], ], dtype=dt) OpArgMngr.add_workload('tril', a) - arr = np.array([[1, 1, np.inf], + arr = np.array([[1, 1, _np.inf], [1, 1, 1], - [np.inf, 1, 1]]) + [_np.inf, 1, 1]]) OpArgMngr.add_workload('tril', arr) OpArgMngr.add_workload('tril', np.zeros((3, 3), dtype=dt)) import mxnet as mx @@ -780,9 +780,9 @@ def _add_workload_triu(): [[1, 1], [0, 0]], ], dtype=dt) OpArgMngr.add_workload('triu', a) - arr = np.array([[1, 1, np.inf], + arr = np.array([[1, 1, _np.inf], [1, 1, 1], - [np.inf, 1, 1]]) + [_np.inf, 1, 1]]) OpArgMngr.add_workload('triu', arr) OpArgMngr.add_workload('triu', np.zeros((3, 3), dtype=dt)) @@ -896,8 +896,8 @@ def _add_workload_einsum(): def _add_workload_expm1(): OpArgMngr.add_workload('expm1', np.random.uniform(size=(4, 1))) OpArgMngr.add_workload('expm1', np.random.uniform(size=(1, 1))) - OpArgMngr.add_workload('expm1', np.array([np.inf])) - OpArgMngr.add_workload('expm1', np.array([-np.inf])) + OpArgMngr.add_workload('expm1', np.array([_np.inf])) + OpArgMngr.add_workload('expm1', np.array([-_np.inf])) OpArgMngr.add_workload('expm1', np.array([0.])) OpArgMngr.add_workload('expm1', np.array([-0.])) @@ -908,10 +908,10 @@ def _add_workload_argmax(): OpArgMngr.add_workload('argmax', np.random.uniform(size=(4, 5, 6, 7, 8)), 2) OpArgMngr.add_workload('argmax', np.random.uniform(size=(4, 5, 6, 7, 8)), 3) OpArgMngr.add_workload('argmax', np.random.uniform(size=(4, 5, 6, 7, 8)), 4) - # OpArgMngr.add_workload('argmax', np.array([0, 1, 2, 3, np.nan])) - # OpArgMngr.add_workload('argmax', np.array([0, 1, 2, np.nan, 3])) - # OpArgMngr.add_workload('argmax', np.array([np.nan, 0, 1, 2, 3])) - # OpArgMngr.add_workload('argmax', np.array([np.nan, 0, np.nan, 2, 3])) + # OpArgMngr.add_workload('argmax', np.array([0, 1, 2, 3, _np.nan])) + # OpArgMngr.add_workload('argmax', np.array([0, 1, 2, _np.nan, 3])) + # OpArgMngr.add_workload('argmax', np.array([_np.nan, 0, 1, 2, 3])) + # OpArgMngr.add_workload('argmax', np.array([_np.nan, 0, _np.nan, 2, 3])) OpArgMngr.add_workload('argmax', np.array([False, False, False, False, True])) OpArgMngr.add_workload('argmax', np.array([False, False, False, True, False])) OpArgMngr.add_workload('argmax', np.array([True, False, False, False, False])) @@ -924,10 +924,10 @@ def _add_workload_argmin(): OpArgMngr.add_workload('argmin', np.random.uniform(size=(4, 5, 6, 7, 8)), 2) OpArgMngr.add_workload('argmin', np.random.uniform(size=(4, 5, 6, 7, 8)), 3) OpArgMngr.add_workload('argmin', np.random.uniform(size=(4, 5, 6, 7, 8)), 4) - # OpArgMngr.add_workload('argmin', np.array([0, 1, 2, 3, np.nan])) - # OpArgMngr.add_workload('argmin', np.array([0, 1, 2, np.nan, 3])) - # OpArgMngr.add_workload('argmin', np.array([np.nan, 0, 1, 2, 3])) - # OpArgMngr.add_workload('argmin', np.array([np.nan, 0, np.nan, 2, 3])) + # OpArgMngr.add_workload('argmin', np.array([0, 1, 2, 3, _np.nan])) + # OpArgMngr.add_workload('argmin', np.array([0, 1, 2, _np.nan, 3])) + # OpArgMngr.add_workload('argmin', np.array([_np.nan, 0, 1, 2, 3])) + # OpArgMngr.add_workload('argmin', np.array([_np.nan, 0, _np.nan, 2, 3])) OpArgMngr.add_workload('argmin', np.array([False, False, False, False, True])) OpArgMngr.add_workload('argmin', np.array([False, False, False, True, False])) OpArgMngr.add_workload('argmin', np.array([True, False, False, False, False])) @@ -1004,7 +1004,7 @@ def _add_workload_clip(): # OpArgMngr.add_workload('clip', np.array([0, 1, 2, 3, 4, 5, 6, 7]), 3) # OpArgMngr.add_workload('clip', np.array([0, 1, 2, 3, 4, 5, 6, 7]), a_min=3) # OpArgMngr.add_workload('clip', np.array([0, 1, 2, 3, 4, 5, 6, 7]), a_max=4) - OpArgMngr.add_workload('clip', np.array([-2., np.nan, 0.5, 3., 0.25, np.nan]), -1, 1) + OpArgMngr.add_workload('clip', np.array([-2., _np.nan, 0.5, 3., 0.25, _np.nan]), -1, 1) def _add_workload_cumsum(): @@ -1311,13 +1311,13 @@ def _add_workload_delete(): def _add_workload_var(array_pool): OpArgMngr.add_workload('var', array_pool['4x1']) - OpArgMngr.add_workload('var', np.array([np.float16(1.)])) + OpArgMngr.add_workload('var', np.array([_np.float16(1.)])) OpArgMngr.add_workload('var', np.array([1])) OpArgMngr.add_workload('var', np.array([1.])) OpArgMngr.add_workload('var', np.array([[1, 2, 3], [4, 5, 6]])) OpArgMngr.add_workload('var', np.array([[1, 2, 3], [4, 5, 6]]), 0) OpArgMngr.add_workload('var', np.array([[1, 2, 3], [4, 5, 6]]), 1) - OpArgMngr.add_workload('var', np.array([np.nan])) + OpArgMngr.add_workload('var', np.array([_np.nan])) OpArgMngr.add_workload('var', np.array([1, -1, 1, -1])) OpArgMngr.add_workload('var', np.array([1,2,3,4], dtype='f8')) @@ -1333,7 +1333,7 @@ def _add_workload_full_like(array_pool): OpArgMngr.add_workload('full_like', array_pool['4x1'], 1) OpArgMngr.add_workload('full_like', np.random.uniform(low=0, high=100, size=(1,3,4), dtype='float64'), 1) OpArgMngr.add_workload('full_like', np.random.uniform(low=0, high=100, size=(9,3,1)), 2, dtype=np.int64) - OpArgMngr.add_workload('full_like', np.random.uniform(low=0, high=100, size=(9,3)), np.nan) + OpArgMngr.add_workload('full_like', np.random.uniform(low=0, high=100, size=(9,3)), _np.nan) OpArgMngr.add_workload('full_like', np.random.uniform(low=0, high=100, size=(2,0)), 0, dtype=np.float32) @@ -1357,13 +1357,13 @@ def _add_workload_meshgrid(): def _add_workload_abs(): OpArgMngr.add_workload('abs', np.random.uniform(size=(11,)).astype(np.float32)) OpArgMngr.add_workload('abs', np.random.uniform(size=(5,)).astype(np.float64)) - OpArgMngr.add_workload('abs', np.array([np.inf, -np.inf, np.nan])) + OpArgMngr.add_workload('abs', np.array([_np.inf, -_np.inf, _np.nan])) def _add_workload_fabs(): OpArgMngr.add_workload('fabs', np.random.uniform(size=(11,)).astype(np.float32)) OpArgMngr.add_workload('fabs', np.random.uniform(size=(5,)).astype(np.float64)) - OpArgMngr.add_workload('fabs', np.array([np.inf, -np.inf, np.nan])) + OpArgMngr.add_workload('fabs', np.array([_np.inf, -_np.inf, _np.nan])) def _add_workload_add(array_pool): @@ -1381,10 +1381,10 @@ def _add_workload_arctan2(): OpArgMngr.add_workload('arctan2', np.array([np.PZERO, np.NZERO]), np.array([1, 1])) OpArgMngr.add_workload('arctan2', np.array([-1, -1]), np.array([np.PZERO, np.NZERO])) OpArgMngr.add_workload('arctan2', np.array([1, 1]), np.array([np.PZERO, np.NZERO])) - OpArgMngr.add_workload('arctan2', np.array([1, -1, 1, -1]), np.array([-np.inf, -np.inf, np.inf, np.inf])) - OpArgMngr.add_workload('arctan2', np.array([np.inf, -np.inf]), np.array([1, 1])) - OpArgMngr.add_workload('arctan2', np.array([np.inf, -np.inf]), np.array([-np.inf, -np.inf])) - OpArgMngr.add_workload('arctan2', np.array([np.inf, -np.inf]), np.array([np.inf, np.inf])) + OpArgMngr.add_workload('arctan2', np.array([1, -1, 1, -1]), np.array([-_np.inf, -_np.inf, _np.inf, _np.inf])) + OpArgMngr.add_workload('arctan2', np.array([_np.inf, -_np.inf]), np.array([1, 1])) + OpArgMngr.add_workload('arctan2', np.array([_np.inf, -_np.inf]), np.array([-_np.inf, -_np.inf])) + OpArgMngr.add_workload('arctan2', np.array([_np.inf, -_np.inf]), np.array([_np.inf, _np.inf])) def _add_workload_copysign(): @@ -1442,7 +1442,7 @@ def _add_workload_interp(): fp0 = np.linspace(0, 1, 5) x0 = np.linspace(0, 1, 50) xp1 = np.array([1, 2, 3, 4]) - fp1 = np.array([1, 2, np.inf, 4]) + fp1 = np.array([1, 2, _np.inf, 4]) x1 = np.array([1, 2, 2.5, 3, 4]) xp2 = np.arange(0, 10, 0.0001) fp2 = np.sin(xp2) @@ -1472,14 +1472,14 @@ def _add_workload_interp(): def _add_workload_hypot(): OpArgMngr.add_workload('hypot', np.array(1), np.array(1)) OpArgMngr.add_workload('hypot', np.array(0), np.array(0)) - OpArgMngr.add_workload('hypot', np.array(np.nan), np.array(np.nan)) - OpArgMngr.add_workload('hypot', np.array(np.nan), np.array(1)) - OpArgMngr.add_workload('hypot', np.array(np.nan), np.array(np.inf)) - OpArgMngr.add_workload('hypot', np.array(np.inf), np.array(np.nan)) - OpArgMngr.add_workload('hypot', np.array(np.inf), np.array(0)) - OpArgMngr.add_workload('hypot', np.array(0), np.array(np.inf)) - OpArgMngr.add_workload('hypot', np.array(np.inf), np.array(np.inf)) - OpArgMngr.add_workload('hypot', np.array(np.inf), np.array(23.0)) + OpArgMngr.add_workload('hypot', np.array(_np.nan), np.array(_np.nan)) + OpArgMngr.add_workload('hypot', np.array(_np.nan), np.array(1)) + OpArgMngr.add_workload('hypot', np.array(_np.nan), np.array(_np.inf)) + OpArgMngr.add_workload('hypot', np.array(_np.inf), np.array(_np.nan)) + OpArgMngr.add_workload('hypot', np.array(_np.inf), np.array(0)) + OpArgMngr.add_workload('hypot', np.array(0), np.array(_np.inf)) + OpArgMngr.add_workload('hypot', np.array(_np.inf), np.array(_np.inf)) + OpArgMngr.add_workload('hypot', np.array(_np.inf), np.array(23.0)) def _add_workload_lcm(): @@ -1673,8 +1673,8 @@ def _signs(dt): for ct in [np.float16, np.float32, np.float64]: fone = np.array(1.0, dtype=ct) fzer = np.array(0.0, dtype=ct) - finf = np.array(np.inf, dtype=ct) - fnan = np.array(np.nan, dtype=ct) + finf = np.array(_np.inf, dtype=ct) + fnan = np.array(_np.nan, dtype=ct) # OpArgMngr.add_workload('remainder', fone, fzer) # failed OpArgMngr.add_workload('remainder', fone, fnan) OpArgMngr.add_workload('remainder', finf, fone) @@ -1734,13 +1734,13 @@ def _add_workload_log(array_pool): def _add_workload_log2(array_pool): OpArgMngr.add_workload('log2', array_pool['4x1']) OpArgMngr.add_workload('log2', np.array(2.**65)) - OpArgMngr.add_workload('log2', np.array(np.inf)) + OpArgMngr.add_workload('log2', np.array(_np.inf)) OpArgMngr.add_workload('log2', np.array(1.)) def _add_workload_log1p(): OpArgMngr.add_workload('log1p', np.array(-1.)) - OpArgMngr.add_workload('log1p', np.array(np.inf)) + OpArgMngr.add_workload('log1p', np.array(_np.inf)) OpArgMngr.add_workload('log1p', np.array(1e-6)) @@ -1749,7 +1749,7 @@ def _add_workload_log10(array_pool): def _add_workload_sqrt(): - OpArgMngr.add_workload('sqrt', np.array([1, np.PZERO, np.NZERO, np.inf, np.nan])) + OpArgMngr.add_workload('sqrt', np.array([1, np.PZERO, np.NZERO, _np.inf, _np.nan])) def _add_workload_square(): @@ -1758,8 +1758,8 @@ def _add_workload_square(): def _add_workload_cbrt(): OpArgMngr.add_workload('cbrt', np.array(-2.5**3, dtype=np.float32)) - OpArgMngr.add_workload('cbrt', np.array([1., 2., -3., np.inf, -np.inf])**3) - OpArgMngr.add_workload('cbrt', np.array([np.inf, -np.inf, np.nan])) + OpArgMngr.add_workload('cbrt', np.array([1., 2., -3., _np.inf, -_np.inf])**3) + OpArgMngr.add_workload('cbrt', np.array([_np.inf, -_np.inf, _np.nan])) def _add_workload_reciprocal(): @@ -1983,8 +1983,8 @@ def _add_workload_equal(array_pool): # TODO(junwu): fp16 does not work yet with TVM generated ops # OpArgMngr.add_workload('equal', np.array([0, 1, 2, 4, 2], dtype=np.float16), np.array([-2, 5, 1, 4, 3], dtype=np.float16)) OpArgMngr.add_workload('equal', np.array([0, 1, 2, 4, 2], dtype=np.float32), np.array([-2, 5, 1, 4, 3], dtype=np.float32)) - # TODO(junwu): mxnet currently does not have a consistent behavior as NumPy in dealing with np.nan - # OpArgMngr.add_workload('equal', np.array([np.nan]), np.array([np.nan])) + # TODO(junwu): mxnet currently does not have a consistent behavior as NumPy in dealing with _np.nan + # OpArgMngr.add_workload('equal', np.array([_np.nan]), np.array([_np.nan])) OpArgMngr.add_workload('equal', array_pool['4x1'], array_pool['1x2']) @@ -1992,8 +1992,8 @@ def _add_workload_not_equal(array_pool): # TODO(junwu): fp16 does not work yet with TVM generated ops # OpArgMngr.add_workload('not_equal', np.array([0, 1, 2, 4, 2], dtype=np.float16), np.array([-2, 5, 1, 4, 3], dtype=np.float16)) OpArgMngr.add_workload('not_equal', np.array([0, 1, 2, 4, 2], dtype=np.float32), np.array([-2, 5, 1, 4, 3], dtype=np.float32)) - # TODO(junwu): mxnet currently does not have a consistent behavior as NumPy in dealing with np.nan - # OpArgMngr.add_workload('not_equal', np.array([np.nan]), np.array([np.nan])) + # TODO(junwu): mxnet currently does not have a consistent behavior as NumPy in dealing with _np.nan + # OpArgMngr.add_workload('not_equal', np.array([_np.nan]), np.array([_np.nan])) OpArgMngr.add_workload('not_equal', array_pool['4x1'], array_pool['1x2']) @@ -2004,8 +2004,8 @@ def _add_workload_greater(array_pool): OpArgMngr.add_workload('greater', array_pool['4x1'], array_pool['1x2']) OpArgMngr.add_workload('greater', array_pool['4x1'], 2) OpArgMngr.add_workload('greater', 2, array_pool['4x1']) - # TODO(junwu): mxnet currently does not have a consistent behavior as NumPy in dealing with np.nan - # OpArgMngr.add_workload('greater', np.array([np.nan]), np.array([np.nan])) + # TODO(junwu): mxnet currently does not have a consistent behavior as NumPy in dealing with _np.nan + # OpArgMngr.add_workload('greater', np.array([_np.nan]), np.array([_np.nan])) def _add_workload_greater_equal(array_pool): @@ -2015,8 +2015,8 @@ def _add_workload_greater_equal(array_pool): OpArgMngr.add_workload('greater_equal', array_pool['4x1'], array_pool['1x2']) OpArgMngr.add_workload('greater_equal', array_pool['4x1'], 2) OpArgMngr.add_workload('greater_equal', 2, array_pool['4x1']) - # TODO(junwu): mxnet currently does not have a consistent behavior as NumPy in dealing with np.nan - # OpArgMngr.add_workload('greater_equal', np.array([np.nan]), np.array([np.nan])) + # TODO(junwu): mxnet currently does not have a consistent behavior as NumPy in dealing with _np.nan + # OpArgMngr.add_workload('greater_equal', np.array([_np.nan]), np.array([_np.nan])) def _add_workload_less(array_pool): @@ -2026,8 +2026,8 @@ def _add_workload_less(array_pool): OpArgMngr.add_workload('less', array_pool['4x1'], array_pool['1x2']) OpArgMngr.add_workload('less', array_pool['4x1'], 2) OpArgMngr.add_workload('less', 2, array_pool['4x1']) - # TODO(junwu): mxnet currently does not have a consistent behavior as NumPy in dealing with np.nan - # OpArgMngr.add_workload('less', np.array([np.nan]), np.array([np.nan])) + # TODO(junwu): mxnet currently does not have a consistent behavior as NumPy in dealing with _np.nan + # OpArgMngr.add_workload('less', np.array([_np.nan]), np.array([_np.nan])) def _add_workload_less_equal(array_pool): @@ -2037,8 +2037,8 @@ def _add_workload_less_equal(array_pool): OpArgMngr.add_workload('less_equal', array_pool['4x1'], array_pool['1x2']) OpArgMngr.add_workload('less_equal', array_pool['4x1'], 2) OpArgMngr.add_workload('less_equal', 2, array_pool['4x1']) - # TODO(junwu): mxnet currently does not have a consistent behavior as NumPy in dealing with np.nan - # OpArgMngr.add_workload('less_equal', np.array([np.nan]), np.array([np.nan])) + # TODO(junwu): mxnet currently does not have a consistent behavior as NumPy in dealing with _np.nan + # OpArgMngr.add_workload('less_equal', np.array([_np.nan]), np.array([_np.nan])) def _add_workload_logical_and(array_pool): @@ -2240,8 +2240,8 @@ def _add_workload_polyval(): def _add_workload_linalg_cond(): A = np.array([[1., 0, 1], [0, -2., 0], [0, 0, 3.]]) - OpArgMngr.add_workload('linalg.cond', A, np.inf) - OpArgMngr.add_workload('linalg.cond', A, -np.inf) + OpArgMngr.add_workload('linalg.cond', A, _np.inf) + OpArgMngr.add_workload('linalg.cond', A, -_np.inf) OpArgMngr.add_workload('linalg.cond', A, 1) OpArgMngr.add_workload('linalg.cond', A, -1) OpArgMngr.add_workload('linalg.cond', A, 'fro') @@ -2286,22 +2286,22 @@ def _add_workload_linalg_multi_dot(): def _add_workload_heaviside(): - x = np.array([[-30.0, -0.1, 0.0, 0.2], [7.5, np.nan, np.inf, -np.inf]], dtype=np.float64) + x = np.array([[-30.0, -0.1, 0.0, 0.2], [7.5, _np.nan, _np.inf, -_np.inf]], dtype=np.float64) OpArgMngr.add_workload('heaviside', x, 0.5) OpArgMngr.add_workload('heaviside', x, 1.0) x = x.astype(np.float32) - OpArgMngr.add_workload('heaviside', x, np.float32(0.5)) - OpArgMngr.add_workload('heaviside', x, np.float32(1.0)) + OpArgMngr.add_workload('heaviside', x, _np.float32(0.5)) + OpArgMngr.add_workload('heaviside', x, _np.float32(1.0)) def _add_workload_spacing(): - OpArgMngr.add_workload('spacing', np.float64(1)) - OpArgMngr.add_workload('spacing', np.float32(1)) - OpArgMngr.add_workload('spacing', np.inf) - OpArgMngr.add_workload('spacing', -np.inf) - OpArgMngr.add_workload('spacing', np.float64(1e30)) - OpArgMngr.add_workload('spacing', np.float32(1e30)) + OpArgMngr.add_workload('spacing', _np.float64(1)) + OpArgMngr.add_workload('spacing', _np.float32(1)) + OpArgMngr.add_workload('spacing', _np.inf) + OpArgMngr.add_workload('spacing', -_np.inf) + OpArgMngr.add_workload('spacing', _np.float64(1e30)) + OpArgMngr.add_workload('spacing', _np.float32(1e30)) def _add_workload_allclose(): @@ -2548,14 +2548,14 @@ def _add_workload_interp(): x0 = np.linspace(0, 1, 50) x1 = 0 x2 = .3 - x3 = np.float32(.3) + x3 = _np.float32(.3) OpArgMngr.add_workload('interp', x0, x, y) OpArgMngr.add_workload('interp', x1, x, y) OpArgMngr.add_workload('interp', x2, x, y) OpArgMngr.add_workload('interp', x3, x, y) x = np.array([1, 2, 2.5, 3, 4]) xp = np.array([1, 2, 3, 4]) - fp = np.array([1, 2, np.inf, 4]) + fp = np.array([1, 2, _np.inf, 4]) OpArgMngr.add_workload('interp', x, xp, fp) @@ -2574,7 +2574,7 @@ def _add_workload_intersect1d(): def _add_workload_isclose(): a = np.array([1e10,1e-7]) b = np.array([1.00001e10,1e-8]) - c = np.array([1.0, np.nan]) + c = np.array([1.0, _np.nan]) d = np.array([0.0, 0.0]) e = np.array([1e-100, 1e-7]) OpArgMngr.add_workload('isclose', a, b) @@ -2633,56 +2633,56 @@ def _add_workload_msort(): def _add_workload_nanargmax(): - a = np.array([[np.nan, 4], [2, 3]]) + a = np.array([[_np.nan, 4], [2, 3]]) OpArgMngr.add_workload('nanargmax', a) OpArgMngr.add_workload('nanargmax', a, axis=0) OpArgMngr.add_workload('nanargmax', a, axis=1) def _add_workload_nanargmin(): - a = np.array([[np.nan, 4], [2, 3]]) + a = np.array([[_np.nan, 4], [2, 3]]) OpArgMngr.add_workload('nanargmin', a) OpArgMngr.add_workload('nanargmin', a, axis=0) OpArgMngr.add_workload('nanargmin', a, axis=1) def _add_workload_nancumprod(): - a = np.array([[1, 2], [3, np.nan]]) + a = np.array([[1, 2], [3, _np.nan]]) OpArgMngr.add_workload('nancumprod', a) OpArgMngr.add_workload('nancumprod', a, axis=0) OpArgMngr.add_workload('nancumprod', a, axis=1) def _add_workload_nancumsum(): - a = np.array([[1, 2], [3, np.nan]]) + a = np.array([[1, 2], [3, _np.nan]]) OpArgMngr.add_workload('nancumsum', a) OpArgMngr.add_workload('nancumsum', a, axis=0) OpArgMngr.add_workload('nancumsum', a, axis=1) def _add_workload_nanmax(): - a = np.array([[1, 2], [3, np.nan]]) + a = np.array([[1, 2], [3, _np.nan]]) OpArgMngr.add_workload('nanmax', a) OpArgMngr.add_workload('nanmax', a, axis=0) OpArgMngr.add_workload('nanmax', a, axis=1) def _add_workload_nanmedian(): - a = np.array([[10.0, np.nan, 4], [3, 2, 1]]) + a = np.array([[10.0, _np.nan, 4], [3, 2, 1]]) OpArgMngr.add_workload('nanmedian', a) OpArgMngr.add_workload('nanmedian', a, axis=0) OpArgMngr.add_workload('nanmedian', a, axis=1) def _add_workload_nanmin(): - a = np.array([[1, 2], [3, np.nan]]) + a = np.array([[1, 2], [3, _np.nan]]) OpArgMngr.add_workload('nanmin', a) OpArgMngr.add_workload('nanmin', a, axis=0) OpArgMngr.add_workload('nanmin', a, axis=1) def _add_workload_nanpercentile(): - a = np.array([[10.0, np.nan, 4], [3, 2, 1]]) + a = np.array([[10.0, _np.nan, 4], [3, 2, 1]]) OpArgMngr.add_workload('nanpercentile', a, 50) OpArgMngr.add_workload('nanpercentile', a, 50, axis=0) OpArgMngr.add_workload('nanpercentile', a, 50, axis=1) @@ -2695,8 +2695,8 @@ def _add_workload_nanpercentile(): def _add_workload_nanprod(): a = 1 - b = np.array([1, np.nan]) - c = np.array([[1, 2], [3, np.nan]]) + b = np.array([1, _np.nan]) + c = np.array([[1, 2], [3, _np.nan]]) OpArgMngr.add_workload('nanprod', a) OpArgMngr.add_workload('nanprod', b) OpArgMngr.add_workload('nanprod', c) @@ -2704,7 +2704,7 @@ def _add_workload_nanprod(): def _add_workload_nanquantile(): - a = np.array([[10.0, np.nan, 4], [3, 2, 1]]) + a = np.array([[10.0, _np.nan, 4], [3, 2, 1]]) OpArgMngr.add_workload('nanquantile', a, 0.4) OpArgMngr.add_workload('nanquantile', a, 0.4, axis=0) OpArgMngr.add_workload('nanquantile', a, 0.4, axis=1) @@ -2717,7 +2717,7 @@ def _add_workload_nanquantile(): def _add_workload_nanstd(): OpArgMngr.add_workload('nanstd', np.random.uniform(size=(4, 1))) - A = np.array([[1, 2, 3], [4, np.nan, 6]]) + A = np.array([[1, 2, 3], [4, _np.nan, 6]]) OpArgMngr.add_workload('nanstd', A) OpArgMngr.add_workload('nanstd', A, 0) OpArgMngr.add_workload('nanstd', A, 1) @@ -2729,8 +2729,8 @@ def _add_workload_nanstd(): def _add_workload_nansum(): a = 1 - b = np.array([1, np.nan]) - c = np.array([[1, 2], [3, np.nan]]) + b = np.array([1, _np.nan]) + c = np.array([[1, 2], [3, _np.nan]]) OpArgMngr.add_workload('nansum', a) OpArgMngr.add_workload('nansum', b) OpArgMngr.add_workload('nansum', c) @@ -2739,7 +2739,7 @@ def _add_workload_nansum(): def _add_workload_nanvar(): OpArgMngr.add_workload('nanvar', np.random.uniform(size=(4, 1))) - A = np.array([[1, 2, 3], [4, np.nan, 6]]) + A = np.array([[1, 2, 3], [4, _np.nan, 6]]) OpArgMngr.add_workload('nanvar', A) OpArgMngr.add_workload('nanvar', A, 0) OpArgMngr.add_workload('nanvar', A, 1) @@ -2960,9 +2960,9 @@ def _add_workload_trapz(): def _add_workload_tril_indices_from(): for dt in ['float16', 'float32', 'float64', 'int32', 'int64', 'int8', 'uint8']: OpArgMngr.add_workload('tril_indices_from', np.ones((2, 2), dtype=dt)) - arr = np.array([[1, 1, np.inf], + arr = np.array([[1, 1, _np.inf], [1, 1, 1], - [np.inf, 1, 1]]) + [_np.inf, 1, 1]]) OpArgMngr.add_workload('tril_indices_from', arr) OpArgMngr.add_workload('tril_indices_from', np.zeros((3, 3), dtype=dt)) diff --git a/tests/python/unittest/test_numpy_ndarray.py b/tests/python/unittest/test_numpy_ndarray.py index 8558c3d561e7..274ad9510113 100644 --- a/tests/python/unittest/test_numpy_ndarray.py +++ b/tests/python/unittest/test_numpy_ndarray.py @@ -801,18 +801,18 @@ def test_setitem_autograd(np_array, index): # Basic indexing # Single int as index 0, - np.int32(0), - np.int64(0), + _np.int32(0), + _np.int64(0), np.array(0, dtype='int32'), np.array(0, dtype='int64'), 5, - np.int32(5), - np.int64(5), + _np.int32(5), + _np.int64(5), np.array(5, dtype='int32'), np.array(5, dtype='int64'), -1, - np.int32(-1), - np.int64(-1), + _np.int32(-1), + _np.int64(-1), np.array(-1, dtype='int32'), np.array(-1, dtype='int64'), # Slicing as index From 291896e0012fa029e9c7b720d57ba7e710792246 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Fri, 5 Nov 2021 16:22:19 -0700 Subject: [PATCH 14/39] fix tests --- ci/jenkins/Jenkins_steps.groovy | 2 +- tests/python/unittest/test_numpy_ndarray.py | 54 ++++++++++----------- 2 files changed, 28 insertions(+), 28 deletions(-) diff --git a/ci/jenkins/Jenkins_steps.groovy b/ci/jenkins/Jenkins_steps.groovy index 69c6a88643ab..e6f40806e273 100644 --- a/ci/jenkins/Jenkins_steps.groovy +++ b/ci/jenkins/Jenkins_steps.groovy @@ -675,7 +675,7 @@ def test_unix_python3_array_api(lib_name) { return ['Python3: Array-API': { node(NODE_LINUX_CPU) { ws('workspace/ut-python3-cpu') { - utils.unpack_and_init(lib_name, mx_lib, true) + utils.unpack_and_init(lib_name, mx_lib, false) python3_ut_array_api('ubuntu_cpu') utils.publish_test_coverage() } diff --git a/tests/python/unittest/test_numpy_ndarray.py b/tests/python/unittest/test_numpy_ndarray.py index 274ad9510113..edec96c34ed8 100644 --- a/tests/python/unittest/test_numpy_ndarray.py +++ b/tests/python/unittest/test_numpy_ndarray.py @@ -622,7 +622,7 @@ def test_nd_no_format(): @use_np @pytest.mark.serial def test_np_ndarray_indexing(): - def np_int(index, int_type=np.int32): + def np_int(index, int_type=_np.int32): """ Helper function for testing indexing that converts slices to slices of ints or None, and tuples to tuples of ints or None. @@ -817,54 +817,54 @@ def test_setitem_autograd(np_array, index): np.array(-1, dtype='int64'), # Slicing as index slice(5), - np_int(slice(5), np.int32), - np_int(slice(5), np.int64), + np_int(slice(5), _np.int32), + np_int(slice(5), _np.int64), slice(1, 5), - np_int(slice(1, 5), np.int32), - np_int(slice(1, 5), np.int64), + np_int(slice(1, 5), _np.int32), + np_int(slice(1, 5), _np.int64), slice(1, 5, 2), slice(1, 2, 2), - np_int(slice(1, 5, 2), np.int32), - np_int(slice(1, 5, 2), np.int64), + np_int(slice(1, 5, 2), _np.int32), + np_int(slice(1, 5, 2), _np.int64), slice(7, 0, -1), np_int(slice(7, 0, -1)), - np_int(slice(7, 0, -1), np.int64), + np_int(slice(7, 0, -1), _np.int64), slice(None, 6), np_int(slice(None, 6)), - np_int(slice(None, 6), np.int64), + np_int(slice(None, 6), _np.int64), slice(None, 6, 3), np_int(slice(None, 6, 3)), - np_int(slice(None, 6, 3), np.int64), + np_int(slice(None, 6, 3), _np.int64), slice(1, None), np_int(slice(1, None)), - np_int(slice(1, None), np.int64), + np_int(slice(1, None), _np.int64), slice(1, None, 3), np_int(slice(1, None, 3)), - np_int(slice(1, None, 3), np.int64), + np_int(slice(1, None, 3), _np.int64), slice(None, None, 2), np_int(slice(None, None, 2)), - np_int(slice(None, None, 2), np.int64), + np_int(slice(None, None, 2), _np.int64), slice(None, None, -1), np_int(slice(None, None, -1)), - np_int(slice(None, None, -1), np.int64), + np_int(slice(None, None, -1), _np.int64), slice(None, None, -2), - np_int(slice(None, None, -2), np.int32), - np_int(slice(None, None, -2), np.int64), + np_int(slice(None, None, -2), _np.int32), + np_int(slice(None, None, -2), _np.int64), # Multiple ints as indices (1, 2, 3), np_int((1, 2, 3)), - np_int((1, 2, 3), np.int64), + np_int((1, 2, 3), _np.int64), (-1, -2, -3), np_int((-1, -2, -3)), - np_int((-1, -2, -3), np.int64), + np_int((-1, -2, -3), _np.int64), (1, 2, 3, 4), np_int((1, 2, 3, 4)), - np_int((1, 2, 3, 4), np.int64), + np_int((1, 2, 3, 4), _np.int64), (-4, -3, -2, -1), (-4, mx.np.array(-3, dtype='int32'), -2, -1), (-4, mx.np.array(-3, dtype='int64'), -2, -1), np_int((-4, -3, -2, -1)), - np_int((-4, -3, -2, -1), np.int64), + np_int((-4, -3, -2, -1), _np.int64), # slice(None) as indices (slice(None), slice(None), 1, 8), (slice(None), slice(None), np.array(1, dtype='int32'), 8), @@ -873,26 +873,26 @@ def test_setitem_autograd(np_array, index): (slice(None), slice(None), 1, -8), (slice(None), slice(None), -1, -8), np_int((slice(None), slice(None), 1, 8)), - np_int((slice(None), slice(None), 1, 8), np.int64), + np_int((slice(None), slice(None), 1, 8), _np.int64), (slice(None), slice(None), 1, 8), np_int((slice(None), slice(None), -1, -8)), - np_int((slice(None), slice(None), -1, -8), np.int64), + np_int((slice(None), slice(None), -1, -8), _np.int64), (slice(None), 2, slice(1, 5), 1), np_int((slice(None), 2, slice(1, 5), 1)), - np_int((slice(None), 2, slice(1, 5), 1), np.int64), + np_int((slice(None), 2, slice(1, 5), 1), _np.int64), # Mixture of ints and slices as indices (slice(None, None, -1), 2, slice(1, 5), 1), np_int((slice(None, None, -1), 2, slice(1, 5), 1)), - np_int((slice(None, None, -1), 2, slice(1, 5), 1), np.int64), + np_int((slice(None, None, -1), 2, slice(1, 5), 1), _np.int64), (slice(None, None, -1), 2, slice(1, 7, 2), 1), np_int((slice(None, None, -1), 2, slice(1, 7, 2), 1)), - np_int((slice(None, None, -1), 2, slice(1, 7, 2), 1), np.int64), + np_int((slice(None, None, -1), 2, slice(1, 7, 2), 1), _np.int64), (slice(1, 8, 2), slice(14, 2, -2), slice(3, 8), slice(0, 7, 3)), np_int((slice(1, 8, 2), slice(14, 2, -2), slice(3, 8), slice(0, 7, 3))), - np_int((slice(1, 8, 2), slice(14, 2, -2), slice(3, 8), slice(0, 7, 3)), np.int64), + np_int((slice(1, 8, 2), slice(14, 2, -2), slice(3, 8), slice(0, 7, 3)), _np.int64), (slice(1, 8, 2), 1, slice(3, 8), 2), np_int((slice(1, 8, 2), 1, slice(3, 8), 2)), - np_int((slice(1, 8, 2), 1, slice(3, 8), 2), np.int64), + np_int((slice(1, 8, 2), 1, slice(3, 8), 2), _np.int64), # Test Ellipsis ('...') (1, Ellipsis, -1), (slice(2), Ellipsis, None, 0), From 92332f4f672f56b270e2fcace49e5b0fca68b6e7 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Sat, 6 Nov 2021 12:16:34 -0700 Subject: [PATCH 15/39] fix elemwise binary --- src/operator/tensor/elemwise_binary_broadcast_op.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/operator/tensor/elemwise_binary_broadcast_op.h b/src/operator/tensor/elemwise_binary_broadcast_op.h index 713e24932f8b..20d874dbd826 100644 --- a/src/operator/tensor/elemwise_binary_broadcast_op.h +++ b/src/operator/tensor/elemwise_binary_broadcast_op.h @@ -293,7 +293,7 @@ void BinaryBroadcastCompute(const nnvm::NodeAttrs& attrs, if (outputs[0].type_flag_ == mshadow::kBool) { LOG(FATAL) << "Operator " << attrs.op->name << " does not support boolean type"; } - MSHADOW_TYPE_SWITCH(outputs[0].type_flag_, DType, { + MSHADOW_TYPE_SWITCH_EXT(outputs[0].type_flag_, DType, { BROADCAST_NDIM_SWITCH(ndim, NDim, { broadcast::BinaryBroadcastComputeImpl(s, req[0], From 3a11d157d007da1ae8057f9a867bc9f0fb221ef2 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Mon, 8 Nov 2021 11:12:53 -0800 Subject: [PATCH 16/39] update asarray --- .github/workflows/os_x_staticbuild.yml | 1 + ci/docker/runtime_functions.sh | 1 + python/mxnet/numpy/multiarray.py | 44 ++++++++++++++++++++------ 3 files changed, 36 insertions(+), 10 deletions(-) diff --git a/.github/workflows/os_x_staticbuild.yml b/.github/workflows/os_x_staticbuild.yml index bc67b31e6830..c6aad0676d0b 100644 --- a/.github/workflows/os_x_staticbuild.yml +++ b/.github/workflows/os_x_staticbuild.yml @@ -69,6 +69,7 @@ jobs: python3 -m pytest --durations=50 --verbose array_api_tests/test_indexing.py python3 -m pytest --durations=50 --verbose array_api_tests/test_constants.py python3 -m pytest --durations=50 --verbose array_api_tests/test_elementwise_functions.py + python3 -m pytest --durations=50 --verbose array_api_tests/test_broadcasting.py python3 -m pytest --durations=50 --verbose \ array_api_tests/test_type_promotion.py::test_elementwise_function_two_arg_bool_type_promotion python3 -m pytest --durations=50 --verbose \ diff --git a/ci/docker/runtime_functions.sh b/ci/docker/runtime_functions.sh index 59d06e7c1416..c731c0d62bef 100755 --- a/ci/docker/runtime_functions.sh +++ b/ci/docker/runtime_functions.sh @@ -876,6 +876,7 @@ unittest_array_api_standardization() { python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_indexing.py python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_elementwise_functions.py python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_constants.py + python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_broadcasting.py python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ array_api_tests/test_type_promotion.py::test_elementwise_function_two_arg_bool_type_promotion python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ diff --git a/python/mxnet/numpy/multiarray.py b/python/mxnet/numpy/multiarray.py index 909025db2dab..c7dadcbf867a 100644 --- a/python/mxnet/numpy/multiarray.py +++ b/python/mxnet/numpy/multiarray.py @@ -48,8 +48,7 @@ from ..device import Device from ..util import set_module, wrap_np_unary_func, wrap_np_binary_func,\ is_np_default_dtype, wrap_ctx_to_device_func,\ - dtype_from_number, wrap_data_api_statical_func,\ - wrap_sort_functions + wrap_sort_functions, wrap_data_api_statical_func from ..device import current_device from ..ndarray import numpy as _mx_nd_np from ..ndarray.numpy import _internal as _npi @@ -13306,15 +13305,40 @@ def asarray(obj, dtype=None, device=None, copy=None): array([[0, 6], [1, 7]]) """ - if isinstance(obj, numeric_types): - dtype = dtype_from_number(obj) if dtype is None else dtype - obj = _np.asarray(obj, dtype=dtype) - elif isinstance(obj, _np.ndarray): - dtype = obj.dtype if dtype is None else dtype - elif isinstance(obj, ndarray): + if device is None: + ctx = current_device() + if isinstance(obj, _np.ndarray): + if is_np_default_dtype(): + dtype = obj.dtype if dtype is None else dtype + else: + dtype = _np.float32 if dtype is None or obj.dtype is _np.float64 else dtype + if isinstance(obj, ndarray): dtype = obj.dtype if dtype is None else dtype - array = _as_mx_np_array(obj, device=device, zero_copy=copy) - return array.astype(dtype) + elif hasattr(obj, '__dlpack__'): + return from_dlpack(obj) + elif isinstance(obj, NDArray): + raise ValueError("If you're trying to create a mxnet.numpy.ndarray " + "from mx.nd.NDArray, please use the zero-copy as_np_ndarray function.") + else: + try: + obj = _np.array(obj) + except Exception as e: + # printing out the error raised by official NumPy's array function + # for transparency on users' side + raise TypeError('{}'.format(str(e))) + if dtype is None: + default_dtype = _np.float64 if is_np_default_dtype() else _np.float32 + dtype = obj.dtype if hasattr(obj, "dtype") else default_dtype + try: + obj = obj.astype(dtype) + except Exception as e: + raise TypeError('{}'.format(str(e))) + ret = empty(obj.shape, dtype=dtype, ctx=ctx) + if len(obj.shape) == 0: + ret[()] = obj + else: + ret[:] = obj + return ret # pylint: disable=redefined-outer-name From 599fdcf8d52beec4f0c3c949ab9f85b5da3743eb Mon Sep 17 00:00:00 2001 From: barry-jin Date: Mon, 8 Nov 2021 13:37:42 -0800 Subject: [PATCH 17/39] Revert "update asarray" This reverts commit 3a11d157d007da1ae8057f9a867bc9f0fb221ef2. --- .github/workflows/os_x_staticbuild.yml | 1 - ci/docker/runtime_functions.sh | 1 - python/mxnet/numpy/multiarray.py | 44 ++++++-------------------- 3 files changed, 10 insertions(+), 36 deletions(-) diff --git a/.github/workflows/os_x_staticbuild.yml b/.github/workflows/os_x_staticbuild.yml index c6aad0676d0b..bc67b31e6830 100644 --- a/.github/workflows/os_x_staticbuild.yml +++ b/.github/workflows/os_x_staticbuild.yml @@ -69,7 +69,6 @@ jobs: python3 -m pytest --durations=50 --verbose array_api_tests/test_indexing.py python3 -m pytest --durations=50 --verbose array_api_tests/test_constants.py python3 -m pytest --durations=50 --verbose array_api_tests/test_elementwise_functions.py - python3 -m pytest --durations=50 --verbose array_api_tests/test_broadcasting.py python3 -m pytest --durations=50 --verbose \ array_api_tests/test_type_promotion.py::test_elementwise_function_two_arg_bool_type_promotion python3 -m pytest --durations=50 --verbose \ diff --git a/ci/docker/runtime_functions.sh b/ci/docker/runtime_functions.sh index c731c0d62bef..59d06e7c1416 100755 --- a/ci/docker/runtime_functions.sh +++ b/ci/docker/runtime_functions.sh @@ -876,7 +876,6 @@ unittest_array_api_standardization() { python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_indexing.py python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_elementwise_functions.py python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_constants.py - python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_broadcasting.py python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ array_api_tests/test_type_promotion.py::test_elementwise_function_two_arg_bool_type_promotion python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ diff --git a/python/mxnet/numpy/multiarray.py b/python/mxnet/numpy/multiarray.py index c7dadcbf867a..909025db2dab 100644 --- a/python/mxnet/numpy/multiarray.py +++ b/python/mxnet/numpy/multiarray.py @@ -48,7 +48,8 @@ from ..device import Device from ..util import set_module, wrap_np_unary_func, wrap_np_binary_func,\ is_np_default_dtype, wrap_ctx_to_device_func,\ - wrap_sort_functions, wrap_data_api_statical_func + dtype_from_number, wrap_data_api_statical_func,\ + wrap_sort_functions from ..device import current_device from ..ndarray import numpy as _mx_nd_np from ..ndarray.numpy import _internal as _npi @@ -13305,40 +13306,15 @@ def asarray(obj, dtype=None, device=None, copy=None): array([[0, 6], [1, 7]]) """ - if device is None: - ctx = current_device() - if isinstance(obj, _np.ndarray): - if is_np_default_dtype(): - dtype = obj.dtype if dtype is None else dtype - else: - dtype = _np.float32 if dtype is None or obj.dtype is _np.float64 else dtype - if isinstance(obj, ndarray): + if isinstance(obj, numeric_types): + dtype = dtype_from_number(obj) if dtype is None else dtype + obj = _np.asarray(obj, dtype=dtype) + elif isinstance(obj, _np.ndarray): dtype = obj.dtype if dtype is None else dtype - elif hasattr(obj, '__dlpack__'): - return from_dlpack(obj) - elif isinstance(obj, NDArray): - raise ValueError("If you're trying to create a mxnet.numpy.ndarray " - "from mx.nd.NDArray, please use the zero-copy as_np_ndarray function.") - else: - try: - obj = _np.array(obj) - except Exception as e: - # printing out the error raised by official NumPy's array function - # for transparency on users' side - raise TypeError('{}'.format(str(e))) - if dtype is None: - default_dtype = _np.float64 if is_np_default_dtype() else _np.float32 - dtype = obj.dtype if hasattr(obj, "dtype") else default_dtype - try: - obj = obj.astype(dtype) - except Exception as e: - raise TypeError('{}'.format(str(e))) - ret = empty(obj.shape, dtype=dtype, ctx=ctx) - if len(obj.shape) == 0: - ret[()] = obj - else: - ret[:] = obj - return ret + elif isinstance(obj, ndarray): + dtype = obj.dtype if dtype is None else dtype + array = _as_mx_np_array(obj, device=device, zero_copy=copy) + return array.astype(dtype) # pylint: disable=redefined-outer-name From ac811b205c7a0495b5f8c1be47f4ff3f8fbd70ca Mon Sep 17 00:00:00 2001 From: barry-jin Date: Mon, 8 Nov 2021 14:05:07 -0800 Subject: [PATCH 18/39] fix precision --- .github/workflows/os_x_staticbuild.yml | 1 + ci/docker/runtime_functions.sh | 1 + python/mxnet/numpy/multiarray.py | 2 ++ python/mxnet/util.py | 2 +- 4 files changed, 5 insertions(+), 1 deletion(-) diff --git a/.github/workflows/os_x_staticbuild.yml b/.github/workflows/os_x_staticbuild.yml index bc67b31e6830..c6aad0676d0b 100644 --- a/.github/workflows/os_x_staticbuild.yml +++ b/.github/workflows/os_x_staticbuild.yml @@ -69,6 +69,7 @@ jobs: python3 -m pytest --durations=50 --verbose array_api_tests/test_indexing.py python3 -m pytest --durations=50 --verbose array_api_tests/test_constants.py python3 -m pytest --durations=50 --verbose array_api_tests/test_elementwise_functions.py + python3 -m pytest --durations=50 --verbose array_api_tests/test_broadcasting.py python3 -m pytest --durations=50 --verbose \ array_api_tests/test_type_promotion.py::test_elementwise_function_two_arg_bool_type_promotion python3 -m pytest --durations=50 --verbose \ diff --git a/ci/docker/runtime_functions.sh b/ci/docker/runtime_functions.sh index 59d06e7c1416..c731c0d62bef 100755 --- a/ci/docker/runtime_functions.sh +++ b/ci/docker/runtime_functions.sh @@ -876,6 +876,7 @@ unittest_array_api_standardization() { python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_indexing.py python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_elementwise_functions.py python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_constants.py + python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_broadcasting.py python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ array_api_tests/test_type_promotion.py::test_elementwise_function_two_arg_bool_type_promotion python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ diff --git a/python/mxnet/numpy/multiarray.py b/python/mxnet/numpy/multiarray.py index 909025db2dab..dcb2c0e9b798 100644 --- a/python/mxnet/numpy/multiarray.py +++ b/python/mxnet/numpy/multiarray.py @@ -13313,6 +13313,8 @@ def asarray(obj, dtype=None, device=None, copy=None): dtype = obj.dtype if dtype is None else dtype elif isinstance(obj, ndarray): dtype = obj.dtype if dtype is None else dtype + elif hasattr(obj, '__dlpack__'): + return from_dlpack(obj) array = _as_mx_np_array(obj, device=device, zero_copy=copy) return array.astype(dtype) diff --git a/python/mxnet/util.py b/python/mxnet/util.py index cf2c2a95e628..ed0ca1449cf8 100644 --- a/python/mxnet/util.py +++ b/python/mxnet/util.py @@ -1349,7 +1349,7 @@ def dtype_from_number(number): else: return _np.int32 else: - if abs(number) > _MAX_VALUE_FLOAT32_REPRESENT_: + if abs(number) > _MAX_VALUE_FLOAT32_REPRESENT_ or abs(number) % 1 != 0: return _np.float64 else: return _np.float64 if is_np_default_dtype() else _np.float32 From 77388efcc31a6505fd01838a0ca7b4ab445b5b34 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Mon, 8 Nov 2021 14:30:00 -0800 Subject: [PATCH 19/39] fix precision --- python/mxnet/util.py | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/python/mxnet/util.py b/python/mxnet/util.py index ed0ca1449cf8..87e91bc6f99c 100644 --- a/python/mxnet/util.py +++ b/python/mxnet/util.py @@ -1349,7 +1349,8 @@ def dtype_from_number(number): else: return _np.int32 else: - if abs(number) > _MAX_VALUE_FLOAT32_REPRESENT_ or abs(number) % 1 != 0: + if abs(number) > _MAX_VALUE_FLOAT32_REPRESENT_ or \ + _np.float32(number) != _np.float64(number): return _np.float64 else: return _np.float64 if is_np_default_dtype() else _np.float32 From ba4ccd5e3feb2e4e0dbaeb3f93b1e6289af52234 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Mon, 8 Nov 2021 17:25:24 -0800 Subject: [PATCH 20/39] fix --- python/mxnet/util.py | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/python/mxnet/util.py b/python/mxnet/util.py index 87e91bc6f99c..cf2c2a95e628 100644 --- a/python/mxnet/util.py +++ b/python/mxnet/util.py @@ -1349,8 +1349,7 @@ def dtype_from_number(number): else: return _np.int32 else: - if abs(number) > _MAX_VALUE_FLOAT32_REPRESENT_ or \ - _np.float32(number) != _np.float64(number): + if abs(number) > _MAX_VALUE_FLOAT32_REPRESENT_: return _np.float64 else: return _np.float64 if is_np_default_dtype() else _np.float32 From 3deb6738e4b56e4d7d473a6c3200f8b0f568ed5a Mon Sep 17 00:00:00 2001 From: barry-jin Date: Mon, 8 Nov 2021 20:46:14 -0800 Subject: [PATCH 21/39] fix floating point exception --- src/operator/mshadow_op.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/operator/mshadow_op.h b/src/operator/mshadow_op.h index 985975b3dcf6..8e0194d662ef 100644 --- a/src/operator/mshadow_op.h +++ b/src/operator/mshadow_op.h @@ -239,7 +239,7 @@ struct floor_divide : public mxnet_op::tunable { typename std::enable_if::value && std::is_integral::value, int>::type = 0> MSHADOW_XINLINE static DType Map(DType a, DType b) { - DType c = static_cast(::floor(a / b)); + DType c = static_cast(::floor(static_cast(a) / static_cast(b))); if ((c * a != b) && ((a < 0) != (b < 0))) { return DType(c - 1); } else { @@ -270,7 +270,7 @@ struct rfloor_divide : public mxnet_op::tunable { typename std::enable_if::value && std::is_integral::value, int>::type = 0> MSHADOW_XINLINE static DType Map(DType a, DType b) { - DType c = static_cast(::floor(b / a)); + DType c = static_cast(::floor(static_cast(b) / static_cast(a))); if ((c * a != b) && ((a < 0) != (b < 0))) { return DType(c - 1); } else { From 780f5ae5e822d55e460c570292831713b76bf7ab Mon Sep 17 00:00:00 2001 From: barry-jin Date: Tue, 9 Nov 2021 10:48:31 -0800 Subject: [PATCH 22/39] fix floor_divide --- src/operator/mshadow_op.h | 14 ++------------ 1 file changed, 2 insertions(+), 12 deletions(-) diff --git a/src/operator/mshadow_op.h b/src/operator/mshadow_op.h index 8e0194d662ef..3fbc5c6901b0 100644 --- a/src/operator/mshadow_op.h +++ b/src/operator/mshadow_op.h @@ -239,12 +239,7 @@ struct floor_divide : public mxnet_op::tunable { typename std::enable_if::value && std::is_integral::value, int>::type = 0> MSHADOW_XINLINE static DType Map(DType a, DType b) { - DType c = static_cast(::floor(static_cast(a) / static_cast(b))); - if ((c * a != b) && ((a < 0) != (b < 0))) { - return DType(c - 1); - } else { - return c; - } + return static_cast(::floor(static_cast(a) / static_cast(b))); } MSHADOW_XINLINE static bool Map(bool a, bool b) { @@ -270,12 +265,7 @@ struct rfloor_divide : public mxnet_op::tunable { typename std::enable_if::value && std::is_integral::value, int>::type = 0> MSHADOW_XINLINE static DType Map(DType a, DType b) { - DType c = static_cast(::floor(static_cast(b) / static_cast(a))); - if ((c * a != b) && ((a < 0) != (b < 0))) { - return DType(c - 1); - } else { - return c; - } + return static_cast(::floor(static_cast(b) / static_cast(a))); } MSHADOW_XINLINE static bool Map(bool a, bool b) { From c176ec27c08f9f8c94f38540908c15b4228799a1 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Tue, 9 Nov 2021 15:27:33 -0800 Subject: [PATCH 23/39] fix dtype_from_number --- python/mxnet/util.py | 14 ++++++++++---- 1 file changed, 10 insertions(+), 4 deletions(-) diff --git a/python/mxnet/util.py b/python/mxnet/util.py index cf2c2a95e628..f99dfd07413e 100644 --- a/python/mxnet/util.py +++ b/python/mxnet/util.py @@ -23,7 +23,7 @@ from struct import calcsize from .base import (_LIB, check_call, c_str, py_str, - numeric_types, integer_types, + numeric_types, integer_types, long, _MAX_VALUE_64_BIT_UNSIGNED_, _MAX_VALUE_64_BIT_SIGNED_, _MAX_VALUE_FLOAT32_REPRESENT_) @@ -1339,7 +1339,7 @@ def dtype_from_number(number): assert isinstance(number, numeric_types),\ "The input number should be either int for float types" import numpy as _np - if isinstance(number, integer_types): + if isinstance(number, (int, long)): if number > _MAX_VALUE_64_BIT_UNSIGNED_: raise OverflowError("Integer out of bounds") if number > _MAX_VALUE_64_BIT_SIGNED_: @@ -1348,8 +1348,14 @@ def dtype_from_number(number): return _np.int64 else: return _np.int32 - else: - if abs(number) > _MAX_VALUE_FLOAT32_REPRESENT_: + elif isinstance(number, float): + if abs(number) > _MAX_VALUE_FLOAT32_REPRESENT_ or \ + ((not _np.isnan(number)) and \ + (_np.float32(number) == int(number)) and \ + (number != int(number))): return _np.float64 else: return _np.float64 if is_np_default_dtype() else _np.float32 + elif isinstance(number, _np.generic): + return number.dtype + raise TypeError('type {} not supported'.format(str(type(number)))) From a4cabe53de1c30534b2fa3a82eb2a31e7649ec06 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Tue, 9 Nov 2021 16:39:29 -0800 Subject: [PATCH 24/39] fix asarray --- python/mxnet/numpy/multiarray.py | 31 +++++++++++++++++++++++++++---- 1 file changed, 27 insertions(+), 4 deletions(-) diff --git a/python/mxnet/numpy/multiarray.py b/python/mxnet/numpy/multiarray.py index dcb2c0e9b798..f2402a950743 100644 --- a/python/mxnet/numpy/multiarray.py +++ b/python/mxnet/numpy/multiarray.py @@ -13310,13 +13310,36 @@ def asarray(obj, dtype=None, device=None, copy=None): dtype = dtype_from_number(obj) if dtype is None else dtype obj = _np.asarray(obj, dtype=dtype) elif isinstance(obj, _np.ndarray): - dtype = obj.dtype if dtype is None else dtype + if is_np_default_dtype(): + dtype = obj.dtype if dtype is None else dtype + else: + dtype = _np.float32 if dtype is None or obj.dtype is _np.float64 else dtype elif isinstance(obj, ndarray): - dtype = obj.dtype if dtype is None else dtype + if dtype is not None: + obj = obj.astype(dtype, copy=copy) + if device is not None: + obj = obj.to_device(device) + return obj elif hasattr(obj, '__dlpack__'): return from_dlpack(obj) - array = _as_mx_np_array(obj, device=device, zero_copy=copy) - return array.astype(dtype) + else: + if dtype is None: + default_dtype = _np.float64 if is_np_default_dtype() else _np.float32 + dtype = obj.dtype if hasattr(obj, "dtype") else default_dtype + try: + obj = _np.array(obj, dtype=dtype) + except Exception as e: + # printing out the error raised by official NumPy's array function + # for transparency on users' side + raise TypeError('{}'.format(str(e))) + if device is None: + device = current_device() + ret = empty(obj.shape, dtype=dtype, device=device) + if len(obj.shape) == 0: + ret[()] = obj + else: + ret[:] = obj + return ret # pylint: disable=redefined-outer-name From 05dc28bd7fd421f057e0458bc43221df34049256 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Tue, 9 Nov 2021 16:45:09 -0800 Subject: [PATCH 25/39] fix asarray docstring --- python/mxnet/numpy/multiarray.py | 25 ++++++++----------------- 1 file changed, 8 insertions(+), 17 deletions(-) diff --git a/python/mxnet/numpy/multiarray.py b/python/mxnet/numpy/multiarray.py index f2402a950743..5a2ac27f7e4c 100644 --- a/python/mxnet/numpy/multiarray.py +++ b/python/mxnet/numpy/multiarray.py @@ -13287,24 +13287,15 @@ def asarray(obj, dtype=None, device=None, copy=None): Examples -------- - >>> a = np.arange(4).reshape(2,2) - >>> a - array([[0, 1], - [2, 3]]) - >>> np.diagonal(a) - array([0, 3]) - >>> np.diagonal(a, 1) - array([1]) + >>> np.asarray([1, 2, 3]) + array([1., 2., 3.]) - >>> a = np.arange(8).reshape(2,2,2) - >>>a - array([[[0, 1], - [2, 3]], - [[4, 5], - [6, 7]]]) - >>> np.diagonal(a, 0, 0, 1) - array([[0, 6], - [1, 7]]) + >>> np.asarray([[1, 2], [3, 4]], dtype=np.int32) + array([[1, 2], + [3, 4]], dtype=int32) + + >>> np.asarray([1.2], device=mx.gpu()) + array([1.2], device=gpu(0)) """ if isinstance(obj, numeric_types): dtype = dtype_from_number(obj) if dtype is None else dtype From 6b77e387d0b0bca11c95e5a32af68f85ef358a60 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Tue, 9 Nov 2021 20:48:04 -0800 Subject: [PATCH 26/39] merge data type functions --- python/mxnet/numpy/__init__.py | 1 + python/mxnet/numpy/fallback.py | 1 - python/mxnet/numpy/type_functions.py | 163 +++++++++++++++++++++++++ python/mxnet/numpy/utils.py | 101 ++++++++++++++- tests/python/unittest/test_numpy_op.py | 30 +++++ 5 files changed, 292 insertions(+), 4 deletions(-) create mode 100644 python/mxnet/numpy/type_functions.py diff --git a/python/mxnet/numpy/__init__.py b/python/mxnet/numpy/__init__.py index 45699f714ed4..1228dac666e8 100644 --- a/python/mxnet/numpy/__init__.py +++ b/python/mxnet/numpy/__init__.py @@ -28,6 +28,7 @@ from .function_base import * # pylint: disable=wildcard-import from .stride_tricks import * # pylint: disable=wildcard-import from .set_functions import * # pylint: disable=wildcard-import +from .type_functions import * # pylint: disable=wildcard-import from .io import * # pylint: disable=wildcard-import from .arrayprint import * # pylint: disable=wildcard-import diff --git a/python/mxnet/numpy/fallback.py b/python/mxnet/numpy/fallback.py index 83bf67372517..c8fc7fbaf7f8 100644 --- a/python/mxnet/numpy/fallback.py +++ b/python/mxnet/numpy/fallback.py @@ -94,7 +94,6 @@ 'pv', 'rate', 'real', - 'result_type', 'roots', 'searchsorted', 'select', diff --git a/python/mxnet/numpy/type_functions.py b/python/mxnet/numpy/type_functions.py new file mode 100644 index 000000000000..6d8125d4a895 --- /dev/null +++ b/python/mxnet/numpy/type_functions.py @@ -0,0 +1,163 @@ +# 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. + +"""Type functions for the numpy module.""" + +from typing import NamedTuple + +import numpy as onp +from .multiarray import ndarray +from .utils import _type_promotion_table + + +__all__ = ['can_cast', 'finfo', 'iinfo', 'result_type'] + +class finfo_obj(NamedTuple): + bits: int + eps: float + max: float + min: float + smallest_normal: float + + +class iinfo_obj(NamedTuple): + bits: int + max: int + min: int + + +def can_cast(from_, to): + """ + Returns True if cast between data types can occur according to + the casting rule. If from is a scalar or array scalar, + also returns True if the scalar value can be cast without + overflow or truncation to an integer. + Parameters + ---------- + from_ : dtype, ndarray or scalar + Data type, scalar, or array to cast from. + to : dtype + Data type to cast to. + Returns + ------- + out : bool + True if cast can occur according to the casting rule. + """ + if isinstance(from_, ndarray): + from_ = from_.asnumpy() + return onp.can_cast(from_, to) + + +def finfo(dtype): + """ + Machine limits for floating-point data types. + Notes + ----- + `finfo` is a standard API in + https://data-apis.org/array-api/latest/API_specification/data_type_functions.html#finfo-type + instead of an official NumPy operator. + Parameters + ---------- + dtype : ndarray, float or dtype + Kind of floating point data-type about which to get information. + Returns + ------- + out : finfo object + an object having the following attributes: + - bits : int + number of bits occupied by the floating-point data type. + - eps : float + difference between 1.0 and the next smallest representable floating-point + number larger than 1.0 according to the IEEE-754 standard. + - max : float + largest representable number. + - min : float + smallest representable number. + - smallest_normal : float + smallest positive floating-point number with full precision. + """ + f_info = onp.finfo(dtype) + return finfo_obj(f_info.bits, float(f_info.eps), + float(f_info.max), float(f_info.min), float(f_info.tiny)) + + +def iinfo(dtype): + """ + Machine limits for floating-point data types. + Notes + ----- + `iinfo` is a standard API in + https://data-apis.org/array-api/latest/API_specification/data_type_functions.html#iinfo-type + instead of an official NumPy operator. + Parameters + ---------- + dtype : ndarray, integer or dtype + The kind of integer data type to get information about. + Returns + ------- + out : iinfo object + an object having the following attributes: + - bits : int + number of bits occupied by the type + - max : int + largest representable number. + - min : int + smallest representable number. + """ + i_info = onp.iinfo(dtype) + return iinfo_obj(i_info.bits, i_info.max, i_info.min) + + +def _get_dtype(array_or_dtype): + """Utility function for result_type""" + if isinstance(array_or_dtype, (ndarray, onp.ndarray)): + return array_or_dtype.dtype + elif isinstance(array_or_dtype, onp.dtype): + return array_or_dtype + else: + raise ValueError("Inputs of result_type must be ndarrays or dtypes") + + +def result_type(*arrays_and_dtypes): + """ + Returns the dtype that results from applying the type promotion rules to the arguments. + Notes + ----- + `result_type` is a standard API in + https://data-apis.org/array-api/latest/API_specification/data_type_functions.html#result-type-arrays-and-dtypes + instead of an official NumPy operator. + Parameters + ---------- + arrays_and_dtypes : mixed ndarrays and dtypes + an arbitrary number of input arrays and/or dtypes. + Returns + ------- + out : dtype + the dtype resulting from an operation involving the input arrays and dtypes. + """ + if len(arrays_and_dtypes) > 0: + ret = _get_dtype(arrays_and_dtypes[0]) + for d in arrays_and_dtypes[1:]: + dd = _get_dtype(arrays_and_dtypes[d]) + if (ret, dd) in _type_promotion_table: + ret = _type_promotion_table[ret, dd] + elif (dd, ret) in _type_promotion_table: + ret = _type_promotion_table[dd, ret] + else: + raise TypeError("Unknown type promotion between {} and {}".format(ret, dd)) + return ret + raise ValueError("at least one array or dtype is required") diff --git a/python/mxnet/numpy/utils.py b/python/mxnet/numpy/utils.py index 35ea80a8e392..378ab749ee74 100644 --- a/python/mxnet/numpy/utils.py +++ b/python/mxnet/numpy/utils.py @@ -23,9 +23,9 @@ __all__ = ['float16', 'float32', 'float64', 'uint8', 'int32', 'int8', 'int64', 'int16', 'uint16', 'uint32', 'uint64', - 'bool', 'bool_', 'pi', 'inf', 'nan', 'PZERO', 'NZERO', 'newaxis', 'finfo', + 'bool', 'bool_', 'pi', 'inf', 'nan', 'PZERO', 'NZERO', 'newaxis', 'e', 'NINF', 'PINF', 'NAN', 'NaN', - '_STR_2_DTYPE_', '_DTYPE_2_STR_'] + '_STR_2_DTYPE_', '_DTYPE_2_STR_', '_type_promotion_table'] py_bool = bool @@ -55,7 +55,6 @@ NaN = onp.NaN newaxis = None -finfo = onp.finfo _STR_2_DTYPE_ = {'float16': float16, 'float32': float32, 'float64': float64, 'float': float64, 'int8': int8, 'int16': int16, 'int32': int32, 'int64': int64, 'int': int64, @@ -77,3 +76,99 @@ def _get_np_op(name): if op is not None: return op raise ValueError('Operator `{}` is not supported by `mxnet.numpy`.'.format(name)) + + +_type_promotion_table = { + # signed integer type promotion + (int8, int8): int8, + (int8, int16): int16, + (int8, int32): int32, + (int8, int64): int64, + (int16, int16): int16, + (int16, int32): int32, + (int16, int64): int64, + (int32, int32): int32, + (int32, int64): int64, + (int64, int64): int64, + # unsigned integer type promotion + (uint8, uint8): uint8, + (uint8, uint16): uint16, + (uint8, uint32): uint32, + (uint8, uint64): uint64, + (uint16, uint16): uint16, + (uint16, uint32): uint32, + (uint16, uint64): uint64, + (uint32, uint32): uint32, + (uint32, uint64): uint64, + (uint64, uint64): uint64, + # mixed signed and unsigned integer type promotion + (int8, uint8): int16, + (int8, uint16): int32, + (int8, uint32): int64, + (int16, uint8): int16, + (int16, uint16): int32, + (int16, uint32): int64, + (int32, uint8): int32, + (int32, uint16): int32, + (int32, uint32): int64, + (int64, uint8): int64, + (int64, uint16): int64, + (int64, uint32): int64, + # float type promotion + (float16, float16): float16, + (float16, float32): float32, + (float16, float64): float64, + (float32, float32): float32, + (float32, float64): float64, + (float64, float64): float64, + # bool type promotion + (bool, bool): bool, + # mixed integer and float16 type promotion + (int8, float16): float16, + (int16, float16): float16, + (int32, float16): float16, + (int64, float16): float16, + (uint8, float16): float16, + (uint16, float16): float16, + (uint32, float16): float16, + (uint64, float16): float16, + # mixed integer and float16 type promotion + (int8, float32): float32, + (int16, float32): float32, + (int32, float32): float32, + (int64, float32): float32, + (uint8, float32): float32, + (uint16, float32): float32, + (uint32, float32): float32, + (uint64, float32): float32, + # mixed integer and float32 type promotion + (int8, float32): float32, + (int16, float32): float32, + (int32, float32): float32, + (int64, float32): float32, + (uint8, float32): float32, + (uint16, float32): float32, + (uint32, float32): float32, + (uint64, float32): float32, + # mixed integer and float64 type promotion + (int8, float64): float64, + (int16, float64): float64, + (int32, float64): float64, + (int64, float64): float64, + (uint8, float64): float64, + (uint16, float64): float64, + (uint32, float64): float32, + (uint64, float64): float64, + # mixed bool and other type promotion + (bool, int8): int8, + (bool, int16): int16, + (bool, int32): int32, + (bool, int64): int64, + (bool, uint8): uint8, + (bool, uint16): uint16, + (bool, uint32): uint32, + (bool, uint64): uint64, + (bool, float16): float16, + (bool, float32): float32, + (bool, float64): float64, +} diff --git a/tests/python/unittest/test_numpy_op.py b/tests/python/unittest/test_numpy_op.py index 0db209c5774f..c0017abe2910 100644 --- a/tests/python/unittest/test_numpy_op.py +++ b/tests/python/unittest/test_numpy_op.py @@ -11239,3 +11239,33 @@ def forward(self, x, *args): assert_almost_equal(deconvOut, deconvRefOut) assert_almost_equal(deconvData.grad, deconvRefGrad) + + +@use_np +@pytest.mark.parametrize('dtype', np.floating_dtypes) +def test_np_finfo(dtype): + mx_finfo_obj = np.finfo(dtype) + np_finfo = onp.finfo(dtype) + assert (mx_finfo_obj.bits, mx_finfo_obj.eps, mx_finfo_obj.max, mx_finfo_obj.min, mx_finfo_obj.smallest_normal) == \ + (np_finfo.bits, np_finfo.eps, np_finfo.max, np_finfo.min, np_finfo.tiny) + + +@use_np +@pytest.mark.parametrize('dtype', np.integer_dtypes) +def test_np_iinfo(dtype): + mx_iinfo_obj = np.iinfo(dtype) + np_iinfo = onp.iinfo(dtype) + assert (mx_iinfo_obj.bits, mx_iinfo_obj.max, mx_iinfo_obj.min) == \ + (np_iinfo.bits, np_iinfo.max, np_iinfo.min) + + +@use_np +@pytest.mark.parametrize('input1', [d for d in np.numeric_dtypes + np.boolean_dtypes] + [np.ones((1,), dtype=d) for d in np.numeric_dtypes + np.boolean_dtypes]) +@pytest.mark.parametrize('input2', [d for d in np.numeric_dtypes + np.boolean_dtypes]) +def test_np_can_cast(input1, input2): + np_input1 = input1 + np_input2 = input2 + if isinstance(input1, np.ndarray): + np_input1 = input1.asnumpy() + assert np.can_cast(input1, input2) == onp.can_cast(np_input1, np_input2) + From 9c790236f998310bc1362e0ad17dcf4adeda0901 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Tue, 9 Nov 2021 23:34:48 -0800 Subject: [PATCH 27/39] add un-func standard tests --- python/mxnet/numpy/utils.py | 29 ++++++- tests/python/unittest/test_numpy_op.py | 103 +++++++++++++++++++++++++ 2 files changed, 131 insertions(+), 1 deletion(-) diff --git a/python/mxnet/numpy/utils.py b/python/mxnet/numpy/utils.py index 378ab749ee74..29b0a607aa5b 100644 --- a/python/mxnet/numpy/utils.py +++ b/python/mxnet/numpy/utils.py @@ -25,7 +25,8 @@ 'int16', 'uint16', 'uint32', 'uint64', 'bool', 'bool_', 'pi', 'inf', 'nan', 'PZERO', 'NZERO', 'newaxis', 'e', 'NINF', 'PINF', 'NAN', 'NaN', - '_STR_2_DTYPE_', '_DTYPE_2_STR_', '_type_promotion_table'] + '_STR_2_DTYPE_', '_DTYPE_2_STR_', '_type_promotion_table', + 'integer_dtypes', 'floating_dtypes', 'boolean_dtypes', 'numeric_dtypes'] py_bool = bool @@ -172,3 +173,29 @@ def _get_np_op(name): (bool, float32): float32, (bool, float64): float64, } + +integer_dtypes = [ + int8, + int16, + int32, + int64, + uint8, + uint16, + uint32, + uint64, +] + +floating_dtypes = [ + float16, + float32, + float64, +] + +numeric_dtypes = [ + *integer_dtypes, + *floating_dtypes, +] + +boolean_dtypes = [ + bool_, +] diff --git a/tests/python/unittest/test_numpy_op.py b/tests/python/unittest/test_numpy_op.py index c0017abe2910..44fda6afad8e 100644 --- a/tests/python/unittest/test_numpy_op.py +++ b/tests/python/unittest/test_numpy_op.py @@ -11269,3 +11269,106 @@ def test_np_can_cast(input1, input2): np_input1 = input1.asnumpy() assert np.can_cast(input1, input2) == onp.can_cast(np_input1, np_input2) + +@use_np +@retry(3) +@pytest.mark.parametrize('func,func2,dtypes,ref_grad,low,high', [ + ('abs', 'abs', 'numeric', lambda x: -1. * (x < 0) + (x > 0), -1.0, 1.0), + ('acos', 'arccos', 'floating-point', lambda x: -1. / (1. - x ** 2.) ** (1. / 2.), -1.0, 1.0), + ('acosh', 'arccosh', 'floating-point', lambda x: 1./(x**2 - 1.)**(1./2.), 2.0, 5.0), + ('asin', 'arcsin', 'floating-point', lambda x: 1. / (1. - x ** 2) ** (1. / 2.), -1.0, 1.0), + ('asinh', 'arcsinh', 'floating-point', lambda x: 1./(x**2 + 1.)**(1./2.), -1.0, 1.0), + ('atan', 'arctan', 'floating-point', lambda x: 1. / (x ** 2. + 1.), -1.0, 1.0), + ('atanh', 'arctanh', 'floating-point', lambda x: -1./(x**2 - 1.), -0.99, 0.99), + ('bitwise_invert', 'invert', 'integer or boolean', None, -5, 5), + ('ceil', 'ceil', 'numeric', None, -10.0, 10.0), + ('cos', 'cos', 'floating-point', lambda x: -onp.sin(x), -1.0, 1.0), + ('cosh', 'cosh', 'floating-point', lambda x: onp.sinh(x), -1.0, 1.0), + ('exp', 'exp', 'floating-point', lambda x: onp.exp(x), -1.0, 1.0), + ('expm1', 'expm1', 'floating-point', lambda x: onp.exp(x), -1.0, 1.0), + ('floor', 'floor', 'numeric', None, -10.0, 10.0), + ('log', 'log', 'floating-point', lambda x: 1.0 / x, 0.1, 5.0), + ('log10', 'log10', 'floating-point', lambda x: 1.0 / (x * onp.log(10)), 0.1, 10.0), + ('log1p', 'log1p', 'floating-point', lambda x: 1.0 / (1.0 + x), -0.9, 5.0), + ('log2', 'log2', 'floating-point', lambda x: 1.0 / (x * onp.log(2)), 0.1, 2.0), + ('logical_not', 'logical_not', 'boolean', None, -1.0, 1.0), + ('negative', 'negative', 'numeric', lambda x: -1. * onp.ones(x.shape), -1.0, 1.0), + ('positive', 'positive', 'numeric', lambda x: onp.ones(x.shape), -1.0, 1.0), + ('sign', 'sign', 'numeric', None, -1.0, 1.0), + ('sin', 'sin', 'floating-point', lambda x: onp.cos(x), -1.0, 1.0), + ('sinh', 'sinh', 'floating-point', lambda x: onp.cosh(x), -1.0, 1.0), + ('sqrt', 'sqrt', 'floating-point', lambda x: 0.5 / onp.sqrt(x), 0.001, 10.0), + ('square', 'square', 'numeric', lambda x: 2.0 * x, -1.0, 1.0), + ('tan', 'tan', 'floating-point', lambda x: onp.tan(x) ** 2 + 1.0, -1.0, 1.0), + ('tanh', 'tanh', 'floating-point', lambda x: 1. - onp.tanh(x) ** 2, -1.0, 1.0), + ('trunc', 'trunc', 'numeric', None, -5.0, 5.0), +]) +@pytest.mark.parametrize('ndim', [2, 3, 4]) +def test_np_standard_unary_funcs(func, func2, dtypes, ref_grad, low, high, ndim): + class TestStandardUnary(HybridBlock): + def __init__(self, func): + super(TestStandardUnary, self).__init__() + self._func = func + + def forward(self, a): + return getattr(np, self._func)(a) + + type_mapping = { + 'floating-point': np.floating_dtypes, + 'numeric': np.numeric_dtypes, + 'integer or boolean': np.integer_dtypes + np.boolean_dtypes, + 'boolean': np.boolean_dtypes, + } + + def array_values(low, high, shape): + for d in np.integer_dtypes + np.boolean_dtypes + np.floating_dtypes: + yield onp.random.uniform(low, high, shape).astype(d), d + + + shapes = [i for i in [rand_shape_nd(ndim, dim=3), (1, 0, 2)]] + for shape in shapes: + for (np_test_data, dtype) in array_values(low, high, shape): + if dtype in type_mapping[dtypes]: + rtol = 1e-2 if dtype == np.float16 else 1e-3 + atol = 1e-4 if dtype == np.float16 else 1e-5 + # get rid of warning: divide by zero + if((func=='log' or func=='log10' or func=='log2') and + (dtype=='int8' or dtype=='uint8' or dtype=='int32' or + dtype=='int64')): + low = 1 + if (func=='arctanh' and dtype=='bool'): + continue + np_func = getattr(onp, func2) + mx_func = TestStandardUnary(func) + mx_test_data = np.array(np_test_data, dtype=dtype) + for hybridize in [True, False]: + if hybridize: + mx_func.hybridize() + if ref_grad: + mx_test_data.attach_grad() + np_out = np_func(np_test_data) + with mx.autograd.record(): + y = mx_func(mx_test_data) + assert y.shape == np_out.shape + assert_almost_equal(y.asnumpy(), np_out, rtol=1e-3, atol=atol) + if np_out.dtype == np.bool_: + assert y.dtype == np.bool_ + + if ref_grad and (dtype == 'float16' or dtype == 'float32' or dtype == 'float64'): + y.backward() + assert_almost_equal(mx_test_data.grad.asnumpy(), ref_grad(np_test_data), rtol=1e-1, atol=1e-2, equal_nan=True) + + np_func = getattr(onp, func2) + mx_out = getattr(mx.np, func)(mx_test_data) + assert mx_out.shape == np_out.shape + assert mx_out.dtype == dtype + assert_almost_equal(mx_out.asnumpy(), np_out, rtol=rtol, atol=1e-5) + + assertRaises(NotImplementedError, getattr(np, func), mx_test_data, where=False) + assertRaises(NotImplementedError, getattr(np, func), mx_test_data, subok=False) + assertRaises(NotImplementedError, getattr(np, func), mx_test_data, dtype=onp.int8) + assertRaises(TypeError, getattr(np, func), mx_test_data, dtype="abcdefg") + assertRaises(NotImplementedError, getattr(np, func), mx_test_data, casting='safe') + assertRaises(TypeError, getattr(np, func), mx_test_data, casting='mxnet') + assertRaises(NotImplementedError, getattr(np, func), mx_test_data, order='C') + assertRaises(NotImplementedError, getattr(np, func), mx_test_data, order='mxnet') From 2e87451f5df173b9389949e680181996c4aa1bc4 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Wed, 10 Nov 2021 08:57:44 -0800 Subject: [PATCH 28/39] support multiple dtypes in gpu copy --- src/ndarray/ndarray_function.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/ndarray/ndarray_function.cu b/src/ndarray/ndarray_function.cu index f6189f939131..3313014ec908 100644 --- a/src/ndarray/ndarray_function.cu +++ b/src/ndarray/ndarray_function.cu @@ -46,7 +46,7 @@ void Copy(const TBlob& from, RunContext ctx) { CHECK_EQ(to->type_flag_, from.type_flag_) << "Source and target must have the same data type when copying across devices."; - MSHADOW_TYPE_SWITCH_WITH_BOOL(to->type_flag_, DType, { + MSHADOW_TYPE_SWITCH_EXT_WITH_BOOL(to->type_flag_, DType, { mshadow::Copy(to->FlatTo1D(), from.FlatTo1D(), ctx.get_stream()); }); } @@ -59,7 +59,7 @@ void Copy(const TBlob& from, RunContext ctx) { CHECK_EQ(to->type_flag_, from.type_flag_) << "Source and target must have the same data type when copying across devices."; - MSHADOW_TYPE_SWITCH_WITH_BOOL(to->type_flag_, DType, { + MSHADOW_TYPE_SWITCH_EXT_WITH_BOOL(to->type_flag_, DType, { mshadow::Copy(to->FlatTo1D(), from.FlatTo1D(), ctx.get_stream()); }); } From 5d88d75a7c8f791c7efe7bf4632f3e40da90637b Mon Sep 17 00:00:00 2001 From: barry-jin Date: Wed, 10 Nov 2021 14:25:21 -0800 Subject: [PATCH 29/39] add type_result tests --- python/mxnet/numpy/type_functions.py | 2 +- python/mxnet/numpy/utils.py | 2 +- src/common/utils.cc | 8 ++++++++ src/common/utils.h | 20 +++++++++++++++----- src/ndarray/ndarray_function-inl.h | 2 +- tests/python/unittest/test_numpy_op.py | 14 ++++++++++++++ 6 files changed, 40 insertions(+), 8 deletions(-) diff --git a/python/mxnet/numpy/type_functions.py b/python/mxnet/numpy/type_functions.py index 6d8125d4a895..bf95f1cc8ef7 100644 --- a/python/mxnet/numpy/type_functions.py +++ b/python/mxnet/numpy/type_functions.py @@ -152,7 +152,7 @@ def result_type(*arrays_and_dtypes): if len(arrays_and_dtypes) > 0: ret = _get_dtype(arrays_and_dtypes[0]) for d in arrays_and_dtypes[1:]: - dd = _get_dtype(arrays_and_dtypes[d]) + dd = _get_dtype(d) if (ret, dd) in _type_promotion_table: ret = _type_promotion_table[ret, dd] elif (dd, ret) in _type_promotion_table: diff --git a/python/mxnet/numpy/utils.py b/python/mxnet/numpy/utils.py index 29b0a607aa5b..21fe1e299d2e 100644 --- a/python/mxnet/numpy/utils.py +++ b/python/mxnet/numpy/utils.py @@ -158,7 +158,7 @@ def _get_np_op(name): (int64, float64): float64, (uint8, float64): float64, (uint16, float64): float64, - (uint32, float64): float32, + (uint32, float64): float64, (uint64, float64): float64, # mixed bool and other type promotion (bool, int8): int8, diff --git a/src/common/utils.cc b/src/common/utils.cc index f400093cc9b5..639ded4ec80e 100644 --- a/src/common/utils.cc +++ b/src/common/utils.cc @@ -117,6 +117,14 @@ MShadowTypeInfo mshadow_type_info(const int type_flag) { return MShadowTypeInfo("float16", 2, sizeof(float)); case kUint8: return MShadowTypeInfo("uint8", sizeof(uint8_t), sizeof(index_t)); + case kUint16: + return MShadowTypeInfo("uint16", sizeof(uint16_t)); + case kUint32: + return MShadowTypeInfo("uint32", sizeof(uint32_t)); + case kUint64: + return MShadowTypeInfo("uint64", sizeof(uint64_t)); + case kInt16: + return MShadowTypeInfo("int16", sizeof(int16_t)); case kInt32: return MShadowTypeInfo("int32", sizeof(int32_t)); case kInt8: diff --git a/src/common/utils.h b/src/common/utils.h index 3965f9fcf35a..269dbbbc2b76 100644 --- a/src/common/utils.h +++ b/src/common/utils.h @@ -983,6 +983,10 @@ inline int type_promotion(const int type1, const int type2) { return mshadow::kUint16; } return mshadow::kUint8; + } else if (type1 == mshadow::kBool) { + return type2; + } else if (type2 == mshadow::kBool) { + return type1; } else if (is_unsigned_int(type1) || is_unsigned_int(type2)) { if (bits_of(type1) < bits_of(type2)) { if (type1 == mshadow::kInt8 && type2 == mshadow::kUint16) { @@ -991,6 +995,10 @@ inline int type_promotion(const int type1, const int type2) { return mshadow::kInt64; } else if (type1 == mshadow::kInt16 && type2 == mshadow::kUint32) { return mshadow::kInt64; + } else if (type2 == mshadow::kUint64) { + LOG(FATAL) << "Unsupported type promotions between " + << mshadow::dtype_string(type1) << " and " + << mshadow::dtype_string(type2); } else { return type2; } @@ -1001,6 +1009,10 @@ inline int type_promotion(const int type1, const int type2) { return mshadow::kInt64; } else if (type2 == mshadow::kInt16 && type1 == mshadow::kUint32) { return mshadow::kInt64; + } else if (type1 == mshadow::kUint64) { + LOG(FATAL) << "Unsupported type promotions between " + << mshadow::dtype_string(type1) << " and " + << mshadow::dtype_string(type2); } else { return type1; } @@ -1015,12 +1027,10 @@ inline int type_promotion(const int type1, const int type2) { return mshadow::kInt64; } } - } else if (type1 == mshadow::kBool) { - return type2; - } else if (type2 == mshadow::kBool) { - return type1; } - LOG(FATAL) << "should not reach here "; + LOG(FATAL) << "Unsupported type promotions between " + << mshadow::dtype_string(type1) << " and " + << mshadow::dtype_string(type2); return -1; } diff --git a/src/ndarray/ndarray_function-inl.h b/src/ndarray/ndarray_function-inl.h index c1d81191dbee..8101bf2a624f 100644 --- a/src/ndarray/ndarray_function-inl.h +++ b/src/ndarray/ndarray_function-inl.h @@ -402,7 +402,7 @@ void EvalRandom(const real_t& mu, template <> void Eval(const real_t& rhs, TBlob* ret, RunContext ctx) { mshadow::Stream* s = ctx.get_stream(); - MSHADOW_TYPE_SWITCH_WITH_BOOL( + MSHADOW_TYPE_SWITCH_EXT_WITH_BOOL( ret->type_flag_, DType, { ret->FlatTo2D(s) = DType(rhs); }); } diff --git a/tests/python/unittest/test_numpy_op.py b/tests/python/unittest/test_numpy_op.py index 44fda6afad8e..89cf463fbdab 100644 --- a/tests/python/unittest/test_numpy_op.py +++ b/tests/python/unittest/test_numpy_op.py @@ -11270,6 +11270,20 @@ def test_np_can_cast(input1, input2): assert np.can_cast(input1, input2) == onp.can_cast(np_input1, np_input2) +@use_np +@pytest.mark.parametrize('nums', [1, 2, 3, 4, 10, 100]) +def test_np_result_type(nums): + PICK_LIST = np.numeric_dtypes + np.boolean_dtypes + [np.ones((1,), dtype=d) for d in np.numeric_dtypes + np.boolean_dtypes] + import random + inputs = [random.choice(PICK_LIST) for _ in range(nums)] + + try: + promoted = np.result_type(*inputs) + except Exception as e: + with pytest.raises(TypeError): + promoted = np.result_type(*inputs) + + @use_np @retry(3) @pytest.mark.parametrize('func,func2,dtypes,ref_grad,low,high', [ From bcb88e321e012486f692c1cb2271a1c6f0c2d066 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Wed, 10 Nov 2021 15:38:28 -0800 Subject: [PATCH 30/39] add binary tests --- .../numpy/np_elemwise_broadcast_op_lae.cc | 2 +- tests/python/unittest/test_numpy_op.py | 117 +++++++++++++++++- 2 files changed, 116 insertions(+), 3 deletions(-) diff --git a/src/operator/numpy/np_elemwise_broadcast_op_lae.cc b/src/operator/numpy/np_elemwise_broadcast_op_lae.cc index 0298d1dbd3ff..651fbf6fe2eb 100644 --- a/src/operator/numpy/np_elemwise_broadcast_op_lae.cc +++ b/src/operator/numpy/np_elemwise_broadcast_op_lae.cc @@ -68,7 +68,7 @@ NNVM_REGISTER_OP(_backward_npi_logaddexp) }) .set_attr( "FCompute", - BinaryBroadcastBackwardUseIn); + NumpyBinaryBackwardUseIn); MXNET_OPERATOR_REGISTER_BINARY(_backward_npi_logaddexp_scalar) .add_arguments(NumpyBinaryScalarParam::__FIELDS__()) diff --git a/tests/python/unittest/test_numpy_op.py b/tests/python/unittest/test_numpy_op.py index 89cf463fbdab..d38ebfde1fa2 100644 --- a/tests/python/unittest/test_numpy_op.py +++ b/tests/python/unittest/test_numpy_op.py @@ -11285,7 +11285,6 @@ def test_np_result_type(nums): @use_np -@retry(3) @pytest.mark.parametrize('func,func2,dtypes,ref_grad,low,high', [ ('abs', 'abs', 'numeric', lambda x: -1. * (x < 0) + (x > 0), -1.0, 1.0), ('acos', 'arccos', 'floating-point', lambda x: -1. / (1. - x ** 2.) ** (1. / 2.), -1.0, 1.0), @@ -11375,7 +11374,7 @@ def array_values(low, high, shape): np_func = getattr(onp, func2) mx_out = getattr(mx.np, func)(mx_test_data) assert mx_out.shape == np_out.shape - assert mx_out.dtype == dtype + assert np.result_type(mx_out) == dtype assert_almost_equal(mx_out.asnumpy(), np_out, rtol=rtol, atol=1e-5) assertRaises(NotImplementedError, getattr(np, func), mx_test_data, where=False) @@ -11386,3 +11385,117 @@ def array_values(low, high, shape): assertRaises(TypeError, getattr(np, func), mx_test_data, casting='mxnet') assertRaises(NotImplementedError, getattr(np, func), mx_test_data, order='C') assertRaises(NotImplementedError, getattr(np, func), mx_test_data, order='mxnet') + + +@use_np +@pytest.mark.parametrize('func,func2,promoted,dtypes,ref_grad_a,ref_grad_b,low,high', [ + ('add', 'add', True, 'numeric', lambda y, x1, x2: onp.ones(y.shape), None, -1.0, 1.0), + ('atan2', 'arctan2', True, 'floating-point', lambda y, x1, x2: x2 / (onp.square(x1) + onp.square(x2)), + lambda y, x1, x2: -x1 / (onp.square(x1) + onp.square(x2)), -1, 1), + ('bitwise_and', 'bitwise_and', True, 'integer or boolean', None, None, -100, 100), + ('bitwise_or', 'bitwise_or', True, 'integer or boolean', None, None, -100, 100), + ('bitwise_xor', 'bitwise_xor', True, 'integer or boolean', None, None, -100, 100), + ('divide', 'divide', True, 'floating-point', lambda y, x1, x2: onp.ones(y.shape) / x2, + lambda y, x1, x2: -x1 / (x2 * x2), 0.1, 1.0), + ('equal', 'equal', False, 'all', None, None, 0.0, 2.0), + ('floor_divide', 'floor_divide', True, 'numeric', lambda y, x1, x2: onp.zeros(y.shape), + lambda y, x1, x2: onp.zeros(y.shape), 2.0, 10.0), + ('greater', 'greater', False, 'numeric', None, None, 0.0, 2.0), + ('greater_equal', 'greater_equal', False, 'numeric', None, None, 0.0, 2.0), + ('less', 'less', False, 'numeric', None, None, 0.0, 2.0), + ('less_equal', 'less_equal', False, 'numeric', None, None, 0.0, 2.0), + ('logaddexp', 'logaddexp', True, 'floating-point', lambda y, x1, x2: onp.exp(x1) / (onp.exp(x1) + onp.exp(x2)), + lambda y, x1, x2: onp.exp(x2) / (onp.exp(x1) + onp.exp(x2)), -10, 10), + ('logical_and', 'logical_and', False, 'boolean', None, None, -100, 100), + ('logical_or', 'logical_or', False, 'boolean', None, None, -100, 100), + ('logical_xor', 'logical_xor', False, 'boolean', None, None, -100, 100), + ('multiply', 'multiply', True, 'numeric', lambda y, x1, x2: onp.broadcast_to(x2, y.shape), + lambda y, x1, x2: onp.broadcast_to(x1, y.shape), -1.0, 1.0), + ('not_equal', 'not_equal', False, 'all', None, None, 0.0, 2.0), + ('pow', 'power', True, 'floating-point', lambda y, x1, x2: onp.power(x1, x2 - 1.0) * x2, + lambda y, x1, x2: onp.power(x1, x2) * onp.log(x1), 1.0, 3.0), + ('subtract', 'subtract', True, 'numeric', lambda y, x1, x2: onp.ones(y.shape), + lambda y, x1, x2: -onp.ones(y.shape), -1.0, 1.0), +]) +@pytest.mark.parametrize('lshape,rshape', [ + ((3, 2), (3, 2)), + ((3, 2), (3, 1)), + ((3, 1), (3, 0)), + ((0, 2), (1, 2)), + ((2, 3, 4), (3, 1)), + ((2, 3), ()), + ((), (2, 3)) +]) +def test_np_standard_binary_funcs(func, func2, promoted, dtypes, ref_grad_a, ref_grad_b, low, high, lshape, rshape): + class TestStandardBinary(HybridBlock): + def __init__(self, func): + super(TestStandardBinary, self).__init__() + self._func = func + + def forward(self, a, b,): + return getattr(np, self._func)(a, b) + + type_mapping = { + 'floating-point': np.floating_dtypes, + 'numeric': np.numeric_dtypes, + 'integer or boolean': np.integer_dtypes + np.boolean_dtypes, + 'boolean': np.boolean_dtypes, + 'all': np.numeric_dtypes + np.boolean_dtypes, + } + + def array_values(low, high, shape): + for d in np.integer_dtypes + np.boolean_dtypes + np.floating_dtypes: + yield onp.random.uniform(low, high, shape).astype(d), d + + + for (left_value, ltype) in array_values(low, high, lshape): + for (right_value, rtype) in array_values(low, high, rshape): + if ltype in type_mapping[dtypes] and rtype in type_mapping[dtypes]: + try: + promote_type = np.result_type(ltype, rtype) + except Exception as e: + # Unkown type promotion between two types + continue + rtol = 1e-2 if ltype == np.float16 or rtype == np.float16 else 1e-3 + atol = 1e-4 if ltype == np.float16 or rtype == np.float16 else 1e-5 + mx_left_value = np.array(left_value, dtype=ltype) + mx_right_value = np.array(right_value, dtype=rtype) + mx_func = TestStandardBinary(func) + np_func = getattr(onp, func2) + for hybridize in [True, False]: + if hybridize: + mx_func.hybridize() + if ref_grad_a: + mx_left_value.attach_grad() + mx_right_value.attach_grad() + np_out = np_func(left_value, right_value) + with mx.autograd.record(): + y = mx_func(mx_left_value, mx_right_value) + assert y.shape == np_out.shape + assert_almost_equal(y.asnumpy(), np_out.astype(y.dtype), rtol=rtol, atol=atol, + use_broadcast=False, equal_nan=True) + + if ref_grad_a and ltype in np.floating_dtypes and rtype in np.floating_dtypes: + y.backward() + assert_almost_equal(mx_left_value.grad.asnumpy(), + collapse_sum_like(ref_grad_a(y.asnumpy(), left_value, right_value), mx_left_value.shape), + rtol=1e-1, atol=1e-2, equal_nan=True, use_broadcast=False) + if ref_grad_b is None: + assert_almost_equal(mx_right_value.grad.asnumpy(), + collapse_sum_like(ref_grad_a(y.asnumpy(), right_value, left_value), mx_right_value.shape), + rtol=1e-1, atol=1e-2, equal_nan=True, use_broadcast=False) + else: + assert_almost_equal(mx_right_value.grad.asnumpy(), + collapse_sum_like(ref_grad_b(y.asnumpy(), left_value, right_value), mx_right_value.shape), + rtol=1e-1, atol=1e-2, equal_nan=True, use_broadcast=False) + + np_out = getattr(onp, func2)(left_value, right_value) + mx_out = getattr(np, func)(mx_left_value, mx_right_value) + assert mx_out.shape == np_out.shape + if promoted: + assert np.result_type(ltype, rtype) == mx_out.dtype + else: + assert mx_out.dtype == np.bool_ + assert_almost_equal(mx_out.asnumpy(), np_out.astype(mx_out.dtype), rtol=1e-3, atol=1e-5, + use_broadcast=False, equal_nan=True) + From a2e221d2c14fdd6df08e936b0c4e19c93e56d10d Mon Sep 17 00:00:00 2001 From: barry-jin Date: Wed, 10 Nov 2021 15:39:36 -0800 Subject: [PATCH 31/39] fix lint --- src/common/utils.h | 13 +++++-------- 1 file changed, 5 insertions(+), 8 deletions(-) diff --git a/src/common/utils.h b/src/common/utils.h index 269dbbbc2b76..7cf54ca2b4fa 100644 --- a/src/common/utils.h +++ b/src/common/utils.h @@ -996,9 +996,8 @@ inline int type_promotion(const int type1, const int type2) { } else if (type1 == mshadow::kInt16 && type2 == mshadow::kUint32) { return mshadow::kInt64; } else if (type2 == mshadow::kUint64) { - LOG(FATAL) << "Unsupported type promotions between " - << mshadow::dtype_string(type1) << " and " - << mshadow::dtype_string(type2); + LOG(FATAL) << "Unsupported type promotions between " << mshadow::dtype_string(type1) + << " and " << mshadow::dtype_string(type2); } else { return type2; } @@ -1010,9 +1009,8 @@ inline int type_promotion(const int type1, const int type2) { } else if (type2 == mshadow::kInt16 && type1 == mshadow::kUint32) { return mshadow::kInt64; } else if (type1 == mshadow::kUint64) { - LOG(FATAL) << "Unsupported type promotions between " - << mshadow::dtype_string(type1) << " and " - << mshadow::dtype_string(type2); + LOG(FATAL) << "Unsupported type promotions between " << mshadow::dtype_string(type1) + << " and " << mshadow::dtype_string(type2); } else { return type1; } @@ -1028,8 +1026,7 @@ inline int type_promotion(const int type1, const int type2) { } } } - LOG(FATAL) << "Unsupported type promotions between " - << mshadow::dtype_string(type1) << " and " + LOG(FATAL) << "Unsupported type promotions between " << mshadow::dtype_string(type1) << " and " << mshadow::dtype_string(type2); return -1; } From b1de51b0aaef69c0c506eb857c9ef5e1ee38dcd9 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Wed, 10 Nov 2021 18:54:09 -0800 Subject: [PATCH 32/39] update --- src/operator/numpy/np_elemwise_broadcast_op.h | 2 +- src/operator/numpy/np_elemwise_broadcast_op_extended.cc | 2 +- tests/python/unittest/test_numpy_op.py | 1 + 3 files changed, 3 insertions(+), 2 deletions(-) diff --git a/src/operator/numpy/np_elemwise_broadcast_op.h b/src/operator/numpy/np_elemwise_broadcast_op.h index fb79dc70080b..29fe12150ff0 100644 --- a/src/operator/numpy/np_elemwise_broadcast_op.h +++ b/src/operator/numpy/np_elemwise_broadcast_op.h @@ -393,7 +393,7 @@ void MixedBinaryBroadcastCompute(const nnvm::NodeAttrs& attrs, BinaryBroadcastCompute( attrs, ctx, {lhs, temp_tblob.reshape(rhs.shape_)}, req, outputs); } else if (rhs.type_flag_ == out.type_flag_) { - MXNET_INT_TYPE_SWITCH(rhs.type_flag_, RType, { + MXNET_INT_TYPE_SWITCH_EXT_WITH_BOOL(rhs.type_flag_, RType, { Tensor temp_tensor = ctx.requested[0].get_space_typed(Shape1(lhs.Size()), s); temp_tblob = TBlob(temp_tensor); diff --git a/src/operator/numpy/np_elemwise_broadcast_op_extended.cc b/src/operator/numpy/np_elemwise_broadcast_op_extended.cc index 33fddab3149b..949aad67ab3e 100644 --- a/src/operator/numpy/np_elemwise_broadcast_op_extended.cc +++ b/src/operator/numpy/np_elemwise_broadcast_op_extended.cc @@ -233,7 +233,7 @@ NNVM_REGISTER_OP(_backward_npi_arctan2) }) .set_attr( "FCompute", - BinaryBroadcastBackwardUseIn); + NumpyBinaryBackwardUseIn); MXNET_OPERATOR_REGISTER_NP_BINARY_SCALAR(_npi_arctan2_scalar) .set_attr("FCompute", BinaryScalarOp::Compute) diff --git a/tests/python/unittest/test_numpy_op.py b/tests/python/unittest/test_numpy_op.py index d38ebfde1fa2..09b87f6e53cc 100644 --- a/tests/python/unittest/test_numpy_op.py +++ b/tests/python/unittest/test_numpy_op.py @@ -11388,6 +11388,7 @@ def array_values(low, high, shape): @use_np +@pytest.mark.flaky @pytest.mark.parametrize('func,func2,promoted,dtypes,ref_grad_a,ref_grad_b,low,high', [ ('add', 'add', True, 'numeric', lambda y, x1, x2: onp.ones(y.shape), None, -1.0, 1.0), ('atan2', 'arctan2', True, 'floating-point', lambda y, x1, x2: x2 / (onp.square(x1) + onp.square(x2)), From f6ff8c3c77e5c44e1fdef6f0681d00e01f6ddcb7 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Wed, 10 Nov 2021 21:48:09 -0800 Subject: [PATCH 33/39] update rtol, atol --- tests/python/unittest/test_numpy_op.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/python/unittest/test_numpy_op.py b/tests/python/unittest/test_numpy_op.py index 09b87f6e53cc..3d4aeb2f6d57 100644 --- a/tests/python/unittest/test_numpy_op.py +++ b/tests/python/unittest/test_numpy_op.py @@ -11497,6 +11497,6 @@ def array_values(low, high, shape): assert np.result_type(ltype, rtype) == mx_out.dtype else: assert mx_out.dtype == np.bool_ - assert_almost_equal(mx_out.asnumpy(), np_out.astype(mx_out.dtype), rtol=1e-3, atol=1e-5, + assert_almost_equal(mx_out.asnumpy(), np_out.astype(mx_out.dtype), rtol=rtol, atol=atol, use_broadcast=False, equal_nan=True) From 811cd6314dc1592869a194a797e70854e805ed9b Mon Sep 17 00:00:00 2001 From: barry-jin Date: Thu, 11 Nov 2021 14:31:21 -0800 Subject: [PATCH 34/39] update rtc types --- src/common/cuda/rtc/util-inl.h | 199 ++++++++++++++++++++++++++++++++- 1 file changed, 198 insertions(+), 1 deletion(-) diff --git a/src/common/cuda/rtc/util-inl.h b/src/common/cuda/rtc/util-inl.h index f294aa0ef2eb..347167ca77bb 100644 --- a/src/common/cuda/rtc/util-inl.h +++ b/src/common/cuda/rtc/util-inl.h @@ -37,6 +37,10 @@ using uint8 = unsigned char; using int8 = char; using int32 = int; using int64 = long long; +using int16 = short; +using uint16 = unsigned short; +using uint32 = unsigned int; +using uint64 = unsigned long long; static_assert(sizeof(float32) == 4, "Size of float32 is expected to be 4B"); static_assert(sizeof(float64) == 8, "Size of float64 is expected to be 8B"); @@ -45,6 +49,10 @@ static_assert(sizeof(uint8) == 1, "Size of uint8 is expected to be 1B"); static_assert(sizeof(int8) == 1, "Size of int8 is expected to be 1B"); static_assert(sizeof(int32) == 4, "Size of int32 is expected to be 4B"); static_assert(sizeof(int64) == 8, "Size of int64 is expected to be 8B"); +static_assert(sizeof(int16) == 2, "Size of int16 is expected to be 2B"); +static_assert(sizeof(uint16) == 2, "Size of uint16 is expected to be 2B"); +static_assert(sizeof(uint32) == 4, "Size of uint32 is expected to be 4B"); +static_assert(sizeof(uint64) == 8, "Size of uint64 is expected to be 8B"); )code" #if MSHADOW_INT64_TENSOR_SIZE == 1 @@ -129,7 +137,11 @@ struct true_type { // is_integral template struct is_integral : false_type {}; template <> struct is_integral : true_type {}; +template <> struct is_integral : true_type {}; +template <> struct is_integral : true_type {}; +template <> struct is_integral : true_type {}; template <> struct is_integral : true_type {}; +template <> struct is_integral : true_type {}; template <> struct is_integral : true_type {}; template <> struct is_integral : true_type {}; template <> struct is_integral : true_type {}; @@ -138,6 +150,9 @@ template <> struct is_integral : true_type {}; // is_unsigned template struct is_unsigned : false_type {}; template <> struct is_unsigned : true_type {}; +template <> struct is_unsigned : true_type {}; +template <> struct is_unsigned : true_type {}; +template <> struct is_unsigned : true_type {}; template <> struct is_unsigned : true_type {}; template <> struct is_unsigned : true_type {}; @@ -211,19 +226,161 @@ struct mixed_type_helper::value>:: template struct mixed_type_helper::value && is_integral::value && + is_unsigned::value && + is_unsigned::value && !is_same::value && - sizeof(T) <= sizeof(U)>::type> { + sizeof(T) < sizeof(U)>::type> { + using type = U; +}; + +template +struct mixed_type_helper::value && + is_integral::value && + !is_unsigned::value && + !is_unsigned::value && + !is_same::value && + sizeof(T) < sizeof(U)>::type> { + using type = U; +}; + +template +struct mixed_type_helper::value && + is_integral::value && + is_unsigned::value && + !is_unsigned::value && + !is_same::value && + sizeof(T) < sizeof(U)>::type> { using type = U; }; template struct mixed_type_helper::value && is_integral::value && + is_unsigned::value && + is_unsigned::value && !is_same::value && sizeof(T) < sizeof(U)>::type> { using type = U; }; +template +struct mixed_type_helper::value && + is_integral::value && + !is_unsigned::value && + !is_unsigned::value && + !is_same::value && + sizeof(T) < sizeof(U)>::type> { + using type = U; +}; + +template +struct mixed_type_helper::value && + is_integral::value && + is_unsigned::value && + !is_unsigned::value && + !is_same::value && + sizeof(T) < sizeof(U)>::type> { + using type = U; +}; + +template +struct mixed_type_helper::value && + is_integral::value && + !is_same::value && + is_same::value>::type> { + using type = U; +}; + +template<> +struct mixed_type_helper { + using type = int16; +}; + +template<> +struct mixed_type_helper { + using type = int16; +}; + +template<> +struct mixed_type_helper { + using type = int32; +}; + +template<> +struct mixed_type_helper { + using type = int32; +}; + +template<> +struct mixed_type_helper { + using type = int64; +}; + +template<> +struct mixed_type_helper { + using type = int64; +}; + +template<> +struct mixed_type_helper { + using type = int32; +}; + +template<> +struct mixed_type_helper { + using type = int32; +}; + +template<> +struct mixed_type_helper { + using type = int64; +}; + +template<> +struct mixed_type_helper { + using type = int64; +}; + +template<> +struct mixed_type_helper { + using type = int64; +}; + +template<> +struct mixed_type_helper { + using type = int64; +}; + +template<> +struct mixed_type_helper { + using type = index_t; +}; + +template<> +struct mixed_type_helper { + using type = index_t; +}; + +template<> +struct mixed_type_helper { + using type = index_t; +}; + +template<> +struct mixed_type_helper { + using type = index_t; +}; + +template<> +struct mixed_type_helper { + using type = uint64; +}; + +template<> +struct mixed_type_helper { + using type = uint64; +}; + template struct mixed_type_helper::value && sizeof(T) < sizeof(bool_t)>::type> { @@ -472,11 +629,31 @@ template<> __device__ inline uint8 MinValue(void) { return 0; } +/*! \brief minimum value of uint16 */ +template<> +__device__ inline uint16 MinValue(void) { + return 0; +} +/*! \brief minimum value of uint32 */ +template<> +__device__ inline uint32 MinValue(void) { + return 0; +} +/*! \brief minimum value of uint64 */ +template<> +__device__ inline uint64 MinValue(void) { + return 0; +} /*! \brief minimum value of int8_t */ template<> __device__ inline int8 MinValue(void) { return -128; } +/*! \brief minimum value of int16 */ +template<> +__device__ inline int16 MinValue(void) { + return -32768; +} /*! \brief minimum value of int32 */ template<> __device__ inline int32 MinValue(void) { @@ -538,11 +715,31 @@ template<> __device__ inline uint8 MaxValue(void) { return 255; } +/*! \brief maximum value of uint16 */ +template<> +__device__ inline uint16 MaxValue(void) { + return 65535; +} +/*! \brief maximum value of uint32 */ +template<> +__device__ inline uint32 MaxValue(void) { + return 4294967295; +} +/*! \brief maximum value of uint64 */ +template<> +__device__ inline uint64 MaxValue(void) { + return 18446744073709551615LL; +} /*! \brief maximum value of int8 */ template<> __device__ inline int8 MaxValue(void) { return 127; } +/*! \brief maximum value of int16 */ +template<> +__device__ inline int16 MaxValue(void) { + return 32767; +} /*! \brief maximum value of int32 */ template<> __device__ inline int32 MaxValue(void) { From 6e248ff3c4a86c21d082e798aaa46c1185bebc8e Mon Sep 17 00:00:00 2001 From: barry-jin Date: Thu, 11 Nov 2021 14:40:41 -0800 Subject: [PATCH 35/39] fix floor,ceil,trunc --- python/mxnet/ndarray/numpy/_op.py | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/python/mxnet/ndarray/numpy/_op.py b/python/mxnet/ndarray/numpy/_op.py index bd6503f55c0a..e0b20abfcf76 100644 --- a/python/mxnet/ndarray/numpy/_op.py +++ b/python/mxnet/ndarray/numpy/_op.py @@ -3757,7 +3757,7 @@ def ceil(x, out=None, **kwargs): >>> a array(4.) """ - if _np.issubdtype(x.dtype, _np.integer): + if isinstance(x, NDArray) and _np.issubdtype(x.dtype, _np.integer): return x return _pure_unary_func_helper(x, _api_internal.ceil, _np.ceil, out=out, **kwargs) @@ -3798,7 +3798,7 @@ def floor(x, out=None, **kwargs): >>> a array(3.) """ - if _np.issubdtype(x.dtype, _np.integer): + if isinstance(x, NDArray) and _np.issubdtype(x.dtype, _np.integer): return x return _pure_unary_func_helper(x, _api_internal.floor, _np.floor, out=out, **kwargs) @@ -3945,7 +3945,7 @@ def trunc(x, out=None, **kwargs): >>> np.trunc(a) array([-1., -1., -0., 0., 1., 1., 2.]) """ - if _np.issubdtype(x.dtype, _np.integer): + if isinstance(x, NDArray) and _np.issubdtype(x.dtype, _np.integer): return x return _pure_unary_func_helper(x, _api_internal.trunc, _np.trunc, out=out, **kwargs) From ce1af6d53893cdf0f00d3b9a33a84e402b097ef5 Mon Sep 17 00:00:00 2001 From: barry-jin Date: Mon, 15 Nov 2021 17:30:47 -0800 Subject: [PATCH 36/39] update rtc type promotion --- src/common/cuda/rtc/util-inl.h | 27 +++++++-------------------- 1 file changed, 7 insertions(+), 20 deletions(-) diff --git a/src/common/cuda/rtc/util-inl.h b/src/common/cuda/rtc/util-inl.h index 347167ca77bb..66e23518b865 100644 --- a/src/common/cuda/rtc/util-inl.h +++ b/src/common/cuda/rtc/util-inl.h @@ -361,26 +361,6 @@ struct mixed_type_helper { using type = index_t; }; -template<> -struct mixed_type_helper { - using type = index_t; -}; - -template<> -struct mixed_type_helper { - using type = index_t; -}; - -template<> -struct mixed_type_helper { - using type = uint64; -}; - -template<> -struct mixed_type_helper { - using type = uint64; -}; - template struct mixed_type_helper::value && sizeof(T) < sizeof(bool_t)>::type> { @@ -399,6 +379,13 @@ struct mixed_type_helper::value && using type = T; }; +template +struct mixed_type_helper::value && + !is_same::value && + sizeof(T) == sizeof(bool_t)>::type> { + using type = T; +}; + template struct multi_mixed_type_helper; From 5fa0297e448ca36c971805984f80258e7e00b96c Mon Sep 17 00:00:00 2001 From: barry-jin Date: Tue, 16 Nov 2021 10:02:36 -0800 Subject: [PATCH 37/39] update tests --- .github/workflows/os_x_staticbuild.yml | 26 +++++++++++++------------- ci/docker/install/requirements | 1 + ci/docker/runtime_functions.sh | 26 +++++++++++++------------- 3 files changed, 27 insertions(+), 26 deletions(-) diff --git a/.github/workflows/os_x_staticbuild.yml b/.github/workflows/os_x_staticbuild.yml index c6aad0676d0b..37b28f3b012e 100644 --- a/.github/workflows/os_x_staticbuild.yml +++ b/.github/workflows/os_x_staticbuild.yml @@ -65,24 +65,24 @@ jobs: git checkout c1dba80a196a03f880d2e0a998a272fb3867b720 export ARRAY_API_TESTS_MODULE=mxnet.numpy pytest export DMLC_LOG_STACK_TRACE_DEPTH=100 - python3 -m pytest --durations=50 --verbose array_api_tests/test_creation_functions.py - python3 -m pytest --durations=50 --verbose array_api_tests/test_indexing.py - python3 -m pytest --durations=50 --verbose array_api_tests/test_constants.py - python3 -m pytest --durations=50 --verbose array_api_tests/test_elementwise_functions.py - python3 -m pytest --durations=50 --verbose array_api_tests/test_broadcasting.py - python3 -m pytest --durations=50 --verbose \ + python3 -m pytest --reruns 3 --durations=50 --verbose array_api_tests/test_creation_functions.py + python3 -m pytest --reruns 3 --durations=50 --verbose array_api_tests/test_indexing.py + python3 -m pytest --reruns 3 --durations=50 --verbose array_api_tests/test_constants.py + python3 -m pytest --reruns 3 --durations=50 --verbose array_api_tests/test_elementwise_functions.py + python3 -m pytest --reruns 3 --durations=50 --verbose array_api_tests/test_broadcasting.py + python3 -m pytest --reruns 3 --durations=50 --verbose \ array_api_tests/test_type_promotion.py::test_elementwise_function_two_arg_bool_type_promotion - python3 -m pytest --durations=50 --verbose \ + python3 -m pytest --reruns 3 --durations=50 --verbose \ array_api_tests/test_type_promotion.py::test_elementwise_function_two_arg_promoted_type_promotion - python3 -m pytest --durations=50 --verbose \ + python3 -m pytest --reruns 3 --durations=50 --verbose \ array_api_tests/test_type_promotion.py::test_elementwise_function_one_arg_bool - python3 -m pytest --durations=50 --verbose \ + python3 -m pytest --reruns 3 --durations=50 --verbose \ array_api_tests/test_type_promotion.py::test_elementwise_function_one_arg_type_promotion - python3 -m pytest --durations=50 --verbose \ + python3 -m pytest --reruns 3 --durations=50 --verbose \ array_api_tests/test_type_promotion.py::test_operator_one_arg_type_promotion - python3 -m pytest --durations=50 --verbose \ + python3 -m pytest --reruns 3 --durations=50 --verbose \ array_api_tests/test_type_promotion.py::test_operator_two_arg_bool_promotion - python3 -m pytest --durations=50 --verbose \ + python3 -m pytest --reruns 3 --durations=50 --verbose \ array_api_tests/test_type_promotion.py::test_operator_two_arg_promoted_promotion - python3 -m pytest --durations=50 --verbose \ + python3 -m pytest --reruns 3 --durations=50 --verbose \ array_api_tests/test_type_promotion.py::test_operator_inplace_two_arg_promoted_promotion diff --git a/ci/docker/install/requirements b/ci/docker/install/requirements index 21f10b92cba8..7b8e2d033591 100644 --- a/ci/docker/install/requirements +++ b/ci/docker/install/requirements @@ -41,6 +41,7 @@ pytest-env==0.6.2 pytest-cov==2.10.1 pytest-xdist==2.1.0 pytest-timeout==1.4.2 +pytest-rerunfailures==10.2 flaky==3.7.0 setuptools==49.6.0 # https://github.com/pypa/setuptools/issues/2352 wheel diff --git a/ci/docker/runtime_functions.sh b/ci/docker/runtime_functions.sh index c731c0d62bef..8ffb49d24141 100755 --- a/ci/docker/runtime_functions.sh +++ b/ci/docker/runtime_functions.sh @@ -872,26 +872,26 @@ unittest_array_api_standardization() { # when cython is enabled export MXNET_ENABLE_CYTHON=0 export DMLC_LOG_STACK_TRACE_DEPTH=100 - python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_creation_functions.py - python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_indexing.py - python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_elementwise_functions.py - python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_constants.py - python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_broadcasting.py - python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ + python3 -m pytest --reruns 3 --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_creation_functions.py + python3 -m pytest --reruns 3 --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_indexing.py + python3 -m pytest --reruns 3 --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_elementwise_functions.py + python3 -m pytest --reruns 3 --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_constants.py + python3 -m pytest --reruns 3 --durations=50 --cov-report xml:tests_api.xml --verbose array_api_tests/test_broadcasting.py + python3 -m pytest --reruns 3 --durations=50 --cov-report xml:tests_api.xml --verbose \ array_api_tests/test_type_promotion.py::test_elementwise_function_two_arg_bool_type_promotion - python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ + python3 -m pytest --reruns 3 --durations=50 --cov-report xml:tests_api.xml --verbose \ array_api_tests/test_type_promotion.py::test_elementwise_function_two_arg_promoted_type_promotion - python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ + python3 -m pytest --reruns 3 --durations=50 --cov-report xml:tests_api.xml --verbose \ array_api_tests/test_type_promotion.py::test_elementwise_function_one_arg_bool - python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ + python3 -m pytest --reruns 3 --durations=50 --cov-report xml:tests_api.xml --verbose \ array_api_tests/test_type_promotion.py::test_elementwise_function_one_arg_type_promotion - python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ + python3 -m pytest --reruns 3 --durations=50 --cov-report xml:tests_api.xml --verbose \ array_api_tests/test_type_promotion.py::test_operator_one_arg_type_promotion - python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ + python3 -m pytest --reruns 3 --durations=50 --cov-report xml:tests_api.xml --verbose \ array_api_tests/test_type_promotion.py::test_operator_two_arg_bool_promotion - python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ + python3 -m pytest --reruns 3 --durations=50 --cov-report xml:tests_api.xml --verbose \ array_api_tests/test_type_promotion.py::test_operator_two_arg_promoted_promotion - python3 -m pytest --durations=50 --cov-report xml:tests_api.xml --verbose \ + python3 -m pytest --reruns 3 --durations=50 --cov-report xml:tests_api.xml --verbose \ array_api_tests/test_type_promotion.py::test_operator_inplace_two_arg_promoted_promotion popd } From 7d77e070c741a0fa1c42415f636cb5200124253f Mon Sep 17 00:00:00 2001 From: barry-jin Date: Wed, 17 Nov 2021 10:25:16 -0800 Subject: [PATCH 38/39] update mod --- src/operator/mshadow_op.h | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) diff --git a/src/operator/mshadow_op.h b/src/operator/mshadow_op.h index 3fbc5c6901b0..e2d6ea9791e4 100644 --- a/src/operator/mshadow_op.h +++ b/src/operator/mshadow_op.h @@ -1004,11 +1004,9 @@ struct mod : public mxnet_op::tunable { return DType(0); } else if (b < DType(0)) { if (a < DType(0)) { - DType ret = DType(-::fmod(-static_cast(a), -static_cast(b))); - if (ret == 0) { - return -ret; - } - return ret; + return DType(-::fmod(-static_cast(a), -static_cast(b))); + } else if (a == DType(0)){ + return -DType(0); } else { DType ret = DType( ::fmod(static_cast(a), -static_cast(b)) + From 5944c4b5c3e1349536416a8e6f44fa41b34c14aa Mon Sep 17 00:00:00 2001 From: barry-jin Date: Wed, 17 Nov 2021 12:51:25 -0800 Subject: [PATCH 39/39] fix lint --- 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 e2d6ea9791e4..41f1aa5d1828 100644 --- a/src/operator/mshadow_op.h +++ b/src/operator/mshadow_op.h @@ -1005,7 +1005,7 @@ struct mod : public mxnet_op::tunable { } else if (b < DType(0)) { if (a < DType(0)) { return DType(-::fmod(-static_cast(a), -static_cast(b))); - } else if (a == DType(0)){ + } else if (a == DType(0)) { return -DType(0); } else { DType ret = DType(