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


_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 depthwise_3x3_kernel(
    const float* __restrict__ input,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int channels,
    int height,
    int width,
    int output_height,
    int output_width,
    int tiles_x,
    int tiles_y
) {
    const int tile_x = blockIdx.x % tiles_x;
    const int tile_y = (blockIdx.x / tiles_x) % tiles_y;
    const int64_t nc = blockIdx.x / (tiles_x * tiles_y);
    const int channel = nc % channels;

    const int x = tile_x * 512 + threadIdx.x * 4;
    const int y = tile_y * 4;
    if (x >= output_width) {
        return;
    }

    const float* plane = input + nc * height * width;
    float* out_plane = output + nc * output_height * output_width;

    float w[9];
    #pragma unroll
    for (int i = 0; i < 9; ++i) {
        w[i] = weight[channel * 9 + i];
    }

    float acc[4][4] = {};

    // Six input rows produce four output rows.
    #pragma unroll
    for (int iy = 0; iy < 6; ++iy) {
        float v[6] = {};
        if (y + iy < height) {
            const float* p = plane + int64_t(y + iy) * width + x;
            const float4 a = *reinterpret_cast<const float4*>(p);
            v[0] = a.x;
            v[1] = a.y;
            v[2] = a.z;
            v[3] = a.w;

            if (x + 4 < width) {
                const float2 b = *reinterpret_cast<const float2*>(p + 4);
                v[4] = b.x;
                v[5] = b.y;
            }
        }

        #pragma unroll
        for (int oy = 0; oy < 4; ++oy) {
            if (iy >= oy && iy < oy + 3) {
                const int k = (iy - oy) * 3;
                #pragma unroll
                for (int ox = 0; ox < 4; ++ox) {
                    acc[oy][ox] = fmaf(v[ox],     w[k],     acc[oy][ox]);
                    acc[oy][ox] = fmaf(v[ox + 1], w[k + 1], acc[oy][ox]);
                    acc[oy][ox] = fmaf(v[ox + 2], w[k + 2], acc[oy][ox]);
                }
            }
        }
    }

    const float b = bias ? bias[channel] : 0.0f;

    #pragma unroll
    for (int oy = 0; oy < 4; ++oy) {
        if (y + oy < output_height) {
            float* p = out_plane + int64_t(y + oy) * output_width + x;

            *reinterpret_cast<float2*>(p) =
                make_float2(acc[oy][0] + b, acc[oy][1] + b);

            if (x + 2 < output_width) {
                *reinterpret_cast<float2*>(p + 2) =
                    make_float2(acc[oy][2] + b, acc[oy][3] + b);
            }
        }
    }
}

torch::Tensor depthwise_3x3_cuda(
    torch::Tensor input,
    torch::Tensor weight,
    torch::Tensor bias
) {
    TORCH_CHECK(input.is_cuda() && weight.is_cuda(),
                "Expected CUDA input and weight");
    TORCH_CHECK(input.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat,
                "Expected float32 input and weight");
    TORCH_CHECK(input.dim() == 4 && input.is_contiguous(),
                "Expected contiguous NCHW input");
    TORCH_CHECK(weight.is_contiguous() && weight.dim() == 4 &&
                weight.size(0) == input.size(1) &&
                weight.size(1) == 1 &&
                weight.size(2) == 3 && weight.size(3) == 3,
                "Expected contiguous depthwise 3x3 weights");
    TORCH_CHECK(input.device() == weight.device(),
                "Input and weight must share a device");

    const c10::cuda::CUDAGuard device_guard(input.device());

    const int64_t batch = input.size(0);
    const int channels = input.size(1);
    const int height = input.size(2);
    const int width = input.size(3);

    TORCH_CHECK(height >= 3 && width >= 4 && width % 4 == 0,
                "Unsupported input dimensions");

    const float* bias_ptr = nullptr;
    if (bias.numel() != 0) {
        TORCH_CHECK(bias.device() == input.device() &&
                    bias.scalar_type() == at::kFloat &&
                    bias.is_contiguous() && bias.numel() == channels,
                    "Invalid bias");
        bias_ptr = bias.data_ptr<float>();
    }

    const int oh = height - 2;
    const int ow = width - 2;
    auto output = torch::empty({batch, channels, oh, ow}, input.options());
    if (output.numel() == 0) {
        return output;
    }

    const int tiles_x = (ow + 511) / 512;
    const int tiles_y = (oh + 3) / 4;
    const int64_t blocks = batch * channels * tiles_x * tiles_y;
    TORCH_CHECK(blocks <= 2147483647LL, "Launch is too large");

    depthwise_3x3_kernel<<<
        static_cast<unsigned int>(blocks),
        128,
        0,
        at::cuda::getCurrentCUDAStream()
    >>>(
        input.data_ptr<float>(),
        weight.data_ptr<float>(),
        bias_ptr,
        output.data_ptr<float>(),
        channels, height, width, oh, ow, tiles_x, tiles_y
    );
    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = load_inline(
    name="depthwise_3x3_tiled_cuda",
    cpp_sources=r"""
        torch::Tensor depthwise_3x3_cuda(
            torch::Tensor input,
            torch::Tensor weight,
            torch::Tensor bias);
    """,
    cuda_sources=_cuda_source,
    functions=["depthwise_3x3_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,
        )
        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.ndim == 4
            and x.is_contiguous()
            and x.shape[1] == conv.in_channels
            and x.shape[2] >= 3
            and x.shape[3] >= 4
            and x.shape[3] % 4 == 0
            and conv.in_channels == conv.out_channels
            and conv.kernel_size == (3, 3)
            and conv.stride == (1, 1)
            and conv.padding == (0, 0)
            and conv.weight.dtype == torch.float32
            and conv.weight.is_contiguous()
            and not torch.is_autocast_enabled()
        ):
            return _extension.depthwise_3x3_cuda(
                x,
                conv.weight,
                conv.bias if conv.bias is not None else self._empty_bias,
            )
        return conv(x)
```