[GitHub] [incubator-mxnet] Kh4L commented on a change in pull request #16268: Add mask target generator operator for Mask-RCNN

[GitBox <gi...@apache.org>]

Kh4L commented on a change in pull request #16268: Add mask target generator operator for Mask-RCNN
URL: https://github.com/apache/incubator-mxnet/pull/16268#discussion_r328631233
 
 

 ##########
 File path: src/operator/contrib/mrcnn_target.cu
 ##########
 @@ -0,0 +1,277 @@
+/*
+ * 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 mrcnn_target.cu
+ * \brief Mask-RCNN target generator
+ * \author Serge Panev
+ */
+
+#include "./mrcnn_target-inl.h"
+
+namespace mxnet {
+namespace op {
+
+using namespace mshadow::cuda;
+
+// The maximum number of blocks to use in the default kernel call.
+constexpr int MAXIMUM_NUM_BLOCKS = 4096;
+
+inline int CUDA_GET_BLOCKS(const int N) {
+  return std::max(
+      std::min(
+          (N + kMaxThreadsPerBlock - 1) / kMaxThreadsPerBlock,
+          MAXIMUM_NUM_BLOCKS),
+      // Use at least 1 block, since CUDA does not allow empty block.
+      1);
+}
+
+// Kernels
+
+template <typename T>
+__device__ T bilinear_interpolate(
+    const T* in_data,
+    const int height,
+    const int width,
+    T y,
+    T x,
+    const int index /* index for debug only*/) {
+  // deal with cases that inverse elements are out of feature map boundary
+  if (y < -1.0 || y > height || x < -1.0 || x > width) {
+    // empty
+    return 0;
+  }
+
+  if (y <= 0) {
+    y = 0;
+  }
+  if (x <= 0) {
+    x = 0;
+  }
+
+  int y_low = static_cast<int>(y);
+  int x_low = static_cast<int>(x);
+  int y_high;
+  int x_high;
+
+  if (y_low >= height - 1) {
+    y_high = y_low = height - 1;
+    y = (T)y_low;
+  } else {
+    y_high = y_low + 1;
+  }
+
+  if (x_low >= width - 1) {
+    x_high = x_low = width - 1;
+    x = (T)x_low;
+  } else {
+    x_high = x_low + 1;
+  }
+
+  T ly = y - y_low;
+  T lx = x - x_low;
+  T hy = 1. - ly, hx = 1. - lx;
+  // do bilinear interpolation
+  T v1 = in_data[y_low * width + x_low];
+  T v2 = in_data[y_low * width + x_high];
+  T v3 = in_data[y_high * width + x_low];
+  T v4 = in_data[y_high * width + x_high];
+  T w1 = hy * hx, w2 = hy * lx, w3 = ly * hx, w4 = ly * lx;
+
+  T val = (w1 * v1 + w2 * v2 + w3 * v3 + w4 * v4);
+
+  return val;
+}
+
+// Modified version of RoIAlignForwardKernel from Caffe (in roi_align.cu)
+// Main modifications:
+// - We don't need position_sensitive neither spatial_scale from the original RoIAlign kernel.
+// - We replace `channels` by `num_classes` and modify the logic consequently (e.g. offset_in_data
+//   does not use `c` anymore).
+template <typename T>
+__device__ void RoIAlignForward(
+    const T* in_data,  // (B, M, H, W)
+    const T* rois,  // (B, N, 4)
+    const T* matches,  // (B, N)
+    const int num_el,
+    const int num_classes,
+    const int height,
+    const int width,
+    const int pooled_height,
+    const int pooled_width,
+    const int sampling_ratio,
+    const int num_rois,
+    const int num_gtmasks,
+    T* out_data) {  // (B, N, C, H, W)
+  // Update kernel
+  for (size_t index = blockIdx.x * blockDim.x + threadIdx.x;
+       index < num_el;
+       index += blockDim.x * gridDim.x) {
+    // (n, c, ph, pw) is an element in the pooled output
+    int pw = index % pooled_width;
+    int ph = (index / pooled_width) % pooled_height;
+    // int c = (index / pooled_width / pooled_height) % num_classes;
+    int n = (index / pooled_width / pooled_height / num_classes) % num_rois;
+    int batch_idx = (index / pooled_width / pooled_height / num_classes / num_rois);
+
+    int roi_batch_ind = matches[batch_idx * num_rois + n];
+
+    const T* offset_rois = rois + batch_idx * (4 * num_rois) + n * 4;
+    // Do not using rounding; this implementation detail is critical
+    T roi_start_w = offset_rois[0];
+    T roi_start_h = offset_rois[1];
+    T roi_end_w = offset_rois[2];
+    T roi_end_h = offset_rois[3];
+
+    // Force malformed ROIs to be 1x1
+    T roi_width = max(roi_end_w - roi_start_w, (T)1.);
+    T roi_height = max(roi_end_h - roi_start_h, (T)1.);
+    T bin_size_h = static_cast<T>(roi_height) / static_cast<T>(pooled_height);
+    T bin_size_w = static_cast<T>(roi_width) / static_cast<T>(pooled_width);
+
+    const T* offset_in_data =
+        in_data + batch_idx * num_gtmasks * height * width
+        + roi_batch_ind * height * width;
+
+    // We use roi_bin_grid to sample the grid and mimic integral
+    int roi_bin_grid_h = (sampling_ratio > 0)
+        ? sampling_ratio
+        : ceil(roi_height / pooled_height);  // e.g., = 2
+    int roi_bin_grid_w =
+        (sampling_ratio > 0) ? sampling_ratio : ceil(roi_width / pooled_width);
+
+    // We do average (integral) pooling inside a bin
+    const T count = roi_bin_grid_h * roi_bin_grid_w;  // e.g. = 4
+
+    T output_val = 0.;
+    for (int iy = 0; iy < roi_bin_grid_h; iy++) {  // e.g., iy = 0, 1
+      const T y = roi_start_h + ph * bin_size_h +
+          static_cast<T>(iy + .5f) * bin_size_h /
+              static_cast<T>(roi_bin_grid_h);  // e.g., 0.5, 1.5
+      for (int ix = 0; ix < roi_bin_grid_w; ix++) {
+        const T x = roi_start_w + pw * bin_size_w +
+            static_cast<T>(ix + .5f) * bin_size_w /
+                static_cast<T>(roi_bin_grid_w);
+
+        T val = bilinear_interpolate(
+            offset_in_data, height, width, y, x, index);
+        output_val += val;
+      }
+    }
+    output_val /= count;
+
+    out_data[index] = output_val;
+  }
+}
+
+
+template<typename DType>
+__global__ void MRCNNTargetKernel(const DType *rois,
+                                  const DType *gt_masks,
+                                  const DType *matches,
+                                  const DType *cls_targets,
+                                  DType* sampled_masks,
+                                  DType* mask_cls,
+                                  const int total_out_el,
+                                  int batch_size,
+                                  int num_classes,
+                                  int num_rois,
+                                  int num_gtmasks,
+                                  int gt_height,
+                                  int gt_width,
+                                  int m) {
+  // computing sampled_masks
+  const int sampling_ratio = 2;
+  RoIAlignForward(gt_masks, rois, matches, total_out_el,
+                  num_classes, gt_height, gt_width, m, m,
+                  sampling_ratio, num_rois, num_gtmasks, sampled_masks);
+  // computing mask_cls
+  int num_masks = batch_size * num_rois * num_classes;
+  int mask_vol = m * m;
+  for (int mask_idx = blockIdx.x; mask_idx < num_masks; mask_idx += gridDim.x) {
+    int cls_idx = mask_idx % num_classes;
+    int roi_idx = (mask_idx / num_classes) % num_rois;
+    int batch_idx = (mask_idx / num_classes / num_rois);
+
+    DType* mask_cls_out = mask_cls + mask_idx * mask_vol;
+
+    DType cls_target = cls_targets[batch_idx * num_rois + roi_idx];
+    for (int mask_pixel = threadIdx.x; mask_pixel < mask_vol; mask_pixel += blockDim.x) {
+      mask_cls_out[mask_pixel] = (cls_target == cls_idx);
 
 Review comment:
   Are you talking about the implicit casting?
   I think it is ok, but I can also change to
   
   ```suggestion
         mask_cls_out[mask_pixel] = static_cast<DType>(static_cast<int>(cls_target) == cls_idx);
   ```

----------------------------------------------------------------
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