```python
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline


_CPP_SOURCE = r"""
#include <torch/extension.h>

torch::Tensor transposed_conv3d_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    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_runtime.h>
#include <mma.h>

namespace {

__global__ void pack_input(
    const float* __restrict__ src,
    float* __restrict__ dst,
    int spatial,
    int64_t count)
{
    for (int64_t i = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
         i < count;
         i += int64_t(blockDim.x) * gridDim.x) {
        int c = i % 32;
        int64_t position = i / 32;
        int s = position % spatial;
        int n = position / spatial;
        dst[i] = src[(int64_t(n) * 32 + c) * spatial + s];
    }
}

__global__ void pack_weight(
    const float* __restrict__ src,
    float* __restrict__ dst)
{
    constexpr int count = 75 * 32 * 64;
    for (int i = blockIdx.x * blockDim.x + threadIdx.x;
         i < count;
         i += blockDim.x * gridDim.x) {
        int co = i % 64;
        int ci = (i / 64) % 32;
        int r = i / (64 * 32);
        dst[i] = src[(ci * 64 + co) * 75 + r];
    }
}

// Each block computes 32 output positions and all 64 output channels.
// TF32 high/residual products retain substantially more precision than
// a single TF32 multiplication.
__global__ void transposed_conv_kernel(
    const float* __restrict__ x,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int depth,
    int height,
    int width,
    int output_positions)
{
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
    using namespace nvcuda;

    __shared__ __align__(32) float a_hi[32 * 32];
    __shared__ __align__(32) float a_lo[32 * 32];
    __shared__ __align__(32) float b_hi[32 * 64];
    __shared__ __align__(32) float b_lo[32 * 64];
    __shared__ __align__(32) float result[32 * 64];

    const int tid = threadIdx.x;
    const int warp = tid / 32;
    const int channel = tid % 32;
    const int tile_start = blockIdx.x * 32;
    const int out_depth = depth + 2;
    const int out_height = height + 4;
    const int out_width = width + 4;

    int row_n[4], row_d[4], row_h[4], row_w[4];
    bool row_valid[4];

    #pragma unroll
    for (int j = 0; j < 4; ++j) {
        int row = tid / 32 + j * 8;
        int position = tile_start + row;
        row_valid[j] = position < output_positions;
        row_w[j] = position % out_width;
        position /= out_width;
        row_h[j] = position % out_height;
        position /= out_height;
        row_d[j] = position % out_depth;
        row_n[j] = position / out_depth;
    }

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

    for (int r = 0; r < 75; ++r) {
        int kd = r / 25;
        int kh = (r / 5) % 5;
        int kw = r % 5;

        #pragma unroll
        for (int j = 0; j < 4; ++j) {
            int id = row_d[j] - kd;
            int ih = row_h[j] - kh;
            int iw = row_w[j] - kw;

            float value = 0.0f;
            if (row_valid[j] &&
                unsigned(id) < unsigned(depth) &&
                unsigned(ih) < unsigned(height) &&
                unsigned(iw) < unsigned(width)) {
                int64_t offset =
                    (((int64_t(row_n[j]) * depth + id) * height + ih)
                     * width + iw) * 32 + channel;
                value = x[offset];
            }

            int index = tid + j * 256;
            float hi = __float_to_tf32(value);
            a_hi[index] = hi;
            a_lo[index] = __float_to_tf32(value - hi);
        }

        #pragma unroll
        for (int j = 0; j < 8; ++j) {
            int index = tid + j * 256;
            float value = weight[r * 32 * 64 + index];
            float hi = __float_to_tf32(value);
            b_hi[index] = hi;
            b_lo[index] = __float_to_tf32(value - hi);
        }

        __syncthreads();

        #pragma unroll
        for (int k = 0; k < 32; k += 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;

            int a_offset = (warp / 4) * 16 * 32 + k;
            int b_offset = k * 64 + (warp % 4) * 16;

            wmma::load_matrix_sync(ah, a_hi + a_offset, 32);
            wmma::load_matrix_sync(al, a_lo + a_offset, 32);
            wmma::load_matrix_sync(bh, b_hi + b_offset, 64);
            wmma::load_matrix_sync(bl, b_lo + b_offset, 64);

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

        __syncthreads();
    }

    wmma::store_matrix_sync(
        result + (warp / 4) * 16 * 64 + (warp % 4) * 16,
        acc, 64, wmma::mem_row_major);

    __syncthreads();

    #pragma unroll
    for (int j = 0; j < 8; ++j) {
        int index = tid + j * 256;
        int position = tile_start + index / 64;
        int co = index % 64;
        if (position < output_positions) {
            float value = result[index];
            if (bias != nullptr)
                value += bias[co];
            output[int64_t(position) * 64 + co] = value;
        }
    }
#endif
}

} // namespace

torch::Tensor transposed_conv3d_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    torch::Tensor bias)
{
    c10::cuda::CUDAGuard guard(x.device());

    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() == 5 && x.size(1) == 32, "Invalid input shape");
    TORCH_CHECK(
        weight.dim() == 5 &&
        weight.size(0) == 32 && weight.size(1) == 64 &&
        weight.size(2) == 3 && weight.size(3) == 5 &&
        weight.size(4) == 5,
        "Invalid weight shape");
    TORCH_CHECK(x.device() == weight.device(), "Device mismatch");

    auto input = x.contiguous();
    auto w = weight.contiguous();
    auto b = bias.contiguous();

    const int n = input.size(0);
    const int d = input.size(2);
    const int h = input.size(3);
    const int width = input.size(4);

    auto packed_input = torch::empty(
        {n, d, h, width, 32}, input.options());
    auto packed_weight = torch::empty(
        {75, 32, 64}, w.options());
    auto output = torch::empty(
        {n, 64, d + 2, h + 4, width + 4},
        input.options().memory_format(at::MemoryFormat::ChannelsLast3d));

    cudaStream_t stream =
        at::cuda::getCurrentCUDAStream(input.get_device());

    int64_t count = input.numel();
    int blocks = int((count + 255) / 256);
    if (blocks > 4096)
        blocks = 4096;

    pack_input<<<blocks, 256, 0, stream>>>(
        input.data_ptr<float>(), packed_input.data_ptr<float>(),
        d * h * width, count);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    pack_weight<<<256, 256, 0, stream>>>(
        w.data_ptr<float>(), packed_weight.data_ptr<float>());
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    int positions = n * (d + 2) * (h + 4) * (width + 4);
    transposed_conv_kernel<<<(positions + 31) / 32, 256, 0, stream>>>(
        packed_input.data_ptr<float>(),
        packed_weight.data_ptr<float>(),
        b.numel() ? b.data_ptr<float>() : nullptr,
        output.data_ptr<float>(),
        d, h, width, positions);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    return output;
}
"""


_extension = None


def _get_extension():
    global _extension
    if _extension is None:
        _extension = load_inline(
            name="custom_transpose3d_tf32_residual_v1",
            cpp_sources=_CPP_SOURCE,
            cuda_sources=_CUDA_SOURCE,
            functions=["transposed_conv3d_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: tuple = (1, 1, 1),
        padding: tuple = (0, 0, 0),
        output_padding: tuple = (0, 0, 0),
        groups: int = 1,
        bias: bool = False,
    ):
        super().__init__()
        self.conv_transpose3d = nn.ConvTranspose3d(
            in_channels,
            out_channels,
            kernel_size,
            stride=stride,
            padding=padding,
            output_padding=output_padding,
            groups=groups,
            bias=bias,
        )
        conv = self.conv_transpose3d
        self._specialized = (
            in_channels == 32
            and out_channels == 64
            and conv.kernel_size == (3, 5, 5)
            and conv.stride == (1, 1, 1)
            and conv.padding == (0, 0, 0)
            and conv.output_padding == (0, 0, 0)
            and groups == 1
        )

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

        if (
            not self._specialized
            or not x.is_cuda
            or x.dtype != torch.float32
            or conv.weight.dtype != torch.float32
            or x.ndim != 5
            or x.numel() == 0
            or torch.is_autocast_enabled()
            or torch.cuda.get_device_capability(x.device)[0] < 8
        ):
            return conv(x)

        bias = conv.bias
        if bias is None:
            bias = conv.weight.new_empty(0)

        return _get_extension().transposed_conv3d_cuda(
            x, conv.weight, bias
        )
```