grok-4.3 · popcorn · pop_l123_st_gpt
compile_fail 22_FusedBallQuery3D.py
model
grok-4.3
turns / max
1 / 1
tool calls
1
tokens (total)
4030
wall clock
613.4s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
—
ref runtime
-1.00 μs
speedup
—
started
2026-05-11T09:35:59Z
compile error: Syntax error in custom generated code or ModelNew not found
timingkernelbench.timing
kernel runtime
—
ref runtime
—
speedup
—
kernel mean / median / std
—
ref mean / median / std
—
reference PyTorch source (the prompt the agent sees)
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]
turn 0 · FINAL · 1 tool call(s) · 613.4s turn · 23.4s LLM · tokens in=1,646 · out=1,568 · total=4,030 · 7,564 chars (click to expand)
turn 0613.4s turn · 23.4s LLM · 590.1s tools · tokens in=1,646 · out=1,568 · total=4,030 · FINAL
reasoning
The task is to replace the PyTorch model's forward with a custom CUDA kernel that's faster but numerically equivalent. I need to create a custom CUDA kernel to replace the PyTorch model's forward pass, focusing on computing max-pooled features over points within a radius r around each query point. The kernel should be numerically equivalent but faster. The computation involves, for each batch b, query q, and point p: check if the squared 3D distance between queries[b,q] and points[b,p] is ≤ r², and if so, consider the features[b,p,:] for max pooling per channel c. If no points are within range, output 0.
assistant```python import torch import torch.nn as nn from torch.utils.cpp_extension import load_inline _cuda_source = r''' #include <torch/extension.h> #include <cuda.h> #include <cuda_runtime.h> __global__ void ball_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 = -1e38f; bool has_any = false; const float* q_ptr = queries + (b * Q + q) * 3; float qx = q_ptr[0]; float qy = q_ptr[1]; float qz = q_ptr[2]; for (int p = 0; p < P; ++p) { const float* p_ptr = points + (b * P + p) * 3; float dx = qx - p_ptr[0]; float dy = qy - p_ptr[1]; float dz = qz - p_ptr[2]; float dist_sq = dx*dx + dy*dy + dz*dz; if (dist_sq <= radius_sq) { float f = features[(b * P + p) * C + c]; if (!has_any || f > max_val) { max_val = f; } has_any = true; } } output[(b * Q + q) * C + c] = has_any ? max_val : 0.0f; } torch::Tensor ball_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::empty({B, Q, C}, features.options()); int total = B * Q * C; int threads = 256; int blocks = (total + threads - 1) / threads; ball_max_pool_kernel<<<blocks, threads>>>( points.data_ptr<float>(), queries.data_ptr<float>(), features.data_ptr<float>(), output.data_ptr<float>(), radius_sq, B, P, Q, C ); return output; } ''' _cpp_source = r''' torch::Tensor ball_max_pool_cuda( torch::Tensor points, torch::Tensor queries, torch::Tensor features, float radius_sq ); ''' _ball_module = load_inline( name='ball_max_pool', cpp_sources=_cpp_source, cuda_sources=_cuda_source, functions=['ball_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 def forward( self, points: torch.Tensor, queries: torch.Tensor, features: torch.Tensor, ) -> torch.Tensor: return _ball_module.ball_max_pool_cuda( points, queries, features, self.radius_sq ) ```
call → submit_kernel
callsubmit_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n_cuda_source = r'''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\n__global__ void ball_max_pool_kernel(\n const float* __restrict__ points,\n const float* __restrict__ queries,\n const float* __restrict__ features,\n float* __restrict__ output,\n float radius_sq,\n int B, int P, int Q, int C\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * Q * C;\n if (idx >= total) return;\n\n int c = idx % C;\n int q = (idx / C) % Q;\n int b = idx / (Q * C);\n\n float max_val = -1e38f;\n bool has_any = false;\n\n const float* q_ptr = queries + (b * Q + q) * 3;\n float qx = q_ptr[0];\n float qy = q_ptr[1];\n float qz = q_ptr[2];\n\n for (int p = 0; p < P; ++p) {\n const float* p_ptr = points + (b * P + p) * 3;\n float dx = qx - p_ptr[0];\n float dy = qy - p_ptr[1];\n float dz = qz - p_ptr[2];\n float dist_sq = dx*dx + dy*dy + dz*dz;\n if (dist_sq <= radius_sq) {\n float f = features[(b * P + p) * C + c];\n if (!has_any || f > max_val) {\n max_val = f;\n }\n has_any = true;\n }\n }\n output[(b * Q + q) * C + c] = has_any ? max_val : 0.0f;\n}\n\ntorch::Tensor ball_max_pool_cuda(\n torch::Tensor points,\n torch::Tensor queries,\n torch::Tensor features,\n float radius_sq\n) {\n int B = points.size(0);\n int P = points.size(1);\n int Q = queries.size(1);\n int C = features.size(2);\n\n auto output = torch::empty({B, Q, C}, features.options());\n\n int total = B * Q * C;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n\n ball_max_pool_kernel<<<blocks, threads>>>(\n points.data_ptr<float>(),\n queries.data_ptr<float>(),\n features.data_ptr<float>(),\n output.data_ptr<float>(),\n radius_sq,\n B, P, Q, C\n );\n\n return output;\n}\n'''\n\n_cpp_source = r'''\ntorch::Tensor ball_max_pool_cuda(\n torch::Tensor points,\n torch::Tensor queries,\n torch::Tensor features,\n float radius_sq\n);\n'''\n\n_ball_module = load_inline(\n name='ball_max_pool',\n cpp_sources=_cpp_source,\n cuda_sources=_cuda_source,\n functions=['ball_max_pool_cuda'],\n verbose=False,\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, in_channels: int, radius: float):\n super().__init__()\n self.in_channels = in_channels\n self.radius_sq = float(radius) ** 2\n\n def forward(\n self,\n points: torch.Tensor,\n queries: torch.Tensor,\n features: torch.Tensor,\n ) -> torch.Tensor:\n return _ball_module.ball_max_pool_cuda(\n points, queries, features, self.radius_sq\n )\n```" }
result ← submit_kernel (fail)
failsubmit_kernel submit_kernel FAILED: kernel did not compile. Syntax error in custom generated code or ModelNew not found
final submitted kernel (114 lines, 2,814 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

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

__global__ void ball_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 = -1e38f;
    bool has_any = false;

    const float* q_ptr = queries + (b * Q + q) * 3;
    float qx = q_ptr[0];
    float qy = q_ptr[1];
    float qz = q_ptr[2];

    for (int p = 0; p < P; ++p) {
        const float* p_ptr = points + (b * P + p) * 3;
        float dx = qx - p_ptr[0];
        float dy = qy - p_ptr[1];
        float dz = qz - p_ptr[2];
        float dist_sq = dx*dx + dy*dy + dz*dz;
        if (dist_sq <= radius_sq) {
            float f = features[(b * P + p) * C + c];
            if (!has_any || f > max_val) {
                max_val = f;
            }
            has_any = true;
        }
    }
    output[(b * Q + q) * C + c] = has_any ? max_val : 0.0f;
}

torch::Tensor ball_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::empty({B, Q, C}, features.options());

    int total = B * Q * C;
    int threads = 256;
    int blocks = (total + threads - 1) / threads;

    ball_max_pool_kernel<<<blocks, threads>>>(
        points.data_ptr<float>(),
        queries.data_ptr<float>(),
        features.data_ptr<float>(),
        output.data_ptr<float>(),
        radius_sq,
        B, P, Q, C
    );

    return output;
}
'''

_cpp_source = r'''
torch::Tensor ball_max_pool_cuda(
    torch::Tensor points,
    torch::Tensor queries,
    torch::Tensor features,
    float radius_sq
);
'''

_ball_module = load_inline(
    name='ball_max_pool',
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=['ball_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

    def forward(
        self,
        points: torch.Tensor,
        queries: torch.Tensor,
        features: torch.Tensor,
    ) -> torch.Tensor:
        return _ball_module.ball_max_pool_cuda(
            points, queries, features, self.radius_sq
        )
```