correct
21_RadialBasisFunctionExpansion.py
started
2026-05-11T10:47:17Z
timingkernelbench.timing
kernel mean / median / std
0.74 / — / 0.00 μs
ref mean / median / std
0.51 / — / 0.01 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
20.00x
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
dram bandwidth
449.52 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 011.0s turn · 11.0s LLM · 0.0s tools · tokens in=1,665 · out=6 · total=2,065
reasoning
The task is to replace the forward computation of this RBF expansion model with a custom CUDA kernel.
I need to replace the forward computation of the PyTorch model with a custom CUDA kernel. The model computes radial basis function features from distances with a cosine cutoff envelope. The input is distances of shape (B, N, N), and output is (*distances.shape, num_rbf).
call → get_gpu_specs
callget_gpu_specs
{}
result ← get_gpu_specs (ok)
okget_gpu_specs
GPU specs for NVIDIA H100 80GB HBM3:
total memory (runtime): 79.1 GB
spec entry: H100
GPU Architecture: Hopper
GPU Memory: 80GB
Memory Bandwidth: 3.35 TB/s
FP64 TFLOPS: 34
FP64 Tensor Core TFLOPS: 67
FP32 TFLOPS: 67
TF32 Tensor Core TFLOPS: 989 with sparsity
BFLOAT16 Tensore Core TFLOPS: 1979 with sparsity
FP16 Tensor Core TFLOPS: 1979 with sparsity
FP8 Tensor Core TFLOPS: 3958 with sparsity
INT8 Tensor Core TOPS: 3958 with sparsity
Register File Size: 64K 32-bit registers per SM
Maximum number of registers per thread: 255
Maximum number of thread blocks per SM: 32
Shared memory capacity per SM: 228 KB
Maximum shared memory per thread block: 227 KB
turn 1175.0s turn · 21.0s LLM · 154.0s tools · tokens in=2,298 · out=627 · total=2,967
reasoning
GPU is H100, high bandwidth.
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\n__global__ void rbf_kernel(\n const float* __restrict__ distances,\n const float* __restrict__ offsets,\n float* __restrict__ output,\n float cutoff,\n float coeff,\n int64_t numel,\n int num_rbf\n) {\n int64_t idx = blockIdx.x * blockDim.x + threadIdx.x;\n if (idx >= numel) return;\n\n float d = distances[idx];\n float env = 0.0f;\n if (d < cutoff) {\n float arg = 3.141592653589793f * d / cutoff;\n env = 0.5f * (cosf(arg) + 1.0f);\n }\n for (int k = 0; k < num_rbf; ++k) {\n float diff = d - offsets[k];\n float rbf = expf(coeff * diff * diff);\n output[idx * num_rbf + k] = rbf * env;\n }\n}\n\ntorch::Tensor rbf_cuda(torch::Tensor distances, torch::Tensor offsets, float cutoff, float coeff) {\n auto sizes = distances.sizes().vec();\n sizes.push_back(offsets.size(0));\n auto output = torch::empty(sizes, distances.options());\n\n int64_t numel = distances.numel();\n int num_rbf = offsets.size(0);\n\n const int block_size = 256;\n int num_blocks = (numel + block_size - 1) / block_size;\n\n rbf_kernel<<<num_blocks, block_size>>>(\n distances.data_ptr<float>(),\n offsets.data_ptr<float>(),\n output.data_ptr<float>(),\n cutoff,\n coeff,\n numel,\n num_rbf\n );\n\n return output;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor rbf_cuda(torch::Tensor distances, torch::Tensor offsets, float cutoff, float coeff);\n\"\"\"\n\nmodule = load_inline(\n name=\"rbf_module\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"rbf_cuda\"],\n verbose=True\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 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 2342.5s turn · 9.4s LLM · 333.0s tools · tokens in=2,984 · out=628 · total=3,619
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\n__global__ void rbf_kernel(\n const float* __restrict__ distances,\n const float* __restrict__ offsets,\n float* __restrict__ output,\n float cutoff,\n float coeff,\n int64_t numel,\n int num_rbf\n) {\n int64_t idx = blockIdx.x * blockDim.x + threadIdx.x;\n if (idx >= numel) return;\n\n float d = distances[idx];\n float env = 0.0f;\n if (d < cutoff) {\n float arg = 3.141592653589793f * d / cutoff;\n env = 0.5f * (cosf(arg) + 1.0f);\n }\n for (int k = 0; k < num_rbf; ++k) {\n float diff = d - offsets[k];\n float rbf = expf(coeff * diff * diff);\n output[idx * num_rbf + k] = rbf * env;\n }\n}\n\ntorch::Tensor rbf_cuda(torch::Tensor distances, torch::Tensor offsets, float cutoff, float coeff) {\n auto sizes = distances.sizes().vec();\n sizes.push_back(offsets.size(0));\n auto output = torch::empty(sizes, distances.options());\n\n int64_t numel = distances.numel();\n int num_rbf = offsets.size(0);\n\n const int block_size = 256;\n int num_blocks = (numel + block_size - 1) / block_size;\n\n rbf_kernel<<<num_blocks, block_size>>>(\n distances.data_ptr<float>(),\n offsets.data_ptr<float>(),\n output.data_ptr<float>(),\n cutoff,\n coeff,\n numel,\n num_rbf\n );\n\n return output;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor rbf_cuda(torch::Tensor distances, torch::Tensor offsets, float cutoff, float coeff);\n\"\"\"\n\nmodule = 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 module.rbf_cuda(distances, self.offsets, self.cutoff, self.coeff)\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=6.89e-10 max_rel_err=1.54e-02
turn 3522.4s turn · 8.2s LLM · 514.2s tools · tokens in=3,679 · out=627 · total=4,356 · FINAL
reasoning
Correct! Now, to check if faster, I could submit, but since it's the final, and it should be faster as it's fused and CUDA.
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\n__global__ void rbf_kernel(\n const float* __restrict__ distances,\n const float* __restrict__ offsets,\n float* __restrict__ output,\n float cutoff,\n float coeff,\n int64_t numel,\n int num_rbf\n) {\n int64_t idx = blockIdx.x * blockDim.x + threadIdx.x;\n if (idx >= numel) return;\n\n float d = distances[idx];\n float env = 0.0f;\n if (d < cutoff) {\n float arg = 3.141592653589793f * d / cutoff;\n env = 0.5f * (cosf(arg) + 1.0f);\n }\n for (int k = 0; k < num_rbf; ++k) {\n float diff = d - offsets[k];\n float rbf = expf(coeff * diff * diff);\n output[idx * num_rbf + k] = rbf * env;\n }\n}\n\ntorch::Tensor rbf_cuda(torch::Tensor distances, torch::Tensor offsets, float cutoff, float coeff) {\n auto sizes = distances.sizes().vec();\n sizes.push_back(offsets.size(0));\n auto output = torch::empty(sizes, distances.options());\n\n int64_t numel = distances.numel();\n int num_rbf = offsets.size(0);\n\n const int block_size = 256;\n int num_blocks = (numel + block_size - 1) / block_size;\n\n rbf_kernel<<<num_blocks, block_size>>>(\n distances.data_ptr<float>(),\n offsets.data_ptr<float>(),\n output.data_ptr<float>(),\n cutoff,\n coeff,\n numel,\n num_rbf\n );\n\n return output;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor rbf_cuda(torch::Tensor distances, torch::Tensor offsets, float cutoff, float coeff);\n\"\"\"\n\nmodule = 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 module.rbf_cuda(distances, self.offsets, self.cutoff, self.coeff)\n"
}
result ← submit_kernel (ok)
oksubmit_kernel
submit_kernel PASSED: (5 / 5) correctness trials all passed.
Kernel runtime: 0.74 μs
Runtime stats: mean=0.74μs median=?μs std=0.00μs
Numerical precision: max_abs_err=2.38e-07 mean_abs_err=6.89e-10 max_rel_err=1.54e-02
Memory: 163.0 MB peak (ref: 293.0 MB, ratio: 0.5563x)
Kernel launches: 1 (ref: 20, fusion ratio: 20.0)
SOL score: 0.1342
final submitted kernel (85 lines, 2,280 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>
__global__ void rbf_kernel(
const float* __restrict__ distances,
const float* __restrict__ offsets,
float* __restrict__ output,
float cutoff,
float coeff,
int64_t numel,
int num_rbf
) {
int64_t idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx >= numel) return;
float d = distances[idx];
float env = 0.0f;
if (d < cutoff) {
float arg = 3.141592653589793f * d / cutoff;
env = 0.5f * (cosf(arg) + 1.0f);
}
for (int k = 0; k < num_rbf; ++k) {
float diff = d - offsets[k];
float rbf = expf(coeff * diff * diff);
output[idx * num_rbf + k] = rbf * env;
}
}
torch::Tensor rbf_cuda(torch::Tensor distances, torch::Tensor offsets, float cutoff, float coeff) {
auto sizes = distances.sizes().vec();
sizes.push_back(offsets.size(0));
auto output = torch::empty(sizes, distances.options());
int64_t numel = distances.numel();
int num_rbf = offsets.size(0);
const int block_size = 256;
int num_blocks = (numel + block_size - 1) / block_size;
rbf_kernel<<<num_blocks, block_size>>>(
distances.data_ptr<float>(),
offsets.data_ptr<float>(),
output.data_ptr<float>(),
cutoff,
coeff,
numel,
num_rbf
);
return output;
}
"""
cpp_source = """
torch::Tensor rbf_cuda(torch::Tensor distances, torch::Tensor offsets, float cutoff, float coeff);
"""
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 module.rbf_cuda(distances, self.offsets, self.cutoff, self.coeff)