```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,
    bool has_bias);
"""

_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>
#include <mma.h>

namespace wmma = nvcuda::wmma;

__global__ void asymmetric_conv_kernel(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int height,
    int width,
    int output_height,
    int output_width)
{
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
    constexpr int BM = 32;
    constexpr int BN = 64;
    constexpr int BK = 32;
    constexpr int K = 32 * 5 * 9;

    __shared__ __align__(32) float storage[2 * BM * BK + 2 * BK * BN];

    float* a_high = storage;
    float* a_low  = a_high + BM * BK;
    float* b_high = a_low + BM * BK;
    float* b_low  = b_high + BK * BN;

    const int tid = threadIdx.x;
    const int warp = tid >> 5;
    const int warp_m = warp & 1;
    const int warp_n = warp >> 1;

    const int spatial_size = output_height * output_width;
    const int spatial_start = blockIdx.x * BM;
    const int batch = blockIdx.y;

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

    for (int k_base = 0; k_base < K; k_base += BK) {
        // A is column-major: adjacent threads load adjacent output positions.
        for (int i = tid; i < BM * BK; i += blockDim.x) {
            const int local_m = i % BM;
            const int k = k_base + i / BM;
            const int position = spatial_start + local_m;

            float value = 0.0f;
            if (position < spatial_size) {
                const int channel = k / 45;
                const int kernel_position = k % 45;
                const int kernel_y = kernel_position / 9;
                const int kernel_x = kernel_position % 9;
                const int output_y = position / output_width;
                const int output_x = position % output_width;

                const size_t input_index =
                    ((size_t(batch) * 32 + channel) * height
                     + output_y + kernel_y) * width
                    + output_x + kernel_x;
                value = x[input_index];
            }

            const float high = wmma::__float_to_tf32(value);
            a_high[i] = high;
            a_low[i] = wmma::__float_to_tf32(value - high);
        }

        // B is column-major, matching contiguous OIHW weight rows.
        for (int i = tid; i < BK * BN; i += blockDim.x) {
            const int k = k_base + i % BK;
            const int channel = i / BK;
            const float value = weight[channel * K + k];
            const float high = wmma::__float_to_tf32(value);
            b_high[i] = high;
            b_low[i] = wmma::__float_to_tf32(value - high);
        }

        __syncthreads();

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

            const int a_offset = warp_m * 16 + kk * BM;
            const int b_offset = kk + warp_n * 16 * BK;

            wmma::load_matrix_sync(ah, a_high + a_offset, BM);
            wmma::load_matrix_sync(al, a_low + a_offset, BM);
            wmma::load_matrix_sync(bh, b_high + b_offset, BK);
            wmma::load_matrix_sync(bl, b_low + b_offset, BK);

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

        __syncthreads();
    }

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

    // Reuse shared storage for a coalesced NCHW output write.
    float* result = storage;
    wmma::store_matrix_sync(
        result + warp_m * 16 + warp_n * 16 * BM,
        main_sum,
        BM,
        wmma::mem_col_major);

    __syncthreads();

    for (int i = tid; i < BM * BN; i += blockDim.x) {
        const int position = spatial_start + i % BM;
        const int channel = i / BM;
        if (position < spatial_size) {
            float value = result[i];
            if (bias != nullptr) {
                value += bias[channel];
            }
            output[(size_t(batch) * BN + channel) * spatial_size
                   + position] = value;
        }
    }
#endif
}

torch::Tensor asymmetric_conv_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    torch::Tensor bias,
    bool has_bias)
{
    const c10::cuda::CUDAGuard guard(x.device());

    TORCH_CHECK(x.is_cuda() && weight.is_cuda(), "CUDA tensors required");
    TORCH_CHECK(
        x.scalar_type() == torch::kFloat32 &&
        weight.scalar_type() == torch::kFloat32,
        "FP32 tensors required");
    TORCH_CHECK(
        x.is_contiguous() && weight.is_contiguous(),
        "Contiguous tensors required");
    TORCH_CHECK(
        x.dim() == 4 && x.size(1) == 32 &&
        weight.dim() == 4 && weight.size(0) == 64 &&
        weight.size(1) == 32 && weight.size(2) == 5 &&
        weight.size(3) == 9,
        "Unsupported convolution dimensions");
    TORCH_CHECK(
        at::cuda::getCurrentDeviceProperties()->major >= 8,
        "This kernel requires compute capability 8.0 or newer");

    const int batch = x.size(0);
    const int height = x.size(2);
    const int width = x.size(3);
    const int output_height = height - 4;
    const int output_width = width - 8;

    TORCH_CHECK(output_height > 0 && output_width > 0,
                "Input is smaller than the convolution kernel");

    auto output = torch::empty(
        {batch, 64, output_height, output_width}, x.options());

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

    const int spatial_size = output_height * output_width;
    const dim3 grid((spatial_size + 31) / 32, batch);
    const cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    asymmetric_conv_kernel<<<grid, 256, 0, stream>>>(
        x.data_ptr<float>(),
        weight.data_ptr<float>(),
        has_bias ? bias.data_ptr<float>() : nullptr,
        output.data_ptr<float>(),
        height,
        width,
        output_height,
        output_width);

    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_corrected",
            cpp_sources=_cpp_source,
            cuda_sources=_cuda_source,
            functions=["asymmetric_conv_cuda"],
            extra_cflags=["-O3"],
            extra_cuda_cflags=["-O3"],
            with_cuda=True,
            verbose=False,
        )
    return _extension


class _AsymmetricConv(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x, weight, bias):
        ctx.save_for_backward(x, weight)
        ctx.has_bias = bias is not None
        return _get_extension().asymmetric_conv_cuda(
            x, weight, weight if bias is None else bias, bias is not None
        )

    @staticmethod
    def backward(ctx, grad_output):
        x, weight = ctx.saved_tensors
        grad_output = grad_output.contiguous()

        grad_x = grad_weight = grad_bias = None
        if ctx.needs_input_grad[0]:
            grad_x = torch.nn.grad.conv2d_input(
                x.shape, weight, grad_output
            )
        if ctx.needs_input_grad[1]:
            grad_weight = torch.nn.grad.conv2d_weight(
                x, weight.shape, grad_output
            )
        if ctx.has_bias and ctx.needs_input_grad[2]:
            grad_bias = grad_output.sum(dim=(0, 2, 3))

        return grad_x, grad_weight, grad_bias


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._device_support = {}

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        conv = self.conv2d

        eligible = (
            x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 4
            and x.is_contiguous()
            and conv.weight.dtype == torch.float32
            and conv.weight.device == x.device
            and conv.weight.is_contiguous()
            and conv.in_channels == 32
            and conv.out_channels == 64
            and conv.kernel_size == (5, 9)
            and conv.stride == (1, 1)
            and conv.padding == (0, 0)
            and conv.dilation == (1, 1)
            and conv.groups == 1
            and x.shape[1] == 32
            and x.shape[2] >= 5
            and x.shape[3] >= 9
            and x.shape[0] <= 65535
            and not torch.is_autocast_enabled()
        )

        if eligible:
            device_index = x.device.index
            if device_index not in self._device_support:
                self._device_support[device_index] = (
                    torch.cuda.get_device_capability(x.device)[0] >= 8
                )
            if self._device_support[device_index]:
                return _AsymmetricConv.apply(x, conv.weight, conv.bias)

        return conv(x)
```