import torch
import torch.nn as nn

class Model(nn.Module):
    """
    Simple model that performs a matrix multiplication, adds a value, applies Swish, Tanh, GELU, and Hardtanh activation functions.
    """
    def __init__(self, in_features, out_features, add_value_shape):
        super(Model, self).__init__()
        self.matmul = nn.Linear(in_features, out_features)
        self.add_value = nn.Parameter(torch.randn(add_value_shape)) 

    def forward(self, x):
        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

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]

import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

_cpp_src = r"""
#include <torch/extension.h>
torch::Tensor fused_forward(torch::Tensor input, torch::Tensor weight, torch::Tensor bias, torch::Tensor add_value);
"""

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

__global__ void add_act_kernel(float* __restrict__ y,
                               const float* __restrict__ bias,
                               const float* __restrict__ addv,
                               int rows, int cols) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y;
    if (col < cols) {
        int idx = row * cols + col;
        float v = y[idx] + bias[col] + addv[col];
        float s = 1.0f / (1.0f + expf(-v));
        float t = tanhf(v * s);
        float g = 0.5f * t * (1.0f + erff(t * 0.70710678118654752440f));
        y[idx] = g;
    }
}

torch::Tensor fused_forward(torch::Tensor input, torch::Tensor weight, torch::Tensor bias, torch::Tensor add_value) {
    auto y = at::mm(input, weight.t());
    const int rows = (int)y.size(0);
    const int cols = (int)y.size(1);
    const int threads = 256;
    dim3 blocks((cols + threads - 1) / threads, rows);
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    add_act_kernel<<<blocks, threads, 0, stream>>>(y.data_ptr<float>(), bias.data_ptr<float>(), add_value.data_ptr<float>(), rows, cols);
    return y;
}
"""

_ext = load_inline(
    name="fused_linear_acts_ext_v2",
    cpp_sources=_cpp_src,
    cuda_sources=_cuda_src,
    functions=["fused_forward"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3", "--use_fast_math"],
    verbose=False,
)

class ModelNew(nn.Module):
    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))

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