```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 transposed_conv_cuda(
    torch::Tensor input,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias);
"""

_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 <cuda_fp16.h>
#include <mma.h>

namespace wmma = nvcuda::wmma;

constexpr int C = 64;
constexpr int N_TILE = 128;
constexpr int K_TILE = 32;
constexpr int THREADS = 256;
constexpr int WEIGHT_COUNT = 9 * C * C;

__global__ void pack_weights_kernel(
    const float* __restrict__ weight,
    __half* __restrict__ packed)
{
    const int index = blockIdx.x * blockDim.x + threadIdx.x;
    if (index < WEIGHT_COUNT) {
        const int ci = index % C;
        const int co = (index / C) % C;
        const int tap = index / (C * C);

        const float value = weight[(ci * C + co) * 9 + tap];
        const __half high = __float2half_rn(value);
        packed[index] = high;
        packed[WEIGHT_COUNT + index] =
            __float2half_rn(value - __half2float(high));
    }
}

union SharedStorage {
    struct {
        __half a_high[C * K_TILE];
        __half a_low[C * K_TILE];
        __half b_high[K_TILE * N_TILE];
        __half b_low[K_TILE * N_TILE];
    } operands;
    float result[C * N_TILE];
};

__global__ __launch_bounds__(THREADS)
void transposed_conv_kernel(
    const float* __restrict__ input,
    const __half* __restrict__ packed,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int height,
    int width)
{
#if __CUDA_ARCH__ >= 700
    __shared__ __align__(32) SharedStorage storage;

    __half* a_high = storage.operands.a_high;
    __half* a_low = storage.operands.a_low;
    __half* b_high = storage.operands.b_high;
    __half* b_low = storage.operands.b_low;

    const int tid = threadIdx.x;
    const int warp = tid >> 5;
    const int warp_m = warp >> 2;
    const int warp_n = warp & 3;

    const int batch = blockIdx.y;
    const int output_width = width + 2;
    const int output_area = (height + 2) * output_width;
    const int input_area = height * width;
    const int position_base = blockIdx.x * N_TILE;

    const int position = position_base + (tid % N_TILE);
    const int oy = position / output_width;
    const int ox = position - oy * output_width;
    const float* batch_input =
        input + int64_t(batch) * C * input_area;

    wmma::fragment<wmma::accumulator, 16, 16, 16, float> acc[2][2];

    #pragma unroll
    for (int m = 0; m < 2; ++m) {
        #pragma unroll
        for (int n = 0; n < 2; ++n) {
            wmma::fill_fragment(acc[m][n], 0.0f);
        }
    }

    // Process adjacent horizontal taps consecutively to reuse input cache lines.
    #pragma unroll 1
    for (int kh = 0; kh < 3; ++kh) {
        const int iy = oy - kh;

        #pragma unroll 1
        for (int ci_base = 0; ci_base < C; ci_base += K_TILE) {
            #pragma unroll 1
            for (int kw = 0; kw < 3; ++kw) {
                const int tap = kh * 3 + kw;
                const int ix = ox - kw;
                const bool valid =
                    position < output_area &&
                    static_cast<unsigned>(iy) <
                        static_cast<unsigned>(height) &&
                    static_cast<unsigned>(ix) <
                        static_cast<unsigned>(width);

                #pragma unroll
                for (int index = tid;
                     index < C * K_TILE;
                     index += THREADS) {
                    const int co = index / K_TILE;
                    const int ci = ci_base + index % K_TILE;
                    const int source = (tap * C + co) * C + ci;
                    a_high[index] = packed[source];
                    a_low[index] = packed[WEIGHT_COUNT + source];
                }

                #pragma unroll
                for (int index = tid;
                     index < K_TILE * N_TILE;
                     index += THREADS) {
                    const int ci = ci_base + index / N_TILE;
                    float value = 0.0f;
                    if (valid) {
                        value = batch_input[
                            ci * input_area + iy * width + ix];
                    }

                    const __half high = __float2half_rn(value);
                    b_high[index] = high;
                    b_low[index] =
                        __float2half_rn(value - __half2float(high));
                }

                __syncthreads();

                #pragma unroll
                for (int step = 0; step < K_TILE; step += 16) {
                    wmma::fragment<
                        wmma::matrix_a, 16, 16, 16,
                        __half, wmma::row_major> ah[2], al[2];
                    wmma::fragment<
                        wmma::matrix_b, 16, 16, 16,
                        __half, wmma::row_major> bh[2], bl[2];

                    #pragma unroll
                    for (int m = 0; m < 2; ++m) {
                        const int offset =
                            (warp_m * 32 + m * 16) * K_TILE + step;
                        wmma::load_matrix_sync(
                            ah[m], a_high + offset, K_TILE);
                        wmma::load_matrix_sync(
                            al[m], a_low + offset, K_TILE);
                    }

                    #pragma unroll
                    for (int n = 0; n < 2; ++n) {
                        const int offset =
                            step * N_TILE + warp_n * 32 + n * 16;
                        wmma::load_matrix_sync(
                            bh[n], b_high + offset, N_TILE);
                        wmma::load_matrix_sync(
                            bl[n], b_low + offset, N_TILE);
                    }

                    #pragma unroll
                    for (int m = 0; m < 2; ++m) {
                        #pragma unroll
                        for (int n = 0; n < 2; ++n) {
                            wmma::mma_sync(
                                acc[m][n], al[m], bh[n], acc[m][n]);
                            wmma::mma_sync(
                                acc[m][n], ah[m], bl[n], acc[m][n]);
                            wmma::mma_sync(
                                acc[m][n], ah[m], bh[n], acc[m][n]);
                        }
                    }
                }

                __syncthreads();
            }
        }
    }

    #pragma unroll
    for (int m = 0; m < 2; ++m) {
        #pragma unroll
        for (int n = 0; n < 2; ++n) {
            const int offset =
                (warp_m * 32 + m * 16) * N_TILE +
                warp_n * 32 + n * 16;
            wmma::store_matrix_sync(
                storage.result + offset,
                acc[m][n],
                N_TILE,
                wmma::mem_row_major);
        }
    }

    __syncthreads();

    #pragma unroll
    for (int index = tid; index < C * N_TILE; index += THREADS) {
        const int co = index / N_TILE;
        const int p = position_base + index % N_TILE;
        if (p < output_area) {
            float value = storage.result[index];
            if (bias != nullptr) {
                value += bias[co];
            }
            output[(int64_t(batch) * C + co) * output_area + p] = value;
        }
    }
#endif
}

torch::Tensor transposed_conv_cuda(
    torch::Tensor input,
    torch::Tensor weight,
    c10::optional<torch::Tensor> bias)
{
    TORCH_CHECK(input.is_cuda() && weight.is_cuda(),
                "Expected CUDA tensors");
    TORCH_CHECK(input.device() == weight.device(), "Device mismatch");
    TORCH_CHECK(input.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat,
                "Expected float32 tensors");
    TORCH_CHECK(input.dim() == 4 && input.size(1) == C &&
                input.is_contiguous(),
                "Expected contiguous NCHW input with 64 channels");
    TORCH_CHECK(weight.dim() == 4 &&
                weight.size(0) == C && weight.size(1) == C &&
                weight.size(2) == 3 && weight.size(3) == 3 &&
                weight.is_contiguous(),
                "Expected contiguous [64, 64, 3, 3] weights");

    c10::cuda::CUDAGuard guard(input.device());
    TORCH_CHECK(at::cuda::getCurrentDeviceProperties()->major >= 7,
                "Tensor cores require compute capability 7.0 or newer");

    const float* bias_ptr = nullptr;
    if (bias.has_value()) {
        const auto& b = bias.value();
        TORCH_CHECK(b.device() == input.device() &&
                    b.scalar_type() == at::kFloat &&
                    b.is_contiguous() && b.numel() == C,
                    "Invalid bias");
        bias_ptr = b.data_ptr<float>();
    }

    const int64_t height64 = input.size(2);
    const int64_t width64 = input.size(3);
    TORCH_CHECK(height64 > 0 && width64 > 0 &&
                input.size(0) <= 65535 &&
                height64 * width64 * C <= 2147483647LL &&
                (height64 + 2) * (width64 + 2) <= 2147483647LL,
                "Unsupported input size");

    auto output = torch::empty(
        {input.size(0), C, height64 + 2, width64 + 2},
        input.options());

    if (input.size(0) == 0) {
        return output;
    }

    auto packed = torch::empty(
        {2 * WEIGHT_COUNT}, weight.options().dtype(at::kHalf));

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    pack_weights_kernel<<<
        (WEIGHT_COUNT + THREADS - 1) / THREADS,
        THREADS, 0, stream>>>(
            weight.data_ptr<float>(),
            reinterpret_cast<__half*>(packed.data_ptr<at::Half>()));
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    const int height = static_cast<int>(height64);
    const int width = static_cast<int>(width64);
    const int64_t output_area = (height64 + 2) * (width64 + 2);

    dim3 grid(
        static_cast<unsigned>((output_area + N_TILE - 1) / N_TILE),
        static_cast<unsigned>(input.size(0)));

    transposed_conv_kernel<<<grid, THREADS, 0, stream>>>(
        input.data_ptr<float>(),
        reinterpret_cast<const __half*>(packed.data_ptr<at::Half>()),
        bias_ptr,
        output.data_ptr<float>(),
        height,
        width);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    return output;
}
"""

_extension = load_inline(
    name="transposed_conv2d_half_residual_64x128_v4",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["transposed_conv_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class _TransposedConv(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x, weight, bias):
        ctx.save_for_backward(x, weight)
        ctx.has_bias = bias is not None
        return _extension.transposed_conv_cuda(x, weight, bias)

    @staticmethod
    def backward(ctx, grad_output):
        x, weight = ctx.saved_tensors
        grad_x = grad_weight = grad_bias = None

        if ctx.needs_input_grad[0]:
            grad_x = F.conv2d(grad_output, weight)

        if ctx.needs_input_grad[1]:
            grad_weight = torch.nn.grad.conv2d_weight(
                grad_output,
                weight.shape,
                x,
                stride=1,
                padding=0,
                dilation=1,
                groups=1,
            )

        if ctx.has_bias and ctx.needs_input_grad[2]:
            grad_bias = grad_output.sum(dim=(0, 2, 3))

        return grad_x, grad_weight, grad_bias


class ModelNew(nn.Module):
    def __init__(
        self,
        in_channels: int,
        out_channels: int,
        kernel_size: int,
        stride: int = 1,
        padding: int = 0,
        output_padding: int = 0,
        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,
            groups=groups,
            bias=bias,
        )

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        conv = self.conv_transpose2d

        if (
            x.is_cuda
            and x.ndim == 4
            and x.dtype == torch.float32
            and x.is_contiguous()
            and x.shape[1] == 64
            and x.shape[0] <= 65535
            and x.shape[2] > 0
            and x.shape[3] > 0
            and x.shape[2] * x.shape[3] * 64 <= 2147483647
            and (x.shape[2] + 2) * (x.shape[3] + 2) <= 2147483647
            and conv.in_channels == 64
            and conv.out_channels == 64
            and conv.kernel_size == (3, 3)
            and conv.stride == (1, 1)
            and conv.padding == (0, 0)
            and conv.output_padding == (0, 0)
            and conv.dilation == (1, 1)
            and conv.groups == 1
            and conv.weight.device == x.device
            and conv.weight.dtype == torch.float32
            and conv.weight.is_contiguous()
            and (
                conv.bias is None
                or (
                    conv.bias.device == x.device
                    and conv.bias.dtype == torch.float32
                    and conv.bias.is_contiguous()
                )
            )
            and not torch.is_autocast_enabled()
            and torch.cuda.get_device_capability(x.device)[0] >= 7
        ):
            return _TransposedConv.apply(x, conv.weight, conv.bias)

        return conv(x)
```