import torch
import torch.nn as nn
import torch.nn.functional as F

class Model(nn.Module):
    """
    Full Hyena operator (order 2): projects input into value (v), query-like (q), and key-like (k)
    branches, applies short depthwise convolutions to q and k, generates a long convolution filter
    via a small MLP, convolves v with the filter via FFT, then applies two rounds of elementwise
    gating. This is the core sequence mixing operator in StripedHyena, HyenaDNA, and Evo
    for subquadratic long-range dependency modeling.

    Args:
        d_model (int): Model dimension.
        seq_len (int): Maximum sequence length.
        order (int): Hyena recurrence order (number of gating rounds, typically 2).
        short_filter_size (int): Kernel size for short depthwise convolutions.
    """
    def __init__(self, d_model: int, seq_len: int, order: int, short_filter_size: int):
        super(Model, self).__init__()
        self.d_model = d_model
        self.seq_len = seq_len
        self.order = order
        self.fft_size = 2 * seq_len

        self.in_proj = nn.Linear(d_model, d_model * (order + 1), bias=True)

        self.short_convs = nn.ModuleList([
            nn.Conv1d(d_model, d_model, short_filter_size,
                      padding=short_filter_size // 2, groups=d_model)
            for _ in range(order)
        ])

        self.filter_mlp = nn.Sequential(
            nn.Linear(1, 64),
            nn.SiLU(),
            nn.Linear(64, d_model),
        )
        t = torch.linspace(0, 1, seq_len).unsqueeze(-1)
        self.register_buffer('filter_positions', t)

        self.out_proj = nn.Linear(d_model, d_model, bias=True)

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        B, L, D = x.shape

        projected = self.in_proj(x)
        chunks = projected.chunk(self.order + 1, dim=-1)
        v = chunks[0].transpose(1, 2)
        gates = [c.transpose(1, 2) for c in chunks[1:]]

        for i in range(self.order):
            gates[i] = self.short_convs[i](gates[i])

        h = self.filter_mlp(self.filter_positions).transpose(0, 1)

        v_padded = F.pad(v, (0, self.seq_len))
        h_padded = F.pad(h, (0, self.seq_len))
        v_f = torch.fft.rfft(v_padded, dim=-1)
        h_f = torch.fft.rfft(h_padded, dim=-1)
        y_f = v_f * h_f.unsqueeze(0)
        y = torch.fft.irfft(y_f, n=self.fft_size, dim=-1)[..., :L]

        for gate in gates:
            y = y * gate

        y = y.transpose(1, 2)
        y = self.out_proj(y)
        return y

batch_size = 16
d_model = 256
seq_len = 8192
order = 2
short_filter_size = 3

def get_inputs():
    return [torch.randn(popcorn_pri.jitter_int(batch_size), seq_len, d_model)]

def get_init_inputs():
    return [d_model, seq_len, order, short_filter_size]

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

_cpp_src = r"""
#include <torch/extension.h>
torch::Tensor pack_v(torch::Tensor projected);
torch::Tensor fuse_gate_transpose(torch::Tensor y, torch::Tensor projected,
                                  torch::Tensor w0, torch::Tensor b0,
                                  torch::Tensor w1, torch::Tensor b1);
"""

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

#define TILE 32

__global__ void pack_v_tiled_kernel(const float* __restrict__ proj,
                                    float* __restrict__ out,
                                    int B, int L, int D, int N) {
    __shared__ float tile[TILE][TILE + 1];
    int d = blockIdx.x * TILE + threadIdx.x;
    int l = blockIdx.y * TILE + threadIdx.y;
    int b = blockIdx.z;
    if (d < D && l < L) {
        tile[threadIdx.y][threadIdx.x] = proj[((long long)b * L + l) * (3 * D) + d];
    }
    __syncthreads();
    int l2 = blockIdx.y * TILE + threadIdx.x;
    int d2 = blockIdx.x * TILE + threadIdx.y;
    if (d2 < D && l2 < L) {
        out[((long long)b * D + d2) * N + l2] = tile[threadIdx.x][threadIdx.y];
    }
}

torch::Tensor pack_v(torch::Tensor projected) {
    int B = (int)projected.size(0);
    int L = (int)projected.size(1);
    int D = (int)projected.size(2) / 3;
    int N = 2 * L;
    auto out = torch::empty({B, D, N}, projected.options());
    cudaMemsetAsync(out.data_ptr<float>(), 0, (size_t)B * D * N * sizeof(float),
                    at::cuda::getCurrentCUDAStream());
    dim3 block(TILE, TILE);
    dim3 grid((D + TILE - 1) / TILE, (L + TILE - 1) / TILE, B);
    pack_v_tiled_kernel<<<grid, block, 0, at::cuda::getCurrentCUDAStream()>>>(
        projected.data_ptr<float>(), out.data_ptr<float>(), B, L, D, N);
    return out;
}

__device__ __forceinline__ float gate_conv3(const float* __restrict__ proj,
                                            const float* __restrict__ w,
                                            const float* __restrict__ bias,
                                            int b, int l, int d,
                                            int L, int D, int gate_idx) {
    const long long base = ((long long)b * L) * (3 * D) + (long long)gate_idx * D + d;
    float acc = bias[d];
    const float* wd = w + (long long)d * 3;
    if (l > 0)     acc += wd[0] * proj[base + (long long)(l - 1) * (3 * D)];
                   acc += wd[1] * proj[base + (long long) l      * (3 * D)];
    if (l + 1 < L) acc += wd[2] * proj[base + (long long)(l + 1) * (3 * D)];
    return acc;
}

__global__ void fuse_gate_transpose_kernel(const float* __restrict__ y,
                                           const float* __restrict__ proj,
                                           const float* __restrict__ w0,
                                           const float* __restrict__ b0,
                                           const float* __restrict__ w1,
                                           const float* __restrict__ b1,
                                           float* __restrict__ out,
                                           int B, int L, int D, int N) {
    long long idx = (long long)blockIdx.x * blockDim.x + threadIdx.x;
    long long total = (long long)B * L * D;
    if (idx >= total) return;
    int d = idx % D;
    long long q = idx / D;
    int l = q % L;
    int b = q / L;
    float yy = y[((long long)b * D + d) * N + l];
    float g0 = gate_conv3(proj, w0, b0, b, l, d, L, D, 1);
    float g1 = gate_conv3(proj, w1, b1, b, l, d, L, D, 2);
    out[idx] = yy * g0 * g1;
}

torch::Tensor fuse_gate_transpose(torch::Tensor y, torch::Tensor projected,
                                  torch::Tensor w0, torch::Tensor b0,
                                  torch::Tensor w1, torch::Tensor b1) {
    int B = (int)projected.size(0);
    int L = (int)projected.size(1);
    int D = (int)projected.size(2) / 3;
    int N = (int)y.size(2);
    auto out = torch::empty({B, L, D}, projected.options());
    int threads = 256;
    long long total = (long long)B * L * D;
    int blocks = (int)((total + threads - 1) / threads);
    fuse_gate_transpose_kernel<<<blocks, threads, 0, at::cuda::getCurrentCUDAStream()>>>(
        y.data_ptr<float>(), projected.data_ptr<float>(),
        w0.data_ptr<float>(), b0.data_ptr<float>(),
        w1.data_ptr<float>(), b1.data_ptr<float>(),
        out.data_ptr<float>(), B, L, D, N);
    return out;
}
"""

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

class ModelNew(nn.Module):
    def __init__(self, d_model: int, seq_len: int, order: int, short_filter_size: int):
        super(ModelNew, self).__init__()
        self.d_model = d_model
        self.seq_len = seq_len
        self.order = order
        self.fft_size = 2 * seq_len
        self.in_proj = nn.Linear(d_model, d_model * (order + 1), bias=True)
        self.short_convs = nn.ModuleList([
            nn.Conv1d(d_model, d_model, short_filter_size,
                      padding=short_filter_size // 2, groups=d_model)
            for _ in range(order)
        ])
        self.filter_mlp = nn.Sequential(nn.Linear(1, 64), nn.SiLU(), nn.Linear(64, d_model))
        self.register_buffer('filter_positions', torch.linspace(0, 1, seq_len).unsqueeze(-1))
        self.out_proj = nn.Linear(d_model, d_model, bias=True)
        self._cached_h_f = None
        self._cached_versions = None

    def _get_h_f(self):
        versions = (self.filter_mlp[0].weight._version, self.filter_mlp[0].bias._version,
                    self.filter_mlp[2].weight._version, self.filter_mlp[2].bias._version,
                    self.filter_positions.device, self.filter_positions.dtype)
        if self._cached_h_f is None or self._cached_versions != versions:
            h = self.filter_mlp(self.filter_positions).transpose(0, 1).contiguous()
            self._cached_h_f = torch.fft.rfft(F.pad(h, (0, self.seq_len)), dim=-1)
            self._cached_versions = versions
        return self._cached_h_f

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        projected = self.in_proj(x)
        v_padded = _ext.pack_v(projected)
        v_f = torch.fft.rfft(v_padded, dim=-1)
        v_f.mul_(self._get_h_f().unsqueeze(0))
        y = torch.fft.irfft(v_f, n=self.fft_size, dim=-1)
        z = _ext.fuse_gate_transpose(y, projected,
                                     self.short_convs[0].weight, self.short_convs[0].bias,
                                     self.short_convs[1].weight, self.short_convs[1].bias)
        return self.out_proj(z)