correct
20_RotaryPositionEmbeddingBio.py
started
2026-05-11T13:36:39Z
timingkernelbench.timing
kernel mean / median / std
0.02 / — / 9.65e-04 μs
ref mean / median / std
0.07 / — / 0.00 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
3.00x
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
dram bandwidth
1327.64 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 0662.1s turn · 3.9s LLM · 658.2s tools · tokens in=1,772 · out=6 · total=2,000
reasoning
The task is to replace the PyTorch Model with a custom CUDA kernel for RoPE (Rotary Position Embedding) applied to query and key tensors.
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 1260.2s turn · 13.6s LLM · 246.6s tools · tokens in=2,233 · out=950 · total=3,279
reasoning
H100 specs: high bandwidth, good for this memory-bound op probably.
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\n__global__ void rope_kernel(\n const float* __restrict__ q,\n const float* __restrict__ k,\n float* __restrict__ q_rot,\n float* __restrict__ k_rot,\n const float* __restrict__ cos,\n const float* __restrict__ sin,\n int B, int H, int N, 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 int half_d = D / 2;\n float c = cos[n * D + d];\n float s = sin[n * D + d];\n\n float val_q = q[idx];\n float val_k = k[idx];\n\n if (d < half_d) {\n // q1 part: q_rot = q1 * c - q2 * s ? Wait, let's check\n int idx2 = idx + half_d;\n float val_q2 = q[idx2];\n float val_k2 = k[idx2];\n q_rot[idx] = val_q * c - val_q2 * s; // because rotate_half gives -x2 for first half? \n k_rot[idx] = val_k * c - val_k2 * s;\n } else {\n int idx1 = idx - half_d;\n float val_q1 = q[idx1];\n float val_k1 = k[idx1];\n q_rot[idx] = val_q1 * s + val_q * c; // for x2 part: x1 * s + x2 * c\n k_rot[idx] = val_k1 * s + val_k * c;\n }\n}\n\nstd::tuple<torch::Tensor, torch::Tensor> apply_rope(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor cos,\n torch::Tensor sin\n) {\n int B = q.size(0);\n int H = q.size(1);\n int N = q.size(2);\n int D = q.size(3);\n\n auto q_rot = torch::empty_like(q);\n auto k_rot = 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_rot.data_ptr<float>(),\n k_rot.data_ptr<float>(),\n cos.data_ptr<float>(),\n sin.data_ptr<float>(),\n B, H, N, D\n );\n\n return std::make_tuple(q_rot, k_rot);\n}\n'''\n\ncpp_source = '''\nstd::tuple<torch::Tensor, torch::Tensor> apply_rope(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor cos,\n torch::Tensor sin\n);\n'''\n\nmodule = load_inline(\n name='rope_cuda',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['apply_rope'],\n verbose=True\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 N = q.shape[2]\n device = q.device\n dtype = self.inv_freq.dtype\n t = torch.arange(N, device=device, dtype=dtype)\n freqs = torch.outer(t, self.inv_freq)\n emb = torch.cat([freqs, freqs], dim=-1)\n cos = emb.cos().contiguous()\n sin = emb.sin().contiguous()\n return module.apply_rope(q, k, cos, sin)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 2669.5s turn · 9.1s LLM · 660.4s tools · tokens in=3,296 · out=896 · total=4,199
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\n__global__ void rope_kernel(\n const float* __restrict__ q,\n const float* __restrict__ k,\n float* __restrict__ q_rot,\n float* __restrict__ k_rot,\n const float* __restrict__ cos,\n const float* __restrict__ sin,\n int B, int H, int N, 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 int half_d = D / 2;\n float c = cos[n * D + d];\n float s = sin[n * D + d];\n\n float val_q = q[idx];\n float val_k = k[idx];\n\n if (d < half_d) {\n int idx2 = idx + half_d;\n float val_q2 = q[idx2];\n float val_k2 = k[idx2];\n q_rot[idx] = val_q * c - val_q2 * s;\n k_rot[idx] = val_k * c - val_k2 * s;\n } else {\n int idx1 = idx - half_d;\n float val_q1 = q[idx1];\n float val_k1 = k[idx1];\n q_rot[idx] = val_q1 * s + val_q * c;\n k_rot[idx] = val_k1 * s + val_k * c;\n }\n}\n\nstd::tuple<torch::Tensor, torch::Tensor> apply_rope(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor cos,\n torch::Tensor sin\n) {\n int B = q.size(0);\n int H = q.size(1);\n int N = q.size(2);\n int D = q.size(3);\n\n auto q_rot = torch::empty_like(q);\n auto k_rot = 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_rot.data_ptr<float>(),\n k_rot.data_ptr<float>(),\n cos.data_ptr<float>(),\n sin.data_ptr<float>(),\n B, H, N, D\n );\n\n return std::make_tuple(q_rot, k_rot);\n}\n'''\n\ncpp_source = '''\nstd::tuple<torch::Tensor, torch::Tensor> apply_rope(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor cos,\n torch::Tensor sin\n);\n'''\n\nmodule = load_inline(\n name='rope_cuda',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['apply_rope'],\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 N = q.shape[2]\n device = q.device\n dtype = self.inv_freq.dtype\n t = torch.arange(N, device=device, dtype=dtype)\n freqs = torch.outer(t, self.inv_freq)\n emb = torch.cat([freqs, freqs], dim=-1)\n cos = emb.cos().contiguous()\n sin = emb.sin().contiguous()\n return module.apply_rope(q, k, cos, sin)\n"
}
result ← run_correctness (ok)
okrun_correctness
run_correctness PASSED: (5 / 5) 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 3901.5s turn · 20.8s LLM · 880.7s tools · tokens in=4,259 · out=895 · total=5,272 · FINAL
reasoning
Correct! Now, to make it faster, I should optimize the kernel: use vectorized loads (float4), better memory access, precompute cos/sin in a better way, perhaps compute cos/sin on the fly or fuse more....
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\n__global__ void rope_kernel(\n const float* __restrict__ q,\n const float* __restrict__ k,\n float* __restrict__ q_rot,\n float* __restrict__ k_rot,\n const float* __restrict__ cos,\n const float* __restrict__ sin,\n int B, int H, int N, 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 int half_d = D / 2;\n float c = cos[n * D + d];\n float s = sin[n * D + d];\n\n float val_q = q[idx];\n float val_k = k[idx];\n\n if (d < half_d) {\n int idx2 = idx + half_d;\n float val_q2 = q[idx2];\n float val_k2 = k[idx2];\n q_rot[idx] = val_q * c - val_q2 * s;\n k_rot[idx] = val_k * c - val_k2 * s;\n } else {\n int idx1 = idx - half_d;\n float val_q1 = q[idx1];\n float val_k1 = k[idx1];\n q_rot[idx] = val_q1 * s + val_q * c;\n k_rot[idx] = val_k1 * s + val_k * c;\n }\n}\n\nstd::tuple<torch::Tensor, torch::Tensor> apply_rope(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor cos,\n torch::Tensor sin\n) {\n int B = q.size(0);\n int H = q.size(1);\n int N = q.size(2);\n int D = q.size(3);\n\n auto q_rot = torch::empty_like(q);\n auto k_rot = 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_rot.data_ptr<float>(),\n k_rot.data_ptr<float>(),\n cos.data_ptr<float>(),\n sin.data_ptr<float>(),\n B, H, N, D\n );\n\n return std::make_tuple(q_rot, k_rot);\n}\n'''\n\ncpp_source = '''\nstd::tuple<torch::Tensor, torch::Tensor> apply_rope(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor cos,\n torch::Tensor sin\n);\n'''\n\nmodule = load_inline(\n name='rope_cuda',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['apply_rope'],\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 N = q.shape[2]\n device = q.device\n dtype = self.inv_freq.dtype\n t = torch.arange(N, device=device, dtype=dtype)\n freqs = torch.outer(t, self.inv_freq)\n emb = torch.cat([freqs, freqs], dim=-1)\n cos = emb.cos().contiguous()\n sin = emb.sin().contiguous()\n return module.apply_rope(q, k, cos, sin)\n"
}
result ← submit_kernel (ok)
oksubmit_kernel
submit_kernel PASSED: (5 / 5) correctness trials all passed.
Kernel runtime: 0.02 μs
Runtime stats: mean=0.02μ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: 15 (ref: 45, fusion ratio: 3.0)
Energy: 237.62 mJ/run (ref: 0.0 mJ, ratio: -1x)
SOL score: 0.3963
final submitted kernel (117 lines, 2,995 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>
__global__ void rope_kernel(
const float* __restrict__ q,
const float* __restrict__ k,
float* __restrict__ q_rot,
float* __restrict__ k_rot,
const float* __restrict__ cos,
const float* __restrict__ sin,
int B, int H, int N, 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);
int half_d = D / 2;
float c = cos[n * D + d];
float s = sin[n * D + d];
float val_q = q[idx];
float val_k = k[idx];
if (d < half_d) {
int idx2 = idx + half_d;
float val_q2 = q[idx2];
float val_k2 = k[idx2];
q_rot[idx] = val_q * c - val_q2 * s;
k_rot[idx] = val_k * c - val_k2 * s;
} else {
int idx1 = idx - half_d;
float val_q1 = q[idx1];
float val_k1 = k[idx1];
q_rot[idx] = val_q1 * s + val_q * c;
k_rot[idx] = val_k1 * s + val_k * c;
}
}
std::tuple<torch::Tensor, torch::Tensor> apply_rope(
torch::Tensor q,
torch::Tensor k,
torch::Tensor cos,
torch::Tensor sin
) {
int B = q.size(0);
int H = q.size(1);
int N = q.size(2);
int D = q.size(3);
auto q_rot = torch::empty_like(q);
auto k_rot = 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_rot.data_ptr<float>(),
k_rot.data_ptr<float>(),
cos.data_ptr<float>(),
sin.data_ptr<float>(),
B, H, N, D
);
return std::make_tuple(q_rot, k_rot);
}
'''
cpp_source = '''
std::tuple<torch::Tensor, torch::Tensor> apply_rope(
torch::Tensor q,
torch::Tensor k,
torch::Tensor cos,
torch::Tensor sin
);
'''
module = load_inline(
name='rope_cuda',
cpp_sources=[cpp_source],
cuda_sources=[cuda_source],
functions=['apply_rope'],
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:
N = q.shape[2]
device = q.device
dtype = self.inv_freq.dtype
t = torch.arange(N, device=device, dtype=dtype)
freqs = torch.outer(t, self.inv_freq)
emb = torch.cat([freqs, freqs], dim=-1)
cos = emb.cos().contiguous()
sin = emb.sin().contiguous()
return module.apply_rope(q, k, cos, sin)