correct
20_RotaryPositionEmbeddingBio.py
started
2026-05-11T21:14:16Z
timingkernelbench.timing
kernel mean / median / std
0.03 / — / 6.97e-04 μs
ref mean / median / std
0.07 / — / 0.01 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
2.50x
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
dram bandwidth
691.76 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 022.0s turn · 22.0s LLM · 0.0s tools · tokens in=2,303 · out=1,047 · total=4,270
reasoning
The task is to replace the RoPE (Rotary Position Embedding) forward computation with a custom CUDA kernel that's faster than PyTorch's reference.
I need to replace the forward computation of the PyTorch model with a custom CUDA kernel for the RoPE operation. The ModelNew class should take the same inputs as the original Model, have the same forward signature, produce identical outputs numerically, and be faster using a custom CUDA kernel.
The task focuses on accelerating the Ro Rotary Position Embedding (RoPE) part, specifically rotating the query and key vectors using precomputed cosine and sine waves based on position and frequency.
assistant```python
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 rope_kernel(
const float* __restrict__ q,
const float* __restrict__ k,
float* __restrict__ q_out,
float* __restrict__ k_out,
const float* __restrict__ cos,
const float* __restrict__ sin,
const int B,
const int H,
const int N,
const int D
) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int total = B * H * N * D;
if (idx >= total) return;
int d = idx % D;
int n = (idx / D) % N;
int h = (idx / (D * N)) % H;
int b = idx / (D * N * H);
float c = cos[n * D + d];
float s = sin[n * D + d];
int half = D / 2;
float val_q, val_k, rot_q, rot_k;
if (d < half) {
val_q = q[idx];
val_k = k[idx];
int d2 = d + half;
rot_q = -q[b * H * N * D + h * N * D + n * D + d2];
rot_k = -k[b * H * N * D + h * N * D + n * D + d2];
q_out[idx] = val_q * c + rot_q * s;
k_out[idx] = val_k * c + rot_k * s;
} else {
val_q = q[idx];
val_k = k[idx];
int d2 = d - half;
rot_q = q[b * H * N * D + h * N * D + n * D + d2];
rot_k = k[b * H * N * D + h * N * D + n * D + d2];
q_out[idx] = val_q * c + rot_q * s;
k_out[idx] = val_k * c + rot_k * s;
}
}
std::vector<torch::Tensor> rope_cuda(
torch::Tensor q,
torch::Tensor k,
torch::Tensor inv_freq,
int max_seq_len
) {
auto B = q.size(0);
auto H = q.size(1);
auto N = q.size(2);
auto D = q.size(3);
auto options = torch::TensorOptions().dtype(q.dtype()).device(q.device());
auto cos = torch::empty({N, D}, options);
auto sin = torch::empty({N, D}, options);
// Build freqs and cos/sin on GPU
auto t = torch::arange(N, options);
auto freqs = t.unsqueeze(1) * inv_freq.unsqueeze(0); // (N, D/2)
auto emb = torch::cat({freqs, freqs}, 1); // (N, D)
cos = torch::cos(emb);
sin = torch::sin(emb);
auto q_out = torch::empty_like(q);
auto k_out = torch::empty_like(k);
int total = B * H * N * D;
int threads = 256;
int blocks = (total + threads - 1) / threads;
rope_kernel<<<blocks, threads>>>(
q.data_ptr<float>(),
k.data_ptr<float>(),
q_out.data_ptr<float>(),
k_out.data_ptr<float>(),
cos.data_ptr<float>(),
sin.data_ptr<float>(),
B, H, N, D
);
return {q_out, k_out};
}
'''
cpp_source = r'''
std::vector<torch::Tensor> rope_cuda(
torch::Tensor q,
torch::Tensor k,
torch::Tensor inv_freq,
int max_seq_len
);
'''
rope_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.max_seq_len = max_seq_len
self.dim = dim
def forward(self, q: torch.Tensor, k: torch.Tensor) -> tuple:
outs = rope_module.rope_cuda(q, k, self.inv_freq, self.max_seq_len)
return outs[0], outs[1]
```
turn 1316.3s turn · 11.0s LLM · 305.4s tools · tokens in=4,338 · out=1,054 · total=5,546
reasoning
The instruction is: "You must call at least one tool this turn: use `compile_kernel` / `run_correctness` to iterate, or call `submit_kernel` with your best `ModelNew` implementation to finalize the ru...
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 = r'''\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 float* __restrict__ q_out,\n float* __restrict__ k_out,\n const float* __restrict__ cos,\n const float* __restrict__ sin,\n const int B,\n const int H,\n const int N,\n const int D\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * H * N * D;\n if (idx >= total) return;\n\n int d = idx % D;\n int n = (idx / D) % N;\n int h = (idx / (D * N)) % H;\n int b = idx / (D * N * H);\n\n float c = cos[n * D + d];\n float s = sin[n * D + d];\n\n int half = D / 2;\n float val_q, val_k, rot_q, rot_k;\n\n if (d < half) {\n val_q = q[idx];\n val_k = k[idx];\n int d2 = d + half;\n rot_q = -q[b * H * N * D + h * N * D + n * D + d2];\n rot_k = -k[b * H * N * D + h * N * D + n * D + d2];\n q_out[idx] = val_q * c + rot_q * s;\n k_out[idx] = val_k * c + rot_k * s;\n } else {\n val_q = q[idx];\n val_k = k[idx];\n int d2 = d - half;\n rot_q = q[b * H * N * D + h * N * D + n * D + d2];\n rot_k = k[b * H * N * D + h * N * D + n * D + d2];\n q_out[idx] = val_q * c + rot_q * s;\n k_out[idx] = val_k * c + rot_k * s;\n }\n}\n\nstd::vector<torch::Tensor> rope_cuda(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor inv_freq,\n int max_seq_len\n) {\n auto B = q.size(0);\n auto H = q.size(1);\n auto N = q.size(2);\n auto D = q.size(3);\n\n auto options = torch::TensorOptions().dtype(q.dtype()).device(q.device());\n auto cos = torch::empty({N, D}, options);\n auto sin = torch::empty({N, D}, options);\n\n // Build freqs and cos/sin on GPU\n auto t = torch::arange(N, options);\n auto freqs = t.unsqueeze(1) * inv_freq.unsqueeze(0); // (N, D/2)\n auto emb = torch::cat({freqs, freqs}, 1); // (N, D)\n cos = torch::cos(emb);\n sin = torch::sin(emb);\n\n auto q_out = torch::empty_like(q);\n auto k_out = torch::empty_like(k);\n\n int total = B * H * N * D;\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 q_out.data_ptr<float>(),\n k_out.data_ptr<float>(),\n cos.data_ptr<float>(),\n sin.data_ptr<float>(),\n B, H, N, D\n );\n\n return {q_out, k_out};\n}\n'''\n\ncpp_source = r'''\nstd::vector<torch::Tensor> rope_cuda(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor inv_freq,\n int max_seq_len\n);\n'''\n\nrope_module = 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.max_seq_len = max_seq_len\n self.dim = dim\n\n def forward(self, q: torch.Tensor, k: torch.Tensor) -> tuple:\n outs = rope_module.rope_cuda(q, k, self.inv_freq, self.max_seq_len)\n return outs[0], outs[1]\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 2169.1s turn · 11.5s LLM · 157.6s tools · tokens in=5,563 · out=1,055 · total=6,633
reasoning
The kernel compiled successfully. Now I should check correctness with run_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 = r'''\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 float* __restrict__ q_out,\n float* __restrict__ k_out,\n const float* __restrict__ cos,\n const float* __restrict__ sin,\n const int B,\n const int H,\n const int N,\n const int D\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * H * N * D;\n if (idx >= total) return;\n\n int d = idx % D;\n int n = (idx / D) % N;\n int h = (idx / (D * N)) % H;\n int b = idx / (D * N * H);\n\n float c = cos[n * D + d];\n float s = sin[n * D + d];\n\n int half = D / 2;\n float val_q, val_k, rot_q, rot_k;\n\n if (d < half) {\n val_q = q[idx];\n val_k = k[idx];\n int d2 = d + half;\n rot_q = -q[b * H * N * D + h * N * D + n * D + d2];\n rot_k = -k[b * H * N * D + h * N * D + n * D + d2];\n q_out[idx] = val_q * c + rot_q * s;\n k_out[idx] = val_k * c + rot_k * s;\n } else {\n val_q = q[idx];\n val_k = k[idx];\n int d2 = d - half;\n rot_q = q[b * H * N * D + h * N * D + n * D + d2];\n rot_k = k[b * H * N * D + h * N * D + n * D + d2];\n q_out[idx] = val_q * c + rot_q * s;\n k_out[idx] = val_k * c + rot_k * s;\n }\n}\n\nstd::vector<torch::Tensor> rope_cuda(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor inv_freq,\n int max_seq_len\n) {\n auto B = q.size(0);\n auto H = q.size(1);\n auto N = q.size(2);\n auto D = q.size(3);\n\n auto options = torch::TensorOptions().dtype(q.dtype()).device(q.device());\n auto cos = torch::empty({N, D}, options);\n auto sin = torch::empty({N, D}, options);\n\n // Build freqs and cos/sin on GPU\n auto t = torch::arange(N, options);\n auto freqs = t.unsqueeze(1) * inv_freq.unsqueeze(0); // (N, D/2)\n auto emb = torch::cat({freqs, freqs}, 1); // (N, D)\n cos = torch::cos(emb);\n sin = torch::sin(emb);\n\n auto q_out = torch::empty_like(q);\n auto k_out = torch::empty_like(k);\n\n int total = B * H * N * D;\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 q_out.data_ptr<float>(),\n k_out.data_ptr<float>(),\n cos.data_ptr<float>(),\n sin.data_ptr<float>(),\n B, H, N, D\n );\n\n return {q_out, k_out};\n}\n'''\n\ncpp_source = r'''\nstd::vector<torch::Tensor> rope_cuda(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor inv_freq,\n int max_seq_len\n);\n'''\n\nrope_module = 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.max_seq_len = max_seq_len\n self.dim = dim\n\n def forward(self, q: torch.Tensor, k: torch.Tensor) -> tuple:\n outs = rope_module.rope_cuda(q, k, self.inv_freq, self.max_seq_len)\n return outs[0], outs[1]\n"
}
result ← run_correctness (ok)
okrun_correctness
run_correctness PASSED: (3 / 3) trials all matched the reference.
Numerical precision: max_abs_err=4.77e-07 mean_abs_err=1.43e-08 max_rel_err=8.85e-03
turn 338.9s turn · 10.8s LLM · 28.1s tools · tokens in=6,693 · out=1,054 · total=7,761 · FINAL
reasoning
Correctness passed with very small errors. Now I can submit it.
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 rope_kernel(\n const float* __restrict__ q,\n const float* __restrict__ k,\n float* __restrict__ q_out,\n float* __restrict__ k_out,\n const float* __restrict__ cos,\n const float* __restrict__ sin,\n const int B,\n const int H,\n const int N,\n const int D\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * H * N * D;\n if (idx >= total) return;\n\n int d = idx % D;\n int n = (idx / D) % N;\n int h = (idx / (D * N)) % H;\n int b = idx / (D * N * H);\n\n float c = cos[n * D + d];\n float s = sin[n * D + d];\n\n int half = D / 2;\n float val_q, val_k, rot_q, rot_k;\n\n if (d < half) {\n val_q = q[idx];\n val_k = k[idx];\n int d2 = d + half;\n rot_q = -q[b * H * N * D + h * N * D + n * D + d2];\n rot_k = -k[b * H * N * D + h * N * D + n * D + d2];\n q_out[idx] = val_q * c + rot_q * s;\n k_out[idx] = val_k * c + rot_k * s;\n } else {\n val_q = q[idx];\n val_k = k[idx];\n int d2 = d - half;\n rot_q = q[b * H * N * D + h * N * D + n * D + d2];\n rot_k = k[b * H * N * D + h * N * D + n * D + d2];\n q_out[idx] = val_q * c + rot_q * s;\n k_out[idx] = val_k * c + rot_k * s;\n }\n}\n\nstd::vector<torch::Tensor> rope_cuda(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor inv_freq,\n int max_seq_len\n) {\n auto B = q.size(0);\n auto H = q.size(1);\n auto N = q.size(2);\n auto D = q.size(3);\n\n auto options = torch::TensorOptions().dtype(q.dtype()).device(q.device());\n auto cos = torch::empty({N, D}, options);\n auto sin = torch::empty({N, D}, options);\n\n // Build freqs and cos/sin on GPU\n auto t = torch::arange(N, options);\n auto freqs = t.unsqueeze(1) * inv_freq.unsqueeze(0); // (N, D/2)\n auto emb = torch::cat({freqs, freqs}, 1); // (N, D)\n cos = torch::cos(emb);\n sin = torch::sin(emb);\n\n auto q_out = torch::empty_like(q);\n auto k_out = torch::empty_like(k);\n\n int total = B * H * N * D;\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 q_out.data_ptr<float>(),\n k_out.data_ptr<float>(),\n cos.data_ptr<float>(),\n sin.data_ptr<float>(),\n B, H, N, D\n );\n\n return {q_out, k_out};\n}\n'''\n\ncpp_source = r'''\nstd::vector<torch::Tensor> rope_cuda(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor inv_freq,\n int max_seq_len\n);\n'''\n\nrope_module = 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.max_seq_len = max_seq_len\n self.dim = dim\n\n def forward(self, q: torch.Tensor, k: torch.Tensor) -> tuple:\n outs = rope_module.rope_cuda(q, k, self.inv_freq, self.max_seq_len)\n return outs[0], outs[1]\n"
}
result ← submit_kernel (ok)
oksubmit_kernel
submit_kernel PASSED: (3 / 3) correctness trials all passed.
Kernel runtime: 0.03 μs
Runtime stats: mean=0.03μs median=?μs std=0.00μs
Numerical precision: max_abs_err=4.77e-07 mean_abs_err=1.43e-08 max_rel_err=8.85e-03
Memory: 41.22 MB peak (ref: 45.13 MB, ratio: 0.9135x)
Kernel launches: 18 (ref: 45, fusion ratio: 2.5)
SOL score: 0.3712
final submitted kernel (128 lines, 3,372 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 rope_kernel(
const float* __restrict__ q,
const float* __restrict__ k,
float* __restrict__ q_out,
float* __restrict__ k_out,
const float* __restrict__ cos,
const float* __restrict__ sin,
const int B,
const int H,
const int N,
const int D
) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int total = B * H * N * D;
if (idx >= total) return;
int d = idx % D;
int n = (idx / D) % N;
int h = (idx / (D * N)) % H;
int b = idx / (D * N * H);
float c = cos[n * D + d];
float s = sin[n * D + d];
int half = D / 2;
float val_q, val_k, rot_q, rot_k;
if (d < half) {
val_q = q[idx];
val_k = k[idx];
int d2 = d + half;
rot_q = -q[b * H * N * D + h * N * D + n * D + d2];
rot_k = -k[b * H * N * D + h * N * D + n * D + d2];
q_out[idx] = val_q * c + rot_q * s;
k_out[idx] = val_k * c + rot_k * s;
} else {
val_q = q[idx];
val_k = k[idx];
int d2 = d - half;
rot_q = q[b * H * N * D + h * N * D + n * D + d2];
rot_k = k[b * H * N * D + h * N * D + n * D + d2];
q_out[idx] = val_q * c + rot_q * s;
k_out[idx] = val_k * c + rot_k * s;
}
}
std::vector<torch::Tensor> rope_cuda(
torch::Tensor q,
torch::Tensor k,
torch::Tensor inv_freq,
int max_seq_len
) {
auto B = q.size(0);
auto H = q.size(1);
auto N = q.size(2);
auto D = q.size(3);
auto options = torch::TensorOptions().dtype(q.dtype()).device(q.device());
auto cos = torch::empty({N, D}, options);
auto sin = torch::empty({N, D}, options);
// Build freqs and cos/sin on GPU
auto t = torch::arange(N, options);
auto freqs = t.unsqueeze(1) * inv_freq.unsqueeze(0); // (N, D/2)
auto emb = torch::cat({freqs, freqs}, 1); // (N, D)
cos = torch::cos(emb);
sin = torch::sin(emb);
auto q_out = torch::empty_like(q);
auto k_out = torch::empty_like(k);
int total = B * H * N * D;
int threads = 256;
int blocks = (total + threads - 1) / threads;
rope_kernel<<<blocks, threads>>>(
q.data_ptr<float>(),
k.data_ptr<float>(),
q_out.data_ptr<float>(),
k_out.data_ptr<float>(),
cos.data_ptr<float>(),
sin.data_ptr<float>(),
B, H, N, D
);
return {q_out, k_out};
}
'''
cpp_source = r'''
std::vector<torch::Tensor> rope_cuda(
torch::Tensor q,
torch::Tensor k,
torch::Tensor inv_freq,
int max_seq_len
);
'''
rope_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.max_seq_len = max_seq_len
self.dim = dim
def forward(self, q: torch.Tensor, k: torch.Tensor) -> tuple:
outs = rope_module.rope_cuda(q, k, self.inv_freq, self.max_seq_len)
return outs[0], outs[1]