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 js_elem_kernel(const float* __restrict__ p, 
                               const float* __restrict__ q, 
                               float* __restrict__ out, 
                               int size) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < size) {
        float eps = 1e-8f;
        float val_p = fabsf(p[idx]) + eps;
        float val_q = fabsf(q[idx]) + eps;
        float val_m = 0.5f * (val_p + val_q);

        float kl_p = val_p * (logf(val_p) - logf(val_m));
        float kl_q = val_q * (logf(val_q) - logf(val_m));

        out[idx] = 0.5f * (kl_p + kl_q);
    }
}

torch::Tensor js_elem_cuda(torch::Tensor p, torch::Tensor q) {
    auto size = p.numel();
    auto out = torch::empty_like(p);

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

    js_elem_kernel<<<grid_size, block_size>>>(
        p.data_ptr<float>(),
        q.data_ptr<float>(),
        out.data_ptr<float>(),
        size
    );

    return out;
}
"""

cpp_source = "torch::Tensor js_elem_cuda(torch::Tensor p, torch::Tensor q);"

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


class ModelNew(nn.Module):
    def __init__(self, num_channels, num_groups=32):
        super(ModelNew, self).__init__()
        self.gn = nn.GroupNorm(num_groups, num_channels)
        self.op = module

    def forward(self, x, y):
        js_elem = self.op.js_elem_cuda(x.contiguous(), y.contiguous())
        out = self.gn(js_elem)
        return out.mean()