```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 dilated_conv_cuda(
    torch::Tensor input,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias,
    int64_t padding_h,
    int64_t padding_w);
"""

_CUDA_SOURCE = r"""
#include <torch/extension.h>
#include <ATen/Context.h>
#include <ATen/ops/cudnn_convolution.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>

// Load all three horizontal phases together so input reads are contiguous.
__global__ __launch_bounds__(256) 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][97];

    const int lane = threadIdx.x;
    const int row = threadIdx.y;
    const int batch = blockIdx.z / 2;
    const int phase_h = blockIdx.z % 2;
    const int spatial = packed_h * packed_w;
    const int first_position = blockIdx.x * 32;
    const int first_channel = blockIdx.y * 32;

    #pragma unroll
    for (int segment = 0; segment < 3; ++segment) {
        const int s = segment * 32 + lane;
        const int position = first_position + s / 3;
        const int iy = (position / packed_w) * 2 + phase_h - padding_h;
        const int ix = (position % packed_w) * 3 + s % 3 - padding_w;

        const bool valid =
            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 && channel < channels) {
                const int64_t offset =
                    ((static_cast<int64_t>(batch) * channels + channel)
                     * height + iy) * width + ix;
                value = input[offset];
            }
            tile[r][s] = value;
        }
    }

    __syncthreads();

    const int channel = first_channel + lane;

    #pragma unroll
    for (int phase_w = 0; phase_w < 3; ++phase_w) {
        const int packed_batch = batch * 6 + phase_h * 3 + phase_w;
        const int64_t base =
            static_cast<int64_t>(packed_batch) * spatial * channels;

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

// Process four output rows per block, amortizing indexing and block overhead.
template<int FixedChannels, bool HasBias>
__global__ __launch_bounds__(256) void unpack4_kernel(
    const float4* __restrict__ input,
    const float4* __restrict__ bias,
    float4* __restrict__ output,
    int channels,
    int packed_h,
    int packed_w,
    int output_h,
    int output_w)
{
    constexpr int Rows = 4;
    const int channel_vectors =
        FixedChannels != 0 ? FixedChannels / 4 : channels / 4;
    const int t = blockIdx.x * blockDim.x + threadIdx.x;
    const int ox = t / channel_vectors;

    if (ox >= output_w)
        return;

    const int cv = t % channel_vectors;
    const int first_y = blockIdx.y * Rows;
    const int batch = blockIdx.z;
    const int phase_w = ox % 3;
    const int packed_x = ox / 3;

    float4 values[Rows];
    float4 b;
    if (HasBias)
        b = bias[cv];

    #pragma unroll
    for (int r = 0; r < Rows; ++r) {
        const int oy = first_y + r;
        if (oy < output_h) {
            const int phase = (oy % 2) * 3 + phase_w;
            const int64_t src =
                (((static_cast<int64_t>(batch) * 6 + phase) * packed_h
                  + oy / 2) * packed_w + packed_x) * channel_vectors + cv;

            float4 value = input[src];
            if (HasBias) {
                value.x += b.x;
                value.y += b.y;
                value.z += b.z;
                value.w += b.w;
            }
            values[r] = value;
        }
    }

    #pragma unroll
    for (int r = 0; r < Rows; ++r) {
        const int oy = first_y + r;
        if (oy < output_h) {
            const int64_t dst =
                ((static_cast<int64_t>(batch) * output_h + oy)
                 * output_w + ox) * channel_vectors + cv;
            output[dst] = values[r];
        }
    }
}

template<bool HasBias>
__global__ void unpack_scalar_kernel(
    const float* __restrict__ input,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int channels,
    int packed_h,
    int packed_w,
    int output_h,
    int output_w)
{
    const int t = blockIdx.x * blockDim.x + threadIdx.x;
    const int ox = t / channels;
    if (ox >= output_w)
        return;

    const int channel = t % channels;
    const int oy = blockIdx.y;
    const int batch = blockIdx.z;
    const int phase = (oy % 2) * 3 + ox % 3;

    const int64_t src =
        (((static_cast<int64_t>(batch) * 6 + phase) * packed_h
          + oy / 2) * packed_w + ox / 3) * channels + channel;

    float value = input[src];
    if (HasBias)
        value += bias[channel];

    const int64_t dst =
        ((static_cast<int64_t>(batch) * output_h + oy)
         * output_w + ox) * channels + channel;
    output[dst] = value;
}

template<bool HasBias>
void launch_unpack(
    const torch::Tensor& packed,
    const float* bias,
    torch::Tensor& output,
    cudaStream_t stream)
{
    const int batch = output.size(0);
    const int channels = output.size(1);
    const int output_h = output.size(2);
    const int output_w = output.size(3);
    const int packed_h = packed.size(2);
    const int packed_w = packed.size(3);
    constexpr int threads = 256;

    if (channels % 4 == 0) {
        const dim3 grid(
            (static_cast<int64_t>(output_w) * (channels / 4)
             + threads - 1) / threads,
            (output_h + 3) / 4,
            batch);

        const auto* src =
            reinterpret_cast<const float4*>(packed.data_ptr<float>());
        const auto* b = reinterpret_cast<const float4*>(bias);
        auto* dst = reinterpret_cast<float4*>(output.data_ptr<float>());

        if (channels == 64) {
            unpack4_kernel<64, HasBias><<<grid, threads, 0, stream>>>(
                src, b, dst, channels,
                packed_h, packed_w, output_h, output_w);
        } else {
            unpack4_kernel<0, HasBias><<<grid, threads, 0, stream>>>(
                src, b, dst, channels,
                packed_h, packed_w, output_h, output_w);
        }
    } else {
        const dim3 grid(
            (static_cast<int64_t>(output_w) * channels
             + threads - 1) / threads,
            output_h,
            batch);

        unpack_scalar_kernel<HasBias><<<grid, threads, 0, stream>>>(
            packed.data_ptr<float>(), bias, output.data_ptr<float>(),
            channels, packed_h, packed_w, output_h, output_w);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
}

torch::Tensor dilated_conv_cuda(
    torch::Tensor input,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias,
    int64_t padding_h,
    int64_t padding_w)
{
    TORCH_CHECK(input.is_cuda() && weight.is_cuda(),
                "Expected CUDA tensors");
    TORCH_CHECK(input.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat,
                "Expected float32 tensors");
    TORCH_CHECK(input.dim() == 4 && input.is_contiguous(),
                "Expected contiguous NCHW input");
    TORCH_CHECK(input.device() == weight.device(),
                "Input and weight must share a device");

    const c10::cuda::CUDAGuard guard(input.device());
    const auto stream = at::cuda::getCurrentCUDAStream();

    const int batch = input.size(0);
    const int channels = input.size(1);
    const int height = input.size(2);
    const int width = input.size(3);
    const int out_channels = weight.size(0);
    const int kernel_h = weight.size(2);
    const int kernel_w = weight.size(3);

    const int packed_h = (height + 2 * padding_h + 1) / 2;
    const int packed_w = (width + 2 * padding_w + 2) / 3;
    const int output_h = height + 2 * padding_h - 2 * (kernel_h - 1);
    const int output_w = width + 2 * padding_w - 3 * (kernel_w - 1);

    TORCH_CHECK(output_h > 0 && output_w > 0,
                "Kernel exceeds padded input size");

    auto options =
        input.options().memory_format(at::MemoryFormat::ChannelsLast);

    auto packed = torch::empty(
        {static_cast<int64_t>(batch) * 6, channels, packed_h, packed_w},
        options);

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

    pack_kernel<<<grid, block, 0, stream>>>(
        input.data_ptr<float>(),
        packed.data_ptr<float>(),
        channels, height, width, packed_h, packed_w,
        padding_h, padding_w);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    auto dense = at::cudnn_convolution(
        packed,
        weight,
        {0, 0},
        {1, 1},
        {1, 1},
        1,
        true,
        at::globalContext().deterministicCuDNN(),
        at::globalContext().allowTF32CuDNN());

    dense = dense.contiguous(at::MemoryFormat::ChannelsLast);

    auto output = torch::empty(
        {batch, out_channels, output_h, output_w}, options);

    if (bias.has_value() && bias->defined()) {
        TORCH_CHECK(
            bias->device() == input.device() &&
            bias->scalar_type() == at::kFloat &&
            bias->numel() == out_channels,
            "Invalid bias");
        auto contiguous_bias = bias->contiguous();
        launch_unpack<true>(
            dense, contiguous_bias.data_ptr<float>(), output, stream);
    } else {
        launch_unpack<false>(dense, nullptr, output, stream);
    }

    return output;
}
"""

_EXTENSION = None


def _get_extension():
    global _EXTENSION
    if _EXTENSION is None:
        _EXTENSION = load_inline(
            name="dilated_conv_space_to_batch_v6",
            cpp_sources=_CPP_SOURCE,
            cuda_sources=_CUDA_SOURCE,
            functions=["dilated_conv_cuda"],
            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 not torch.backends.cudnn.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 (
            not 0 < output_h <= 65535
            or output_w <= 0
            or output_w * conv.out_channels >= 2**31
            or (x.shape[2] + 2 * ph) * (x.shape[3] + 2 * pw) >= 2**31
        ):
            return conv(x)

        return _get_extension().dilated_conv_cuda(
            x.contiguous(),
            conv.weight,
            conv.bias,
            ph,
            pw,
        )
```