```python
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline


_cpp_source = r"""
#include <torch/extension.h>
#include <c10/util/Optional.h>

torch::Tensor depthwise3x3_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 <c10/util/Optional.h>
#include <cuda_runtime.h>

template <bool HasBias>
__global__ void depthwise3x3_kernel(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int channels,
    int out_channels,
    int height,
    int width,
    int out_height,
    int out_width)
{
    constexpr int TileWidth = 128;
    constexpr int TileHeight = 16;
    constexpr int SharedWidth = TileWidth + 2;
    constexpr int SharedHeight = TileHeight + 2;

    __shared__ float tile[SharedHeight][SharedWidth];

    const int nc = blockIdx.z;
    const int batch = nc / out_channels;
    const int oc = nc % out_channels;
    const int ic = oc / (out_channels / channels);

    const int base_y = blockIdx.y * TileHeight;
    const int base_x = blockIdx.x * TileWidth;
    const int tid = threadIdx.y * 32 + threadIdx.x;

    const float* input_plane =
        x + (static_cast<int64_t>(batch) * channels + ic) * height * width;

    for (int i = tid; i < SharedHeight * SharedWidth; i += 256) {
        const int sy = i / SharedWidth;
        const int sx = i % SharedWidth;
        const int iy = base_y + sy;
        const int ix = base_x + sx;

        tile[sy][sx] =
            (iy < height && ix < width)
            ? input_plane[static_cast<int64_t>(iy) * width + ix]
            : 0.0f;
    }

    float filter[9];
    #pragma unroll
    for (int i = 0; i < 9; ++i) {
        filter[i] = weight[static_cast<int64_t>(oc) * 9 + i];
    }

    const float bias_value = HasBias ? bias[oc] : 0.0f;

    __syncthreads();

    float* output_plane =
        output + static_cast<int64_t>(nc) * out_height * out_width;

    const int sy = threadIdx.y * 2;
    const int oy = base_y + sy;

    #pragma unroll
    for (int segment = 0; segment < 4; ++segment) {
        const int sx = threadIdx.x + segment * 32;
        const int ox = base_x + sx;

        float sum0 = 0.0f;
        float sum1 = 0.0f;

        #pragma unroll
        for (int ky = 0; ky < 3; ++ky) {
            #pragma unroll
            for (int kx = 0; kx < 3; ++kx) {
                const float w = filter[ky * 3 + kx];
                sum0 = fmaf(tile[sy + ky][sx + kx], w, sum0);
                sum1 = fmaf(tile[sy + ky + 1][sx + kx], w, sum1);
            }
        }

        if (ox < out_width) {
            if (oy < out_height) {
                output_plane[static_cast<int64_t>(oy) * out_width + ox] =
                    sum0 + bias_value;
            }
            if (oy + 1 < out_height) {
                output_plane[static_cast<int64_t>(oy + 1) * out_width + ox] =
                    sum1 + bias_value;
            }
        }
    }
}

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

    const int batch = static_cast<int>(x.size(0));
    const int channels = static_cast<int>(x.size(1));
    const int height = static_cast<int>(x.size(2));
    const int width = static_cast<int>(x.size(3));
    const int out_channels = static_cast<int>(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;
    }

    const dim3 block(32, 8);
    const dim3 grid(
        (out_width + 127) / 128,
        (out_height + 15) / 16,
        batch * out_channels);

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

    if (bias.has_value()) {
        depthwise3x3_kernel<true><<<grid, block, 0, stream>>>(
            x.data_ptr<float>(),
            weight.data_ptr<float>(),
            bias.value().data_ptr<float>(),
            output.data_ptr<float>(),
            channels, out_channels, height, width, out_height, out_width);
    } else {
        depthwise3x3_kernel<false><<<grid, block, 0, stream>>>(
            x.data_ptr<float>(),
            weight.data_ptr<float>(),
            nullptr,
            output.data_ptr<float>(),
            channels, out_channels, height, width, out_height, out_width);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_depthwise_cuda = load_inline(
    name="depthwise3x3_tiled_cuda_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 x.shape[1] == conv.in_channels
            and x.shape[2] >= 3
            and x.shape[3] >= 3
            and x.shape[0] * conv.out_channels <= 65535
            and (x.shape[2] - 2 + 15) // 16 <= 65535
            and (
                bias is None
                or (
                    bias.device == x.device
                    and bias.dtype == torch.float32
                    and bias.is_contiguous()
                )
            )
        ):
            return _depthwise_cuda.depthwise3x3_cuda(x, weight, bias)

        return conv(x)
```