grok-4.3 · popcorn · pop_l123_st_gpt
correct 21_RadialBasisFunctionExpansion.py
model
grok-4.3
turns / max
1 / 1
tool calls
1
tokens (total)
2302
wall clock
284.7s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
0.74 μs
ref runtime
0.53 μs
speedup
0.72x
started
2026-05-11T08:53:45Z
timingkernelbench.timing
kernel runtime
0.74 μs
ref runtime
0.53 μs
speedup
0.72x
kernel mean / median / std
0.74 / — / 0.01 μs
ref mean / median / std
0.53 / — / 0.06 μs
numerical precisioneval_kernel_against_ref
max abs error
2.38e-07
mean abs error
6.89e-10
max rel error
0.02
mean rel error
7.23e-08
gpu memoryextended_metrics.measure_memory
kernel peak
130.00 MB
ref peak
260.00 MB
kernel / ref
0.50x
kernel launches / fusiontorch.profiler
kernel launches
1
ref launches
20
fusion ratio (ref / kernel)
20.00x
total cuda time
43 μs
energyNVML (pynvml)
kernel mJ / run
1833.72
ref mJ / run
229.14
ratio (ref / kernel)
0.12x
avg power
2483.02 W
ref avg power
451.61 W
SOL (speed-of-light)nsight_hardware_counters
SOL score (0–1)
0.135
dram util %
13.5%
compute util %
0.0%
bottleneck
unknown
dominant pipe
None
arithmetic intensity
—
ridge point
20.00
roofline / occupancynsight_hardware_counters
dram bandwidth
452.84 GB/s
dram util %
13.5%
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
85.4%
bottleneck
unknown
dominant pipe %
—
arithmetic intensity
—
ridge point
20.00
L1 hit %
—
L2 hit %
94.3%
ld sectors / req
1.05
st sectors / req
—
regs / thread
29
smem / block
1024 B
block size
256
peak bw
3350.00 GB/s
peak fp32 tflops
67.00
peak fp16 tflops
1979.00
warp stalls (top 5)nsight
long_scoreboard
11246.9%
mio_throttle
906.1%
short_scoreboard
194.9%
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 0284.7s turn · 29.7s LLM · 255.1s tools · tokens in=1,298 · out=675 · total=2,302 · FINAL
reasoning
The task is to replace the forward computation of a PyTorch RBF expansion model with a custom CUDA kernel that's faster but numerically equivalent. The task is to create a new PyTorch model class called ModelNew that replaces the forward computation of an existing RBF expansion model with a custom CUDA kernel for better performance, while keeping it numerically equivalent. The computation involves computing Radial Basis Function (RBF) features for distances with a cosine cutoff envelope.
call → submit_kernel
callsubmit_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nimport math\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#include <math.h>\n\n__global__ void rbf_kernel(\n const float* __restrict__ distances,\n float* __restrict__ output,\n const float* __restrict__ offsets,\n float coeff,\n float cutoff,\n int num_rbf,\n int total_elements\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n if (idx >= total_elements) return;\n\n float d = distances[idx];\n float d_scaled = d;\n bool mask = d_scaled < cutoff;\n float env = 0.0f;\n if (mask) {\n float arg = 3.141592653589793f * d_scaled / cutoff;\n env = 0.5f * (cosf(arg) + 1.0f);\n }\n\n float* out_ptr = output + idx * num_rbf;\n for (int i = 0; i < num_rbf; i++) {\n float diff = d_scaled - offsets[i];\n float rbf_val = expf(coeff * diff * diff);\n out_ptr[i] = rbf_val * env;\n }\n}\n\ntorch::Tensor rbf_expansion_cuda(torch::Tensor distances, torch::Tensor offsets, float coeff, float cutoff, int num_rbf) {\n auto sizes = distances.sizes();\n std::vector<int64_t> out_sizes(sizes.begin(), sizes.end());\n out_sizes.push_back(num_rbf);\n auto output = torch::empty(out_sizes, distances.options());\n\n int total_elements = distances.numel();\n int threads = 256;\n int blocks = (total_elements + threads - 1) / threads;\n\n rbf_kernel<<<blocks, threads>>>(\n distances.data_ptr<float>(),\n output.data_ptr<float>(),\n offsets.data_ptr<float>(),\n coeff,\n cutoff,\n num_rbf,\n total_elements\n );\n\n return output;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor rbf_expansion_cuda(torch::Tensor distances, torch::Tensor offsets, float coeff, float cutoff, int num_rbf);\n'''\n\nrbf_module = load_inline(\n name='rbf_cuda',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['rbf_expansion_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_expansion_cuda(distances, self.offsets, self.coeff, self.cutoff, self.num_rbf)\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.01μs Numerical precision: max_abs_err=2.38e-07 mean_abs_err=6.89e-10 max_rel_err=1.54e-02 Memory: 130.0 MB peak (ref: 260.0 MB, ratio: 0.5x) Kernel launches: 1 (ref: 20, fusion ratio: 20.0) Energy: 1833.72 mJ/run (ref: 229.14 mJ, ratio: 0.125x) SOL score: 0.1352
final submitted kernel (89 lines, 2,517 chars) — full source, untruncated
import torch
import torch.nn as nn
import math
from torch.utils.cpp_extension import load_inline

cuda_source = '''
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <math.h>

__global__ void rbf_kernel(
    const float* __restrict__ distances,
    float* __restrict__ output,
    const float* __restrict__ offsets,
    float coeff,
    float cutoff,
    int num_rbf,
    int total_elements
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx >= total_elements) return;

    float d = distances[idx];
    float d_scaled = d;
    bool mask = d_scaled < cutoff;
    float env = 0.0f;
    if (mask) {
        float arg = 3.141592653589793f * d_scaled / cutoff;
        env = 0.5f * (cosf(arg) + 1.0f);
    }

    float* out_ptr = output + idx * num_rbf;
    for (int i = 0; i < num_rbf; i++) {
        float diff = d_scaled - offsets[i];
        float rbf_val = expf(coeff * diff * diff);
        out_ptr[i] = rbf_val * env;
    }
}

torch::Tensor rbf_expansion_cuda(torch::Tensor distances, torch::Tensor offsets, float coeff, float cutoff, int num_rbf) {
    auto sizes = distances.sizes();
    std::vector<int64_t> out_sizes(sizes.begin(), sizes.end());
    out_sizes.push_back(num_rbf);
    auto output = torch::empty(out_sizes, distances.options());

    int total_elements = distances.numel();
    int threads = 256;
    int blocks = (total_elements + threads - 1) / threads;

    rbf_kernel<<<blocks, threads>>>(
        distances.data_ptr<float>(),
        output.data_ptr<float>(),
        offsets.data_ptr<float>(),
        coeff,
        cutoff,
        num_rbf,
        total_elements
    );

    return output;
}
'''

cpp_source = '''
torch::Tensor rbf_expansion_cuda(torch::Tensor distances, torch::Tensor offsets, float coeff, float cutoff, int num_rbf);
'''

rbf_module = load_inline(
    name='rbf_cuda',
    cpp_sources=[cpp_source],
    cuda_sources=[cuda_source],
    functions=['rbf_expansion_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_expansion_cuda(distances, self.offsets, self.coeff, self.cutoff, self.num_rbf)