```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 tiled_conv_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    torch::Tensor bias,
    bool has_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_fp16.h>
#include <cuda_runtime.h>
#include <cublas_v2.h>
#include <algorithm>
#include <cstdint>

static void check_blas(cublasStatus_t status) {
    TORCH_CHECK(status == CUBLAS_STATUS_SUCCESS,
                "cuBLAS operation failed: ", static_cast<int>(status));
}

// Represent each FP32 value with a leading FP16 value and an FP16 residual.
__global__ void split_weights(
    const float* __restrict__ weight,
    half* __restrict__ high,
    half* __restrict__ low,
    int64_t count)
{
    for (int64_t i = static_cast<int64_t>(blockIdx.x) * blockDim.x
                     + threadIdx.x;
         i < count;
         i += static_cast<int64_t>(blockDim.x) * gridDim.x) {
        float value = weight[i];
        half leading = __float2half_rn(value);
        high[i] = leading;
        low[i] = __float2half_rn(value - __half2float(leading));
    }
}

// Fuse dilation, padding, tiled im2col, and precision splitting.
__global__ void lower_input(
    const float* __restrict__ x,
    half* __restrict__ high,
    half* __restrict__ low,
    int batch,
    int channels,
    int height,
    int width,
    int kernel_h,
    int kernel_w,
    int output_w,
    int stride_h,
    int stride_w,
    int pad_h,
    int pad_w,
    int dilation_h,
    int dilation_w,
    int start,
    int columns,
    int reduction)
{
    int64_t count = static_cast<int64_t>(batch) * reduction * columns;

    for (int64_t i = static_cast<int64_t>(blockIdx.x) * blockDim.x
                     + threadIdx.x;
         i < count;
         i += static_cast<int64_t>(blockDim.x) * gridDim.x) {
        int column = static_cast<int>(i % columns);
        int64_t row = i / columns;
        int k = static_cast<int>(row % reduction);
        int n = static_cast<int>(row / reduction);

        int kw = k % kernel_w;
        int kh = (k / kernel_w) % kernel_h;
        int channel = k / (kernel_h * kernel_w);

        int position = start + column;
        int oh = position / output_w;
        int ow = position % output_w;
        int ih = oh * stride_h - pad_h + kh * dilation_h;
        int iw = ow * stride_w - pad_w + kw * dilation_w;

        float value = 0.0f;
        if (static_cast<unsigned>(ih) < static_cast<unsigned>(height) &&
            static_cast<unsigned>(iw) < static_cast<unsigned>(width)) {
            int64_t offset =
                ((static_cast<int64_t>(n) * channels + channel) * height
                 + ih) * width + iw;
            value = x[offset];
        }

        half leading = __float2half_rn(value);
        high[i] = leading;
        low[i] = __float2half_rn(value - __half2float(leading));
    }
}

__global__ void add_bias(
    float* __restrict__ output,
    const float* __restrict__ bias,
    int64_t count,
    int channels,
    int spatial)
{
    for (int64_t i = static_cast<int64_t>(blockIdx.x) * blockDim.x
                     + threadIdx.x;
         i < count;
         i += static_cast<int64_t>(blockDim.x) * gridDim.x) {
        int channel = static_cast<int>((i / spatial) % channels);
        output[i] += bias[channel];
    }
}

static int blocks_for(int64_t count) {
    return static_cast<int>(
        std::min<int64_t>((count + 255) / 256, 65535));
}

class MathModeGuard {
public:
    explicit MathModeGuard(cublasHandle_t handle) : handle_(handle) {
        check_blas(cublasGetMathMode(handle_, &previous_));
        check_blas(cublasSetMathMode(
            handle_,
            static_cast<cublasMath_t>(
                CUBLAS_DEFAULT_MATH |
                CUBLAS_MATH_DISALLOW_REDUCED_PRECISION_REDUCTION)));
    }

    ~MathModeGuard() {
        cublasSetMathMode(handle_, previous_);
    }

private:
    cublasHandle_t handle_;
    cublasMath_t previous_;
};

torch::Tensor tiled_conv_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    torch::Tensor bias,
    bool has_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 FP32 tensors");
    TORCH_CHECK(x.dim() == 4 && weight.dim() == 4,
                "Expected four-dimensional tensors");
    TORCH_CHECK(x.device() == weight.device(),
                "Input and weight must share a device");

    c10::cuda::CUDAGuard device_guard(x.device());

    x = x.contiguous();
    weight = weight.contiguous();

    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 kernel_h = static_cast<int>(weight.size(2));
    const int kernel_w = static_cast<int>(weight.size(3));

    TORCH_CHECK(weight.size(1) == channels, "Channel mismatch");

    const int64_t available_h =
        height + 2 * pad_h - dilation_h * (kernel_h - 1) - 1;
    const int64_t available_w =
        width + 2 * pad_w - dilation_w * (kernel_w - 1) - 1;

    TORCH_CHECK(available_h >= 0 && available_w >= 0,
                "Kernel exceeds padded input");

    const int output_h = static_cast<int>(available_h / stride_h + 1);
    const int output_w = static_cast<int>(available_w / stride_w + 1);
    const int spatial = output_h * output_w;
    const int reduction = channels * kernel_h * kernel_w;

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

    if (batch == 0) {
        return output;
    }

    // Keep temporary storage bounded independently of image dimensions.
    constexpr int64_t workspace_budget = 384LL * 1024 * 1024;
    int64_t tile_limit =
        workspace_budget / (4LL * batch * reduction);
    tile_limit = std::max<int64_t>(256, (tile_limit / 256) * 256);
    const int tile = static_cast<int>(
        std::min<int64_t>(spatial, std::min<int64_t>(16384, tile_limit)));

    auto half_options = x.options().dtype(at::kHalf);
    auto packed_weights = torch::empty(
        {2, out_channels, reduction}, half_options);
    auto workspace = torch::empty(
        {2, batch, reduction, tile}, half_options);

    const int64_t weight_count =
        static_cast<int64_t>(out_channels) * reduction;
    const int64_t workspace_plane =
        static_cast<int64_t>(batch) * reduction * tile;

    half* weight_high =
        reinterpret_cast<half*>(packed_weights.data_ptr<at::Half>());
    half* weight_low = weight_high + weight_count;
    half* input_high =
        reinterpret_cast<half*>(workspace.data_ptr<at::Half>());
    half* input_low = input_high + workspace_plane;

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

    split_weights<<<blocks_for(weight_count), 256, 0, stream>>>(
        weight.data_ptr<float>(), weight_high, weight_low, weight_count);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    cublasHandle_t handle = at::cuda::getCurrentCUDABlasHandle();
    check_blas(cublasSetStream(handle, stream));
    MathModeGuard math_guard(handle);

    const float alpha = 1.0f;
    const float zero = 0.0f;
    const float one = 1.0f;
    const long long output_batch_stride =
        static_cast<long long>(out_channels) * spatial;

    for (int start = 0; start < spatial; start += tile) {
        const int columns = std::min(tile, spatial - start);
        const int64_t count =
            static_cast<int64_t>(batch) * reduction * columns;

        lower_input<<<blocks_for(count), 256, 0, stream>>>(
            x.data_ptr<float>(), input_high, input_low,
            batch, channels, height, width,
            kernel_h, kernel_w, output_w,
            static_cast<int>(stride_h), static_cast<int>(stride_w),
            static_cast<int>(pad_h), static_cast<int>(pad_w),
            static_cast<int>(dilation_h), static_cast<int>(dilation_w),
            start, columns, reduction);
        C10_CUDA_KERNEL_LAUNCH_CHECK();

        const long long input_batch_stride =
            static_cast<long long>(reduction) * columns;
        float* destination = output.data_ptr<float>() + start;

        // Column-major interpretation writes directly into NCHW output.
        auto gemm = [&](const half* a, const half* b, const float* beta) {
            check_blas(cublasGemmStridedBatchedEx(
                handle,
                CUBLAS_OP_N, CUBLAS_OP_N,
                columns, out_channels, reduction,
                &alpha,
                a, CUDA_R_16F, columns, input_batch_stride,
                b, CUDA_R_16F, reduction, 0,
                beta,
                destination, CUDA_R_32F, spatial, output_batch_stride,
                batch,
                CUBLAS_COMPUTE_32F,
                CUBLAS_GEMM_DEFAULT_TENSOR_OP));
        };

        // FP32 accumulation with first-order residual correction.
        gemm(input_high, weight_high, &zero);
        gemm(input_low, weight_high, &one);
        gemm(input_high, weight_low, &one);
    }

    if (has_bias) {
        TORCH_CHECK(bias.is_cuda() &&
                    bias.device() == x.device() &&
                    bias.scalar_type() == at::kFloat &&
                    bias.numel() == out_channels,
                    "Invalid bias");
        bias = bias.contiguous();
        add_bias<<<blocks_for(output.numel()), 256, 0, stream>>>(
            output.data_ptr<float>(), bias.data_ptr<float>(),
            output.numel(), out_channels, spatial);
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }

    return output;
}
"""


_conv_extension = load_inline(
    name="asymmetric_dilated_conv_split_gemm_v1",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["tiled_conv_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    extra_ldflags=["-lcublas"],
    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.is_autocast_enabled()
        ):
            return conv(x)

        return _conv_extension.tiled_conv_cuda(
            x,
            conv.weight,
            conv.bias if conv.bias is not None else conv.weight,
            conv.bias is not None,
            conv.stride[0],
            conv.stride[1],
            conv.padding[0],
            conv.padding[1],
            conv.dilation[0],
            conv.dilation[1],
        )
```