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

constexpr int TILE_W = 128;
constexpr int TILE_H = 8;
constexpr int SHARED_W = TILE_W + 6;
constexpr int SHARED_H = TILE_H + 2;

__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
) {
    __shared__ float tile[SHARED_H][SHARED_W];
    __shared__ float filter[21];
    __shared__ float channel_bias;

    const int tx = 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* src =
        input + static_cast<int64_t>(plane) * height * width;
    float* dst =
        output + static_cast<int64_t>(plane) * out_height * out_width;

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

    if (tx < 21)
        filter[tx] = weight[channel * 21 + tx];
    if (tx == 0)
        channel_bias = bias ? bias[channel] : 0.0f;

    __syncthreads();

    const int ox = x0 + tx;
    if (ox >= out_width)
        return;

    float sums[TILE_H];

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

    #pragma unroll
    for (int kh = 0; kh < 3; ++kh) {
        #pragma unroll
        for (int kw = 0; kw < 7; ++kw) {
            const float w = filter[kh * 7 + kw];

            #pragma unroll
            for (int row = 0; row < TILE_H; ++row)
                sums[row] = fmaf(tile[row + kh][tx + kw], w, sums[row]);
        }
    }

    #pragma unroll
    for (int row = 0; row < TILE_H; ++row) {
        const int oy = y0 + row;
        if (oy < out_height)
            dst[static_cast<int64_t>(oy) * out_width + ox] =
                sums[row] + channel_bias;
    }
}

torch::Tensor depthwise_3x7_cuda(
    torch::Tensor input,
    torch::Tensor weight,
    torch::Tensor bias,
    bool has_bias
) {
    const c10::cuda::CUDAGuard device_guard(input.device());

    const int64_t batch = input.size(0);
    const int channels = static_cast<int>(input.size(1));
    const int height = static_cast<int>(input.size(2));
    const int width = static_cast<int>(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 dim3 grid(
        (out_width + TILE_W - 1) / TILE_W,
        (out_height + TILE_H - 1) / TILE_H,
        static_cast<unsigned int>(batch * channels)
    );

    depthwise_3x7_kernel<<<
        grid, TILE_W, 0, at::cuda::getCurrentCUDAStream()
    >>>(
        input.data_ptr<float>(),
        weight.data_ptr<float>(),
        has_bias ? bias.data_ptr<float>() : nullptr,
        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,
            torch::Tensor bias,
            bool has_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

        use_custom = (
            x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 4
            and weight.dtype == torch.float32
            and weight.device == x.device
            and (bias is None or (
                bias.dtype == torch.float32 and bias.device == x.device
            ))
            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 x.shape[2] >= 3
            and x.shape[3] >= 7
            and x.shape[0] * x.shape[1] <= 65535
            and (x.shape[2] - 2 + 7) // 8 <= 65535
            and not torch.is_autocast_enabled()
        )

        if not use_custom:
            return conv(x)

        return _extension.depthwise_3x7_cuda(
            x.contiguous(),
            weight.contiguous(),
            bias.contiguous() if bias is not None else weight,
            bias is not None,
        )
```