```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 to_channels_last_64(torch::Tensor x);

torch::Tensor fused_relu_group_norm_128(
    torch::Tensor x,
    torch::Tensor weight,
    torch::Tensor bias,
    double eps);
"""

_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>

namespace {

constexpr int THREADS = 256;
constexpr int CHANNELS = 128;
constexpr int GROUPS = 8;
constexpr int CHANNEL_VECTORS = 32;
constexpr int SPATIAL_TILE = 128;
constexpr int OUTPUT_VECTORS = 2048;

__device__ __forceinline__ float warp_sum(float v) {
    #pragma unroll
    for (int offset = 16; offset > 0; offset >>= 1)
        v += __shfl_down_sync(0xffffffff, v, offset);
    return v;
}

__device__ __forceinline__ float4 relu4(float4 v) {
    v.x = v.x < 0.0f ? 0.0f : v.x;
    v.y = v.y < 0.0f ? 0.0f : v.y;
    v.z = v.z < 0.0f ? 0.0f : v.z;
    v.w = v.w < 0.0f ? 0.0f : v.w;
    return v;
}

__global__ __launch_bounds__(THREADS)
void channels_last_64_kernel(
    const float* __restrict__ input,
    float* __restrict__ output,
    int spatial)
{
    __shared__ float tile[64][33];

    const int lane = threadIdx.x & 31;
    const int warp = threadIdx.x >> 5;
    const int start = blockIdx.x * 32;
    const int64_t base = static_cast<int64_t>(blockIdx.y) * 64 * spatial;

    #pragma unroll
    for (int j = 0; j < 8; ++j) {
        const int c = warp + j * 8;
        if (start + lane < spatial)
            tile[c][lane] = input[base + static_cast<int64_t>(c) * spatial
                                 + start + lane];
    }
    __syncthreads();

    #pragma unroll
    for (int j = 0; j < 4; ++j) {
        const int s = warp + j * 8;
        if (start + s < spatial) {
            const int64_t index =
                base + static_cast<int64_t>(start + s) * 64 + lane;
            output[index] = tile[lane][s];
            output[index + 32] = tile[lane + 32][s];
        }
    }
}

__global__ __launch_bounds__(THREADS)
void partial_stats_kernel(
    const float4* __restrict__ x,
    float2* __restrict__ partial,
    int spatial,
    int chunks)
{
    __shared__ float shared_sum[64];
    __shared__ float shared_square[64];

    const int tid = threadIdx.x;
    const int channel_vector = tid & 31;
    const int spatial_lane = tid >> 5;
    const int n = blockIdx.y;
    const int chunk = blockIdx.x;
    const int64_t base =
        static_cast<int64_t>(n) * spatial * CHANNEL_VECTORS;

    float sum = 0.0f;
    float square = 0.0f;

    #pragma unroll
    for (int j = 0; j < SPATIAL_TILE / 8; ++j) {
        const int s = chunk * SPATIAL_TILE + spatial_lane + j * 8;
        if (s < spatial) {
            const float4 v = relu4(
                x[base + static_cast<int64_t>(s) * CHANNEL_VECTORS
                  + channel_vector]);
            sum += (v.x + v.y) + (v.z + v.w);
            square += (v.x * v.x + v.y * v.y)
                    + (v.z * v.z + v.w * v.w);
        }
    }

    #pragma unroll
    for (int offset = 2; offset > 0; offset >>= 1) {
        sum += __shfl_down_sync(0xffffffff, sum, offset, 4);
        square += __shfl_down_sync(0xffffffff, square, offset, 4);
    }

    if ((tid & 3) == 0) {
        shared_sum[tid >> 2] = sum;
        shared_square[tid >> 2] = square;
    }
    __syncthreads();

    if (tid < 64) {
        const int group = tid >> 3;
        const int lane = tid & 7;
        sum = shared_sum[lane * GROUPS + group];
        square = shared_square[lane * GROUPS + group];

        #pragma unroll
        for (int offset = 4; offset > 0; offset >>= 1) {
            sum += __shfl_down_sync(0xffffffff, sum, offset, 8);
            square += __shfl_down_sync(0xffffffff, square, offset, 8);
        }

        if (lane == 0)
            partial[(static_cast<int64_t>(n) * GROUPS + group) * chunks
                    + chunk] = make_float2(sum, square);
    }
}

__global__ __launch_bounds__(THREADS)
void finish_stats_kernel(
    const float2* __restrict__ partial,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ scale,
    float* __restrict__ shift,
    int chunks,
    int spatial,
    float eps)
{
    __shared__ float shared_sum[8];
    __shared__ float shared_square[8];

    const int ng = blockIdx.x;
    const int tid = threadIdx.x;
    const int lane = tid & 31;
    const int warp = tid >> 5;

    float sum = 0.0f;
    float square = 0.0f;
    for (int i = tid; i < chunks; i += THREADS) {
        const float2 p = partial[static_cast<int64_t>(ng) * chunks + i];
        sum += p.x;
        square += p.y;
    }

    sum = warp_sum(sum);
    square = warp_sum(square);

    if (lane == 0) {
        shared_sum[warp] = sum;
        shared_square[warp] = square;
    }
    __syncthreads();

    if (warp == 0) {
        sum = warp_sum(lane < 8 ? shared_sum[lane] : 0.0f);
        square = warp_sum(lane < 8 ? shared_square[lane] : 0.0f);

        sum = __shfl_sync(0xffffffff, sum, 0);
        square = __shfl_sync(0xffffffff, square, 0);

        if (lane < 16) {
            const float reciprocal = 1.0f / (static_cast<float>(spatial) * 16.0f);
            const float mean = sum * reciprocal;
            const float variance = fmaxf(square * reciprocal - mean * mean, 0.0f);
            const float rstd = rsqrtf(variance + eps);
            const int channel = (ng % GROUPS) * 16 + lane;
            const float a = rstd * weight[channel];

            scale[ng * 16 + lane] = a;
            shift[ng * 16 + lane] = fmaf(-mean, a, bias[channel]);
        }
    }
}

__global__ __launch_bounds__(THREADS)
void normalize_kernel(
    const float4* __restrict__ x,
    float4* __restrict__ output,
    const float4* __restrict__ scale,
    const float4* __restrict__ shift,
    int vectors_per_batch)
{
    const int tid = threadIdx.x;
    const int n = blockIdx.y;
    const int channel_vector = tid & 31;

    const float4 a = scale[n * CHANNEL_VECTORS + channel_vector];
    const float4 b = shift[n * CHANNEL_VECTORS + channel_vector];
    const int64_t base = static_cast<int64_t>(n) * vectors_per_batch;
    const int start = blockIdx.x * OUTPUT_VECTORS + tid;

    #pragma unroll
    for (int j = 0; j < OUTPUT_VECTORS / THREADS; ++j) {
        const int i = start + j * THREADS;
        if (i < vectors_per_batch) {
            float4 v = relu4(x[base + i]);
            v.x = fmaf(v.x, a.x, b.x);
            v.y = fmaf(v.y, a.y, b.y);
            v.z = fmaf(v.z, a.z, b.z);
            v.w = fmaf(v.w, a.w, b.w);
            output[base + i] = v;
        }
    }
}

} // namespace

torch::Tensor to_channels_last_64(torch::Tensor x)
{
    TORCH_CHECK(x.is_cuda() && x.scalar_type() == at::kFloat &&
                x.dim() == 5 && x.size(1) == 64 && x.is_contiguous(),
                "Expected contiguous CUDA float32 NCDHW with 64 channels");
    const c10::cuda::CUDAGuard guard(x.device());

    auto output = torch::empty(
        x.sizes(), x.options().memory_format(at::MemoryFormat::ChannelsLast3d));
    if (x.numel() == 0)
        return output;

    const int spatial = static_cast<int>(x.size(2) * x.size(3) * x.size(4));
    const cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    channels_last_64_kernel<<<
        dim3((spatial + 31) / 32, static_cast<unsigned int>(x.size(0))),
        THREADS, 0, stream>>>(
        x.data_ptr<float>(), output.data_ptr<float>(), spatial);
    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}

torch::Tensor fused_relu_group_norm_128(
    torch::Tensor x,
    torch::Tensor weight,
    torch::Tensor bias,
    double eps)
{
    TORCH_CHECK(x.is_cuda() && x.scalar_type() == at::kFloat &&
                x.dim() == 5 && x.size(1) == CHANNELS &&
                x.is_contiguous(at::MemoryFormat::ChannelsLast3d),
                "Expected channels-last CUDA float32 input with 128 channels");
    TORCH_CHECK(weight.is_cuda() && bias.is_cuda() &&
                weight.device() == x.device() && bias.device() == x.device() &&
                weight.scalar_type() == at::kFloat &&
                bias.scalar_type() == at::kFloat &&
                weight.numel() == CHANNELS && bias.numel() == CHANNELS &&
                weight.is_contiguous() && bias.is_contiguous(),
                "Invalid group normalization parameters");

    const c10::cuda::CUDAGuard guard(x.device());
    const int batch = static_cast<int>(x.size(0));
    const int spatial = static_cast<int>(x.size(2) * x.size(3) * x.size(4));
    const int chunks = (spatial + SPATIAL_TILE - 1) / SPATIAL_TILE;
    const int vectors_per_batch = spatial * CHANNEL_VECTORS;

    auto output = torch::empty_like(x);
    auto partial = torch::empty({batch * GROUPS, chunks, 2}, x.options());
    auto affine = torch::empty({2, batch, CHANNELS}, x.options());
    float* scale = affine.data_ptr<float>();
    float* shift = scale + static_cast<int64_t>(batch) * CHANNELS;

    const cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    partial_stats_kernel<<<dim3(chunks, batch), THREADS, 0, stream>>>(
        reinterpret_cast<const float4*>(x.data_ptr<float>()),
        reinterpret_cast<float2*>(partial.data_ptr<float>()),
        spatial, chunks);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    finish_stats_kernel<<<batch * GROUPS, THREADS, 0, stream>>>(
        reinterpret_cast<const float2*>(partial.data_ptr<float>()),
        weight.data_ptr<float>(), bias.data_ptr<float>(),
        scale, shift, chunks, spatial, static_cast<float>(eps));
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    normalize_kernel<<<
        dim3((vectors_per_batch + OUTPUT_VECTORS - 1) / OUTPUT_VECTORS, batch),
        THREADS, 0, stream>>>(
        reinterpret_cast<const float4*>(x.data_ptr<float>()),
        reinterpret_cast<float4*>(output.data_ptr<float>()),
        reinterpret_cast<const float4*>(scale),
        reinterpret_cast<const float4*>(shift),
        vectors_per_batch);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    return output;
}
"""

_fused_ops = load_inline(
    name="relu_groupnorm_ndhwc_tiled_affine_v4",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["to_channels_last_64", "fused_relu_group_norm_128"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class ModelNew(nn.Module):
    def __init__(
        self, in_channels, out_channels, kernel_size, groups, bias=False
    ):
        super().__init__()
        self.conv_transpose = nn.ConvTranspose3d(
            in_channels, out_channels, kernel_size, bias=bias
        ).to(memory_format=torch.channels_last_3d)
        self.relu = nn.ReLU()
        self.group_norm = nn.GroupNorm(groups, out_channels)

    def forward(self, x):
        if x.is_cuda and not x.is_contiguous(memory_format=torch.channels_last_3d):
            if (
                x.dtype == torch.float32
                and x.ndim == 5
                and x.shape[1] == 64
                and x.is_contiguous()
                and 0 < x.shape[0] <= 65535
                and 0 < x[0].numel() < 2**31
            ):
                x = _fused_ops.to_channels_last_64(x)
            else:
                x = x.contiguous(memory_format=torch.channels_last_3d)

        x = self.conv_transpose(x)
        norm = self.group_norm

        if (
            x.is_cuda
            and x.dtype == torch.float32
            and x.shape[1] == 128
            and norm.num_groups == 8
            and 0 < x.shape[0] <= 65535
            and 0 < x[0].numel() < 2**31
            and norm.weight.dtype == torch.float32
            and norm.bias.dtype == torch.float32
        ):
            return _fused_ops.fused_relu_group_norm_128(
                x.contiguous(memory_format=torch.channels_last_3d),
                norm.weight.contiguous(),
                norm.bias.contiguous(),
                norm.eps,
            )

        return F.group_norm(
            F.relu(x), norm.num_groups, norm.weight, norm.bias, norm.eps
        )
```