From 53c0d99c76ac64ff511c800bf5d5ff566c3e26df Mon Sep 17 00:00:00 2001
From: JiangZhaoh <jiang_zhaoh@163.com>
Date: Tue, 26 Nov 2019 06:37:43 +0000
Subject: [PATCH 1/5] add_op_delete

---
 python/mxnet/ndarray/numpy/_op.py             |  68 +++-
 python/mxnet/numpy/multiarray.py              |  56 ++-
 python/mxnet/numpy_dispatch_protocol.py       |   1 +
 python/mxnet/symbol/numpy/_symbol.py          |  46 ++-
 src/operator/numpy/np_delete_op-inl.h         | 340 ++++++++++++++++++
 src/operator/numpy/np_delete_op.cc            |  79 ++++
 src/operator/numpy/np_delete_op.cu            |  16 +
 .../unittest/test_numpy_interoperability.py   |  22 ++
 tests/python/unittest/test_numpy_op.py        |  75 ++++
 9 files changed, 700 insertions(+), 3 deletions(-)
 create mode 100644 src/operator/numpy/np_delete_op-inl.h
 create mode 100644 src/operator/numpy/np_delete_op.cc
 create mode 100644 src/operator/numpy/np_delete_op.cu

diff --git a/python/mxnet/ndarray/numpy/_op.py b/python/mxnet/ndarray/numpy/_op.py
index 02e42145fb18..b14e25c68ac6 100644
--- a/python/mxnet/ndarray/numpy/_op.py
+++ b/python/mxnet/ndarray/numpy/_op.py
@@ -28,7 +28,7 @@
 from . import _internal as _npi
 from ..ndarray import NDArray
 
-__all__ = ['shape', 'zeros', 'zeros_like', 'ones', 'ones_like', 'full', 'full_like', 'invert',
+__all__ = ['shape', 'zeros', 'zeros_like', 'ones', 'ones_like', 'full', 'full_like', 'invert', 'delete',
            'add', 'subtract', 'multiply', 'divide', 'mod', 'remainder', 'power', 'bitwise_not',
            'arctan2', 'sin', 'cos', 'tan', 'sinh', 'cosh', 'tanh', 'log10', 'sqrt', 'cbrt', 'abs',
            'absolute', 'exp', 'expm1', 'arcsin', 'arccos', 'arctan', 'sign', 'log', 'degrees', 'log2',
@@ -913,6 +913,72 @@ def mod(x1, x2, out=None, **kwargs):
     return _ufunc_helper(x1, x2, _npi.mod, _np.mod, _npi.mod_scalar, _npi.rmod_scalar, out)
 
 
+@set_module('mxnet.ndarray.numpy')
+def delete(arr, obj, axis=None):
+    """
+    Return a new array with sub-arrays along an axis deleted. For a one
+    dimensional array, this returns those entries not returned by
+    `arr[obj]`.
+
+    Parameters
+    ----------
+    arr : ndarray
+      Input array.
+    obj : slice, int or ndarray of ints
+      Indicate indices of sub-arrays to remove along the specified axis.
+    axis : int, optional
+      The axis along which to delete the subarray defined by `obj`.
+      If `axis` is None, `obj` is applied to the flattened array.
+
+    Returns
+    -------
+    out : ndarray
+        A copy of `arr` with the elements specified by `obj` removed. Note
+        that `delete` does not occur in-place. If `axis` is None, `out` is
+        a flattened array.
+
+    See Also
+    --------
+    insert : Insert elements into an array.
+    append : Append elements at the end of an array.
+
+    Examples
+    --------
+    >>> arr = np.array([[1,2,3,4], [5,6,7,8], [9,10,11,12]])
+    >>> arr
+    array([[ 1.,  2.,  3.,  4.],
+           [ 5.,  6.,  7.,  8.],
+           [ 9., 10., 11., 12.]])
+
+    >>> np.delete(arr, 1, 0)
+    array([[ 1.,  2.,  3.,  4.],
+           [ 9., 10., 11., 12.]])
+
+    >>> np.delete(arr, slice(None, None, 2), 1)
+    array([[ 2.,  4.],
+           [ 6.,  8.],
+           [10., 12.]])
+
+    >>> np.delete(arr, np.array([1,3,5]), None)
+    array([ 1.,  3.,  5.,  7.,  8.,  9., 10., 11., 12.])
+    >>> np.delete(arr, np.array([1,1,5]), None)
+    array([ 1.,  3.,  4.,  5.,  7.,  8.,  9., 10., 11., 12.])
+    """
+    if not isinstance(arr, NDArray):
+        raise TypeError("'arr' can not support type {}".format(str(type(arr))))
+    if isinstance(obj, slice):
+        start = obj.start
+        stop = obj.stop
+        step = 1 if obj.step is None else obj.step
+        return _npi.delete(arr, start=start, stop=stop, step=step, axis=axis)
+    elif isinstance(obj, integer_types):
+        return _npi.delete(arr, int_ind=obj, axis=axis)
+    elif isinstance(obj, NDArray):
+        return _npi.delete(arr, obj, axis=axis)
+    else:
+        raise TypeError("'obj' can not support type {}".format(str(type(obj))))
+
+
 @set_module('mxnet.ndarray.numpy')
 @wrap_np_binary_func
 def remainder(x1, x2, out=None):
diff --git a/python/mxnet/numpy/multiarray.py b/python/mxnet/numpy/multiarray.py
index 4910b4d6b925..7cca23649d70 100644
--- a/python/mxnet/numpy/multiarray.py
+++ b/python/mxnet/numpy/multiarray.py
@@ -47,7 +47,7 @@
 from ..ndarray.ndarray import _storage_type
 
 __all__ = ['ndarray', 'empty', 'array', 'shape', 'zeros', 'zeros_like', 'ones', 'ones_like', 'full', 'full_like',
-           'add', 'subtract', 'multiply', 'divide', 'mod', 'remainder', 'power', 'bitwise_not',
+           'add', 'subtract', 'multiply', 'divide', 'mod', 'remainder', 'power', 'bitwise_not', 'delete',
            'arctan2', 'sin', 'cos', 'tan', 'sinh', 'cosh', 'tanh', 'log10', 'invert',
            'sqrt', 'cbrt', 'abs', 'absolute', 'exp', 'expm1', 'arcsin', 'arccos', 'arctan', 'sign', 'log',
            'degrees', 'log2', 'log1p', 'rint', 'radians', 'reciprocal', 'square', 'negative', 'histogram',
@@ -5848,6 +5848,60 @@ def std(a, axis=None, dtype=None, out=None, ddof=0, keepdims=False):  # pylint:
     return _npi.std(a, axis=axis, dtype=dtype, ddof=ddof, keepdims=keepdims, out=out)
 
 
+@set_module('mxnet.numpy')
+def delete(arr, obj, axis=None):
+    """
+    Return a new array with sub-arrays along an axis deleted. For a one
+    dimensional array, this returns those entries not returned by
+    `arr[obj]`.
+
+    Parameters
+    ----------
+    arr : ndarray
+      Input array.
+    obj : slice, int or ndarray of ints
+      Indicate indices of sub-arrays to remove along the specified axis.
+    axis : int, optional
+      The axis along which to delete the subarray defined by `obj`.
+      If `axis` is None, `obj` is applied to the flattened array.
+
+    Returns
+    -------
+    out : ndarray
+        A copy of `arr` with the elements specified by `obj` removed. Note
+        that `delete` does not occur in-place. If `axis` is None, `out` is
+        a flattened array.
+
+    See Also
+    --------
+    insert : Insert elements into an array.
+    append : Append elements at the end of an array.
+
+    Examples
+    --------
+    >>> arr = np.array([[1,2,3,4], [5,6,7,8], [9,10,11,12]])
+    >>> arr
+    array([[ 1.,  2.,  3.,  4.],
+           [ 5.,  6.,  7.,  8.],
+           [ 9., 10., 11., 12.]])
+
+    >>> np.delete(arr, 1, 0)
+    array([[ 1.,  2.,  3.,  4.],
+           [ 9., 10., 11., 12.]])
+
+    >>> np.delete(arr, slice(None, None, 2), 1)
+    array([[ 2.,  4.],
+           [ 6.,  8.],
+           [10., 12.]])
+
+    >>> np.delete(arr, np.array([1,3,5]), None)
+    array([ 1.,  3.,  5.,  7.,  8.,  9., 10., 11., 12.])
+    >>> np.delete(arr, np.array([1,1,5]), None)
+    array([ 1.,  3.,  4.,  5.,  7.,  8.,  9., 10., 11., 12.])
+    """
+    return _mx_nd_np.delete(arr, obj, axis=axis)
+
+
 @set_module('mxnet.numpy')
 def var(a, axis=None, dtype=None, out=None, ddof=0, keepdims=False):  # pylint: disable=too-many-arguments
     """
diff --git a/python/mxnet/numpy_dispatch_protocol.py b/python/mxnet/numpy_dispatch_protocol.py
index c7e9dd1398eb..43161409838a 100644
--- a/python/mxnet/numpy_dispatch_protocol.py
+++ b/python/mxnet/numpy_dispatch_protocol.py
@@ -125,6 +125,7 @@ def _run_with_array_ufunc_proto(*args, **kwargs):
     'transpose',
     'unique',
     'unravel_index',
+    'delete',
     'var',
     'vdot',
     'vstack',
diff --git a/python/mxnet/symbol/numpy/_symbol.py b/python/mxnet/symbol/numpy/_symbol.py
index 6efc333cc16c..24672dfcfce8 100644
--- a/python/mxnet/symbol/numpy/_symbol.py
+++ b/python/mxnet/symbol/numpy/_symbol.py
@@ -36,7 +36,7 @@
 except ImportError:
     from builtins import slice as py_slice
 
-__all__ = ['zeros', 'zeros_like', 'ones', 'ones_like', 'full_like', 'bitwise_not', 'invert',
+__all__ = ['zeros', 'zeros_like', 'ones', 'ones_like', 'full_like', 'bitwise_not', 'invert', 'delete',
            'add', 'subtract', 'multiply', 'divide', 'mod', 'remainder', 'power', 'arctan2',
            'sin', 'cos', 'tan', 'sinh', 'cosh', 'tanh', 'log10', 'sqrt', 'cbrt', 'abs', 'absolute', 'exp',
            'expm1', 'arcsin', 'arccos', 'arctan', 'sign', 'log', 'degrees', 'log2', 'log1p',
@@ -3161,6 +3161,50 @@ def arange(start, stop=None, step=1, dtype=None, ctx=None):
     return _npi.arange(start=start, stop=stop, step=step, dtype=dtype, ctx=ctx)
 
 
+@set_module('mxnet.symbol.numpy')
+def delete(arr, obj, axis=None):
+    """
+    Return a new array with sub-arrays along an axis deleted. For a one
+    dimensional array, this returns those entries not returned by
+    `arr[obj]`.
+
+    Parameters
+    ----------
+    arr : _Symbol
+      Input array.
+    obj : slice, scaler or _Symbol of ints
+      Indicate indices of sub-arrays to remove along the specified axis.
+    axis : scaler, optional
+      The axis along which to delete the subarray defined by `obj`.
+      If `axis` is None, `obj` is applied to the flattened array.
+
+    Returns
+    -------
+    out : _Symbol
+        A copy of `arr` with the elements specified by `obj` removed. Note
+        that `delete` does not occur in-place. If `axis` is None, `out` is
+        a flattened array.
+
+    See Also
+    --------
+    insert : Insert elements into an array.
+    append : Append elements at the end of an array.
+    """
+    if not isinstance(arr, Symbol):
+        raise TypeError("'arr' can not support type {}".format(str(type(arr))))
+    if isinstance(obj, slice):
+        start = obj.start
+        stop = obj.stop
+        step = 1 if obj.step is None else obj.step
+        return _npi.delete(arr, start=start, stop=stop, step=step, axis=axis)
+    elif isinstance(obj, integer_types):
+        return _npi.delete(arr, int_ind=obj, axis=axis)
+    elif isinstance(obj, Symbol):
+        return _npi.delete(arr, obj, axis=axis)
+    else:
+        raise TypeError("'obj' can not support type {}".format(str(type(obj))))
+
+
 # pylint: disable=redefined-outer-name
 @set_module('mxnet.symbol.numpy')
 def split(ary, indices_or_sections, axis=0):
diff --git a/src/operator/numpy/np_delete_op-inl.h b/src/operator/numpy/np_delete_op-inl.h
new file mode 100644
index 000000000000..be14a69c6552
--- /dev/null
+++ b/src/operator/numpy/np_delete_op-inl.h
@@ -0,0 +1,340 @@
+/*!
+ *  Copyright (c) 2019 by Contributors
+ * \file np_delete_op-inl.h
+ * \brief Function definition of delete operators
+ */
+#ifndef MXNET_OPERATOR_NUMPY_NP_DELETE_OP_INL_H_
+#define MXNET_OPERATOR_NUMPY_NP_DELETE_OP_INL_H_
+
+#include <vector>
+#include <memory>
+#include "../../common/utils.h"
+#include "../tensor/sort_op.h"
+#include "../operator_common.h"
+#include "../tensor/broadcast_reduce_op.h"
+#ifdef __CUDACC__
+#include <thrust/device_ptr.h>
+#include <thrust/device_vector.h>
+#include <thrust/functional.h>
+#include <thrust/scan.h>
+#include <thrust/sequence.h>
+#include <thrust/sort.h>
+#endif
+
+namespace mxnet {
+namespace op {
+
+struct NumpyDeleteParam : public dmlc::Parameter<NumpyDeleteParam> {
+  dmlc::optional<int> start;
+  dmlc::optional<int> stop;
+  dmlc::optional<int> step;
+  dmlc::optional<int> int_ind;
+  dmlc::optional<int> axis;
+  DMLC_DECLARE_PARAMETER(NumpyDeleteParam) {
+    DMLC_DECLARE_FIELD(start)
+    .set_default(dmlc::optional<int>())
+    .describe("If 'obj' is slice, 'start' is one of it's arguments.");
+    DMLC_DECLARE_FIELD(stop)
+    .set_default(dmlc::optional<int>())
+    .describe("If 'obj' is slice, 'stop' is one of it's arguments.");
+    DMLC_DECLARE_FIELD(step)
+    .set_default(dmlc::optional<int>())
+    .describe("If 'obj' is slice, 'step' is one of it's arguments.");
+    DMLC_DECLARE_FIELD(int_ind)
+    .set_default(dmlc::optional<int>())
+    .describe("If 'obj' is int, 'int_ind' is the index before which"
+              "'values' is inserted");
+    DMLC_DECLARE_FIELD(axis)
+    .set_default(dmlc::optional<int>())
+    .describe("Axis along which to insert `values`.");
+  }
+};
+
+namespace delete_ {
+enum DeleteOpInputs {kArr, kObj};
+enum DeleteOpOutputs {kOut};
+}  // namespace delete_
+
+template<int req>
+struct Copy {
+  template<typename DType>
+  MSHADOW_XINLINE static void Map(int i, DType* out_data, const DType* in_data) {
+    KERNEL_ASSIGN(out_data[i], req, in_data[i]);
+  }
+};
+
+struct SliceToIndices {
+  template<typename IType>
+  MSHADOW_XINLINE static void Map(int i, IType* indices, int start, int step) {
+    indices[i] = start + i * step;
+  }
+};
+
+struct ObjToIndices {
+  template<typename IType>
+  MSHADOW_XINLINE static void Map(int i, IType* indices, const IType* obj) {
+    indices[i] = obj[i];
+  }
+};
+
+template<typename IType>
+struct AssignNum {
+  MSHADOW_XINLINE static void Map(int i, IType* output, const IType data) {
+    output[i] = data;
+  }
+};
+
+struct IsDeleteCal {
+  template<typename IType>
+  MSHADOW_XINLINE static void Map(int i, int N, bool* is_delete, const IType* indices) {
+    if ((static_cast<int64_t>(indices[i]) >= 0) &&
+      (static_cast<int64_t>(indices[i]) < N)) {
+      is_delete[static_cast<int64_t>(indices[i])] = true;
+    }
+  }
+};
+
+struct OutPosCal {
+  /*!
+   * \brief map the index from input to output
+   */
+  MSHADOW_XINLINE static void Map(int i, int64_t* out_pos, const bool* is_delete) {
+    if (!is_delete[i]) {
+      int cnt = 0;
+      for ( int j = 0; j < i; ++j) {
+        if (!is_delete[j]) {
+          cnt++;
+        }
+      }
+      out_pos[i] = cnt;
+    }
+  }
+};
+
+template<int ndim>
+inline mshadow::Shape<ndim> GetStride(const mxnet::TShape& shape) {
+  mshadow::Shape<ndim>stride;
+  size_t tmp = 1;
+  for (int i = shape.ndim() - 1; i >= 0; --i) {
+    stride[i] = tmp;
+    tmp *= shape[i];
+  }
+  return stride;
+}
+
+template<int ndim>
+inline mshadow::Shape<ndim> GetKernelShape(const mxnet::TShape& shape) {
+  mshadow::Shape<ndim>k_shape;
+  for (int i = 0 ; i < shape.ndim() ; ++i) {
+    k_shape[i] = shape[i];
+  }
+  return k_shape;
+}
+
+template<int req>
+struct DeleteImpl {
+  /*!
+   * \brief delete a sub-array from input along an axis according to 'is_delete'.
+   * \tparam xpu - cpu or gpu.
+   * \param out_data - output: a new array with sub-arrays along an axis deleted.
+   * \param in_arr - input: 'arr', original array.
+   * \param is_delete - mark where will be deleted or be reminded in 'arr'
+   * \param out_pos - if is_delete[i] is 'false', out_pos[i] indicates its.
+   * \param arrshape - the shape of 'arr'.
+   * \param arr_stride - the stride of 'arr'.
+   * \param out_stride - the stride of 'out_data'.
+   * \param out_ndim - the ndim of 'out_data'.
+   * \param axis - delete sub-array along this axis
+   */
+  template<typename DType>
+  MSHADOW_XINLINE static void Map(int i, DType* out_data,
+                                  const DType* in_arr,
+                                  const bool* is_delete,
+                                  const int64_t* out_pos,
+                                  const mshadow::Shape<10> arrshape,
+                                  const mshadow::Shape<10> arr_stride,
+                                  const mshadow::Shape<10> out_stride,
+                                  const int out_ndim, const int axis) {
+    const int64_t arr_head = i / arr_stride[axis];
+    const int64_t arr_mid = arr_head % arrshape[axis];
+    mshadow::Shape<10> arr_idx;  // i -> position in in_arr's shape
+    for (int j = 0; j < out_ndim; ++j) {
+      const int64_t head = i / arr_stride[j];
+      const int64_t mid = head % arrshape[j];
+      arr_idx[j] = mid;
+    }
+    if (!is_delete[arr_mid]) {
+      arr_idx[axis] = out_pos[arr_mid];
+      int64_t dest_idx = 0;
+      for (int j =0; j < out_ndim; ++j) {
+        dest_idx += out_stride[j] * arr_idx[j];
+      }
+      KERNEL_ASSIGN(out_data[dest_idx], req, in_arr[i]);
+    }
+  }
+};
+
+template<typename xpu>
+void NumpyDeleteCompute(const nnvm::NodeAttrs& attrs,
+                           const OpContext &ctx,
+                           const std::vector<NDArray> &inputs,
+                           const std::vector<OpReqType> &req,
+                           const std::vector<NDArray> &outputs) {
+  using namespace mshadow;
+  using namespace mxnet_op;
+
+  const NumpyDeleteParam& param = nnvm::get<NumpyDeleteParam>(attrs.parsed);
+  CHECK_EQ(inputs.size(),
+          (param.step.has_value() || param.int_ind.has_value()) ? 1U : 2U);
+  CHECK_EQ(outputs.size(), 1U);
+  CHECK_EQ(req.size(), 1U);
+  mshadow::Stream<xpu> *s = ctx.get_stream<xpu>();
+
+  int ndim = inputs[delete_::kArr].shape().ndim();
+  int axis = param.axis.has_value() ? param.axis.value() : -1;
+  NDArray arr;
+
+  if (!param.axis.has_value()) {
+    arr = inputs[delete_::kArr].Reshape(Shape1(inputs[delete_::kArr].shape().Size()));
+    ndim = 1;
+    axis = -1;
+  } else {
+    arr = inputs[delete_::kArr];
+  }
+
+  if (ndim == 0) {
+    const_cast<NDArray &>(outputs[delete_::kOut]).Init(arr.shape());
+    MSHADOW_TYPE_SWITCH(outputs[delete_::kOut].dtype(), DType, {
+      MXNET_ASSIGN_REQ_SWITCH(req[delete_::kOut], req_type, {
+        Kernel<Copy<req_type>, xpu>::Launch(
+            s, outputs[delete_::kOut].shape().Size(),
+            outputs[delete_::kOut].data().dptr<DType>(),
+            inputs[delete_::kArr].data().dptr<DType>());
+      });
+    });
+    return;
+  }
+
+  axis = CheckAxis(axis, ndim);
+  int N = (arr.shape())[axis];
+  mxnet::TShape newshape(arr.shape());
+  int start = 0, stop = 0, step = 0;
+  size_t numtodel = 0;
+  int index = 0;
+
+  if (param.step.has_value()) {
+    step = param.step.value();
+    CHECK_NE(step, 0) << "'step' can not equal to 0.";
+    if (param.stop.has_value()) {
+      stop = param.stop.value();
+      stop += (stop < 0) ? N : 0;
+      stop = (stop < 0) ? ((step < 0) ? -1 : 0) : stop;
+      stop = (stop >= N) ? ((step < 0) ? N - 1 : N) : stop;
+    } else {
+      stop = (step > 0) ? N : -1;
+    }
+    if (param.start.has_value()) {
+      start = param.start.value();
+      start += (start < 0) ? N : 0;
+      start = (start < 0) ? ((step < 0) ? -1 : 0) : start;
+      start = (start >= N) ? ((step < 0) ? N - 1 : N) : start;
+    } else {
+      start = (step > 0) ? 0 : N - 1;
+    }
+    if (step > 0 && stop >= start) {
+      numtodel = static_cast<size_t>((stop - start + step - 1) / step);
+    } else if (step < 0 && stop <= start) {
+      numtodel = static_cast<size_t>((stop - start + step + 1) / step);
+    }
+    if (numtodel == 0) {
+      const_cast<NDArray &>(outputs[delete_::kOut]).Init(arr.shape());
+      MSHADOW_TYPE_SWITCH(outputs[delete_::kOut].dtype(), DType, {
+        MXNET_ASSIGN_REQ_SWITCH(req[delete_::kOut], req_type, {
+          Kernel<Copy<req_type>, xpu>::Launch(
+                s, outputs[delete_::kOut].shape().Size(),
+                outputs[delete_::kOut].data().dptr<DType>(),
+                inputs[delete_::kArr].data().dptr<DType>());
+        });
+      });
+      return;
+    }
+    newshape[axis] -= numtodel;
+    const_cast<NDArray &>(outputs[delete_::kOut]).Init(newshape);
+  } else if (param.int_ind.has_value()) {
+    index = param.int_ind.value();
+    CHECK((index >= -1 * N) && (index < N))
+      << "index " << index
+      << " is out of bounds for axis " << axis
+      << " with size " << N << "\n";
+    index += ((index < 0) ? N : 0);
+    numtodel = static_cast<size_t>(1);
+    newshape[axis] -= 1;
+    const_cast<NDArray &>(outputs[delete_::kOut]).Init(newshape);
+  } else {
+    numtodel = inputs[delete_::kObj].shape().Size();
+  }
+
+  MSHADOW_TYPE_SWITCH((inputs.size() == 2U) ?
+                       inputs[delete_::kObj].dtype() :
+                       mshadow::DataType<int64_t>::kFlag, IType, {
+    size_t temp_mem_size = sizeof(int64_t) * arr.shape()[axis] +
+                           sizeof(IType) * numtodel +
+                           sizeof(bool) * arr.shape()[axis];
+    Tensor<xpu, 1, char> temp_mem =
+            ctx.requested[0].get_space_typed<xpu, 1, char>(Shape1(temp_mem_size), s);
+    int64_t* out_pos_ptr = reinterpret_cast<int64_t*>(temp_mem.dptr_);
+    IType* indices_ptr = reinterpret_cast<IType*>
+                         (temp_mem.dptr_ + sizeof(int64_t) * arr.shape()[axis]);
+    bool* is_delete_ptr = reinterpret_cast<bool*>
+                          (temp_mem.dptr_ + sizeof(int64_t) * arr.shape()[axis]
+                          + sizeof(IType) * numtodel);
+    if (param.step.has_value()) {
+      Kernel<SliceToIndices, xpu>::Launch(s, numtodel,
+                                          indices_ptr, start, step);
+    } else if (param.int_ind.has_value()) {
+      Kernel<AssignNum<IType>, xpu>::Launch(s, numtodel, indices_ptr, index);
+    } else {
+      Kernel<ObjToIndices, xpu>::Launch(s, numtodel, indices_ptr,
+                                        inputs[delete_::kObj].data().dptr<IType>());
+    }
+    Kernel<AssignNum<bool>, xpu>::Launch(s, arr.shape()[axis], is_delete_ptr, false);
+    Kernel<IsDeleteCal, xpu>::Launch(s, numtodel, N, is_delete_ptr, indices_ptr);
+    Kernel<OutPosCal, xpu>::Launch(s, arr.shape()[axis], out_pos_ptr, is_delete_ptr);
+    if (inputs.size() == 2U) {
+      IType* input_obj = inputs[delete_::kObj].data().dptr<IType>();
+      #ifndef __CUDACC__
+        std::vector<bool>vec_is_delete(is_delete_ptr, is_delete_ptr + arr.shape()[axis]);
+      #else
+        thrust::device_ptr<bool>is_delete_dev(is_delete_ptr);
+        thrust::device_vector<bool>vec_is_delete(is_delete_dev, is_delete_dev + arr.shape()[axis]);
+      #endif
+      numtodel = 0;
+      for (int i = 0; i < arr.shape()[axis]; ++i) {
+        if (vec_is_delete[i]) {
+          numtodel++;
+        }
+      }
+      newshape[axis] -= numtodel;
+      const_cast<NDArray &>(outputs[delete_::kOut]).Init(newshape);
+    }
+    mshadow::Shape<10> arr_strides = GetStride<10>(arr.shape());
+    mshadow::Shape<10> out_strides = GetStride<10>(newshape);
+    mshadow::Shape<10> k_arrshape = GetKernelShape<10>(arr.shape());
+    MSHADOW_TYPE_SWITCH(outputs[delete_::kOut].dtype(), DType, {
+      MXNET_ASSIGN_REQ_SWITCH(req[delete_::kOut], req_type, {
+        Kernel<DeleteImpl<req_type>, xpu>::Launch(s, arr.shape().Size(),
+                                                  outputs[delete_::kOut].data().dptr<DType>(),
+                                                  arr.data().dptr<DType>(),
+                                                  is_delete_ptr, out_pos_ptr,
+                                                  k_arrshape,
+                                                  arr_strides, out_strides,
+                                                  newshape.ndim(), axis);
+      });
+    });
+  });
+}
+
+}  // namespace op
+}  // namespace mxnet
+
+#endif  // MXNET_OPERATOR_NUMPY_NP_DELETE_OP_INL_H_
diff --git a/src/operator/numpy/np_delete_op.cc b/src/operator/numpy/np_delete_op.cc
new file mode 100644
index 000000000000..a85a5eb4f407
--- /dev/null
+++ b/src/operator/numpy/np_delete_op.cc
@@ -0,0 +1,79 @@
+/*!
+ *  Copyright (c) 2019 by Contributors
+ * \file np_delete_op.cc
+ * \brief CPU Implementation of numpy insert operations
+ */
+
+#include <vector>
+#include "./np_delete_op-inl.h"
+
+namespace mxnet {
+namespace op {
+
+DMLC_REGISTER_PARAMETER(NumpyDeleteParam);
+
+bool NumpyDeleteType(const nnvm::NodeAttrs& attrs,
+                     std::vector<int> *in_type,
+                     std::vector<int> *out_type) {
+  const NumpyDeleteParam& param = nnvm::get<NumpyDeleteParam>(attrs.parsed);
+  int insize = (param.step.has_value() || param.int_ind.has_value()) ? 1 : 2;
+  CHECK_EQ(in_type->size(), insize);
+  CHECK_EQ(out_type->size(), 1U);
+  if (insize == 3) {
+    CHECK_NE((*in_type)[1], -1) << "Index type must be set for insert operator\n";
+    CHECK(((*in_type)[1] == mshadow::DataType<int64_t>::kFlag)
+         || ((*in_type)[1] == mshadow::DataType<int32_t>::kFlag))
+      << "Index type only support int32 or int64.\n";
+  }
+  TYPE_ASSIGN_CHECK(*out_type, 0, (*in_type)[0]);
+  TYPE_ASSIGN_CHECK(*in_type, 0, (*out_type)[0]);
+  return (*in_type)[0] != -1;
+}
+
+inline bool NumpyDeleteStorageType(const nnvm::NodeAttrs& attrs,
+                                   const int dev_mask,
+                                   DispatchMode* dispatch_mode,
+                                   std::vector<int> *in_attrs,
+                                   std::vector<int> *out_attrs) {
+  const NumpyDeleteParam& param = nnvm::get<NumpyDeleteParam>(attrs.parsed);
+  unsigned int insize = (param.step.has_value() || param.int_ind.has_value()) ? 1U : 2U;
+  CHECK_EQ(in_attrs->size(), insize);
+  CHECK_EQ(out_attrs->size(), 1U);
+  for (int &attr : *in_attrs) {
+    CHECK_EQ(attr, kDefaultStorage) << "Only default storage is supported";
+  }
+  for (int &attr : *out_attrs) {
+    attr = kDefaultStorage;
+  }
+  *dispatch_mode = DispatchMode::kFComputeEx;
+  return true;
+}
+
+NNVM_REGISTER_OP(_npi_delete)
+.describe(R"code(Delete values along the given axis before the given indices.)code" ADD_FILELINE)
+.set_attr_parser(ParamParser<NumpyDeleteParam>)
+.set_num_inputs([](const NodeAttrs& attrs) {
+  const NumpyDeleteParam& params = nnvm::get<NumpyDeleteParam>(attrs.parsed);
+  return (params.step.has_value() || params.int_ind.has_value()) ? 1U : 2U;
+})
+.set_num_outputs(1)
+.set_attr<nnvm::FListInputNames>("FListInputNames",
+  [](const NodeAttrs& attrs) {
+    const NumpyDeleteParam& params = nnvm::get<NumpyDeleteParam>(attrs.parsed);
+    return (params.step.has_value() || params.int_ind.has_value()) ?
+            std::vector<std::string>{"arr"} :
+            std::vector<std::string>{"arr", "obj"};
+})
+.set_attr<nnvm::FInferType>("FInferType", NumpyDeleteType)
+.set_attr<mxnet::FComputeEx>("FComputeEx<cpu>", NumpyDeleteCompute<cpu>)
+.set_attr<FInferStorageType>("FInferStorageType", NumpyDeleteStorageType)
+.set_attr<FResourceRequest>("FResourceRequest",
+  [](const NodeAttrs& attrs) {
+    return std::vector<ResourceRequest>{ResourceRequest::kTempSpace};
+  })
+.add_argument("arr", "NDArray-or-Symbol", "Input ndarray")
+.add_argument("obj", "NDArray-or-Symbol", "Input ndarray")
+.add_arguments(NumpyDeleteParam::__FIELDS__());
+
+}  // namespace op
+}  // namespace mxnet
\ No newline at end of file
diff --git a/src/operator/numpy/np_delete_op.cu b/src/operator/numpy/np_delete_op.cu
new file mode 100644
index 000000000000..5f66f0207eae
--- /dev/null
+++ b/src/operator/numpy/np_delete_op.cu
@@ -0,0 +1,16 @@
+/*!
+ *  Copyright (c) 2019 by Contributors
+ * \file np_delete_op.cu
+ * \brief GPU Implementation of numpy delete operations
+ */
+
+#include "./np_delete_op-inl.h"
+
+namespace mxnet {
+namespace op {
+
+NNVM_REGISTER_OP(_npi_delete)
+.set_attr<FComputeEx>("FComputeEx<gpu>", NumpyDeleteCompute<gpu>);
+
+}
+}
\ No newline at end of file
diff --git a/tests/python/unittest/test_numpy_interoperability.py b/tests/python/unittest/test_numpy_interoperability.py
index fcdf547bfbec..0cc4f7cac96d 100644
--- a/tests/python/unittest/test_numpy_interoperability.py
+++ b/tests/python/unittest/test_numpy_interoperability.py
@@ -855,6 +855,27 @@ def _add_workload_unique():
     # OpArgMngr.add_workload('unique', np.arange(10, dtype=np.uint8).reshape(-1, 2).astype(bool), axis=1)
 
 
+def _add_workload_delete():
+    a = np.arange(5)
+    nd_a = np.arange(5).repeat(2).reshape(1, 5, 2)
+    lims = [-6, -2, 0, 1, 2, 4, 5]
+    steps = [-3, -1, 1, 3]
+    for start in lims:
+        for stop in lims:
+            for step in steps:
+                s = slice(start, stop, step)
+                OpArgMngr.add_workload('delete', a, s)
+                OpArgMngr.add_workload('delete', nd_a, s, axis=1)
+    OpArgMngr.add_workload('delete', a, np.array([]), axis=0)
+    OpArgMngr.add_workload('delete', a, 0)
+    OpArgMngr.add_workload('delete', a, np.array([]))
+    OpArgMngr.add_workload('delete', a, np.array([0, 1]))
+    OpArgMngr.add_workload('delete', a, slice(1, 2))
+    OpArgMngr.add_workload('delete', a, slice(1, -2))
+    k = np.arange(10).reshape(2, 5)
+    OpArgMngr.add_workload('delete', k, slice(60, None), axis=1)
+
+
 def _add_workload_var(array_pool):
     OpArgMngr.add_workload('var', array_pool['4x1'])
     OpArgMngr.add_workload('var', np.array([np.float16(1.)]))
@@ -1482,6 +1503,7 @@ def _prepare_workloads():
     _add_workload_tile()
     _add_workload_transpose()
     _add_workload_unique()
+    _add_workload_delete()
     _add_workload_var(array_pool)
     _add_workload_zeros_like(array_pool)
     _add_workload_linalg_norm()
diff --git a/tests/python/unittest/test_numpy_op.py b/tests/python/unittest/test_numpy_op.py
index af9228d45991..99d104df1198 100644
--- a/tests/python/unittest/test_numpy_op.py
+++ b/tests/python/unittest/test_numpy_op.py
@@ -2754,6 +2754,81 @@ def hybrid_forward(self, F, x):
             assert same(ret_mx.asnumpy(), ret_np)
 
 
+@with_seed()
+@use_np
+def test_np_delete():
+    class TestDelete(HybridBlock):
+        def __init__(self, obj, axis=None):
+            super(TestDelete, self).__init__()
+            self._obj = obj
+            self._axis = axis
+
+        def hybrid_forward(self, F, a):
+            return F.np.delete(a, self._obj, axis=self._axis)
+    
+    def GetSize(shp):
+        if len(shp) == 0:
+            return 0
+        else:
+            res = 1
+            shp_list = list(shp)
+            for x in shp:
+                res *= x
+            return res
+
+    def GetDimSize(shp, axis):
+        if axis is None:
+            return GetSize(shp)
+        shp_list = list(shp)
+        return shp_list[axis]
+
+    shape = [(), (0, ), (1, ), (2, 3), (2, 1, 4, 5)]
+    config = []
+    for shp in shape:
+        for ax in range(-1 * len(shp), len(shp), 2):
+            #test slice
+            for st in [-5, -2, 0, 2, 5, None]:
+                for ed in [-5, -2, 0, 2, 5, None]:
+                    for stp in [-5, -2, 2, 5, None]:
+                        config.append(tuple([shp, slice(st, ed, stp), None]))
+                        config.append(tuple([shp, slice(st, ed, stp), ax]))
+            #test iteger
+            for idx in range(-1 * GetDimSize(shp, ax), GetDimSize(shp, ax)):
+                config.append(tuple([shp, idx, ax]))
+            #test ndarray indices
+            idx =  _np.random.randint(-1 * shp[ax], shp[ax] + 1, size = (4)).tolist()
+            config.append(tuple([shp, idx, ax]))
+
+    for arr_shape, obj, axis in config:
+        for objtype in ['int32', 'int64']:
+            if type(obj) == list:
+                obj_mxnp = np.array(obj, dtype=objtype)
+                obj_onp = _np.array(obj, dtype=objtype)
+            elif type(obj) == slice:
+                obj_mxnp = obj
+                obj_onp = obj
+            else:
+                obj_mxnp = (_np.int32(obj) if objtype == 'int32' else _np.int64(obj)) 
+                obj_onp = (_np.int32(obj) if objtype == 'int32' else _np.int64(obj)) 
+            test_delete = TestDelete(obj=obj_mxnp, axis=axis)
+
+            a = mx.nd.random.uniform(-1.0, 1.0, shape=arr_shape).as_np_ndarray()
+            a.attach_grad()
+            expected_ret = _np.delete(a.asnumpy(), obj_onp, axis=axis)
+
+            with mx.autograd.record():
+                y = test_delete(a)
+
+            assert y.shape == expected_ret.shape
+            assert_almost_equal(y.asnumpy(), expected_ret, rtol=1e-3, atol=1e-5)
+
+            #test imperative
+            mx_out = np.delete(a, obj_mxnp, axis=axis)
+            np_out = _np.delete(a.asnumpy(), obj_onp, axis=axis)
+
+            assert_almost_equal(mx_out.asnumpy(), np_out, rtol=1e-3, atol=1e-5)
+
+
 @with_seed()
 @use_np
 def test_np_argmin_argmax():

From 73bb4ee954b19d5cc17cd111b8744bbdaa9b924d Mon Sep 17 00:00:00 2001
From: JiangZhaoh <jiang_zhaoh@163.com>
Date: Mon, 9 Dec 2019 10:28:56 +0000
Subject: [PATCH 2/5] modify code

---
 python/mxnet/ndarray/numpy/_op.py     |  5 --
 python/mxnet/numpy/multiarray.py      |  5 --
 python/mxnet/symbol/numpy/_symbol.py  |  5 --
 src/operator/numpy/np_delete_op-inl.h | 85 +++++++++------------------
 src/operator/numpy/np_delete_op.cc    |  2 +-
 src/operator/numpy/np_delete_op.cu    |  2 +-
 6 files changed, 31 insertions(+), 73 deletions(-)

diff --git a/python/mxnet/ndarray/numpy/_op.py b/python/mxnet/ndarray/numpy/_op.py
index b14e25c68ac6..874ddc97abae 100644
--- a/python/mxnet/ndarray/numpy/_op.py
+++ b/python/mxnet/ndarray/numpy/_op.py
@@ -937,11 +937,6 @@ def delete(arr, obj, axis=None):
         that `delete` does not occur in-place. If `axis` is None, `out` is
         a flattened array.
 
-    See Also
-    --------
-    insert : Insert elements into an array.
-    append : Append elements at the end of an array.
-
     Examples
     --------
     >>> arr = np.array([[1,2,3,4], [5,6,7,8], [9,10,11,12]])
diff --git a/python/mxnet/numpy/multiarray.py b/python/mxnet/numpy/multiarray.py
index 7cca23649d70..689a5b1322c6 100644
--- a/python/mxnet/numpy/multiarray.py
+++ b/python/mxnet/numpy/multiarray.py
@@ -5872,11 +5872,6 @@ def delete(arr, obj, axis=None):
         that `delete` does not occur in-place. If `axis` is None, `out` is
         a flattened array.
 
-    See Also
-    --------
-    insert : Insert elements into an array.
-    append : Append elements at the end of an array.
-
     Examples
     --------
     >>> arr = np.array([[1,2,3,4], [5,6,7,8], [9,10,11,12]])
diff --git a/python/mxnet/symbol/numpy/_symbol.py b/python/mxnet/symbol/numpy/_symbol.py
index 24672dfcfce8..bda1c52c796c 100644
--- a/python/mxnet/symbol/numpy/_symbol.py
+++ b/python/mxnet/symbol/numpy/_symbol.py
@@ -3184,11 +3184,6 @@ def delete(arr, obj, axis=None):
         A copy of `arr` with the elements specified by `obj` removed. Note
         that `delete` does not occur in-place. If `axis` is None, `out` is
         a flattened array.
-
-    See Also
-    --------
-    insert : Insert elements into an array.
-    append : Append elements at the end of an array.
     """
     if not isinstance(arr, Symbol):
         raise TypeError("'arr' can not support type {}".format(str(type(arr))))
diff --git a/src/operator/numpy/np_delete_op-inl.h b/src/operator/numpy/np_delete_op-inl.h
index be14a69c6552..68ae09b5a42f 100644
--- a/src/operator/numpy/np_delete_op-inl.h
+++ b/src/operator/numpy/np_delete_op-inl.h
@@ -8,9 +8,11 @@
 
 #include <vector>
 #include <memory>
+#include <algorithm>
 #include "../../common/utils.h"
 #include "../tensor/sort_op.h"
 #include "../operator_common.h"
+#include "../mxnet_op.h"
 #include "../tensor/broadcast_reduce_op.h"
 #ifdef __CUDACC__
 #include <thrust/device_ptr.h>
@@ -55,14 +57,6 @@ enum DeleteOpInputs {kArr, kObj};
 enum DeleteOpOutputs {kOut};
 }  // namespace delete_
 
-template<int req>
-struct Copy {
-  template<typename DType>
-  MSHADOW_XINLINE static void Map(int i, DType* out_data, const DType* in_data) {
-    KERNEL_ASSIGN(out_data[i], req, in_data[i]);
-  }
-};
-
 struct SliceToIndices {
   template<typename IType>
   MSHADOW_XINLINE static void Map(int i, IType* indices, int start, int step) {
@@ -70,13 +64,6 @@ struct SliceToIndices {
   }
 };
 
-struct ObjToIndices {
-  template<typename IType>
-  MSHADOW_XINLINE static void Map(int i, IType* indices, const IType* obj) {
-    indices[i] = obj[i];
-  }
-};
-
 template<typename IType>
 struct AssignNum {
   MSHADOW_XINLINE static void Map(int i, IType* output, const IType data) {
@@ -88,7 +75,7 @@ struct IsDeleteCal {
   template<typename IType>
   MSHADOW_XINLINE static void Map(int i, int N, bool* is_delete, const IType* indices) {
     if ((static_cast<int64_t>(indices[i]) >= 0) &&
-      (static_cast<int64_t>(indices[i]) < N)) {
+        (static_cast<int64_t>(indices[i]) < N)) {
       is_delete[static_cast<int64_t>(indices[i])] = true;
     }
   }
@@ -101,7 +88,7 @@ struct OutPosCal {
   MSHADOW_XINLINE static void Map(int i, int64_t* out_pos, const bool* is_delete) {
     if (!is_delete[i]) {
       int cnt = 0;
-      for ( int j = 0; j < i; ++j) {
+      for (int j = 0; j < i; ++j) {
         if (!is_delete[j]) {
           cnt++;
         }
@@ -176,10 +163,10 @@ struct DeleteImpl {
 
 template<typename xpu>
 void NumpyDeleteCompute(const nnvm::NodeAttrs& attrs,
-                           const OpContext &ctx,
-                           const std::vector<NDArray> &inputs,
-                           const std::vector<OpReqType> &req,
-                           const std::vector<NDArray> &outputs) {
+                        const OpContext &ctx,
+                        const std::vector<NDArray> &inputs,
+                        const std::vector<OpReqType> &req,
+                        const std::vector<NDArray> &outputs) {
   using namespace mshadow;
   using namespace mxnet_op;
 
@@ -204,14 +191,7 @@ void NumpyDeleteCompute(const nnvm::NodeAttrs& attrs,
 
   if (ndim == 0) {
     const_cast<NDArray &>(outputs[delete_::kOut]).Init(arr.shape());
-    MSHADOW_TYPE_SWITCH(outputs[delete_::kOut].dtype(), DType, {
-      MXNET_ASSIGN_REQ_SWITCH(req[delete_::kOut], req_type, {
-        Kernel<Copy<req_type>, xpu>::Launch(
-            s, outputs[delete_::kOut].shape().Size(),
-            outputs[delete_::kOut].data().dptr<DType>(),
-            inputs[delete_::kArr].data().dptr<DType>());
-      });
-    });
+    mxnet_op::copy(s, outputs[delete_::kOut].data(), inputs[delete_::kArr].data());
     return;
   }
 
@@ -248,14 +228,7 @@ void NumpyDeleteCompute(const nnvm::NodeAttrs& attrs,
     }
     if (numtodel == 0) {
       const_cast<NDArray &>(outputs[delete_::kOut]).Init(arr.shape());
-      MSHADOW_TYPE_SWITCH(outputs[delete_::kOut].dtype(), DType, {
-        MXNET_ASSIGN_REQ_SWITCH(req[delete_::kOut], req_type, {
-          Kernel<Copy<req_type>, xpu>::Launch(
-                s, outputs[delete_::kOut].shape().Size(),
-                outputs[delete_::kOut].data().dptr<DType>(),
-                inputs[delete_::kArr].data().dptr<DType>());
-        });
-      });
+      mxnet_op::copy(s, outputs[delete_::kOut].data(), inputs[delete_::kArr].data());
       return;
     }
     newshape[axis] -= numtodel;
@@ -274,9 +247,9 @@ void NumpyDeleteCompute(const nnvm::NodeAttrs& attrs,
     numtodel = inputs[delete_::kObj].shape().Size();
   }
 
-  MSHADOW_TYPE_SWITCH((inputs.size() == 2U) ?
-                       inputs[delete_::kObj].dtype() :
-                       mshadow::DataType<int64_t>::kFlag, IType, {
+  MSHADOW_TYPE_SWITCH(((inputs.size() == 2U) ?
+                        inputs[delete_::kObj].dtype() :
+                        mshadow::DataType<int64_t>::kFlag), IType, {
     size_t temp_mem_size = sizeof(int64_t) * arr.shape()[axis] +
                            sizeof(IType) * numtodel +
                            sizeof(bool) * arr.shape()[axis];
@@ -286,27 +259,27 @@ void NumpyDeleteCompute(const nnvm::NodeAttrs& attrs,
     IType* indices_ptr = reinterpret_cast<IType*>
                          (temp_mem.dptr_ + sizeof(int64_t) * arr.shape()[axis]);
     bool* is_delete_ptr = reinterpret_cast<bool*>
-                          (temp_mem.dptr_ + sizeof(int64_t) * arr.shape()[axis]
-                          + sizeof(IType) * numtodel);
+                          (temp_mem.dptr_ + sizeof(int64_t) * arr.shape()[axis] +
+                           sizeof(IType) * numtodel);
     if (param.step.has_value()) {
       Kernel<SliceToIndices, xpu>::Launch(s, numtodel,
                                           indices_ptr, start, step);
     } else if (param.int_ind.has_value()) {
       Kernel<AssignNum<IType>, xpu>::Launch(s, numtodel, indices_ptr, index);
     } else {
-      Kernel<ObjToIndices, xpu>::Launch(s, numtodel, indices_ptr,
-                                        inputs[delete_::kObj].data().dptr<IType>());
+      mxnet_op::copy(s,
+        TBlob(indices_ptr, inputs[delete_::kObj].shape(), inputs[delete_::kObj].data().dev_mask()),
+        inputs[delete_::kObj].data());
     }
-    Kernel<AssignNum<bool>, xpu>::Launch(s, arr.shape()[axis], is_delete_ptr, false);
+    mxnet_op::Kernel<mxnet_op::set_zero, xpu>::Launch(s, arr.shape()[axis], is_delete_ptr);
     Kernel<IsDeleteCal, xpu>::Launch(s, numtodel, N, is_delete_ptr, indices_ptr);
     Kernel<OutPosCal, xpu>::Launch(s, arr.shape()[axis], out_pos_ptr, is_delete_ptr);
     if (inputs.size() == 2U) {
-      IType* input_obj = inputs[delete_::kObj].data().dptr<IType>();
-      #ifndef __CUDACC__
-        std::vector<bool>vec_is_delete(is_delete_ptr, is_delete_ptr + arr.shape()[axis]);
-      #else
+      #ifdef __CUDACC__
         thrust::device_ptr<bool>is_delete_dev(is_delete_ptr);
         thrust::device_vector<bool>vec_is_delete(is_delete_dev, is_delete_dev + arr.shape()[axis]);
+      #else
+        std::vector<bool>vec_is_delete(is_delete_ptr, is_delete_ptr + arr.shape()[axis]);
       #endif
       numtodel = 0;
       for (int i = 0; i < arr.shape()[axis]; ++i) {
@@ -322,13 +295,13 @@ void NumpyDeleteCompute(const nnvm::NodeAttrs& attrs,
     mshadow::Shape<10> k_arrshape = GetKernelShape<10>(arr.shape());
     MSHADOW_TYPE_SWITCH(outputs[delete_::kOut].dtype(), DType, {
       MXNET_ASSIGN_REQ_SWITCH(req[delete_::kOut], req_type, {
-        Kernel<DeleteImpl<req_type>, xpu>::Launch(s, arr.shape().Size(),
-                                                  outputs[delete_::kOut].data().dptr<DType>(),
-                                                  arr.data().dptr<DType>(),
-                                                  is_delete_ptr, out_pos_ptr,
-                                                  k_arrshape,
-                                                  arr_strides, out_strides,
-                                                  newshape.ndim(), axis);
+        Kernel<DeleteImpl<req_type>, xpu>::Launch(
+          s, arr.shape().Size(),
+          outputs[delete_::kOut].data().dptr<DType>(),
+          arr.data().dptr<DType>(),
+          is_delete_ptr, out_pos_ptr,
+          k_arrshape, arr_strides, out_strides,
+          newshape.ndim(), axis);
       });
     });
   });
diff --git a/src/operator/numpy/np_delete_op.cc b/src/operator/numpy/np_delete_op.cc
index a85a5eb4f407..11d0ea62a03f 100644
--- a/src/operator/numpy/np_delete_op.cc
+++ b/src/operator/numpy/np_delete_op.cc
@@ -76,4 +76,4 @@ NNVM_REGISTER_OP(_npi_delete)
 .add_arguments(NumpyDeleteParam::__FIELDS__());
 
 }  // namespace op
-}  // namespace mxnet
\ No newline at end of file
+}  // namespace mxnet
diff --git a/src/operator/numpy/np_delete_op.cu b/src/operator/numpy/np_delete_op.cu
index 5f66f0207eae..7abb824fec3e 100644
--- a/src/operator/numpy/np_delete_op.cu
+++ b/src/operator/numpy/np_delete_op.cu
@@ -13,4 +13,4 @@ NNVM_REGISTER_OP(_npi_delete)
 .set_attr<FComputeEx>("FComputeEx<gpu>", NumpyDeleteCompute<gpu>);
 
 }
-}
\ No newline at end of file
+}

From 71ecb71462455e0713860d3af8e5b3958f6dc009 Mon Sep 17 00:00:00 2001
From: JiangZhaoh <jiang_zhaoh@163.com>
Date: Mon, 16 Dec 2019 07:57:18 +0000
Subject: [PATCH 3/5] reuse existing function / add licenses / modify alignment

---
 src/operator/numpy/np_delete_op-inl.h | 120 +++++++++++---------------
 src/operator/numpy/np_delete_op.cc    |  23 ++++-
 src/operator/numpy/np_delete_op.cu    |  19 ++++
 3 files changed, 89 insertions(+), 73 deletions(-)

diff --git a/src/operator/numpy/np_delete_op-inl.h b/src/operator/numpy/np_delete_op-inl.h
index 68ae09b5a42f..bcdb126f233f 100644
--- a/src/operator/numpy/np_delete_op-inl.h
+++ b/src/operator/numpy/np_delete_op-inl.h
@@ -1,3 +1,22 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements.  See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership.  The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License.  You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied.  See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
 /*!
  *  Copyright (c) 2019 by Contributors
  * \file np_delete_op-inl.h
@@ -11,6 +30,7 @@
 #include <algorithm>
 #include "../../common/utils.h"
 #include "../tensor/sort_op.h"
+#include "../tensor/init_op.h"
 #include "../operator_common.h"
 #include "../mxnet_op.h"
 #include "../tensor/broadcast_reduce_op.h"
@@ -64,19 +84,11 @@ struct SliceToIndices {
   }
 };
 
-template<typename IType>
-struct AssignNum {
-  MSHADOW_XINLINE static void Map(int i, IType* output, const IType data) {
-    output[i] = data;
-  }
-};
-
 struct IsDeleteCal {
   template<typename IType>
   MSHADOW_XINLINE static void Map(int i, int N, bool* is_delete, const IType* indices) {
-    if ((static_cast<int64_t>(indices[i]) >= 0) &&
-        (static_cast<int64_t>(indices[i]) < N)) {
-      is_delete[static_cast<int64_t>(indices[i])] = true;
+    if ((indices[i] >= 0) && (indices[i] < N)) {
+      is_delete[static_cast<int>(indices[i])] = true;
     }
   }
 };
@@ -98,39 +110,16 @@ struct OutPosCal {
   }
 };
 
-template<int ndim>
-inline mshadow::Shape<ndim> GetStride(const mxnet::TShape& shape) {
-  mshadow::Shape<ndim>stride;
-  size_t tmp = 1;
-  for (int i = shape.ndim() - 1; i >= 0; --i) {
-    stride[i] = tmp;
-    tmp *= shape[i];
-  }
-  return stride;
-}
-
-template<int ndim>
-inline mshadow::Shape<ndim> GetKernelShape(const mxnet::TShape& shape) {
-  mshadow::Shape<ndim>k_shape;
-  for (int i = 0 ; i < shape.ndim() ; ++i) {
-    k_shape[i] = shape[i];
-  }
-  return k_shape;
-}
-
-template<int req>
+template<int req, int ndim>
 struct DeleteImpl {
   /*!
    * \brief delete a sub-array from input along an axis according to 'is_delete'.
-   * \tparam xpu - cpu or gpu.
    * \param out_data - output: a new array with sub-arrays along an axis deleted.
    * \param in_arr - input: 'arr', original array.
    * \param is_delete - mark where will be deleted or be reminded in 'arr'
    * \param out_pos - if is_delete[i] is 'false', out_pos[i] indicates its.
    * \param arrshape - the shape of 'arr'.
-   * \param arr_stride - the stride of 'arr'.
    * \param out_stride - the stride of 'out_data'.
-   * \param out_ndim - the ndim of 'out_data'.
    * \param axis - delete sub-array along this axis
    */
   template<typename DType>
@@ -138,24 +127,14 @@ struct DeleteImpl {
                                   const DType* in_arr,
                                   const bool* is_delete,
                                   const int64_t* out_pos,
-                                  const mshadow::Shape<10> arrshape,
-                                  const mshadow::Shape<10> arr_stride,
-                                  const mshadow::Shape<10> out_stride,
-                                  const int out_ndim, const int axis) {
-    const int64_t arr_head = i / arr_stride[axis];
-    const int64_t arr_mid = arr_head % arrshape[axis];
-    mshadow::Shape<10> arr_idx;  // i -> position in in_arr's shape
-    for (int j = 0; j < out_ndim; ++j) {
-      const int64_t head = i / arr_stride[j];
-      const int64_t mid = head % arrshape[j];
-      arr_idx[j] = mid;
-    }
-    if (!is_delete[arr_mid]) {
-      arr_idx[axis] = out_pos[arr_mid];
-      int64_t dest_idx = 0;
-      for (int j =0; j < out_ndim; ++j) {
-        dest_idx += out_stride[j] * arr_idx[j];
-      }
+                                  const mshadow::Shape<ndim> arrshape,
+                                  const mshadow::Shape<ndim> out_stride,
+                                  const int axis) {
+    // i -> position in in_arr's shape
+    mshadow::Shape<ndim> arr_idx = mxnet_op::unravel(i, arrshape);
+    if (!is_delete[arr_idx[axis]]) {
+      arr_idx[axis] = out_pos[arr_idx[axis]];
+      int64_t dest_idx = mxnet_op::dot(arr_idx, out_stride);
       KERNEL_ASSIGN(out_data[dest_idx], req, in_arr[i]);
     }
   }
@@ -248,24 +227,23 @@ void NumpyDeleteCompute(const nnvm::NodeAttrs& attrs,
   }
 
   MSHADOW_TYPE_SWITCH(((inputs.size() == 2U) ?
-                        inputs[delete_::kObj].dtype() :
-                        mshadow::DataType<int64_t>::kFlag), IType, {
+                      inputs[delete_::kObj].dtype() :
+                      mshadow::DataType<int64_t>::kFlag), IType, {
     size_t temp_mem_size = sizeof(int64_t) * arr.shape()[axis] +
                            sizeof(IType) * numtodel +
                            sizeof(bool) * arr.shape()[axis];
     Tensor<xpu, 1, char> temp_mem =
-            ctx.requested[0].get_space_typed<xpu, 1, char>(Shape1(temp_mem_size), s);
+      ctx.requested[0].get_space_typed<xpu, 1, char>(Shape1(temp_mem_size), s);
     int64_t* out_pos_ptr = reinterpret_cast<int64_t*>(temp_mem.dptr_);
     IType* indices_ptr = reinterpret_cast<IType*>
                          (temp_mem.dptr_ + sizeof(int64_t) * arr.shape()[axis]);
     bool* is_delete_ptr = reinterpret_cast<bool*>
                           (temp_mem.dptr_ + sizeof(int64_t) * arr.shape()[axis] +
-                           sizeof(IType) * numtodel);
+                          sizeof(IType) * numtodel);
     if (param.step.has_value()) {
-      Kernel<SliceToIndices, xpu>::Launch(s, numtodel,
-                                          indices_ptr, start, step);
+      Kernel<SliceToIndices, xpu>::Launch(s, numtodel, indices_ptr, start, step);
     } else if (param.int_ind.has_value()) {
-      Kernel<AssignNum<IType>, xpu>::Launch(s, numtodel, indices_ptr, index);
+      Fill(s, TBlob(indices_ptr, Shape1(numtodel), xpu::kDevMask), kWriteTo, index);
     } else {
       mxnet_op::copy(s,
         TBlob(indices_ptr, inputs[delete_::kObj].shape(), inputs[delete_::kObj].data().dev_mask()),
@@ -290,18 +268,18 @@ void NumpyDeleteCompute(const nnvm::NodeAttrs& attrs,
       newshape[axis] -= numtodel;
       const_cast<NDArray &>(outputs[delete_::kOut]).Init(newshape);
     }
-    mshadow::Shape<10> arr_strides = GetStride<10>(arr.shape());
-    mshadow::Shape<10> out_strides = GetStride<10>(newshape);
-    mshadow::Shape<10> k_arrshape = GetKernelShape<10>(arr.shape());
-    MSHADOW_TYPE_SWITCH(outputs[delete_::kOut].dtype(), DType, {
-      MXNET_ASSIGN_REQ_SWITCH(req[delete_::kOut], req_type, {
-        Kernel<DeleteImpl<req_type>, xpu>::Launch(
-          s, arr.shape().Size(),
-          outputs[delete_::kOut].data().dptr<DType>(),
-          arr.data().dptr<DType>(),
-          is_delete_ptr, out_pos_ptr,
-          k_arrshape, arr_strides, out_strides,
-          newshape.ndim(), axis);
+    MXNET_NDIM_SWITCH(newshape.ndim(), ndim, {
+      mshadow::Shape<ndim> out_strides = mxnet_op::calc_stride(newshape.get<ndim>());
+      MSHADOW_TYPE_SWITCH(outputs[delete_::kOut].dtype(), DType, {
+        MXNET_ASSIGN_REQ_SWITCH(req[delete_::kOut], req_type, {
+          Kernel<DeleteImpl<req_type, ndim>, xpu>::Launch(
+            s, arr.shape().Size(),
+            outputs[delete_::kOut].data().dptr<DType>(),
+            arr.data().dptr<DType>(),
+            is_delete_ptr, out_pos_ptr,
+            arr.shape().get<ndim>(),
+            out_strides, axis);
+        });
       });
     });
   });
diff --git a/src/operator/numpy/np_delete_op.cc b/src/operator/numpy/np_delete_op.cc
index 11d0ea62a03f..48840bf9d230 100644
--- a/src/operator/numpy/np_delete_op.cc
+++ b/src/operator/numpy/np_delete_op.cc
@@ -1,3 +1,22 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements.  See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership.  The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License.  You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied.  See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
 /*!
  *  Copyright (c) 2019 by Contributors
  * \file np_delete_op.cc
@@ -21,8 +40,8 @@ bool NumpyDeleteType(const nnvm::NodeAttrs& attrs,
   CHECK_EQ(out_type->size(), 1U);
   if (insize == 3) {
     CHECK_NE((*in_type)[1], -1) << "Index type must be set for insert operator\n";
-    CHECK(((*in_type)[1] == mshadow::DataType<int64_t>::kFlag)
-         || ((*in_type)[1] == mshadow::DataType<int32_t>::kFlag))
+    CHECK(((*in_type)[1] == mshadow::DataType<int64_t>::kFlag) ||
+          ((*in_type)[1] == mshadow::DataType<int32_t>::kFlag))
       << "Index type only support int32 or int64.\n";
   }
   TYPE_ASSIGN_CHECK(*out_type, 0, (*in_type)[0]);
diff --git a/src/operator/numpy/np_delete_op.cu b/src/operator/numpy/np_delete_op.cu
index 7abb824fec3e..599d01788138 100644
--- a/src/operator/numpy/np_delete_op.cu
+++ b/src/operator/numpy/np_delete_op.cu
@@ -1,3 +1,22 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements.  See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership.  The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License.  You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied.  See the License for the
+ * specific language governing permissions and limitations
+ * under the License.ø
+ */
+
 /*!
  *  Copyright (c) 2019 by Contributors
  * \file np_delete_op.cu

From 410faa4c5b72e83d8a6d842d9db4a6a5e447bdc5 Mon Sep 17 00:00:00 2001
From: JiangZhaoh <jiang_zhaoh@163.com>
Date: Tue, 24 Dec 2019 16:57:04 +0800
Subject: [PATCH 4/5] add comment / abstarct basic function

---
 src/operator/numpy/np_delete_op-inl.h | 128 +++++++++++++++++---------
 1 file changed, 84 insertions(+), 44 deletions(-)

diff --git a/src/operator/numpy/np_delete_op-inl.h b/src/operator/numpy/np_delete_op-inl.h
index bcdb126f233f..8a9fcb30840e 100644
--- a/src/operator/numpy/np_delete_op-inl.h
+++ b/src/operator/numpy/np_delete_op-inl.h
@@ -73,11 +73,15 @@ struct NumpyDeleteParam : public dmlc::Parameter<NumpyDeleteParam> {
 };
 
 namespace delete_ {
+
 enum DeleteOpInputs {kArr, kObj};
 enum DeleteOpOutputs {kOut};
 }  // namespace delete_
 
 struct SliceToIndices {
+  /*!
+   * \brief transfer slice to indices array
+   */
   template<typename IType>
   MSHADOW_XINLINE static void Map(int i, IType* indices, int start, int step) {
     indices[i] = start + i * step;
@@ -85,6 +89,13 @@ struct SliceToIndices {
 };
 
 struct IsDeleteCal {
+  /*!
+   * \brief indicate which indices need to be deleted in input
+   * \param N used to check indices legality 
+   * \param is_delete if is_delete[i] == False, index i needn't to be deleted from output
+   *                  if is_delete[i] == True, index i need to be deleted from output
+   * \param indices the indices need to be deleted
+   */
   template<typename IType>
   MSHADOW_XINLINE static void Map(int i, int N, bool* is_delete, const IType* indices) {
     if ((indices[i] >= 0) && (indices[i] < N)) {
@@ -95,7 +106,10 @@ struct IsDeleteCal {
 
 struct OutPosCal {
   /*!
-   * \brief map the index from input to output
+   * \brief map the index from input to output. e.g.
+   * \example original_position 0 1 2 3 4
+   *          is_delete         F T T F F
+   *          out_position      0 - - 1 2
    */
   MSHADOW_XINLINE static void Map(int i, int64_t* out_pos, const bool* is_delete) {
     if (!is_delete[i]) {
@@ -111,7 +125,7 @@ struct OutPosCal {
 };
 
 template<int req, int ndim>
-struct DeleteImpl {
+struct DeleteKernel {
   /*!
    * \brief delete a sub-array from input along an axis according to 'is_delete'.
    * \param out_data - output: a new array with sub-arrays along an axis deleted.
@@ -140,6 +154,47 @@ struct DeleteImpl {
   }
 };
 
+/*!
+ * /brief equals to numpy's slice.indices(range)
+ * /param pstart - slice.start
+ * /param pstep - slice.step
+ * /param pstop - slice.stop
+ * /return start - slice.indices(range).start
+ * /return stop - slice.indices(range).stop
+ * /return step - slice.indices(range).step
+ * /return tot - total number of slice.indices(range)
+ */
+inline void SliceIndices(const dmlc::optional<int>& pstart,
+                         const dmlc::optional<int>& pstop,
+                         const dmlc::optional<int>& pstep,
+                         const int range,
+                         int* start, int* stop, int* step,
+                         size_t* tot) {
+  *step = pstep.has_value() ? pstep.value() : 1;
+  CHECK_NE(*step, 0) << "'step' can not equal to 0.";
+  if (pstop.has_value()) {
+    *stop = pstop.value();
+    *stop += (*stop < 0) ? range : 0;
+    *stop = (*stop < 0) ? ((*step < 0) ? -1 : 0) : *stop;
+    *stop = (*stop >= range) ? ((*step < 0) ? range - 1 : range) : *stop;
+  } else {
+    *stop = (*step > 0) ? range : -1;
+  }
+  if (pstart.has_value()) {
+    *start = pstart.value();
+    *start += (*start < 0) ? range : 0;
+    *start = (*start < 0) ? ((*step < 0) ? -1 : 0) : *start;
+    *start = (*start >= range) ? ((*step < 0) ? range - 1 : range) : *start;
+  } else {
+    *start = (*step > 0) ? 0 : range - 1;
+  }
+  if (*step > 0 && *stop >= *start) {
+    *tot = static_cast<size_t>((*stop - *start + *step - 1) / *step);
+  } else if (*step < 0 && *stop <= *start) {
+    *tot = static_cast<size_t>((*stop - *start + *step + 1) / *step);
+  }
+}
+
 template<typename xpu>
 void NumpyDeleteCompute(const nnvm::NodeAttrs& attrs,
                         const OpContext &ctx,
@@ -158,7 +213,7 @@ void NumpyDeleteCompute(const nnvm::NodeAttrs& attrs,
 
   int ndim = inputs[delete_::kArr].shape().ndim();
   int axis = param.axis.has_value() ? param.axis.value() : -1;
-  NDArray arr;
+  NDArray arr;  // original array
 
   if (!param.axis.has_value()) {
     arr = inputs[delete_::kArr].Reshape(Shape1(inputs[delete_::kArr].shape().Size()));
@@ -176,43 +231,25 @@ void NumpyDeleteCompute(const nnvm::NodeAttrs& attrs,
 
   axis = CheckAxis(axis, ndim);
   int N = (arr.shape())[axis];
-  mxnet::TShape newshape(arr.shape());
+  mxnet::TShape outshape(arr.shape());
+  // if obj is slice, they're obj's arguments
   int start = 0, stop = 0, step = 0;
+  // total number to be deleted
   size_t numtodel = 0;
+  // if obj is scaler, index is it's value
   int index = 0;
 
-  if (param.step.has_value()) {
-    step = param.step.value();
-    CHECK_NE(step, 0) << "'step' can not equal to 0.";
-    if (param.stop.has_value()) {
-      stop = param.stop.value();
-      stop += (stop < 0) ? N : 0;
-      stop = (stop < 0) ? ((step < 0) ? -1 : 0) : stop;
-      stop = (stop >= N) ? ((step < 0) ? N - 1 : N) : stop;
-    } else {
-      stop = (step > 0) ? N : -1;
-    }
-    if (param.start.has_value()) {
-      start = param.start.value();
-      start += (start < 0) ? N : 0;
-      start = (start < 0) ? ((step < 0) ? -1 : 0) : start;
-      start = (start >= N) ? ((step < 0) ? N - 1 : N) : start;
-    } else {
-      start = (step > 0) ? 0 : N - 1;
-    }
-    if (step > 0 && stop >= start) {
-      numtodel = static_cast<size_t>((stop - start + step - 1) / step);
-    } else if (step < 0 && stop <= start) {
-      numtodel = static_cast<size_t>((stop - start + step + 1) / step);
-    }
+  if (param.step.has_value()) {  // obj is slice
+    SliceIndices(param.start, param.stop, param.step,
+                 N, &start, &stop, &step, &numtodel);
     if (numtodel == 0) {
       const_cast<NDArray &>(outputs[delete_::kOut]).Init(arr.shape());
       mxnet_op::copy(s, outputs[delete_::kOut].data(), inputs[delete_::kArr].data());
       return;
     }
-    newshape[axis] -= numtodel;
-    const_cast<NDArray &>(outputs[delete_::kOut]).Init(newshape);
-  } else if (param.int_ind.has_value()) {
+    outshape[axis] -= numtodel;
+    const_cast<NDArray &>(outputs[delete_::kOut]).Init(outshape);
+  } else if (param.int_ind.has_value()) {  // obj is scaler
     index = param.int_ind.value();
     CHECK((index >= -1 * N) && (index < N))
       << "index " << index
@@ -220,13 +257,13 @@ void NumpyDeleteCompute(const nnvm::NodeAttrs& attrs,
       << " with size " << N << "\n";
     index += ((index < 0) ? N : 0);
     numtodel = static_cast<size_t>(1);
-    newshape[axis] -= 1;
-    const_cast<NDArray &>(outputs[delete_::kOut]).Init(newshape);
-  } else {
+    outshape[axis] -= 1;
+    const_cast<NDArray &>(outputs[delete_::kOut]).Init(outshape);
+  } else {  // obj is tensor
     numtodel = inputs[delete_::kObj].shape().Size();
   }
 
-  MSHADOW_TYPE_SWITCH(((inputs.size() == 2U) ?
+  MSHADOW_TYPE_SWITCH(((inputs.size() == 2U) ?  // obj is tensor
                       inputs[delete_::kObj].dtype() :
                       mshadow::DataType<int64_t>::kFlag), IType, {
     size_t temp_mem_size = sizeof(int64_t) * arr.shape()[axis] +
@@ -240,19 +277,22 @@ void NumpyDeleteCompute(const nnvm::NodeAttrs& attrs,
     bool* is_delete_ptr = reinterpret_cast<bool*>
                           (temp_mem.dptr_ + sizeof(int64_t) * arr.shape()[axis] +
                           sizeof(IType) * numtodel);
-    if (param.step.has_value()) {
+    if (param.step.has_value()) {  // obj is slice, transfer slice to tensor
       Kernel<SliceToIndices, xpu>::Launch(s, numtodel, indices_ptr, start, step);
-    } else if (param.int_ind.has_value()) {
+    } else if (param.int_ind.has_value()) {  // obj is scaler, copy it to tensor
       Fill(s, TBlob(indices_ptr, Shape1(numtodel), xpu::kDevMask), kWriteTo, index);
-    } else {
+    } else {  // obj is tensor, copy it to a unified tensor
       mxnet_op::copy(s,
         TBlob(indices_ptr, inputs[delete_::kObj].shape(), inputs[delete_::kObj].data().dev_mask()),
         inputs[delete_::kObj].data());
     }
     mxnet_op::Kernel<mxnet_op::set_zero, xpu>::Launch(s, arr.shape()[axis], is_delete_ptr);
+    // mark which position need to be deleted from input arr
     Kernel<IsDeleteCal, xpu>::Launch(s, numtodel, N, is_delete_ptr, indices_ptr);
+    // calculate output data's original position in input arr
     Kernel<OutPosCal, xpu>::Launch(s, arr.shape()[axis], out_pos_ptr, is_delete_ptr);
-    if (inputs.size() == 2U) {
+    if (inputs.size() == 2U) {  // obj is tensor
+      // get total number of nonredundant indices
       #ifdef __CUDACC__
         thrust::device_ptr<bool>is_delete_dev(is_delete_ptr);
         thrust::device_vector<bool>vec_is_delete(is_delete_dev, is_delete_dev + arr.shape()[axis]);
@@ -265,14 +305,14 @@ void NumpyDeleteCompute(const nnvm::NodeAttrs& attrs,
           numtodel++;
         }
       }
-      newshape[axis] -= numtodel;
-      const_cast<NDArray &>(outputs[delete_::kOut]).Init(newshape);
+      outshape[axis] -= numtodel;
+      const_cast<NDArray &>(outputs[delete_::kOut]).Init(outshape);
     }
-    MXNET_NDIM_SWITCH(newshape.ndim(), ndim, {
-      mshadow::Shape<ndim> out_strides = mxnet_op::calc_stride(newshape.get<ndim>());
+    MXNET_NDIM_SWITCH(outshape.ndim(), ndim, {
+      mshadow::Shape<ndim> out_strides = mxnet_op::calc_stride(outshape.get<ndim>());
       MSHADOW_TYPE_SWITCH(outputs[delete_::kOut].dtype(), DType, {
         MXNET_ASSIGN_REQ_SWITCH(req[delete_::kOut], req_type, {
-          Kernel<DeleteImpl<req_type, ndim>, xpu>::Launch(
+          Kernel<DeleteKernel<req_type, ndim>, xpu>::Launch(
             s, arr.shape().Size(),
             outputs[delete_::kOut].data().dptr<DType>(),
             arr.data().dptr<DType>(),

From c3ff130fbc7a1fd94a24d21416f06040c426a4c9 Mon Sep 17 00:00:00 2001
From: JiangZhaoh <jiang_zhaoh@163.com>
Date: Thu, 26 Dec 2019 16:21:15 +0800
Subject: [PATCH 5/5] fix windows build error - remove '#' from
 MSHADOW_TYPE_SWITCH

---
 src/operator/numpy/np_delete_op-inl.h | 72 ++++++++++++++++-----------
 1 file changed, 44 insertions(+), 28 deletions(-)

diff --git a/src/operator/numpy/np_delete_op-inl.h b/src/operator/numpy/np_delete_op-inl.h
index 8a9fcb30840e..a144833f3294 100644
--- a/src/operator/numpy/np_delete_op-inl.h
+++ b/src/operator/numpy/np_delete_op-inl.h
@@ -263,6 +263,9 @@ void NumpyDeleteCompute(const nnvm::NodeAttrs& attrs,
     numtodel = inputs[delete_::kObj].shape().Size();
   }
 
+  char* out_pos_ptr = NULL;
+  char* indices_ptr = NULL;
+  char* is_delete_ptr = NULL;
   MSHADOW_TYPE_SWITCH(((inputs.size() == 2U) ?  // obj is tensor
                       inputs[delete_::kObj].dtype() :
                       mshadow::DataType<int64_t>::kFlag), IType, {
@@ -271,43 +274,55 @@ void NumpyDeleteCompute(const nnvm::NodeAttrs& attrs,
                            sizeof(bool) * arr.shape()[axis];
     Tensor<xpu, 1, char> temp_mem =
       ctx.requested[0].get_space_typed<xpu, 1, char>(Shape1(temp_mem_size), s);
-    int64_t* out_pos_ptr = reinterpret_cast<int64_t*>(temp_mem.dptr_);
-    IType* indices_ptr = reinterpret_cast<IType*>
-                         (temp_mem.dptr_ + sizeof(int64_t) * arr.shape()[axis]);
-    bool* is_delete_ptr = reinterpret_cast<bool*>
-                          (temp_mem.dptr_ + sizeof(int64_t) * arr.shape()[axis] +
-                          sizeof(IType) * numtodel);
+    out_pos_ptr = temp_mem.dptr_;
+    indices_ptr = out_pos_ptr + sizeof(int64_t) * arr.shape()[axis];
+    is_delete_ptr = indices_ptr + sizeof(IType) * numtodel;
     if (param.step.has_value()) {  // obj is slice, transfer slice to tensor
-      Kernel<SliceToIndices, xpu>::Launch(s, numtodel, indices_ptr, start, step);
+      Kernel<SliceToIndices, xpu>::Launch(
+        s, numtodel, reinterpret_cast<IType*>(indices_ptr), start, step);
     } else if (param.int_ind.has_value()) {  // obj is scaler, copy it to tensor
-      Fill(s, TBlob(indices_ptr, Shape1(numtodel), xpu::kDevMask), kWriteTo, index);
+      Fill(s, TBlob(reinterpret_cast<IType*>(indices_ptr),
+           Shape1(numtodel), xpu::kDevMask), kWriteTo, index);
     } else {  // obj is tensor, copy it to a unified tensor
       mxnet_op::copy(s,
-        TBlob(indices_ptr, inputs[delete_::kObj].shape(), inputs[delete_::kObj].data().dev_mask()),
+        TBlob(reinterpret_cast<IType*>(indices_ptr), inputs[delete_::kObj].shape(),
+              inputs[delete_::kObj].data().dev_mask()),
         inputs[delete_::kObj].data());
     }
-    mxnet_op::Kernel<mxnet_op::set_zero, xpu>::Launch(s, arr.shape()[axis], is_delete_ptr);
+    mxnet_op::Kernel<mxnet_op::set_zero, xpu>::Launch(
+      s, arr.shape()[axis], reinterpret_cast<bool*>(is_delete_ptr));
     // mark which position need to be deleted from input arr
-    Kernel<IsDeleteCal, xpu>::Launch(s, numtodel, N, is_delete_ptr, indices_ptr);
+    Kernel<IsDeleteCal, xpu>::Launch(
+      s, numtodel, N, reinterpret_cast<bool*>(is_delete_ptr),
+      reinterpret_cast<IType*>(indices_ptr));
     // calculate output data's original position in input arr
-    Kernel<OutPosCal, xpu>::Launch(s, arr.shape()[axis], out_pos_ptr, is_delete_ptr);
-    if (inputs.size() == 2U) {  // obj is tensor
-      // get total number of nonredundant indices
-      #ifdef __CUDACC__
-        thrust::device_ptr<bool>is_delete_dev(is_delete_ptr);
-        thrust::device_vector<bool>vec_is_delete(is_delete_dev, is_delete_dev + arr.shape()[axis]);
-      #else
-        std::vector<bool>vec_is_delete(is_delete_ptr, is_delete_ptr + arr.shape()[axis]);
-      #endif
-      numtodel = 0;
-      for (int i = 0; i < arr.shape()[axis]; ++i) {
-        if (vec_is_delete[i]) {
-          numtodel++;
-        }
+    Kernel<OutPosCal, xpu>::Launch(
+      s, arr.shape()[axis], reinterpret_cast<int64_t*>(out_pos_ptr),
+      reinterpret_cast<bool*>(is_delete_ptr));
+  });
+
+  if (inputs.size() == 2U) {  // obj is tensor
+    // get total number of nonredundant indices
+    #ifdef __CUDACC__
+      thrust::device_ptr<bool>is_delete_dev(reinterpret_cast<bool*>(is_delete_ptr));
+      thrust::device_vector<bool>vec_is_delete(is_delete_dev, is_delete_dev + arr.shape()[axis]);
+    #else
+      std::vector<bool>vec_is_delete(reinterpret_cast<bool*>(is_delete_ptr),
+                                     reinterpret_cast<bool*>(is_delete_ptr) + arr.shape()[axis]);
+    #endif
+    numtodel = 0;
+    for (int i = 0; i < arr.shape()[axis]; ++i) {
+      if (vec_is_delete[i]) {
+        numtodel++;
       }
-      outshape[axis] -= numtodel;
-      const_cast<NDArray &>(outputs[delete_::kOut]).Init(outshape);
     }
+    outshape[axis] -= numtodel;
+    const_cast<NDArray &>(outputs[delete_::kOut]).Init(outshape);
+  }
+
+  MSHADOW_TYPE_SWITCH(((inputs.size() == 2U) ?  // obj is tensor
+                      inputs[delete_::kObj].dtype() :
+                      mshadow::DataType<int64_t>::kFlag), IType, {
     MXNET_NDIM_SWITCH(outshape.ndim(), ndim, {
       mshadow::Shape<ndim> out_strides = mxnet_op::calc_stride(outshape.get<ndim>());
       MSHADOW_TYPE_SWITCH(outputs[delete_::kOut].dtype(), DType, {
@@ -316,7 +331,8 @@ void NumpyDeleteCompute(const nnvm::NodeAttrs& attrs,
             s, arr.shape().Size(),
             outputs[delete_::kOut].data().dptr<DType>(),
             arr.data().dptr<DType>(),
-            is_delete_ptr, out_pos_ptr,
+            reinterpret_cast<bool*>(is_delete_ptr),
+            reinterpret_cast<int64_t*>(out_pos_ptr),
             arr.shape().get<ndim>(),
             out_strides, axis);
         });