```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_pair(
    float2& sum,
    const float2 a,
    const float2 b,
    const float2 c,
    const float2 d,
    const float* __restrict__ w)
{
    sum.x = fmaf(a.x, w[0], sum.x);
    sum.y = fmaf(a.y, w[0], sum.y);
    sum.x = fmaf(a.y, w[1], sum.x);
    sum.y = fmaf(b.x, w[1], sum.y);
    sum.x = fmaf(b.x, w[2], sum.x);
    sum.y = fmaf(b.y, w[2], sum.y);
    sum.x = fmaf(b.y, w[3], sum.x);
    sum.y = fmaf(c.x, w[3], sum.y);
    sum.x = fmaf(c.x, w[4], sum.x);
    sum.y = fmaf(c.y, w[4], sum.y);
    sum.x = fmaf(c.y, w[5], sum.x);
    sum.y = fmaf(d.x, w[5], sum.y);
    sum.x = fmaf(d.x, w[6], sum.x);
    sum.y = fmaf(d.y, w[6], sum.y);
}

template<int ROWS>
__global__ __launch_bounds__(128, 8)
void depthwise_3x7_streaming_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;
    constexpr int INPUT_ROWS = ROWS + 2;

    __shared__ float4 tile[INPUT_ROWS * 64];
    __shared__ float shared_weight[21];

    const int tid = threadIdx.x;
    const int plane = blockIdx.y;
    const int channel = plane % channels;
    const int y0 = blockIdx.x * ROWS;
    const int out_height = height - 2;

    const float4* __restrict__ source =
        reinterpret_cast<const float4*>(
            input + static_cast<int64_t>(plane) * height * WIDTH
                  + static_cast<int64_t>(y0) * WIDTH);

    #pragma unroll
    for (int i = 0; i < INPUT_ROWS / 2; ++i) {
        const int index = i * 128 + tid;
        const int row = index / 64;
        float4 value = make_float4(0.f, 0.f, 0.f, 0.f);
        if (y0 + row < height)
            value = source[index];
        tile[index] = value;
    }

    if (tid < 21)
        shared_weight[tid] = weight[channel * 21 + tid];

    __syncthreads();

    if (tid >= OUT_WIDTH / 2)
        return;

    float w[21];
    #pragma unroll
    for (int k = 0; k < 21; ++k)
        w[k] = shared_weight[k];

    const float bias_value = bias ? __ldg(bias + channel) : 0.f;
    const float2* __restrict__ shared2 =
        reinterpret_cast<const float2*>(tile);

    float2* __restrict__ destination =
        reinterpret_cast<float2*>(
            output + static_cast<int64_t>(plane) * out_height * OUT_WIDTH
                   + static_cast<int64_t>(y0) * OUT_WIDTH);

    // Retain only the two unfinished output rows in registers.
    float2 current = make_float2(0.f, 0.f);
    float2 next = make_float2(0.f, 0.f);

    #pragma unroll
    for (int r = 0; r < INPUT_ROWS; ++r) {
        const float2* row = shared2 + r * 128 + tid;
        const float2 a = row[0];
        const float2 b = row[1];
        const float2 c = row[2];
        const float2 d = row[3];

        if (r >= 2) {
            accumulate_pair(current, a, b, c, d, w + 14);
            if (y0 + r - 2 < out_height) {
                destination[(r - 2) * 125 + tid] = make_float2(
                    current.x + bias_value,
                    current.y + bias_value);
            }
        }

        current = next;
        next = make_float2(0.f, 0.f);

        if (r >= 1 && r < ROWS + 1)
            accumulate_pair(current, a, b, c, d, w + 7);

        if (r < ROWS)
            accumulate_pair(next, a, b, c, d, w);
    }
}

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 = 14;
    const dim3 grid(
        static_cast<unsigned>((height - 2 + ROWS - 1) / ROWS),
        static_cast<unsigned>(batch * channels));

    depthwise_3x7_streaming_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_streaming_rows14_v5",
    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)
```