```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 fused_groupnorm_swish_cuda(
    torch::Tensor x,
    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>

__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 swish(float value) {
    return value / (1.0f + __expf(-value));
}

template<bool Group32>
__global__ void fused_groupnorm_swish_kernel(
    const float* __restrict__ x,
    const float* __restrict__ gamma,
    const float* __restrict__ beta,
    const float* __restrict__ multiply_weight,
    float* __restrict__ output,
    int64_t total_groups,
    int groups_per_row,
    int group_size,
    float eps)
{
    const int lane = threadIdx.x & 31;
    const int64_t group =
        static_cast<int64_t>(blockIdx.x) * (blockDim.x / 32)
        + threadIdx.x / 32;

    if (group >= total_groups)
        return;

    const int64_t input_offset = group * group_size;
    const int channel_offset = (group % groups_per_row) * group_size;

    if (Group32) {
        const float value = x[input_offset + lane];
        const float mean = warp_sum(value) * (1.0f / 32.0f);
        const float centered = value - mean;
        const float variance =
            warp_sum(centered * centered) * (1.0f / 32.0f);
        const float inverse_std = rsqrtf(variance + eps);

        const int channel = channel_offset + lane;
        const float normalized =
            (centered * inverse_std) * gamma[channel] + beta[channel];
        const float multiplied =
            swish(normalized) * multiply_weight[channel];

        output[input_offset + lane] = swish(multiplied);
    } else {
        float sum = 0.0f;
        for (int i = lane; i < group_size; i += 32)
            sum += x[input_offset + i];

        const float mean = warp_sum(sum) / group_size;

        float squared_sum = 0.0f;
        for (int i = lane; i < group_size; i += 32) {
            const float centered = x[input_offset + i] - mean;
            squared_sum += centered * centered;
        }

        const float variance = warp_sum(squared_sum) / group_size;
        const float inverse_std = rsqrtf(variance + eps);

        for (int i = lane; i < group_size; i += 32) {
            const int channel = channel_offset + i;
            const float normalized =
                ((x[input_offset + i] - mean) * inverse_std)
                * gamma[channel] + beta[channel];
            const float multiplied =
                swish(normalized) * multiply_weight[channel];

            output[input_offset + i] = swish(multiplied);
        }
    }
}

torch::Tensor fused_groupnorm_swish_cuda(
    torch::Tensor x,
    torch::Tensor gamma,
    torch::Tensor beta,
    torch::Tensor multiply_weight,
    int64_t groups,
    double eps)
{
    TORCH_CHECK(x.is_cuda() && x.dim() == 2 && x.is_contiguous(),
                "x must be a contiguous CUDA matrix");
    TORCH_CHECK(x.scalar_type() == at::kFloat,
                "x must have float32 dtype");
    TORCH_CHECK(groups > 0 && x.size(1) % groups == 0,
                "Invalid group count");

    const int64_t channels = x.size(1);
    for (const auto& parameter : {gamma, beta, multiply_weight}) {
        TORCH_CHECK(parameter.device() == x.device() &&
                    parameter.scalar_type() == at::kFloat &&
                    parameter.is_contiguous() &&
                    parameter.numel() == channels,
                    "Parameters must be contiguous float32 tensors "
                    "on the input device with one element per channel");
    }

    const c10::cuda::CUDAGuard device_guard(x.device());
    auto output = torch::empty_like(x);
    const int64_t total_groups = x.size(0) * groups;

    if (total_groups == 0)
        return output;

    constexpr int threads = 256;
    constexpr int warps_per_block = threads / 32;
    const int blocks =
        (total_groups + warps_per_block - 1) / warps_per_block;
    const int group_size = channels / groups;
    const auto stream = at::cuda::getCurrentCUDAStream(x.get_device());

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

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_fused_ops = load_inline(
    name="fused_groupnorm_double_swish_v1",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["fused_groupnorm_swish_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    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):
        x = self.gemm(x)
        parameters = (
            self.group_norm.weight,
            self.group_norm.bias,
            self.multiply_weight,
        )
        vector_weight = (
            self.multiply_weight.shape == (x.shape[-1],)
            or self.multiply_weight.shape == (1, x.shape[-1])
        )
        needs_autograd = torch.is_grad_enabled() and (
            x.requires_grad or any(p.requires_grad for p in parameters)
        )

        if (
            x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 2
            and x.is_contiguous()
            and vector_weight
            and not needs_autograd
            and all(
                p.dtype == torch.float32
                and p.device == x.device
                and p.is_contiguous()
                for p in parameters
            )
        ):
            return _fused_ops.fused_groupnorm_swish_cuda(
                x,
                self.group_norm.weight,
                self.group_norm.bias,
                self.multiply_weight,
                self.group_norm.num_groups,
                self.group_norm.eps,
            )

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