```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_eight(
    float (&sum)[8],
    const float (&values)[14],
    const float* __restrict__ weights)
{
    #pragma unroll
    for (int k = 0; k < 7; ++k) {
        const float w = weights[k];
        #pragma unroll
        for (int j = 0; j < 8; ++j)
            sum[j] = fmaf(values[j + k], w, sum[j]);
    }
}

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

    const int lane = threadIdx.x & 31;
    const int row_group = threadIdx.x >> 5;
    const int y0 = blockIdx.x * (4 * ROWS) + row_group * ROWS;
    const int out_height = height - 2;

    if (y0 >= out_height)
        return;

    const int plane = blockIdx.y;
    const int channel = plane % channels;
    const int column = lane * 8;

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

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

    float sums[ROWS][8];
    #pragma unroll
    for (int r = 0; r < ROWS; ++r) {
        #pragma unroll
        for (int j = 0; j < 8; ++j)
            sums[r][j] = 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);
        float4 c = make_float4(0.f, 0.f, 0.f, 0.f);
        float2 d = 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 (lane < 31) {
                c = __ldg(reinterpret_cast<const float4*>(row + 8));
                d = __ldg(reinterpret_cast<const float2*>(row + 12));
            }
        }

        const float values[14] = {
            a.x, a.y, a.z, a.w,
            b.x, b.y, b.z, b.w,
            c.x, c.y, c.z, c.w,
            d.x, d.y
        };

        if (r < ROWS)
            accumulate_eight(sums[r], values, weights);
        if (r >= 1 && r < ROWS + 1)
            accumulate_eight(sums[r - 1], values, weights + 7);
        if (r >= 2)
            accumulate_eight(sums[r - 2], values, weights + 14);
    }

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

    #pragma unroll
    for (int r = 0; r < ROWS; ++r) {
        if (y0 + r < out_height) {
            float4* dst = reinterpret_cast<float4*>(
                destination + r * WIDTH);
            dst[0] = make_float4(
                sums[r][0] + bv,
                sums[r][1] + bv,
                sums[r][2] + bv,
                sums[r][3] + bv);
            dst[1] = make_float4(
                sums[r][4] + bv,
                sums[r][5] + bv,
                sums[r][6] + bv,
                sums[r][7] + 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>();
    }

    // Padded rows permit aligned vector stores throughout the output.
    auto storage = torch::empty(
        {batch, channels, height - 2, 256}, input.options());

    if (storage.numel() != 0) {
        constexpr int ROWS = 4;
        const dim3 grid(
            static_cast<unsigned>(
                (height - 2 + 4 * ROWS - 1) / (4 * ROWS)),
            static_cast<unsigned>(batch * channels));

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

        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }

    return storage.narrow(3, 0, 250);
}
"""

_depthwise_extension = load_inline(
    name="depthwise_3x7_eight_aligned_v4",
    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)
```