```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 dilated_conv_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias,
    int64_t stride_h, int64_t stride_w,
    int64_t pad_h, int64_t pad_w,
    int64_t dilation_h, int64_t dilation_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.h>
#include <cuda_runtime.h>
#include <mma.h>

namespace wmma = nvcuda::wmma;

template<int KH, int KW, int DH, int DW>
__global__ void implicit_conv_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,
    int kernel_h, int kernel_w,
    int stride_h, int stride_w,
    int pad_h, int pad_w,
    int dilation_h, int dilation_w)
{
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
    constexpr int BM = 64;
    constexpr int BN = 64;
    constexpr int BK = 32;

    const int kh_size = KH ? KH : kernel_h;
    const int kw_size = KW ? KW : kernel_w;
    const int dh = DH ? DH : dilation_h;
    const int dw = DW ? DW : dilation_w;

    const int reduction_size = channels * kh_size * kw_size;
    const int spatial_size = out_height * out_width;
    const int batch = blockIdx.z;
    const int first_channel = blockIdx.y * BM;
    const int first_spatial = blockIdx.x * BN;
    const int tid = threadIdx.x;
    const int warp = tid >> 5;
    const int warp_row = (warp >> 1) * 16;
    const int warp_col = (warp & 1) * 32;

    __shared__ __align__(32) float storage[8192];
    float* a_hi = storage;
    float* a_lo = storage + 2048;
    float* b_hi = storage + 4096;
    float* b_lo = storage + 6144;

    wmma::fragment<wmma::accumulator, 16, 16, 8, float> acc0;
    wmma::fragment<wmma::accumulator, 16, 16, 8, float> acc1;
    wmma::fill_fragment(acc0, 0.0f);
    wmma::fill_fragment(acc1, 0.0f);

    // Each thread loads one spatial column across multiple reduction rows.
    const int spatial = first_spatial + (tid & 63);
    const int input_y_base = (spatial / out_width) * stride_h - pad_h;
    const int input_x_base = (spatial % out_width) * stride_w - pad_w;

    for (int k_base = 0; k_base < reduction_size; k_base += BK) {
        #pragma unroll
        for (int i = tid; i < BM * BK; i += 256) {
            const int oc = first_channel + i / BK;
            const int k = k_base + i % BK;

            float a = 0.0f;
            if (oc < out_channels && k < reduction_size) {
                a = weight[static_cast<int64_t>(oc) * reduction_size + k];
            }

            const float ah = wmma::__float_to_tf32(a);
            a_hi[i] = ah;
            a_lo[i] = wmma::__float_to_tf32(a - ah);
        }

        #pragma unroll
        for (int i = tid; i < BK * BN; i += 256) {
            const int k = k_base + i / BN;
            float b = 0.0f;

            if (spatial < spatial_size && k < reduction_size) {
                const int ic = k / (kh_size * kw_size);
                const int kernel_position = k % (kh_size * kw_size);
                const int ky = kernel_position / kw_size;
                const int kx = kernel_position % kw_size;
                const int iy = input_y_base + ky * dh;
                const int ix = input_x_base + kx * dw;

                if (static_cast<unsigned>(iy) < static_cast<unsigned>(height) &&
                    static_cast<unsigned>(ix) < static_cast<unsigned>(width)) {
                    const int64_t offset =
                        ((static_cast<int64_t>(batch) * channels + ic) *
                         height + iy) * width + ix;
                    b = x[offset];
                }
            }

            const float bh = wmma::__float_to_tf32(b);
            b_hi[i] = bh;
            b_lo[i] = wmma::__float_to_tf32(b - bh);
        }

        __syncthreads();

        #pragma unroll
        for (int kk = 0; kk < BK; kk += 8) {
            wmma::fragment<
                wmma::matrix_a, 16, 16, 8,
                wmma::precision::tf32, wmma::row_major> ah, al;
            wmma::fragment<
                wmma::matrix_b, 16, 16, 8,
                wmma::precision::tf32, wmma::row_major> bh, bl;

            wmma::load_matrix_sync(ah, a_hi + warp_row * BK + kk, BK);
            wmma::load_matrix_sync(al, a_lo + warp_row * BK + kk, BK);

            wmma::load_matrix_sync(bh, b_hi + kk * BN + warp_col, BN);
            wmma::load_matrix_sync(bl, b_lo + kk * BN + warp_col, BN);

            // Three TF32 products retain substantially more FP32 precision
            // than a single TF32 multiplication.
            wmma::mma_sync(acc0, al, bh, acc0);
            wmma::mma_sync(acc0, ah, bl, acc0);
            wmma::mma_sync(acc0, ah, bh, acc0);

            wmma::load_matrix_sync(bh, b_hi + kk * BN + warp_col + 16, BN);
            wmma::load_matrix_sync(bl, b_lo + kk * BN + warp_col + 16, BN);

            wmma::mma_sync(acc1, al, bh, acc1);
            wmma::mma_sync(acc1, ah, bl, acc1);
            wmma::mma_sync(acc1, ah, bh, acc1);
        }

        __syncthreads();
    }

    // Reuse the operand storage for coalesced output writes.
    wmma::store_matrix_sync(
        storage + warp_row * BN + warp_col,
        acc0, BN, wmma::mem_row_major);
    wmma::store_matrix_sync(
        storage + warp_row * BN + warp_col + 16,
        acc1, BN, wmma::mem_row_major);

    __syncthreads();

    #pragma unroll
    for (int i = tid; i < BM * BN; i += 256) {
        const int oc = first_channel + i / BN;
        const int p = first_spatial + i % BN;
        if (oc < out_channels && p < spatial_size) {
            float value = storage[i];
            if (bias != nullptr) {
                value += bias[oc];
            }
            output[
                (static_cast<int64_t>(batch) * out_channels + oc) *
                spatial_size + p] = value;
        }
    }
#endif
}

torch::Tensor dilated_conv_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias,
    int64_t stride_h, int64_t stride_w,
    int64_t pad_h, int64_t pad_w,
    int64_t dilation_h, int64_t dilation_w)
{
    TORCH_CHECK(x.is_cuda() && weight.is_cuda(), "Expected CUDA tensors");
    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat, "Expected float32 tensors");
    TORCH_CHECK(x.is_contiguous() && weight.is_contiguous(),
                "Expected contiguous tensors");
    TORCH_CHECK(x.dim() == 4 && weight.dim() == 4, "Expected 4D tensors");
    TORCH_CHECK(x.device() == weight.device(), "Device mismatch");
    TORCH_CHECK(x.size(1) == weight.size(1), "Channel mismatch");

    const c10::cuda::CUDAGuard guard(x.device());
    TORCH_CHECK(at::cuda::getCurrentDeviceProperties()->major >= 8,
                "This kernel requires compute capability 8.0 or newer");

    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 out_channels = weight.size(0);
    const int kh = weight.size(2);
    const int kw = weight.size(3);

    const int effective_h = dilation_h * (kh - 1) + 1;
    const int effective_w = dilation_w * (kw - 1) + 1;
    TORCH_CHECK(height + 2 * pad_h >= effective_h &&
                width + 2 * pad_w >= effective_w, "Kernel exceeds padded input");

    const int oh = (height + 2 * pad_h - effective_h) / stride_h + 1;
    const int ow = (width + 2 * pad_w - effective_w) / stride_w + 1;
    auto output = torch::empty({batch, out_channels, oh, ow}, x.options());

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

    const float* bias_ptr = nullptr;
    if (bias.has_value() && bias->defined()) {
        TORCH_CHECK(bias->device() == x.device() &&
                    bias->scalar_type() == at::kFloat &&
                    bias->is_contiguous() &&
                    bias->numel() == out_channels, "Invalid bias");
        bias_ptr = bias->data_ptr<float>();
    }

    const dim3 grid(
        (oh * ow + 63) / 64,
        (out_channels + 63) / 64,
        batch);
    const cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if (kh == 5 && kw == 9 && dilation_h == 2 && dilation_w == 3) {
        implicit_conv_kernel<5, 9, 2, 3><<<grid, 256, 0, stream>>>(
            x.data_ptr<float>(), weight.data_ptr<float>(), bias_ptr,
            output.data_ptr<float>(),
            channels, out_channels, height, width, oh, ow, kh, kw,
            stride_h, stride_w, pad_h, pad_w, dilation_h, dilation_w);
    } else {
        implicit_conv_kernel<0, 0, 0, 0><<<grid, 256, 0, stream>>>(
            x.data_ptr<float>(), weight.data_ptr<float>(), bias_ptr,
            output.data_ptr<float>(),
            channels, out_channels, height, width, oh, ow, kh, kw,
            stride_h, stride_w, pad_h, pad_w, dilation_h, dilation_w);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = load_inline(
    name="dilated_asymmetric_conv_tf32x3",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["dilated_conv_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: 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,
        )

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        conv = self.conv2d
        if (
            not x.is_cuda
            or x.dtype != torch.float32
            or conv.weight.dtype != torch.float32
            or x.ndim != 4
            or torch.cuda.get_device_capability(x.device)[0] < 8
            or torch.is_autocast_enabled()
            or isinstance(conv.padding, str)
        ):
            return F.conv2d(
                x,
                conv.weight,
                conv.bias,
                conv.stride,
                conv.padding,
                conv.dilation,
                conv.groups,
            )

        return _extension.dilated_conv_cuda(
            x.contiguous(),
            conv.weight.contiguous(),
            conv.bias,
            conv.stride[0],
            conv.stride[1],
            conv.padding[0],
            conv.padding[1],
            conv.dilation[0],
            conv.dilation[1],
        )
```