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):
    """
    Fused radius (ball) neighborhood query over a 3D point map with per-query feature
    aggregation—common in navigation / manipulation perception (PointNet++-style set
    abstraction, local maps around end-effector or goal queries on LiDAR / depth).

    Single forward fusion target:
      1) pairwise Euclidean distances between each query center and every point,
      2) membership in the closed ball of radius ``r``,
      3) channel-wise **max** pooling of point features over in-ball neighbors.

    Points with no neighbors inside the radius contribute a **zero** feature vector
    for that query (instead of ``-inf`` max).
    """

    def __init__(self, in_channels: int, radius: float):
        super().__init__()
        self.in_channels = in_channels
        self.register_buffer("radius_sq", torch.tensor(float(radius) ** 2))

    def forward(
        self,
        points: torch.Tensor,
        queries: torch.Tensor,
        features: torch.Tensor,
    ) -> torch.Tensor:
        """
        Args:
            points: (B, P, 3) point coordinates (e.g. world or sensor frame).
            queries: (B, Q, 3) ball centers (e.g. navigation waypoints, grasp samples).
            features: (B, P, C) per-point descriptors (geometry or lifted RGB-D).
        Returns:
            (B, Q, C) max-pooled features over neighbors with ||p - q||_2 <= r.
        """
        B, P, _ = points.shape
        _, Q, _ = queries.shape
        _, P2, C = features.shape
        assert P == P2 and C == self.in_channels

        # Squared distances (B, Q, P)
        dist_sq = torch.cdist(queries, points, p=2.0) ** 2
        in_ball = dist_sq <= self.radius_sq

        # (B, 1, P, C) broadcasts with (B, Q, P) for masked max over P
        feats = features.unsqueeze(1)
        neg_large = torch.finfo(features.dtype).min
        masked = torch.where(
            in_ball.unsqueeze(-1),
            feats,
            torch.full_like(feats, neg_large),
        )
        pooled, _ = masked.max(dim=2)

        empty = in_ball.sum(dim=2) == 0
        if empty.any():
            pooled = pooled.masked_fill(empty.unsqueeze(-1), 0.0)

        return pooled


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

    def __init__(self, in_channels: int, radius: float):
        super().__init__()
        self._impl = Model(in_channels, radius)

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


batch_size = 4
num_points = 512
num_queries = 64
in_channels = 32
radius = 0.12


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


def get_init_inputs():
    return [in_channels, radius]

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

_cpp_src = r'''
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
void launcher(const float*, const float*, const float*, float*, int, int, int, int, float, cudaStream_t);

torch::Tensor ball_max_forward(torch::Tensor points, torch::Tensor queries, torch::Tensor features, double radius_sq) {
    int B = points.size(0), P = points.size(1), Q = queries.size(1), C = features.size(2);
    auto out = torch::empty({B, Q, C}, features.options());
    launcher(points.data_ptr<float>(), queries.data_ptr<float>(), features.data_ptr<float>(), out.data_ptr<float>(),
             B, P, Q, C, (float)radius_sq, at::cuda::getCurrentCUDAStream());
    return out;
}
'''

_cuda_src = r'''
#include <cuda_runtime.h>
#include <float.h>

__device__ __forceinline__ void atomicMaxFloat(float* addr, float val) {
    int* ai = (int*)addr;
    int old = *ai, assumed;
    while (val > __int_as_float(old)) {
        assumed = old;
        old = atomicCAS(ai, assumed, __float_as_int(val));
        if (old == assumed) break;
    }
}

__global__ __launch_bounds__(512, 1)
void ball_kernel(const float* __restrict__ points,
                 const float* __restrict__ queries,
                 const float* __restrict__ features,
                 float* __restrict__ out,
                 int B, int P, int Q, int C, float r2) {
    __shared__ float smax[32];
    __shared__ int sfound;
    int tid = threadIdx.x;
    int w = blockIdx.x;
    int b = w / Q;
    int q = w - b * Q;

    if (tid < 32) smax[tid] = -FLT_MAX;
    if (tid == 0) sfound = 0;
    __syncthreads();

    const float* qptr = queries + ((b * Q + q) * 3);
    float qx = qptr[0], qy = qptr[1], qz = qptr[2];
    const float* pbase = points + b * P * 3;
    const float* fbase = features + b * P * C;

    int p = tid;
    if (p < P) {
        const float* pp = pbase + p * 3;
        float dx = pp[0] - qx, dy = pp[1] - qy, dz = pp[2] - qz;
        float d2 = dx*dx + dy*dy + dz*dz;
        if (d2 <= r2) {
            sfound = 1;
            const float* fp = fbase + p * C;
            #pragma unroll
            for (int c = 0; c < 32; ++c) atomicMaxFloat(&smax[c], fp[c]);
        }
    }

    __syncthreads();
    if (tid < C && tid < 32) out[(b * Q + q) * C + tid] = sfound ? smax[tid] : 0.0f;
}

void launcher(const float* p, const float* q, const float* f, float* o,
              int B, int P, int Q, int C, float r, cudaStream_t s) {
    ball_kernel<<<B * Q, 512, 0, s>>>(p, q, f, o, B, P, Q, C, r);
}
'''

_ext = load_inline(
    name='ball_max_final_ext',
    cpp_sources=_cpp_src,
    cuda_sources=_cuda_src,
    functions=['ball_max_forward'],
    extra_cflags=['-O3'],
    extra_cuda_cflags=['-O3', '--use_fast_math'],
    verbose=False,
)

class ModelNew(nn.Module):
    def __init__(self, in_channels: int, radius: float):
        super().__init__()
        self.in_channels = in_channels
        self.radius_sq = float(radius) * float(radius)

    def forward(self, points: torch.Tensor, queries: torch.Tensor, features: torch.Tensor) -> torch.Tensor:
        return _ext.ball_max_forward(points.contiguous(), queries.contiguous(), features.contiguous(), self.radius_sq)