```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 x,
    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)
void depthwise_3x7_vector_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 * (WIDTH / 4)];
    __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__ input4 =
        reinterpret_cast<const float4*>(
            input + static_cast<int64_t>(plane) * height * WIDTH);

    #pragma unroll
    for (int i = 0; i < INPUT_ROWS / 2; ++i) {
        const int row = 2 * i + tid / 64;
        const int col = tid % 64;
        const int iy = y0 + row;
        float4 value = make_float4(0.f, 0.f, 0.f, 0.f);
        if (iy < height)
            value = input4[static_cast<int64_t>(iy) * 64 + col];
        tile[row * 64 + col] = 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];

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

    const float2* shared2 = reinterpret_cast<const float2*>(tile);

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

        if (r < ROWS)
            accumulate_pair(sums[r], a, b, c, d, w);

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

        if (r >= 2)
            accumulate_pair(sums[r - 2], a, b, c, d, w + 14);
    }

    const float bias_value = bias ? bias[channel] : 0.f;
    float2* __restrict__ output2 =
        reinterpret_cast<float2*>(
            output + static_cast<int64_t>(plane) * out_height * OUT_WIDTH);

    #pragma unroll
    for (int r = 0; r < ROWS; ++r) {
        if (y0 + r < out_height) {
            float2 value = sums[r];
            value.x += bias_value;
            value.y += bias_value;
            output2[static_cast<int64_t>(y0 + r) * (OUT_WIDTH / 2) + tid]
                = value;
        }
    }
}

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

    const c10::cuda::CUDAGuard guard(x.device());

    const int64_t batch = x.size(0);
    const int64_t channels = x.size(1);
    const int64_t height = x.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() == x.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}, x.options());

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

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

    depthwise_3x7_vector_kernel<ROWS>
        <<<grid, 128, 0, at::cuda::getCurrentCUDAStream()>>>(
            x.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_vector_rows16",
    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.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)
```