```python
import functools

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,
    int64_t sh, int64_t sw,
    int64_t ph, int64_t pw,
    int64_t dh, int64_t dw);
"""


_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 BM = 64;
constexpr int BN = 64;
constexpr int BK = 32;

union __align__(32) SharedStorage {
    struct {
        float ah[BM * BK];
        float al[BM * BK];
        float bh[BK * BN];
        float bl[BK * BN];
    } operands;
    float result[BM * BN];
};

__global__ void convolution_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 kh, int kw,
    int out_height, int out_width,
    int sh, int sw, int ph, int pw, int dh, int dw)
{
#if __CUDA_ARCH__ >= 800
    using namespace nvcuda;

    __shared__ SharedStorage storage;

    const int tid = threadIdx.x;
    const int warp = tid / 32;
    const int tile_m = blockIdx.y * BM;
    const int tile_n = blockIdx.x * BN;
    const int batch = blockIdx.z;
    const int spatial = out_height * out_width;
    const int kernel_area = kh * kw;
    const int reduction = channels * kernel_area;

    const int warp_m = (warp / 2) * 16;
    const int warp_n = (warp % 2) * 16;

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

    const int input_col = tid % BN;
    const int output_position = tile_n + input_col;
    const int base_h = (output_position / out_width) * sh - ph;
    const int base_w = (output_position % out_width) * sw - pw;

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

    for (int k_base = 0; k_base < reduction; k_base += BK) {
        #pragma unroll
        for (int i = tid; i < BM * BK; i += 256) {
            const int row = i / BK;
            const int k = k_base + i % BK;
            float value = 0.0f;

            if (tile_m + row < out_channels && k < reduction) {
                value = weight[
                    static_cast<int64_t>(tile_m + row) * reduction + k];
            }

            const float high = __float_to_tf32(value);
            storage.operands.ah[i] = high;
            storage.operands.al[i] = __float_to_tf32(value - high);
        }

        #pragma unroll
        for (int i = tid; i < BK * BN; i += 256) {
            const int k = k_base + i / BN;
            float value = 0.0f;

            if (k < reduction && output_position < spatial) {
                const int channel = k / kernel_area;
                const int kernel_position = k % kernel_area;
                const int input_h = base_h + (kernel_position / kw) * dh;
                const int input_w = base_w + (kernel_position % kw) * dw;

                if (static_cast<unsigned>(input_h) <
                        static_cast<unsigned>(height) &&
                    static_cast<unsigned>(input_w) <
                        static_cast<unsigned>(width)) {
                    value = batch_input[
                        (static_cast<int64_t>(channel) * height + input_h)
                        * width + input_w];
                }
            }

            const float high = __float_to_tf32(value);
            storage.operands.bh[i] = high;
            storage.operands.bl[i] = __float_to_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::fragment<
                wmma::matrix_b, 16, 16, 8,
                wmma::precision::tf32, wmma::row_major> bh, bl;

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

            wmma::load_matrix_sync(
                bh, storage.operands.bh + kk * BN + warp_n, BN);
            wmma::load_matrix_sync(
                bl, storage.operands.bl + kk * BN + warp_n, BN);

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

            wmma::load_matrix_sync(
                bh, storage.operands.bh + kk * BN + warp_n + 32, BN);
            wmma::load_matrix_sync(
                bl, storage.operands.bl + kk * BN + warp_n + 32, BN);

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

        __syncthreads();
    }

    wmma::store_matrix_sync(
        storage.result + warp_m * BN + warp_n,
        acc0, BN, wmma::mem_row_major);
    wmma::store_matrix_sync(
        storage.result + warp_m * BN + warp_n + 32,
        acc1, BN, wmma::mem_row_major);

    __syncthreads();

    #pragma unroll
    for (int i = tid; i < BM * BN; i += 256) {
        const int channel = tile_m + i / BN;
        const int position = tile_n + i % BN;

        if (channel < out_channels && position < spatial) {
            float value = storage.result[i];
            if (bias != nullptr) {
                value += bias[channel];
            }
            output[
                (static_cast<int64_t>(batch) * out_channels + channel)
                * spatial + position] = value;
        }
    }
#endif
}

}  // namespace

torch::Tensor asymmetric_conv_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias,
    int64_t sh, int64_t sw,
    int64_t ph, int64_t pw,
    int64_t dh, int64_t dw)
{
    TORCH_CHECK(x.is_cuda() && weight.is_cuda(), "Expected CUDA tensors");
    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(x.device() == weight.device(), "Device mismatch");
    TORCH_CHECK(x.size(1) == weight.size(1), "Channel mismatch");

    const c10::cuda::CUDAGuard device_guard(x.device());
    TORCH_CHECK(
        at::cuda::getCurrentDeviceProperties()->major >= 8,
        "This kernel requires an Ampere or newer GPU");

    const int64_t oh =
        (x.size(2) + 2 * ph - dh * (weight.size(2) - 1) - 1) / sh + 1;
    const int64_t ow =
        (x.size(3) + 2 * pw - dw * (weight.size(3) - 1) - 1) / sw + 1;
    TORCH_CHECK(oh > 0 && ow > 0, "Invalid convolution output size");

    const float* bias_ptr = nullptr;
    if (bias.has_value()) {
        TORCH_CHECK(
            bias->device() == x.device() &&
            bias->scalar_type() == torch::kFloat32 &&
            bias->is_contiguous() &&
            bias->numel() == weight.size(0),
            "Invalid bias");
        bias_ptr = bias->data_ptr<float>();
    }

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

    const dim3 grid(
        (oh * ow + BN - 1) / BN,
        (weight.size(0) + BM - 1) / BM,
        x.size(0));

    convolution_kernel<<<
        grid, 256, 0, at::cuda::getCurrentCUDAStream()>>>(
        x.data_ptr<float>(),
        weight.data_ptr<float>(),
        bias_ptr,
        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);

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""


@functools.lru_cache(maxsize=1)
def _extension():
    return load_inline(
        name="asymmetric_dilated_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,
    )


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.ndim != 4
            or x.dtype != torch.float32
            or conv.weight.dtype != torch.float32
            or torch.is_autocast_enabled()
            or isinstance(conv.padding, str)
        ):
            return conv(x)

        if torch.cuda.get_device_capability(x.device)[0] < 8:
            return conv(x)

        return _extension().asymmetric_conv_cuda(
            x.contiguous(),
            conv.weight.contiguous(),
            conv.bias,
            conv.stride[0],
            conv.stride[1],
            conv.padding[0],
            conv.padding[1],
            conv.dilation[0],
            conv.dilation[1],
        )
```