```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 float4 c,
    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);

    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_four_kernel(
    const float* __restrict__ input,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int channels,
    int height)
{
    constexpr int TILE_ROWS = 2 * ROWS;
    constexpr int INPUT_ROWS = TILE_ROWS + 2;

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

    const int tid = threadIdx.x;
    const int col = tid & 63;
    const int group_row = (tid >> 6) * ROWS;
    const int plane = blockIdx.y;
    const int channel = plane % channels;
    const int tile_y = blockIdx.x * TILE_ROWS;
    const int out_height = height - 2;

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

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

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

    __syncthreads();

    if (col == 63)
        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 int y0 = tile_y + group_row;

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

    float4 current = make_float4(0.f, 0.f, 0.f, 0.f);
    float4 next = make_float4(0.f, 0.f, 0.f, 0.f);

    #pragma unroll
    for (int r = 0; r < ROWS + 2; ++r) {
        const float4* row = tile + (group_row + r) * 64 + col;
        const float4 a = row[0];
        const float4 b = row[1];
        float4 c = make_float4(0.f, 0.f, 0.f, 0.f);
        if (col < 62)
            c = row[2];

        if (r >= 2) {
            accumulate_four(current, a, b, c, w + 14);
            const int oy = y0 + r - 2;
            if (oy < out_height) {
                const int64_t offset =
                    static_cast<int64_t>(oy) * 125 + col * 2;
                destination[offset] = make_float2(
                    current.x + bias_value, current.y + bias_value);
                if (col < 62)
                    destination[offset + 1] = make_float2(
                        current.z + bias_value, current.w + bias_value);
            }
        }

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

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

        if (r < ROWS)
            accumulate_four(next, a, b, c, 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 = 8;
    const dim3 grid(
        static_cast<unsigned>((height - 2 + 2 * ROWS - 1) / (2 * ROWS)),
        static_cast<unsigned>(batch * channels));

    depthwise_3x7_four_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_four_streaming_v6",
    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)
```