Optimize nearest_interp forward

paddle/fluid/operators/interpolate_v2_op.cu

@@ -16,13 +16,35 @@
 #include "paddle/fluid/platform/device/gpu/gpu_device_function.h"
 #include "paddle/fluid/platform/device/gpu/gpu_launch_config.h"
 #include "paddle/fluid/platform/device/gpu/gpu_primitives.h"
+#include "paddle/fluid/platform/fast_divmod.h"
 
 namespace paddle {
 namespace operators {
 
 using framework::Tensor;
 using DataLayout = framework::DataLayout;
 
+struct FastDivModForInterpolate {
+ public:
+  platform::FastDivMod channels_;
+  platform::FastDivMod output_w_;
+  platform::FastDivMod outimg_w_;
+  platform::FastDivMod out_size_;
+  platform::FastDivMod outimgw_mul_chann_;
+
+  explicit HOSTDEVICE FastDivModForInterpolate(const int channels,
+                                               const int output_w,
+                                               const int outimg_w,
+                                               const int out_size,
+                                               const int outimgw_mul_chann) {
+    channels_ = platform::FastDivMod(channels);
+    output_w_ = platform::FastDivMod(output_w);
+    outimg_w_ = platform::FastDivMod(outimg_w);
+    out_size_ = platform::FastDivMod(out_size);
+    outimgw_mul_chann_ = platform::FastDivMod(outimgw_mul_chann);
+  }
+};
+
 template <typename T>
 __global__ void KeNearestNeighborInterpFw(
     const T* in, const size_t in_img_h, const size_t in_img_w,
@@ -33,29 +55,37 @@ __global__ void KeNearestNeighborInterpFw(
   int nthreads = output_h * output_w;
   int tid = blockIdx.x * blockDim.x + threadIdx.x;
   int stride = blockDim.x * gridDim.x;
+  int in_img_size = in_img_h * in_img_w;
+  int out_img_size = out_img_h * out_img_w;
+
   for (; tid < nthreads; tid += stride) {
-    int out_id_h = tid / output_w;
-    int out_id_w = tid % output_w;
-    int in_img_size = input_w / num_channels;
-    int out_img_size = output_w / num_channels;
+    int outimgw_mul_channels = out_img_w * num_channels;
+    FastDivModForInterpolate divmods(num_channels, output_w, out_img_w,
+                                     out_img_size, outimgw_mul_channels);
+
+    auto out_id_divmod = divmods.output_w_.Divmod(tid);
+    int out_id_h = out_id_divmod.val[0];
+    int out_id_w = out_id_divmod.val[1];
 
     int channel_id, out_img_idy, out_img_idx;
     if (data_layout == DataLayout::kNCHW) {
-      channel_id = out_id_w / out_img_size;
-      out_img_idy = (out_id_w % out_img_size) / out_img_w;
-      out_img_idx = tid % out_img_w;
+      auto channel_divmod = divmods.out_size_.Divmod(out_id_w);
+      channel_id = channel_divmod.val[0];
+      out_img_idy = divmods.outimg_w_.Divmod(channel_divmod.val[1]).val[0];
+      out_img_idx = divmods.outimg_w_.Divmod(tid).val[1];
     } else {
-      out_img_idy = out_id_w / (out_img_w * num_channels);
-      out_img_idx = out_id_w % (out_img_w * num_channels) / num_channels;
-      channel_id = tid % num_channels;
+      channel_id = divmods.channels_.Divmod(tid).val[1];
+      auto outimg_id_divmod = divmods.outimgw_mul_chann_.Divmod(out_id_w);
+      out_img_idy = outimg_id_divmod.val[0];
+      out_img_idx = divmods.channels_.Divmod(outimg_id_divmod.val[1]).val[0];
     }
 
-    int in_img_idy = (align_corners)
-                         ? static_cast<int>(ratio_h * out_img_idy + 0.5)
-                         : static_cast<int>(ratio_h * out_img_idy);
-    int in_img_idx = (align_corners)
-                         ? static_cast<int>(ratio_w * out_img_idx + 0.5)
-                         : static_cast<int>(ratio_w * out_img_idx);
+    int in_img_idx = ratio_w * out_img_idx;
+    int in_img_idy = ratio_h * out_img_idy;
+    if (align_corners) {
+      in_img_idx += 0.5;
+      in_img_idy += 0.5;
+    }
 
     if (data_layout == DataLayout::kNCHW) {
       out[tid] = in[out_id_h * input_w + channel_id * in_img_size +

paddle/fluid/platform/device/gpu/gpu_launch_config.h

@@ -131,8 +131,6 @@ inline GpuLaunchConfig GetGpuLaunchConfig2D(
   return config;
 }
 
-// TODO(wangchaochaohu): 3D will add later
-
 }  // namespace platform
 }  // namespace paddle