   config.thread_per_block = dim3(block_x, block_y, block_z);
   config.block_per_grid = dim3(grid_x, grid_y, grid_z);
   return config;
+}
+
+template <typename T>
+__forceinline__ __device__ void PreCalculatorForLinearInterpInputIndex(
+    int* in_img_idx, int* w_id, T* w1lambda, T* w2lambda, T src_w,
+    const int in_img_w) {
+  src_w = (src_w > 0) ? src_w : 0.f;
+  *in_img_idx = static_cast<int>(src_w);
+  *w_id = (*in_img_idx < in_img_w - 1) ? 1 : 0;
+  *w1lambda = src_w - *in_img_idx;
+  *w2lambda = 1.f - *w1lambda;
 }
 
 struct FastDivModForInterpolate {

    T w1lambda, h1lambda, w2lambda, h2lambda;
    T src_w = ratio_w * (out_img_idx + align_type_value) - align_type_value;
    T src_h = ratio_h * (out_img_idy + align_type_value) - align_type_value;
    PreCalculatorForInputIndex(&in_img_idx, &in_img_idy, &w_id, &h_id,
                               &w1lambda, &h1lambda, &w2lambda, &h2lambda,
                               src_w, src_h, in_w, in_h);

    // top_left_index is just input_index.
    int input_index = out_id_h * in_chw + channel_id * in_img_size +

      T w1lambda, h1lambda, w2lambda, h2lambda;
      T src_w = ratio_w * (out_img_idx + align_type_value) - align_type_value;
      T src_h = ratio_h * (out_img_idy + align_type_value) - align_type_value;

      PreCalculatorForLinearInterpInputIndex(&in_img_idx, &w_id, &w1lambda,
                                             &w2lambda, src_w, in_w);
      PreCalculatorForLinearInterpInputIndex(&in_img_idy, &h_id, &h1lambda,
                                             &h2lambda, src_h, in_h);

      T* in_pos = &in[out_id_h * in_chw + in_img_idy * in_w * num_channels +
                      in_img_idx * num_channels + channel_id];
