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

cuda_source = """
#include <torch/extension.h>
#include <cuda_runtime.h>

__global__ void kl_elem_kernel(const float* __restrict__ x, 
                               const float* __restrict__ y, 
                               float* __restrict__ out, 
                               int size) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < size) {
        float eps = 1e-8f;
        float val_x = fabsf(x[idx]) + eps;
        float val_y = fabsf(y[idx]) + eps;

        // KL element: y * log(y / x) = y * (log(y) - log(x))
        out[idx] = val_y * (logf(val_y) - logf(val_x));
    }
}

torch::Tensor kl_elem_cuda(torch::Tensor x, torch::Tensor y) {
    auto size = x.numel();
    auto out = torch::empty_like(x);

    const int block_size = 256;
    int grid_size = (size + block_size - 1) / block_size;

    kl_elem_kernel<<<grid_size, block_size>>>(
        x.data_ptr<float>(),
        y.data_ptr<float>(),
        out.data_ptr<float>(),
        size
    );

    return out;
}
"""

cpp_source = "torch::Tensor kl_elem_cuda(torch::Tensor x, torch::Tensor y);"

module = load_inline(
    name="kullbackleibler_layernorm_ext",
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=["kl_elem_cuda"],
    verbose=False,
    with_cuda=True
)


class ModelNew(nn.Module):
    def __init__(self, dim):
        super(ModelNew, self).__init__()
        self.ln = nn.LayerNorm(dim)
        self.op = module

    def forward(self, x, y):
        kl_elem = self.op.kl_elem_cuda(x.contiguous(), y.contiguous())
        out = self.ln(kl_elem)
        return out.mean()