import math

import torch
import torch.nn as nn

# Default ModelNew is PyTorch-only. run_and_check: use check_kernel=False with backend=cuda
# until the candidate file contains a real CUDA kernel (see kernel_static_checker).


class Model(nn.Module):
    """
    **Batched forward kinematics (FK)** for a **planar serial** revolute chain.

    Fixed link lengths ``L_j`` (buffer), joint angles ``θ_j`` (batch ``B``). For each
    link, the heading is ``Σ_{k≤j} θ_k``; joint positions are the cumulative sum of
    ``L_j [cos(heading), sin(heading)]``—fully **vectorized** over ``B`` and ``J``.

    Returns positions of the **end of each link** in the base frame ``(B, J, 2)``
    (last row is the end-effector in the plane). Typical fusion target for
    batched geometry in calibration, whole-body control, and VLA motion heads.
    """

    def __init__(self, num_joints: int, link_scale: float = 1.0):
        super().__init__()
        assert num_joints >= 1
        self.num_joints = num_joints
        lengths = torch.ones(num_joints, dtype=torch.float32) * float(link_scale)
        self.register_buffer("link_lengths", lengths)

    def forward(self, joint_angles: torch.Tensor) -> torch.Tensor:
        """
        Args:
            joint_angles: (B, J) radians, one revolute per link.
        Returns:
            positions: (B, J, 2) world-frame (x, y) at each link end; ``[..., -1, :]`` is EE.
        """
        B, J = joint_angles.shape
        assert J == self.num_joints

        heading = torch.cumsum(joint_angles, dim=-1)
        c = torch.cos(heading)
        s = torch.sin(heading)
        direction = torch.stack((c, s), dim=-1)
        displacements = direction * self.link_lengths.view(1, J, 1)
        return torch.cumsum(displacements, dim=1)


class ModelNew(nn.Module):
    """KernelBench candidate entry point; default delegates to ``Model``."""

    def __init__(self, num_joints: int, link_scale: float = 1.0):
        super().__init__()
        self._impl = Model(num_joints, link_scale=link_scale)

    def forward(self, joint_angles: torch.Tensor) -> torch.Tensor:
        return self._impl(joint_angles)


batch_size = 32
num_joints = 7
link_scale = 0.15


def get_inputs():
    p = popcorn_pri
    return [torch.empty(p.jitter_int(batch_size), num_joints).uniform_(-math.pi, math.pi)]


def get_init_inputs():
    return [num_joints, link_scale]

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

cuda_source = """
#include <cuda.h>
#include <cuda_runtime.h>
#include <cmath>

// Optimized kernel - keep it simple for small batch sizes
// Use sincosf which is faster than separate sin/cos calls
__global__ void __launch_bounds__(64) forward_kinematics_kernel(
    const float* __restrict__ joint_angles,
    const float* __restrict__ link_lengths,
    float* __restrict__ output,
    int batch_size,
    int num_joints
) {
    int batch_idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (batch_idx >= batch_size) return;
    
    float heading = 0.0f;
    float x = 0.0f;
    float y = 0.0f;
    
    // Pre-compute strides
    const float* angles = joint_angles + batch_idx * num_joints;
    float* pos = output + batch_idx * num_joints * 2;
    
    // Unrolled loop
    for (int j = 0; j < num_joints; j++) {
        heading += angles[j];
        float c = cosf(heading);
        float s = sinf(heading);
        x += link_lengths[j] * c;
        y += link_lengths[j] * s;
        pos[j * 2] = x;
        pos[j * 2 + 1] = y;
    }
}

torch::Tensor forward_kinematics_cuda(torch::Tensor joint_angles, torch::Tensor link_lengths) {
    int batch_size = joint_angles.size(0);
    int num_joints = joint_angles.size(1);
    
    auto output = torch::empty({batch_size, num_joints, 2}, joint_angles.options());
    
    const int block_size = 64;
    const int num_blocks = (batch_size + block_size - 1) / block_size;
    
    forward_kinematics_kernel<<<num_blocks, block_size>>>(
        joint_angles.data_ptr<float>(),
        link_lengths.data_ptr<float>(),
        output.data_ptr<float>(),
        batch_size,
        num_joints
    );
    
    return output;
}
"""

ext = load_inline(
    name="forward_kinematics_ext",
    cpp_sources=["torch::Tensor forward_kinematics_cuda(torch::Tensor joint_angles, torch::Tensor link_lengths);"],
    cuda_sources=[cuda_source],
    functions=["forward_kinematics_cuda"],
    with_cuda=True
)

class ModelNew(nn.Module):
    def __init__(self, num_joints: int, link_scale: float = 1.0):
        super().__init__()
        self.num_joints = num_joints
        lengths = torch.ones(num_joints, dtype=torch.float32) * float(link_scale)
        self.register_buffer("link_lengths", lengths)
        self.ext = ext
    
    def forward(self, joint_angles: torch.Tensor) -> torch.Tensor:
        return self.ext.forward_kinematics_cuda(joint_angles, self.link_lengths)


batch_size = 32
num_joints = 7
link_scale = 0.15


def get_inputs():
    return [torch.empty(batch_size, num_joints).uniform_(-math.pi, math.pi)]


def get_init_inputs():
    return [num_joints, link_scale]