FW-GLM-5-1 · popcorn · pop_l123_st_gpt
correct 21_RadialBasisFunctionExpansion.py
model
FW-GLM-5-1
turns / max
1 / 1
tool calls
1
tokens (total)
4621
wall clock
460.1s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
0.85 μs
ref runtime
0.59 μs
speedup
0.69x
started
2026-05-11T09:26:00Z
timingkernelbench.timing
kernel runtime
0.85 μs
ref runtime
0.59 μs
speedup
0.69x
kernel mean / median / std
0.85 / — / 0.01 μs
ref mean / median / std
0.59 / — / 0.02 μs
numerical precisioneval_kernel_against_ref
max abs error
4.29e-05
mean abs error
2.90e-07
max rel error
29803.23
mean rel error
0.00
gpu memoryextended_metrics.measure_memory
kernel peak
163.08 MB
ref peak
293.08 MB
kernel / ref
0.56x
kernel launches / fusiontorch.profiler
kernel launches
1
ref launches
20
fusion ratio (ref / kernel)
20.00x
total cuda time
23 μs
SOL (speed-of-light)nsight_hardware_counters
SOL score (0–1)
0.162
dram util %
16.2%
compute util %
0.0%
bottleneck
unknown
dominant pipe
None
arithmetic intensity
—
ridge point
20.00
roofline / occupancynsight_hardware_counters
dram bandwidth
543.59 GB/s
dram util %
16.2%
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
87.1%
bottleneck
unknown
dominant pipe %
—
arithmetic intensity
—
ridge point
20.00
L1 hit %
—
L2 hit %
93.8%
ld sectors / req
4.00
st sectors / req
—
regs / thread
32
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
9759.5%
mio_throttle
4468.8%
short_scoreboard
786.5%
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 0 · FINAL · 1 tool call(s) · 460.1s turn · 41.5s LLM · tokens in=1,248 · out=3,373 · total=4,621 · 19,736 chars (click to expand)
turn 0460.1s turn · 41.5s LLM · 418.6s tools · tokens in=1,248 · out=3,373 · total=4,621 · FINAL
reasoning
Let me analyze the task: 1. I need to implement a custom CUDA kernel for RBF (Radial Basis Function) expansion 2. The forward computation involves: - Computing a cosine cutoff envelope - Computing RBF features using Gaussian basis functions - Multiplying them together Let me understand the computation: - Input: `distances` tensor of shape (B, N, N) where B is batch size, N is num_atoms - Output: shape (B, N, N, num_rbf) The computation: 1. `d_scaled = distances.unsqueeze(-1)` - adds a dimension, shape becomes (B, N, N, 1) 2. `envelope = 0.5 * (torch.cos(math.pi * d_scaled / self.cutoff) + 1.0)` 3. `envelope = envelope * (d_scaled < self.cutoff).float()` 4. `rbf = torch.exp(self.coeff * (d_scaled - self.offsets) ** 2)` 5. `return rbf * envelope` The offsets are `torch.linspace(0.0, cutoff, num_rbf)` - so they go from 0 to cutoff. `self.coeff = -0.5 / (offsets[1] - offsets[0]).item() ** 2` Let me think about the CUDA kernel design: - Each thread will process one element of the input distances tensor and compute all num_rbf output values - Or we could have each thread process one output element (one distance, one rbf index) Given the output shape is (B, N, N, num_rbf), and the computation for each (b, i, j, k) depends only on distances[b,i,j] and offsets[k], we have good parallelism. For a grid: - Grid dimension: we can use 3D grid for (B, N, N) and have each thread compute all num_rbf outputs - Or use 4D grid (B, N, N, num_rbf) where each thread computes one output I'm focusing on the computational characteristics. Each distance computation involves a cosine envelope and multiple RBF calculations. The kernel needs to efficiently handle the Gaussian basis function generation across multiple offsets. The core operation involves parallel computation of envelope and RBF values, with careful attention to the cutoff condition and exponential decay. Vectorized memory access and efficient thread utilization will be crucial for performance. The input tensor spans 524,288 elements across batch, neighbor, and atom dimensions, with 64 radial basis functions to compute for each distance. Two primary kernel strategies emerge: a compact single-thread-per-distance approach versus a thread-per-output-element method. The single-thread approach minimizes redundant distance loads but risks performance from loop overhead, while the thread-per-output strategy maximizes parallelism but increases global memory bandwidth pressure. Shared memory optimization offers a promising alternative, potentially reducing global memory access overhead while maintaining high thread utilization. The offsets' repeated use across threads suggests potential caching benefits. Key optimization vectors include memory access patterns, loop vectorization, and strategic thread block configurations to balance computational efficiency and memory bandwidth constraints. The chosen CUDA kernel configuration maps a three-dimensional grid directly to batch, row, and column indices, with each thread computing multiple radial basis function (RBF) outputs. This approach leverages constant memory for offsets and enables coalesced memory writes through careful output tensor indexing. The implementation focuses on efficient memory access and computational parallelism, with potential for further optimization through careful selection of CUDA runtime parameters. I'll explore two primary optimization strategies: constant memory caching and vectorized memory operations, targeting improved performance across different hardware generations. The core kernel design focuses on efficiently computing radial basis function (RBF) values with a cosine envelope, carefully managing thread indexing and memory access patterns. Key considerations include minimizing redundant computations and maximizing memory bandwidth utilization. The implementation requires precise handling of input distances, offsets, and output tensor dimensions, ensuring numerical stability and computational efficiency across varying input sizes. Performance optimization centers on leveraging constant memory for small, read-only data like offsets and precomputed coefficients, which can significantly reduce global memory access overhead. I'll explore vectorized memory writes using float4, which could potentially improve bandwidth utilization. However, this approach requires careful register management and might introduce complexity that outweighs potential gains. The current kernel implementation seems promising, with straightforward indexing and efficient computation of Gaussian basis functions. Careful memory access patterns and minimal divergence remain key optimization strategies. I'll focus on exploring shared memory reduction techniques and carefully evaluating vectorization trade-offs for the radial basis function computation. The computation reveals a nuanced scaling pattern across RBF indices, with potential optimization opportunities through careful memory access and precomputation strategies. The offset range and coefficient calculations suggest interesting computational symmetry that could inform efficient kernel design. By strategically repositioning the envelope computation relative to the RBF loop, there's potential to minimize redundant calculations and improve overall computational efficiency. The exponential decay characteristics and distance scaling parameters provide a structured approach to optimizing the underlying mathematical transformation. The current implementation suggests a methodical approach to managing computational complexity while maintaining precision across the radial basis function evaluation. I'll focus on creating a streamlined kernel that efficiently computes radial basis functions with precise scaling and envelope calculations, leveraging CUDA's parallel processing capabilities to maximize performance across thread blocks and dimensions. The kernel design aims to minimize computational overhead by strategically positioning critical calculations and utilizing vectorized memory access patterns where possible. I'll precompute the squared distance and envelope factor outside the radial basis function loop, reducing redundant calculations and improving overall kernel efficiency. By extracting these values before the inner loop, I can streamline the computation and potentially reduce instruction overhead. The approach involves calculating the envelope and scaled distance squared once, then applying these precomputed values across all radial basis function iterations. I'll leverage CUDA's constant memory and vectorized memory access to optimize performance, targeting efficient computation across multiple radial basis functions. The kernel will compute radial basis function values with a cosine envelope, using precomputed coefficients and offsets stored in constant memory. By processing four elements simultaneously, I can maximize memory bandwidth and computational efficiency. The implementation strategically handles edge cases for vectorized processing, ensuring all 64 elements are computed accurately across different thread configurations. I'll use Python's torch extension mechanism to compile the CUDA kernel inline, creating a flexible and high-performance implementation. The kernel computes radial basis function values efficiently by leveraging CUDA's parallel processing capabilities. Each thread handles a specific batch, row, and column index, calculating scaled distance coefficients and applying a cosine envelope cutoff. The core computation involves precomputing distance scaling factors and applying exponential transformations across multiple radial basis function centers, optimizing memory access and computational efficiency. I'll focus on optimizing memory access patterns, ensuring coalesced writes by carefully organizing thread blocks and grid dimensions to maximize parallel processing performance. The current memory write strategy reveals a potential performance bottleneck. Threads within a warp are writing to non-contiguous memory locations, causing scattered memory access patterns that reduce bandwidth efficiency. This strided access means memory writes are fragmented, increasing latency and reducing throughput. The key observation is that thread mapping directly impacts memory access patterns. By reorganizing thread blocks and grid dimensions, I can potentially align thread writes to contiguous memory segments, dramatically improving memory coalescing and overall kernel performance. A potential optimization involves restructuring the thread and block configuration to ensure sequential memory writes across the radial basis function dimension, potentially transforming the current strided pattern into a more efficient, contiguous access model. The core strategy focuses on maximizing memory bandwidth and computational efficiency by carefully organizing thread block dimensions and memory access patterns. Exploring vectorized memory operations and potential shared memory techniques could provide incremental performance improvements. The critical observation centers on how threads write to the output tensor, ensuring writes align with memory coalescing principles and minimize potential performance bottlenecks. Computational complexity reveals interesting parallelism opportunities, with roughly 2 million threads potentially available for processing. The small constant data size suggests constant memory optimization could be highly effective. Kernel design focuses on efficient RBF computation, leveraging constant memory for offsets and carefully managing thread indexing to maximize computational throughput. I'll compute scaled distances, applying a cosine envelope and exponential decay based on predefined coefficients and offsets. The approach ensures precise radial basis function calculations across multidimensional data. Each thread processes a specific spatial location, generating a feature vector representing distance transformations. The implementation leverages GPU parallelism to efficiently compute localized representations across batch, row, and column dimensions. The kernel maps input distances through exponential and cosine functions, creating a compact feature encoding that captures spatial proximity characteristics. Performance considerations suggest strategic memory access and thread block sizing could optimize computational efficiency. I'll configure thread blocks with 128 threads per block, which balances occupancy and resource utilization. The grid dimensions will dynamically adapt based on input batch size and spatial dimensions, ensuring comprehensive coverage of computational space. The kernel launch strategy leverages CUDA's 3D grid capabilities, mapping thread blocks across batch, row, and column dimensions to maximize parallel processing potential. Kernel implementation focuses on computing radial basis function transformations, utilizing constant memory for pre-stored offsets and handling edge cases through warp-level processing techniques. After evaluating vectorization trade-offs, the decision is to maintain the simple kernel approach, prioritizing code clarity and avoiding complex indexing logic that might introduce potential bugs. The implementation leverages constant memory optimization and straightforward computation, accepting potential minor performance variations compared to vectorized approaches. Checking offsets memory transfer to CUDA constant memory reveals careful handling of tensor storage and CUDA runtime requirements for correct memory allocation and copying. The kernel computes radial basis function values by scaling distances, applying a cosine envelope cutoff, and calculating exponential decay based on predefined coefficients. Each thread calculates contributions for multiple RBF basis functions efficiently using constant memory offsets. The implementation leverages CUDA's parallel processing capabilities to transform distance calculations across multiple dimensions, optimizing computational performance through strategic memory management and vectorized operations. I'll ensure thread safety by checking index boundaries and systematically processing each radial basis function coefficient, maintaining computational precision while minimizing redundant calculations. The memory writes are strategically structured to ensure contiguous access across threads, enabling efficient coalesced memory operations within each warp's computational block. This approach optimizes memory bandwidth and reduces potential access latency. The output tensor's layout reveals a systematic thread mapping strategy where each thread handles a specific slice of radial basis function computations, carefully aligning memory access patterns for maximum performance. The block size optimization reveals nuanced trade-offs between register pressure, shared memory constraints, and theoretical occupancy. While 128 threads per block seems conservative, it minimizes potential register spill risks. Exploring 256 threads could improve parallelism if hardware permits, but requires careful performance validation. Ultimately, selecting an optimal block dimension demands empirical testing across different GPU architectures and computational kernels. The kernel optimizes RBF calculations with constant memory, cutoff envelope, and loop unrolling, ensuring efficient GPU performance through carefully designed memory access and computation strategies.
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\n# Compile CUDA kernel at module load time\ncuda_source = r'''\n#include <cuda.h>\n#include <cuda_runtime.h>\n\n// Constant memory for offsets (64 floats easily fits in constant cache)\n__constant__ float d_offsets[64];\n__constant__ float d_coeff;\n__constant__ float d_cutoff;\n__constant__ float d_pi_over_cutoff;\n\n// Optimized kernel: each thread computes one distance -> all RBF outputs\n// Memory access pattern: strided reads (efficient with cache), coalesced writes\n// Constant memory broadcasts offsets to all threads efficiently\n__global__ void rbf_kernel(const float* __restrict__ distances,\n float* __restrict__ output,\n int total_distances,\n int num_rbf) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n \n if (idx >= total_distances) return;\n \n float d = distances[idx];\n float d_scaled = d / d_cutoff;\n \n // Cosine envelope with cutoff\n float envelope = 0.0f;\n if (d < d_cutoff) {\n envelope = 0.5f * (cosf(d_pi_over_cutoff * d) + 1.0f);\n }\n \n // Precompute for RBF: coeff * d^2 and d\n float coeff_d_sq = d_coeff * d * d;\n float coeff_d = d_coeff * 2.0f * d;\n \n // Compute all RBF outputs with loop unrolling hint\n #pragma unroll 16\n for (int k = 0; k < num_rbf; ++k) {\n float offset = d_offsets[k];\n // Expand: coeff * (d - offset)^2 = coeff * d^2 - 2*coeff*d*offset + coeff*offset^2\n float rbf = expf(coeff_d_sq - coeff_d * offset + d_coeff * offset * offset);\n output[idx * num_rbf + k] = rbf * envelope;\n }\n}\n\n// Host wrapper\ntorch::Tensor rbf_forward(torch::Tensor distances, torch::Tensor offsets, \n float coeff, float cutoff, int num_rbf) {\n // Ensure inputs are contiguous and on GPU\n auto distances_cont = distances.contiguous();\n TORCH_CHECK(distances_cont.dim() >= 1, \"distances must have at least 1 dimension\");\n TORCH_CHECK(distances_cont.scalar_type() == torch::kFloat32, \"distances must be float32\");\n TORCH_CHECK(distances_cont.is_cuda(), \"distances must be on CUDA\");\n \n // Compute total number of distance elements\n int total_distances = 1;\n for (int i = 0; i < distances_cont.dim(); ++i) {\n total_distances *= distances_cont.size(i);\n }\n \n // Create output tensor with shape (*distances.shape, num_rbf)\n std::vector<int64_t> out_shape(distances_cont.sizes().vec());\n out_shape.push_back(num_rbf);\n auto output = torch::empty(out_shape, distances_cont.options());\n \n // Copy offsets to constant memory\n cudaMemcpyToSymbol(d_offsets, offsets.data_ptr<float>(), num_rbf * sizeof(float));\n cudaMemcpyToSymbol(d_coeff, &coeff, sizeof(float));\n cudaMemcpyToSymbol(d_cutoff, &cutoff, sizeof(float));\n float pi_over_cutoff = M_PI / cutoff;\n cudaMemcpyToSymbol(d_pi_over_cutoff, &pi_over_cutoff, sizeof(float));\n \n // Launch kernel\n int threads = 256;\n int blocks = (total_distances + threads - 1) / threads;\n \n rbf_kernel<<<blocks, threads>>>(\n distances_cont.data_ptr<float>(),\n output.data_ptr<float>(),\n total_distances,\n num_rbf\n );\n \n return output;\n}\n'''\n\ncpp_source = r'''\ntorch::Tensor rbf_forward(torch::Tensor distances, torch::Tensor offsets,\n float coeff, float cutoff, int num_rbf);\n'''\n\n# Load and compile the CUDA extension\nrbf_extension = load_inline(\n name='rbf_extension',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['rbf_forward'],\n extra_cuda_cflags=['-O3', '-use_fast_math'],\n verbose=False\n)\n\n\nclass ModelNew(nn.Module):\n \"\"\"\n CUDA-optimized RBF expansion kernel.\n Replaces PyTorch reference with custom kernel for improved performance.\n \"\"\"\n \n def __init__(self, num_rbf=64, cutoff=5.0):\n super().__init__()\n self.num_rbf = num_rbf\n self.cutoff = cutoff\n \n offsets = torch.linspace(0.0, cutoff, num_rbf)\n self.register_buffer(\"offsets\", offsets)\n \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 \"\"\"\n Args:\n distances: (B, N, N) - pairwise distance matrix (or any shape)\n Returns:\n rbf_features: (*distances.shape, num_rbf) - expanded features\n \"\"\"\n return rbf_extension.rbf_forward(\n distances, \n self.offsets, \n self.coeff, \n self.cutoff, \n self.num_rbf\n )" }
result ← submit_kernel (ok)
oksubmit_kernel submit_kernel PASSED: (5 / 5) correctness trials all passed. Kernel runtime: 0.85 μs Runtime stats: mean=0.85μs median=?μs std=0.01μs Numerical precision: max_abs_err=4.29e-05 mean_abs_err=2.90e-07 max_rel_err=2.98e+04 Memory: 163.08 MB peak (ref: 293.08 MB, ratio: 0.5564x) Kernel launches: 1 (ref: 20, fusion ratio: 20.0) SOL score: 0.1623
final submitted kernel (139 lines, 4,677 chars) — full source, untruncated
import torch
import torch.nn as nn
import math
from torch.utils.cpp_extension import load_inline

# Compile CUDA kernel at module load time
cuda_source = r'''
#include <cuda.h>
#include <cuda_runtime.h>

// Constant memory for offsets (64 floats easily fits in constant cache)
__constant__ float d_offsets[64];
__constant__ float d_coeff;
__constant__ float d_cutoff;
__constant__ float d_pi_over_cutoff;

// Optimized kernel: each thread computes one distance -> all RBF outputs
// Memory access pattern: strided reads (efficient with cache), coalesced writes
// Constant memory broadcasts offsets to all threads efficiently
__global__ void rbf_kernel(const float* __restrict__ distances,
                           float* __restrict__ output,
                           int total_distances,
                           int num_rbf) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx >= total_distances) return;
    
    float d = distances[idx];
    float d_scaled = d / d_cutoff;
    
    // Cosine envelope with cutoff
    float envelope = 0.0f;
    if (d < d_cutoff) {
        envelope = 0.5f * (cosf(d_pi_over_cutoff * d) + 1.0f);
    }
    
    // Precompute for RBF: coeff * d^2 and d
    float coeff_d_sq = d_coeff * d * d;
    float coeff_d = d_coeff * 2.0f * d;
    
    // Compute all RBF outputs with loop unrolling hint
    #pragma unroll 16
    for (int k = 0; k < num_rbf; ++k) {
        float offset = d_offsets[k];
        // Expand: coeff * (d - offset)^2 = coeff * d^2 - 2*coeff*d*offset + coeff*offset^2
        float rbf = expf(coeff_d_sq - coeff_d * offset + d_coeff * offset * offset);
        output[idx * num_rbf + k] = rbf * envelope;
    }
}

// Host wrapper
torch::Tensor rbf_forward(torch::Tensor distances, torch::Tensor offsets, 
                          float coeff, float cutoff, int num_rbf) {
    // Ensure inputs are contiguous and on GPU
    auto distances_cont = distances.contiguous();
    TORCH_CHECK(distances_cont.dim() >= 1, "distances must have at least 1 dimension");
    TORCH_CHECK(distances_cont.scalar_type() == torch::kFloat32, "distances must be float32");
    TORCH_CHECK(distances_cont.is_cuda(), "distances must be on CUDA");
    
    // Compute total number of distance elements
    int total_distances = 1;
    for (int i = 0; i < distances_cont.dim(); ++i) {
        total_distances *= distances_cont.size(i);
    }
    
    // Create output tensor with shape (*distances.shape, num_rbf)
    std::vector<int64_t> out_shape(distances_cont.sizes().vec());
    out_shape.push_back(num_rbf);
    auto output = torch::empty(out_shape, distances_cont.options());
    
    // Copy offsets to constant memory
    cudaMemcpyToSymbol(d_offsets, offsets.data_ptr<float>(), num_rbf * sizeof(float));
    cudaMemcpyToSymbol(d_coeff, &coeff, sizeof(float));
    cudaMemcpyToSymbol(d_cutoff, &cutoff, sizeof(float));
    float pi_over_cutoff = M_PI / cutoff;
    cudaMemcpyToSymbol(d_pi_over_cutoff, &pi_over_cutoff, sizeof(float));
    
    // Launch kernel
    int threads = 256;
    int blocks = (total_distances + threads - 1) / threads;
    
    rbf_kernel<<<blocks, threads>>>(
        distances_cont.data_ptr<float>(),
        output.data_ptr<float>(),
        total_distances,
        num_rbf
    );
    
    return output;
}
'''

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

# Load and compile the CUDA extension
rbf_extension = load_inline(
    name='rbf_extension',
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=['rbf_forward'],
    extra_cuda_cflags=['-O3', '-use_fast_math'],
    verbose=False
)


class ModelNew(nn.Module):
    """
    CUDA-optimized RBF expansion kernel.
    Replaces PyTorch reference with custom kernel for improved performance.
    """
    
    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 any shape)
        Returns:
            rbf_features: (*distances.shape, num_rbf) - expanded features
        """
        return rbf_extension.rbf_forward(
            distances, 
            self.offsets, 
            self.coeff, 
            self.cutoff, 
            self.num_rbf
        )