turn 0667.0s turn · 39.4s LLM · 627.6s tools · tokens in=1,440 · out=2,204 · total=4,459 · FINAL
reasoning
The task is to replace the forward computation of the PyTorch model with a custom CUDA kernel. The model does voxel-grid pooling for 3-D molecular point clouds.
I need to replace the forward computation of the PyTorch model with a custom CUDA kernel for voxel-grid pooling on 3-D molecular point clouds. The goal is to speed it up by fusing the linear transformation and the scatter-add operation into a single CUDA kernel.
The original process involves a linear transform on features, centering coordinates coordinates, shifting and discretizing to indices, then for each feature dimension scatter-adding features into a voxel grid using index_put_ with accumulate=True. The loop over feature dimensions is inefficient, especially with feat_dim=32.
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 voxel_pool_kernel(
const float* __restrict__ coords,
const float* __restrict__ features,
const float* __restrict__ weight,
const float* __restrict__ bias,
float* __restrict__ voxel_grid,
int B, int N, int F, int G, float voxel_size, float inv_voxel_size
) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int total = B * N;
if (idx >= total) return;
int b = idx / N;
int n = idx % N;
// Compute center for batch b - need to load all coords for batch or precompute? For simplicity, recompute mean per batch (small N)
float cx = 0.0f, cy = 0.0f, cz = 0.0f;
for (int i = 0; i < N; ++i) {
int base = (b * N + i) * 3;
cx += coords[base];
cy += coords[base + 1];
cz += coords[base + 2];
}
cx /= N; cy /= N; cz /= N;
int base_c = (b * N + n) * 3;
float x = (coords[base_c] - cx) * inv_voxel_size + G * 0.5f;
float y = (coords[base_c + 1] - cy) * inv_voxel_size + G * 0.5f;
float z = (coords[base_c + 2] - cz) * inv_voxel_size + G * 0.5f;
int ix = (int)floorf(x);
int iy = (int)floorf(y);
int iz = (int)floorf(z);
if (ix < 0) ix = 0; if (ix >= G) ix = G-1;
if (iy < 0) iy = 0; if (iy >= G) iy = G-1;
if (iz < 0) iz = 0; if (iz >= G) iz = G-1;
// Linear projection for this point's features
const float* feat_ptr = features + (b * N + n) * F;
float* out_ptr_base = voxel_grid + b * F * G * G * G;
for (int d = 0; d < F; ++d) {
float val = bias[d];
const float* w_row = weight + d * F;
for (int k = 0; k < F; ++k) {
val += feat_ptr[k] * w_row[k];
}
float* voxel_addr = out_ptr_base + d * G*G*G + ix * G*G + iy * G + iz;
atomicAdd(voxel_addr, val);
}
}
torch::Tensor voxel_pool_cuda(
torch::Tensor coords,
torch::Tensor features,
torch::Tensor weight,
torch::Tensor bias,
int grid_size,
float voxel_size
) {
int B = coords.size(0);
int N = coords.size(1);
int F = features.size(2);
int G = grid_size;
float inv_voxel_size = 1.0f / voxel_size;
auto voxel_grid = torch::zeros({B, F, G, G, G}, coords.options());
int total = B * N;
int threads = 256;
int blocks = (total + threads - 1) / threads;
voxel_pool_kernel<<<blocks, threads>>>(
coords.data_ptr<float>(),
features.data_ptr<float>(),
weight.data_ptr<float>(),
bias.data_ptr<float>(),
voxel_grid.data_ptr<float>(),
B, N, F, G, voxel_size, inv_voxel_size
);
return voxel_grid;
}
'''
cpp_source = '''
torch::Tensor voxel_pool_cuda(torch::Tensor coords, torch::Tensor features, torch::Tensor weight, torch::Tensor bias, int grid_size, float voxel_size);
'''
mod = load_inline(
name='voxel_pool_mod',
cpp_sources=[cpp_source],
cuda_sources=[cuda_source],
functions=['voxel_pool_cuda'],
verbose=False
)
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:
return mod.voxel_pool_cuda(
coords, features,
self.linear.weight, self.linear.bias,
self.grid_size, self.voxel_size
)
```
This is the submission. The kernel fuses linear projection and atomic scatter-add into voxels, computing batch mean on-the-fly (acceptable for N=256). It should be faster than the looped index_put_.
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.h>\n#include <cuda_runtime.h>\n\n__global__ void voxel_pool_kernel(\n const float* __restrict__ coords,\n const float* __restrict__ features,\n const float* __restrict__ weight,\n const float* __restrict__ bias,\n float* __restrict__ voxel_grid,\n int B, int N, int F, int G, float voxel_size, float inv_voxel_size\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 // Compute center for batch b\n float cx = 0.0f, cy = 0.0f, cz = 0.0f;\n for (int i = 0; i < N; ++i) {\n int base = (b * N + i) * 3;\n cx += coords[base];\n cy += coords[base + 1];\n cz += coords[base + 2];\n }\n cx /= N; cy /= N; cz /= N;\n\n int base_c = (b * N + n) * 3;\n float x = (coords[base_c] - cx) * inv_voxel_size + G * 0.5f;\n float y = (coords[base_c + 1] - cy) * inv_voxel_size + G * 0.5f;\n float z = (coords[base_c + 2] - cz) * inv_voxel_size + G * 0.5f;\n\n int ix = (int)floorf(x);\n int iy = (int)floorf(y);\n int iz = (int)floorf(z);\n if (ix < 0) ix = 0; if (ix >= G) ix = G-1;\n if (iy < 0) iy = 0; if (iy >= G) iy = G-1;\n if (iz < 0) iz = 0; if (iz >= G) iz = G-1;\n\n // Linear projection for this point's features\n const float* feat_ptr = features + (b * N + n) * F;\n float* out_ptr_base = voxel_grid + b * F * G * G * G;\n\n for (int d = 0; d < F; ++d) {\n float val = bias[d];\n const float* w_row = weight + d * F;\n for (int k = 0; k < F; ++k) {\n val += feat_ptr[k] * w_row[k];\n }\n float* voxel_addr = out_ptr_base + d * G*G*G + ix * G*G + iy * G + iz;\n atomicAdd(voxel_addr, val);\n }\n}\n\ntorch::Tensor voxel_pool_cuda(\n torch::Tensor coords,\n torch::Tensor features,\n torch::Tensor weight,\n torch::Tensor bias,\n int grid_size,\n float voxel_size\n) {\n int B = coords.size(0);\n int N = coords.size(1);\n int F = features.size(2);\n int G = grid_size;\n float inv_voxel_size = 1.0f / voxel_size;\n\n auto voxel_grid = torch::zeros({B, F, G, G, G}, coords.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 coords.data_ptr<float>(),\n features.data_ptr<float>(),\n weight.data_ptr<float>(),\n bias.data_ptr<float>(),\n voxel_grid.data_ptr<float>(),\n B, N, F, G, voxel_size, inv_voxel_size\n );\n\n return voxel_grid;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor voxel_pool_cuda(torch::Tensor coords, torch::Tensor features, torch::Tensor weight, torch::Tensor bias, int grid_size, float voxel_size);\n'''\n\nmod = load_inline(\n name='voxel_pool_mod',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['voxel_pool_cuda'],\n verbose=False\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 return mod.voxel_pool_cuda(\n coords, features,\n self.linear.weight, self.linear.bias,\n self.grid_size, self.voxel_size\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