```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: (linear bias add) + (add_value add) + Swish + Tanh + GELU + Hardtanh
# in a single memory pass. The GEMM itself is left to cuBLAS (FP32 / TF32).
fused_source = r"""
#include <torch/extension.h>
#include <cuda_runtime.h>

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

__global__ void fused_bias_act_kernel_vec4(
    const float4* __restrict__ x,
    const float4* __restrict__ b1,
    const float4* __restrict__ b2,
    float4* __restrict__ out,
    int n4, int row4)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n4) {
        int j = i % row4;
        float4 v = x[i];
        float4 p = b1[j];
        float4 q = b2[j];
        v.x = act_chain(v.x + p.x + q.x);
        v.y = act_chain(v.y + p.y + q.y);
        v.z = act_chain(v.z + p.z + q.z);
        v.w = act_chain(v.w + p.w + q.w);
        out[i] = v;
    }
}

__global__ void fused_bias_act_kernel(
    const float* __restrict__ x,
    const float* __restrict__ b1,
    const float* __restrict__ b2,
    float* __restrict__ out,
    int n, int row)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        int j = i % row;
        out[i] = act_chain(x[i] + b1[j] + b2[j]);
    }
}

torch::Tensor fused_bias_act(torch::Tensor x, torch::Tensor b1, torch::Tensor b2) {
    TORCH_CHECK(x.is_cuda() && b1.is_cuda() && b2.is_cuda(), "Inputs must be CUDA tensors");
    TORCH_CHECK(x.scalar_type() == at::kFloat, "Input must be float32");
    auto xc = x.contiguous();
    auto b1c = b1.contiguous();
    auto b2c = b2.contiguous();
    auto out = torch::empty_like(xc);

    int64_t n = xc.numel();
    int64_t row = b1c.numel();
    if (n == 0) return out;

    const int threads = 256;
    if (row % 4 == 0 && (n % 4 == 0)) {
        int64_t n4 = n / 4;
        int64_t row4 = row / 4;
        int blocks = (int)((n4 + threads - 1) / threads);
        fused_bias_act_kernel_vec4<<<blocks, threads>>>(
            reinterpret_cast<const float4*>(xc.data_ptr<float>()),
            reinterpret_cast<const float4*>(b1c.data_ptr<float>()),
            reinterpret_cast<const float4*>(b2c.data_ptr<float>()),
            reinterpret_cast<float4*>(out.data_ptr<float>()),
            (int)n4, (int)row4);
    } else {
        int blocks = (int)((n + threads - 1) / threads);
        fused_bias_act_kernel<<<blocks, threads>>>(
            xc.data_ptr<float>(),
            b1c.data_ptr<float>(),
            b2c.data_ptr<float>(),
            out.data_ptr<float>(),
            (int)n, (int)row);
    }
    return out;
}
"""

fused_cpp_source = (
    "torch::Tensor fused_bias_act(torch::Tensor x, torch::Tensor b1, torch::Tensor b2);"
)

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


class ModelNew(nn.Module):
    """
    Optimized model: cuBLAS GEMM (FP32 with TF32 tensor cores allowed) followed by a
    single fused CUDA kernel performing bias-add + value-add + Swish + Tanh + GELU + Hardtanh.
    """
    def __init__(self, in_features, out_features, add_value_shape):
        super(ModelNew, self).__init__()
        # TF32 tensor-core math is explicitly permitted for FP32 workloads.
        torch.backends.cuda.matmul.allow_tf32 = True
        torch.backends.cudnn.allow_tf32 = True

        self.matmul = nn.Linear(in_features, out_features)
        self.add_value = nn.Parameter(torch.randn(add_value_shape))
        self.fused = fused_op

    def forward(self, x):
        # GEMM in FP32/TF32 via cuBLAS (bias handled by the fused epilogue kernel).
        x = torch.matmul(x, self.matmul.weight.t())
        # Single fused kernel: + linear bias, + add_value, Swish, Tanh, GELU, Hardtanh.
        return self.fused.fused_bias_act(x, self.matmul.bias, self.add_value)
```