```python
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline


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

namespace wmma = nvcuda::wmma;

__device__ __forceinline__ float round_tf32(float x) {
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
    unsigned bits;
    asm volatile("cvt.rn.tf32.f32 %0, %1;"
                 : "=r"(bits) : "f"(x));
    return __uint_as_float(bits);
#else
    return x;
#endif
}

// Implicit GEMM: [64, 1440] @ [1440, output_height * output_width].
// Each block computes 64 output channels and 128 spatial positions.
__global__ __launch_bounds__(256)
void asymmetric_conv_kernel(
    const float* __restrict__ input,
    const float* __restrict__ weight,
    float* __restrict__ output,
    int height,
    int width,
    int output_width,
    int spatial_size)
{
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
    constexpr int K = 32 * 5 * 9;
    constexpr int BK = 32;
    constexpr int BN = 128;

    __shared__ union {
        struct {
            float ah[64 * BK];
            float al[64 * BK];
            float bh[BK * BN];
            float bl[BK * BN];
        } operands;
        float result[64 * BN];
    } storage;

    const int tid = threadIdx.x;
    const int warp = tid / 32;
    const int channel_start = (warp / 2) * 16;
    const int column_start = (warp % 2) * 64;
    const int spatial_start = blockIdx.x * BN;

    const float* batch_input =
        input + static_cast<int64_t>(blockIdx.y) * 32 * height * width;

    wmma::fragment<wmma::accumulator, 16, 16, 8, float> accum[4];

    #pragma unroll
    for (int j = 0; j < 4; ++j)
        wmma::fill_fragment(accum[j], 0.0f);

    for (int kb = 0; kb < K; kb += BK) {
        for (int i = tid; i < 64 * BK; i += 256) {
            const int channel = i / BK;
            const int k = kb + i % BK;
            const float value = weight[channel * K + k];
            const float high = round_tf32(value);
            storage.operands.ah[i] = high;
            storage.operands.al[i] = round_tf32(value - high);
        }

        for (int i = tid; i < BK * BN; i += 256) {
            const int k = kb + i / BN;
            const int spatial = spatial_start + i % BN;

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

                const int64_t offset =
                    (static_cast<int64_t>(channel) * height
                     + output_y + kernel_y) * width
                    + output_x + kernel_x;

                value = batch_input[offset];
            }

            const float high = round_tf32(value);
            storage.operands.bh[i] = high;
            storage.operands.bl[i] = round_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::row_major
            > ah, al;

            wmma::load_matrix_sync(
                ah, storage.operands.ah + channel_start * BK + kk, BK);
            wmma::load_matrix_sync(
                al, storage.operands.al + channel_start * BK + kk, BK);

            #pragma unroll
            for (int j = 0; j < 4; ++j) {
                wmma::fragment<
                    wmma::matrix_b, 16, 16, 8,
                    wmma::precision::tf32, wmma::row_major
                > bh, bl;

                const int offset = kk * BN + column_start + j * 16;
                wmma::load_matrix_sync(
                    bh, storage.operands.bh + offset, BN);
                wmma::load_matrix_sync(
                    bl, storage.operands.bl + offset, BN);

                // Three TF32 products retain substantially more precision
                // than a single TF32 product.
                wmma::mma_sync(accum[j], ah, bh, accum[j]);
                wmma::mma_sync(accum[j], ah, bl, accum[j]);
                wmma::mma_sync(accum[j], al, bh, accum[j]);
            }
        }

        __syncthreads();
    }

    #pragma unroll
    for (int j = 0; j < 4; ++j) {
        wmma::store_matrix_sync(
            storage.result + channel_start * BN + column_start + j * 16,
            accum[j], BN, wmma::mem_row_major);
    }

    __syncthreads();

    for (int i = tid; i < 64 * BN; i += 256) {
        const int channel = i / BN;
        const int spatial = spatial_start + i % BN;
        if (spatial < spatial_size) {
            const int64_t offset =
                (static_cast<int64_t>(blockIdx.y) * 64 + channel)
                * spatial_size + spatial;
            output[offset] = storage.result[i];
        }
    }
#endif
}

torch::Tensor asymmetric_conv_cuda(
    torch::Tensor input,
    torch::Tensor weight)
{
    TORCH_CHECK(input.is_cuda() && weight.is_cuda(),
                "Expected CUDA tensors");
    TORCH_CHECK(input.scalar_type() == torch::kFloat32 &&
                weight.scalar_type() == torch::kFloat32,
                "Expected float32 tensors");
    TORCH_CHECK(input.is_contiguous() && weight.is_contiguous(),
                "Expected contiguous tensors");
    TORCH_CHECK(input.dim() == 4 && input.size(1) == 32,
                "Expected NCHW input with 32 channels");
    TORCH_CHECK(weight.dim() == 4 &&
                weight.size(0) == 64 && weight.size(1) == 32 &&
                weight.size(2) == 5 && weight.size(3) == 9,
                "Expected weight shape [64, 32, 5, 9]");
    TORCH_CHECK(input.device() == weight.device(),
                "Input and weight must use the same device");
    TORCH_CHECK(input.size(2) >= 5 && input.size(3) >= 9,
                "Input is smaller than the kernel");

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

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

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

    if (batch == 0)
        return output;

    const dim3 grid((spatial_size + 127) / 128, batch);
    const auto stream = c10::cuda::getCurrentCUDAStream();

    asymmetric_conv_kernel<<<grid, 256, 0, stream.stream()>>>(
        input.data_ptr<float>(),
        weight.data_ptr<float>(),
        output.data_ptr<float>(),
        height, width, output_width, spatial_size);

    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_tf32x3",
            cpp_sources="""
                torch::Tensor asymmetric_conv_cuda(
                    torch::Tensor input, torch::Tensor weight);
            """,
            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,
        )

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        conv = self.conv2d
        supported = (
            x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 4
            and x.is_contiguous()
            and x.shape[1] == 32
            and x.shape[2] >= 5
            and x.shape[3] >= 9
            and x.shape[0] <= 65535
            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 conv.bias is None
            and conv.weight.dtype == torch.float32
            and conv.weight.device == x.device
            and conv.weight.is_contiguous()
            and not torch.is_autocast_enabled()
        )

        if supported and torch.cuda.get_device_capability(x.device)[0] >= 8:
            return _get_extension().asymmetric_conv_cuda(x, conv.weight)

        return conv(x)
```