```python
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline


_cpp_source = r"""
#include <torch/extension.h>

torch::Tensor depthwise_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias,
    int64_t stride_h, int64_t stride_w,
    int64_t padding_h, int64_t padding_w,
    int64_t dilation_h, int64_t dilation_w);
"""

_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/AccumulateType.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>

__global__ void depthwise_3x7_kernel(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int channels, int height, int width,
    int out_h, int out_w)
{
    const int ox = blockIdx.x * 128 + threadIdx.x;
    const int oy = (blockIdx.y * 2 + threadIdx.y) * 4;
    const int plane = blockIdx.z;

    if (ox >= out_w || oy >= out_h)
        return;

    const int channel = plane % channels;
    const float* w = weight + channel * 21;
    const float* input =
        x + static_cast<int64_t>(plane) * height * width
          + static_cast<int64_t>(oy) * width + ox;

    float a0 = 0.0f;
    float a1 = 0.0f;
    float a2 = 0.0f;
    float a3 = 0.0f;

    // Each thread computes four vertically adjacent outputs.
    // Six input rows cover all four convolution windows.
    #pragma unroll
    for (int k = 0; k < 7; ++k) {
        const float w0 = w[k];
        const float w1 = w[7 + k];
        const float w2 = w[14 + k];

        const float v0 = input[k];
        const float v1 = input[width + k];
        const float v2 = input[2 * width + k];
        const float v3 =
            oy + 3 < height ? input[3 * width + k] : 0.0f;
        const float v4 =
            oy + 4 < height ? input[4 * width + k] : 0.0f;
        const float v5 =
            oy + 5 < height ? input[5 * width + k] : 0.0f;

        a0 = fmaf(v0, w0, a0);
        a0 = fmaf(v1, w1, a0);
        a0 = fmaf(v2, w2, a0);

        a1 = fmaf(v1, w0, a1);
        a1 = fmaf(v2, w1, a1);
        a1 = fmaf(v3, w2, a1);

        a2 = fmaf(v2, w0, a2);
        a2 = fmaf(v3, w1, a2);
        a2 = fmaf(v4, w2, a2);

        a3 = fmaf(v3, w0, a3);
        a3 = fmaf(v4, w1, a3);
        a3 = fmaf(v5, w2, a3);
    }

    const float b = bias ? bias[channel] : 0.0f;
    const int64_t offset =
        static_cast<int64_t>(plane) * out_h * out_w
        + static_cast<int64_t>(oy) * out_w + ox;

    output[offset] = a0 + b;
    if (oy + 1 < out_h)
        output[offset + out_w] = a1 + b;
    if (oy + 2 < out_h)
        output[offset + 2 * out_w] = a2 + b;
    if (oy + 3 < out_h)
        output[offset + 3 * out_w] = a3 + b;
}

template <typename scalar_t>
__global__ void depthwise_generic_kernel(
    const scalar_t* __restrict__ x,
    const scalar_t* __restrict__ weight,
    const scalar_t* __restrict__ bias,
    scalar_t* __restrict__ output,
    int64_t total,
    int64_t channels, int64_t height, int64_t width,
    int64_t out_h, int64_t out_w,
    int64_t kernel_h, int64_t kernel_w,
    int64_t stride_h, int64_t stride_w,
    int64_t padding_h, int64_t padding_w,
    int64_t dilation_h, int64_t dilation_w)
{
    using acc_t = at::acc_type<scalar_t, true>;

    for (int64_t index =
             static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
         index < total;
         index += static_cast<int64_t>(blockDim.x) * gridDim.x)
    {
        const int64_t ox = index % out_w;
        const int64_t oy = (index / out_w) % out_h;
        const int64_t plane = index / (out_h * out_w);
        const int64_t channel = plane % channels;

        const int64_t iy0 = oy * stride_h - padding_h;
        const int64_t ix0 = ox * stride_w - padding_w;
        const int64_t input_base = plane * height * width;
        const int64_t weight_base = channel * kernel_h * kernel_w;

        acc_t sum = acc_t(0);
        for (int64_t ky = 0; ky < kernel_h; ++ky) {
            const int64_t iy = iy0 + ky * dilation_h;
            if (iy < 0 || iy >= height)
                continue;

            for (int64_t kx = 0; kx < kernel_w; ++kx) {
                const int64_t ix = ix0 + kx * dilation_w;
                if (ix >= 0 && ix < width) {
                    sum += static_cast<acc_t>(
                               x[input_base + iy * width + ix])
                         * static_cast<acc_t>(
                               weight[weight_base + ky * kernel_w + kx]);
                }
            }
        }

        if (bias)
            sum += static_cast<acc_t>(bias[channel]);

        output[index] = static_cast<scalar_t>(sum);
    }
}

torch::Tensor depthwise_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias,
    int64_t stride_h, int64_t stride_w,
    int64_t padding_h, int64_t padding_w,
    int64_t dilation_h, int64_t dilation_w)
{
    TORCH_CHECK(x.is_cuda() && weight.is_cuda(),
                "Input and weight must be CUDA tensors");
    TORCH_CHECK(x.dim() == 4 && weight.dim() == 4,
                "Expected four-dimensional input and weight");
    TORCH_CHECK(x.is_contiguous() && weight.is_contiguous(),
                "Input and weight must be contiguous");
    TORCH_CHECK(x.device() == weight.device() &&
                x.scalar_type() == weight.scalar_type(),
                "Input and weight must have matching devices and dtypes");
    TORCH_CHECK(weight.size(0) == x.size(1) && weight.size(1) == 1,
                "Expected depthwise convolution weights");
    TORCH_CHECK(stride_h > 0 && stride_w > 0 &&
                dilation_h > 0 && dilation_w > 0 &&
                padding_h >= 0 && padding_w >= 0,
                "Invalid convolution parameters");

    const c10::cuda::CUDAGuard guard(x.device());

    const int64_t n = x.size(0);
    const int64_t c = x.size(1);
    const int64_t h = x.size(2);
    const int64_t w = x.size(3);
    const int64_t kh = weight.size(2);
    const int64_t kw = weight.size(3);
    const int64_t effective_h = dilation_h * (kh - 1) + 1;
    const int64_t effective_w = dilation_w * (kw - 1) + 1;

    TORCH_CHECK(kh > 0 && kw > 0 && c > 0 &&
                h + 2 * padding_h >= effective_h &&
                w + 2 * padding_w >= effective_w,
                "Kernel is larger than padded input");

    const int64_t oh = (h + 2 * padding_h - effective_h) / stride_h + 1;
    const int64_t ow = (w + 2 * padding_w - effective_w) / stride_w + 1;

    const bool has_bias = bias.has_value() && bias->defined();
    if (has_bias) {
        TORCH_CHECK(bias->device() == x.device() &&
                    bias->scalar_type() == x.scalar_type() &&
                    bias->is_contiguous() &&
                    bias->dim() == 1 && bias->numel() == c,
                    "Invalid bias");
    }

    auto output = torch::empty({n, c, oh, ow}, x.options());
    if (output.numel() == 0)
        return output;

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if (x.scalar_type() == at::kFloat &&
        kh == 3 && kw == 7 &&
        stride_h == 1 && stride_w == 1 &&
        padding_h == 0 && padding_w == 0 &&
        dilation_h == 1 && dilation_w == 1 &&
        n * c <= 65535 &&
        h <= 2147483647 && w <= 2147483647 &&
        (oh + 7) / 8 <= 65535)
    {
        const dim3 threads(128, 2);
        const dim3 blocks((ow + 127) / 128, (oh + 7) / 8, n * c);

        depthwise_3x7_kernel<<<blocks, threads, 0, stream>>>(
            x.data_ptr<float>(),
            weight.data_ptr<float>(),
            has_bias ? bias->data_ptr<float>() : nullptr,
            output.data_ptr<float>(),
            static_cast<int>(c), static_cast<int>(h), static_cast<int>(w),
            static_cast<int>(oh), static_cast<int>(ow));
    } else {
        const int64_t total = output.numel();
        const int blocks = static_cast<int>(
            std::min<int64_t>((total + 255) / 256, 65535));

        AT_DISPATCH_FLOATING_TYPES_AND2(
            at::ScalarType::Half, at::ScalarType::BFloat16,
            x.scalar_type(), "depthwise_generic_cuda", [&] {
                depthwise_generic_kernel<scalar_t>
                    <<<blocks, 256, 0, stream>>>(
                        x.data_ptr<scalar_t>(),
                        weight.data_ptr<scalar_t>(),
                        has_bias ? bias->data_ptr<scalar_t>() : nullptr,
                        output.data_ptr<scalar_t>(),
                        total, c, h, w, oh, ow, kh, kw,
                        stride_h, stride_w, padding_h, padding_w,
                        dilation_h, dilation_w);
            });
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_depthwise_extension = load_inline(
    name="depthwise_asymmetric_3x7_v1",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["depthwise_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class ModelNew(nn.Module):
    def __init__(
        self,
        in_channels: int,
        out_channels: int,
        kernel_size_h: int,
        kernel_size_w: int,
        stride_h: int = 1,
        stride_w: int = 1,
        padding_h: int = 0,
        padding_w: int = 0,
        dilation_h: int = 1,
        dilation_w: int = 1,
        groups: int = 1,
        bias: bool = False,
    ):
        super().__init__()
        self.conv2d = nn.Conv2d(
            in_channels,
            in_channels,
            (kernel_size_h, kernel_size_w),
            stride=(stride_h, stride_w),
            padding=(padding_h, padding_w),
            dilation=(dilation_h, dilation_w),
            groups=in_channels,
            bias=bias,
        )

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        if not x.is_cuda:
            return self.conv2d(x)

        conv = self.conv2d
        return _depthwise_extension.depthwise_cuda(
            x.contiguous(),
            conv.weight.contiguous(),
            conv.bias,
            conv.stride[0],
            conv.stride[1],
            conv.padding[0],
            conv.padding[1],
            conv.dilation[0],
            conv.dilation[1],
        )
```