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

os.environ.setdefault("MAX_JOBS", "1")

_CPP = r"""
#include <ATen/ATen.h>

at::Tensor asymmetric_conv_forward(
    at::Tensor x, at::Tensor weight, at::Tensor bias,
    int64_t sh, int64_t sw, int64_t ph, int64_t pw,
    int64_t dh, int64_t dw, int64_t groups);
"""

_CUDA = r"""
#include <ATen/ATen.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAStream.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>
#include <mma.h>

namespace {

constexpr int BM = 64;
constexpr int BN = 64;
constexpr int BK = 32;
constexpr int AS = 40;
constexpr int BS = 72;

struct __align__(32) Operands {
    float ah[BM * AS];
    float al[BM * AS];
    float bh[BK * BS];
    float bl[BK * BS];
};

union __align__(32) Shared {
    Operands operands;
    float output[BM * BN];
};

__device__ __forceinline__ float tf32_round(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
}

template<bool FIXED>
__global__ __launch_bounds__(256)
void conv_kernel(
    const float* __restrict__ x,
    const float* __restrict__ w,
    const float* __restrict__ bias,
    float* __restrict__ y,
    int channels, int height, int width, int outputs,
    int kernel_h, int kernel_w, int output_h, int output_w,
    int stride_h, int stride_w, int padding_h, int padding_w,
    int dilation_h, int dilation_w, int groups)
{
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
    using namespace nvcuda;

    __shared__ Shared smem;

    const int C = FIXED ? 32 : channels;
    const int H = FIXED ? 512 : height;
    const int W = FIXED ? 512 : width;
    const int O = FIXED ? 64 : outputs;
    const int KH = FIXED ? 5 : kernel_h;
    const int KW = FIXED ? 9 : kernel_w;
    const int OH = FIXED ? 508 : output_h;
    const int OW = FIXED ? 504 : output_w;
    const int SH = FIXED ? 1 : stride_h;
    const int SW = FIXED ? 1 : stride_w;
    const int PH = FIXED ? 0 : padding_h;
    const int PW = FIXED ? 0 : padding_w;
    const int DH = FIXED ? 1 : dilation_h;
    const int DW = FIXED ? 1 : dilation_w;
    const int G = FIXED ? 1 : groups;

    const int CG = C / G;
    const int OG = O / G;
    const int K = CG * KH * KW;
    const int spatial = OH * OW;

    const int group = FIXED ? 0 : blockIdx.z % G;
    const int batch = FIXED ? blockIdx.z : blockIdx.z / G;
    const int mbase = blockIdx.y * BM;
    const int nbase = blockIdx.x * BN;

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

    typedef wmma::fragment<
        wmma::accumulator, 16, 16, 8, float> Acc;
    typedef wmma::fragment<
        wmma::matrix_a, 16, 16, 8,
        wmma::precision::tf32, wmma::row_major> FragA;
    typedef wmma::fragment<
        wmma::matrix_b, 16, 16, 8,
        wmma::precision::tf32, wmma::row_major> FragB;

    Acc acc0, acc1, correction0, correction1;
    wmma::fill_fragment(acc0, 0.0f);
    wmma::fill_fragment(acc1, 0.0f);
    wmma::fill_fragment(correction0, 0.0f);
    wmma::fill_fragment(correction1, 0.0f);

    for (int base = 0; base < K; base += BK) {
        #pragma unroll
        for (int i = tid; i < BM * BK; i += 256) {
            const int row = i / BK;
            const int k = i % BK;
            const int oc = mbase + row;
            const int r = base + k;
            float v = 0.0f;
            if (oc < OG && r < K) {
                v = w[(int64_t)(group * OG + oc) * K + r];
            }
            const float hi = tf32_round(v);
            smem.operands.ah[row * AS + k] = hi;
            smem.operands.al[row * AS + k] = tf32_round(v - hi);
        }

        #pragma unroll
        for (int i = tid; i < BK * BN; i += 256) {
            const int k = i / BN;
            const int col = i % BN;
            const int r = base + k;
            const int p = nbase + col;

            float v = 0.0f;
            if (r < K && p < spatial) {
                const int ic = r / (KH * KW);
                const int kr = r % (KH * KW);
                const int kh = kr / KW;
                const int kw = kr % KW;
                const int ih = (p / OW) * SH - PH + kh * DH;
                const int iw = (p % OW) * SW - PW + kw * DW;

                if (FIXED ||
                    ((unsigned)ih < (unsigned)H &&
                     (unsigned)iw < (unsigned)W)) {
                    const int64_t offset =
                        ((int64_t)batch * C + group * CG + ic) * H * W
                        + (int64_t)ih * W + iw;
                    v = x[offset];
                }
            }
            const float hi = tf32_round(v);
            smem.operands.bh[k * BS + col] = hi;
            smem.operands.bl[k * BS + col] = tf32_round(v - hi);
        }

        __syncthreads();

        #pragma unroll
        for (int kk = 0; kk < BK; kk += 8) {
            FragA ah, al;
            FragB bh, bl;

            wmma::load_matrix_sync(
                ah, smem.operands.ah + wm * AS + kk, AS);
            wmma::load_matrix_sync(
                al, smem.operands.al + wm * AS + kk, AS);

            wmma::load_matrix_sync(
                bh, smem.operands.bh + kk * BS + wn, BS);
            wmma::load_matrix_sync(
                bl, smem.operands.bl + kk * BS + wn, BS);

            wmma::mma_sync(acc0, ah, bh, acc0);
            wmma::mma_sync(correction0, al, bh, correction0);
            wmma::mma_sync(correction0, ah, bl, correction0);

            wmma::load_matrix_sync(
                bh, smem.operands.bh + kk * BS + wn + 16, BS);
            wmma::load_matrix_sync(
                bl, smem.operands.bl + kk * BS + wn + 16, BS);

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

        __syncthreads();
    }

    #pragma unroll
    for (int i = 0; i < acc0.num_elements; ++i) {
        acc0.x[i] += correction0.x[i];
        acc1.x[i] += correction1.x[i];
    }

    wmma::store_matrix_sync(
        smem.output + wm * BN + wn,
        acc0, BN, wmma::mem_row_major);
    wmma::store_matrix_sync(
        smem.output + wm * BN + wn + 16,
        acc1, BN, wmma::mem_row_major);
    __syncthreads();

    #pragma unroll
    for (int i = tid; i < BM * BN; i += 256) {
        const int oc = mbase + i / BN;
        const int p = nbase + i % BN;
        if (oc < OG && p < spatial) {
            const int global_oc = group * OG + oc;
            float v = smem.output[i];
            if (bias != nullptr) {
                v += bias[global_oc];
            }
            y[((int64_t)batch * O + global_oc) * spatial + p] = v;
        }
    }
#endif
}

} // namespace

at::Tensor asymmetric_conv_forward(
    at::Tensor x, at::Tensor weight, at::Tensor bias,
    int64_t sh, int64_t sw, int64_t ph, int64_t pw,
    int64_t dh, int64_t dw, int64_t groups)
{
    TORCH_CHECK(x.is_cuda() && weight.is_cuda(),
                "Input and weight must be CUDA tensors");
    TORCH_CHECK(x.device() == weight.device(), "Device mismatch");
    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat,
                "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(sh > 0 && sw > 0 && dh > 0 && dw > 0 &&
                ph >= 0 && pw >= 0, "Invalid convolution parameters");
    TORCH_CHECK(groups > 0 &&
                x.size(1) == weight.size(1) * groups &&
                weight.size(0) % groups == 0,
                "Invalid groups");

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

    const int64_t nh = x.size(2) + 2 * ph
                       - dh * (weight.size(2) - 1) - 1;
    const int64_t nw = x.size(3) + 2 * pw
                       - dw * (weight.size(3) - 1) - 1;
    TORCH_CHECK(nh >= 0 && nw >= 0, "Kernel exceeds padded input");
    const int64_t oh = nh / sh + 1;
    const int64_t ow = nw / sw + 1;

    const bool has_bias = bias.numel() != 0;
    if (has_bias) {
        TORCH_CHECK(bias.is_cuda() && bias.device() == x.device() &&
                    bias.scalar_type() == at::kFloat &&
                    bias.is_contiguous() &&
                    bias.numel() == weight.size(0),
                    "Invalid bias");
    }

    auto output = at::empty(
        {x.size(0), weight.size(0), oh, ow}, x.options());
    if (output.numel() == 0) {
        return output;
    }

    TORCH_CHECK(x.size(0) * groups <= 65535 &&
                (weight.size(0) / groups + BM - 1) / BM <= 65535,
                "Convolution grid is too large");

    const dim3 grid(
        (oh * ow + BN - 1) / BN,
        (weight.size(0) / groups + BM - 1) / BM,
        x.size(0) * groups);
    const cudaStream_t stream =
        c10::cuda::getCurrentCUDAStream(x.get_device()).stream();

    const float* bp = has_bias ? bias.data_ptr<float>() : nullptr;

    const bool fixed =
        x.size(1) == 32 && x.size(2) == 512 && x.size(3) == 512 &&
        weight.size(0) == 64 &&
        weight.size(2) == 5 && weight.size(3) == 9 &&
        sh == 1 && sw == 1 && ph == 0 && pw == 0 &&
        dh == 1 && dw == 1 && groups == 1;

    if (fixed) {
        conv_kernel<true><<<grid, 256, 0, stream>>>(
            x.data_ptr<float>(), weight.data_ptr<float>(), bp,
            output.data_ptr<float>(),
            32, 512, 512, 64, 5, 9, 508, 504,
            1, 1, 0, 0, 1, 1, 1);
    } else {
        conv_kernel<false><<<grid, 256, 0, stream>>>(
            x.data_ptr<float>(), weight.data_ptr<float>(), bp,
            output.data_ptr<float>(),
            x.size(1), x.size(2), x.size(3), weight.size(0),
            weight.size(2), weight.size(3), oh, ow,
            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_split_tf32_v2",
            cpp_sources=_CPP,
            cuda_sources=_CUDA,
            functions=["asymmetric_conv_forward"],
            extra_cflags=["-O2"],
            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 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)

        return _get_extension().asymmetric_conv_forward(
            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,
        )
```