 template class MaxPool2dWithIndexGradFunctor<phi::GPUContext, double, int>;
 
 template <typename T1, typename T2>
-__global__ void KernelMaxPool3DWithIdx(const int nthreads,
+__global__ void KernelMaxPool3DWithIdx(const int ncd,
                                        const T1* input_data,
                                        const int channels,
                                        const int input_depth,

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


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

    int ncd = batch_size * input_channels * output_depth;

    int thread_x = 32;
    int thread_y = 8;
    int thread_z = 1;
    dim3 threads(thread_x, thread_y, thread_z);
    std::array<int, 3> max_grid_dim = context.GetCUDAMaxGridDimSize();
    int block_x = (output_width + threads.x - 1) / threads.x;
    int block_y = (output_height + threads.y - 1) / threads.y;
    int block_z = (ncd > max_grid_dim[2] * threads.z)
                      ? max_grid_dim[2]
                      : (ncd + threads.z - 1) / threads.z;
    dim3 grid(block_x, block_y, block_z);

    auto pool_divmods_output = FastDivModForPooling3D(
        input_channels, output_width, output_height, output_depth);

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