grok-4.3 · popcorn · pop_l123_all_gpt
correct 42_BatchedForwardKinematics.py
model
grok-4.3
turns / max
4 / 15
tool calls
4
tokens (total)
15285
wall clock
1212.8s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
0.01 μs
ref runtime
0.02 μs
speedup
3.11x
started
2026-05-11T22:32:28Z
timingkernelbench.timing
kernel runtime
0.01 μs
ref runtime
0.02 μs
speedup
3.11x
kernel mean / median / std
0.01 / — / 1.48e-04 μs
ref mean / median / std
0.02 / — / 6.11e-04 μs
numerical precisioneval_kernel_against_ref
max abs error
2.64e-07
mean abs error
1.32e-08
max rel error
1.22e-04
mean rel error
3.06e-07
gpu memoryextended_metrics.measure_memory
kernel peak
33.00 MB
ref peak
33.01 MB
kernel / ref
1.00x
kernel launches / fusiontorch.profiler
kernel launches
1
ref launches
14
fusion ratio (ref / kernel)
14.00x
total cuda time
22 μs
SOL (speed-of-light)nsight_hardware_counters
SOL score (0–1)
5.000e-04
dram util %
0.1%
compute util %
0.0%
bottleneck
unknown
dominant pipe
None
arithmetic intensity
—
ridge point
20.00
roofline / occupancynsight_hardware_counters
dram bandwidth
1.52 GB/s
dram util %
0.0%
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
4.1%
bottleneck
unknown
dominant pipe %
—
arithmetic intensity
—
ridge point
20.00
L1 hit %
—
L2 hit %
99.0%
ld sectors / req
14.50
st sectors / req
—
regs / thread
24
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
short_scoreboard
664.2%
long_scoreboard
374.3%
reference PyTorch source (the prompt the agent sees)
import math

import torch
import torch.nn as nn

# Default ModelNew is PyTorch-only. run_and_check: use check_kernel=False with backend=cuda
# until the candidate file contains a real CUDA kernel (see kernel_static_checker).


class Model(nn.Module):
    """
    **Batched forward kinematics (FK)** for a **planar serial** revolute chain.

    Fixed link lengths ``L_j`` (buffer), joint angles ``θ_j`` (batch ``B``). For each
    link, the heading is ``Σ_{k≤j} θ_k``; joint positions are the cumulative sum of
    ``L_j [cos(heading), sin(heading)]``—fully **vectorized** over ``B`` and ``J``.

    Returns positions of the **end of each link** in the base frame ``(B, J, 2)``
    (last row is the end-effector in the plane). Typical fusion target for
    batched geometry in calibration, whole-body control, and VLA motion heads.
    """

    def __init__(self, num_joints: int, link_scale: float = 1.0):
        super().__init__()
        assert num_joints >= 1
        self.num_joints = num_joints
        lengths = torch.ones(num_joints, dtype=torch.float32) * float(link_scale)
        self.register_buffer("link_lengths", lengths)

    def forward(self, joint_angles: torch.Tensor) -> torch.Tensor:
        """
        Args:
            joint_angles: (B, J) radians, one revolute per link.
        Returns:
            positions: (B, J, 2) world-frame (x, y) at each link end; ``[..., -1, :]`` is EE.
        """
        B, J = joint_angles.shape
        assert J == self.num_joints

        heading = torch.cumsum(joint_angles, dim=-1)
        c = torch.cos(heading)
        s = torch.sin(heading)
        direction = torch.stack((c, s), dim=-1)
        displacements = direction * self.link_lengths.view(1, J, 1)
        return torch.cumsum(displacements, dim=1)


class ModelNew(nn.Module):
    """KernelBench candidate entry point; default delegates to ``Model``."""

    def __init__(self, num_joints: int, link_scale: float = 1.0):
        super().__init__()
        self._impl = Model(num_joints, link_scale=link_scale)

    def forward(self, joint_angles: torch.Tensor) -> torch.Tensor:
        return self._impl(joint_angles)


batch_size = 32
num_joints = 7
link_scale = 0.15


def get_inputs():
    p = popcorn_pri
    return [torch.empty(p.jitter_int(batch_size), num_joints).uniform_(-math.pi, math.pi)]


def get_init_inputs():
    return [num_joints, link_scale]
turn 0291.1s turn · 4.7s LLM · 286.4s tools · tokens in=2,376 · out=6 · total=2,675
reasoning
The task is to replace the forward computation of this PyTorch model with a custom CUDA kernel. The model computes batched forward kinematics for a planar serial revolute chain.
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 1226.5s turn · 6.6s LLM · 219.9s tools · tokens in=2,908 · out=596 · total=3,558
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\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 fk_kernel(const float* __restrict__ angles, float* __restrict__ positions, \n const float* __restrict__ lengths, int B, int J) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= B) return;\n \n float cum_heading = 0.0f;\n float cum_x = 0.0f;\n float cum_y = 0.0f;\n \n for (int j = 0; j < J; j++) {\n cum_heading += angles[b * J + j];\n float c = cosf(cum_heading);\n float s = sinf(cum_heading);\n float dx = c * lengths[j];\n float dy = s * lengths[j];\n cum_x += dx;\n cum_y += dy;\n positions[b * J * 2 + j * 2 + 0] = cum_x;\n positions[b * J * 2 + j * 2 + 1] = cum_y;\n }\n}\n\ntorch::Tensor forward_kinematics_cuda(torch::Tensor joint_angles, torch::Tensor link_lengths) {\n int B = joint_angles.size(0);\n int J = joint_angles.size(1);\n \n auto options = torch::TensorOptions().dtype(torch::kFloat32).device(joint_angles.device());\n auto positions = torch::empty({B, J, 2}, options);\n \n int threads = 256;\n int blocks = (B + threads - 1) / threads;\n \n fk_kernel<<<blocks, threads>>>(\n joint_angles.data_ptr<float>(),\n positions.data_ptr<float>(),\n link_lengths.data_ptr<float>(),\n B, J\n );\n \n return positions;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor forward_kinematics_cuda(torch::Tensor joint_angles, torch::Tensor link_lengths);\n\"\"\"\n\nmodule = load_inline(\n name=\"fk_cuda\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"forward_kinematics_cuda\"],\n verbose=True,\n extra_cflags=[\"-O3\"]\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, num_joints: int, link_scale: float = 1.0):\n super().__init__()\n self.num_joints = num_joints\n lengths = torch.ones(num_joints, dtype=torch.float32) * float(link_scale)\n self.register_buffer(\"link_lengths\", lengths)\n\n def forward(self, joint_angles: torch.Tensor) -> torch.Tensor:\n return module.forward_kinematics_cuda(joint_angles, self.link_lengths)\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 2518.3s turn · 4.3s LLM · 514.0s tools · tokens in=3,575 · out=597 · total=4,179
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\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 fk_kernel(const float* __restrict__ angles, float* __restrict__ positions, \n const float* __restrict__ lengths, int B, int J) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= B) return;\n \n float cum_heading = 0.0f;\n float cum_x = 0.0f;\n float cum_y = 0.0f;\n \n for (int j = 0; j < J; j++) {\n cum_heading += angles[b * J + j];\n float c = cosf(cum_heading);\n float s = sinf(cum_heading);\n float dx = c * lengths[j];\n float dy = s * lengths[j];\n cum_x += dx;\n cum_y += dy;\n positions[b * J * 2 + j * 2 + 0] = cum_x;\n positions[b * J * 2 + j * 2 + 1] = cum_y;\n }\n}\n\ntorch::Tensor forward_kinematics_cuda(torch::Tensor joint_angles, torch::Tensor link_lengths) {\n int B = joint_angles.size(0);\n int J = joint_angles.size(1);\n \n auto options = torch::TensorOptions().dtype(torch::kFloat32).device(joint_angles.device());\n auto positions = torch::empty({B, J, 2}, options);\n \n int threads = 256;\n int blocks = (B + threads - 1) / threads;\n \n fk_kernel<<<blocks, threads>>>(\n joint_angles.data_ptr<float>(),\n positions.data_ptr<float>(),\n link_lengths.data_ptr<float>(),\n B, J\n );\n \n return positions;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor forward_kinematics_cuda(torch::Tensor joint_angles, torch::Tensor link_lengths);\n\"\"\"\n\nmodule = load_inline(\n name=\"fk_cuda\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"forward_kinematics_cuda\"],\n verbose=False,\n extra_cflags=[\"-O3\"]\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, num_joints: int, link_scale: float = 1.0):\n super().__init__()\n self.num_joints = num_joints\n lengths = torch.ones(num_joints, dtype=torch.float32) * float(link_scale)\n self.register_buffer(\"link_lengths\", lengths)\n\n def forward(self, joint_angles: torch.Tensor) -> torch.Tensor:\n return module.forward_kinematics_cuda(joint_angles, self.link_lengths)\n" }
result ← run_correctness (ok)
okrun_correctness run_correctness PASSED: (3 / 3) trials all matched the reference. Numerical precision: max_abs_err=2.64e-07 mean_abs_err=1.32e-08 max_rel_err=1.22e-04
turn 3177.0s turn · 4.5s LLM · 172.5s tools · tokens in=4,239 · out=596 · total=4,873 · FINAL
reasoning
Correct! Now, to check if faster, I should submit, but let's profile to see.
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 = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <math.h>\n\n__global__ void fk_kernel(const float* __restrict__ angles, float* __restrict__ positions, \n const float* __restrict__ lengths, int B, int J) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= B) return;\n \n float cum_heading = 0.0f;\n float cum_x = 0.0f;\n float cum_y = 0.0f;\n \n for (int j = 0; j < J; j++) {\n cum_heading += angles[b * J + j];\n float c = cosf(cum_heading);\n float s = sinf(cum_heading);\n float dx = c * lengths[j];\n float dy = s * lengths[j];\n cum_x += dx;\n cum_y += dy;\n positions[b * J * 2 + j * 2 + 0] = cum_x;\n positions[b * J * 2 + j * 2 + 1] = cum_y;\n }\n}\n\ntorch::Tensor forward_kinematics_cuda(torch::Tensor joint_angles, torch::Tensor link_lengths) {\n int B = joint_angles.size(0);\n int J = joint_angles.size(1);\n \n auto options = torch::TensorOptions().dtype(torch::kFloat32).device(joint_angles.device());\n auto positions = torch::empty({B, J, 2}, options);\n \n int threads = 256;\n int blocks = (B + threads - 1) / threads;\n \n fk_kernel<<<blocks, threads>>>(\n joint_angles.data_ptr<float>(),\n positions.data_ptr<float>(),\n link_lengths.data_ptr<float>(),\n B, J\n );\n \n return positions;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor forward_kinematics_cuda(torch::Tensor joint_angles, torch::Tensor link_lengths);\n\"\"\"\n\nmodule = load_inline(\n name=\"fk_cuda\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"forward_kinematics_cuda\"],\n verbose=False,\n extra_cflags=[\"-O3\"]\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, num_joints: int, link_scale: float = 1.0):\n super().__init__()\n self.num_joints = num_joints\n lengths = torch.ones(num_joints, dtype=torch.float32) * float(link_scale)\n self.register_buffer(\"link_lengths\", lengths)\n\n def forward(self, joint_angles: torch.Tensor) -> torch.Tensor:\n return module.forward_kinematics_cuda(joint_angles, self.link_lengths)\n" }
result ← submit_kernel (ok)
oksubmit_kernel submit_kernel PASSED: (3 / 3) correctness trials all passed. Kernel runtime: 0.01 μs Runtime stats: mean=0.01μs median=?μs std=0.00μs Numerical precision: max_abs_err=2.64e-07 mean_abs_err=1.32e-08 max_rel_err=1.22e-04 Memory: 33.0 MB peak (ref: 33.01 MB, ratio: 0.9998x) Kernel launches: 1 (ref: 14, fusion ratio: 14.0) SOL score: 0.0005
final submitted kernel (76 lines, 2,252 chars) — full source, untruncated
import torch
import torch.nn as nn
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 fk_kernel(const float* __restrict__ angles, float* __restrict__ positions, 
                          const float* __restrict__ lengths, int B, int J) {
    int b = blockIdx.x * blockDim.x + threadIdx.x;
    if (b >= B) return;
    
    float cum_heading = 0.0f;
    float cum_x = 0.0f;
    float cum_y = 0.0f;
    
    for (int j = 0; j < J; j++) {
        cum_heading += angles[b * J + j];
        float c = cosf(cum_heading);
        float s = sinf(cum_heading);
        float dx = c * lengths[j];
        float dy = s * lengths[j];
        cum_x += dx;
        cum_y += dy;
        positions[b * J * 2 + j * 2 + 0] = cum_x;
        positions[b * J * 2 + j * 2 + 1] = cum_y;
    }
}

torch::Tensor forward_kinematics_cuda(torch::Tensor joint_angles, torch::Tensor link_lengths) {
    int B = joint_angles.size(0);
    int J = joint_angles.size(1);
    
    auto options = torch::TensorOptions().dtype(torch::kFloat32).device(joint_angles.device());
    auto positions = torch::empty({B, J, 2}, options);
    
    int threads = 256;
    int blocks = (B + threads - 1) / threads;
    
    fk_kernel<<<blocks, threads>>>(
        joint_angles.data_ptr<float>(),
        positions.data_ptr<float>(),
        link_lengths.data_ptr<float>(),
        B, J
    );
    
    return positions;
}
"""

cpp_source = """
torch::Tensor forward_kinematics_cuda(torch::Tensor joint_angles, torch::Tensor link_lengths);
"""

module = load_inline(
    name="fk_cuda",
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=["forward_kinematics_cuda"],
    verbose=False,
    extra_cflags=["-O3"]
)

class ModelNew(nn.Module):
    def __init__(self, num_joints: int, link_scale: float = 1.0):
        super().__init__()
        self.num_joints = num_joints
        lengths = torch.ones(num_joints, dtype=torch.float32) * float(link_scale)
        self.register_buffer("link_lengths", lengths)

    def forward(self, joint_angles: torch.Tensor) -> torch.Tensor:
        return module.forward_kinematics_cuda(joint_angles, self.link_lengths)