     dim3 grid(input_channels, batch_size, 1);
     int filter_multiplier = output_channels / input_channels;
 
-#define check_case(c_filter_multiplier, c_stride)                       \
+#define check_case(c_filter_multiplier, c_stride, c_filter)             \
   if (c_filter_multiplier == 0 ||                                       \
       filter_multiplier == c_filter_multiplier &&                       \
-          stride_height == stride_width && stride_height == c_stride) { \
+          stride_height == stride_width && stride_height == c_stride && \
+          (ksize_height == ksize_width && ksize_height == c_filter ||   \
+           c_filter == -1)) {                                           \
     KernelDepthwiseConvInputGradSp<                                     \
-        T, c_filter_multiplier,                                         \
-        c_stride><<<grid, threads, 0, context.stream()>>>(              \
+        T, c_filter_multiplier, c_stride,                               \
+        c_filter><<<grid, threads, 0, context.stream()>>>(              \
         output_grad_data, filter_data, batch_size, output_channels,     \
         output_height, output_width, input_channels, input_height,      \
         input_width, filter_multiplier, ksize_height, ksize_width,      \

  }
}

template <typename T, int c_filter_multiplier, int c_stride>
__global__ void KernelDepthwiseConvSp(
    const T* const input_data, const T* const filter_data, const int batch_size,
    const int output_channels, const int output_height, const int output_width,
    const int input_channels, const int input_height, const int input_width,
    const int filter_multiplier, const int filter_height,
    const int filter_width, const int stride_height, const int stride_width,
    const int padding_height, const int padding_width, const int dilate_height,
    const int dilate_width, T* const output_data) {
  if (c_filter_multiplier == 0)
    KernelDepthwiseConv<T>(input_data, filter_data, batch_size, output_channels,
                           output_height, output_width, input_channels,
                           input_height, input_width, filter_multiplier,
                           filter_height, filter_width, stride_height,
                           stride_width, padding_height, padding_width,
                           dilate_height, dilate_width, output_data);

  else
    KernelDepthwiseConv<T>(input_data, filter_data, batch_size, output_channels,
                           output_height, output_width, input_channels,
                           input_height, input_width, c_filter_multiplier,
                           filter_height, filter_height, c_stride, c_stride,
                           padding_height, padding_width, dilate_height,
                           dilate_width, output_data);
}

// CUDA kernel to compute the depthwise convolution backprop w.r.t input.
template <typename T>
__device__ __inline__ void KernelDepthwiseConvInputGrad(
    const T* const output_grad_data, const T* const filter_data,
    const int batch_size, const int output_channels, const int output_height,
    const int output_width, const int input_channels, const int input_height,
    const int input_width, const int filter_multiplier, const int filter_height,
    const int filter_width, const int stride_height, const int stride_width,
    const int padding_height, const int padding_width, const int dilate_height,
    const int dilate_width, T* const input_grad_data) {
  for (int w_in = threadIdx.x; w_in < input_width; w_in += blockDim.x) {
    for (int h_in = threadIdx.y; h_in < input_height; h_in += blockDim.y) {
      const int batch = blockIdx.y;

    dim3 threads(std::min(output_width, thread), blocks, 1);
    dim3 grid(output_channels, batch_size, 1);
    int filter_multiplier = output_channels / input_channels;
#define check_case(c_filter_multiplier, c_stride, c_filter)                  \
  if (c_filter_multiplier == 0 ||                                            \
      filter_multiplier == c_filter_multiplier &&                            \
          stride_height == stride_width && stride_height == c_stride &&      \
          (ksize_height == ksize_width && ksize_height == c_filter ||        \
           c_filter == -1)) {                                                \
    KernelDepthwiseConvSp<T, c_filter_multiplier, c_stride,                  \
                          c_filter><<<grid, threads, 0, context.stream()>>>( \
        input_data, filter_data, batch_size, output_channels, output_height, \
        output_width, input_channels, input_height, input_width,             \
        filter_multiplier, ksize_height, ksize_width, stride_height,         \
