grok-4.3 · popcorn · pop_l123_default_gpt
correct 20_RotaryPositionEmbeddingBio.py
model
grok-4.3
turns / max
4 / 12
tool calls
4
tokens (total)
14750
wall clock
2493.3s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
0.02 μs
ref runtime
0.07 μs
speedup
2.72x
started
2026-05-11T13:36:39Z
timingkernelbench.timing
kernel runtime
0.02 μs
ref runtime
0.07 μs
speedup
2.72x
kernel mean / median / std
0.02 / — / 9.65e-04 μs
ref mean / median / std
0.07 / — / 0.00 μs
numerical precisioneval_kernel_against_ref
max abs error
4.77e-07
mean abs error
1.43e-08
max rel error
0.01
mean rel error
5.94e-08
gpu memoryextended_metrics.measure_memory
kernel peak
41.22 MB
ref peak
45.13 MB
kernel / ref
0.91x
kernel launches / fusiontorch.profiler
kernel launches
15
ref launches
45
fusion ratio (ref / kernel)
3.00x
total cuda time
128 μs
energyNVML (pynvml)
kernel mJ / run
237.62
ref mJ / run
0.00
ratio (ref / kernel)
—
avg power
1615.75 W
ref avg power
0.00 W
SOL (speed-of-light)nsight_hardware_counters
SOL score (0–1)
0.396
dram util %
39.6%
compute util %
25.1%
bottleneck
memory-bound
dominant pipe
FP32
arithmetic intensity
—
ridge point
20.00
roofline / occupancynsight_hardware_counters
dram bandwidth
1327.64 GB/s
dram util %
39.6%
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
98.8%
bottleneck
memory-bound
dominant pipe %
25.1%
arithmetic intensity
—
ridge point
20.00
L1 hit %
—
L2 hit %
536.0%
ld sectors / req
54.50
st sectors / req
—
regs / thread
144
smem / block
6144 B
block size
832
peak bw
3350.00 GB/s
peak fp32 tflops
67.00
peak fp16 tflops
1979.00
warp stalls (top 5)nsight
long_scoreboard
3680.3%
short_scoreboard
2025.6%
mio_throttle
379.9%
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)