import torch
import torch.nn as nn
import torch.nn.functional as F

# 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).


def _quaternion_to_rotation_matrix(q: torch.Tensor) -> torch.Tensor:
    """
    Unit quaternion ``q = (w, x, y, z)`` (scalar-first) to rotation matrix ``R``,
    shape ``(B, 3, 3)``, with ``p' = p @ R.T`` for row vectors ``p`` (``N×3``).
    """
    q = F.normalize(q, dim=-1, eps=1e-8)
    w, x, y, z = q.unbind(dim=-1)

    xx = x * x
    yy = y * y
    zz = z * z
    xy = x * y
    xz = x * z
    yz = y * z
    wx = w * x
    wy = w * y
    wz = w * z

    m00 = 1.0 - 2.0 * (yy + zz)
    m01 = 2.0 * (xy - wz)
    m02 = 2.0 * (xz + wy)
    m10 = 2.0 * (xy + wz)
    m11 = 1.0 - 2.0 * (xx + zz)
    m12 = 2.0 * (yz - wx)
    m20 = 2.0 * (xz - wy)
    m21 = 2.0 * (yz + wx)
    m22 = 1.0 - 2.0 * (xx + yy)

    row0 = torch.stack((m00, m01, m02), dim=-1)
    row1 = torch.stack((m10, m11, m12), dim=-1)
    row2 = torch.stack((m20, m21, m22), dim=-1)
    return torch.stack((row0, row1, row2), dim=-2)


class Model(nn.Module):
    """
    **Fused vectorized pose transform**: rotation from unit quaternion ``q`` and
    translation ``t`` applied to a batch of 3D points.

    One forward fuses **quaternion → ``3×3``** (Hamilton, scalar-first ``w,x,y,z``)
    with **batched matmul** and **broadcast translation**—typical for LiDAR / camera
    extrinsics, hand-eye calibration, and world ↔ body frame changes in VLA stacks.

    ``points' = points @ R.T + t`` with ``points`` row-wise ``(B, N, 3)``.
    """

    def __init__(self):
        super().__init__()

    def forward(
        self,
        points: torch.Tensor,
        quat: torch.Tensor,
        trans: torch.Tensor,
    ) -> torch.Tensor:
        """
        Args:
            points: (B, N, 3) points in the **source** frame.
            quat: (B, 4) unit quaternion ``(w, x, y, z)`` mapping source → target rotation.
            trans: (B, 3) translation in the **target** frame (same convention as ``R``).
        Returns:
            (B, N, 3) points in the **target** frame.
        """
        B, N, _ = points.shape
        assert quat.shape == (B, 4) and trans.shape == (B, 3)

        R = _quaternion_to_rotation_matrix(quat)
        return torch.matmul(points, R.transpose(-1, -2)) + trans.unsqueeze(1)


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

    def __init__(self):
        super().__init__()
        self._impl = Model()

    def forward(
        self,
        points: torch.Tensor,
        quat: torch.Tensor,
        trans: torch.Tensor,
    ) -> torch.Tensor:
        return self._impl(points, quat, trans)


batch_size = 8
num_points = 512


def get_inputs():
    p = popcorn_pri
    points = torch.randn(p.jitter_int(batch_size), p.jitter_int(num_points), 3)
    quat = torch.randn(p.jitter_int(batch_size), 4)
    trans = torch.randn(p.jitter_int(batch_size), 3)
    return [points, quat, trans]


def get_init_inputs():
    return []

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

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

__device__ __forceinline__ void quaternion_to_rotation_matrix(const float* q, float* R) {
    float qw = q[0], qx = q[1], qy = q[2], qz = q[3];
    float norm = sqrtf(qw*qw + qx*qx + qy*qy + qz*qz);
    qw /= norm; qx /= norm; qy /= norm; qz /= norm;
    
    float xx = qx * qx;
    float yy = qy * qy;
    float zz = qz * qz;
    float xy = qx * qy;
    float xz = qx * qz;
    float yz = qy * qz;
    float wx = qw * qx;
    float wy = qw * qy;
    float wz = qw * qz;
    
    R[0] = 1.0f - 2.0f * (yy + zz);
    R[1] = 2.0f * (xy - wz);
    R[2] = 2.0f * (xz + wy);
    R[3] = 2.0f * (xy + wz);
    R[4] = 1.0f - 2.0f * (xx + zz);
    R[5] = 2.0f * (yz - wx);
    R[6] = 2.0f * (xz - wy);
    R[7] = 2.0f * (yz + wx);
    R[8] = 1.0f - 2.0f * (xx + yy);
}

__global__ void fused_transform_kernel(
    const float* __restrict__ points,
    const float* __restrict__ quats,
    const float* __restrict__ trans,
    float* __restrict__ output,
    int B, int N
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int total = B * N;
    if (idx >= total) return;
    
    int b = idx / N;
    int i = idx % N;
    
    float R[9];
    quaternion_to_rotation_matrix(quats + b * 4, R);
    
    float px = points[(b * N + i) * 3 + 0];
    float py = points[(b * N + i) * 3 + 1];
    float pz = points[(b * N + i) * 3 + 2];
    
    float tx = trans[b * 3 + 0];
    float ty = trans[b * 3 + 1];
    float tz = trans[b * 3 + 2];
    
    float ox = px * R[0] + py * R[1] + pz * R[2] + tx;
    float oy = px * R[3] + py * R[4] + pz * R[5] + ty;
    float oz = px * R[6] + py * R[7] + pz * R[8] + tz;
    
    output[(b * N + i) * 3 + 0] = ox;
    output[(b * N + i) * 3 + 1] = oy;
    output[(b * N + i) * 3 + 2] = oz;
}

torch::Tensor fused_transform_cuda(torch::Tensor points, torch::Tensor quat, torch::Tensor trans) {
    int B = points.size(0);
    int N = points.size(1);
    
    auto output = torch::empty({B, N, 3}, points.options());
    
    int total = B * N;
    int threads = 256;
    int blocks = (total + threads - 1) / threads;
    
    fused_transform_kernel<<<blocks, threads>>>(
        points.data_ptr<float>(),
        quat.data_ptr<float>(),
        trans.data_ptr<float>(),
        output.data_ptr<float>(),
        B, N
    );
    
    return output;
}
"""

cpp_source = """
torch::Tensor fused_transform_cuda(torch::Tensor points, torch::Tensor quat, torch::Tensor trans);
"""

module = load_inline(
    name="fused_transform",
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=["fused_transform_cuda"],
    verbose=False
)

class ModelNew(nn.Module):
    def __init__(self):
        super().__init__()
    
    def forward(self, points, quat, trans):
        return module.fused_transform_cuda(points, quat, trans)