```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,
    c10::optional<torch::Tensor> 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.h>
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include <mma.h>

namespace {

constexpr int K = 32 * 5 * 9;
constexpr int TILE_M = 64;
constexpr int TILE_N = 64;
constexpr int TILE_K = 32;
constexpr int OUTPUT_LD = 68;

// Each float is represented by a leading half and a residual half.
// Three tensor-core products retain substantially more precision than
// a single half-precision convolution.
__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 __CUDA_ARCH__ >= 700
    using namespace nvcuda;

    __shared__ __align__(32) half a_hi[TILE_M * TILE_K];
    __shared__ __align__(32) half a_lo[TILE_M * TILE_K];
    __shared__ __align__(32) half b_hi[TILE_N * TILE_K];
    __shared__ __align__(32) half b_lo[TILE_N * TILE_K];
    __shared__ __align__(32) float result[TILE_M * OUTPUT_LD];

    const int tid = threadIdx.x;
    const int warp = tid >> 5;
    const int warp_m = (warp >> 1) * 16;
    const int warp_n = (warp & 1) * 32;

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

    const float* batch_input =
        x + static_cast<int64_t>(batch) * 32 * height * width;

    wmma::fragment<
        wmma::accumulator, 16, 16, 16, float
    > accum0, accum1;

    wmma::fill_fragment(accum0, 0.0f);
    wmma::fill_fragment(accum1, 0.0f);

    for (int k_start = 0; k_start < K; k_start += TILE_K) {
        #pragma unroll
        for (int i = tid; i < TILE_M * TILE_K; i += 256) {
            const int local_m = i / TILE_K;
            const int k = k_start + i % TILE_K;
            const int spatial = spatial_start + local_m;

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

                value = batch_input[
                    (static_cast<int64_t>(channel) * height +
                     y + kernel_y) * width + x_pos + kernel_x
                ];
            }

            const half leading = __float2half_rn(value);
            a_hi[i] = leading;
            a_lo[i] = __float2half_rn(value - __half2float(leading));
        }

        #pragma unroll
        for (int i = tid; i < TILE_N * TILE_K; i += 256) {
            const int channel = i / TILE_K;
            const int k = k_start + i % TILE_K;
            const float value = weight[channel * K + k];

            const half leading = __float2half_rn(value);
            b_hi[i] = leading;
            b_lo[i] = __float2half_rn(value - __half2float(leading));
        }

        __syncthreads();

        #pragma unroll
        for (int kk = 0; kk < TILE_K; kk += 16) {
            wmma::fragment<
                wmma::matrix_a, 16, 16, 16, half, wmma::row_major
            > ah, al;

            wmma::fragment<
                wmma::matrix_b, 16, 16, 16, half, wmma::col_major
            > bh, bl;

            wmma::load_matrix_sync(
                ah, a_hi + warp_m * TILE_K + kk, TILE_K);
            wmma::load_matrix_sync(
                al, a_lo + warp_m * TILE_K + kk, TILE_K);

            wmma::load_matrix_sync(
                bh, b_hi + warp_n * TILE_K + kk, TILE_K);
            wmma::load_matrix_sync(
                bl, b_lo + warp_n * TILE_K + kk, TILE_K);

            wmma::mma_sync(accum0, ah, bh, accum0);
            wmma::mma_sync(accum0, ah, bl, accum0);
            wmma::mma_sync(accum0, al, bh, accum0);

            wmma::load_matrix_sync(
                bh, b_hi + (warp_n + 16) * TILE_K + kk, TILE_K);
            wmma::load_matrix_sync(
                bl, b_lo + (warp_n + 16) * TILE_K + kk, TILE_K);

            wmma::mma_sync(accum1, ah, bh, accum1);
            wmma::mma_sync(accum1, ah, bl, accum1);
            wmma::mma_sync(accum1, al, bh, accum1);
        }

        __syncthreads();
    }

    wmma::store_matrix_sync(
        result + warp_m * OUTPUT_LD + warp_n,
        accum0, OUTPUT_LD, wmma::mem_row_major);

    wmma::store_matrix_sync(
        result + warp_m * OUTPUT_LD + warp_n + 16,
        accum1, OUTPUT_LD, wmma::mem_row_major);

    __syncthreads();

    #pragma unroll
    for (int i = tid; i < TILE_M * TILE_N; i += 256) {
        const int channel = i / TILE_M;
        const int local_m = i % TILE_M;
        const int spatial = spatial_start + local_m;

        if (spatial < spatial_size) {
            float value = result[local_m * OUTPUT_LD + channel];
            if (bias != nullptr) {
                value += bias[channel];
            }

            output[
                (static_cast<int64_t>(batch) * 64 + channel) *
                spatial_size + spatial
            ] = value;
        }
    }
#endif
}

} // namespace

torch::Tensor asymmetric_conv_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias)
{
    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.device() == weight.device(),
                "Input and weight must be on the same device");
    TORCH_CHECK(x.dim() == 4 && x.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)");

    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 kernel");
    TORCH_CHECK(
        at::cuda::getDeviceProperties(x.get_device())->major >= 7,
        "The custom kernel requires compute capability 7.0 or newer");

    const float* bias_ptr = nullptr;
    if (bias.has_value()) {
        const auto& b = bias.value();
        TORCH_CHECK(
            b.device() == x.device() &&
            b.scalar_type() == at::kFloat &&
            b.is_contiguous() &&
            b.numel() == 64,
            "Bias must be a contiguous CUDA float32 tensor of length 64");
        bias_ptr = b.data_ptr<float>();
    }

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

    if (x.size(0) == 0) {
        return output;
    }

    const dim3 grid(
        (output_height * output_width + TILE_M - 1) / TILE_M,
        static_cast<unsigned int>(x.size(0)));

    asymmetric_conv_kernel<<<
        grid, 256, 0, at::cuda::getCurrentCUDAStream()
    >>>(
        x.data_ptr<float>(),
        weight.data_ptr<float>(),
        bias_ptr,
        output.data_ptr<float>(),
        height,
        width,
        output_height,
        output_width);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = load_inline(
    name="asymmetric_conv_5x9_split_half_v1",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["asymmetric_conv_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


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._specialized = (
            in_channels == 32
            and out_channels == 64
            and self.conv2d.kernel_size == (5, 9)
            and self.conv2d.stride == (1, 1)
            and self.conv2d.padding == (0, 0)
            and self.conv2d.dilation == (1, 1)
            and groups == 1
        )
        self._device_support = {}

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

        if (
            not self._specialized
            or not x.is_cuda
            or x.dtype != torch.float32
            or weight.dtype != torch.float32
            or not x.is_contiguous()
            or not weight.is_contiguous()
            or x.ndim != 4
            or x.shape[0] > 65535
            or torch.is_autocast_enabled()
        ):
            return self.conv2d(x)

        device_index = x.device.index
        if device_index not in self._device_support:
            self._device_support[device_index] = (
                torch.cuda.get_device_capability(device_index)[0] >= 7
            )

        if not self._device_support[device_index]:
            return self.conv2d(x)

        return _extension.asymmetric_conv_cuda(
            x, weight, self.conv2d.bias
        )
```