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_runtime.h>
#include <float.h>

__global__ void fps_kernel(const float* points, int64_t* indices, float* min_dist_global, int B, int N, int M) {
    int b = blockIdx.x;
    float* min_dist = min_dist_global + b * N;
    const float* bp = points + (size_t)b * N * 3;
    int64_t* bi = indices + (size_t)b * M;

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

    int farthest = 0;
    if (threadIdx.x == 0) bi[0] = 0;
    __syncthreads();

    for (int m = 1; m < M; m++) {
        float cx = bp[farthest*3], cy = bp[farthest*3+1], cz = bp[farthest*3+2];
        for (int p = threadIdx.x; p < N; p += blockDim.x) {
            float dx = bp[p*3] - cx, dy = bp[p*3+1] - cy, dz = bp[p*3+2] - cz;
            float d2 = dx*dx + dy*dy + dz*dz;
            if (d2 < min_dist[p]) min_dist[p] = d2;
        }
        __syncthreads();
        if (threadIdx.x == 0) min_dist[farthest] = -1.0f;
        __syncthreads();

        if (threadIdx.x == 0) {
            float lmax = -FLT_MAX; int lidx = 0;
            for (int p = 0; p < N; p++) {
                if (min_dist[p] > lmax) { lmax = min_dist[p]; lidx = p; }
            }
            farthest = lidx;
            bi[m] = farthest;
        }
        __syncthreads();
    }
}

torch::Tensor fps_cuda(torch::Tensor pts, int M) {
    int B = pts.size(0), N = pts.size(1);
    auto idx = torch::empty({B, M}, torch::TensorOptions().dtype(torch::kInt64).device(pts.device()));
    auto min_dist = torch::empty({B, N}, pts.options());
    int tpb = 256;
    fps_kernel<<<B, tpb>>>(pts.data_ptr<float>(), idx.data_ptr<int64_t>(), min_dist.data_ptr<float>(), B, N, M);
    return idx;
}
'''

cpp_src = 'torch::Tensor fps_cuda(torch::Tensor, int);'

mod = load_inline(name='m4', cpp_sources=[cpp_src], cuda_sources=[cuda_source], functions=['fps_cuda'])

class ModelNew(nn.Module):
    def __init__(self, num_samples):
        super().__init__()
        self.M = num_samples
    def forward(self, x):
        return mod.fps_cuda(x.contiguous(), self.M)