polishing after qingqing's comments

Author: wangyang59

paddle/fluid/operators/bilinear_interp_op.cc

@@ -44,10 +44,13 @@ class BilinearInterpOpMaker : public framework::OpProtoAndCheckerMaker {
   BilinearInterpOpMaker(OpProto* proto, OpAttrChecker* op_checker)
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X",
-             "The input tensor of bilinear interpolation, 4-D with NCHW shape");
-    AddOutput("Out", "The output tensor with the same shape as X");
-    AddAttr<int>("out_h", "output height of bilinear interpolation op.");
-    AddAttr<int>("out_w", "output weight of bilinear interpolation op.");
+             "(Tensor) The input tensor of bilinear interpolation, "
+             "This is a 4-D tensor with shape of (N x C x h x w)");
+    AddOutput("Out",
+              "(Tensor) The dimension of output is (N x C x out_h x out_w]");
+
+    AddAttr<int>("out_h", "(int) output height of bilinear interpolation op.");
+    AddAttr<int>("out_w", "(int) output width of bilinear interpolation op.");
     AddComment(R"DOC(
           Bilinear interpolation is an extension of linear interpolation for 
           interpolating functions of two variables (e.g. H-direction and 

paddle/fluid/operators/bilinear_interp_op.cu

@@ -9,15 +9,90 @@
    See the License for the specific language governing permissions and
    limitations under the License. */
 
-#include "paddle/fluid/framework/op_registry.h"
-#include "paddle/fluid/operators/bilinear_interp_op.cu.h"
-#include "paddle/fluid/operators/math/math_function.h"
+#include "paddle/fluid/operators/bilinear_interp_op.h"
+#include "paddle/fluid/platform/cuda_helper.h"
 
 namespace paddle {
 namespace operators {
 
 using framework::Tensor;
 
+template <typename T>
+__global__ void KeBilinearInterpFw(
+    const T* in, const size_t in_img_h, const size_t in_img_w,
+    const size_t input_h, const size_t input_w, T* out, const size_t out_img_h,
+    const size_t out_img_w, const size_t output_h, const size_t output_w,
+    const size_t num_channels, const T ratio_h, const T ratioW) {
+  int nthreads = output_h * output_w;
+  int tid = blockIdx.x * blockDim.x + threadIdx.x;
+  if (tid < nthreads) {
+    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 channel_id = out_id_w / out_img_size;
+
+    int out_img_idy = (out_id_w % out_img_size) / out_img_w;
+    int in_img_idy = ratio_h * out_img_idy;
+    int h_id = (in_img_idy < in_img_h - 1) ? 1 : 0;
+    T h1lambda = ratio_h * out_img_idy - in_img_idy;
+    T h2lambda = 1.f - h1lambda;
+
+    int out_img_idx = tid % out_img_w;
+    int in_img_idx = ratioW * out_img_idx;
+    int w_id = (in_img_idx < in_img_w - 1) ? 1 : 0;
+    T w1lambda = ratioW * out_img_idx - in_img_idx;
+    T w2lambda = 1.f - w1lambda;
+
+    const T* in_pos = &in[out_id_h * input_w + channel_id * in_img_size +
+                          in_img_idy * in_img_w + in_img_idx];
+
+    // bilinear interpolation
+    out[out_id_h * output_w + out_id_w] =
+        h2lambda * (w2lambda * in_pos[0] + w1lambda * in_pos[w_id]) +
+        h1lambda * (w2lambda * in_pos[h_id * in_img_w] +
+                    w1lambda * in_pos[h_id * in_img_w + w_id]);
+  }
+}
+
+template <typename T>
+__global__ void KeBilinearInterpBw(
+    T* in, const size_t in_img_h, const size_t in_img_w, const size_t input_h,
+    const size_t input_w, const T* out, const size_t out_img_h,
+    const size_t out_img_w, const size_t output_h, const size_t output_w,
+    const size_t num_channels, const T ratio_h, const T ratioW) {
+  int nthreads = output_h * output_w;
+  int tid = blockIdx.x * blockDim.x + threadIdx.x;
+  if (tid < nthreads) {
+    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 channel_id = out_id_w / out_img_size;
+
+    int out_img_idy = (out_id_w % out_img_size) / out_img_w;
+    int in_img_idy = ratio_h * out_img_idy;
+    int h_id = (in_img_idy < in_img_h - 1) ? 1 : 0;
+    T h1lambda = ratio_h * out_img_idy - in_img_idy;
+    T h2lambda = 1.f - h1lambda;
+
+    int out_img_idx = tid % out_img_w;
+    int in_img_idx = ratioW * out_img_idx;
+    int w_id = (in_img_idx < in_img_w - 1) ? 1 : 0;
+    T w1lambda = ratioW * out_img_idx - in_img_idx;
+    T w2lambda = 1.f - w1lambda;
+
+    T* in_pos = &in[out_id_h * input_w + channel_id * in_img_size +
+                    in_img_idy * in_img_w + in_img_idx];
+    const T* out_pos = &out[out_id_h * output_w + out_id_w];
+    atomicAdd(&in_pos[0], h2lambda * w2lambda * out_pos[0]);
+    atomicAdd(&in_pos[w_id], h2lambda * w1lambda * out_pos[0]);
+    atomicAdd(&in_pos[h_id * in_img_w], h1lambda * w2lambda * out_pos[0]);
+    atomicAdd(&in_pos[h_id * in_img_w + w_id],
+              h1lambda * w1lambda * out_pos[0]);
+  }
+}
+
 template <typename T>
 class BilinearInterpOpCUDAKernel : public framework::OpKernel<T> {
  public:

paddle/fluid/operators/bilinear_interp_op.cu.h

@@ -84,5 +84,19 @@ def init_test_case(self):
         self.out_w = 12
 
 
+class TestCase2(TestBilinearInterpOp):
+    def init_test_case(self):
+        self.input_shape = [16, 3, 512, 1024]
+        self.out_h = 128
+        self.out_w = 256
+
+
+class TestCase2(TestBilinearInterpOp):
+    def init_test_case(self):
+        self.input_shape = [8, 1, 256, 128]
+        self.out_h = 1024
+        self.out_w = 1024
+
+
 if __name__ == "__main__":
     unittest.main()