```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 transpose_conv_weight_cuda(
    torch::Tensor weight, int64_t groups);

torch::Tensor transpose_conv_weight_backward_cuda(
    torch::Tensor grad, int64_t groups);
"""

_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>

template <typename scalar_t, bool BACKWARD>
__global__ void transform_weight_kernel(
    const scalar_t* __restrict__ src,
    scalar_t* __restrict__ dst,
    int64_t count,
    int64_t input_channels_per_group,
    int64_t output_channels_per_group,
    int64_t kernel_h,
    int64_t kernel_w)
{
    for (int64_t index =
             static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
         index < count;
         index += static_cast<int64_t>(blockDim.x) * gridDim.x)
    {
        // Destination forward-convolution weight has physical OHWI layout.
        int64_t remaining = index;
        const int64_t ic = remaining % input_channels_per_group;
        remaining /= input_channels_per_group;
        const int64_t kw = remaining % kernel_w;
        remaining /= kernel_w;
        const int64_t kh = remaining % kernel_h;
        const int64_t oc = remaining / kernel_h;

        const int64_t group = oc / output_channels_per_group;
        const int64_t local_oc = oc % output_channels_per_group;
        const int64_t original_ic =
            group * input_channels_per_group + ic;

        const int64_t original_index =
            ((original_ic * output_channels_per_group + local_oc)
                * kernel_h + (kernel_h - 1 - kh))
                * kernel_w + (kernel_w - 1 - kw);

        if (BACKWARD) {
            dst[original_index] = src[index];
        } else {
            dst[index] = src[original_index];
        }
    }
}

torch::Tensor transpose_conv_weight_cuda(
    torch::Tensor weight, int64_t groups)
{
    TORCH_CHECK(weight.is_cuda(), "weight must be a CUDA tensor");
    TORCH_CHECK(weight.dim() == 4, "weight must have four dimensions");
    TORCH_CHECK(groups > 0 && weight.size(0) % groups == 0,
                "invalid groups");

    const c10::cuda::CUDAGuard guard(weight.device());
    weight = weight.contiguous();

    const int64_t icg = weight.size(0) / groups;
    const int64_t ocg = weight.size(1);
    const int64_t kh = weight.size(2);
    const int64_t kw = weight.size(3);

    auto output = torch::empty(
        {ocg * groups, icg, kh, kw},
        weight.options().memory_format(at::MemoryFormat::ChannelsLast));

    const int64_t count = weight.numel();
    if (count == 0) {
        return output;
    }

    constexpr int threads = 256;
    const int blocks = static_cast<int>(
        std::min<int64_t>((count + threads - 1) / threads, 4096));
    const auto stream = at::cuda::getCurrentCUDAStream();

    AT_DISPATCH_FLOATING_TYPES_AND2(
        at::ScalarType::Half, at::ScalarType::BFloat16,
        weight.scalar_type(), "transpose_conv_weight_cuda", [&] {
            transform_weight_kernel<scalar_t, false>
                <<<blocks, threads, 0, stream>>>(
                    weight.data_ptr<scalar_t>(),
                    output.data_ptr<scalar_t>(),
                    count, icg, ocg, kh, kw);
        });

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}

torch::Tensor transpose_conv_weight_backward_cuda(
    torch::Tensor grad, int64_t groups)
{
    TORCH_CHECK(grad.is_cuda(), "gradient must be a CUDA tensor");
    TORCH_CHECK(grad.dim() == 4, "gradient must have four dimensions");
    TORCH_CHECK(groups > 0 && grad.size(0) % groups == 0,
                "invalid groups");

    const c10::cuda::CUDAGuard guard(grad.device());
    grad = grad.contiguous(at::MemoryFormat::ChannelsLast);

    const int64_t icg = grad.size(1);
    const int64_t ocg = grad.size(0) / groups;
    const int64_t kh = grad.size(2);
    const int64_t kw = grad.size(3);

    auto output = torch::empty(
        {icg * groups, ocg, kh, kw},
        grad.options().memory_format(at::MemoryFormat::Contiguous));

    const int64_t count = grad.numel();
    if (count == 0) {
        return output;
    }

    constexpr int threads = 256;
    const int blocks = static_cast<int>(
        std::min<int64_t>((count + threads - 1) / threads, 4096));
    const auto stream = at::cuda::getCurrentCUDAStream();

    AT_DISPATCH_FLOATING_TYPES_AND2(
        at::ScalarType::Half, at::ScalarType::BFloat16,
        grad.scalar_type(), "transpose_conv_weight_backward_cuda", [&] {
            transform_weight_kernel<scalar_t, true>
                <<<blocks, threads, 0, stream>>>(
                    grad.data_ptr<scalar_t>(),
                    output.data_ptr<scalar_t>(),
                    count, icg, ocg, kh, kw);
        });

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_extension = None


def _get_extension():
    global _extension
    if _extension is None:
        _extension = load_inline(
            name="asymmetric_transpose_conv_weight_v1",
            cpp_sources=_cpp_source,
            cuda_sources=_cuda_source,
            functions=[
                "transpose_conv_weight_cuda",
                "transpose_conv_weight_backward_cuda",
            ],
            extra_cflags=["-O3"],
            extra_cuda_cflags=["-O3"],
            verbose=False,
        )
    return _extension


class _TransformWeight(torch.autograd.Function):
    @staticmethod
    def forward(ctx, weight, groups):
        ctx.groups = groups
        return _get_extension().transpose_conv_weight_cuda(weight, groups)

    @staticmethod
    def backward(ctx, grad_output):
        return (
            _get_extension().transpose_conv_weight_backward_cuda(
                grad_output, ctx.groups
            ),
            None,
        )


class ModelNew(nn.Module):
    def __init__(
        self,
        in_channels: int,
        out_channels: int,
        kernel_size: tuple,
        stride: tuple = (1, 1),
        padding: tuple = (0, 0),
        output_padding: tuple = (0, 0),
        dilation: tuple = (1, 1),
        groups: int = 1,
        bias: bool = False,
    ):
        super().__init__()
        self.conv_transpose2d = nn.ConvTranspose2d(
            in_channels,
            out_channels,
            kernel_size,
            stride=stride,
            padding=padding,
            output_padding=output_padding,
            dilation=dilation,
            groups=groups,
            bias=bias,
        )

        conv = self.conv_transpose2d
        self._forward_padding = tuple(
            d * (k - 1) - p
            for d, k, p in zip(
                conv.dilation, conv.kernel_size, conv.padding
            )
        )
        self._use_forward_convolution = (
            conv.stride == (1, 1)
            and conv.output_padding == (0, 0)
            and min(self._forward_padding) >= 0
        )

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        conv = self.conv_transpose2d
        if (
            not self._use_forward_convolution
            or not x.is_cuda
            or x.ndim != 4
            or conv.weight.dtype
            not in (torch.float16, torch.bfloat16, torch.float32, torch.float64)
        ):
            return conv(x)

        # At stride one, transposed convolution equals forward convolution
        # with exchanged channel axes, reversed spatial axes, and full padding.
        weight = _TransformWeight.apply(conv.weight, conv.groups)
        x = x.contiguous(memory_format=torch.channels_last)

        return F.conv2d(
            x,
            weight,
            conv.bias,
            stride=1,
            padding=self._forward_padding,
            dilation=conv.dilation,
            groups=conv.groups,
        )
```