```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 depthwise3x3_cuda(
    torch::Tensor input,
    torch::Tensor weight,
    torch::Tensor bias,
    bool has_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>
#include <cstdint>

__global__ void depthwise3x3_kernel(
    const float* __restrict__ input,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int channels,
    int height,
    int width,
    int out_height,
    int out_width)
{
    // Each thread computes two adjacent columns across four output rows.
    const int x = (blockIdx.x * blockDim.x + threadIdx.x) * 2;
    const int y = (blockIdx.y * blockDim.y + threadIdx.y) * 4;
    if (x >= out_width || y >= out_height)
        return;

    const int plane = blockIdx.z;
    const int channel = plane % channels;

    const float* src =
        input + static_cast<int64_t>(plane) * height * width
              + static_cast<int64_t>(y) * width + x;
    float* dst =
        output + static_cast<int64_t>(plane) * out_height * out_width
               + static_cast<int64_t>(y) * out_width + x;

    const float* filter = weight + channel * 9;
    const float w0 = filter[0];
    const float w1 = filter[1];
    const float w2 = filter[2];
    const float w3 = filter[3];
    const float w4 = filter[4];
    const float w5 = filter[5];
    const float w6 = filter[6];
    const float w7 = filter[7];
    const float w8 = filter[8];
    const float b = bias ? bias[channel] : 0.0f;

    float2 a0 = *reinterpret_cast<const float2*>(src);
    float2 a1 = *reinterpret_cast<const float2*>(src + 2);
    float2 b0 = *reinterpret_cast<const float2*>(src + width);
    float2 b1 = *reinterpret_cast<const float2*>(src + width + 2);

    #pragma unroll
    for (int row = 0; row < 4; ++row) {
        if (y + row < out_height) {
            const float* next = src + static_cast<int64_t>(row + 2) * width;
            const float2 c0 = *reinterpret_cast<const float2*>(next);
            const float2 c1 = *reinterpret_cast<const float2*>(next + 2);

            float left = a0.x * w0;
            left = fmaf(a0.y, w1, left);
            left = fmaf(a1.x, w2, left);
            left = fmaf(b0.x, w3, left);
            left = fmaf(b0.y, w4, left);
            left = fmaf(b1.x, w5, left);
            left = fmaf(c0.x, w6, left);
            left = fmaf(c0.y, w7, left);
            left = fmaf(c1.x, w8, left);

            float right = a0.y * w0;
            right = fmaf(a1.x, w1, right);
            right = fmaf(a1.y, w2, right);
            right = fmaf(b0.y, w3, right);
            right = fmaf(b1.x, w4, right);
            right = fmaf(b1.y, w5, right);
            right = fmaf(c0.y, w6, right);
            right = fmaf(c1.x, w7, right);
            right = fmaf(c1.y, w8, right);

            *reinterpret_cast<float2*>(
                dst + static_cast<int64_t>(row) * out_width
            ) = make_float2(left + b, right + b);

            a0 = b0;
            a1 = b1;
            b0 = c0;
            b1 = c1;
        }
    }
}

torch::Tensor depthwise3x3_cuda(
    torch::Tensor input,
    torch::Tensor weight,
    torch::Tensor bias,
    bool has_bias)
{
    const c10::cuda::CUDAGuard device_guard(input.device());

    const int64_t batch = input.size(0);
    const int channels = static_cast<int>(input.size(1));
    const int height = static_cast<int>(input.size(2));
    const int width = static_cast<int>(input.size(3));
    const int out_height = height - 2;
    const int out_width = width - 2;

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

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

    const dim3 block(32, 8);
    const dim3 grid(
        (out_width + 63) / 64,
        (out_height + 31) / 32,
        static_cast<unsigned int>(batch * channels));

    depthwise3x3_kernel<<<
        grid, block, 0, at::cuda::getCurrentCUDAStream()
    >>>(
        input.data_ptr<float>(),
        weight.data_ptr<float>(),
        has_bias ? bias.data_ptr<float>() : nullptr,
        output.data_ptr<float>(),
        channels,
        height,
        width,
        out_height,
        out_width);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_depthwise_extension = load_inline(
    name="depthwise3x3_tiled_float2_v1",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["depthwise3x3_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
        weight = conv.weight
        bias = conv.bias

        if (
            x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 4
            and x.is_contiguous()
            and weight.device == x.device
            and weight.dtype == torch.float32
            and weight.is_contiguous()
            and conv.kernel_size == (3, 3)
            and conv.stride == (1, 1)
            and conv.padding == (0, 0)
            and conv.dilation == (1, 1)
            and conv.in_channels == conv.out_channels == x.shape[1]
            and x.shape[2] >= 3
            and x.shape[3] >= 4
            and x.shape[3] % 2 == 0
            and x.shape[0] * x.shape[1] <= 65535
            and (x.shape[2] - 2 + 31) // 32 <= 65535
            and (
                bias is None
                or (
                    bias.device == x.device
                    and bias.dtype == torch.float32
                    and bias.is_contiguous()
                )
            )
        ):
            return _depthwise_extension.depthwise3x3_cuda(
                x,
                weight,
                bias if bias is not None else weight,
                bias is not None,
            )

        return conv(x)
```