correct
21_RadialBasisFunctionExpansion.py
started
2026-05-11T17:36:20Z
timingkernelbench.timing
kernel mean / median / std
0.15 / — / 4.42e-04 μs
ref mean / median / std
0.52 / — / 0.02 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
10.00x
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
dram bandwidth
767.65 GB/s
warp stalls (top 5)nsight
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn
import math
class Model(nn.Module):
"""
Radial Basis Function (RBF) expansion of interatomic distances.
Standard featurization in molecular graph neural networks (SchNet,
DimeNet, GemNet) that converts scalar distances into a feature vector
of Gaussian-basis evaluations.
"""
def __init__(self, num_rbf=64, cutoff=5.0):
super().__init__()
self.num_rbf = num_rbf
self.cutoff = cutoff
offsets = torch.linspace(0.0, cutoff, num_rbf)
self.register_buffer("offsets", offsets)
coeff = -0.5 / (offsets[1] - offsets[0]).item() ** 2
self.coeff = coeff
def forward(self, distances: torch.Tensor) -> torch.Tensor:
"""
Args:
distances: (B, N, N) – pairwise distance matrix (or (B, E) edge list)
Returns:
rbf_features: (*distances.shape, num_rbf) – expanded features
"""
# Cosine cutoff envelope
d_scaled = distances.unsqueeze(-1) # (..., 1)
envelope = 0.5 * (torch.cos(math.pi * d_scaled / self.cutoff) + 1.0)
envelope = envelope * (d_scaled < self.cutoff).float()
rbf = torch.exp(self.coeff * (d_scaled - self.offsets) ** 2)
return rbf * envelope
num_atoms = 256
num_rbf = 64
cutoff = 5.0
batch_size = 8
def get_inputs():
p = popcorn_pri
coords = torch.randn(p.jitter_int(batch_size), p.jitter_int(num_atoms), 3)
diff = coords.unsqueeze(2) - coords.unsqueeze(1)
distances = torch.sqrt((diff ** 2).sum(-1) + 1e-8)
return [distances]
def get_init_inputs():
return [num_rbf, cutoff]
turn 094.1s turn · 19.4s LLM · 74.7s tools · tokens in=2,196 · out=686 · total=3,441
reasoning
The task is to replace the forward computation of this RBF expansion model with a custom CUDA kernel that's faster than PyTorch.
I need to replace the forward computation of the PyTorch model with a custom CUDA kernel. The model computes radial basis function expansions for distances with a cosine cutoff envelope. The input is distances of shape (B, N, N) or (*, ), and it computes d_scaled = distances.
call → compile_kernel
callcompile_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\nimport math\n\ncuda_source = '''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <math.h>\n\n__global__ void rbf_kernel(\n const float* __restrict__ distances,\n const float* __restrict__ offsets,\n float* __restrict__ output,\n int total_elements,\n int num_rbf,\n float cutoff,\n float coeff\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total_out = total_elements * num_rbf;\n if (idx >= total_out) return;\n\n int elem_idx = idx / num_rbf;\n int rbf_idx = idx % num_rbf;\n\n float d = distances[elem_idx];\n float off = offsets[rbf_idx];\n float d_scaled = d;\n\n float envelope = 0.5f * (cosf(M_PI * d_scaled / cutoff) + 1.0f);\n envelope *= (d_scaled < cutoff) ? 1.0f : 0.0f;\n\n float rbf_val = expf(coeff * (d_scaled - off) * (d_scaled - off));\n output[idx] = rbf_val * envelope;\n}\n\ntorch::Tensor rbf_cuda(torch::Tensor distances, torch::Tensor offsets, float cutoff, float coeff) {\n int total_elements = distances.numel();\n int num_rbf = offsets.size(0);\n auto output = torch::empty({total_elements, num_rbf}, distances.options());\n \n int total_out = total_elements * num_rbf;\n int threads = 256;\n int blocks = (total_out + threads - 1) / threads;\n\n rbf_kernel<<<blocks, threads>>>(\n distances.data_ptr<float>(),\n offsets.data_ptr<float>(),\n output.data_ptr<float>(),\n total_elements,\n num_rbf,\n cutoff,\n coeff\n );\n\n // reshape to original shape + num_rbf\n std::vector<int64_t> new_shape(distances.sizes().begin(), distances.sizes().end());\n new_shape.push_back(num_rbf);\n return output.view(new_shape);\n}\n'''\n\ncpp_source = '''\ntorch::Tensor rbf_cuda(torch::Tensor distances, torch::Tensor offsets, float cutoff, float coeff);\n'''\n\nrbf_module = load_inline(\n name='rbf_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['rbf_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, num_rbf=64, cutoff=5.0):\n super().__init__()\n self.num_rbf = num_rbf\n self.cutoff = cutoff\n offsets = torch.linspace(0.0, cutoff, num_rbf)\n self.register_buffer(\"offsets\", offsets)\n coeff = -0.5 / (offsets[1] - offsets[0]).item() ** 2\n self.coeff = coeff\n\n def forward(self, distances: torch.Tensor) -> torch.Tensor:\n return rbf_module.rbf_cuda(distances, self.offsets, self.cutoff, self.coeff)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 129.4s turn · 5.1s LLM · 24.3s tools · tokens in=3,458 · out=687 · total=4,152
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\nimport math\n\ncuda_source = '''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <math.h>\n\n__global__ void rbf_kernel(\n const float* __restrict__ distances,\n const float* __restrict__ offsets,\n float* __restrict__ output,\n int total_elements,\n int num_rbf,\n float cutoff,\n float coeff\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total_out = total_elements * num_rbf;\n if (idx >= total_out) return;\n\n int elem_idx = idx / num_rbf;\n int rbf_idx = idx % num_rbf;\n\n float d = distances[elem_idx];\n float off = offsets[rbf_idx];\n float d_scaled = d;\n\n float envelope = 0.5f * (cosf(M_PI * d_scaled / cutoff) + 1.0f);\n envelope *= (d_scaled < cutoff) ? 1.0f : 0.0f;\n\n float rbf_val = expf(coeff * (d_scaled - off) * (d_scaled - off));\n output[idx] = rbf_val * envelope;\n}\n\ntorch::Tensor rbf_cuda(torch::Tensor distances, torch::Tensor offsets, float cutoff, float coeff) {\n int total_elements = distances.numel();\n int num_rbf = offsets.size(0);\n auto output = torch::empty({total_elements, num_rbf}, distances.options());\n \n int total_out = total_elements * num_rbf;\n int threads = 256;\n int blocks = (total_out + threads - 1) / threads;\n\n rbf_kernel<<<blocks, threads>>>(\n distances.data_ptr<float>(),\n offsets.data_ptr<float>(),\n output.data_ptr<float>(),\n total_elements,\n num_rbf,\n cutoff,\n coeff\n );\n\n // reshape to original shape + num_rbf\n std::vector<int64_t> new_shape(distances.sizes().begin(), distances.sizes().end());\n new_shape.push_back(num_rbf);\n return output.view(new_shape);\n}\n'''\n\ncpp_source = '''\ntorch::Tensor rbf_cuda(torch::Tensor distances, torch::Tensor offsets, float cutoff, float coeff);\n'''\n\nrbf_module = load_inline(\n name='rbf_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['rbf_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, num_rbf=64, cutoff=5.0):\n super().__init__()\n self.num_rbf = num_rbf\n self.cutoff = cutoff\n offsets = torch.linspace(0.0, cutoff, num_rbf)\n self.register_buffer(\"offsets\", offsets)\n coeff = -0.5 / (offsets[1] - offsets[0]).item() ** 2\n self.coeff = coeff\n\n def forward(self, distances: torch.Tensor) -> torch.Tensor:\n return rbf_module.rbf_cuda(distances, self.offsets, self.cutoff, self.coeff)\n"
}
result ← run_correctness (ok)
okrun_correctness
run_correctness PASSED: (3 / 3) trials all matched the reference.
Numerical precision: max_abs_err=2.38e-07 mean_abs_err=1.29e-09 max_rel_err=1.54e-02
turn 286.1s turn · 6.7s LLM · 79.4s tools · tokens in=4,212 · out=686 · total=4,953 · FINAL
reasoning
Correctness passed. Now, to check if it's faster, I should profile or just submit since it should be faster than PyTorch's multiple ops.
call → submit_kernel
callsubmit_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\nimport math\n\ncuda_source = '''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <math.h>\n\n__global__ void rbf_kernel(\n const float* __restrict__ distances,\n const float* __restrict__ offsets,\n float* __restrict__ output,\n int total_elements,\n int num_rbf,\n float cutoff,\n float coeff\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total_out = total_elements * num_rbf;\n if (idx >= total_out) return;\n\n int elem_idx = idx / num_rbf;\n int rbf_idx = idx % num_rbf;\n\n float d = distances[elem_idx];\n float off = offsets[rbf_idx];\n float d_scaled = d;\n\n float envelope = 0.5f * (cosf(M_PI * d_scaled / cutoff) + 1.0f);\n envelope *= (d_scaled < cutoff) ? 1.0f : 0.0f;\n\n float rbf_val = expf(coeff * (d_scaled - off) * (d_scaled - off));\n output[idx] = rbf_val * envelope;\n}\n\ntorch::Tensor rbf_cuda(torch::Tensor distances, torch::Tensor offsets, float cutoff, float coeff) {\n int total_elements = distances.numel();\n int num_rbf = offsets.size(0);\n auto output = torch::empty({total_elements, num_rbf}, distances.options());\n \n int total_out = total_elements * num_rbf;\n int threads = 256;\n int blocks = (total_out + threads - 1) / threads;\n\n rbf_kernel<<<blocks, threads>>>(\n distances.data_ptr<float>(),\n offsets.data_ptr<float>(),\n output.data_ptr<float>(),\n total_elements,\n num_rbf,\n cutoff,\n coeff\n );\n\n // reshape to original shape + num_rbf\n std::vector<int64_t> new_shape(distances.sizes().begin(), distances.sizes().end());\n new_shape.push_back(num_rbf);\n return output.view(new_shape);\n}\n'''\n\ncpp_source = '''\ntorch::Tensor rbf_cuda(torch::Tensor distances, torch::Tensor offsets, float cutoff, float coeff);\n'''\n\nrbf_module = load_inline(\n name='rbf_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['rbf_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, num_rbf=64, cutoff=5.0):\n super().__init__()\n self.num_rbf = num_rbf\n self.cutoff = cutoff\n offsets = torch.linspace(0.0, cutoff, num_rbf)\n self.register_buffer(\"offsets\", offsets)\n coeff = -0.5 / (offsets[1] - offsets[0]).item() ** 2\n self.coeff = coeff\n\n def forward(self, distances: torch.Tensor) -> torch.Tensor:\n return rbf_module.rbf_cuda(distances, self.offsets, self.cutoff, self.coeff)\n"
}
result ← submit_kernel (ok)
oksubmit_kernel
submit_kernel PASSED: (3 / 3) correctness trials all passed.
Kernel runtime: 0.15 μs
Runtime stats: mean=0.15μs median=?μs std=0.00μs
Numerical precision: max_abs_err=2.38e-07 mean_abs_err=1.29e-09 max_rel_err=1.54e-02
Memory: 130.0 MB peak (ref: 260.0 MB, ratio: 0.5x)
Kernel launches: 2 (ref: 20, fusion ratio: 10.0)
SOL score: 0.4016
final submitted kernel (89 lines, 2,558 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline
import math
cuda_source = '''
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <math.h>
__global__ void rbf_kernel(
const float* __restrict__ distances,
const float* __restrict__ offsets,
float* __restrict__ output,
int total_elements,
int num_rbf,
float cutoff,
float coeff
) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int total_out = total_elements * num_rbf;
if (idx >= total_out) return;
int elem_idx = idx / num_rbf;
int rbf_idx = idx % num_rbf;
float d = distances[elem_idx];
float off = offsets[rbf_idx];
float d_scaled = d;
float envelope = 0.5f * (cosf(M_PI * d_scaled / cutoff) + 1.0f);
envelope *= (d_scaled < cutoff) ? 1.0f : 0.0f;
float rbf_val = expf(coeff * (d_scaled - off) * (d_scaled - off));
output[idx] = rbf_val * envelope;
}
torch::Tensor rbf_cuda(torch::Tensor distances, torch::Tensor offsets, float cutoff, float coeff) {
int total_elements = distances.numel();
int num_rbf = offsets.size(0);
auto output = torch::empty({total_elements, num_rbf}, distances.options());
int total_out = total_elements * num_rbf;
int threads = 256;
int blocks = (total_out + threads - 1) / threads;
rbf_kernel<<<blocks, threads>>>(
distances.data_ptr<float>(),
offsets.data_ptr<float>(),
output.data_ptr<float>(),
total_elements,
num_rbf,
cutoff,
coeff
);
// reshape to original shape + num_rbf
std::vector<int64_t> new_shape(distances.sizes().begin(), distances.sizes().end());
new_shape.push_back(num_rbf);
return output.view(new_shape);
}
'''
cpp_source = '''
torch::Tensor rbf_cuda(torch::Tensor distances, torch::Tensor offsets, float cutoff, float coeff);
'''
rbf_module = load_inline(
name='rbf_module',
cpp_sources=[cpp_source],
cuda_sources=[cuda_source],
functions=['rbf_cuda'],
verbose=False
)
class ModelNew(nn.Module):
def __init__(self, num_rbf=64, cutoff=5.0):
super().__init__()
self.num_rbf = num_rbf
self.cutoff = cutoff
offsets = torch.linspace(0.0, cutoff, num_rbf)
self.register_buffer("offsets", offsets)
coeff = -0.5 / (offsets[1] - offsets[0]).item() ** 2
self.coeff = coeff
def forward(self, distances: torch.Tensor) -> torch.Tensor:
return rbf_module.rbf_cuda(distances, self.offsets, self.cutoff, self.coeff)