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 ppo_loss_kernel(
    const float* __restrict__ new_log_probs,
    const float* __restrict__ old_log_probs,
    const float* __restrict__ advantages,
    float* __restrict__ output,
    float clip_param,
    int n
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float ratio = expf(new_log_probs[idx] - old_log_probs[idx]);
        float adv = advantages[idx];

        float surr1 = ratio * adv;

        float low = 1.0f - clip_param;
        float high = 1.0f + clip_param;
        float ratio_clipped = fminf(fmaxf(ratio, low), high);

        float surr2 = ratio_clipped * adv;

        output[idx] = -fminf(surr1, surr2);
    }
}

torch::Tensor ppo_loss_cuda(torch::Tensor new_log_probs, torch::Tensor old_log_probs, torch::Tensor advantages, float clip_param) {
    int n = new_log_probs.numel();
    
    auto output = at::empty_like(new_log_probs);

    int threads = 256;
    int blocks = (n + threads - 1) / threads;

    ppo_loss_kernel<<<blocks, threads>>>(
        new_log_probs.data_ptr<float>(),
        old_log_probs.data_ptr<float>(),
        advantages.data_ptr<float>(),
        output.data_ptr<float>(),
        clip_param,
        n
    );

    return output.mean();
}
"""

cpp_source = """
torch::Tensor ppo_loss_cuda(torch::Tensor new_log_probs, torch::Tensor old_log_probs, torch::Tensor advantages, float clip_param);
"""

ppo_loss = load_inline(
    name="ppo_loss",
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=["ppo_loss_cuda"],
    verbose=False
)


class ModelNew(nn.Module):
    def __init__(self, clip_param=0.2):
        super(ModelNew, self).__init__()
        self.clip_param = clip_param

    def forward(self, new_log_probs, old_log_probs, advantages):
        return ppo_loss.ppo_loss_cuda(new_log_probs, old_log_probs, advantages, self.clip_param)