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

cuda_source = r'''
#include <cuda_runtime.h>
#include <math.h>

__global__ void ball_query_max_pool_kernel(
    const float* __restrict__ points,
    const float* __restrict__ queries,
    const float* __restrict__ features,
    float* __restrict__ output,
    float radius_sq,
    int B, int P, int Q, int C
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int total = B * Q * C;
    if (idx >= total) return;
    
    int c = idx % C;
    int q = (idx / C) % Q;
    int b = idx / (Q * C);
    
    float max_val = -INFINITY;
    bool has_neighbor = false;
    
    float qx = queries[b * Q * 3 + q * 3];
    float qy = queries[b * Q * 3 + q * 3 + 1];
    float qz = queries[b * Q * 3 + q * 3 + 2];
    
    for (int p = 0; p < P; p++) {
        float px = points[b * P * 3 + p * 3];
        float py = points[b * P * 3 + p * 3 + 1];
        float pz = points[b * P * 3 + p * 3 + 2];
        
        float dx = px - qx;
        float dy = py - qy;
        float dz = pz - qz;
        float dist_sq = dx*dx + dy*dy + dz*dz;
        
        if (dist_sq <= radius_sq) {
            float feat = features[b * P * C + p * C + c];
            if (feat > max_val) max_val = feat;
            has_neighbor = true;
        }
    }
    
    if (!has_neighbor) max_val = 0.0f;
    output[b * Q * C + q * C + c] = max_val;
}
'''

cpp_source = r'''
#include <torch/extension.h>
#include <cuda_runtime.h>

void ball_query_max_pool_kernel(
    const float* points,
    const float* queries,
    const float* features,
    float* output,
    float radius_sq,
    int B, int P, int Q, int C
);

torch::Tensor ball_query_max_pool_cuda(
    torch::Tensor points,
    torch::Tensor queries,
    torch::Tensor features,
    float radius_sq
) {
    int B = points.size(0);
    int P = points.size(1);
    int Q = queries.size(1);
    int C = features.size(2);
    
    auto output = torch::zeros({B, Q, C}, features.options());
    
    int total_threads = B * Q * C;
    int block_size = 256;
    int num_blocks = (total_threads + block_size - 1) / block_size;
    
    ball_query_max_pool_kernel<<<num_blocks, block_size>>>(
        points.data_ptr<float>(),
        queries.data_ptr<float>(),
        features.data_ptr<float>(),
        output.data_ptr<float>(),
        radius_sq,
        B, P, Q, C
    );
    
    return output;
}
'''

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

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

    def forward(self, points, queries, features):
        return self.module.ball_query_max_pool_cuda(
            points.contiguous(),
            queries.contiguous(),
            features.contiguous(),
            self.radius_sq
        )