                                        const int padding_width,
                                        bool adaptive,
                                        T1* output_data,
-                                       T2* mask_data) {
-  for (int index = blockIdx.x * blockDim.x + threadIdx.x; index < nthreads;
-       index += blockDim.x * gridDim.x) {
-    int pw = index % output_width;
-    int ph = (index / output_width) % output_height;
-    int pd = (index / output_width / output_height) % output_depth;
-    int c = (index / output_width / output_height / output_depth) % channels;
-    int batch_idx =
-        index / output_width / output_height / output_depth / channels;
-
-    int dstart, dend;
-    int hstart, hend;
-    int wstart, wend;
-    if (adaptive) {
-      dstart = AdaptStartIndex(pd, input_depth, output_depth);
-      dend = AdaptEndIndex(pd, input_depth, output_depth);
-
-      hstart = AdaptStartIndex(ph, input_height, output_height);
-      hend = AdaptEndIndex(ph, input_height, output_height);
-
-      wstart = AdaptStartIndex(pw, input_width, output_width);
-      wend = AdaptEndIndex(pw, input_width, output_width);
-    } else {
-      dstart = pd * stride_depth - padding_depth;
-      hstart = ph * stride_height - padding_height;
-      wstart = pw * stride_width - padding_width;
-      dend = min(dstart + ksize_depth, input_depth);
-      hend = min(hstart + ksize_height, input_height);
-      wend = min(wstart + ksize_width, input_width);
-      dstart = max(dstart, 0);
-      hstart = max(hstart, 0);
-      wstart = max(wstart, 0);
-    }
-
-    T1 ele = -FLT_MAX;
-    int max_index = -1;
-    input_data +=
-        (batch_idx * channels + c) * input_depth * input_height * input_width;
-
-    for (int d = dstart; d < dend; ++d) {
-      for (int h = hstart; h < hend; ++h) {
-        for (int w = wstart; w < wend; ++w) {
-          if (ele < input_data[(d * input_height + h) * input_width + w]) {
-            max_index = (d * input_height + h) * input_width + w;
-            ele = input_data[max_index];
+                                       T2* mask_data,
+                                       FastDivModForPooling3D divmods_output) {
+  int w_offset, h_offset, d_offset, nc_offset;
+  int dstart, dend, hstart, hend, wstart, wend;
+  const T1* input_data_cur;
+
+  w_offset = blockIdx.x * blockDim.x + threadIdx.x;
+  h_offset = blockIdx.y * blockDim.y + threadIdx.y;
+
+  if (w_offset < output_width && h_offset < output_height) {
+    for (int index_z = blockIdx.z * blockDim.z + threadIdx.z; index_z < ncd;
+         index_z += gridDim.z * blockDim.z) {
+      auto output_depth_divmod = divmods_output.depth.Divmod(index_z);
+      d_offset = output_depth_divmod.val[1];
+      nc_offset = output_depth_divmod.val[0];
+      int output_index =
+          nc_offset * output_depth * output_height * output_width +
+          d_offset * output_height * output_width + h_offset * output_width +
+          w_offset;
+      int input_offset = nc_offset * input_depth * input_height * input_width;
+      input_data_cur = input_data + input_offset;
+
+      if (adaptive) {
+        dstart = AdaptStartIndex(d_offset, input_depth, output_depth);
+        dend = AdaptEndIndex(d_offset, input_depth, output_depth);
+
+        hstart = AdaptStartIndex(h_offset, input_height, output_height);
+        hend = AdaptEndIndex(h_offset, input_height, output_height);
+
+        wstart = AdaptStartIndex(w_offset, input_width, output_width);
+        wend = AdaptEndIndex(w_offset, input_width, output_width);
+      } else {
+        dstart = d_offset * stride_depth - padding_depth;
+        hstart = h_offset * stride_height - padding_height;
+        wstart = w_offset * stride_width - padding_width;
+        dend = min(dstart + ksize_depth, input_depth);
+        hend = min(hstart + ksize_height, input_height);
+        wend = min(wstart + ksize_width, input_width);
+        dstart = max(dstart, 0);
+        hstart = max(hstart, 0);
+        wstart = max(wstart, 0);
+      }
+
+      T1 ele = -FLT_MAX;
+      int max_index = -1;
+      for (int d = dstart; d < dend; ++d) {
+        for (int h = hstart; h < hend; ++h) {
+          for (int w = wstart; w < wend; ++w) {
+            if (ele <
+                input_data_cur[(d * input_height + h) * input_width + w]) {
+              max_index = (d * input_height + h) * input_width + w;
+              ele = input_data_cur[max_index];
+            }
           }
         }
       }
-    }
-    output_data[index] = ele;
-    mask_data[index] = max_index;
+      output_data[output_index] = ele;
+      mask_data[output_index] = max_index;
+    }
   }
 }
 

                                                 padding_width,
                                                 adaptive,
                                                 output_data,
                                                 mask_data);
  }
};


    channel = paddle::platform::FastDivMod(channels);
    width = paddle::platform::FastDivMod(output_width);
    height = paddle::platform::FastDivMod(output_height);
  }
};

struct FastDivModForPooling3D {
 public:
  paddle::platform::FastDivMod channel;
  paddle::platform::FastDivMod width;
  paddle::platform::FastDivMod height;
  paddle::platform::FastDivMod depth;

  explicit HOSTDEVICE FastDivModForPooling3D(const int channels,
                                             const int output_width,
                                             const int output_height,
                                             const int output_depth) {
    channel = paddle::platform::FastDivMod(channels);
    width = paddle::platform::FastDivMod(output_width);
    height = paddle::platform::FastDivMod(output_height);
    depth = paddle::platform::FastDivMod(output_depth);
  }
};

