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 torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

_cpp_src = r'''
#include <torch/extension.h>
torch::Tensor fk_forward(torch::Tensor joint_angles, double link_scale);
'''

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

__global__ void fk_kernel7_warp(const float* __restrict__ angles, float* __restrict__ out, int B, float link_scale) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int b = tid >> 5;
    int lane = threadIdx.x & 31;
    if (b >= B) return;

    float v = (lane < 7) ? angles[b * 7 + lane] : 0.0f;
    const unsigned mask = 0xffffffffu;

    #pragma unroll
    for (int off = 1; off < 8; off <<= 1) {
        float u = __shfl_up_sync(mask, v, off);
        if (lane >= off) v += u;
    }

    float s = 0.0f, c = 0.0f;
    if (lane < 7) __sincosf(v, &s, &c);

    float x = (lane < 7) ? link_scale * c : 0.0f;
    float y = (lane < 7) ? link_scale * s : 0.0f;

    #pragma unroll
    for (int off = 1; off < 8; off <<= 1) {
        float ux = __shfl_up_sync(mask, x, off);
        float uy = __shfl_up_sync(mask, y, off);
        if (lane >= off) { x += ux; y += uy; }
    }

    if (lane < 7) {
        int idx = (b * 7 + lane) * 2;
        out[idx] = x;
        out[idx + 1] = y;
    }
}

__global__ void fk_kernel_generic(const float* __restrict__ angles, float* __restrict__ out, int B, int J, float link_scale) {
    int b = blockIdx.x * blockDim.x + threadIdx.x;
    if (b >= B) return;
    const float* a = angles + (long long)b * J;
    float heading = 0.0f, x = 0.0f, y = 0.0f;
    for (int j = 0; j < J; ++j) {
        heading += a[j];
        float s, c;
        __sincosf(heading, &s, &c);
        x += link_scale * c;
        y += link_scale * s;
        long long idx = ((long long)b * J + j) * 2;
        out[idx] = x;
        out[idx + 1] = y;
    }
}

torch::Tensor fk_forward(torch::Tensor joint_angles, double link_scale) {
    const int B = (int)joint_angles.size(0);
    const int J = (int)joint_angles.size(1);
    auto out = torch::empty({B, J, 2}, joint_angles.options());
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if (J == 7) {
        const int threads = 128;
        const int blocks = (B * 32 + threads - 1) / threads;
        fk_kernel7_warp<<<blocks, threads, 0, stream>>>(joint_angles.data_ptr<float>(), out.data_ptr<float>(), B, (float)link_scale);
    } else {
        const int threads = 128;
        const int blocks = (B + threads - 1) / threads;
        fk_kernel_generic<<<blocks, threads, 0, stream>>>(joint_angles.data_ptr<float>(), out.data_ptr<float>(), B, J, (float)link_scale);
    }
    return out;
}
'''

_fk_ext = load_inline(
    name='fk_planar_ext_final',
    cpp_sources=_cpp_src,
    cuda_sources=_cuda_src,
    functions=['fk_forward'],
    extra_cuda_cflags=['-O3', '--use_fast_math'],
    verbose=False,
)

class ModelNew(nn.Module):
    def __init__(self, num_joints: int, link_scale: float = 1.0):
        super().__init__()
        assert num_joints >= 1
        self.num_joints = num_joints
        self.link_scale = float(link_scale)
        self.register_buffer('link_lengths', torch.ones(num_joints, dtype=torch.float32) * self.link_scale)

    def forward(self, joint_angles: torch.Tensor) -> torch.Tensor:
        return _fk_ext.fk_forward(joint_angles, self.link_scale)