```python
# Copyright (c) Meta Platforms, Inc. and affiliates.
# All rights reserved.
#
# This source code is licensed under the license found in the
# LICENSE file in the root directory of this source tree.

import torch
import torch.nn as nn
import torch.nn.functional as F
from torch.utils.cpp_extension import load_inline

# Fused CUDA kernel:
#   y = Linear(x)                       (cuBLAS via F.linear - already near peak)
#   y = hardtanh(gelu(tanh(swish(y + add_value))), -1, 1)
# All pointwise ops (add + swish + tanh + gelu + hardtanh) are fused into a
# single memory-bound kernel using float4 vectorized loads/stores.
fused_ops_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <cstdint>

__device__ __forceinline__ float act_chain(float v) {
    // Swish: v * sigmoid(v)
    float sig = 1.0f / (1.0f + expf(-v));
    v = v * sig;
    // Tanh
    v = tanhf(v);
    // GELU (exact, erf-based, matches F.gelu default)
    v = 0.5f * v * (1.0f + erff(v * 0.70710678118654752440f));
    // Hardtanh [-1, 1]
    v = fminf(fmaxf(v, -1.0f), 1.0f);
    return v;
}

__global__ void fused_ops_vec4_kernel(const float4* __restrict__ x,
                                      const float4* __restrict__ addv,
                                      float4* __restrict__ out,
                                      int n4, int of4) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n4) {
        float4 a = addv[i % of4];
        float4 v = x[i];
        v.x = act_chain(v.x + a.x);
        v.y = act_chain(v.y + a.y);
        v.z = act_chain(v.z + a.z);
        v.w = act_chain(v.w + a.w);
        out[i] = v;
    }
}

__global__ void fused_ops_scalar_kernel(const float* __restrict__ x,
                                        const float* __restrict__ addv,
                                        float* __restrict__ out,
                                        long n, int of) {
    long i = (long)blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        out[i] = act_chain(x[i] + addv[i % of]);
    }
}

torch::Tensor fused_ops_cuda(torch::Tensor x, torch::Tensor add_value) {
    TORCH_CHECK(x.is_cuda(), "x must be a CUDA tensor");
    TORCH_CHECK(add_value.is_cuda(), "add_value must be a CUDA tensor");
    TORCH_CHECK(x.scalar_type() == torch::kFloat32, "x must be float32");
    TORCH_CHECK(add_value.scalar_type() == torch::kFloat32, "add_value must be float32");

    auto x_c = x.contiguous();
    auto add_c = add_value.contiguous();
    auto out = torch::empty_like(x_c);

    long n = x_c.numel();
    int of = (int)add_c.numel();
    if (n == 0) {
        return out;
    }

    const int block = 256;
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    bool vec_ok = (of % 4 == 0) && (n % 4 == 0) &&
                  (reinterpret_cast<uintptr_t>(x_c.data_ptr<float>()) % 16 == 0) &&
                  (reinterpret_cast<uintptr_t>(add_c.data_ptr<float>()) % 16 == 0) &&
                  (reinterpret_cast<uintptr_t>(out.data_ptr<float>()) % 16 == 0);

    if (vec_ok) {
        long n4 = n / 4;
        int of4 = of / 4;
        long grid = (n4 + block - 1) / block;
        fused_ops_vec4_kernel<<<(unsigned int)grid, block, 0, stream>>>(
            reinterpret_cast<const float4*>(x_c.data_ptr<float>()),
            reinterpret_cast<const float4*>(add_c.data_ptr<float>()),
            reinterpret_cast<float4*>(out.data_ptr<float>()),
            (int)n4, of4);
    } else {
        long grid = (n + block - 1) / block;
        fused_ops_scalar_kernel<<<(unsigned int)grid, block, 0, stream>>>(
            x_c.data_ptr<float>(),
            add_c.data_ptr<float>(),
            out.data_ptr<float>(),
            n, of);
    }

    return out;
}
"""

fused_ops_cpp_source = "torch::Tensor fused_ops_cuda(torch::Tensor x, torch::Tensor add_value);"

fused_ops = load_inline(
    name="fused_add_swish_tanh_gelu_hardtanh",
    cpp_sources=fused_ops_cpp_source,
    cuda_sources=fused_ops_cuda_source,
    functions=["fused_ops_cuda"],
    verbose=False,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3", "--use_fast_math=false"],
)


class ModelNew(nn.Module):
    """
    Optimized model:
      - GEMM + bias kept on cuBLAS (F.linear), which is near-peak for this size.
      - add_value + Swish + Tanh + GELU + Hardtanh fused into one custom
        vectorized CUDA kernel, eliminating 4 intermediate global-memory
        round-trips and multiple kernel launches.
    """
    def __init__(self, in_features, out_features, add_value_shape):
        super(ModelNew, self).__init__()
        self.matmul = nn.Linear(in_features, out_features)
        self.add_value = nn.Parameter(torch.randn(add_value_shape))
        self.fused_ops = fused_ops

    def forward(self, x):
        if x.is_cuda and x.dtype == torch.float32:
            y = F.linear(x, self.matmul.weight, self.matmul.bias)
            return self.fused_ops.fused_ops_cuda(y, self.add_value)
        # Fallback reference path (CPU / non-float32)
        x = self.matmul(x)
        x = x + self.add_value
        x = torch.sigmoid(x) * x  # Swish
        x = torch.tanh(x)
        x = F.gelu(x)  # GELU
        x = F.hardtanh(x, min_val=-1, max_val=1)  # Hardtanh
        return x
```