solin319 commented on a change in pull request #7393: add depthwise 
convolution's gpu version optimization
URL: https://github.com/apache/incubator-mxnet/pull/7393#discussion_r132133186
 
 

 ##########
 File path: src/operator/depthwise_convolution-inl.h
 ##########
 @@ -0,0 +1,346 @@
+/*
+ * 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 depthwise_convolution-inl.h
+ * \brief CUDA depthwise convolution code
+ * \author shuqian...@hobot.cc
+*/
+#ifndef MXNET_OPERATOR_DEPTHWISE_CONVOLUTION_INL_H_
+#define MXNET_OPERATOR_DEPTHWISE_CONVOLUTION_INL_H_
+#include <algorithm>
+#include <vector>
+#include "./convolution-inl.h"
+#include "../common/cuda_utils.h"
+
+#if MXNET_USE_CUDA
+#include <cub/cub.cuh>
+#include "./depthwise_convolution_tf.cuh"
+
+namespace mxnet {
+namespace op {
+using namespace tf::depthwise_conv;
+template<typename DType>
+class DepthwiseConvolutionOp : public Operator {
+ public:
+  explicit DepthwiseConvolutionOp(const ConvolutionParam& param,
+                                  const std::vector<TShape>& in_shape,
+                                  const std::vector<TShape>& out_shape) {
+    args_.batch = in_shape[conv::kData][0];
+    args_.in_channel = in_shape[conv::kData][1];
+    args_.in_height = in_shape[conv::kData][2];
+    args_.in_width = in_shape[conv::kData][3];
+    args_.filter_height = in_shape[conv::kWeight][2];
+    args_.filter_width = in_shape[conv::kWeight][3];
+    args_.stride_height = param.stride[0];
+    args_.stride_width = param.stride[1];
+    args_.pad_height = param.pad[0];
+    args_.pad_width = param.pad[1];
+    args_.out_channel = out_shape[conv::kOut][1];
+    args_.out_height = out_shape[conv::kOut][2];
+    args_.out_width = out_shape[conv::kOut][3];
+    bias_term_ = !param.no_bias;
+  }
+
+  ~DepthwiseConvolutionOp() {}
+
+  virtual void Forward(const OpContext &ctx,
+                       const std::vector<TBlob> &in_data,
+                       const std::vector<OpReqType> &req,
+                       const std::vector<TBlob> &out_data,
+                       const std::vector<TBlob> &aux_args);
+
+  virtual void Backward(const OpContext &ctx,
+                        const std::vector<TBlob> &out_grad,
+                        const std::vector<TBlob> &in_data,
+                        const std::vector<TBlob> &out_data,
+                        const std::vector<OpReqType> &req,
+                        const std::vector<TBlob> &in_grad,
+                        const std::vector<TBlob> &aux_args);
+
+ private:
+  DepthwiseArgs args_;
+  bool bias_term_;
+};  // class DepthwiseConvolutionOp
+
+namespace depthwise_conv {
+namespace cuda {
+template<typename DType, int kFilterWidth, int kFilterHeight>
+__global__ void __launch_bounds__(1024, 2)
+DepthwiseConv2dBackwardFilterKernel(const DepthwiseArgs args,
 
 Review comment:
   Does this logic faster than the logic in tensorflow?  There are total 4 
loops in this logic.
   Why there is no cross - border judgment during filter * input in this logic. 
Just like 
   ' if (in_r_start >= 0 && in_c_start >= 0 && in_r_end < in_rows && in_c_end < 
in_cols)'
 
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on GitHub and use the
URL above to go to the specific comment.
 
For queries about this service, please contact Infrastructure at:
us...@infra.apache.org


With regards,
Apache Git Services

Reply via email to