```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 depthwise3_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<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>

template<int ROWS, bool HAS_BIAS>
__global__ void depthwise3_kernel(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int channels,
    int out_channels,
    int multiplier,
    int height,
    int width,
    int out_height,
    int out_width)
{
    const int ox = 2 * (blockIdx.x * blockDim.x + threadIdx.x);
    const int oy = blockIdx.y * ROWS;
    if (ox >= out_width) return;

    const int nc = blockIdx.z;
    const int n = nc / out_channels;
    const int oc = nc % out_channels;
    const int ic = oc / multiplier;

    const float* w = weight + static_cast<int64_t>(oc) * 9;
    const float w0 = w[0], w1 = w[1], w2 = w[2];
    const float w3 = w[3], w4 = w[4], w5 = w[5];
    const float w6 = w[6], w7 = w[7], w8 = w[8];
    const float b = HAS_BIAS ? bias[oc] : 0.0f;

    const float* src =
        x + (static_cast<int64_t>(n) * channels + ic) * height * width
          + static_cast<int64_t>(oy) * width + ox;

    float2 top01 = *reinterpret_cast<const float2*>(src);
    float2 top23 = *reinterpret_cast<const float2*>(src + 2);
    float2 mid01 = *reinterpret_cast<const float2*>(src + width);
    float2 mid23 = *reinterpret_cast<const float2*>(src + width + 2);

    float* dst =
        output + static_cast<int64_t>(nc) * out_height * out_width
               + static_cast<int64_t>(oy) * out_width + ox;

    #pragma unroll
    for (int r = 0; r < ROWS; ++r) {
        if (oy + r < out_height) {
            const float* bottom = src + static_cast<int64_t>(r + 2) * width;
            const float2 bot01 = *reinterpret_cast<const float2*>(bottom);
            const float2 bot23 = *reinterpret_cast<const float2*>(bottom + 2);

            float a = top01.x * w0;
            a = fmaf(top01.y, w1, a);
            a = fmaf(top23.x, w2, a);
            a = fmaf(mid01.x, w3, a);
            a = fmaf(mid01.y, w4, a);
            a = fmaf(mid23.x, w5, a);
            a = fmaf(bot01.x, w6, a);
            a = fmaf(bot01.y, w7, a);
            a = fmaf(bot23.x, w8, a);

            float z = top01.y * w0;
            z = fmaf(top23.x, w1, z);
            z = fmaf(top23.y, w2, z);
            z = fmaf(mid01.y, w3, z);
            z = fmaf(mid23.x, w4, z);
            z = fmaf(mid23.y, w5, z);
            z = fmaf(bot01.y, w6, z);
            z = fmaf(bot23.x, w7, z);
            z = fmaf(bot23.y, w8, z);

            if (HAS_BIAS) {
                a += b;
                z += b;
            }

            *reinterpret_cast<float2*>(
                dst + static_cast<int64_t>(r) * out_width
            ) = make_float2(a, z);

            top01 = mid01;
            top23 = mid23;
            mid01 = bot01;
            mid23 = bot23;
        }
    }
}

torch::Tensor depthwise3_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias)
{
    const c10::cuda::CUDAGuard device_guard(x.device());

    const int batch = x.size(0);
    const int channels = x.size(1);
    const int height = x.size(2);
    const int width = x.size(3);
    const int out_channels = weight.size(0);
    const int out_height = height - 2;
    const int out_width = width - 2;

    auto output = torch::empty(
        {batch, out_channels, out_height, out_width}, x.options());

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

    constexpr int threads = 128;
    constexpr int rows = 8;
    const dim3 grid(
        (out_width + 2 * threads - 1) / (2 * threads),
        (out_height + rows - 1) / rows,
        batch * out_channels);

    const auto stream = at::cuda::getCurrentCUDAStream();

    if (bias.has_value()) {
        depthwise3_kernel<rows, true><<<grid, threads, 0, stream>>>(
            x.data_ptr<float>(),
            weight.data_ptr<float>(),
            bias->data_ptr<float>(),
            output.data_ptr<float>(),
            channels, out_channels, out_channels / channels,
            height, width, out_height, out_width);
    } else {
        depthwise3_kernel<rows, false><<<grid, threads, 0, stream>>>(
            x.data_ptr<float>(),
            weight.data_ptr<float>(),
            nullptr,
            output.data_ptr<float>(),
            channels, out_channels, out_channels / channels,
            height, width, out_height, out_width);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_depthwise_extension = load_inline(
    name="depthwise3_paired_rows_cuda",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["depthwise3_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: int,
        stride: int = 1,
        padding: int = 0,
        bias: bool = False,
    ):
        super().__init__()
        self.conv2d = nn.Conv2d(
            in_channels,
            out_channels,
            kernel_size=(kernel_size, kernel_size),
            stride=stride,
            padding=padding,
            groups=in_channels,
            bias=bias,
        )

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        conv = self.conv2d
        if (
            x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 4
            and x.is_contiguous()
            and x.data_ptr() % 8 == 0
            and conv.weight.is_cuda
            and conv.weight.device == x.device
            and conv.weight.dtype == torch.float32
            and conv.weight.is_contiguous()
            and (
                conv.bias is None
                or (
                    conv.bias.device == x.device
                    and conv.bias.dtype == torch.float32
                    and conv.bias.is_contiguous()
                )
            )
            and conv.kernel_size == (3, 3)
            and conv.stride == (1, 1)
            and conv.padding == (0, 0)
            and x.shape[1] == conv.in_channels
            and x.shape[2] >= 3
            and x.shape[3] >= 4
            and x.shape[3] % 2 == 0
            and x.shape[0] * conv.out_channels <= 65535
            and (x.shape[2] - 2 + 7) // 8 <= 65535
            and not torch.is_autocast_enabled()
        ):
            return _depthwise_extension.depthwise3_cuda(
                x, conv.weight, conv.bias
            )
        return conv(x)
```