```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,
    torch::Tensor bias,
    int64_t stride_h,
    int64_t stride_w,
    int64_t pad_h,
    int64_t pad_w,
    int64_t dilation_h,
    int64_t dilation_w,
    int64_t groups,
    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.h>
#include <cuda_runtime.h>
#include <mma.h>

namespace {

constexpr int TILE_M = 32;
constexpr int TILE_N = 64;
constexpr int TILE_K = 32;
constexpr int STRIDE_A = 40;
constexpr int STRIDE_B = 72;

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

__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 kernel_h,
    int kernel_w,
    int out_h,
    int out_w,
    int stride_h,
    int stride_w,
    int pad_h,
    int pad_w,
    int dilation_h,
    int dilation_w,
    int groups,
    bool has_bias) {

#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
    using namespace nvcuda;

    __shared__ __align__(32) union SharedStorage {
        struct {
            float a_hi[TILE_M * STRIDE_A];
            float a_lo[TILE_M * STRIDE_A];
            float b_hi[TILE_K * STRIDE_B];
            float b_lo[TILE_K * STRIDE_B];
        } operands;
        float result[TILE_M * TILE_N];
    } shared;

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

    const int channels_per_group = channels / groups;
    const int outputs_per_group = out_channels / groups;
    const int group = blockIdx.z % groups;
    const int batch = blockIdx.z / groups;

    const int row_base = blockIdx.y * TILE_M;
    const int column_base = blockIdx.x * TILE_N;
    const int spatial_size = out_h * out_w;
    const int kernel_area = kernel_h * kernel_w;
    const int reduction_size = channels_per_group * kernel_area;

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

    for (int k_base = 0; k_base < reduction_size; k_base += TILE_K) {
        for (int i = tid; i < TILE_M * TILE_K; i += blockDim.x) {
            const int row = i / TILE_K;
            const int k = i % TILE_K;
            const int output_channel = row_base + row;
            const int reduction_index = k_base + k;

            float value = 0.0f;
            if (output_channel < outputs_per_group &&
                reduction_index < reduction_size) {
                const int global_channel =
                    group * outputs_per_group + output_channel;
                value = weight[
                    static_cast<int64_t>(global_channel) * reduction_size +
                    reduction_index];
            }

            const float hi = round_tf32(value);
            shared.operands.a_hi[row * STRIDE_A + k] = hi;
            shared.operands.a_lo[row * STRIDE_A + k] =
                round_tf32(value - hi);
        }

        for (int i = tid; i < TILE_K * TILE_N; i += blockDim.x) {
            const int k = i / TILE_N;
            const int column = i % TILE_N;
            const int reduction_index = k_base + k;
            const int spatial_index = column_base + column;

            float value = 0.0f;
            if (reduction_index < reduction_size &&
                spatial_index < spatial_size) {
                const int input_channel = reduction_index / kernel_area;
                const int kernel_index = reduction_index % kernel_area;
                const int kh = kernel_index / kernel_w;
                const int kw = kernel_index % kernel_w;

                const int oh = spatial_index / out_w;
                const int ow = spatial_index % out_w;
                const int ih = oh * stride_h - pad_h + kh * dilation_h;
                const int iw = ow * stride_w - pad_w + kw * dilation_w;

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

            const float hi = round_tf32(value);
            shared.operands.b_hi[k * STRIDE_B + column] = hi;
            shared.operands.b_lo[k * STRIDE_B + column] =
                round_tf32(value - hi);
        }

        __syncthreads();

        #pragma unroll
        for (int kk = 0; kk < TILE_K; kk += 8) {
            wmma::fragment<
                wmma::matrix_a, 16, 16, 8,
                wmma::precision::tf32, wmma::row_major> a_hi, a_lo;
            wmma::fragment<
                wmma::matrix_b, 16, 16, 8,
                wmma::precision::tf32, wmma::row_major> b_hi, b_lo;

            wmma::load_matrix_sync(
                a_hi,
                shared.operands.a_hi + warp_m * STRIDE_A + kk,
                STRIDE_A);
            wmma::load_matrix_sync(
                a_lo,
                shared.operands.a_lo + warp_m * STRIDE_A + kk,
                STRIDE_A);

            wmma::load_matrix_sync(
                b_hi,
                shared.operands.b_hi + kk * STRIDE_B + warp_n,
                STRIDE_B);
            wmma::load_matrix_sync(
                b_lo,
                shared.operands.b_lo + kk * STRIDE_B + warp_n,
                STRIDE_B);

            wmma::mma_sync(acc0, a_hi, b_hi, acc0);
            wmma::mma_sync(acc0, a_lo, b_hi, acc0);
            wmma::mma_sync(acc0, a_hi, b_lo, acc0);

            wmma::load_matrix_sync(
                b_hi,
                shared.operands.b_hi + kk * STRIDE_B + warp_n + 16,
                STRIDE_B);
            wmma::load_matrix_sync(
                b_lo,
                shared.operands.b_lo + kk * STRIDE_B + warp_n + 16,
                STRIDE_B);

            wmma::mma_sync(acc1, a_hi, b_hi, acc1);
            wmma::mma_sync(acc1, a_lo, b_hi, acc1);
            wmma::mma_sync(acc1, a_hi, b_lo, acc1);
        }

        __syncthreads();
    }

    wmma::store_matrix_sync(
        shared.result + warp_m * TILE_N + warp_n,
        acc0, TILE_N, wmma::mem_row_major);
    wmma::store_matrix_sync(
        shared.result + warp_m * TILE_N + warp_n + 16,
        acc1, TILE_N, wmma::mem_row_major);

    __syncthreads();

    for (int i = tid; i < TILE_M * TILE_N; i += blockDim.x) {
        const int row = i / TILE_N;
        const int column = i % TILE_N;
        const int output_channel = row_base + row;
        const int spatial_index = column_base + column;

        if (output_channel < outputs_per_group &&
            spatial_index < spatial_size) {
            const int global_channel =
                group * outputs_per_group + output_channel;
            float value = shared.result[i];
            if (has_bias) {
                value += bias[global_channel];
            }

            output[
                (static_cast<int64_t>(batch) * out_channels +
                 global_channel) * spatial_size + spatial_index] = value;
        }
    }
#endif
}

} // namespace

torch::Tensor asymmetric_conv_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    torch::Tensor bias,
    int64_t stride_h,
    int64_t stride_w,
    int64_t pad_h,
    int64_t pad_w,
    int64_t dilation_h,
    int64_t dilation_w,
    int64_t groups,
    bool has_bias) {

    TORCH_CHECK(x.is_cuda() && weight.is_cuda(), "Expected CUDA tensors");
    TORCH_CHECK(x.device() == weight.device(), "Device mismatch");
    TORCH_CHECK(x.scalar_type() == torch::kFloat32 &&
                weight.scalar_type() == torch::kFloat32,
                "Expected float32 tensors");
    TORCH_CHECK(x.dim() == 4 && weight.dim() == 4,
                "Expected four-dimensional tensors");
    TORCH_CHECK(x.is_contiguous() && weight.is_contiguous(),
                "Expected contiguous tensors");
    TORCH_CHECK(groups > 0 &&
                x.size(1) == weight.size(1) * groups &&
                weight.size(0) % groups == 0,
                "Invalid convolution groups");

    c10::cuda::CUDAGuard guard(x.device());
    TORCH_CHECK(at::cuda::getCurrentDeviceProperties()->major >= 8,
                "This kernel requires compute capability 8.0 or newer");

    const int64_t out_h =
        (x.size(2) + 2 * pad_h -
         dilation_h * (weight.size(2) - 1) - 1) / stride_h + 1;
    const int64_t out_w =
        (x.size(3) + 2 * pad_w -
         dilation_w * (weight.size(3) - 1) - 1) / stride_w + 1;

    TORCH_CHECK(out_h > 0 && out_w > 0, "Invalid output size");
    TORCH_CHECK(x.size(0) * groups <= 65535, "Batch/group count too large");

    if (has_bias) {
        TORCH_CHECK(bias.device() == x.device() &&
                    bias.scalar_type() == torch::kFloat32 &&
                    bias.is_contiguous() &&
                    bias.numel() == weight.size(0),
                    "Invalid bias");
    }

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

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

    const dim3 grid(
        (out_h * out_w + TILE_N - 1) / TILE_N,
        (weight.size(0) / groups + TILE_M - 1) / TILE_M,
        x.size(0) * groups);

    asymmetric_conv_kernel<<<
        grid, 128, 0, at::cuda::getCurrentCUDAStream()>>>(
        x.data_ptr<float>(),
        weight.data_ptr<float>(),
        has_bias ? bias.data_ptr<float>() : nullptr,
        output.data_ptr<float>(),
        x.size(1), x.size(2), x.size(3),
        weight.size(0), weight.size(2), weight.size(3),
        out_h, out_w,
        stride_h, stride_w,
        pad_h, pad_w,
        dilation_h, dilation_w,
        groups, has_bias);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = None


def _get_extension():
    global _extension
    if _extension is None:
        _extension = load_inline(
            name="asymmetric_conv_tf32_split_v1",
            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,
        )

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

        bias = conv.bias if conv.bias is not None else conv.weight

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