FW-GLM-5-1 · popcorn · pop_l123_st_gpt
correct 24_VoxelGridPooling.py
model
FW-GLM-5-1
turns / max
1 / 1
tool calls
1
tokens (total)
3379
wall clock
1471.1s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
0.05 μs
ref runtime
9.71 μs
speedup
188.91x
started
2026-05-11T09:25:17Z
timingkernelbench.timing
kernel runtime
0.05 μs
ref runtime
9.71 μs
speedup
188.91x
kernel mean / median / std
0.05 / — / 6.04e-04 μs
ref mean / median / std
9.71 / — / 0.02 μs
numerical precisioneval_kernel_against_ref
max abs error
2.38e-07
mean abs error
2.34e-14
max rel error
2.10e-07
mean rel error
2.97e-14
gpu memoryextended_metrics.measure_memory
kernel peak
49.31 MB
ref peak
49.85 MB
kernel / ref
0.99x
kernel launches / fusiontorch.profiler
kernel launches
21
ref launches
4484
fusion ratio (ref / kernel)
213.52x
total cuda time
265 μs
energyNVML (pynvml)
kernel mJ / run
458.28
ref mJ / run
5724.58
ratio (ref / kernel)
12.49x
avg power
3137.07 W
ref avg power
543.90 W
SOL (speed-of-light)heuristic_from_speedup
SOL score (0–1)
1.000
dram util %
—
compute util %
—
bottleneck
unknown
dominant pipe
—
arithmetic intensity
—
ridge point
—
roofline / occupancyheuristic
dram bandwidth
—
dram util %
—
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
—
bottleneck
—
dominant pipe %
—
arithmetic intensity
—
ridge point
—
L1 hit %
—
L2 hit %
—
ld sectors / req
—
st sectors / req
—
regs / thread
—
smem / block
—
block size
—
peak bw
—
peak fp32 tflops
—
peak fp16 tflops
—
flagssubmit_kernel guard
excessive speedup
flagged
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn


class Model(nn.Module):
    """
    Voxel-grid pooling for 3-D molecular point clouds.  Discretizes
    continuous 3-D coordinates into a regular grid and pools per-atom
    features into each voxel.  Used in 3-D-CNN approaches for molecular
    property prediction and ligand docking (e.g., PointVoxel-CNN,
    3D-Infomax).
    """

    def __init__(self, feat_dim, grid_size=32, voxel_size=1.0):
        super().__init__()
        self.feat_dim = feat_dim
        self.grid_size = grid_size
        self.voxel_size = voxel_size
        self.linear = nn.Linear(feat_dim, feat_dim)

    def forward(
        self, coords: torch.Tensor, features: torch.Tensor
    ) -> torch.Tensor:
        """
        Args:
            coords:   (B, N, 3) – atom coordinates
            features: (B, N, feat_dim) – per-atom features
        Returns:
            voxel_grid: (B, feat_dim, G, G, G) – pooled voxel features
        """
        B, N, _ = coords.shape
        G = self.grid_size

        features = self.linear(features)

        # Center coordinates and discretize
        center = coords.mean(dim=1, keepdim=True)
        shifted = (coords - center) / self.voxel_size + G / 2.0
        indices = shifted.long().clamp(0, G - 1)  # (B, N, 3)

        voxel_grid = coords.new_zeros(B, self.feat_dim, G, G, G)
        ix = indices[:, :, 0]
        iy = indices[:, :, 1]
        iz = indices[:, :, 2]

        batch_idx = torch.arange(B, device=coords.device).unsqueeze(1).expand(-1, N)
        for d in range(self.feat_dim):
            feat_d = features[:, :, d]
            voxel_grid[:, d].index_put_(
                (batch_idx.reshape(-1), ix.reshape(-1), iy.reshape(-1), iz.reshape(-1)),
                feat_d.reshape(-1),
                accumulate=True,
            )

        return voxel_grid


feat_dim = 32
grid_size = 32
voxel_size = 1.0
num_atoms = 256
batch_size = 4


def get_inputs():
    p = popcorn_pri
    bs = p.jitter_int(batch_size)
    na = p.jitter_int(num_atoms)
    coords = torch.randn(bs, na, 3) * 5.0
    features = torch.randn(bs, na, feat_dim)
    return [coords, features]


def get_init_inputs():
    return [feat_dim, grid_size, voxel_size]
turn 0 · FINAL · 1 tool call(s) · 1471.1s turn · 20.8s LLM · tokens in=1,391 · out=1,988 · total=3,379 · 9,177 chars (click to expand)
turn 01471.1s turn · 20.8s LLM · 1450.3s tools · tokens in=1,391 · out=1,988 · total=3,379 · FINAL
reasoning
Let me analyze this task. I need to create a custom CUDA kernel that replaces the PyTorch model for voxel-grid pooling. The reference implementation: 1. Takes coordinates (B, N, 3) and features (B, N, feat_dim) 2. Applies a linear transformation to features 3. Centers coordinates and discretizes them into a grid 4. Accumulates features into voxel grid (B, feat_dim, G, G, G) The key operation is the scatter-accumulate into a voxel grid. This is the bottleneck that needs to be optimized with a CUDA kernel. Let me design the kernel: 1. The linear layer can remain in PyTorch 2. The centering and discretization can be done in PyTorch or CUDA 3. The main optimization target is the scatter-accumulate operation For the CUDA kernel: - Each thread can handle one atom's contribution - We need to handle atomic accumulation since multiple atoms can map to the same voxel - Grid size is 32x32x32, which is small enough that we can have good caching Let me write the implementation: ```python import torch import torch.nn as nn from torch.utils.cpp_extension import load_inline # Define the CUDA kernel cuda_source = """ #include <cuda.h> #include <cuda_runtime.h> __global__ void voxel_pool_kernel( const float* __restrict__ features, const long* __restrict__ indices, float* __restrict__ voxel_grid, int B, int N, int feat_dim, int G ) { // Each thread handles one (batch, atom, feature_dim) element int idx = blockIdx.x * blockDim.x + threadIdx.x; int total_elements = B * N * feat_dim; if (idx < total_elements) { int d = idx % feat_dim; int atom_idx = (idx / feat_dim) % N; int b = idx / (feat_dim * N); int ix = indices[(b * N + atom_idx) * 3 + 0]; int iy = indices[(b * N + atom_idx) * 3 + 1]; int iz = indices[(b * N + atom_idx) * 3 + 2]; float val = features[idx]; // Atomic add to voxel grid int voxel_idx = ((b * feat_dim + d) * G + ix) * G * G + iy * G + iz; atomicAdd(&voxel_grid[voxel_idx], val); } } torch::Tensor voxel_pool_cuda( torch::Tensor features, torch::Tensor indices, int grid_size ) { int B = features.size(0); int N = features.size(1); int feat_dim = features.size(2); int G = grid_size; auto voxel_grid = torch::zeros({B, feat_dim, G, G, G}, features.options()); int total_elements = B * N * feat_dim; const int threads = 256; const int blocks = (total_elements + threads - 1) / threads; voxel_pool_kernel<<<blocks, threads>>>( features.data_ptr<float>(), indices.data_ptr<long>(), voxel_grid.data_ptr<float>(), B, N, feat_dim, G ); return voxel_grid; } """ cpp_source = """ torch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, int grid_size); """ # Load the extension voxel_ops = load_inline( name='voxel_ops', cpp_sources=cpp_source, cuda_sources=cuda_source, functions=['voxel_pool_cuda'], extra_cuda_cflags=['-O3'], verbose=True ) class ModelNew(nn.Module): def __init__(self, feat_dim, grid_size=32, voxel_size=1.0): super().__init__() self.feat_dim = feat_dim self.grid_size = grid_size self.voxel_size = voxel_size self.linear = nn.Linear(feat_dim, feat_dim) ``` I need to reconsider the voxel indexing formula. The grid dimensions are B × feat_dim × G × G × G, where the spatial indices (ix, iy, iz) map to positions [2], [3], [4] in the tensor. The linearized index should be: b * feat_dim * G * G * G + d * G * G * G + ix * G * G + iy * G + iz. This works because PyTorch flattens dimensions in row-major order, so the first dimension (ix) strides by G², followed by iy striding by G, and iz as the final offset. Checking with explicit indices: voxel_grid[b, d, ix, iy, iz] unpacks to b*feat_dim*G³ + d*G³ + ix*G² + iy*G + iz, which matches my formula. Now I can write out the complete kernel implementation with the corrected indexing.
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 kernel for voxel pooling\ncuda_source = r\"\"\"\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <ATen/ATen.h>\n\n__global__ void voxel_pool_kernel(\n const float* __restrict__ features,\n const long* __restrict__ indices,\n float* __restrict__ voxel_grid,\n int B, int N, int feat_dim, int G\n) {\n // Each thread handles one (batch, atom, feature_dim) element\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total_elements = B * N * feat_dim;\n \n if (idx < total_elements) {\n int d = idx % feat_dim;\n int atom_idx = (idx / feat_dim) % N;\n int b = idx / (feat_dim * N);\n \n int ix = indices[(b * N + atom_idx) * 3 + 0];\n int iy = indices[(b * N + atom_idx) * 3 + 1];\n int iz = indices[(b * N + atom_idx) * 3 + 2];\n \n float val = features[idx];\n \n // Linear index into voxel_grid[b, d, ix, iy, iz]\n // Shape is [B, feat_dim, G, G, G]\n long voxel_idx = ((((long)b * feat_dim + d) * G + ix) * G + iy) * G + iz;\n atomicAdd(&voxel_grid[voxel_idx], val);\n }\n}\n\ntorch::Tensor voxel_pool_cuda(\n torch::Tensor features,\n torch::Tensor indices,\n int grid_size\n) {\n int B = features.size(0);\n int N = features.size(1);\n int feat_dim = features.size(2);\n int G = grid_size;\n \n auto voxel_grid = torch::zeros({B, feat_dim, G, G, G}, features.options());\n \n int total_elements = B * N * feat_dim;\n const int threads = 256;\n const int blocks = (total_elements + threads - 1) / threads;\n \n voxel_pool_kernel<<<blocks, threads>>>(\n features.data_ptr<float>(),\n indices.data_ptr<long>(),\n voxel_grid.data_ptr<float>(),\n B, N, feat_dim, G\n );\n \n return voxel_grid;\n}\n\"\"\"\n\ncpp_source = r\"\"\"\ntorch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, int grid_size);\n\"\"\"\n\n# Load the CUDA extension\nvoxel_ops = load_inline(\n name='voxel_ops_v2',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['voxel_pool_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n \"\"\"\n Voxel-grid pooling for 3-D molecular point clouds. Uses custom CUDA kernel\n for the scatter-accumulate operation.\n \"\"\"\n def __init__(self, feat_dim, grid_size=32, voxel_size=1.0):\n super().__init__()\n self.feat_dim = feat_dim\n self.grid_size = grid_size\n self.voxel_size = voxel_size\n self.linear = nn.Linear(feat_dim, feat_dim)\n\n def forward(\n self, coords: torch.Tensor, features: torch.Tensor\n ) -> torch.Tensor:\n \"\"\"\n Args:\n coords: (B, N, 3) - atom coordinates\n features: (B, N, feat_dim) - per-atom features\n Returns:\n voxel_grid: (B, feat_dim, G, G, G) - pooled voxel features\n \"\"\"\n B, N, _ = coords.shape\n G = self.grid_size\n\n # Linear transformation (keep in PyTorch)\n features = self.linear(features)\n\n # Center coordinates and discretize (keep in PyTorch)\n center = coords.mean(dim=1, keepdim=True)\n shifted = (coords - center) / self.voxel_size + G / 2.0\n indices = shifted.long().clamp(0, G - 1) # (B, N, 3)\n\n # Use custom CUDA kernel for scatter-accumulate\n voxel_grid = voxel_ops.voxel_pool_cuda(features, indices, G)\n\n return voxel_grid" }
result ← submit_kernel (ok)
oksubmit_kernel submit_kernel PASSED: (5 / 5) correctness trials all passed. Kernel runtime: 0.05 μs Runtime stats: mean=0.05μs median=?μs std=0.00μs Numerical precision: max_abs_err=2.38e-07 mean_abs_err=2.34e-14 max_rel_err=2.10e-07 Memory: 49.31 MB peak (ref: 49.85 MB, ratio: 0.9891x) Kernel launches: 21 (ref: 4484, fusion ratio: 213.5238) Energy: 458.28 mJ/run (ref: 5724.58 mJ, ratio: 12.4914x) SOL score: 1.0 Flagged for excessive speedup — this submission may have been rejected by automated review.
final submitted kernel (116 lines, 3,480 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

# CUDA kernel for voxel pooling
cuda_source = r"""
#include <cuda.h>
#include <cuda_runtime.h>
#include <ATen/ATen.h>

__global__ void voxel_pool_kernel(
    const float* __restrict__ features,
    const long* __restrict__ indices,
    float* __restrict__ voxel_grid,
    int B, int N, int feat_dim, int G
) {
    // Each thread handles one (batch, atom, feature_dim) element
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int total_elements = B * N * feat_dim;
    
    if (idx < total_elements) {
        int d = idx % feat_dim;
        int atom_idx = (idx / feat_dim) % N;
        int b = idx / (feat_dim * N);
        
        int ix = indices[(b * N + atom_idx) * 3 + 0];
        int iy = indices[(b * N + atom_idx) * 3 + 1];
        int iz = indices[(b * N + atom_idx) * 3 + 2];
        
        float val = features[idx];
        
        // Linear index into voxel_grid[b, d, ix, iy, iz]
        // Shape is [B, feat_dim, G, G, G]
        long voxel_idx = ((((long)b * feat_dim + d) * G + ix) * G + iy) * G + iz;
        atomicAdd(&voxel_grid[voxel_idx], val);
    }
}

torch::Tensor voxel_pool_cuda(
    torch::Tensor features,
    torch::Tensor indices,
    int grid_size
) {
    int B = features.size(0);
    int N = features.size(1);
    int feat_dim = features.size(2);
    int G = grid_size;
    
    auto voxel_grid = torch::zeros({B, feat_dim, G, G, G}, features.options());
    
    int total_elements = B * N * feat_dim;
    const int threads = 256;
    const int blocks = (total_elements + threads - 1) / threads;
    
    voxel_pool_kernel<<<blocks, threads>>>(
        features.data_ptr<float>(),
        indices.data_ptr<long>(),
        voxel_grid.data_ptr<float>(),
        B, N, feat_dim, G
    );
    
    return voxel_grid;
}
"""

cpp_source = r"""
torch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, int grid_size);
"""

# Load the CUDA extension
voxel_ops = load_inline(
    name='voxel_ops_v2',
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=['voxel_pool_cuda'],
    extra_cuda_cflags=['-O3'],
    verbose=False
)

class ModelNew(nn.Module):
    """
    Voxel-grid pooling for 3-D molecular point clouds. Uses custom CUDA kernel
    for the scatter-accumulate operation.
    """
    def __init__(self, feat_dim, grid_size=32, voxel_size=1.0):
        super().__init__()
        self.feat_dim = feat_dim
        self.grid_size = grid_size
        self.voxel_size = voxel_size
        self.linear = nn.Linear(feat_dim, feat_dim)

    def forward(
        self, coords: torch.Tensor, features: torch.Tensor
    ) -> torch.Tensor:
        """
        Args:
            coords:   (B, N, 3) - atom coordinates
            features: (B, N, feat_dim) - per-atom features
        Returns:
            voxel_grid: (B, feat_dim, G, G, G) - pooled voxel features
        """
        B, N, _ = coords.shape
        G = self.grid_size

        # Linear transformation (keep in PyTorch)
        features = self.linear(features)

        # Center coordinates and discretize (keep in PyTorch)
        center = coords.mean(dim=1, keepdim=True)
        shifted = (coords - center) / self.voxel_size + G / 2.0
        indices = shifted.long().clamp(0, G - 1)  # (B, N, 3)

        # Use custom CUDA kernel for scatter-accumulate
        voxel_grid = voxel_ops.voxel_pool_cuda(features, indices, G)

        return voxel_grid