```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__ input,
    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 = 32;
    constexpr int TILE_H = 16;
    constexpr int SHARED_W = TILE_W + 6;
    constexpr int SHARED_H = TILE_H + 2;

    __shared__ float tile[SHARED_H][SHARED_W];
    __shared__ float filter[21];

    const int tx = threadIdx.x;
    const int ty = threadIdx.y;
    const int tid = ty * 32 + tx;
    const int nc = blockIdx.z;
    const int channel = nc % channels;
    const int origin_x = blockIdx.x * TILE_W;
    const int origin_y = blockIdx.y * TILE_H;

    const float* plane =
        input + static_cast<int64_t>(nc) * height * width;

    for (int i = tid; i < SHARED_H * SHARED_W; i += 128) {
        const int sy = i / SHARED_W;
        const int sx = i % SHARED_W;
        const int iy = origin_y + sy;
        const int ix = origin_x + sx;
        tile[sy][sx] = (iy < height && ix < width)
            ? plane[static_cast<int64_t>(iy) * width + ix]
            : 0.0f;
    }

    if (tid < 21) {
        filter[tid] = weight[channel * 21 + tid];
    }
    __syncthreads();

    const int local_y = ty * 4;
    float acc0 = 0.0f;
    float acc1 = 0.0f;
    float acc2 = 0.0f;
    float acc3 = 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];

        const float v0 = tile[local_y + 0][tx + kx];
        const float v1 = tile[local_y + 1][tx + kx];
        const float v2 = tile[local_y + 2][tx + kx];
        const float v3 = tile[local_y + 3][tx + kx];
        const float v4 = tile[local_y + 4][tx + kx];
        const float v5 = tile[local_y + 5][tx + kx];

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

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

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

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

    const int ox = origin_x + tx;
    const int oy = origin_y + local_y;
    const float b = bias != nullptr ? bias[channel] : 0.0f;

    if (ox < out_width) {
        const int64_t base =
            (static_cast<int64_t>(nc) * out_height + oy) * out_width + ox;

        if (oy + 0 < out_height)
            output[base] = acc0 + b;
        if (oy + 1 < out_height)
            output[base + out_width] = acc1 + b;
        if (oy + 2 < out_height)
            output[base + 2LL * out_width] = acc2 + b;
        if (oy + 3 < out_height)
            output[base + 3LL * out_width] = acc3 + b;
    }
}

torch::Tensor depthwise_3x7_cuda(
    torch::Tensor input,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias)
{
    const c10::cuda::CUDAGuard device_guard(input.device());

    const int batch = input.size(0);
    const int channels = input.size(1);
    const int height = input.size(2);
    const int width = input.size(3);
    const int out_height = height - 2;
    const int out_width = width - 6;

    auto output = torch::empty(
        {batch, channels, out_height, out_width}, input.options());

    if (output.numel() == 0)
        return output;

    const float* bias_ptr =
        bias.has_value() ? bias.value().data_ptr<float>() : nullptr;

    const dim3 block(32, 4);
    const dim3 grid(
        (out_width + 31) / 32,
        (out_height + 15) / 16,
        batch * channels);

    depthwise_3x7_kernel<<<
        grid, block, 0, at::cuda::getCurrentCUDAStream()>>>(
        input.data_ptr<float>(),
        weight.data_ptr<float>(),
        bias_ptr,
        output.data_ptr<float>(),
        channels,
        height,
        width,
        out_height,
        out_width);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

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

torch::Tensor depthwise_3x7_cuda(
    torch::Tensor input,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias);
"""

_depthwise_extension = load_inline(
    name="depthwise_asymmetric_3x7_tiled",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["depthwise_3x7_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._specialized = (
            (kernel_size_h, kernel_size_w) == (3, 7)
            and (stride_h, stride_w) == (1, 1)
            and (padding_h, padding_w) == (0, 0)
            and (dilation_h, dilation_w) == (1, 1)
        )

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

        if (
            self._specialized
            and x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 4
            and x.is_contiguous()
            and weight.is_contiguous()
            and weight.dtype == x.dtype
            and weight.device == x.device
            and x.shape[1] == self.conv2d.in_channels
            and x.shape[2] >= 3
            and x.shape[3] >= 7
            and x.shape[0] * x.shape[1] <= 65535
            and (x.shape[2] - 2 + 15) // 16 <= 65535
            and (
                bias is None
                or (
                    bias.dtype == x.dtype
                    and bias.device == x.device
                    and bias.is_contiguous()
                )
            )
        ):
            return _depthwise_extension.depthwise_3x7_cuda(x, weight, bias)

        return self.conv2d(x)
```