```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_forward(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias,
    int64_t sh, int64_t sw,
    int64_t ph, int64_t pw,
    int64_t dh, int64_t dw);
"""

_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>

constexpr int TILE_W = 128;
constexpr int TILE_H = 8;

// Each thread computes eight vertically adjacent outputs. Shared memory
// reuses input pixels across both horizontal and vertical kernel positions.
__global__ void depthwise_3x7(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ out,
    int channels, int height, int width,
    int out_h, int out_w, int tiles_w)
{
    __shared__ float tile[TILE_H + 2][TILE_W + 6];

    const int tx = threadIdx.x;
    const int nc = blockIdx.y;
    const int channel = nc % channels;
    const int start_y = (blockIdx.x / tiles_w) * TILE_H;
    const int start_x = (blockIdx.x % tiles_w) * TILE_W;

    const float* input = x + int64_t(nc) * height * width;
    const float* filter = weight + channel * 21;

    for (int i = tx; i < (TILE_H + 2) * (TILE_W + 6);
         i += TILE_W) {
        const int r = i / (TILE_W + 6);
        const int c = i % (TILE_W + 6);
        const int iy = start_y + r;
        const int ix = start_x + c;
        tile[r][c] = (iy < height && ix < width)
            ? input[int64_t(iy) * width + ix] : 0.0f;
    }
    __syncthreads();

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

    #pragma unroll
    for (int kx = 0; kx < 7; ++kx) {
        const float w0 = filter[kx];
        const float w1 = filter[7 + kx];
        const float w2 = filter[14 + kx];

        #pragma unroll
        for (int r = 0; r < TILE_H + 2; ++r) {
            const float value = tile[r][tx + kx];
            if (r < TILE_H)
                accum[r] = fmaf(value, w0, accum[r]);
            if (r >= 1 && r < TILE_H + 1)
                accum[r - 1] = fmaf(value, w1, accum[r - 1]);
            if (r >= 2)
                accum[r - 2] = fmaf(value, w2, accum[r - 2]);
        }
    }

    const int ox = start_x + tx;
    const float b = bias ? bias[channel] : 0.0f;
    if (ox < out_w) {
        #pragma unroll
        for (int r = 0; r < TILE_H; ++r) {
            const int oy = start_y + r;
            if (oy < out_h)
                out[(int64_t(nc) * out_h + oy) * out_w + ox]
                    = accum[r] + b;
        }
    }
}

__global__ void depthwise_generic(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ out,
    int64_t total,
    int channels, int height, int width,
    int out_h, int out_w, int kh, int kw,
    int sh, int sw, int ph, int pw, int dh, int dw)
{
    for (int64_t index = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         index < total;
         index += int64_t(blockDim.x) * gridDim.x) {
        const int ox = index % out_w;
        const int oy = (index / out_w) % out_h;
        const int64_t nc = index / (int64_t(out_h) * out_w);
        const int channel = nc % channels;

        const float* input = x + nc * height * width;
        const float* filter = weight + int64_t(channel) * kh * kw;
        const int y0 = oy * sh - ph;
        const int x0 = ox * sw - pw;

        float sum = 0.0f;
        for (int ky = 0; ky < kh; ++ky) {
            const int iy = y0 + ky * dh;
            if (iy < 0 || iy >= height)
                continue;
            for (int kx = 0; kx < kw; ++kx) {
                const int ix = x0 + kx * dw;
                if (ix >= 0 && ix < width)
                    sum = fmaf(
                        input[int64_t(iy) * width + ix],
                        filter[ky * kw + kx], sum);
            }
        }
        out[index] = sum + (bias ? bias[channel] : 0.0f);
    }
}

torch::Tensor depthwise_forward(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<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 4D tensors");
    TORCH_CHECK(x.device() == weight.device(), "Device mismatch");
    TORCH_CHECK(sh > 0 && sw > 0 && dh > 0 && dw > 0,
                "Stride and dilation must be positive");

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

    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(weight.size(0) == c && weight.size(1) == 1,
                "Expected depthwise convolution weights");
    TORCH_CHECK(h + 2 * ph >= dh * (kh - 1) + 1 &&
                w + 2 * pw >= dw * (kw - 1) + 1,
                "Kernel is larger than the padded input");

    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;
    auto output = torch::empty({n, c, oh, ow}, x.options());

    torch::Tensor contiguous_bias;
    const float* bias_ptr = nullptr;
    if (bias.has_value()) {
        TORCH_CHECK(bias->device() == x.device() &&
                    bias->scalar_type() == at::kFloat &&
                    bias->numel() == c, "Invalid bias");
        contiguous_bias = bias->contiguous();
        bias_ptr = contiguous_bias.data_ptr<float>();
    }

    const int64_t total = output.numel();
    if (total == 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) {
        const int tiles_w = (ow + TILE_W - 1) / TILE_W;
        const int tiles_h = (oh + TILE_H - 1) / TILE_H;
        const dim3 grid(tiles_w * tiles_h, n * c);

        depthwise_3x7<<<grid, TILE_W, 0, stream>>>(
            x.data_ptr<float>(), weight.data_ptr<float>(), bias_ptr,
            output.data_ptr<float>(), c, h, w, oh, ow, tiles_w);
    } else {
        const int threads = 256;
        const int64_t needed = (total + threads - 1) / threads;
        const int blocks = needed < 65535 ? int(needed) : 65535;

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

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_depthwise_ext = load_inline(
    name="depthwise_asymmetric_tiled_cuda",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["depthwise_forward"],
    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:
        conv = self.conv2d
        if (
            not x.is_cuda
            or x.dtype != torch.float32
            or conv.weight.dtype != torch.float32
            or x.ndim != 4
            or torch.is_autocast_enabled()
        ):
            return conv(x)

        return _depthwise_ext.depthwise_forward(
            x,
            conv.weight,
            conv.bias,
            *conv.stride,
            *conv.padding,
            *conv.dilation,
        )
```