     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);
   }
 };
 

    T1* output_data = context.template Alloc<T1>(output);
    T2* mask_data = context.template Alloc<T2>(mask);

    int nthreads = batch_size * output_channels * output_depth * output_height *
                   output_width;
    int thread_num = 1024;
#ifdef WITH_NV_JETSON
    backends::gpu::ChangeThreadNum(context, &thread_num);
#endif

    int blocks = (nthreads + thread_num - 1) / thread_num;
    dim3 threads(thread_num, 1);
    dim3 grid(blocks, 1);

    KernelMaxPool3DWithIdx<T1, T2>
        <<<grid, threads, 0, context.stream()>>>(nthreads,
                                                 input_data,
                                                 input_channels,
                                                 input_depth,

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

