```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 = 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 base_y = blockIdx.y * TILE_H;
    const int base_x = blockIdx.x * TILE_W;

    const int64_t input_base = static_cast<int64_t>(plane) * height * width;

    for (int i = tid; i < SHARED_H * SHARED_W; i += 256) {
        const int sy = i / SHARED_W;
        const int sx = i % SHARED_W;
        const int iy = base_y + sy;
        const int ix = base_x + sx;

        tile[sy][sx] = (iy < height && ix < width)
            ? input[input_base + static_cast<int64_t>(iy) * width + ix]
            : 0.0f;
    }

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

    const int tx = tid % TILE_W;
    const int ty = (tid / TILE_W) * 4;
    const int ox = base_x + tx;
    const int oy = base_y + ty;

    if (ox >= out_width || oy >= out_height) {
        return;
    }

    float r0 = 0.0f;
    float r1 = 0.0f;
    float r2 = 0.0f;
    float r3 = 0.0f;

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

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

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

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

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

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

    const float b = bias == nullptr ? 0.0f : bias[channel];
    const int64_t output_base =
        (static_cast<int64_t>(plane) * out_height + oy) * out_width + ox;

    output[output_base] = r0 + b;
    if (oy + 1 < out_height)
        output[output_base + out_width] = r1 + b;
    if (oy + 2 < out_height)
        output[output_base + 2LL * out_width] = r2 + b;
    if (oy + 3 < out_height)
        output[output_base + 3LL * out_width] = r3 + 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()
    );

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

    const dim3 grid(
        (out_width + 127) / 128,
        (out_height + 7) / 8,
        batch * channels
    );

    depthwise_3x7_kernel<<<
        grid, 256, 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;
}
"""

_extension = load_inline(
    name="depthwise_asymmetric_3x7_tiled_v1",
    cpp_sources=r"""
#include <torch/extension.h>
torch::Tensor depthwise_3x7_cuda(
    torch::Tensor input,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias
);
""",
    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,
        )

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

        if (
            x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 4
            and x.is_contiguous()
            and weight.device == x.device
            and weight.dtype == torch.float32
            and weight.is_contiguous()
            and conv.kernel_size == (3, 7)
            and conv.stride == (1, 1)
            and conv.padding == (0, 0)
            and conv.dilation == (1, 1)
            and x.shape[1] == conv.in_channels
            and 0 < x.shape[0] * x.shape[1] <= 65535
            and 3 <= x.shape[2] <= 524282
            and 7 <= x.shape[3] < 2147483520
            and not torch.is_autocast_enabled()
            and (
                bias is None
                or (
                    bias.device == x.device
                    and bias.dtype == torch.float32
                    and bias.is_contiguous()
                )
            )
        ):
            return _extension.depthwise_3x7_cuda(x, weight, bias)

        return conv(x)
```