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):
    """
    Farthest Point Sampling (FPS) on 3D point coordinates (B, N, 3).

    Used in navigation / manipulation perception (PointNet++, PVCNN, voxelization
    pipelines) to subsample a dense cloud while spreading support across the shape.

    Greedy construction (reference fusion target): fix the first index to **0** for
    determinism, then repeatedly choose the point whose **minimum** squared distance
    to the set of already chosen points is **largest** (classical FPS).
    """

    def __init__(self, num_samples: int):
        super().__init__()
        assert num_samples >= 1
        self.num_samples = num_samples

    def forward(self, points: torch.Tensor) -> torch.Tensor:
        """
        Args:
            points: (B, N, 3) point coordinates.
        Returns:
            indices: (B, M) int64, each row is M chosen indices in [0, N).
        """
        B, N, _ = points.shape
        M = self.num_samples
        assert M <= N, "FPS cannot sample more points than exist"
        device = points.device
        dtype = points.dtype

        batch_idx = torch.arange(B, device=device)
        indices = torch.empty(B, M, dtype=torch.long, device=device)

        # min squared distance from each point to the current chosen set (large init)
        min_dist_sq = torch.full((B, N), float("inf"), device=device, dtype=dtype)
        # First center: point index 0 (deterministic)
        farthest = torch.zeros(B, dtype=torch.long, device=device)

        for i in range(M):
            indices[:, i] = farthest
            center = points[batch_idx, farthest]
            dist_sq = ((points - center.unsqueeze(1)) ** 2).sum(dim=-1)
            min_dist_sq = torch.minimum(min_dist_sq, dist_sq)
            # Exclude the just-chosen index from being selected again
            min_dist_sq[batch_idx, farthest] = -1.0
            farthest = min_dist_sq.argmax(dim=-1)

        return indices


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

    def __init__(self, num_samples: int):
        super().__init__()
        self._impl = Model(num_samples)

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


batch_size = 4
num_points = 1024
num_samples = 128


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


def get_init_inputs():
    return [num_samples]

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>
#include <cfloat>

__global__ void fps_kernel(
    const float* __restrict__ points,
    int64_t* __restrict__ indices,
    float* __restrict__ min_dist_sq,
    const int B,
    const int N,
    const int M
) {
    int b = blockIdx.x;
    if (b >= B) return;

    const float* pts = points + (size_t)b * N * 3;
    int64_t* idx = indices + (size_t)b * M;
    float* min_dist = min_dist_sq + (size_t)b * N;

    extern __shared__ char smem[];
    float* sdata = (float*)smem;
    int* sidx = (int*)(sdata + blockDim.x);

    for (int i = threadIdx.x; i < N; i += blockDim.x) {
        min_dist[i] = FLT_MAX;
    }
    __syncthreads();

    __shared__ int farthest_shared;
    if (threadIdx.x == 0) farthest_shared = 0;
    __syncthreads();

    for (int i = 0; i < M; ++i) {
        int farthest = farthest_shared;
        idx[i] = farthest;
        float cx = pts[farthest * 3];
        float cy = pts[farthest * 3 + 1];
        float cz = pts[farthest * 3 + 2];

        for (int j = threadIdx.x; j < N; j += blockDim.x) {
            float dx = pts[j * 3] - cx;
            float dy = pts[j * 3 + 1] - cy;
            float dz = pts[j * 3 + 2] - cz;
            float dist = dx*dx + dy*dy + dz*dz;
            if (dist < min_dist[j]) min_dist[j] = dist;
        }
        __syncthreads();

        if (threadIdx.x == 0) {
            min_dist[farthest] = -1.0f;
        }
        __syncthreads();

        float local_max = -1.0f;
        int local_idx = 0;
        for (int j = threadIdx.x; j < N; j += blockDim.x) {
            float val = min_dist[j];
            if (val > local_max) {
                local_max = val;
                local_idx = j;
            }
        }

        sdata[threadIdx.x] = local_max;
        sidx[threadIdx.x] = local_idx;
        __syncthreads();

        for (int s = blockDim.x / 2; s > 0; s >>= 1) {
            if (threadIdx.x < s) {
                if (sdata[threadIdx.x + s] > sdata[threadIdx.x]) {
                    sdata[threadIdx.x] = sdata[threadIdx.x + s];
                    sidx[threadIdx.x] = sidx[threadIdx.x + s];
                }
            }
            __syncthreads();
        }

        if (threadIdx.x == 0) {
            farthest_shared = sidx[0];
        }
        __syncthreads();
    }
}

torch::Tensor fps_cuda(torch::Tensor points, int num_samples) {
    int B = points.size(0);
    int N = points.size(1);
    int M = num_samples;
    auto options = torch::TensorOptions().dtype(torch::kInt64).device(points.device());
    auto indices = torch::empty({B, M}, options);
    auto min_dist = torch::empty({B, N}, points.options());
    int threads = 256;
    size_t smem = threads * (sizeof(float) + sizeof(int));
    fps_kernel<<<B, threads, smem>>>(
        points.data_ptr<float>(),
        indices.data_ptr<int64_t>(),
        min_dist.data_ptr<float>(),
        B, N, M
    );
    return indices;
}
"""

cpp_source = """
torch::Tensor fps_cuda(torch::Tensor points, int num_samples);
"""

module = load_inline(
    name="fps_module",
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=["fps_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3"]
)

class ModelNew(nn.Module):
    def __init__(self, num_samples: int):
        super().__init__()
        self.num_samples = num_samples

    def forward(self, points: torch.Tensor) -> torch.Tensor:
        return module.fps_cuda(points, self.num_samples)