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


_cuda_source = r"""
#include <torch/extension.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_height, int out_width)
{
    constexpr int TILE_W = 128;
    constexpr int TILE_H = 8;
    constexpr int SHARED_W = TILE_W + 6;
    constexpr int SHARED_H = TILE_H + 2;

    __shared__ float tile[SHARED_H * SHARED_W];
    __shared__ float weights[21];

    const int tid = threadIdx.x;
    const int plane = blockIdx.z;
    const int channel = plane % channels;
    const int x0 = blockIdx.x * TILE_W;
    const int y0 = blockIdx.y * TILE_H;

    const float* input_plane =
        x + static_cast<int64_t>(plane) * height * width;

    for (int i = tid; i < SHARED_H * SHARED_W; i += TILE_W) {
        const int iy = y0 + i / SHARED_W;
        const int ix = x0 + i % SHARED_W;
        tile[i] = (iy < height && ix < width)
            ? input_plane[static_cast<int64_t>(iy) * width + ix]
            : 0.0f;
    }

    if (tid < 21)
        weights[tid] = weight[channel * 21 + tid];

    __syncthreads();

    float sums[TILE_H];

    #pragma unroll
    for (int r = 0; r < TILE_H; ++r)
        sums[r] = 0.0f;

    // Each shared-memory input value contributes to up to three output rows.
    #pragma unroll
    for (int iy = 0; iy < SHARED_H; ++iy) {
        #pragma unroll
        for (int kx = 0; kx < 7; ++kx) {
            const float value = tile[iy * SHARED_W + tid + kx];

            if (iy < TILE_H)
                sums[iy] = fmaf(value, weights[kx], sums[iy]);

            if (iy >= 1 && iy < TILE_H + 1)
                sums[iy - 1] =
                    fmaf(value, weights[7 + kx], sums[iy - 1]);

            if (iy >= 2)
                sums[iy - 2] =
                    fmaf(value, weights[14 + kx], sums[iy - 2]);
        }
    }

    const int ox = x0 + tid;
    const float b = bias ? bias[channel] : 0.0f;

    if (ox < out_width) {
        float* output_plane =
            output + static_cast<int64_t>(plane) * out_height * out_width;

        #pragma unroll
        for (int r = 0; r < TILE_H; ++r) {
            if (y0 + r < out_height)
                output_plane[static_cast<int64_t>(y0 + r) * out_width + ox]
                    = sums[r] + b;
        }
    }
}

__global__ void depthwise_generic_kernel(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int64_t total,
    int channels, int height, int width,
    int out_height, int out_width,
    int kh, int kw, int sh, int sw,
    int ph, int pw, int dh, int dw)
{
    const int64_t spatial = static_cast<int64_t>(out_height) * out_width;

    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 int ox = index % out_width;
        const int oy = (index / out_width) % out_height;
        const int64_t plane = index / spatial;
        const int channel = plane % channels;

        const float* input_plane =
            x + plane * static_cast<int64_t>(height) * width;
        const float* kernel =
            weight + static_cast<int64_t>(channel) * kh * kw;

        float sum = 0.0f;

        for (int ky = 0; ky < kh; ++ky) {
            const int iy = oy * sh - ph + ky * dh;
            if (iy < 0 || iy >= height)
                continue;

            for (int kx = 0; kx < kw; ++kx) {
                const int ix = ox * sw - pw + kx * dw;
                if (ix >= 0 && ix < width)
                    sum = fmaf(
                        input_plane[static_cast<int64_t>(iy) * width + ix],
                        kernel[ky * kw + kx],
                        sum);
            }
        }

        output[index] = sum + (bias ? bias[channel] : 0.0f);
    }
}

torch::Tensor depthwise_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    torch::Tensor bias,
    int64_t sh, int64_t sw,
    int64_t ph, int64_t pw,
    int64_t dh, int64_t dw)
{
    TORCH_CHECK(x.is_cuda() && weight.is_cuda(), "Expected CUDA tensors");
    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat,
                "Expected float32 tensors");
    TORCH_CHECK(x.dim() == 4 && weight.dim() == 4,
                "Expected four-dimensional tensors");
    TORCH_CHECK(x.is_contiguous() && weight.is_contiguous(),
                "Expected contiguous tensors");
    TORCH_CHECK(x.device() == weight.device(), "Device mismatch");
    TORCH_CHECK(weight.size(0) == x.size(1) && weight.size(1) == 1,
                "Invalid depthwise convolution weight");
    TORCH_CHECK(sh > 0 && sw > 0 && dh > 0 && dw > 0 &&
                ph >= 0 && pw >= 0, "Invalid convolution parameters");

    const c10::cuda::CUDAGuard device_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);

    TORCH_CHECK(h + 2 * ph >= dh * (kh - 1) + 1 &&
                w + 2 * pw >= dw * (kw - 1) + 1,
                "Kernel exceeds padded input dimensions");

    const int64_t oh = (h + 2 * ph - dh * (kh - 1) - 1) / sh + 1;
    const int64_t ow = (w + 2 * pw - dw * (kw - 1) - 1) / sw + 1;

    const float* bias_ptr = nullptr;
    if (bias.numel() != 0) {
        TORCH_CHECK(bias.device() == x.device() &&
                    bias.scalar_type() == at::kFloat &&
                    bias.is_contiguous() && bias.numel() == c,
                    "Invalid bias");
        bias_ptr = bias.data_ptr<float>();
    }

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

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

    if (kh == 3 && kw == 7 &&
        sh == 1 && sw == 1 &&
        ph == 0 && pw == 0 &&
        dh == 1 && dw == 1 &&
        n * c <= 65535 && (oh + 7) / 8 <= 65535)
    {
        dim3 grid(
            static_cast<unsigned>((ow + 127) / 128),
            static_cast<unsigned>((oh + 7) / 8),
            static_cast<unsigned>(n * c));

        depthwise_3x7_kernel<<<grid, 128, 0, stream>>>(
            x.data_ptr<float>(), weight.data_ptr<float>(),
            bias_ptr, output.data_ptr<float>(),
            c, h, w, oh, ow);
    } else {
        const int64_t needed = (output.numel() + 255) / 256;
        const int blocks = static_cast<int>(needed < 65535 ? needed : 65535);

        depthwise_generic_kernel<<<blocks, 256, 0, stream>>>(
            x.data_ptr<float>(), weight.data_ptr<float>(),
            bias_ptr, output.data_ptr<float>(), output.numel(),
            c, h, w, oh, ow, kh, kw, sh, sw, ph, pw, dh, dw);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

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

torch::Tensor depthwise_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    torch::Tensor bias,
    int64_t sh, int64_t sw,
    int64_t ph, int64_t pw,
    int64_t dh, int64_t dw);
"""

_depthwise_extension = load_inline(
    name="depthwise_asymmetric_tiled_cuda",
    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,
        )
        self.register_buffer("_empty_bias", torch.empty(0), persistent=False)

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        conv = self.conv2d

        if (
            not x.is_cuda
            or x.dtype != torch.float32
            or conv.weight.dtype != torch.float32
            or x.dim() != 4
            or torch.is_autocast_enabled()
        ):
            return conv(x)

        return _depthwise_extension.depthwise_cuda(
            x.contiguous(),
            conv.weight.contiguous(),
            conv.bias if conv.bias is not None else self._empty_bias,
            conv.stride[0],
            conv.stride[1],
            conv.padding[0],
            conv.padding[1],
            conv.dilation[0],
            conv.dilation[1],
        )
```