```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,
    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_fp16.h>
#include <cuda_runtime.h>
#include <mma.h>

namespace {

constexpr int BM = 32;
constexpr int BN = 64;
constexpr int BK = 32;

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

template<int FixedKH, int FixedKW>
__global__ void implicit_conv(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int ci, int co, int height, int width,
    int kernel_h, int kernel_w,
    int out_h, int out_w,
    int sh, int sw, int ph, int pw, int dh, int dw)
{
#if __CUDA_ARCH__ >= 700
    using namespace nvcuda;

    __shared__ SharedStorage shared;

    const int kh = FixedKH ? FixedKH : kernel_h;
    const int kw = FixedKW ? FixedKW : kernel_w;
    const int kernel_area = kh * kw;
    const int reduction_size = ci * kernel_area;
    const int spatial_size = out_h * out_w;

    const int tid = threadIdx.x;
    const int warp = tid / 32;
    const int warp_m = warp / 4;
    const int warp_n = warp % 4;
    const int channel_base = blockIdx.y * BM;
    const int spatial_base = blockIdx.x * BN;
    const int batch = blockIdx.z;

    wmma::fragment<wmma::accumulator, 16, 16, 16, float> accum;
    wmma::fill_fragment(accum, 0.0f);

    for (int kb = 0; kb < reduction_size; kb += BK) {
        for (int i = tid; i < BM * BK; i += blockDim.x) {
            const int oc = channel_base + i / BK;
            const int k = kb + i % BK;
            float value = 0.0f;
            if (oc < co && k < reduction_size) {
                value = weight[
                    static_cast<int64_t>(oc) * reduction_size + k];
            }
            const half high = __float2half_rn(value);
            shared.operands.ah[i] = high;
            shared.operands.al[i] =
                __float2half_rn(value - __half2float(high));
        }

        for (int i = tid; i < BK * BN; i += blockDim.x) {
            const int k = kb + i / BN;
            const int position = spatial_base + i % BN;
            float value = 0.0f;

            if (k < reduction_size && position < spatial_size) {
                const int ic = k / kernel_area;
                const int kernel_position = k % kernel_area;
                const int ky = kernel_position / kw;
                const int kx = kernel_position % kw;

                const int iy = (position / out_w) * sh - ph + ky * dh;
                const int ix = (position % out_w) * sw - pw + kx * dw;

                if (static_cast<unsigned>(iy) <
                        static_cast<unsigned>(height) &&
                    static_cast<unsigned>(ix) <
                        static_cast<unsigned>(width)) {
                    const int64_t offset =
                        ((static_cast<int64_t>(batch) * ci + ic) *
                         height + iy) * width + ix;
                    value = x[offset];
                }
            }

            const half high = __float2half_rn(value);
            shared.operands.bh[i] = high;
            shared.operands.bl[i] =
                __float2half_rn(value - __half2float(high));
        }

        __syncthreads();

        #pragma unroll
        for (int kk = 0; kk < BK; 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::row_major
            > bh, bl;

            const int a_offset = warp_m * 16 * BK + kk;
            const int b_offset = kk * BN + warp_n * 16;

            wmma::load_matrix_sync(
                ah, shared.operands.ah + a_offset, BK);
            wmma::load_matrix_sync(
                al, shared.operands.al + a_offset, BK);
            wmma::load_matrix_sync(
                bh, shared.operands.bh + b_offset, BN);
            wmma::load_matrix_sync(
                bl, shared.operands.bl + b_offset, BN);

            // Two-component FP16 representation with FP32 accumulation.
            wmma::mma_sync(accum, ah, bh, accum);
            wmma::mma_sync(accum, ah, bl, accum);
            wmma::mma_sync(accum, al, bh, accum);
        }

        __syncthreads();
    }

    wmma::store_matrix_sync(
        shared.output + warp_m * 16 * BN + warp_n * 16,
        accum, BN, wmma::mem_row_major);

    __syncthreads();

    for (int i = tid; i < BM * BN; i += blockDim.x) {
        const int oc = channel_base + i / BN;
        const int position = spatial_base + i % BN;
        if (oc < co && position < spatial_size) {
            float value = shared.output[i];
            if (bias != nullptr) {
                value += bias[oc];
            }
            output[
                (static_cast<int64_t>(batch) * co + oc) *
                spatial_size + 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() == 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.device() == weight.device(),
                "Input and weight must share a device");
    TORCH_CHECK(x.size(1) == weight.size(1),
                "Input channel count does not match weight");

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

    const auto* properties = at::cuda::getCurrentDeviceProperties();
    TORCH_CHECK(properties->major >= 7,
                "This kernel requires compute capability 7.0 or newer");

    x = x.contiguous();
    weight = weight.contiguous();

    const int batch = x.size(0);
    const int ci = x.size(1);
    const int height = x.size(2);
    const int width = x.size(3);
    const int co = weight.size(0);
    const int kh = weight.size(2);
    const int kw = weight.size(3);

    TORCH_CHECK(sh > 0 && sw > 0 && dh > 0 && dw > 0,
                "Stride and dilation must be positive");
    TORCH_CHECK(ph >= 0 && pw >= 0, "Padding must be nonnegative");

    const int64_t effective_h = height + 2 * ph - dh * (kh - 1) - 1;
    const int64_t effective_w = width + 2 * pw - dw * (kw - 1) - 1;
    TORCH_CHECK(effective_h >= 0 && effective_w >= 0,
                "Kernel exceeds padded input");

    const int out_h = effective_h / sh + 1;
    const int out_w = effective_w / sw + 1;

    auto output = torch::empty(
        {batch, co, out_h, out_w}, x.options());

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

    torch::Tensor contiguous_bias;
    const float* bias_ptr = nullptr;
    if (bias.has_value()) {
        TORCH_CHECK(bias->device() == x.device() &&
                    bias->scalar_type() == at::kFloat &&
                    bias->numel() == co,
                    "Invalid bias");
        contiguous_bias = bias->contiguous();
        bias_ptr = contiguous_bias.data_ptr<float>();
    }

    const dim3 grid(
        (out_h * out_w + BN - 1) / BN,
        (co + BM - 1) / BM,
        batch);
    const cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if (kh == 5 && kw == 9) {
        implicit_conv<5, 9><<<grid, 256, 0, stream>>>(
            x.data_ptr<float>(), weight.data_ptr<float>(),
            bias_ptr, output.data_ptr<float>(),
            ci, co, height, width, kh, kw, out_h, out_w,
            sh, sw, ph, pw, dh, dw);
    } else {
        implicit_conv<0, 0><<<grid, 256, 0, stream>>>(
            x.data_ptr<float>(), weight.data_ptr<float>(),
            bias_ptr, output.data_ptr<float>(),
            ci, co, height, width, kh, kw, out_h, out_w,
            sh, sw, ph, pw, dh, dw);
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = load_inline(
    name="asymmetric_dilated_implicit_conv_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.dtype != torch.float32
            or conv.weight.dtype != torch.float32
            or torch.is_autocast_enabled()
            or torch.cuda.get_device_capability(x.device)[0] < 7
        ):
            return F.conv2d(
                x,
                conv.weight,
                conv.bias,
                conv.stride,
                conv.padding,
                conv.dilation,
            )

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