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

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

// Implicit GEMM: 32 output channels by 64 spatial positions.
// Two TF32 components retain substantially more precision than plain TF32.
__global__ void asymmetric_conv_kernel(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int channels, int height, int width,
    int out_channels, int out_height, int out_width,
    int channels_per_group, int outputs_per_group,
    int sh, int sw, int ph, int pw, int dh, int dw,
    int groups)
{
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
    __shared__ __align__(32) float a_hi[32 * 32];
    __shared__ __align__(32) float a_lo[32 * 32];
    __shared__ __align__(32) float b_hi[32 * 64];
    __shared__ __align__(32) float b_lo[32 * 64];
    __shared__ __align__(32) float result[32 * 64];

    const int tid = threadIdx.x;
    const int warp = tid / 32;
    const int tile_row = (warp / 2) * 16;
    const int tile_col = (warp % 2) * 32;

    const int batch = blockIdx.z / groups;
    const int group = blockIdx.z % groups;
    const int output_channel_base = blockIdx.y * 32;
    const int position_base = blockIdx.x * 64;
    const int positions = out_height * out_width;
    const int reduction = channels_per_group * 45;

    using AFragment = wmma::fragment<
        wmma::matrix_a, 16, 16, 8,
        wmma::precision::tf32, wmma::row_major>;
    using BFragment = wmma::fragment<
        wmma::matrix_b, 16, 16, 8,
        wmma::precision::tf32, wmma::row_major>;
    using CFragment = wmma::fragment<
        wmma::accumulator, 16, 16, 8, float>;

    CFragment main_sum[2], correction[2];

    #pragma unroll
    for (int j = 0; j < 2; ++j) {
        wmma::fill_fragment(main_sum[j], 0.0f);
        wmma::fill_fragment(correction[j], 0.0f);
    }

    for (int k_base = 0; k_base < reduction; k_base += 32) {
        for (int index = tid; index < 32 * 32; index += 128) {
            const int row = index / 32;
            const int k = k_base + index % 32;
            const int oc = output_channel_base + row;

            float value = 0.0f;
            if (oc < outputs_per_group && k < reduction) {
                const int global_oc = group * outputs_per_group + oc;
                value = weight[(int64_t)global_oc * reduction + k];
            }

            const float high = __float_to_tf32(value);
            a_hi[index] = high;
            a_lo[index] = __float_to_tf32(value - high);
        }

        for (int index = tid; index < 32 * 64; index += 128) {
            const int k = k_base + index / 64;
            const int position = position_base + index % 64;

            float value = 0.0f;
            if (k < reduction && position < positions) {
                const int ic = k / 45;
                const int kernel_position = k % 45;
                const int kh = kernel_position / 9;
                const int kw = kernel_position % 9;

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

                if ((unsigned)ih < (unsigned)height &&
                    (unsigned)iw < (unsigned)width) {
                    const int global_ic = group * channels_per_group + ic;
                    const int64_t offset =
                        (((int64_t)batch * channels + global_ic) * height
                         + ih) * width + iw;
                    value = x[offset];
                }
            }

            const float high = __float_to_tf32(value);
            b_hi[index] = high;
            b_lo[index] = __float_to_tf32(value - high);
        }

        __syncthreads();

        #pragma unroll
        for (int kk = 0; kk < 32; kk += 8) {
            AFragment ah, al;
            wmma::load_matrix_sync(
                ah, a_hi + tile_row * 32 + kk, 32);
            wmma::load_matrix_sync(
                al, a_lo + tile_row * 32 + kk, 32);

            #pragma unroll
            for (int j = 0; j < 2; ++j) {
                BFragment bh, bl;
                const int offset = kk * 64 + tile_col + j * 16;

                wmma::load_matrix_sync(bh, b_hi + offset, 64);
                wmma::load_matrix_sync(bl, b_lo + offset, 64);

                wmma::mma_sync(
                    main_sum[j], ah, bh, main_sum[j]);
                wmma::mma_sync(
                    correction[j], ah, bl, correction[j]);
                wmma::mma_sync(
                    correction[j], al, bh, correction[j]);
                wmma::mma_sync(
                    correction[j], al, bl, correction[j]);
            }
        }

        __syncthreads();
    }

    #pragma unroll
    for (int j = 0; j < 2; ++j) {
        #pragma unroll
        for (int i = 0; i < main_sum[j].num_elements; ++i) {
            main_sum[j].x[i] += correction[j].x[i];
        }

        wmma::store_matrix_sync(
            result + tile_row * 64 + tile_col + j * 16,
            main_sum[j], 64, wmma::mem_row_major);
    }

    __syncthreads();

    for (int index = tid; index < 32 * 64; index += 128) {
        const int oc = output_channel_base + index / 64;
        const int position = position_base + index % 64;

        if (oc < outputs_per_group && position < positions) {
            const int global_oc = group * outputs_per_group + oc;
            float value = result[index];
            if (bias != nullptr) {
                value += bias[global_oc];
            }

            output[
                ((int64_t)batch * out_channels + global_oc) * positions
                + position
            ] = value;
        }
    }
#endif
}

torch::Tensor asymmetric_conv_cuda(
    torch::Tensor x, torch::Tensor weight, torch::Tensor bias,
    int64_t sh, int64_t sw, int64_t ph, int64_t pw,
    int64_t dh, int64_t dw, int64_t groups)
{
    const c10::cuda::CUDAGuard device_guard(x.device());

    TORCH_CHECK(x.is_cuda() && weight.is_cuda(),
                "Input and weight must be CUDA tensors");
    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat,
                "Input and weight must be float32");
    TORCH_CHECK(x.is_contiguous() && weight.is_contiguous(),
                "Input and weight must be contiguous");
    TORCH_CHECK(x.dim() == 4 && weight.dim() == 4,
                "Expected four-dimensional tensors");
    TORCH_CHECK(weight.size(2) == 5 && weight.size(3) == 9,
                "Expected a 5x9 kernel");
    TORCH_CHECK(x.device() == weight.device(),
                "Input and weight must share a device");
    TORCH_CHECK(groups > 0 && sh > 0 && sw > 0 && dh > 0 && dw > 0,
                "Invalid convolution parameters");
    TORCH_CHECK(
        x.size(1) == weight.size(1) * groups &&
        weight.size(0) % groups == 0,
        "Invalid channel or group configuration");
    TORCH_CHECK(at::cuda::getCurrentDeviceProperties()->major >= 8,
                "This kernel requires compute capability 8.0 or newer");

    const int64_t height_numerator = x.size(2) + 2 * ph - 4 * dh - 1;
    const int64_t width_numerator = x.size(3) + 2 * pw - 8 * dw - 1;
    TORCH_CHECK(height_numerator >= 0 && width_numerator >= 0,
                "Kernel exceeds padded input dimensions");

    const int64_t oh = height_numerator / sh + 1;
    const int64_t ow = width_numerator / sw + 1;

    auto output = torch::empty(
        {x.size(0), weight.size(0), oh, ow}, x.options());

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

    const float* bias_ptr = nullptr;
    if (bias.numel() != 0) {
        TORCH_CHECK(
            bias.device() == x.device() &&
            bias.scalar_type() == at::kFloat &&
            bias.is_contiguous() &&
            bias.numel() == weight.size(0),
            "Invalid bias tensor");
        bias_ptr = bias.data_ptr<float>();
    }

    const int outputs_per_group = weight.size(0) / groups;
    dim3 grid(
        (oh * ow + 63) / 64,
        (outputs_per_group + 31) / 32,
        x.size(0) * groups);

    asymmetric_conv_kernel<<<
        grid, 128, 0, at::cuda::getCurrentCUDAStream()
    >>>(
        x.data_ptr<float>(),
        weight.data_ptr<float>(),
        bias_ptr,
        output.data_ptr<float>(),
        x.size(1), x.size(2), x.size(3),
        weight.size(0), oh, ow,
        weight.size(1), outputs_per_group,
        sh, sw, ph, pw, dh, dw, groups);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = None


def _get_extension():
    global _extension
    if _extension is None:
        _extension = load_inline(
            name="asymmetric_conv_5x9_tf32_components",
            cpp_sources=_cpp_source,
            cuda_sources=_cuda_source,
            functions=["asymmetric_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: int = 0,
        dilation: int = 1,
        groups: int = 1,
        bias: bool = False,
    ):
        super().__init__()
        self.conv2d = nn.Conv2d(
            in_channels,
            out_channels,
            kernel_size,
            stride=stride,
            padding=padding,
            dilation=dilation,
            groups=groups,
            bias=bias,
        )
        self.register_buffer("_empty_bias", torch.empty(0), persistent=False)

    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 conv.kernel_size != (5, 9)
            or isinstance(conv.padding, str)
            or x.ndim != 4
            or x.shape[0] * conv.groups > 65535
            or torch.is_autocast_enabled()
        ):
            return conv(x)

        if torch.cuda.get_device_capability(x.device)[0] < 8:
            return conv(x)

        return _get_extension().asymmetric_conv_cuda(
            x.contiguous(),
            conv.weight.contiguous(),
            conv.bias if conv.bias is not None else self._empty_bias,
            conv.stride[0],
            conv.stride[1],
            conv.padding[0],
            conv.padding[1],
            conv.dilation[0],
            conv.dilation[1],
            conv.groups,
        )
```