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

_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;

__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)
{
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
    // Implicit GEMM: [64 output channels, K] x [K, 64 positions].
    extern __shared__ float shared[];
    float* shared_a = shared;
    float* shared_b = shared + 64 * 32;

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

    const int channel_base = blockIdx.y * 64;
    const int position_base = blockIdx.x * 64;
    const int batch = blockIdx.z;

    const int positions = out_h * out_w;
    const int kernel_area = kernel_h * kernel_w;
    const int reduction_size = channels * kernel_area;

    using AFragment = wmma::fragment<
        wmma::matrix_a, 16, 16, 8,
        wmma::precision::tf32, wmma::row_major>;
    using BFragment = wmma::fragment<
        wmma::matrix_b, 16, 16, 8,
        wmma::precision::tf32, wmma::row_major>;
    using CFragment = wmma::fragment<
        wmma::accumulator, 16, 16, 8, float>;

    CFragment acc0, acc1;
    wmma::fill_fragment(acc0, 0.0f);
    wmma::fill_fragment(acc1, 0.0f);

    for (int reduction_base = 0;
         reduction_base < reduction_size;
         reduction_base += 32)
    {
        for (int index = tid; index < 64 * 32; index += 256) {
            const int row = index / 32;
            const int k = reduction_base + index % 32;
            const int oc = channel_base + row;

            shared_a[index] =
                (oc < out_channels && k < reduction_size)
                ? weight[(int64_t)oc * reduction_size + k]
                : 0.0f;
        }

        for (int index = tid; index < 32 * 64; index += 256) {
            const int k = reduction_base + index / 64;
            const int position = position_base + index % 64;

            float value = 0.0f;
            if (k < reduction_size && position < positions) {
                const int ic = k / kernel_area;
                const int tap = k - ic * kernel_area;
                const int kh = tap / kernel_w;
                const int kw = tap - kh * kernel_w;

                const int oh = position / out_w;
                const int ow = position - oh * 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 ((unsigned)ih < (unsigned)height &&
                    (unsigned)iw < (unsigned)width) {
                    const int64_t offset =
                        (((int64_t)batch * channels + ic) * height + ih)
                        * width + iw;
                    value = x[offset];
                }
            }
            shared_b[index] = value;
        }

        __syncthreads();

        #pragma unroll
        for (int kk = 0; kk < 32; kk += 8) {
            AFragment a0_hi, a0_lo, a1_hi, a1_lo;
            BFragment b_hi, b_lo;

            wmma::load_matrix_sync(
                a0_hi, shared_a + warp_m * 32 + kk, 32);
            wmma::load_matrix_sync(
                a1_hi, shared_a + (warp_m + 16) * 32 + kk, 32);
            wmma::load_matrix_sync(
                b_hi, shared_b + kk * 64 + warp_n, 64);

            // Correct the leading TF32 products with both residual terms.
            #pragma unroll
            for (int i = 0; i < AFragment::num_elements; ++i) {
                const float v0 = a0_hi.x[i];
                const float v1 = a1_hi.x[i];
                const float h0 = __float_to_tf32(v0);
                const float h1 = __float_to_tf32(v1);
                a0_hi.x[i] = h0;
                a1_hi.x[i] = h1;
                a0_lo.x[i] = __float_to_tf32(v0 - h0);
                a1_lo.x[i] = __float_to_tf32(v1 - h1);
            }

            #pragma unroll
            for (int i = 0; i < BFragment::num_elements; ++i) {
                const float value = b_hi.x[i];
                const float high = __float_to_tf32(value);
                b_hi.x[i] = high;
                b_lo.x[i] = __float_to_tf32(value - high);
            }

            wmma::mma_sync(acc0, a0_lo, b_hi, acc0);
            wmma::mma_sync(acc0, a0_hi, b_lo, acc0);
            wmma::mma_sync(acc0, a0_hi, b_hi, acc0);

            wmma::mma_sync(acc1, a1_lo, b_hi, acc1);
            wmma::mma_sync(acc1, a1_hi, b_lo, acc1);
            wmma::mma_sync(acc1, a1_hi, b_hi, acc1);
        }

        __syncthreads();
    }

    // Reuse the input staging memory for coalesced output stores.
    wmma::store_matrix_sync(
        shared + warp_m * 64 + warp_n,
        acc0, 64, wmma::mem_row_major);
    wmma::store_matrix_sync(
        shared + (warp_m + 16) * 64 + warp_n,
        acc1, 64, wmma::mem_row_major);

    __syncthreads();

    for (int index = tid; index < 64 * 64; index += 256) {
        const int oc = channel_base + index / 64;
        const int position = position_base + index % 64;

        if (oc < out_channels && position < positions) {
            float value = shared[index];
            if (bias != nullptr) {
                value += bias[oc];
            }
            output[((int64_t)batch * out_channels + oc) * positions
                   + position] = value;
        }
    }
#endif
}

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)
{
    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.dim() == 4 && weight.dim() == 4,
                "Expected four-dimensional tensors");
    TORCH_CHECK(x.size(1) == weight.size(1),
                "Input channel count does not match weight");

    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 dimensions");

    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 + 63) / 64,
        (weight.size(0) + 63) / 64,
        x.size(0));

    asymmetric_conv_kernel<<<
        grid, 256, 4096 * sizeof(float),
        c10::cuda::getCurrentCUDAStream().stream()>>>(
        x.data_ptr<float>(),
        weight.data_ptr<float>(),
        bias.numel() ? 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);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_EXTENSION = None


def _get_extension():
    global _EXTENSION
    if _EXTENSION is None:
        _EXTENSION = load_inline(
            name="asymmetric_dilated_conv_tf32_residual_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: tuple = (0, 0),
        dilation: tuple = (1, 1),
        bias: bool = False,
    ):
        super().__init__()
        self.conv2d = nn.Conv2d(
            in_channels,
            out_channels,
            kernel_size,
            stride=stride,
            padding=padding,
            dilation=dilation,
            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 torch.cuda.get_device_capability(x.device)[0] < 8
            or isinstance(conv.padding, str)
            or torch.is_autocast_enabled()
        ):
            return conv(x)

        bias = conv.bias
        if bias is None:
            bias = x.new_empty(0)
        else:
            bias = bias.contiguous()

        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],
        )
```