```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
from torch.utils.cpp_extension import load_inline

# Custom CUDA extension:
#  - Uses cuBLAS (at::addmm) for the large GEMM + bias (near-peak FP32 throughput).
#  - Fuses the entire elementwise chain (add_value + Swish + Tanh + GELU + Hardtanh)
#    into a single vectorized (float4) kernel, eliminating 5+ intermediate
#    global-memory round-trips.
cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>

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

__global__ void fused_act_kernel_vec4(const float4* __restrict__ in,
                                      const float4* __restrict__ addv,
                                      float4* __restrict__ out,
                                      const int total4,
                                      const int cols4) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < total4) {
        int c = idx % cols4;
        float4 a = addv[c];
        float4 v = in[idx];
        v.x = fused_act(v.x + a.x);
        v.y = fused_act(v.y + a.y);
        v.z = fused_act(v.z + a.z);
        v.w = fused_act(v.w + a.w);
        out[idx] = v;
    }
}

__global__ void fused_act_kernel(const float* __restrict__ in,
                                 const float* __restrict__ addv,
                                 float* __restrict__ out,
                                 const int total,
                                 const int cols) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < total) {
        int c = idx % cols;
        out[idx] = fused_act(in[idx] + addv[c]);
    }
}

torch::Tensor fused_forward(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 inputs must be CUDA tensors");
    TORCH_CHECK(x.scalar_type() == at::kFloat, "Only float32 is supported");

    // y = x @ weight^T + bias  (cuBLAS, bias broadcast over rows)
    auto y = at::addmm(bias, x, weight.t());

    auto addv = add_value.contiguous();
    const int total = (int)y.numel();
    const int cols = (int)y.size(-1);
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    const int block = 256;

    if ((cols & 3) == 0) {
        int total4 = total >> 2;
        int cols4 = cols >> 2;
        int blocks = (total4 + block - 1) / block;
        fused_act_kernel_vec4<<<blocks, block, 0, stream>>>(
            reinterpret_cast<const float4*>(y.data_ptr<float>()),
            reinterpret_cast<const float4*>(addv.data_ptr<float>()),
            reinterpret_cast<float4*>(y.data_ptr<float>()),
            total4, cols4);
    } else {
        int blocks = (total + block - 1) / block;
        fused_act_kernel<<<blocks, block, 0, stream>>>(
            y.data_ptr<float>(), addv.data_ptr<float>(), y.data_ptr<float>(),
            total, cols);
    }
    return y;
}
"""

cpp_source = "torch::Tensor fused_forward(torch::Tensor x, torch::Tensor weight, torch::Tensor bias, torch::Tensor add_value);"

fused_mod = load_inline(
    name="fused_linear_swish_tanh_gelu_hardtanh",
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=["fused_forward"],
    verbose=False,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model: cuBLAS GEMM + a single fused CUDA kernel for
    (add_value + 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 = fused_mod

    def forward(self, x):
        return self.fused.fused_forward(
            x, self.matmul.weight, self.matmul.bias, self.add_value
        )
```