```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 at::linear)
#   y = hardtanh(gelu(tanh(swish(y + add_value))))   (single fused kernel)
#
# The elementwise chain (add, sigmoid, mul, tanh, gelu, hardtanh) normally
# launches 6 separate kernels with full memory round-trips each. Here they are
# fused into a single vectorized (float4) kernel applied directly as an
# epilogue to the matmul output.
fused_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <cstdint>

__device__ __forceinline__ float fused_activation(float v) {
    // Swish: v * sigmoid(v)
    float sig = 1.0f / (1.0f + expf(-v));
    v = v * sig;
    // Tanh
    v = tanhf(v);
    // GELU (exact erf formulation): 0.5 * v * (1 + erf(v / sqrt(2)))
    v = 0.5f * v * (1.0f + erff(v * 0.70710678118654752440f));
    // Hardtanh clamp to [-1, 1]
    v = fminf(fmaxf(v, -1.0f), 1.0f);
    return v;
}

__global__ void fused_activation_vec4_kernel(float4* __restrict__ data,
                                             const float4* __restrict__ add_value,
                                             int n4,
                                             int width4) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n4) {
        int col = idx - (idx / width4) * width4;
        float4 xv = data[idx];
        float4 av = add_value[col];
        xv.x = fused_activation(xv.x + av.x);
        xv.y = fused_activation(xv.y + av.y);
        xv.z = fused_activation(xv.z + av.z);
        xv.w = fused_activation(xv.w + av.w);
        data[idx] = xv;
    }
}

__global__ void fused_activation_kernel(float* __restrict__ data,
                                        const float* __restrict__ add_value,
                                        int n,
                                        int width) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        int col = idx - (idx / width) * width;
        data[idx] = fused_activation(data[idx] + add_value[col]);
    }
}

torch::Tensor fused_linear_activation_cuda(torch::Tensor x,
                                           torch::Tensor weight,
                                           torch::Tensor bias,
                                           torch::Tensor add_value) {
    TORCH_CHECK(x.is_cuda() && weight.is_cuda() && bias.is_cuda() && add_value.is_cuda(),
                "All tensors must be CUDA tensors");
    TORCH_CHECK(x.scalar_type() == at::kFloat, "Only float32 is supported");

    auto x_c = x.contiguous();
    auto w_c = weight.contiguous();
    auto b_c = bias.contiguous();
    auto add_c = add_value.contiguous();

    // cuBLAS matmul + bias (fused epilogue), output is a fresh contiguous tensor.
    auto y = at::linear(x_c, w_c, b_c);

    const int64_t n64 = y.numel();
    const int64_t width64 = add_c.numel();
    if (n64 == 0) {
        return y;
    }
    const int n = static_cast<int>(n64);
    const int width = static_cast<int>(width64);

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

    float* y_ptr = y.data_ptr<float>();
    const float* add_ptr = add_c.data_ptr<float>();

    const bool aligned =
        ((reinterpret_cast<uintptr_t>(y_ptr) & 0xF) == 0) &&
        ((reinterpret_cast<uintptr_t>(add_ptr) & 0xF) == 0);

    if ((width % 4 == 0) && aligned) {
        const int width4 = width / 4;
        const int n4 = n / 4;
        const int grid = (n4 + block - 1) / block;
        fused_activation_vec4_kernel<<<grid, block, 0, stream>>>(
            reinterpret_cast<float4*>(y_ptr),
            reinterpret_cast<const float4*>(add_ptr),
            n4, width4);
    } else {
        const int grid = (n + block - 1) / block;
        fused_activation_kernel<<<grid, block, 0, stream>>>(
            y_ptr, add_ptr, n, width);
    }
    return y;
}
"""

fused_cpp_source = (
    "torch::Tensor fused_linear_activation_cuda(torch::Tensor x, torch::Tensor weight, "
    "torch::Tensor bias, torch::Tensor add_value);"
)

fused_op = load_inline(
    name="fused_linear_activation",
    cpp_sources=fused_cpp_source,
    cuda_sources=fused_cuda_source,
    functions=["fused_linear_activation_cuda"],
    verbose=False,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model: cuBLAS Linear followed by a single fused CUDA kernel that
    performs add + Swish + Tanh + GELU + Hardtanh in one pass over the data.
    """
    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_op = fused_op

    def forward(self, x):
        if x.is_cuda and x.dtype == torch.float32:
            return self.fused_op.fused_linear_activation_cuda(
                x, self.matmul.weight, self.matmul.bias, self.add_value
            )
        # Fallback to the reference PyTorch implementation.
        x = self.matmul(x)
        x = x + self.add_value
        x = torch.sigmoid(x) * x  # Swish
        x = torch.tanh(x)
        x = torch.nn.functional.gelu(x)  # GELU
        x = torch.nn.functional.hardtanh(x, min_val=-1, max_val=1)  # Hardtanh
        return x
```