```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_3x7_cuda(
    torch::Tensor input,
    torch::Tensor weight,
    torch::Tensor bias);
"""

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

__device__ __forceinline__ void accumulate_four(
    float4& s,
    const float4 a,
    const float4 b,
    const float* __restrict__ w)
{
    s.x = fmaf(a.x, w[0], s.x);
    s.y = fmaf(a.y, w[0], s.y);
    s.z = fmaf(a.z, w[0], s.z);
    s.w = fmaf(a.w, w[0], s.w);

    s.x = fmaf(a.y, w[1], s.x);
    s.y = fmaf(a.z, w[1], s.y);
    s.z = fmaf(a.w, w[1], s.z);
    s.w = fmaf(b.x, w[1], s.w);

    s.x = fmaf(a.z, w[2], s.x);
    s.y = fmaf(a.w, w[2], s.y);
    s.z = fmaf(b.x, w[2], s.z);
    s.w = fmaf(b.y, w[2], s.w);

    s.x = fmaf(a.w, w[3], s.x);
    s.y = fmaf(b.x, w[3], s.y);
    s.z = fmaf(b.y, w[3], s.z);
    s.w = fmaf(b.z, w[3], s.w);

    s.x = fmaf(b.x, w[4], s.x);
    s.y = fmaf(b.y, w[4], s.y);
    s.z = fmaf(b.z, w[4], s.z);
    s.w = fmaf(b.w, w[4], s.w);
}

__device__ __forceinline__ void accumulate_tail(
    float4& s,
    const float4 b,
    const float2 c,
    const float* __restrict__ w)
{
    s.x = fmaf(b.y, w[5], s.x);
    s.y = fmaf(b.z, w[5], s.y);
    s.z = fmaf(b.w, w[5], s.z);
    s.w = fmaf(c.x, w[5], s.w);

    s.x = fmaf(b.z, w[6], s.x);
    s.y = fmaf(b.w, w[6], s.y);
    s.z = fmaf(c.x, w[6], s.z);
    s.w = fmaf(c.y, w[6], s.w);
}

template<int ROWS>
__global__ __launch_bounds__(128)
void depthwise_3x7_register_kernel(
    const float* __restrict__ input,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int channels,
    int height)
{
    constexpr int WIDTH = 256;
    constexpr int OUT_WIDTH = 250;

    const int column4 = threadIdx.x & 63;
    const int y0 = blockIdx.x * (2 * ROWS)
                 + (threadIdx.x >> 6) * ROWS;
    const int out_height = height - 2;

    if (column4 >= 63 || y0 >= out_height)
        return;

    const int plane = blockIdx.y;
    const int channel = plane % channels;

    const float* __restrict__ source =
        input + static_cast<int64_t>(plane) * height * WIDTH
              + static_cast<int64_t>(y0) * WIDTH + column4 * 4;

    float w[21];
    #pragma unroll
    for (int k = 0; k < 21; ++k)
        w[k] = __ldg(weight + channel * 21 + k);

    float4 sums[ROWS];
    #pragma unroll
    for (int r = 0; r < ROWS; ++r)
        sums[r] = make_float4(0.f, 0.f, 0.f, 0.f);

    #pragma unroll
    for (int r = 0; r < ROWS + 2; ++r) {
        float4 a = make_float4(0.f, 0.f, 0.f, 0.f);
        float4 b = make_float4(0.f, 0.f, 0.f, 0.f);
        float2 c = make_float2(0.f, 0.f);

        if (y0 + r < height) {
            const float* row = source + r * WIDTH;
            a = __ldg(reinterpret_cast<const float4*>(row));
            b = __ldg(reinterpret_cast<const float4*>(row + 4));
            if (column4 < 62)
                c = __ldg(reinterpret_cast<const float2*>(row + 8));
        }

        if (r < ROWS) {
            accumulate_four(sums[r], a, b, w);
            accumulate_tail(sums[r], b, c, w);
        }
        if (r >= 1 && r < ROWS + 1) {
            accumulate_four(sums[r - 1], a, b, w + 7);
            accumulate_tail(sums[r - 1], b, c, w + 7);
        }
        if (r >= 2) {
            accumulate_four(sums[r - 2], a, b, w + 14);
            accumulate_tail(sums[r - 2], b, c, w + 14);
        }
    }

    const float bv = bias ? __ldg(bias + channel) : 0.f;
    float* __restrict__ destination =
        output + static_cast<int64_t>(plane) * out_height * OUT_WIDTH
               + static_cast<int64_t>(y0) * OUT_WIDTH + column4 * 4;

    #pragma unroll
    for (int r = 0; r < ROWS; ++r) {
        if (y0 + r < out_height) {
            float2* dst =
                reinterpret_cast<float2*>(destination + r * OUT_WIDTH);
            dst[0] = make_float2(sums[r].x + bv, sums[r].y + bv);
            if (column4 < 62)
                dst[1] = make_float2(sums[r].z + bv, sums[r].w + bv);
        }
    }
}

torch::Tensor depthwise_3x7_cuda(
    torch::Tensor input,
    torch::Tensor weight,
    torch::Tensor bias)
{
    TORCH_CHECK(input.is_cuda() && weight.is_cuda(),
                "Input and weight must be CUDA tensors");
    TORCH_CHECK(input.device() == weight.device(), "Device mismatch");
    TORCH_CHECK(input.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat,
                "Expected float32 tensors");
    TORCH_CHECK(input.is_contiguous() && weight.is_contiguous(),
                "Expected contiguous tensors");
    TORCH_CHECK(input.dim() == 4 &&
                input.size(2) >= 3 && input.size(3) == 256,
                "Expected NCHW input with width 256 and height >= 3");
    TORCH_CHECK(weight.dim() == 4 &&
                weight.size(0) == input.size(1) &&
                weight.size(1) == 1 &&
                weight.size(2) == 3 &&
                weight.size(3) == 7,
                "Expected depthwise 3x7 weights");

    const c10::cuda::CUDAGuard guard(input.device());
    const int64_t batch = input.size(0);
    const int64_t channels = input.size(1);
    const int64_t height = input.size(2);

    TORCH_CHECK(batch * channels <= 65535, "Too many input planes");

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

    auto output = torch::empty(
        {batch, channels, height - 2, 250}, input.options());

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

    constexpr int ROWS = 8;
    const dim3 grid(
        static_cast<unsigned>((height - 2 + 2 * ROWS - 1) / (2 * ROWS)),
        static_cast<unsigned>(batch * channels));

    depthwise_3x7_register_kernel<ROWS>
        <<<grid, 128, 0, at::cuda::getCurrentCUDAStream()>>>(
            input.data_ptr<float>(),
            weight.data_ptr<float>(),
            bias_ptr,
            output.data_ptr<float>(),
            static_cast<int>(channels),
            static_cast<int>(height));

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

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

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        conv = self.conv2d
        if (
            x.is_cuda
            and x.dtype == torch.float32
            and x.dim() == 4
            and x.is_contiguous()
            and x.shape[1] == conv.in_channels
            and x.shape[2] >= 3
            and x.shape[3] == 256
            and x.shape[0] * x.shape[1] <= 65535
            and conv.weight.device == x.device
            and conv.weight.dtype == torch.float32
            and conv.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 not torch.is_autocast_enabled()
        ):
            return _depthwise_extension.depthwise_3x7_cuda(
                x,
                conv.weight,
                conv.bias if conv.bias is not None else self._empty_bias,
            )
        return conv(x)
```