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

# Custom CUDA kernel that fuses the entire pointwise chain:
#   add_value -> Swish (x * sigmoid(x)) -> Tanh -> exact GELU -> Hardtanh[-1, 1]
# The matmul (linear) is left to cuBLAS, which is already near-optimal for
# these sizes; everything after it is fused into a single memory-bound pass.
fused_ops_cuda_source = r"""
#include <torch/extension.h>
#include <cuda_runtime.h>

__device__ __forceinline__ float fused_op(float v) {
    // Swish: v * sigmoid(v)
    v = v / (1.0f + expf(-v));
    // Tanh
    v = tanhf(v);
    // GELU (exact, erf-based): 0.5 * v * (1 + erf(v / sqrt(2)))
    v = 0.5f * v * (1.0f + erff(v * 0.70710678118654752440f));
    // Hardtanh with min_val=-1, max_val=1
    v = fminf(fmaxf(v, -1.0f), 1.0f);
    return v;
}

__global__ void fused_ops_kernel_vec4(const float4* __restrict__ x,
                                      const float4* __restrict__ add_val,
                                      float4* __restrict__ out,
                                      int n4, int cols4) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n4) {
        float4 v = x[idx];
        float4 a = add_val[idx % cols4];
        v.x = fused_op(v.x + a.x);
        v.y = fused_op(v.y + a.y);
        v.z = fused_op(v.z + a.z);
        v.w = fused_op(v.w + a.w);
        out[idx] = v;
    }
}

__global__ void fused_ops_kernel_scalar(const float* __restrict__ x,
                                        const float* __restrict__ add_val,
                                        float* __restrict__ out,
                                        long size, int cols) {
    long idx = (long)blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < size) {
        float v = x[idx] + add_val[idx % cols];
        out[idx] = fused_op(v);
    }
}

torch::Tensor fused_ops_forward(torch::Tensor x, torch::Tensor add_val) {
    auto out = torch::empty_like(x);
    long size = x.numel();
    int cols = x.size(-1);
    const int block_size = 256;

    if ((size % 4) == 0 && (cols % 4) == 0) {
        int n4 = (int)(size / 4);
        int cols4 = cols / 4;
        int num_blocks = (n4 + block_size - 1) / block_size;
        fused_ops_kernel_vec4<<<num_blocks, block_size>>>(
            reinterpret_cast<const float4*>(x.data_ptr<float>()),
            reinterpret_cast<const float4*>(add_val.data_ptr<float>()),
            reinterpret_cast<float4*>(out.data_ptr<float>()),
            n4, cols4);
    } else {
        long num_blocks = (size + block_size - 1) / block_size;
        fused_ops_kernel_scalar<<<num_blocks, block_size>>>(
            x.data_ptr<float>(), add_val.data_ptr<float>(),
            out.data_ptr<float>(), size, cols);
    }
    return out;
}
"""

fused_ops_cpp_source = (
    "torch::Tensor fused_ops_forward(torch::Tensor x, torch::Tensor add_val);"
)

# Compile the inline CUDA extension
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_forward"],
    verbose=False,
    extra_cflags=[""],
    extra_ldflags=[""],
)


class ModelNew(nn.Module):
    """
    Optimized model: cuBLAS matmul + a single fused CUDA kernel for
    bias-add/Swish/Tanh/GELU/Hardtanh.
    """
    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):
        x = F.linear(x, self.matmul.weight, self.matmul.bias)
        x = self.fused_ops.fused_ops_forward(
            x.contiguous(), self.add_value.contiguous()
        )
        return x
```