```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 asymmetric_conv_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias,
    int64_t sh, int64_t sw,
    int64_t ph, int64_t pw,
    int64_t dh, int64_t dw);
"""

_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 <cuda_fp16.h>
#include <cublas_v2.h>
#include <algorithm>
#include <climits>

namespace {

constexpr float RESIDUAL_SCALE = 4096.0f;
constexpr float INV_RESIDUAL_SCALE = 1.0f / RESIDUAL_SCALE;

__device__ __forceinline__ void split_float(
    float value, __half& high, __half& low) {
    high = __float2half_rn(value);
    low = __float2half_rn(
        (value - __half2float(high)) * RESIDUAL_SCALE);
}

__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) {
        __half h, l;
        split_float(weight[i], h, l);
        high[i] = h;
        low[i] = l;
    }
}

__global__ void im2col_split(
    const float* __restrict__ x,
    __half* __restrict__ high,
    __half* __restrict__ low,
    int H, int W,
    int KH, int KW,
    int OW, int K,
    int start, int positions,
    int sh, int sw,
    int ph, int pw,
    int dh, int dw) {
    const int64_t count = static_cast<int64_t>(K) * positions;

    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) {
        const int k = static_cast<int>(i / positions);
        const int p = static_cast<int>(i - static_cast<int64_t>(k) * positions)
                      + start;

        const int kw = k % KW;
        const int t = k / KW;
        const int kh = t % KH;
        const int channel = t / KH;

        const int oh = p / OW;
        const int ow = p - oh * OW;
        const int ih = oh * sh - ph + kh * dh;
        const int iw = ow * sw - pw + kw * dw;

        float value = 0.0f;
        if (static_cast<unsigned>(ih) < static_cast<unsigned>(H) &&
            static_cast<unsigned>(iw) < static_cast<unsigned>(W)) {
            value = x[(static_cast<int64_t>(channel) * H + ih) * W + iw];
        }

        __half h, l;
        split_float(value, h, l);
        high[i] = h;
        low[i] = l;
    }
}

__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) {
        output[i] += bias[(i / spatial) % channels];
    }
}

inline void check_blas(cublasStatus_t status) {
    TORCH_CHECK(
        status == CUBLAS_STATUS_SUCCESS,
        "cuBLAS GEMM failed with status ", static_cast<int>(status));
}

inline void gemm(
    cublasHandle_t handle,
    const __half* columns,
    const __half* weights,
    float* output,
    int positions,
    int channels,
    int K,
    int output_stride,
    float alpha,
    float beta) {
    check_blas(cublasGemmEx(
        handle,
        CUBLAS_OP_N,
        CUBLAS_OP_N,
        positions,
        channels,
        K,
        &alpha,
        columns,
        CUDA_R_16F,
        positions,
        weights,
        CUDA_R_16F,
        K,
        &beta,
        output,
        CUDA_R_32F,
        output_stride,
        CUBLAS_COMPUTE_32F,
        CUBLAS_GEMM_DEFAULT_TENSOR_OP));
}

}  // namespace

torch::Tensor asymmetric_conv_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias,
    int64_t sh, int64_t sw,
    int64_t ph, int64_t pw,
    int64_t dh, int64_t dw) {
    TORCH_CHECK(x.is_cuda() && weight.is_cuda(), "Expected CUDA tensors");
    TORCH_CHECK(
        x.scalar_type() == torch::kFloat32 &&
        weight.scalar_type() == torch::kFloat32,
        "Expected float32 tensors");
    TORCH_CHECK(x.dim() == 4 && weight.dim() == 4, "Expected 4D tensors");
    TORCH_CHECK(x.device() == weight.device(), "Device mismatch");
    TORCH_CHECK(x.is_contiguous() && weight.is_contiguous(),
                "Expected contiguous tensors");
    TORCH_CHECK(x.size(1) == weight.size(1), "Channel mismatch");
    TORCH_CHECK(sh > 0 && sw > 0 && dh > 0 && dw > 0,
                "Stride and dilation must be positive");
    TORCH_CHECK(ph >= 0 && pw >= 0, "Padding must be nonnegative");

    const c10::cuda::CUDAGuard guard(x.device());

    const int64_t B = x.size(0);
    const int64_t C = x.size(1);
    const int64_t H = x.size(2);
    const int64_t W = x.size(3);
    const int64_t OC = weight.size(0);
    const int64_t KH = weight.size(2);
    const int64_t KW = weight.size(3);

    const int64_t nh = H + 2 * ph - dh * (KH - 1) - 1;
    const int64_t nw = W + 2 * pw - dw * (KW - 1) - 1;
    TORCH_CHECK(nh >= 0 && nw >= 0, "Kernel exceeds padded input");

    const int64_t OH = nh / sh + 1;
    const int64_t OW = nw / sw + 1;
    const int64_t spatial64 = OH * OW;
    const int64_t K64 = C * KH * KW;

    TORCH_CHECK(
        H <= INT_MAX && W <= INT_MAX &&
        OC <= INT_MAX && spatial64 <= INT_MAX && K64 <= INT_MAX,
        "Tensor dimensions exceed supported indexing limits");

    auto output = torch::empty({B, OC, OH, OW}, x.options());
    if (output.numel() == 0) {
        return output;
    }

    const int spatial = static_cast<int>(spatial64);
    const int K = static_cast<int>(K64);
    const int chunk = std::min(spatial, 65536);

    auto half_options = x.options().dtype(torch::kFloat16);
    auto packed_weights = torch::empty({2, OC, K64}, half_options);
    auto columns = torch::empty({2, K64, chunk}, half_options);

    auto* wh = reinterpret_cast<__half*>(
        packed_weights.data_ptr<at::Half>());
    auto* wl = wh + OC * K64;
    auto* ch = reinterpret_cast<__half*>(columns.data_ptr<at::Half>());
    auto* cl = ch + K64 * chunk;

    const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    const cublasHandle_t handle = at::cuda::getCurrentCUDABlasHandle();

    const int weight_blocks = static_cast<int>(
        std::min<int64_t>((weight.numel() + 255) / 256, 65535));

    split_weights<<<weight_blocks, 256, 0, stream>>>(
        weight.data_ptr<float>(), wh, wl, weight.numel());
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    for (int64_t batch = 0; batch < B; ++batch) {
        const float* xb = x.data_ptr<float>() + batch * C * H * W;
        float* yb = output.data_ptr<float>() + batch * OC * spatial64;

        for (int start = 0; start < spatial; start += chunk) {
            const int positions = std::min(chunk, spatial - start);
            const int blocks = static_cast<int>(
                std::min<int64_t>(
                    (K64 * positions + 255) / 256, 65535));

            im2col_split<<<blocks, 256, 0, stream>>>(
                xb, ch, cl,
                static_cast<int>(H), static_cast<int>(W),
                static_cast<int>(KH), static_cast<int>(KW),
                static_cast<int>(OW), K,
                start, positions,
                static_cast<int>(sh), static_cast<int>(sw),
                static_cast<int>(ph), static_cast<int>(pw),
                static_cast<int>(dh), static_cast<int>(dw));
            C10_CUDA_KERNEL_LAUNCH_CHECK();

            float* destination = yb + start;

            // Accumulate the two residual products before the main product.
            gemm(handle, cl, wh, destination,
                 positions, static_cast<int>(OC), K, spatial,
                 INV_RESIDUAL_SCALE, 0.0f);

            gemm(handle, ch, wl, destination,
                 positions, static_cast<int>(OC), K, spatial,
                 INV_RESIDUAL_SCALE, 1.0f);

            gemm(handle, ch, wh, destination,
                 positions, static_cast<int>(OC), K, spatial,
                 1.0f, 1.0f);
        }
    }

    if (bias.has_value()) {
        const auto& b = bias.value();
        TORCH_CHECK(
            b.is_cuda() && b.device() == x.device() &&
            b.scalar_type() == torch::kFloat32 &&
            b.is_contiguous() && b.numel() == OC,
            "Invalid bias");

        const int blocks = static_cast<int>(
            std::min<int64_t>((output.numel() + 255) / 256, 65535));
        add_bias<<<blocks, 256, 0, stream>>>(
            output.data_ptr<float>(), b.data_ptr<float>(),
            output.numel(), static_cast<int>(OC), spatial);
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }

    return output;
}
"""

_extension = load_inline(
    name="asymmetric_dilated_conv_split_gemm_v1",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["asymmetric_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 torch.is_autocast_enabled()
        ):
            return conv(x)

        return _extension.asymmetric_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],
        )
```