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


_CPP_SOURCE = r"""
#include <torch/extension.h>

torch::Tensor pack_dilated_input(
    torch::Tensor x, int64_t padding_h, int64_t padding_w);

torch::Tensor unpack_dilated_output(
    torch::Tensor x, int64_t batch, int64_t output_h, int64_t output_w);
"""

_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>

// Split the two spatial dilation dimensions into six batch entries.
// Shared memory simultaneously converts NCHW into channels-last storage.
__global__ void pack_kernel(
    const float* __restrict__ input,
    float* __restrict__ output,
    int channels, int height, int width,
    int packed_h, int packed_w,
    int padding_h, int padding_w)
{
    __shared__ float tile[32][33];

    const int lane = threadIdx.x;
    const int row = threadIdx.y;
    const int packed_batch = blockIdx.z;
    const int batch = packed_batch / 6;
    const int phase_h = (packed_batch % 6) / 3;
    const int phase_w = packed_batch % 3;

    const int spatial = packed_h * packed_w;
    const int first_position = blockIdx.x * 32;
    const int first_channel = blockIdx.y * 32;
    const int position = first_position + lane;

    int iy = 0;
    int ix = 0;
    if (position < spatial) {
        iy = (position / packed_w) * 2 + phase_h - padding_h;
        ix = (position % packed_w) * 3 + phase_w - padding_w;
    }

    const bool valid_position =
        position < spatial &&
        static_cast<unsigned>(iy) < static_cast<unsigned>(height) &&
        static_cast<unsigned>(ix) < static_cast<unsigned>(width);

    #pragma unroll
    for (int r = row; r < 32; r += 8) {
        const int channel = first_channel + r;
        float value = 0.0f;
        if (valid_position && channel < channels) {
            const int64_t offset =
                ((static_cast<int64_t>(batch) * channels + channel)
                 * height + iy) * width + ix;
            value = input[offset];
        }
        tile[r][lane] = value;
    }

    __syncthreads();

    const int channel = first_channel + lane;
    #pragma unroll
    for (int r = row; r < 32; r += 8) {
        const int p = first_position + r;
        if (p < spatial && channel < channels) {
            const int64_t offset =
                (static_cast<int64_t>(packed_batch) * spatial + p)
                * channels + channel;
            output[offset] = tile[lane][r];
        }
    }
}

// Interleave the six convolution results, crop excess right/bottom
// positions, and produce contiguous NCHW output in one pass.
__global__ void unpack_kernel(
    const float* __restrict__ input,
    float* __restrict__ output,
    int channels,
    int packed_h, int packed_w,
    int output_h, int output_w)
{
    __shared__ float tile[32][33];

    const int lane = threadIdx.x;
    const int row = threadIdx.y;
    const int batch = blockIdx.z;
    const int spatial = output_h * output_w;
    const int first_position = blockIdx.x * 32;
    const int first_channel = blockIdx.y * 32;
    const int channel = first_channel + lane;

    #pragma unroll
    for (int r = row; r < 32; r += 8) {
        const int position = first_position + r;
        float value = 0.0f;
        if (position < spatial && channel < channels) {
            const int oy = position / output_w;
            const int ox = position % output_w;
            const int phase = (oy % 2) * 3 + ox % 3;
            const int64_t packed_position =
                (static_cast<int64_t>(batch * 6 + phase) * packed_h
                 + oy / 2) * packed_w + ox / 3;
            value = input[packed_position * channels + channel];
        }
        tile[r][lane] = value;
    }

    __syncthreads();

    const int position = first_position + lane;
    #pragma unroll
    for (int r = row; r < 32; r += 8) {
        const int c = first_channel + r;
        if (position < spatial && c < channels) {
            const int64_t offset =
                (static_cast<int64_t>(batch) * channels + c)
                * spatial + position;
            output[offset] = tile[lane][r];
        }
    }
}

torch::Tensor pack_dilated_input(
    torch::Tensor x, int64_t padding_h, int64_t padding_w)
{
    TORCH_CHECK(x.is_cuda() && x.scalar_type() == at::kFloat,
                "Expected a CUDA float32 tensor");
    TORCH_CHECK(x.dim() == 4 && x.is_contiguous(),
                "Expected contiguous NCHW input");

    const c10::cuda::CUDAGuard 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 packed_h = (height + 2 * padding_h + 1) / 2;
    const int packed_w = (width + 2 * padding_w + 2) / 3;

    auto output = torch::empty(
        {static_cast<int64_t>(batch) * 6, channels, packed_h, packed_w},
        x.options().memory_format(at::MemoryFormat::ChannelsLast));

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

    const dim3 block(32, 8);
    const dim3 grid(
        (static_cast<int64_t>(packed_h) * packed_w + 31) / 32,
        (channels + 31) / 32,
        batch * 6);

    pack_kernel<<<grid, block, 0, at::cuda::getCurrentCUDAStream()>>>(
        x.data_ptr<float>(), output.data_ptr<float>(),
        channels, height, width, packed_h, packed_w,
        padding_h, padding_w);
    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}

torch::Tensor unpack_dilated_output(
    torch::Tensor x, int64_t batch, int64_t output_h, int64_t output_w)
{
    TORCH_CHECK(x.is_cuda() && x.scalar_type() == at::kFloat,
                "Expected a CUDA float32 tensor");
    TORCH_CHECK(x.dim() == 4 &&
                x.is_contiguous(at::MemoryFormat::ChannelsLast),
                "Expected channels-last input");
    TORCH_CHECK(x.size(0) == batch * 6, "Invalid packed batch size");

    const c10::cuda::CUDAGuard guard(x.device());
    const int channels = x.size(1);
    const int packed_h = x.size(2);
    const int packed_w = x.size(3);

    auto output = torch::empty(
        {batch, static_cast<int64_t>(channels), output_h, output_w},
        x.options().memory_format(at::MemoryFormat::Contiguous));

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

    const dim3 block(32, 8);
    const dim3 grid(
        (output_h * output_w + 31) / 32,
        (channels + 31) / 32,
        batch);

    unpack_kernel<<<grid, block, 0, at::cuda::getCurrentCUDAStream()>>>(
        x.data_ptr<float>(), output.data_ptr<float>(),
        channels, packed_h, packed_w, output_h, output_w);
    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_EXTENSION = None


def _get_extension():
    global _EXTENSION
    if _EXTENSION is None:
        _EXTENSION = load_inline(
            name="dilated_conv_space_to_batch_v4",
            cpp_sources=_CPP_SOURCE,
            cuda_sources=_CUDA_SOURCE,
            functions=["pack_dilated_input", "unpack_dilated_output"],
            extra_cflags=["-O3"],
            extra_cuda_cflags=["-O3"],
            verbose=False,
        )
    return _EXTENSION


class ModelNew(nn.Module):
    def __init__(
        self,
        in_channels: int,
        out_channels: int,
        kernel_size: tuple,
        stride: int = 1,
        padding: tuple = (0, 0),
        dilation: tuple = (1, 1),
        bias: bool = False,
    ):
        super().__init__()
        self.conv2d = nn.Conv2d(
            in_channels,
            out_channels,
            kernel_size,
            stride=stride,
            padding=padding,
            dilation=dilation,
            bias=bias,
        )
        self.conv2d.to(memory_format=torch.channels_last)

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        conv = self.conv2d
        if (
            not x.is_cuda
            or x.ndim != 4
            or x.dtype != torch.float32
            or conv.weight.dtype != torch.float32
            or x.device != conv.weight.device
            or torch.is_autocast_enabled()
            or conv.stride != (1, 1)
            or conv.dilation != (2, 3)
            or isinstance(conv.padding, str)
            or x.shape[1] != conv.in_channels
            or not 0 < x.shape[0] <= 10922
            or max(conv.in_channels, conv.out_channels) > 65535 * 32
        ):
            return conv(x)

        ph, pw = conv.padding
        kh, kw = conv.kernel_size
        output_h = x.shape[2] + 2 * ph - 2 * (kh - 1)
        output_w = x.shape[3] + 2 * pw - 3 * (kw - 1)

        if output_h <= 0 or output_w <= 0:
            return conv(x)

        extension = _get_extension()
        packed = extension.pack_dilated_input(x.contiguous(), ph, pw)

        # Dilation becomes an ordinary dense convolution over six phases.
        result = F.conv2d(
            packed,
            conv.weight,
            conv.bias,
            stride=1,
            padding=0,
            dilation=1,
        )

        return extension.unpack_dilated_output(
            result.contiguous(memory_format=torch.channels_last),
            x.shape[0],
            output_h,
            output_w,
        )
```