grok-4.3 · popcorn · pop_l123_all_gpt
correct 20_RotaryPositionEmbeddingBio.py
model
grok-4.3
turns / max
4 / 15
tool calls
3
tokens (total)
24210
wall clock
546.4s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
0.03 μs
ref runtime
0.07 μs
speedup
2.83x
started
2026-05-11T21:14:16Z
timingkernelbench.timing
kernel runtime
0.03 μs
ref runtime
0.07 μs
speedup
2.83x
kernel mean / median / std
0.03 / — / 6.97e-04 μs
ref mean / median / std
0.07 / — / 0.01 μ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
6.40e-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
18
ref launches
45
fusion ratio (ref / kernel)
2.50x
total cuda time
351 μs
SOL (speed-of-light)nsight_hardware_counters
SOL score (0–1)
0.371
dram util %
20.6%
compute util %
37.1%
bottleneck
compute-bound
dominant pipe
FP32
arithmetic intensity
—
ridge point
20.00
roofline / occupancynsight_hardware_counters
dram bandwidth
691.76 GB/s
dram util %
20.6%
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
105.6%
bottleneck
compute-bound
dominant pipe %
37.1%
arithmetic intensity
—
ridge point
20.00
L1 hit %
—
L2 hit %
536.3%
ld sectors / req
54.50
st sectors / req
—
regs / thread
148
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
2429.4%
short_scoreboard
2005.5%
mio_throttle
411.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 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]