```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: (add bias value) + Swish + Tanh + GELU + Hardtanh
# The GEMM itself is left to cuBLAS (torch.nn.Linear), which is already
# near-optimal for these sizes. All pointwise ops are fused into a single
# memory-bound pass (with float4 vectorization) instead of 5+ separate passes.
fused_cuda_source = r"""
#include <torch/extension.h>
#include <cuda_runtime.h>
#include <cstdint>

#define CHECK_CUDA(x) TORCH_CHECK(x.is_cuda(), #x " must be a CUDA tensor")
#define CHECK_CONTIGUOUS(x) TORCH_CHECK(x.is_contiguous(), #x " must be contiguous")
#define CHECK_FLOAT(x) TORCH_CHECK(x.scalar_type() == at::kFloat, #x " must be float32")

__device__ __forceinline__ float fused_op(float x, float add) {
    // add the per-column bias value
    x = x + add;
    // Swish: x * sigmoid(x) == x / (1 + exp(-x))
    float s = x / (1.0f + expf(-x));
    // Tanh
    float t = tanhf(s);
    // GELU (exact erf formulation, matches F.gelu default):
    // 0.5 * t * (1 + erf(t / sqrt(2)))
    float g = 0.5f * t * (1.0f + erff(t * 0.70710678118654752440f));
    // Hardtanh with min=-1, max=1
    return fminf(fmaxf(g, -1.0f), 1.0f);
}

__global__ void fused_activation_kernel_vec4(const float4* __restrict__ x,
                                             const float4* __restrict__ add,
                                             float4* __restrict__ out,
                                             const int64_t n4,
                                             const int64_t add_len4) {
    int64_t idx = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
    int64_t stride = (int64_t)blockDim.x * gridDim.x;
    for (; idx < n4; idx += stride) {
        float4 xv = x[idx];
        float4 av = add[idx % add_len4];
        float4 ov;
        ov.x = fused_op(xv.x, av.x);
        ov.y = fused_op(xv.y, av.y);
        ov.z = fused_op(xv.z, av.z);
        ov.w = fused_op(xv.w, av.w);
        out[idx] = ov;
    }
}

__global__ void fused_activation_kernel_scalar(const float* __restrict__ x,
                                               const float* __restrict__ add,
                                               float* __restrict__ out,
                                               const int64_t n,
                                               const int64_t add_len) {
    int64_t idx = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
    int64_t stride = (int64_t)blockDim.x * gridDim.x;
    for (; idx < n; idx += stride) {
        out[idx] = fused_op(x[idx], add[idx % add_len]);
    }
}

torch::Tensor fused_add_activations_cuda(torch::Tensor x, torch::Tensor add_value) {
    CHECK_CUDA(x);
    CHECK_CUDA(add_value);
    CHECK_CONTIGUOUS(x);
    CHECK_CONTIGUOUS(add_value);
    CHECK_FLOAT(x);
    CHECK_FLOAT(add_value);

    auto out = torch::empty_like(x);
    const int64_t n = x.numel();
    const int64_t add_len = add_value.numel();
    if (n == 0) {
        return out;
    }

    const int block_size = 256;
    const int64_t max_blocks = 65535;

    // Vectorized fast path requires the per-row broadcast length to be a
    // multiple of 4 so each row maps to a whole number of float4 elements.
    if ((add_len % 4 == 0) && (n % 4 == 0) &&
        (reinterpret_cast<uintptr_t>(x.data_ptr<float>()) % 16 == 0) &&
        (reinterpret_cast<uintptr_t>(add_value.data_ptr<float>()) % 16 == 0) &&
        (reinterpret_cast<uintptr_t>(out.data_ptr<float>()) % 16 == 0)) {
        const int64_t n4 = n / 4;
        const int64_t add_len4 = add_len / 4;
        int64_t num_blocks = (n4 + block_size - 1) / block_size;
        if (num_blocks > max_blocks) num_blocks = max_blocks;
        fused_activation_kernel_vec4<<<num_blocks, block_size>>>(
            reinterpret_cast<const float4*>(x.data_ptr<float>()),
            reinterpret_cast<const float4*>(add_value.data_ptr<float>()),
            reinterpret_cast<float4*>(out.data_ptr<float>()),
            n4, add_len4);
    } else {
        int64_t num_blocks = (n + block_size - 1) / block_size;
        if (num_blocks > max_blocks) num_blocks = max_blocks;
        fused_activation_kernel_scalar<<<num_blocks, block_size>>>(
            x.data_ptr<float>(),
            add_value.data_ptr<float>(),
            out.data_ptr<float>(),
            n, add_len);
    }

    return out;
}
"""

fused_cpp_source = (
    "torch::Tensor fused_add_activations_cuda(torch::Tensor x, torch::Tensor add_value);"
)

# Compile the inline CUDA extension
fused_activation_ops = load_inline(
    name="fused_add_activations_ext",
    cpp_sources=fused_cpp_source,
    cuda_sources=fused_cuda_source,
    functions=["fused_add_activations_cuda"],
    verbose=False,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model: keeps the GEMM on cuBLAS (nn.Linear) and replaces the
    add + Swish + Tanh + GELU + Hardtanh chain with a single fused CUDA kernel.
    """

    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_activation_ops

    def forward(self, x):
        x = self.matmul(x)
        x = self.fused_ops.fused_add_activations_cuda(x, self.add_value)
        return x
```