incorrect
20_RotaryPositionEmbeddingBio.py
started
2026-05-11T09:50:04Z
correctness issue: Output mismatch
timingkernelbench.timing
kernel mean / median / std
—
ref mean / median / std
0.07 / — / 9.64e-04 μs
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
dram bandwidth
971.59 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):
"""
Rotary Position Embedding (RoPE) applied to protein residue sequences.
Used in ESM-2 and similar protein language models. Encodes absolute
position via rotation of query/key pairs in 2-D subspaces, preserving
relative-position information in dot-product attention.
"""
def __init__(self, dim, max_seq_len=2048, base=10000.0):
super().__init__()
inv_freq = 1.0 / (base ** (torch.arange(0, dim, 2).float() / dim))
self.register_buffer("inv_freq", inv_freq)
self.max_seq_len = max_seq_len
self.dim = dim
def _build_cos_sin(self, seq_len: int, device: torch.device):
t = torch.arange(seq_len, device=device, dtype=self.inv_freq.dtype)
freqs = torch.outer(t, self.inv_freq)
emb = torch.cat([freqs, freqs], dim=-1)
return emb.cos(), emb.sin()
@staticmethod
def _rotate_half(x: torch.Tensor) -> torch.Tensor:
x1, x2 = x.chunk(2, dim=-1)
return torch.cat([-x2, x1], dim=-1)
def forward(self, q: torch.Tensor, k: torch.Tensor) -> tuple:
"""
Args:
q: (B, H, N, D) – queries
k: (B, H, N, D) – keys
Returns:
(q_rot, k_rot) each (B, H, N, D) with RoPE applied
"""
N = q.shape[2]
cos, sin = self._build_cos_sin(N, q.device) # (N, D)
cos = cos.unsqueeze(0).unsqueeze(0)
sin = sin.unsqueeze(0).unsqueeze(0)
q_rot = q * cos + self._rotate_half(q) * sin
k_rot = k * cos + self._rotate_half(k) * sin
return q_rot, k_rot
dim = 64
num_heads = 8
seq_len = 256
batch_size = 4
def get_inputs():
p = popcorn_pri
bs = p.jitter_int(batch_size)
nh = p.jitter_int(num_heads)
sl = p.jitter_int(seq_len)
q = torch.randn(bs, nh, sl, dim)
k = torch.randn(bs, nh, sl, dim)
return [q, k]
def get_init_inputs():
return [dim]
turn 0320.4s turn · 15.3s LLM · 305.1s tools · tokens in=1,405 · out=842 · total=2,862 · FINAL
reasoning
The task is to replace the forward computation of a PyTorch model with a custom CUDA kernel for Rotary Position Embedding (RoPE).
I need to create a custom CUDA kernel for Rotary Position Embedding (RoPE) to speed up the forward computation in a PyTorch model. The task involves replacing the standard PyTorch RoPE implementation with a faster version using a custom kernel, compiled via `torch.utils.cpp_extension.load_inline`.
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 rope_kernel(\n const float* __restrict__ q,\n const float* __restrict__ k,\n const float* __restrict__ inv_freq,\n float* __restrict__ q_out,\n float* __restrict__ k_out,\n int batch_size,\n int num_heads,\n int seq_len,\n int dim,\n int dim_half\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = batch_size * num_heads * seq_len * dim;\n if (idx >= total) return;\n\n int d = idx % dim;\n int n = (idx / dim) % seq_len;\n int h = (idx / (dim * seq_len)) % num_heads;\n int b = idx / (dim * seq_len * num_heads);\n\n float freq = inv_freq[d / 2];\n float angle = n * freq;\n float cos_val = cosf(angle);\n float sin_val = sinf(angle);\n\n int offset = ((b * num_heads + h) * seq_len + n) * dim + d;\n float q_val = q[offset];\n float k_val = k[offset];\n\n if (d % 2 == 0) {\n int pair_offset = offset + 1;\n float q_pair = q[pair_offset];\n float k_pair = k[pair_offset];\n q_out[offset] = q_val * cos_val - q_pair * sin_val;\n k_out[offset] = k_val * cos_val - k_pair * sin_val;\n } else {\n int pair_offset = offset - 1;\n float q_pair = q[pair_offset];\n float k_pair = k[pair_offset];\n q_out[offset] = q_val * cos_val + q_pair * sin_val;\n k_out[offset] = k_val * cos_val + k_pair * sin_val;\n }\n}\n\nstd::vector<torch::Tensor> rope_cuda(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor inv_freq\n) {\n int batch_size = q.size(0);\n int num_heads = q.size(1);\n int seq_len = q.size(2);\n int dim = q.size(3);\n int dim_half = dim / 2;\n\n auto q_out = torch::empty_like(q);\n auto k_out = torch::empty_like(k);\n\n int total = batch_size * num_heads * seq_len * dim;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n\n rope_kernel<<<blocks, threads>>>(\n q.data_ptr<float>(),\n k.data_ptr<float>(),\n inv_freq.data_ptr<float>(),\n q_out.data_ptr<float>(),\n k_out.data_ptr<float>(),\n batch_size, num_heads, seq_len, dim, dim_half\n );\n\n return {q_out, k_out};\n}\n'''\n\ncpp_source = '''\nstd::vector<torch::Tensor> rope_cuda(torch::Tensor q, torch::Tensor k, torch::Tensor inv_freq);\n'''\n\nmodule = load_inline(\n name='rope_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['rope_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim, max_seq_len=2048, base=10000.0):\n super().__init__()\n inv_freq = 1.0 / (base ** (torch.arange(0, dim, 2).float() / dim))\n self.register_buffer(\"inv_freq\", inv_freq)\n self.dim = dim\n\n def forward(self, q: torch.Tensor, k: torch.Tensor) -> tuple:\n return module.rope_cuda(q, k, self.inv_freq)\n"
}
result ← submit_kernel (fail)
failsubmit_kernel
submit_kernel FAILED: correctness check did not pass ((0 / 5) trials).
correctness_issue: Output mismatch
max_difference: ['8.537722', '8.351884', '8.241857', '11.210609', '8.204172', '8.851664', '9.117870', '8.540900', '8.596428', '8.158693']
avg_difference: ['0.852953', '0.856102', '0.855250', '0.853945', '0.856452', '0.853398', '0.854333', '0.853263', '0.853908', '0.853664']
final submitted kernel (108 lines, 2,966 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 rope_kernel(
const float* __restrict__ q,
const float* __restrict__ k,
const float* __restrict__ inv_freq,
float* __restrict__ q_out,
float* __restrict__ k_out,
int batch_size,
int num_heads,
int seq_len,
int dim,
int dim_half
) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int total = batch_size * num_heads * seq_len * dim;
if (idx >= total) return;
int d = idx % dim;
int n = (idx / dim) % seq_len;
int h = (idx / (dim * seq_len)) % num_heads;
int b = idx / (dim * seq_len * num_heads);
float freq = inv_freq[d / 2];
float angle = n * freq;
float cos_val = cosf(angle);
float sin_val = sinf(angle);
int offset = ((b * num_heads + h) * seq_len + n) * dim + d;
float q_val = q[offset];
float k_val = k[offset];
if (d % 2 == 0) {
int pair_offset = offset + 1;
float q_pair = q[pair_offset];
float k_pair = k[pair_offset];
q_out[offset] = q_val * cos_val - q_pair * sin_val;
k_out[offset] = k_val * cos_val - k_pair * sin_val;
} else {
int pair_offset = offset - 1;
float q_pair = q[pair_offset];
float k_pair = k[pair_offset];
q_out[offset] = q_val * cos_val + q_pair * sin_val;
k_out[offset] = k_val * cos_val + k_pair * sin_val;
}
}
std::vector<torch::Tensor> rope_cuda(
torch::Tensor q,
torch::Tensor k,
torch::Tensor inv_freq
) {
int batch_size = q.size(0);
int num_heads = q.size(1);
int seq_len = q.size(2);
int dim = q.size(3);
int dim_half = dim / 2;
auto q_out = torch::empty_like(q);
auto k_out = torch::empty_like(k);
int total = batch_size * num_heads * seq_len * dim;
int threads = 256;
int blocks = (total + threads - 1) / threads;
rope_kernel<<<blocks, threads>>>(
q.data_ptr<float>(),
k.data_ptr<float>(),
inv_freq.data_ptr<float>(),
q_out.data_ptr<float>(),
k_out.data_ptr<float>(),
batch_size, num_heads, seq_len, dim, dim_half
);
return {q_out, k_out};
}
'''
cpp_source = '''
std::vector<torch::Tensor> rope_cuda(torch::Tensor q, torch::Tensor k, torch::Tensor inv_freq);
'''
module = load_inline(
name='rope_module',
cpp_sources=[cpp_source],
cuda_sources=[cuda_source],
functions=['rope_cuda'],
verbose=False
)
class ModelNew(nn.Module):
def __init__(self, dim, max_seq_len=2048, base=10000.0):
super().__init__()
inv_freq = 1.0 / (base ** (torch.arange(0, dim, 2).float() / dim))
self.register_buffer("inv_freq", inv_freq)
self.dim = dim
def forward(self, q: torch.Tensor, k: torch.Tensor) -> tuple:
return module.rope_cuda(q, k, self.inv_freq)