You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@mxnet.apache.org by GitBox <gi...@apache.org> on 2020/04/17 10:08:39 UTC
[GitHub] [incubator-mxnet] JiangZhaoh opened a new pull request #18089: Add
npx op 'index_add'
JiangZhaoh opened a new pull request #18089: Add npx op 'index_add'
URL: https://github.com/apache/incubator-mxnet/pull/18089
## Description ##
Forward for index_add.
## Checklist ##
### Essentials ###
Please feel free to remove inapplicable items for your PR.
- [ ] The PR title starts with [MXNET-$JIRA_ID], where $JIRA_ID refers to the relevant [JIRA issue](https://issues.apache.org/jira/projects/MXNET/issues) created (except PRs with tiny changes)
- [ ] Changes are complete (i.e. I finished coding on this PR)
- [ ] All changes have test coverage:
- Unit tests are added for small changes to verify correctness (e.g. adding a new operator)
- Nightly tests are added for complicated/long-running ones (e.g. changing distributed kvstore)
- Build tests will be added for build configuration changes (e.g. adding a new build option with NCCL)
- [ ] Code is well-documented:
- For user-facing API changes, API doc string has been updated.
- For new C++ functions in header files, their functionalities and arguments are documented.
- For new examples, README.md is added to explain the what the example does, the source of the dataset, expected performance on test set and reference to the original paper if applicable
- Check the API doc at https://mxnet-ci-doc.s3-accelerate.dualstack.amazonaws.com/PR-$PR_ID/$BUILD_ID/index.html
- [ ] To the best of my knowledge, examples are either not affected by this change, or have been fixed to be compatible with this change
### Changes ###
- [ ] Feature1, tests, (and when applicable, API doc)
- [ ] Feature2, tests, (and when applicable, API doc)
## Comments ##
- If this change is a backward incompatible change, why must this change be made.
- Interesting edge cases to note here
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] mxnet-bot commented on issue #18089: Add npx op
'index_add'
Posted by GitBox <gi...@apache.org>.
mxnet-bot commented on issue #18089: Add npx op 'index_add'
URL: https://github.com/apache/incubator-mxnet/pull/18089#issuecomment-615162641
Hey @JiangZhaoh , Thanks for submitting the PR
All tests are already queued to run once. If tests fail, you can trigger one or more tests again with the following commands:
- To trigger all jobs: @mxnet-bot run ci [all]
- To trigger specific jobs: @mxnet-bot run ci [job1, job2]
***
**CI supported jobs**: [website, centos-gpu, centos-cpu, sanity, windows-cpu, unix-cpu, unix-gpu, windows-gpu, clang, edge, miscellaneous]
***
_Note_:
Only following 3 categories can trigger CI :PR Author, MXNet Committer, Jenkins Admin.
All CI tests must pass before the PR can be merged.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] sxjscience commented on a change in pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
sxjscience commented on a change in pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#discussion_r429055166
##########
File path: src/operator/tensor/index_add_backward.cu
##########
@@ -0,0 +1,119 @@
+/*
+ * 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.
+ */
+
+/*!
+ * \file index_add.cu
+ * \brief GPU implementation of index_add operator
+ */
+
+#include <cub/cub.cuh>
+#include "./index_add-inl.h"
+#include "../tensor/util/tensor_util-inl.cuh"
+#include "../tensor/util/tensor_util-inl.h"
+
+namespace mxnet {
+namespace op {
+
+template<typename xpu, typename DType>
+void IndexAddOpBackwardACalc(mshadow::Stream<xpu> *s,
+ DType* grad_a, const DType* ograd,
+ const size_t* stride,
+ const size_t tail_size, const int ind_num,
+ const int ind_ndim, const int32_t* ind_vec,
+ const int req, const int out_ndim) {
+ using namespace mxnet_op;
+ using namespace mshadow;
+ size_t* d_stride;
+ size_t shape_size = sizeof(size_t) * out_ndim;
+ cudaMalloc(reinterpret_cast<void**>(&d_stride), shape_size);
+ cudaMemcpy(d_stride, stride, shape_size, cudaMemcpyHostToDevice);
+ Kernel<IndexAddBackwardAKernel<DType>, xpu>::Launch(
+ s, ind_num, grad_a, ograd,
+ d_stride, tail_size,
+ ind_num, ind_ndim, ind_vec, req);
+ cudaFree(d_stride);
Review comment:
Instead of cudaMalloc + cudaFree. Use the memory management system to allocate temporary memory.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] sxjscience commented on a change in pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
sxjscience commented on a change in pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#discussion_r430831424
##########
File path: src/operator/tensor/index_add-inl.h
##########
@@ -0,0 +1,275 @@
+/*
+ * 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.
+ */
+
+/*!
+ * \file index_add-inl.h
+ * \brief Function definition of index_add operator
+*/
+#ifndef MXNET_OPERATOR_TENSOR_INDEX_ADD_INL_H_
+#define MXNET_OPERATOR_TENSOR_INDEX_ADD_INL_H_
+
+#include <mxnet/operator_util.h>
+#include <vector>
+#include <algorithm>
+#include "../mxnet_op.h"
+#include "../operator_common.h"
+#include "../elemwise_op_common.h"
+
+namespace mxnet {
+namespace op {
+
+inline bool IndexModifyOpShape(const nnvm::NodeAttrs& attrs,
+ mxnet::ShapeVector* in_attrs,
+ mxnet::ShapeVector* out_attrs) {
+ CHECK_EQ(in_attrs->size(), 3U);
+ CHECK_EQ(out_attrs->size(), 1U);
+ SHAPE_ASSIGN_CHECK(*out_attrs, 0, (*in_attrs)[0]);
+ return true;
+}
+
+inline bool IndexModifyOpType(const nnvm::NodeAttrs& attrs,
+ std::vector<int>* in_attrs,
+ std::vector<int>* out_attrs) {
+ CHECK_EQ(in_attrs->size(), 3U);
+ CHECK_EQ(out_attrs->size(), 1U);
+ CHECK_NE((*in_attrs)[0], -1);
+ CHECK_NE((*in_attrs)[1], -1);
+ CHECK_NE((*in_attrs)[2], -1);
+ CHECK_EQ((*in_attrs)[0], (*in_attrs)[2])
+ << "index_add(a, ind, val) only support a.dtype == val.dtype";
+ CHECK((*in_attrs)[1] == mshadow::kInt64 ||
+ (*in_attrs)[1] == mshadow::kInt32)
+ << "'ind' only support int dtype.";
+ TYPE_ASSIGN_CHECK(*out_attrs, 0, (*in_attrs)[0]);
+ return (*out_attrs)[0] != -1;
+}
+
+template<typename xpu, typename DType>
+void IndexAddForwardCalc(mshadow::Stream<xpu> *s,
+ const int ind_num, DType* out,
+ const DType* val,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_tail_shape,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_pre_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_shape,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_shape,
+ const int a_tail_size,
+ const int ind_ndim, const int* ind,
+ const int a_ndim);
+
+template<typename xpu>
+void IndexAddOpForward(const nnvm::NodeAttrs& attrs,
+ const OpContext& ctx,
+ const std::vector<TBlob>& inputs,
+ const std::vector<OpReqType>& req,
+ const std::vector<TBlob>& outputs) {
+ using namespace mxnet_op;
+ using namespace mshadow;
+ CHECK_EQ(inputs.size(), 3U);
+ CHECK_EQ(outputs.size(), 1U);
+ Stream<xpu> *s = ctx.get_stream<xpu>();
+ const TBlob a = inputs[0];
+ TBlob ind = inputs[1];
+ TBlob val = inputs[2];
+ TBlob out = outputs[0];
+ CHECK_GT(a.shape_.ndim(), 0) << "The first input is saclar, please use '+' instead.";
+ int a_ndim = a.shape_.ndim();
+ CHECK_LE(a_ndim, MXNET_SPECIAL_MAX_NDIM)
+ << "ndim should less than "<< MXNET_SPECIAL_MAX_NDIM
+ << "but get " << a_ndim <<"\n";
+ int val_ndim = val.shape_.ndim();
+ if (val_ndim == 0) {
+ val.shape_ = Shape1(1);
+ val_ndim = 1;
+ }
+ // ind=np.array([]), ind.shape_.ndim() = 1
+ // ind=np.array(1), ind.shape_.ndim() = 0
+ // ind=np.array([[0,0],[0,1]]), ind.shape_.ndim() = 2
+ CHECK_NE(ind.shape_.Size(), 0) << "Param 'ind' is []. Please just use op 'add' instead.\n";
+ CHECK_LE(ind.shape_.ndim(), 2) << "'ind' array allow 2 dimension at most.";
+ if (ind.shape_.ndim() == 0) {
+ ind.shape_ = Shape2(1, 1);
+ } else if (ind.shape_.ndim() == 1) {
+ ind.shape_ = Shape2(1, ind.shape_[0]);
+ }
+ int ind_ndim = ind.shape_[0];
+ int ind_num = ind.shape_[1];
+ CHECK_LE(ind_ndim, a_ndim) << "IndexError: too many indices for array.";
+
+ // check 'val' broadcast legality
+ CHECK_LE(val_ndim, a_ndim - ind_ndim + 1)
+ << "The ndim of param 'val' is " << val_ndim
+ << ", but it should less than or equal to " << a_ndim - ind_ndim + 1;
+ for (int i = a_ndim - 1, j = val_ndim - 1; j >= 0 ; --i, --j) {
+ if ((j == 0) && (val_ndim == a_ndim - ind_ndim + 1)) {
+ // val_ndim == a_ndim - ind_ndim + 1, check the first dim of input 'val'
+ CHECK(val.shape_[j] == ind_num || val.shape_[j] == 1)
+ << "can not broadcast from " << val.shape_[j] << " to " << ind_num;
+ } else {
+ CHECK(val.shape_[j] == a.shape_[i] || val.shape_[j] == 1)
+ << "can not broadcast from " << val.shape_[j] << " to " << a.shape_[i]
+ << " in axis " << i;
+ }
+ }
+ int a_tail_size = static_cast<int>(a.shape_.ProdShape(ind_ndim, a_ndim));
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM>a_shape, val_shape;
+ for (int i = MXNET_SPECIAL_MAX_NDIM - 1, j = a_ndim - 1; i >= 0; --i, --j) {
+ a_shape[i] = (j >= 0) ? a.shape_[j] : 1;
+ }
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM>a_pre_shape(a_shape);
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM>a_tail_shape(a_shape);
+
+ int seg = MXNET_SPECIAL_MAX_NDIM - a_ndim;
+ for (int i = seg; i < ind_ndim + seg; ++i) {
+ a_tail_shape[i] = 1;
+ }
+ for (int i = ind_ndim + seg; i < a_ndim + seg; ++i) {
+ a_pre_shape[i] = 1;
+ }
+ for (int i = MXNET_SPECIAL_MAX_NDIM - 1, j = val_ndim - 1; i >= 0; --i, --j) {
+ val_shape[i] = (j >= 0) ? val.shape_[j] : 1;
+ }
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM>a_pre_stride = calc_stride(a_pre_shape);
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM>val_stride = calc_stride(val_shape);
+ mxnet_op::copy(s, out, a);
+ Tensor<xpu, 1, char> temp_mem = ctx.requested[0].get_space_typed<xpu, 1, char>(
+ Shape1(sizeof(int) * ind_ndim * ind_num), s);
+ TBlob t_ind = TBlob(reinterpret_cast<int*>(temp_mem.dptr_),
+ Shape1(ind_ndim * ind_num), xpu::kDevMask);
+ mxnet_op::copy(s, t_ind, ind);
+ MSHADOW_TYPE_SWITCH(a.type_flag_, DType, {
+ IndexAddForwardCalc<xpu, DType>(s, ind_num,
+ out.dptr<DType>(), val.dptr<DType>(),
+ a_tail_shape, a_pre_stride,
+ val_stride, val_shape, a_shape,
+ a_tail_size, ind_ndim,
+ t_ind.dptr<int>(), a_ndim);
+ });
+}
+
+template<typename DType>
+struct IndexAddBackwardAKernel {
+ MSHADOW_XINLINE static void Map(size_t i, DType* grad_a,
+ const DType* ograd,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> stride,
+ const int tail_size, const int ind_num, const int ind_ndim,
+ const int32_t* ind_vec, const int req, const int out_ndim) {
+ index_t id = 0;
+ int seg = MXNET_SPECIAL_MAX_NDIM - out_ndim;
+ for (int dim = 0; dim < ind_ndim; ++dim) {
+ id += stride[seg + dim] * ind_vec[dim * ind_num + i];
+ }
+ for (int _i = 0; _i < tail_size; ++_i) {
+ KERNEL_ASSIGN(grad_a[id + _i], req, ograd[id + _i]);
+ }
+ }
+};
+
+template<typename xpu, typename DType>
+void IndexAddOpBackwardACalc(mshadow::Stream<xpu> *s,
+ DType* grad_a, const DType* ograd,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> stride,
+ const int tail_size, const int ind_num,
+ const int ind_ndim, const int32_t* ind_vec,
+ const int req, const int out_ndim);
+
+template<typename xpu, typename DType>
+void IndexAddOpBackwardValCalc(mshadow::Stream<xpu> *s,
+ DType* grad_val, const DType* ograd,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> ograd_tail_shape,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> ograd_pre_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_shape,
+ const int tail_size, const int ind_num,
+ const int ind_ndim, const int32_t* ind_vec,
+ const int out_ndim);
+
+template<typename xpu>
+void IndexAddOpBackward(const nnvm::NodeAttrs& attrs,
+ const OpContext& ctx,
+ const std::vector<TBlob>& inputs,
+ const std::vector<OpReqType>& req,
+ const std::vector<TBlob>& outputs) {
+ using namespace mshadow;
+ using namespace mxnet_op;
+ if (req[0] == kNullOp && req[2] == kNullOp) {
+ return;
+ }
+ CHECK_EQ(inputs.size(), 4);
+ CHECK_EQ(outputs.size(), 3);
+ const TBlob& ograd = inputs[0];
+ const TBlob& a = inputs[1];
+ TBlob ind = inputs[2];
+ const TBlob& val = inputs[3];
+ const TBlob& grad_a = outputs[0];
+ const TBlob& grad_val = outputs[2];
+ mshadow::Stream<xpu> *s = ctx.get_stream<xpu>();
+ // get the number of 'ind' index
+ if (ind.shape_.ndim() == 0) {
+ ind.shape_ = Shape2(1, 1);
+ } else if (ind.shape_.ndim() == 1) {
+ ind.shape_ = Shape2(1, ind.shape_[0]);
+ }
+ int ind_ndim = ind.shape_[0];
+ int ind_num = ind.shape_[1];
+ // broadcast 'ind'
+ int ndim = ograd.shape_.ndim();
+ int tail_size = static_cast<int>(ograd.shape_.ProdShape(ind_ndim, ndim));
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM>ograd_shape, val_shape;
+ for (int i = MXNET_SPECIAL_MAX_NDIM - 1, j = ndim - 1; i >= 0; --i, --j) {
+ ograd_shape[i] = (j >= 0) ? ograd.shape_[j] : 1;
+ }
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM>ograd_stride = mxnet_op::calc_stride(ograd_shape);
Review comment:
nit: space: `mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> ograd_stride`
##########
File path: src/operator/tensor/index_add_backward.cc
##########
@@ -0,0 +1,107 @@
+/*
+ * 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.
+ */
+
+/*!
+ * \file index_add-inl.cc
+ * \brief CPU implementation of index_add operator
+*/
+#include <vector>
+#include "./index_add-inl.h"
+
+namespace mxnet {
+namespace op {
+
+template<typename xpu, typename DType>
+void IndexAddOpBackwardACalc(mshadow::Stream<xpu> *s,
+ DType* grad_a, const DType* ograd,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> stride,
+ const int tail_size, const int ind_num,
+ const int ind_ndim, const int32_t* ind_vec,
+ const int req, const int out_ndim) {
+ using namespace mxnet_op;
+ using namespace mshadow;
+ Kernel<IndexAddBackwardAKernel<DType>, xpu>::Launch(
+ s, ind_num, grad_a, ograd, stride, tail_size, ind_num, ind_ndim, ind_vec, req, out_ndim);
+}
+
+template<typename DType>
+struct IndexAddBackwardValCPUKernel {
+ MSHADOW_XINLINE static void Map(size_t i, DType* grad_val,
+ const DType* ograd,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> ograd_tail_shape,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> ograd_pre_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_shape,
+ const int ograd_tail_size, const int ind_num,
+ const int ind_ndim, const int32_t* ind_vec,
+ const int out_ndim) {
+ index_t id = 0;
+ int seg = MXNET_SPECIAL_MAX_NDIM - out_ndim;
+ for (int dim = 0; dim < ind_ndim; ++dim) {
+ id += ograd_pre_stride[seg + dim] * ind_vec[dim * ind_num + i];
+ }
+ id *= ograd_tail_size;
+ #pragma omp parallel for
+ for (int _i = 0; _i < ograd_tail_size; ++_i) {
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> ograd_tail_id =
+ mxnet_op::unravel(_i, ograd_tail_shape);
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_id;
+ for (int _j = seg; _j < seg + out_ndim; ++_j) {
+ val_id[_j] = (val_shape[_j] == 1) ? 0 : ograd_tail_id[_j];
+ }
+ val_id[seg + ind_ndim - 1] = (val_shape[seg + ind_ndim - 1] == 1) ? 0 : i;
+ index_t val_dest = mxnet_op::dot(val_id, val_stride);
+ #pragma omp critical
+ {
+ grad_val[val_dest] += ograd[id + _i];
+ }
+ }
+ }
+};
+
+template<typename xpu, typename DType>
+void IndexAddOpBackwardValCalc(mshadow::Stream<xpu> *s,
+ DType* grad_val, const DType* ograd,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> ograd_tail_shape,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> ograd_pre_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_shape,
+ const int tail_size, const int ind_num,
+ const int ind_ndim, const int32_t* ind_vec,
+ const int out_ndim) {
+ using namespace mxnet_op;
+ using namespace mshadow;
+ Kernel<IndexAddBackwardValCPUKernel<DType>, xpu>::Launch(
+ s, ind_num, grad_val, ograd, ograd_tail_shape, ograd_pre_stride,
+ val_stride, val_shape, tail_size, ind_num, ind_ndim, ind_vec, out_ndim);
+}
+
+NNVM_REGISTER_OP(_backward_index_add)
+.set_num_inputs(4)
+.set_num_outputs(3)
Review comment:
We may just use 2 here.
##########
File path: src/operator/tensor/index_add_forward.cc
##########
@@ -0,0 +1,117 @@
+/*
+ * 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.
+ */
+
+/*!
+ * \file index_add-inl.cc
+ * \brief CPU implementation of index_add operator
+*/
+#include <vector>
+#include "./index_add-inl.h"
+
+namespace mxnet {
+namespace op {
+template<typename DType>
+struct IndexAddForwardCPUKernel {
+ MSHADOW_XINLINE static void Map(size_t i, DType* out,
+ const DType* val,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_tail_shape,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_pre_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_shape,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_shape,
+ const int a_tail_size, const int ind_num,
+ const int ind_ndim, const int* ind,
+ const int a_ndim) {
+ index_t id = 0;
+ int seg = MXNET_SPECIAL_MAX_NDIM - a_ndim;
+ for (int dim = 0; dim < ind_ndim; ++dim) {
+ CHECK_LT(ind[dim * ind_num + i], a_shape[seg + dim])
+ << "IndexError: index " << ind[dim * ind_num + i]
+ << " is out of bounds for axis " << dim
+ << " with size " << a_shape[seg + dim];
+ CHECK_GE(ind[dim * ind_num + i], 0)
+ << "IndexError: index " << ind[dim * ind_num + i]
+ << " should be greater or equal to 0.";
+ id += a_pre_stride[seg + dim] * ind[dim * ind_num + i];
+ }
+ id *= a_tail_size;
+ #pragma omp parallel for
+ for (int _i = 0; _i < a_tail_size; ++_i) {
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_tail_id = mxnet_op::unravel(_i, a_tail_shape);
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_id;
+ for (int _j = seg; _j < seg + a_ndim; ++_j) {
+ val_id[_j] = (val_shape[_j] == 1) ? 0 : a_tail_id[_j];
+ }
+ val_id[seg + ind_ndim - 1] = (val_shape[seg + ind_ndim - 1] == 1) ? 0 : i;
+ index_t val_dest = mxnet_op::dot(val_id, val_stride);
+ #pragma omp critical
+ {
+ out[id + _i] += val[val_dest];
+ }
+ }
+ }
+};
+
+template<typename xpu, typename DType>
+void IndexAddForwardCalc(mshadow::Stream<xpu> *s,
+ const int ind_num, DType* out,
+ const DType* val,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_tail_shape,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_pre_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_shape,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_shape,
+ const int a_tail_size,
+ const int ind_ndim, const int* ind,
+ const int a_ndim) {
+ using namespace mxnet_op;
+ using namespace mshadow;
+ Kernel<IndexAddForwardCPUKernel<DType>, xpu>::Launch(
+ s, ind_num, out, val,
+ a_tail_shape, a_pre_stride,
+ val_stride, val_shape, a_shape,
+ a_tail_size, ind_num,
+ ind_ndim, ind, a_ndim);
+}
+
+
+
+NNVM_REGISTER_OP(_npx_index_add)
+.describe(R"code(This operators implements the "+=" mimic function.
+)code" ADD_FILELINE)
+.set_num_inputs(3)
+.set_num_outputs(1)
+.set_attr<nnvm::FListInputNames>("FListInputNames",
+ [](const NodeAttrs& attrs) {
+ return std::vector<std::string>{"a", "ind", "val"};
+ })
+.set_attr<mxnet::FInferShape>("FInferShape", IndexModifyOpShape)
+.set_attr<nnvm::FInferType>("FInferType", IndexModifyOpType)
+.set_attr<FCompute>("FCompute<cpu>", IndexAddOpForward<cpu>)
+.set_attr<FResourceRequest>("FResourceRequest",
+ [](const NodeAttrs& attrs) {
+ return std::vector<ResourceRequest>{ResourceRequest::kTempSpace};
+ })
+.set_attr<nnvm::FGradient>("FGradient", ElemwiseGradUseIn{"_backward_index_add"})
Review comment:
You may separately register `_backward_index_add_a`, `_backward_index_add_val` and then register the gradient like this:
```c++
.set_attr<nnvm::FGradient>("FGradient",
[](const nnvm::ObjectPtr& n, const std::vector<nnvm::NodeEntry>& ograds) {
auto a_grad = MakeNode("_backward_index_add_a", n->attrs.name + "_backward_a",
{ograds[0], n->inputs[0]}, nullptr, &n);
auto idx_grad = MakeNode("zeros_like", n->attrs.name + "_backward_indices",
{n->inputs[1]}, nullptr, &n);
auto val_grad = MakeNode("_backward_index_add_val", n->attrs.name + "_backward_val",
{ograds[0], n->inputs[2]}, nullptr, &n);
std::vector<nnvm::NodeEntry> ret;
ret.emplace_back(a_grad);
ret.emplace_back(idx_grad);
ret.emplace_back(val_grad);
return ret;
})
```
##########
File path: src/operator/tensor/index_add_forward.cc
##########
@@ -0,0 +1,117 @@
+/*
+ * 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.
+ */
+
+/*!
+ * \file index_add-inl.cc
+ * \brief CPU implementation of index_add operator
+*/
+#include <vector>
+#include "./index_add-inl.h"
+
+namespace mxnet {
+namespace op {
+template<typename DType>
+struct IndexAddForwardCPUKernel {
+ MSHADOW_XINLINE static void Map(size_t i, DType* out,
+ const DType* val,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_tail_shape,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_pre_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_shape,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_shape,
+ const int a_tail_size, const int ind_num,
+ const int ind_ndim, const int* ind,
+ const int a_ndim) {
+ index_t id = 0;
+ int seg = MXNET_SPECIAL_MAX_NDIM - a_ndim;
+ for (int dim = 0; dim < ind_ndim; ++dim) {
+ CHECK_LT(ind[dim * ind_num + i], a_shape[seg + dim])
+ << "IndexError: index " << ind[dim * ind_num + i]
+ << " is out of bounds for axis " << dim
+ << " with size " << a_shape[seg + dim];
+ CHECK_GE(ind[dim * ind_num + i], 0)
+ << "IndexError: index " << ind[dim * ind_num + i]
+ << " should be greater or equal to 0.";
+ id += a_pre_stride[seg + dim] * ind[dim * ind_num + i];
+ }
+ id *= a_tail_size;
+ #pragma omp parallel for
+ for (int _i = 0; _i < a_tail_size; ++_i) {
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_tail_id = mxnet_op::unravel(_i, a_tail_shape);
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_id;
+ for (int _j = seg; _j < seg + a_ndim; ++_j) {
+ val_id[_j] = (val_shape[_j] == 1) ? 0 : a_tail_id[_j];
+ }
+ val_id[seg + ind_ndim - 1] = (val_shape[seg + ind_ndim - 1] == 1) ? 0 : i;
+ index_t val_dest = mxnet_op::dot(val_id, val_stride);
Review comment:
Is it possible to merge them into one API call?
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] sxjscience commented on issue #18089: Add npx op
'index_add'
Posted by GitBox <gi...@apache.org>.
sxjscience commented on issue #18089: Add npx op 'index_add'
URL: https://github.com/apache/incubator-mxnet/pull/18089#issuecomment-615405466
Is it using the new FFI?
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
With regards,
Apache Git Services
[GitHub] [incubator-mxnet] sxjscience commented on a change in pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
sxjscience commented on a change in pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#discussion_r430856261
##########
File path: src/operator/tensor/index_add_backward.cc
##########
@@ -0,0 +1,107 @@
+/*
+ * 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.
+ */
+
+/*!
+ * \file index_add-inl.cc
+ * \brief CPU implementation of index_add operator
+*/
+#include <vector>
+#include "./index_add-inl.h"
+
+namespace mxnet {
+namespace op {
+
+template<typename xpu, typename DType>
+void IndexAddOpBackwardACalc(mshadow::Stream<xpu> *s,
+ DType* grad_a, const DType* ograd,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> stride,
+ const int tail_size, const int ind_num,
+ const int ind_ndim, const int32_t* ind_vec,
+ const int req, const int out_ndim) {
+ using namespace mxnet_op;
+ using namespace mshadow;
+ Kernel<IndexAddBackwardAKernel<DType>, xpu>::Launch(
+ s, ind_num, grad_a, ograd, stride, tail_size, ind_num, ind_ndim, ind_vec, req, out_ndim);
+}
+
+template<typename DType>
+struct IndexAddBackwardValCPUKernel {
+ MSHADOW_XINLINE static void Map(size_t i, DType* grad_val,
+ const DType* ograd,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> ograd_tail_shape,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> ograd_pre_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_shape,
+ const int ograd_tail_size, const int ind_num,
+ const int ind_ndim, const int32_t* ind_vec,
+ const int out_ndim) {
+ index_t id = 0;
+ int seg = MXNET_SPECIAL_MAX_NDIM - out_ndim;
+ for (int dim = 0; dim < ind_ndim; ++dim) {
+ id += ograd_pre_stride[seg + dim] * ind_vec[dim * ind_num + i];
+ }
+ id *= ograd_tail_size;
+ #pragma omp parallel for
+ for (int _i = 0; _i < ograd_tail_size; ++_i) {
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> ograd_tail_id =
+ mxnet_op::unravel(_i, ograd_tail_shape);
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_id;
+ for (int _j = seg; _j < seg + out_ndim; ++_j) {
+ val_id[_j] = (val_shape[_j] == 1) ? 0 : ograd_tail_id[_j];
+ }
+ val_id[seg + ind_ndim - 1] = (val_shape[seg + ind_ndim - 1] == 1) ? 0 : i;
+ index_t val_dest = mxnet_op::dot(val_id, val_stride);
+ #pragma omp critical
+ {
+ grad_val[val_dest] += ograd[id + _i];
+ }
+ }
+ }
+};
+
+template<typename xpu, typename DType>
+void IndexAddOpBackwardValCalc(mshadow::Stream<xpu> *s,
+ DType* grad_val, const DType* ograd,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> ograd_tail_shape,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> ograd_pre_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_shape,
+ const int tail_size, const int ind_num,
+ const int ind_ndim, const int32_t* ind_vec,
+ const int out_ndim) {
+ using namespace mxnet_op;
+ using namespace mshadow;
+ Kernel<IndexAddBackwardValCPUKernel<DType>, xpu>::Launch(
+ s, ind_num, grad_val, ograd, ograd_tail_shape, ograd_pre_stride,
+ val_stride, val_shape, tail_size, ind_num, ind_ndim, ind_vec, out_ndim);
+}
Review comment:
It's not necessary to use Kernel::Launch.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] JiangZhaoh commented on pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
JiangZhaoh commented on pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#issuecomment-631412311
For `npx.index_add(a, ind, val)`, now it supports int32 or int64 tensor `ind`, and both int or float tensor `a` and `val`. But it requires the dtype of 'a' and 'val' keep same.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] sxjscience commented on a change in pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
sxjscience commented on a change in pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#discussion_r431592402
##########
File path: src/operator/tensor/index_add-inl.h
##########
@@ -0,0 +1,262 @@
+/*
+ * 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.
+ */
+
+/*!
+ * \file index_add-inl.h
+ * \brief Function definition of index_add operator
+*/
+#ifndef MXNET_OPERATOR_TENSOR_INDEX_ADD_INL_H_
+#define MXNET_OPERATOR_TENSOR_INDEX_ADD_INL_H_
+
+#include <mxnet/operator_util.h>
+#include <vector>
+#include <algorithm>
+#include "../mxnet_op.h"
+#include "../operator_common.h"
+#include "../elemwise_op_common.h"
+
+namespace mxnet {
+namespace op {
+
+inline bool IndexModifyOpShape(const nnvm::NodeAttrs& attrs,
+ mxnet::ShapeVector* in_attrs,
+ mxnet::ShapeVector* out_attrs) {
+ CHECK_EQ(in_attrs->size(), 3U);
+ CHECK_EQ(out_attrs->size(), 1U);
+ SHAPE_ASSIGN_CHECK(*out_attrs, 0, (*in_attrs)[0]);
+ return true;
+}
+
+inline bool IndexModifyOpType(const nnvm::NodeAttrs& attrs,
+ std::vector<int>* in_attrs,
+ std::vector<int>* out_attrs) {
+ CHECK_EQ(in_attrs->size(), 3U);
+ CHECK_EQ(out_attrs->size(), 1U);
+ CHECK_NE((*in_attrs)[0], -1);
+ CHECK_NE((*in_attrs)[1], -1);
+ CHECK_NE((*in_attrs)[2], -1);
+ CHECK_EQ((*in_attrs)[0], (*in_attrs)[2])
+ << "index_add(a, ind, val) only support a.dtype == val.dtype";
+ CHECK((*in_attrs)[1] == mshadow::kInt64 ||
+ (*in_attrs)[1] == mshadow::kInt32)
+ << "'ind' only support int dtype.";
+ TYPE_ASSIGN_CHECK(*out_attrs, 0, (*in_attrs)[0]);
+ return (*out_attrs)[0] != -1;
+}
+
+template<typename xpu, typename DType>
+void IndexAddForwardCalc(mshadow::Stream<xpu> *s,
+ const int ind_num, DType* out,
+ const DType* val,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_tail_shape,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_pre_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_shape,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_shape,
+ const int a_tail_size,
+ const int ind_ndim, const int* ind,
+ const int a_ndim);
+
+template<typename xpu>
+void IndexAddOpForward(const nnvm::NodeAttrs& attrs,
+ const OpContext& ctx,
+ const std::vector<TBlob>& inputs,
+ const std::vector<OpReqType>& req,
+ const std::vector<TBlob>& outputs) {
+ using namespace mxnet_op;
+ using namespace mshadow;
+ CHECK_EQ(inputs.size(), 3U);
+ CHECK_EQ(outputs.size(), 1U);
+ Stream<xpu> *s = ctx.get_stream<xpu>();
+ const TBlob a = inputs[0];
+ TBlob ind = inputs[1];
+ TBlob val = inputs[2];
+ TBlob out = outputs[0];
+ CHECK_GT(a.shape_.ndim(), 0) << "The first input is saclar, please use '+' instead.";
+ int a_ndim = a.shape_.ndim();
+ CHECK_LE(a_ndim, MXNET_SPECIAL_MAX_NDIM)
+ << "ndim should less than "<< MXNET_SPECIAL_MAX_NDIM
+ << "but get " << a_ndim <<"\n";
+ int val_ndim = val.shape_.ndim();
+ if (val_ndim == 0) {
+ val.shape_ = Shape1(1);
+ val_ndim = 1;
+ }
+ // ind=np.array([]), ind.shape_.ndim() = 1
+ // ind=np.array(1), ind.shape_.ndim() = 0
+ // ind=np.array([[0,0],[0,1]]), ind.shape_.ndim() = 2
+ CHECK_NE(ind.shape_.Size(), 0) << "Param 'ind' is []. Please just use op 'add' instead.\n";
+ CHECK_LE(ind.shape_.ndim(), 2) << "'ind' array allow 2 dimension at most.";
+ if (ind.shape_.ndim() == 0) {
+ ind.shape_ = Shape2(1, 1);
+ } else if (ind.shape_.ndim() == 1) {
+ ind.shape_ = Shape2(1, ind.shape_[0]);
+ }
+ int ind_ndim = ind.shape_[0];
+ int ind_num = ind.shape_[1];
+ CHECK_LE(ind_ndim, a_ndim) << "IndexError: too many indices for array.";
+
+ // check 'val' broadcast legality
+ CHECK_LE(val_ndim, a_ndim - ind_ndim + 1)
+ << "The ndim of param 'val' is " << val_ndim
+ << ", but it should less than or equal to " << a_ndim - ind_ndim + 1;
+ for (int i = a_ndim - 1, j = val_ndim - 1; j >= 0 ; --i, --j) {
+ if ((j == 0) && (val_ndim == a_ndim - ind_ndim + 1)) {
+ // val_ndim == a_ndim - ind_ndim + 1, check the first dim of input 'val'
+ CHECK(val.shape_[j] == ind_num || val.shape_[j] == 1)
+ << "can not broadcast from " << val.shape_[j] << " to " << ind_num;
+ } else {
+ CHECK(val.shape_[j] == a.shape_[i] || val.shape_[j] == 1)
+ << "can not broadcast from " << val.shape_[j] << " to " << a.shape_[i]
+ << " in axis " << i;
+ }
+ }
+ int a_tail_size = static_cast<int>(a.shape_.ProdShape(ind_ndim, a_ndim));
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_shape, val_shape;
+ for (int i = MXNET_SPECIAL_MAX_NDIM - 1, j = a_ndim - 1; i >= 0; --i, --j) {
+ a_shape[i] = (j >= 0) ? a.shape_[j] : 1;
+ }
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_pre_shape(a_shape);
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_tail_shape(a_shape);
+
+ int seg = MXNET_SPECIAL_MAX_NDIM - a_ndim;
+ for (int i = seg; i < ind_ndim + seg; ++i) {
+ a_tail_shape[i] = 1;
+ }
+ for (int i = ind_ndim + seg; i < a_ndim + seg; ++i) {
+ a_pre_shape[i] = 1;
+ }
+ for (int i = MXNET_SPECIAL_MAX_NDIM - 1, j = val_ndim - 1; i >= 0; --i, --j) {
+ val_shape[i] = (j >= 0) ? val.shape_[j] : 1;
+ }
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_pre_stride = calc_stride(a_pre_shape);
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_stride = calc_stride(val_shape);
+ mxnet_op::copy(s, out, a);
+ Tensor<xpu, 1, char> temp_mem = ctx.requested[0].get_space_typed<xpu, 1, char>(
+ Shape1(sizeof(int) * ind_ndim * ind_num), s);
Review comment:
Is it possible to directly allocate `int` typed storage?
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] sxjscience commented on a change in pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
sxjscience commented on a change in pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#discussion_r429054854
##########
File path: src/operator/tensor/index_add_backward.cu
##########
@@ -0,0 +1,119 @@
+/*
+ * 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.
+ */
+
+/*!
+ * \file index_add.cu
+ * \brief GPU implementation of index_add operator
+ */
+
+#include <cub/cub.cuh>
+#include "./index_add-inl.h"
+#include "../tensor/util/tensor_util-inl.cuh"
+#include "../tensor/util/tensor_util-inl.h"
+
+namespace mxnet {
+namespace op {
+
+template<typename xpu, typename DType>
+void IndexAddOpBackwardACalc(mshadow::Stream<xpu> *s,
+ DType* grad_a, const DType* ograd,
+ const size_t* stride,
+ const size_t tail_size, const int ind_num,
+ const int ind_ndim, const int32_t* ind_vec,
+ const int req, const int out_ndim) {
+ using namespace mxnet_op;
+ using namespace mshadow;
+ size_t* d_stride;
+ size_t shape_size = sizeof(size_t) * out_ndim;
+ cudaMalloc(reinterpret_cast<void**>(&d_stride), shape_size);
Review comment:
We may not use cudaMalloc here. Allocate additional temporary storage to do this.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] mxnet-bot commented on pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
mxnet-bot commented on pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#issuecomment-633556769
Jenkins CI successfully triggered : [windows-cpu, windows-gpu, unix-gpu]
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] ZheyuYe commented on pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
ZheyuYe commented on pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#issuecomment-628344520
Related to #17823
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] sxjscience commented on a change in pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
sxjscience commented on a change in pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#discussion_r433035858
##########
File path: src/operator/tensor/index_add_forward.cc
##########
@@ -0,0 +1,132 @@
+/*
+ * 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.
+ */
+
+/*!
+ * \file index_add-inl.cc
+ * \brief CPU implementation of index_add operator
+*/
+#include <vector>
+#include "./index_add-inl.h"
+
+namespace mxnet {
+namespace op {
+template<typename DType>
+struct IndexAddForwardCPUKernel {
+ MSHADOW_XINLINE static void Map(size_t i, DType* out,
+ const DType* val,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_tail_shape,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_pre_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_shape,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_shape,
+ const int a_tail_size, const int ind_num,
+ const int ind_ndim, const int* ind,
+ const int a_ndim, const int seg) {
+ index_t id = 0;
+ for (int dim = 0; dim < ind_ndim; ++dim) {
+ CHECK_LT(ind[dim * ind_num + i], a_shape[seg + dim])
+ << "IndexError: index " << ind[dim * ind_num + i]
+ << " is out of bounds for axis " << dim
+ << " with size " << a_shape[seg + dim];
+ CHECK_GE(ind[dim * ind_num + i], 0)
+ << "IndexError: index " << ind[dim * ind_num + i]
+ << " should be greater or equal to 0.";
+ id += a_pre_stride[seg + dim] * ind[dim * ind_num + i];
+ }
+ id *= a_tail_size;
+ for (int _i = 0; _i < a_tail_size; ++_i) {
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_tail_id = mxnet_op::unravel(_i, a_tail_shape);
+ mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_id;
+ for (int _j = 0; _j < seg; ++_j) {
+ val_id[_j] = 0;
+ }
+ for (int _j = seg; _j < seg + a_ndim; ++_j) {
+ val_id[_j] = (val_shape[_j] == 1) ? 0 : a_tail_id[_j];
+ }
+ val_id[seg + ind_ndim - 1] = (val_shape[seg + ind_ndim - 1] == 1) ? 0 : i;
+ index_t val_dest = mxnet_op::dot(val_id, val_stride);
+ #pragma omp critical
+ {
+ out[id + _i] += val[val_dest];
+ }
+ }
+ }
+};
+
+template<typename xpu, typename DType>
+void IndexAddForwardCalc(mshadow::Stream<xpu> *s,
+ const int ind_num, DType* out,
+ const DType* val,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_tail_shape,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_pre_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_stride,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> val_shape,
+ const mshadow::Shape<MXNET_SPECIAL_MAX_NDIM> a_shape,
+ const int a_tail_size,
+ const int ind_ndim, const int* ind,
+ const int a_ndim) {
+ using namespace mxnet_op;
+ using namespace mshadow;
+ int seg = MXNET_SPECIAL_MAX_NDIM - a_ndim;
+ Kernel<IndexAddForwardCPUKernel<DType>, xpu>::Launch(
+ s, ind_num, out, val,
+ a_tail_shape, a_pre_stride,
+ val_stride, val_shape, a_shape,
+ a_tail_size, ind_num,
+ ind_ndim, ind, a_ndim, seg);
+}
+
+
+
+NNVM_REGISTER_OP(_npx_index_add)
+.describe(R"code(This operators implements the "+=" mimic function.
+)code" ADD_FILELINE)
+.set_num_inputs(3)
+.set_num_outputs(1)
+.set_attr<nnvm::FListInputNames>("FListInputNames",
+ [](const NodeAttrs& attrs) {
+ return std::vector<std::string>{"a", "ind", "val"};
+ })
+.set_attr<mxnet::FInferShape>("FInferShape", IndexModifyOpShape)
+.set_attr<nnvm::FInferType>("FInferType", IndexModifyOpType)
+.set_attr<FCompute>("FCompute<cpu>", IndexAddOpForward<cpu>)
+.set_attr<FResourceRequest>("FResourceRequest",
+ [](const NodeAttrs& attrs) {
+ return std::vector<ResourceRequest>{ResourceRequest::kTempSpace};
+ })
+.set_attr<nnvm::FGradient>("FGradient",
+ [](const nnvm::ObjectPtr& n, const std::vector<nnvm::NodeEntry>& ograds) {
+ auto a_grad = MakeNode("_copy", n->attrs.name + "_backward_a",
+ {ograds[0]}, nullptr, &n);
+ auto idx_grad = MakeNode("zeros_like", n->attrs.name + "_backward_indices",
+ {n->inputs[1]}, nullptr, &n);
+ auto val_grad = MakeNode("_backward_index_add_val", n->attrs.name + "_backward_val",
+ {ograds[0], n->inputs[1]}, nullptr, &n);
+ std::vector<nnvm::NodeEntry> ret;
+ ret.emplace_back(a_grad);
+ ret.emplace_back(idx_grad);
+ ret.emplace_back(val_grad);
+ return ret;
Review comment:
@yzhliu I've suggested @JiangZhaoh to write the gradient op by reusing the forward op. This has multiple benefits:
1) We can do more optimization in the graph pass
2) It's easy to support higher-order gradient.
I think we may add the document about such usage since not many people are aware of that.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] JiangZhaoh removed a comment on pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
JiangZhaoh removed a comment on pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#issuecomment-633954573
run ci [unix-gpu, unix-cpu, centos-gpu]
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] sxjscience commented on a change in pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
sxjscience commented on a change in pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#discussion_r431595029
##########
File path: tests/python/unittest/test_numpy_op.py
##########
@@ -1323,6 +1323,154 @@ def hybrid_forward(self, F, a):
expected_grad[basic_index] = 1
assert same(a.grad.asnumpy(), expected_grad)
+@with_seed()
+@use_np
+def test_npx_index_add():
+ class TestIndexAdd(HybridBlock):
+ def __init__(self):
+ super(TestIndexAdd, self).__init__()
+
+ def hybrid_forward(self, F, a, ind, val):
+ return F.npx.index_add(a, ind, val)
+
+ def index_add_forward(a, ind, val, ind_ndim, ind_num):
+ if val.dtype != a.dtype:
+ val = val.astype(a.dtype)
+ ind_arr = ind.transpose()
+ if ind_arr.ndim == 0:
+ ind_arr = _np.array([ind_arr])
+ for i in range(ind_arr.shape[0]):
+ t_ind = ind_arr[i]
+ t_ind = tuple(t_ind.tolist()) if type(t_ind) is _np.ndarray else t_ind.tolist()
+ if val.ndim + ind_ndim > a.ndim:
+ t_val = val[tuple([0 if val.shape[0]==1 else i])]
+ if type(t_val) is _np.ndarray and t_val.shape[0] == 1:
+ a[t_ind] += _np.squeeze(t_val, axis=0)
+ else:
+ a[t_ind] += t_val
+ else:
+ a[t_ind] += val
+ return a
+
+ def index_add_bwd(out_grad, a_grad, ind, val_grad, ind_ndim, ind_num, grad_req_a, grad_req_val):
+ if grad_req_a == 'add':
+ init_a_grad = _np.array(a_grad)
+ if grad_req_val == 'add':
+ init_val_grad = _np.array(val_grad)
+ a_grad = _np.zeros(a_grad.shape).astype(a_grad.dtype)
+ val_grad = _np.zeros(val_grad.shape).astype(val_grad.dtype)
+
+ ind_arr = ind.transpose()
+ if ind_arr.ndim == 0:
+ ind_arr = _np.array([ind_arr])
+ for i in range(ind_arr.shape[0]):
+ t_ind = ind_arr[i]
+ t_ind = tuple(ind_arr[i].tolist()) if type(ind_arr[i]) is _np.ndarray else ind_arr[i].tolist()
+ a_grad[t_ind] = out_grad[t_ind]
+ if val_grad.ndim + ind_ndim > a_grad.ndim:
+ idx = 0 if val_grad.shape[0]==1 else i
+ t_grad = out_grad[t_ind]
+ t_grad_shape = _np.array(t_grad.shape)
+ val_grad_shape = _np.array(val_grad[idx].shape)
+ if type(val_grad[idx]) is not _np.ndarray:
+ t_grad = _np.sum(t_grad)
+ else:
+ is_not_equal = t_grad_shape - val_grad_shape
+ if _np.any(is_not_equal):
+ broadcast_dim = _np.nonzero(_np.where(is_not_equal, 1, 0))
+ t_grad = _np.sum(t_grad, axis=tuple(broadcast_dim[0].reshape(1, -1)[0]), keepdims=True)
+ val_grad[idx] += t_grad
+ else:
+ t_grad = out_grad[t_ind]
+ if type(val_grad) is not _np.ndarray or val_grad.shape == ():
+ t_grad = _np.sum(t_grad)
+ else:
+ if type(t_grad) is _np.ndarray:
+ ext_dim = t_grad.ndim() - val_grad.ndim()
+ if ext_dim:
+ t_grad = _np.sum(t_grad, axis=tuple(_np.arange(ext_dim)))
+ t_grad_shape = _np.array(t_grad.shape)
+ val_grad_shape = _np.array(val_grad.shape)
+ is_not_equal = t_grad_shape - val_grad_shape
+ if _np.any(is_not_equal):
+ broadcast_dim = _np.nonzero(_np.where(is_not_equal, 1, 0))
+ t_grad = _np.sum(t_grad, axis=tuple(broadcast_dim.reshape(1, -1)[0]), keepdims=True)
+ val_grad += t_grad
+ if grad_req_a == 'add':
+ a_grad += init_a_grad
+ if grad_req_val == 'add':
+ val_grad += init_val_grad
+ return a_grad, val_grad
+
+ # a.shape, ind.shape, val.shape, ind_ndim, ind_num
+ configs = [((2, ), np.array(1, dtype=_np.int32), (1, ), 1, 1)]
+ shape = tuple(_np.random.randint(1, 6, size=(4))) # a.shape
+ for ind_ndim in range(1, 5): # ind.shape: (ind_ndim, ind_num)
+ ind_num = _np.random.randint(1, 7)
+ ind = []
+ for ind_dim in range(ind_ndim):
+ ind.append(_np.random.randint(0, shape[ind_dim], size=(ind_num)))
+ ind = _np.array(ind).astype(_np.int32)
+ # case: val is scalar
+ configs.append(tuple([shape, ind, (), ind_ndim, ind_num]))
+ for val_ndim in range(1, 5 - ind_ndim):
+ val_shape = [1 if _np.random.randint(0, 5)==0 else ind_num]
+ for val_dim in range(ind_ndim, 4):
+ val_shape.append(1 if _np.random.randint(0, 5)==0 else shape[val_dim])
+ # case: val is tensor
+ configs.append(tuple([shape, ind, tuple(val_shape), ind_ndim, ind_num]))
+
+ dtypes = ['float32', 'float64', 'int32', 'int64']
+ grad_req = ['write', 'null']
Review comment:
Should we also test add?
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] JiangZhaoh commented on pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
JiangZhaoh commented on pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#issuecomment-633556688
@mxnet-bot run ci [unix-gpu, windows-cpu, windows-gpu]
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] JiangZhaoh commented on pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
JiangZhaoh commented on pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#issuecomment-633954573
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] sxjscience merged pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
sxjscience merged pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] mxnet-bot commented on pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
mxnet-bot commented on pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#issuecomment-634010813
Jenkins CI successfully triggered : [unix-cpu, unix-gpu, centos-gpu]
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] JiangZhaoh edited a comment on pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
JiangZhaoh edited a comment on pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#issuecomment-631412311
For `npx.index_add(a, ind, val)`, now it supports int32 or int64 tensor `ind`, and both int or float tensor `a` and `val`. But it requires the dtype of `a` and `val` keep same.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] sxjscience commented on pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
sxjscience commented on pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#issuecomment-637207570
Thanks @JiangZhaoh , this is merged.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] sxjscience commented on pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
sxjscience commented on pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#issuecomment-637002441
@mxnet-bot run ci [centos-gpu] [unix-cpu] [unix-gpu] [windows-cpu] [windows-gpu]
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] sxjscience commented on pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
sxjscience commented on pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#issuecomment-632499329
LGTM overall. Needs to fix the CudaMalloc + cudaFree.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] leezu commented on pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
leezu commented on pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#issuecomment-637006616
@mxnet-bot run ci [unix-cpu, unix-gpu, windows-cpu, windows-gpu]
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] ZheyuYe commented on pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
ZheyuYe commented on pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#issuecomment-628374562
It seems that this operation currently only supports floating type. In addition, it would be great if the data types of indices can be extended to `np.ndarray` instead of only `tuple` and `list` as
```python
import mxnet as mx
a = mx.np.zeros((2,6))
val = mx.np.arange(6)
x = [1,0,0,0,0,0]
y = [0,2,3,4,0,1]
# indices = [list, list]
mx.npx.index_add(a, val, [x, y])
# indices = (list, list)
mx.npx.index_add(a, val, (x, y))
# indices = [np.ndarray, np.ndarray]
x = mx.np.array(x).astype('int32')
y = mx.np.array(y).astype('int32')
mx.npx.index_add(a, val, [x, y])
>>> MXNetError: MXNetError: Invalid Parameter format for ind expect tuple of <Shape(tuple)>
but value='[array([1, 0, 0, 0, 0, 0], dtype=int32), array([0, 2, 3, 4, 0, 1], dtype=int32)]', in
operator _npx_index_add(name="", ind="[array([1, 0, 0, 0, 0, 0], dtype=int32), array([0, 2, 3, 4, 0, 1], dtype=int32)]")
```
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] mxnet-bot commented on pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
mxnet-bot commented on pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#issuecomment-637006669
Jenkins CI successfully triggered : [unix-cpu, windows-gpu, windows-cpu, unix-gpu]
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-mxnet] mxnet-bot commented on pull request #18089: Add npx op 'index_add'
Posted by GitBox <gi...@apache.org>.
mxnet-bot commented on pull request #18089:
URL: https://github.com/apache/incubator-mxnet/pull/18089#issuecomment-637002494
Jenkins CI successfully triggered : [centos-gpu]
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org