grok-4.3 · popcorn · pop_l123_all_gpt
correct 24_VoxelGridPooling.py
model
grok-4.3
turns / max
5 / 15
tool calls
5
tokens (total)
25468
wall clock
1141.5s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
0.14 μs
ref runtime
12.80 μs
speedup
92.09x
started
2026-05-11T22:12:38Z
timingkernelbench.timing
kernel runtime
0.14 μs
ref runtime
12.80 μs
speedup
92.09x
kernel mean / median / std
0.14 / — / 0.06 μs
ref mean / median / std
12.80 / — / 0.69 μs
numerical precisioneval_kernel_against_ref
max abs error
0.00
mean abs error
0.00
max rel error
0.00
mean rel error
0.00
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
39
ref launches
4484
fusion ratio (ref / kernel)
114.97x
total cuda time
836 μs
SOL (speed-of-light)nsight_hardware_counters
SOL score (0–1)
0.190
dram util %
8.1%
compute util %
19.0%
bottleneck
compute-bound
dominant pipe
FP32
arithmetic intensity
—
ridge point
20.00
roofline / occupancynsight_hardware_counters
dram bandwidth
269.64 GB/s
dram util %
8.0%
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
182.5%
bottleneck
compute-bound
dominant pipe %
19.0%
arithmetic intensity
—
ridge point
20.00
L1 hit %
—
L2 hit %
1228.1%
ld sectors / req
97.86
st sectors / req
—
regs / thread
336
smem / block
19968 B
block size
2496
peak bw
3350.00 GB/s
peak fp32 tflops
67.00
peak fp16 tflops
1979.00
warp stalls (top 5)nsight
long_scoreboard
16127.0%
short_scoreboard
4650.1%
mio_throttle
2773.6%
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 0238.3s turn · 11.8s LLM · 226.5s tools · tokens in=2,338 · out=578 · total=3,488
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 3-D molecular point clouds. The task involves replacing the forward computation of a PyTorch model that performs voxel-grid pooling on 3D molecular point clouds with a custom CUDA kernel.
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, // (B*N, D)\n const long* __restrict__ indices, // (B*N, 3)\n const long* __restrict__ batch_idx, // (B*N)\n float* __restrict__ voxel_grid, // (B, D, G, G, G)\n int B, int N, int D, 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 = batch_idx[idx];\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 for (int d = 0; d < D; d++) {\n float val = features[idx * D + d];\n atomicAdd(&voxel_grid[(b * D + d) * G * G * G + (ix * G + iy) * G + iz], val);\n }\n}\n\ntorch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, torch::Tensor batch_idx, int G) {\n int B = features.size(0) / indices.size(0) * indices.size(0); wait, better way.\n auto sizes = features.sizes();\n int BN = features.size(0);\n int D = features.size(1);\n int B_ = batch_idx.size(0) / indices.size(0) wait no.\n // Assume features is (B*N, D), indices (B*N, 3), batch_idx (B*N)\n int B = batch_idx.max().item<int>() + 1; // or pass B\n // Better to pass B\n // Let's redesign.\n return torch::zeros({1}, features.options());\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor voxel_pool(torch::Tensor features, torch::Tensor indices, torch::Tensor batch_idx, int G, int B);\n\"\"\"\n\nmodule = load_inline(\n name='voxel_pool',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['voxel_pool'],\n verbose=True,\n extra_cflags=['-O3']\n)\nprint(\"Compiled successfully\")\nprint(module)" }
result ← compile_kernel (fail)
failcompile_kernel compile_kernel FAILED: RuntimeError. RuntimeError: Error building extension 'voxel_pool'
turn 1303.5s turn · 5.4s LLM · 298.1s tools · tokens in=3,513 · out=555 · total=4,119
reasoning
The compile failed because my CUDA code was incomplete and had errors. I need to fix it properly.
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 = r'''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n\n__global__ void voxel_pool_kernel(\n const float* __restrict__ features, // (BN, D)\n const long* __restrict__ indices, // (BN, 3)\n const long* __restrict__ batch_idx, // (BN)\n float* __restrict__ voxel_grid, // (B, D, G, G, G)\n int BN, int D, int G\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n if (idx >= BN) return;\n\n long b = batch_idx[idx];\n long ix = indices[idx * 3];\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 for (int d = 0; d < D; ++d) {\n float val = features[idx * D + d];\n int offset = ((b * D + d) * G + ix) * G * G + iy * G + iz;\n atomicAdd(&voxel_grid[offset], val);\n }\n}\n\ntorch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, torch::Tensor batch_idx, int G, int B) {\n int BN = features.size(0);\n int D = features.size(1);\n auto options = torch::TensorOptions().dtype(torch::kFloat32).device(features.device());\n auto voxel_grid = torch::zeros({B, D, G, G, G}, options);\n\n int threads = 256;\n int blocks = (BN + threads - 1) / threads;\n\n voxel_pool_kernel<<<blocks, threads>>>(\n features.data_ptr<float>(),\n indices.data_ptr<long>(),\n batch_idx.data_ptr<long>(),\n voxel_grid.data_ptr<float>(),\n BN, D, G\n );\n\n return voxel_grid;\n}\n'''\n\ncpp_source = r'''\ntorch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, torch::Tensor batch_idx, int G, int B);\n'''\n\nmodule = load_inline(\n name='voxel_pool',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['voxel_pool_cuda'],\n verbose=True,\n extra_cflags=['-O3']\n)\nprint(\"Compiled successfully\")" }
result ← compile_kernel (fail)
failcompile_kernel compile_kernel FAILED: ModelNew class not found or syntax error prevented execution. Compiled successfully
turn 2162.6s turn · 9.1s LLM · 153.5s tools · tokens in=4,144 · out=844 · total=5,058
reasoning
It compiled, but the output says "ModelNew class not found", because I only have the compilation snippet, not the full ModelNew.
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 = r'''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n\n__global__ void voxel_pool_kernel(\n const float* __restrict__ features, // (BN, D)\n const long* __restrict__ indices, // (BN, 3)\n const long* __restrict__ batch_idx, // (BN)\n float* __restrict__ voxel_grid, // (B, D, G, G, G)\n int BN, int D, int G\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n if (idx >= BN) return;\n\n long b = batch_idx[idx];\n long ix = indices[idx * 3];\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 for (int d = 0; d < D; ++d) {\n float val = features[idx * D + d];\n int offset = ((b * D + d) * G + ix) * G * G + iy * G + iz;\n atomicAdd(&voxel_grid[offset], val);\n }\n}\n\ntorch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, torch::Tensor batch_idx, int G, int B) {\n int BN = features.size(0);\n int D = features.size(1);\n auto options = torch::TensorOptions().dtype(torch::kFloat32).device(features.device());\n auto voxel_grid = torch::zeros({B, D, G, G, G}, options);\n\n int threads = 256;\n int blocks = (BN + threads - 1) / threads;\n\n voxel_pool_kernel<<<blocks, threads>>>(\n features.data_ptr<float>(),\n indices.data_ptr<long>(),\n batch_idx.data_ptr<long>(),\n voxel_grid.data_ptr<float>(),\n BN, D, G\n );\n\n return voxel_grid;\n}\n'''\n\ncpp_source = r'''\ntorch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, torch::Tensor batch_idx, int G, int B);\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_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) # (B, N, D)\n\n # Center coordinates and discretize\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 # Prepare for kernel: flatten to (B*N, ...)\n BN = B * N\n features_flat = features.reshape(BN, -1).contiguous()\n indices_flat = indices.reshape(BN, 3).contiguous()\n batch_idx = torch.arange(B, device=coords.device).unsqueeze(1).expand(B, N).reshape(BN).contiguous()\n\n return voxel_pool.voxel_pool_cuda(features_flat, indices_flat, batch_idx, G, B)\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 3105.4s turn · 6.4s LLM · 99.0s tools · tokens in=5,075 · out=845 · total=5,927
reasoning
Compiled successfully. Now check 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 = r'''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n\n__global__ void voxel_pool_kernel(\n const float* __restrict__ features, // (BN, D)\n const long* __restrict__ indices, // (BN, 3)\n const long* __restrict__ batch_idx, // (BN)\n float* __restrict__ voxel_grid, // (B, D, G, G, G)\n int BN, int D, int G\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n if (idx >= BN) return;\n\n long b = batch_idx[idx];\n long ix = indices[idx * 3];\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 for (int d = 0; d < D; ++d) {\n float val = features[idx * D + d];\n int offset = ((b * D + d) * G + ix) * G * G + iy * G + iz;\n atomicAdd(&voxel_grid[offset], val);\n }\n}\n\ntorch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, torch::Tensor batch_idx, int G, int B) {\n int BN = features.size(0);\n int D = features.size(1);\n auto options = torch::TensorOptions().dtype(torch::kFloat32).device(features.device());\n auto voxel_grid = torch::zeros({B, D, G, G, G}, options);\n\n int threads = 256;\n int blocks = (BN + threads - 1) / threads;\n\n voxel_pool_kernel<<<blocks, threads>>>(\n features.data_ptr<float>(),\n indices.data_ptr<long>(),\n batch_idx.data_ptr<long>(),\n voxel_grid.data_ptr<float>(),\n BN, D, G\n );\n\n return voxel_grid;\n}\n'''\n\ncpp_source = r'''\ntorch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, torch::Tensor batch_idx, int G, int B);\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_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) # (B, N, D)\n\n # Center coordinates and discretize\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 # Prepare for kernel: flatten to (B*N, ...)\n BN = B * N\n features_flat = features.reshape(BN, -1).contiguous()\n indices_flat = indices.reshape(BN, 3).contiguous()\n batch_idx = torch.arange(B, device=coords.device).unsqueeze(1).expand(B, N).reshape(BN).contiguous()\n\n return voxel_pool.voxel_pool_cuda(features_flat, indices_flat, batch_idx, G, B)\n" }
result ← run_correctness (ok)
okrun_correctness run_correctness PASSED: (3 / 3) trials all matched the reference. Numerical precision: max_abs_err=0.00e+00 mean_abs_err=0.00e+00 max_rel_err=0.00e+00
turn 4331.6s turn · 6.8s LLM · 324.9s tools · tokens in=5,987 · out=844 · total=6,876 · FINAL
reasoning
Correctness passed perfectly. Now, to submit it, but I should check if it's faster, though since the task requires faster than reference, and the reference has a Python loop, this should be.
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 = r'''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n\n__global__ void voxel_pool_kernel(\n const float* __restrict__ features, // (BN, D)\n const long* __restrict__ indices, // (BN, 3)\n const long* __restrict__ batch_idx, // (BN)\n float* __restrict__ voxel_grid, // (B, D, G, G, G)\n int BN, int D, int G\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n if (idx >= BN) return;\n\n long b = batch_idx[idx];\n long ix = indices[idx * 3];\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 for (int d = 0; d < D; ++d) {\n float val = features[idx * D + d];\n int offset = ((b * D + d) * G + ix) * G * G + iy * G + iz;\n atomicAdd(&voxel_grid[offset], val);\n }\n}\n\ntorch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, torch::Tensor batch_idx, int G, int B) {\n int BN = features.size(0);\n int D = features.size(1);\n auto options = torch::TensorOptions().dtype(torch::kFloat32).device(features.device());\n auto voxel_grid = torch::zeros({B, D, G, G, G}, options);\n\n int threads = 256;\n int blocks = (BN + threads - 1) / threads;\n\n voxel_pool_kernel<<<blocks, threads>>>(\n features.data_ptr<float>(),\n indices.data_ptr<long>(),\n batch_idx.data_ptr<long>(),\n voxel_grid.data_ptr<float>(),\n BN, D, G\n );\n\n return voxel_grid;\n}\n'''\n\ncpp_source = r'''\ntorch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, torch::Tensor batch_idx, int G, int B);\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_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) # (B, N, D)\n\n # Center coordinates and discretize\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 # Prepare for kernel: flatten to (B*N, ...)\n BN = B * N\n features_flat = features.reshape(BN, -1).contiguous()\n indices_flat = indices.reshape(BN, 3).contiguous()\n batch_idx = torch.arange(B, device=coords.device).unsqueeze(1).expand(B, N).reshape(BN).contiguous()\n\n return voxel_pool.voxel_pool_cuda(features_flat, indices_flat, batch_idx, G, B)\n" }
result ← submit_kernel (ok)
oksubmit_kernel submit_kernel PASSED: (3 / 3) correctness trials all passed. Kernel runtime: 0.14 μs Runtime stats: mean=0.14μs median=?μs std=0.06μs Numerical precision: max_abs_err=0.00e+00 mean_abs_err=0.00e+00 max_rel_err=0.00e+00 Memory: 49.31 MB peak (ref: 49.85 MB, ratio: 0.9892x) Kernel launches: 39 (ref: 4484, fusion ratio: 114.9744) SOL score: 0.19 Flagged for excessive speedup — this submission may have been rejected by automated review.
final submitted kernel (93 lines, 2,996 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_runtime.h>

__global__ void voxel_pool_kernel(
    const float* __restrict__ features,   // (BN, D)
    const long* __restrict__ indices,     // (BN, 3)
    const long* __restrict__ batch_idx,   // (BN)
    float* __restrict__ voxel_grid,       // (B, D, G, G, G)
    int BN, int D, int G
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx >= BN) return;

    long b = batch_idx[idx];
    long ix = indices[idx * 3];
    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;

    for (int d = 0; d < D; ++d) {
        float val = features[idx * D + d];
        int offset = ((b * D + d) * G + ix) * G * G + iy * G + iz;
        atomicAdd(&voxel_grid[offset], val);
    }
}

torch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, torch::Tensor batch_idx, int G, int B) {
    int BN = features.size(0);
    int D = features.size(1);
    auto options = torch::TensorOptions().dtype(torch::kFloat32).device(features.device());
    auto voxel_grid = torch::zeros({B, D, G, G, G}, options);

    int threads = 256;
    int blocks = (BN + threads - 1) / threads;

    voxel_pool_kernel<<<blocks, threads>>>(
        features.data_ptr<float>(),
        indices.data_ptr<long>(),
        batch_idx.data_ptr<long>(),
        voxel_grid.data_ptr<float>(),
        BN, D, G
    );

    return voxel_grid;
}
'''

cpp_source = r'''
torch::Tensor voxel_pool_cuda(torch::Tensor features, torch::Tensor indices, torch::Tensor batch_idx, int G, int B);
'''

voxel_pool = load_inline(
    name='voxel_pool',
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=['voxel_pool_cuda'],
    verbose=False,
    extra_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)  # (B, N, D)

        # 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)

        # Prepare for kernel: flatten to (B*N, ...)
        BN = B * N
        features_flat = features.reshape(BN, -1).contiguous()
        indices_flat = indices.reshape(BN, 3).contiguous()
        batch_idx = torch.arange(B, device=coords.device).unsqueeze(1).expand(B, N).reshape(BN).contiguous()

        return voxel_pool.voxel_pool_cuda(features_flat, indices_flat, batch_idx, G, B)