```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 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 GROUP_VECTORS = 4;
constexpr int SPATIAL_TILE = 128;
constexpr int OUTPUT_VECTORS = 2048;

__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;
}

__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;
}

// Each warp reads all channels of a spatial position, so all input
// transactions are contiguous even when reducing separate groups.
__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 batch_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[batch_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,
    float2* __restrict__ stats,
    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);

        if (lane == 0) {
            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);
            stats[ng] = make_float2(mean, rsqrtf(variance + eps));
        }
    }
}

__global__ __launch_bounds__(THREADS)
void normalize_kernel(
    const float4* __restrict__ x,
    float4* __restrict__ output,
    const float2* __restrict__ stats,
    const float4* __restrict__ weight,
    const float4* __restrict__ bias,
    int vectors_per_batch)
{
    const int tid = threadIdx.x;
    const int n = blockIdx.y;
    const int channel_vector = tid & 31;
    const int group = channel_vector >> 2;

    const float2 st = stats[n * GROUPS + group];
    const float4 w = weight[channel_vector];
    const float4 b = bias[channel_vector];
    const float4 scale = make_float4(
        st.y * w.x, st.y * w.y, st.y * w.z, st.y * w.w);

    const int64_t batch_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[batch_base + i]);
            v.x = (v.x - st.x) * scale.x + b.x;
            v.y = (v.y - st.x) * scale.y + b.y;
            v.z = (v.z - st.x) * scale.z + b.z;
            v.w = (v.w - st.x) * scale.w + b.w;
            output[batch_base + i] = v;
        }
    }
}

} // namespace

torch::Tensor fused_relu_group_norm_128(
    torch::Tensor x,
    torch::Tensor weight,
    torch::Tensor bias,
    double eps)
{
    const c10::cuda::CUDAGuard guard(x.device());

    TORCH_CHECK(x.is_cuda() && x.scalar_type() == at::kFloat,
                "Expected a CUDA float32 input");
    TORCH_CHECK(x.dim() == 5 && x.size(1) == CHANNELS &&
                x.is_contiguous(at::MemoryFormat::ChannelsLast3d),
                "Expected channels-last 3D input with 128 channels");

    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 stats = torch::empty({batch * GROUPS, 2}, x.options());

    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>()),
        reinterpret_cast<float2*>(stats.data_ptr<float>()),
        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 float2*>(stats.data_ptr<float>()),
        reinterpret_cast<const float4*>(weight.data_ptr<float>()),
        reinterpret_cast<const float4*>(bias.data_ptr<float>()),
        vectors_per_batch);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    return output;
}
"""

_fused_ops = load_inline(
    name="relu_groupnorm_ndhwc_coalesced_v3",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["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:
            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 x.numel() > 0
            and 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
        )
```