correct
42_BatchedForwardKinematics.py
started
2026-05-11T10:12:49Z
timingkernelbench.timing
kernel mean / median / std
0.01 / — / 1.45e-04 μs
ref mean / median / std
0.02 / — / 6.29e-04 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
14.00x
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
warp stalls (top 5)nsight
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 0782.1s turn · 9.4s LLM · 772.8s tools · tokens in=1,478 · out=583 · total=2,331 · FINAL
reasoning
The task is to replace the forward computation of a PyTorch model with a custom CUDA kernel. The model computes batched forward kinematics for a planar serial revolute chain.
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 = r'''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <math.h>\n\n__global__ void fk_kernel(\n const float* __restrict__ joint_angles,\n const float* __restrict__ link_lengths,\n float* __restrict__ positions,\n int B, int J\n) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= B) return;\n\n float heading = 0.0f;\n float px = 0.0f, py = 0.0f;\n for (int j = 0; j < J; ++j) {\n heading += joint_angles[b * J + j];\n float c = cosf(heading);\n float s = sinf(heading);\n float dx = c * link_lengths[j];\n float dy = s * link_lengths[j];\n px += dx;\n py += dy;\n positions[b * J * 2 + j * 2 + 0] = px;\n positions[b * J * 2 + j * 2 + 1] = py;\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 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 fk_kernel<<<blocks, threads>>>(\n joint_angles.data_ptr<float>(),\n link_lengths.data_ptr<float>(),\n positions.data_ptr<float>(),\n B, J\n );\n return positions;\n}\n'''\n\ncpp_source = r'''\ntorch::Tensor forward_kinematics_cuda(torch::Tensor joint_angles, torch::Tensor link_lengths);\n'''\n\nfk_module = 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)\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 fk_module.forward_kinematics_cuda(joint_angles, self.link_lengths)\n"
}
result ← submit_kernel (ok)
oksubmit_kernel
submit_kernel PASSED: (5 / 5) 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.33e-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 (74 lines, 2,180 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline
cuda_source = r'''
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <math.h>
__global__ void fk_kernel(
const float* __restrict__ joint_angles,
const float* __restrict__ link_lengths,
float* __restrict__ positions,
int B, int J
) {
int b = blockIdx.x * blockDim.x + threadIdx.x;
if (b >= B) return;
float heading = 0.0f;
float px = 0.0f, py = 0.0f;
for (int j = 0; j < J; ++j) {
heading += joint_angles[b * J + j];
float c = cosf(heading);
float s = sinf(heading);
float dx = c * link_lengths[j];
float dy = s * link_lengths[j];
px += dx;
py += dy;
positions[b * J * 2 + j * 2 + 0] = px;
positions[b * J * 2 + j * 2 + 1] = py;
}
}
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>(),
link_lengths.data_ptr<float>(),
positions.data_ptr<float>(),
B, J
);
return positions;
}
'''
cpp_source = r'''
torch::Tensor forward_kinematics_cuda(torch::Tensor joint_angles, torch::Tensor link_lengths);
'''
fk_module = load_inline(
name='fk_cuda',
cpp_sources=cpp_source,
cuda_sources=cuda_source,
functions=['forward_kinematics_cuda'],
verbose=False
)
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 fk_module.forward_kinematics_cuda(joint_angles, self.link_lengths)