```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 linear_norm_swish_cuda(
    torch::Tensor input,
    torch::Tensor weight,
    torch::Tensor bias,
    torch::Tensor gamma,
    torch::Tensor beta,
    torch::Tensor multiply_weight,
    int64_t groups,
    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>
#include <cublas_v2.h>

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

__device__ __forceinline__ float finish_value(
    float value, float mean, float inv_std,
    float gamma, float beta, float multiplier
) {
    float normalized = ((value - mean) * inv_std) * gamma + beta;
    float activated = normalized / (1.0f + __expf(-normalized));
    float weighted = activated * multiplier;
    return weighted / (1.0f + __expf(-weighted));
}

template<bool GROUP_SIZE_32>
__global__ void fused_norm_kernel(
    const float* __restrict__ input,
    const float* __restrict__ bias,
    const float* __restrict__ gamma,
    const float* __restrict__ beta,
    const float* __restrict__ multiplier,
    float* __restrict__ output,
    int channels,
    int groups,
    int group_size,
    int total_groups,
    float eps
) {
    int lane = threadIdx.x & 31;
    int group = blockIdx.x * (blockDim.x / 32) + threadIdx.x / 32;
    if (group >= total_groups) return;

    int channel_start = (group % groups) * group_size;
    int base = (group / groups) * channels + channel_start;

    if (GROUP_SIZE_32) {
        int channel = channel_start + lane;
        float value = input[base + lane] + bias[channel];
        float mean = warp_sum(value) * (1.0f / 32.0f);
        float centered = value - mean;
        float variance = warp_sum(centered * centered) * (1.0f / 32.0f);
        float inv_std = rsqrtf(variance + eps);

        output[base + lane] = finish_value(
            value, mean, inv_std,
            gamma[channel], beta[channel], multiplier[channel]);
    } else {
        float sum = 0.0f;
        for (int i = lane; i < group_size; i += 32) {
            sum += input[base + i] + bias[channel_start + i];
        }
        float mean = warp_sum(sum) / float(group_size);

        float squared_sum = 0.0f;
        for (int i = lane; i < group_size; i += 32) {
            float value = input[base + i] + bias[channel_start + i];
            float centered = value - mean;
            squared_sum += centered * centered;
        }
        float inv_std = rsqrtf(
            warp_sum(squared_sum) / float(group_size) + eps);

        for (int i = lane; i < group_size; i += 32) {
            int channel = channel_start + i;
            float value = input[base + i] + bias[channel];
            output[base + i] = finish_value(
                value, mean, inv_std,
                gamma[channel], beta[channel], multiplier[channel]);
        }
    }
}

torch::Tensor linear_norm_swish_cuda(
    torch::Tensor input,
    torch::Tensor weight,
    torch::Tensor bias,
    torch::Tensor gamma,
    torch::Tensor beta,
    torch::Tensor multiply_weight,
    int64_t groups,
    double eps
) {
    TORCH_CHECK(input.is_cuda(), "Expected CUDA input");
    const c10::cuda::CUDAGuard guard(input.device());

    TORCH_CHECK(input.dim() == 2 && weight.dim() == 2);
    TORCH_CHECK(input.scalar_type() == at::kFloat);
    TORCH_CHECK(input.is_contiguous());

    const int batch = static_cast<int>(input.size(0));
    const int channels = static_cast<int>(weight.size(0));
    const int features = static_cast<int>(input.size(1));

    TORCH_CHECK(weight.size(1) == features);
    TORCH_CHECK(groups > 0 && channels % groups == 0);

    for (const auto& tensor : {weight, bias, gamma, beta, multiply_weight}) {
        TORCH_CHECK(tensor.device() == input.device());
        TORCH_CHECK(tensor.scalar_type() == at::kFloat);
        TORCH_CHECK(tensor.is_contiguous());
    }
    TORCH_CHECK(bias.numel() == channels);
    TORCH_CHECK(gamma.numel() == channels);
    TORCH_CHECK(beta.numel() == channels);
    TORCH_CHECK(multiply_weight.numel() == channels);

    auto intermediate = torch::empty({batch, channels}, input.options());
    auto output = torch::empty_like(intermediate);
    if (batch == 0) return output;

    auto stream = at::cuda::getCurrentCUDAStream();
    auto handle = at::cuda::getCurrentCUDABlasHandle();

    const float alpha = 1.0f;
    const float gemm_beta = 0.0f;

    // Row-major X @ W.T is column-major W @ X.T.
    // Explicit TF32 compute avoids falling back to the FP32 SIMT GEMM.
    cublasStatus_t status = cublasGemmEx(
        handle,
        CUBLAS_OP_T,
        CUBLAS_OP_N,
        channels,
        batch,
        features,
        &alpha,
        weight.data_ptr<float>(),
        CUDA_R_32F,
        features,
        input.data_ptr<float>(),
        CUDA_R_32F,
        features,
        &gemm_beta,
        intermediate.data_ptr<float>(),
        CUDA_R_32F,
        channels,
        CUBLAS_COMPUTE_32F_FAST_TF32,
        CUBLAS_GEMM_DEFAULT_TENSOR_OP
    );
    TORCH_CHECK(
        status == CUBLAS_STATUS_SUCCESS,
        "cublasGemmEx failed with status ", static_cast<int>(status));

    const int group_size = channels / static_cast<int>(groups);
    const int total_groups = batch * static_cast<int>(groups);
    constexpr int threads = 256;
    const int blocks = (total_groups + 7) / 8;

    if (group_size == 32) {
        fused_norm_kernel<true><<<blocks, threads, 0, stream>>>(
            intermediate.data_ptr<float>(),
            bias.data_ptr<float>(),
            gamma.data_ptr<float>(),
            beta.data_ptr<float>(),
            multiply_weight.data_ptr<float>(),
            output.data_ptr<float>(),
            channels, static_cast<int>(groups), group_size,
            total_groups, static_cast<float>(eps));
    } else {
        fused_norm_kernel<false><<<blocks, threads, 0, stream>>>(
            intermediate.data_ptr<float>(),
            bias.data_ptr<float>(),
            gamma.data_ptr<float>(),
            beta.data_ptr<float>(),
            multiply_weight.data_ptr<float>(),
            output.data_ptr<float>(),
            channels, static_cast<int>(groups), group_size,
            total_groups, static_cast<float>(eps));
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_fused_ops = load_inline(
    name="tf32_linear_fused_groupnorm_swish_v2",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["linear_norm_swish_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3", "--fmad=false"],
    extra_ldflags=["-lcublas"],
    verbose=False,
)


class ModelNew(nn.Module):
    def __init__(
        self, in_features, out_features, num_groups, multiply_weight_shape
    ):
        super().__init__()
        self.gemm = nn.Linear(in_features, out_features)
        self.group_norm = nn.GroupNorm(num_groups, out_features)
        self.multiply_weight = nn.Parameter(torch.randn(multiply_weight_shape))

    def forward(self, x):
        if (
            x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 2
            and x.shape[1] > 0
            and self.gemm.weight.dtype == torch.float32
            and self.gemm.bias.dtype == torch.float32
            and self.group_norm.weight.dtype == torch.float32
            and self.group_norm.bias.dtype == torch.float32
            and self.multiply_weight.dtype == torch.float32
            and self.multiply_weight.shape == (self.gemm.out_features,)
        ):
            return _fused_ops.linear_norm_swish_cuda(
                x.contiguous(),
                self.gemm.weight.contiguous(),
                self.gemm.bias.contiguous(),
                self.group_norm.weight.contiguous(),
                self.group_norm.bias.contiguous(),
                self.multiply_weight.contiguous(),
                self.group_norm.num_groups,
                self.group_norm.eps,
            )

        x = self.gemm(x)
        x = self.group_norm(x)
        x = x * torch.sigmoid(x)
        x = x * self.multiply_weight
        return x * torch.sigmoid(x)
```