compile_fail
24_VoxelGridPooling.py
started
2026-05-11T14:27:46Z
compile error: Syntax error in custom generated code or ModelNew not found
timingkernelbench.timing
kernel mean / median / std
—
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 018.8s turn · 18.8s LLM · 0.0s tools · tokens in=1,807 · out=708 · total=3,012
reasoning
The task is to replace the forward computation of this PyTorch model with a custom CUDA kernel. The model does voxel-grid pooling for 3D molecular point clouds.
The task involves replacing the forward pass of a PyTorch model that performs voxel-grid pooling on 3D molecular point clouds with a custom CUDA kernel. The new ModelNew class must take the same inputs as the original Model, have the same forward signature, and produce identical numerical outputs.
assistant```python
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline
cuda_source = '''
#include <torch/extension.h>
#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 F, int G
) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int total = B * N;
if (idx >= total) return;
int b = idx / N;
int n = idx % N;
long ix = indices[idx * 3 + 0];
long iy = indices[idx * 3 + 1];
long iz = indices[idx * 3 + 2];
if (ix < 0 || ix >= G || iy < 0 || iy >= G || iz < 0 || iz >= G) return;
const float* feat_ptr = features + (b * N + n) * F;
for (int d = 0; d < F; ++d) {
float val = feat_ptr[d];
atomicAdd(&voxel_grid[((b * F + d) * G + ix) * G * G + iy * G + iz], val);
}
}
torch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, int G) {
int B = features.size(0);
int N = features.size(1);
int F = features.size(2);
auto voxel_grid = torch::zeros({B, F, G, G, G}, features.options());
int total = B * N;
int threads = 256;
int blocks = (total + threads - 1) / threads;
voxel_pool_kernel<<<blocks, threads>>>(
features.data_ptr<float>(),
indices.data_ptr<long>(),
voxel_grid.data_ptr<float>(),
B, N, F, G
);
return voxel_grid;
}
'''
cpp_source = '''
torch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, int G);
'''
voxel_pool = load_inline(
name='voxel_pool',
cpp_sources=[cpp_source],
cuda_sources=[cuda_source],
functions=['voxel_pool_cuda'],
verbose=False,
extra_cuda_cflags=['-O3']
)
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)
def forward(self, coords: torch.Tensor, features: torch.Tensor) -> torch.Tensor:
B, N, _ = coords.shape
G = self.grid_size
features = self.linear(features)
center = coords.mean(dim=1, keepdim=True)
shifted = (coords - center) / self.voxel_size + G / 2.0
indices = shifted.long().clamp(0, G - 1)
indices_flat = indices.reshape(B * N, 3).contiguous()
return voxel_pool.voxel_pool_cuda(features, indices_flat, G)
```
turn 1354.9s turn · 6.0s LLM · 348.9s tools · tokens in=3,080 · out=715 · total=3,894
reasoning
The instruction says: "You must call at least one tool this turn: use `compile_kernel` / `run_correctness` to iterate, or call `submit_kernel` with your best `ModelNew` implementation to finalize the ...
call → compile_kernel
callcompile_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = '''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.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 F, int G\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * N;\n if (idx >= total) return;\n\n int b = idx / N;\n int n = idx % N;\n\n long ix = indices[idx * 3 + 0];\n long iy = indices[idx * 3 + 1];\n long iz = indices[idx * 3 + 2];\n\n if (ix < 0 || ix >= G || iy < 0 || iy >= G || iz < 0 || iz >= G) return;\n\n const float* feat_ptr = features + (b * N + n) * F;\n for (int d = 0; d < F; ++d) {\n float val = feat_ptr[d];\n atomicAdd(&voxel_grid[((b * F + d) * G + ix) * G * G + iy * G + iz], val);\n }\n}\n\ntorch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, int G) {\n int B = features.size(0);\n int N = features.size(1);\n int F = features.size(2);\n\n auto voxel_grid = torch::zeros({B, F, G, G, G}, features.options());\n\n int total = B * N;\n int threads = 256;\n int blocks = (total + 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, F, G\n );\n\n return voxel_grid;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, int G);\n'''\n\nvoxel_pool = load_inline(\n name='voxel_pool',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['voxel_pool_cuda'],\n verbose=False,\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\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(self, coords: torch.Tensor, features: torch.Tensor) -> torch.Tensor:\n B, N, _ = coords.shape\n G = self.grid_size\n\n features = self.linear(features)\n\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)\n\n indices_flat = indices.reshape(B * N, 3).contiguous()\n return voxel_pool.voxel_pool_cuda(features, indices_flat, G)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 2497.8s turn · 6.0s LLM · 491.7s tools · tokens in=3,911 · out=716 · total=4,640
reasoning
Compilation succeeded. Now I should check correctness with run_correctness.
call → run_correctness
callrun_correctness
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = '''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.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 F, int G\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * N;\n if (idx >= total) return;\n\n int b = idx / N;\n int n = idx % N;\n\n long ix = indices[idx * 3 + 0];\n long iy = indices[idx * 3 + 1];\n long iz = indices[idx * 3 + 2];\n\n if (ix < 0 || ix >= G || iy < 0 || iy >= G || iz < 0 || iz >= G) return;\n\n const float* feat_ptr = features + (b * N + n) * F;\n for (int d = 0; d < F; ++d) {\n float val = feat_ptr[d];\n atomicAdd(&voxel_grid[((b * F + d) * G + ix) * G * G + iy * G + iz], val);\n }\n}\n\ntorch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, int G) {\n int B = features.size(0);\n int N = features.size(1);\n int F = features.size(2);\n\n auto voxel_grid = torch::zeros({B, F, G, G, G}, features.options());\n\n int total = B * N;\n int threads = 256;\n int blocks = (total + 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, F, G\n );\n\n return voxel_grid;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, int G);\n'''\n\nvoxel_pool = load_inline(\n name='voxel_pool',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['voxel_pool_cuda'],\n verbose=False,\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\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(self, coords: torch.Tensor, features: torch.Tensor) -> torch.Tensor:\n B, N, _ = coords.shape\n G = self.grid_size\n\n features = self.linear(features)\n\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)\n\n indices_flat = indices.reshape(B * N, 3).contiguous()\n return voxel_pool.voxel_pool_cuda(features, indices_flat, G)\n"
}
result ← run_correctness (ok)
okrun_correctness
run_correctness PASSED: (5 / 5) trials all matched the reference.
Numerical precision: max_abs_err=2.38e-07 mean_abs_err=4.05e-14 max_rel_err=3.78e-07
turn 381.1s turn · 5.5s LLM · 75.6s tools · tokens in=4,700 · out=716 · total=5,443 · FINAL
reasoning
Correctness passed perfectly. Now I can submit it, as it's correct and should be faster due to the custom kernel avoiding Python loops.
call → submit_kernel
callsubmit_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = '''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.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 F, int G\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * N;\n if (idx >= total) return;\n\n int b = idx / N;\n int n = idx % N;\n\n long ix = indices[idx * 3 + 0];\n long iy = indices[idx * 3 + 1];\n long iz = indices[idx * 3 + 2];\n\n if (ix < 0 || ix >= G || iy < 0 || iy >= G || iz < 0 || iz >= G) return;\n\n const float* feat_ptr = features + (b * N + n) * F;\n for (int d = 0; d < F; ++d) {\n float val = feat_ptr[d];\n atomicAdd(&voxel_grid[((b * F + d) * G + ix) * G * G + iy * G + iz], val);\n }\n}\n\ntorch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, int G) {\n int B = features.size(0);\n int N = features.size(1);\n int F = features.size(2);\n\n auto voxel_grid = torch::zeros({B, F, G, G, G}, features.options());\n\n int total = B * N;\n int threads = 256;\n int blocks = (total + 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, F, G\n );\n\n return voxel_grid;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, int G);\n'''\n\nvoxel_pool = load_inline(\n name='voxel_pool',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['voxel_pool_cuda'],\n verbose=False,\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\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(self, coords: torch.Tensor, features: torch.Tensor) -> torch.Tensor:\n B, N, _ = coords.shape\n G = self.grid_size\n\n features = self.linear(features)\n\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)\n\n indices_flat = indices.reshape(B * N, 3).contiguous()\n return voxel_pool.voxel_pool_cuda(features, indices_flat, G)\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 (91 lines, 2,536 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline
cuda_source = '''
#include <torch/extension.h>
#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 F, int G
) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int total = B * N;
if (idx >= total) return;
int b = idx / N;
int n = idx % N;
long ix = indices[idx * 3 + 0];
long iy = indices[idx * 3 + 1];
long iz = indices[idx * 3 + 2];
if (ix < 0 || ix >= G || iy < 0 || iy >= G || iz < 0 || iz >= G) return;
const float* feat_ptr = features + (b * N + n) * F;
for (int d = 0; d < F; ++d) {
float val = feat_ptr[d];
atomicAdd(&voxel_grid[((b * F + d) * G + ix) * G * G + iy * G + iz], val);
}
}
torch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, int G) {
int B = features.size(0);
int N = features.size(1);
int F = features.size(2);
auto voxel_grid = torch::zeros({B, F, G, G, G}, features.options());
int total = B * N;
int threads = 256;
int blocks = (total + threads - 1) / threads;
voxel_pool_kernel<<<blocks, threads>>>(
features.data_ptr<float>(),
indices.data_ptr<long>(),
voxel_grid.data_ptr<float>(),
B, N, F, G
);
return voxel_grid;
}
'''
cpp_source = '''
torch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, int G);
'''
voxel_pool = load_inline(
name='voxel_pool',
cpp_sources=[cpp_source],
cuda_sources=[cuda_source],
functions=['voxel_pool_cuda'],
verbose=False,
extra_cuda_cflags=['-O3']
)
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)
def forward(self, coords: torch.Tensor, features: torch.Tensor) -> torch.Tensor:
B, N, _ = coords.shape
G = self.grid_size
features = self.linear(features)
center = coords.mean(dim=1, keepdim=True)
shifted = (coords - center) / self.voxel_size + G / 2.0
indices = shifted.long().clamp(0, G - 1)
indices_flat = indices.reshape(B * N, 3).contiguous()
return voxel_pool.voxel_pool_cuda(features, indices_flat, G)
```