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

# ---------------------------------------------------------------------------
# Fused epilogue kernel:
#   y = y + add_value
#   y = y * sigmoid(y)        (Swish)
#   y = tanh(y)               (Tanh)
#   y = gelu(y)               (exact erf GELU, PyTorch default)
#   y = clamp(y, -1, 1)       (Hardtanh)
#
# The GEMM itself is executed with cuBLAS (torch.addmm), which is near-peak
# for this shape; all pointwise ops are fused into a single memory-bound
# kernel, eliminating 5 extra kernel launches and ~5 round-trips to HBM.
# ---------------------------------------------------------------------------

fused_epilogue_cpp_source = (
    "torch::Tensor fused_epilogue_cuda(torch::Tensor y, torch::Tensor add_value);"
)

fused_epilogue_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>

__device__ __forceinline__ float full_epilogue(float v) {
    // Swish: v * sigmoid(v) = v / (1 + exp(-v))
    float sv = v / (1.0f + expf(-v));
    // Tanh
    float t = tanhf(sv);
    // Exact (erf) GELU: 0.5 * t * (1 + erf(t / sqrt(2)))
    float g = 0.5f * t * (1.0f + erff(t * 0.70710678118654752440f));
    // Hardtanh on [-1, 1]
    return fminf(fmaxf(g, -1.0f), 1.0f);
}

__global__ void fused_epilogue_vec4_kernel(const float4* __restrict__ y,
                                           const float4* __restrict__ add_value,
                                           float4* __restrict__ out,
                                           long nvec,
                                           int row_vec) {
    long i = (long)blockIdx.x * blockDim.x + threadIdx.x;
    if (i < nvec) {
        int col = (int)(i % row_vec);
        float4 a = add_value[col];
        float4 v = y[i];
        float4 o;
        o.x = full_epilogue(v.x + a.x);
        o.y = full_epilogue(v.y + a.y);
        o.z = full_epilogue(v.z + a.z);
        o.w = full_epilogue(v.w + a.w);
        out[i] = o;
    }
}

__global__ void fused_epilogue_scalar_kernel(const float* __restrict__ y,
                                             const float* __restrict__ add_value,
                                             float* __restrict__ out,
                                             long total,
                                             int row_size) {
    long i = (long)blockIdx.x * blockDim.x + threadIdx.x;
    if (i < total) {
        int col = (int)(i % row_size);
        out[i] = full_epilogue(y[i] + add_value[col]);
    }
}

torch::Tensor fused_epilogue_cuda(torch::Tensor y, torch::Tensor add_value) {
    TORCH_CHECK(y.is_cuda(), "y must be a CUDA tensor");
    TORCH_CHECK(add_value.is_cuda(), "add_value must be a CUDA tensor");
    TORCH_CHECK(y.scalar_type() == at::kFloat, "y must be float32");
    TORCH_CHECK(add_value.scalar_type() == at::kFloat, "add_value must be float32");

    auto y_c = y.contiguous();
    auto add_c = add_value.contiguous();
    auto out = torch::empty_like(y_c);

    long total = y_c.numel();
    long row_size = add_c.numel();
    TORCH_CHECK(row_size > 0, "add_value must be non-empty");
    TORCH_CHECK(total % row_size == 0, "row size mismatch");

    if (total == 0) {
        return out;
    }

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

    if ((total % 4 == 0) && (row_size % 4 == 0)) {
        long nvec = total / 4;
        int row_vec = (int)(row_size / 4);
        unsigned int num_blocks = (unsigned int)((nvec + block_size - 1) / block_size);
        fused_epilogue_vec4_kernel<<<num_blocks, block_size, 0, stream>>>(
            reinterpret_cast<const float4*>(y_c.data_ptr<float>()),
            reinterpret_cast<const float4*>(add_c.data_ptr<float>()),
            reinterpret_cast<float4*>(out.data_ptr<float>()),
            nvec, row_vec);
    } else {
        unsigned int num_blocks = (unsigned int)((total + block_size - 1) / block_size);
        fused_epilogue_scalar_kernel<<<num_blocks, block_size, 0, stream>>>(
            y_c.data_ptr<float>(), add_c.data_ptr<float>(), out.data_ptr<float>(),
            total, (int)row_size);
    }

    return out;
}
"""

fused_epilogue = load_inline(
    name="fused_swish_tanh_gelu_hardtanh",
    cpp_sources=fused_epilogue_cpp_source,
    cuda_sources=fused_epilogue_cuda_source,
    functions=["fused_epilogue_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model: cuBLAS GEMM (addmm) + a single fused CUDA kernel that
    applies the bias-add, Swish, Tanh, GELU and Hardtanh in one pass.
    """
    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_epilogue

    def forward(self, x):
        lead_shape = x.shape[:-1]
        x2 = x.reshape(-1, self.matmul.in_features)
        if not x2.is_contiguous():
            x2 = x2.contiguous()
        # GEMM + linear bias via cuBLAS.
        y = torch.addmm(self.matmul.bias, x2, self.matmul.weight.t())
        # Fused: + add_value, Swish, Tanh, GELU, Hardtanh.
        out = self._fused.fused_epilogue_cuda(y, self.add_value)
        return out.view(lead_shape + (self.matmul.out_features,))


batch_size = 1024
in_features = 8192
out_features = 8192
add_value_shape = (out_features,)


def get_inputs():
    return [torch.rand(batch_size, in_features)]


def get_init_inputs():
    return [in_features, out_features, add_value_shape]
```