correct
3_FusedCrossAttentionRoPEDropout.py
started
2026-05-04T08:36:35Z
timingkernelbench.timing
kernel mean / median / std
0.10 / — / 0.02 μs
ref mean / median / std
0.39 / — / 0.06 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
3.45x
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
dram bandwidth
1682.18 GB/s
warp stalls (top 5)nsight
reference PyTorch source (the prompt the agent sees)
import math
import torch
import torch.nn as nn
import torch.nn.functional as F
# Default ModelNew is PyTorch-only. run_and_check: use check_kernel=False with backend=cuda
# until the candidate file contains a real CUDA kernel (see kernel_static_checker).
class Model(nn.Module):
"""
Single-layer multi-head cross-attention with RoPE on Q and K, attention
dropout, and output projection—intended as a fusion target for robotics
stacks (e.g. proprio / action queries attending to vision or memory).
Pipeline (conceptually one fused kernel): linear Q from query stream,
linear K/V from key–value stream → reshape heads → apply RoPE separately
along query positions (T_q) and key positions (T_k) → scaled dot-product,
softmax, dropout on weights → aggregate values → output projection.
"""
def __init__(
self,
dim: int,
num_heads: int,
dropout: float = 0.0,
rope_base: float = 10000.0,
):
super().__init__()
assert dim % num_heads == 0
self.dim = dim
self.num_heads = num_heads
self.head_dim = dim // num_heads
assert self.head_dim % 2 == 0, "RoPE requires an even head dimension"
self.q_proj = nn.Linear(dim, dim, bias=False)
self.k_proj = nn.Linear(dim, dim, bias=False)
self.v_proj = nn.Linear(dim, dim, bias=False)
self.out_proj = nn.Linear(dim, dim, bias=False)
self.dropout = nn.Dropout(dropout)
self.scale = 1.0 / math.sqrt(self.head_dim)
inv_freq = 1.0 / (
rope_base ** (torch.arange(0, self.head_dim, 2).float() / self.head_dim)
)
self.register_buffer("inv_freq", inv_freq, persistent=False)
@staticmethod
def _rotate_half(x: torch.Tensor) -> torch.Tensor:
x1, x2 = x.chunk(2, dim=-1)
return torch.cat([-x2, x1], dim=-1)
def _rope(self, x: torch.Tensor, seq_len: int) -> torch.Tensor:
"""Apply RoPE along the sequence dimension (last-but-one after head). x: (B, H, L, d)."""
t = torch.arange(seq_len, device=x.device, dtype=self.inv_freq.dtype)
freqs = torch.outer(t, self.inv_freq)
emb = torch.cat([freqs, freqs], dim=-1)
cos = emb.cos().view(1, 1, seq_len, self.head_dim)
sin = emb.sin().view(1, 1, seq_len, self.head_dim)
return x * cos + self._rotate_half(x) * sin
def forward(self, x_query: torch.Tensor, x_kv: torch.Tensor) -> torch.Tensor:
"""
Args:
x_query: (B, T_q, dim) — e.g. robot or decoder tokens
x_kv: (B, T_k, dim) — e.g. encoder / memory tokens
Returns:
(B, T_q, dim)
"""
B, T_q, _ = x_query.shape
_, T_k, _ = x_kv.shape
h, d = self.num_heads, self.head_dim
q = self.q_proj(x_query).view(B, T_q, h, d).transpose(1, 2)
k = self.k_proj(x_kv).view(B, T_k, h, d).transpose(1, 2)
v = self.v_proj(x_kv).view(B, T_k, h, d).transpose(1, 2)
q = self._rope(q, T_q)
k = self._rope(k, T_k)
attn = torch.matmul(q, k.transpose(-2, -1)) * self.scale
attn = F.softmax(attn, dim=-1)
attn = self.dropout(attn)
out = torch.matmul(attn, v)
out = out.transpose(1, 2).contiguous().view(B, T_q, self.dim)
return self.out_proj(out)
class ModelNew(nn.Module):
"""
KernelBench candidate entry point (`eval_kernel_against_ref` / `run_and_check`).
Replace this class with a fused CUDA/Triton (etc.) implementation; the default
delegates to `Model` so you can smoke-test the harness with the same file for
both `ref_arch_src_path` and `kernel_src_path`.
"""
def __init__(
self,
dim: int,
num_heads: int,
dropout: float = 0.0,
rope_base: float = 10000.0,
):
super().__init__()
self._impl = Model(dim, num_heads, dropout, rope_base)
def forward(self, x_query: torch.Tensor, x_kv: torch.Tensor) -> torch.Tensor:
return self._impl(x_query, x_kv)
batch_size = 4
seq_len_q = 64
seq_len_kv = 128
dim = 256
num_heads = 8
# Use 0.0 so `run_and_check` / multi-trial correctness matches: the harness does not
# re-seed RNG between the reference forward and `ModelNew` forward, so attention
# dropout > 0 would make outputs differ even for identical math.
dropout_p = 0.0
def get_inputs():
p = popcorn_pri
bs = p.jitter_int(batch_size)
tq = p.jitter_int(seq_len_q)
tk = p.jitter_int(seq_len_kv)
x_query = torch.randn(bs, tq, dim)
x_kv = torch.randn(bs, tk, dim)
return [x_query, x_kv]
def get_init_inputs():
return [dim, num_heads, dropout_p]
turn 011.4s turn · 11.4s LLM · 0.0s tools
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 1 · 1 tool call(s) · 39.0s turn · 39.0s LLM · 8,933 chars (click to expand)
turn 139.0s turn · 39.0s LLM · 0.0s tools
call → compile_kernel
callcompile_kernel
{
"kernel_code": "import math\nimport torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cfloat>\n\n__global__ void rope_kernel(const float* __restrict__ x, float* __restrict__ y,\n const float* __restrict__ inv_freq,\n int B, int T, int D, int H, int d) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * T * D;\n if (idx >= total) return;\n int j_global = idx % D;\n int j = j_global % d;\n int half = d >> 1;\n int pair = j < half ? j : j - half;\n int base = idx - j;\n float c, s;\n int t = (idx / D) % T;\n sincosf((float)t * inv_freq[pair], &s, &c);\n float a = x[idx];\n float b = (j < half) ? x[base + j + half] : x[base + j - half];\n y[idx] = (j < half) ? (a * c - b * s) : (a * c + b * s);\n}\n\n__global__ void attn_kernel(const float* __restrict__ q,\n const float* __restrict__ k,\n const float* __restrict__ v,\n float* __restrict__ out,\n int B, int Tq, int Tk, int D, int H, int d,\n float scale) {\n extern __shared__ float scores[];\n int tq = blockIdx.x;\n int h = blockIdx.y;\n int b = blockIdx.z;\n int tx = threadIdx.x;\n int qbase = (b * Tq + tq) * D + h * d;\n\n if (tx < Tk) {\n int kbase = (b * Tk + tx) * D + h * d;\n float sum = 0.0f;\n #pragma unroll\n for (int j = 0; j < 32; ++j) {\n if (j < d) sum += q[qbase + j] * k[kbase + j];\n }\n scores[tx] = sum * scale;\n }\n __syncthreads();\n\n // Serial softmax per row; Tk is small (128 in benchmark).\n if (tx == 0) {\n float m = -FLT_MAX;\n for (int i = 0; i < Tk; ++i) m = fmaxf(m, scores[i]);\n float denom = 0.0f;\n for (int i = 0; i < Tk; ++i) {\n float e = __expf(scores[i] - m);\n scores[i] = e;\n denom += e;\n }\n float inv = 1.0f / denom;\n for (int i = 0; i < Tk; ++i) scores[i] *= inv;\n }\n __syncthreads();\n\n if (tx < d) {\n float acc = 0.0f;\n for (int i = 0; i < Tk; ++i) {\n int vbase = (b * Tk + i) * D + h * d;\n acc += scores[i] * v[vbase + tx];\n }\n out[qbase + tx] = acc;\n }\n}\n\ntorch::Tensor attn_rope_forward(torch::Tensor q, torch::Tensor k, torch::Tensor v,\n torch::Tensor inv_freq, int64_t num_heads) {\n int B = q.size(0);\n int Tq = q.size(1);\n int D = q.size(2);\n int Tk = k.size(1);\n int H = (int)num_heads;\n int d = D / H;\n auto qrot = torch::empty_like(q);\n auto krot = torch::empty_like(k);\n auto out = torch::empty_like(q);\n cudaStream_t stream = at::cuda::getCurrentCUDAStream();\n int threads = 256;\n int n_q = B * Tq * D;\n int n_k = B * Tk * D;\n rope_kernel<<<(n_q + threads - 1) / threads, threads, 0, stream>>>(\n q.data_ptr<float>(), qrot.data_ptr<float>(), inv_freq.data_ptr<float>(), B, Tq, D, H, d);\n rope_kernel<<<(n_k + threads - 1) / threads, threads, 0, stream>>>(\n k.data_ptr<float>(), krot.data_ptr<float>(), inv_freq.data_ptr<float>(), B, Tk, D, H, d);\n dim3 grid(Tq, H, B);\n int attn_threads = 128;\n if (Tk > 128) attn_threads = 256;\n attn_kernel<<<grid, attn_threads, Tk * sizeof(float), stream>>>(\n qrot.data_ptr<float>(), krot.data_ptr<float>(), v.data_ptr<float>(), out.data_ptr<float>(),\n B, Tq, Tk, D, H, d, rsqrtf((float)d));\n return out;\n}\n'''\n\n_ext = load_inline(name='cross_attn_rope_ext_v1', cpp_sources='', cuda_sources=_cuda_src,\n functions=['attn_rope_forward'], extra_cuda_cflags=['-O3', '--use_fast_math'], verbose=False)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim: int, num_heads: int, dropout: float = 0.0, rope_base: float = 10000.0):\n super().__init__()\n assert dim % num_heads == 0\n self.dim = dim\n self.num_heads = num_heads\n self.head_dim = dim // num_heads\n assert self.head_dim % 2 == 0\n self.q_proj = nn.Linear(dim, dim, bias=False)\n self.k_proj = nn.Linear(dim, dim, bias=False)\n self.v_proj = nn.Linear(dim, dim, bias=False)\n self.out_proj = nn.Linear(dim, dim, bias=False)\n self.dropout = nn.Dropout(dropout)\n self.scale = 1.0 / math.sqrt(self.head_dim)\n inv_freq = 1.0 / (rope_base ** (torch.arange(0, self.head_dim, 2).float() / self.head_dim))\n self.register_buffer('inv_freq', inv_freq, persistent=False)\n def forward(self, x_query: torch.Tensor, x_kv: torch.Tensor) -> torch.Tensor:\n q = self.q_proj(x_query).contiguous()\n k = self.k_proj(x_kv).contiguous()\n v = self.v_proj(x_kv).contiguous()\n out = _ext.attn_rope_forward(q, k, v, self.inv_freq, self.num_heads)\n return self.out_proj(out)\n"
}
result ← compile_kernel (fail)
failcompile_kernel
compile_kernel FAILED: RuntimeError.
RuntimeError: Error building extension 'cross_attn_rope_ext_v1': [1/3] /opt/conda/bin/x86_64-conda-linux-gnu-c++ -MMD -MF main.o.d -DTORCH_EXTENSION_NAME=cross_attn_rope_ext_v1 -DTORCH_API_INCLUDE_EXTENSION_H -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /usr/local/cuda/include -isystem /opt/conda/include/python3.10 -fPIC -std=c++17 -c /scratch/abaner/KernelBench/runs/l12_all_gpt/popcorn/gpt-5.5-priority/level_2_problem_3_cache/k_af272e1140b1/cross_attn_rope_ext_v1/main.cpp -o main.o
FAILED: [code=1] main.o
/opt/conda/bin/x86_64-conda-linux-gnu-c++ -MMD -MF main.o.d -DTORCH_EXTENSION_NAME=cross_attn_rope_ext_v1 -DTORCH_API_INCLUDE_EXTENSION_H -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /usr/local/cuda/include -isystem /opt/conda/include/python3.10 -fPIC -std=c++17 -c /scratch/abaner/KernelBench/runs/l12_all_gpt/popcorn/gpt-5.5-priority/level_2_problem_3_cache/k_af272e1140b1/cross_attn_rope_ext_v1/main.cpp -o main.o
/scratch/abaner/KernelBench/runs/l12_all_gpt/popcorn/gpt-5.5-priority/level_2_problem_3_cache/k_af272e1140b1/cross_attn_rope_ext_v1/main.cpp: In function 'void pybind11_init_cross_attn_rope_ext_v1(pybind11::module_&)':
/scratch/abaner/KernelBench/runs/l12_all_gpt/popcorn/gpt-5.5-priority/level_2_problem_3_cache/k_af272e1140b1/cross_attn_rope_ext_v1/main.cpp:4:56: error: 'attn_rope_forward' was not declared in this scope
4 | m.def("attn_rope_forward", torch::wrap_pybind_function(attn_rope_forward), "attn_rope_forward");
| ^~~~~~~~~~~~~~~~~
[2/3] /usr/local/cuda/bin/nvcc -MD -MF cuda.cuda.o.d -ccbin /opt/conda/bin/x86_64-conda-linux-gnu-cc -DTORCH_EXTENSION_NAME=cross_attn_rope_ext_v1 -DTORCH_API_INCLUDE_EXTENSION_H -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /usr/local/cuda/include -isystem /opt/conda/include/python3.10 -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr -gencode=arch=compute_90,code=compute_90 -gencode=arch=compute_90,code=sm_90 --compiler-options '-fPIC' -O3 --use_fast_math -std=c++17 -c /scratch/abaner/KernelBench/runs/l12_all_gpt/popcorn/gpt-5.5-priority/level_2_problem_3_cache/k_af272e1140b1/cross_attn_rope_ext_v1/cuda.cu -o cuda.cuda.o
nvcc warning : incompatible redefinition for option 'compiler-bindir', the last value of this option was used
ninja: build stopped: subcommand failed.
turn 2 · 1 tool call(s) · 10.9s turn · 10.9s LLM · 6,926 chars (click to expand)
turn 210.9s turn · 10.9s LLM · 0.0s tools
reasoning
**Considering C++ Prototypes**
I'm thinking about the need for a declaration in my C++ source files. It's important because a proper prototype helps define the functions before they're used. I’ll make sure to include the appropriate function signature that matches the implementation later. It feels a bit tricky, but I know it's essential for avoiding errors and ensuring everything compiles smoothly. I should check what parameters and return types I need to specify.
**Preparing code variables**
I’m getting ready to compile my code variables. I think it’s important to set everything up properly, but I also want to be efficient and reuse any variables when I can. This way, I won’t create unnecessary duplicate code. So, I’ll focus on ensuring that things are organized, and easy to manage later. Let’s see how I can structure it for the best outcome!
call → compile_kernel
callcompile_kernel
{
"kernel_code": "import math\nimport torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n_cpp_src = 'torch::Tensor attn_rope_forward(torch::Tensor q, torch::Tensor k, torch::Tensor v, torch::Tensor inv_freq, int64_t num_heads);'\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cfloat>\n\n__global__ void rope_kernel(const float* __restrict__ x, float* __restrict__ y,\n const float* __restrict__ inv_freq,\n int B, int T, int D, int H, int d) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * T * D;\n if (idx >= total) return;\n int j_global = idx % D;\n int j = j_global % d;\n int half = d >> 1;\n int pair = j < half ? j : j - half;\n int base = idx - j;\n float c, s;\n int t = (idx / D) % T;\n sincosf((float)t * inv_freq[pair], &s, &c);\n float a = x[idx];\n float b = (j < half) ? x[base + j + half] : x[base + j - half];\n y[idx] = (j < half) ? (a * c - b * s) : (a * c + b * s);\n}\n\n__global__ void attn_kernel(const float* __restrict__ q,\n const float* __restrict__ k,\n const float* __restrict__ v,\n float* __restrict__ out,\n int B, int Tq, int Tk, int D, int H, int d,\n float scale) {\n extern __shared__ float scores[];\n int tq = blockIdx.x;\n int h = blockIdx.y;\n int b = blockIdx.z;\n int tx = threadIdx.x;\n int qbase = (b * Tq + tq) * D + h * d;\n\n if (tx < Tk) {\n int kbase = (b * Tk + tx) * D + h * d;\n float sum = 0.0f;\n for (int j = 0; j < d; ++j) sum += q[qbase + j] * k[kbase + j];\n scores[tx] = sum * scale;\n }\n __syncthreads();\n\n if (tx == 0) {\n float m = -FLT_MAX;\n for (int i = 0; i < Tk; ++i) m = fmaxf(m, scores[i]);\n float denom = 0.0f;\n for (int i = 0; i < Tk; ++i) {\n float e = __expf(scores[i] - m);\n scores[i] = e;\n denom += e;\n }\n float inv = 1.0f / denom;\n for (int i = 0; i < Tk; ++i) scores[i] *= inv;\n }\n __syncthreads();\n\n if (tx < d) {\n float acc = 0.0f;\n for (int i = 0; i < Tk; ++i) {\n int vbase = (b * Tk + i) * D + h * d;\n acc += scores[i] * v[vbase + tx];\n }\n out[qbase + tx] = acc;\n }\n}\n\ntorch::Tensor attn_rope_forward(torch::Tensor q, torch::Tensor k, torch::Tensor v,\n torch::Tensor inv_freq, int64_t num_heads) {\n int B = q.size(0);\n int Tq = q.size(1);\n int D = q.size(2);\n int Tk = k.size(1);\n int H = (int)num_heads;\n int d = D / H;\n auto qrot = torch::empty_like(q);\n auto krot = torch::empty_like(k);\n auto out = torch::empty_like(q);\n cudaStream_t stream = at::cuda::getCurrentCUDAStream();\n int threads = 256;\n int n_q = B * Tq * D;\n int n_k = B * Tk * D;\n rope_kernel<<<(n_q + threads - 1) / threads, threads, 0, stream>>>(\n q.data_ptr<float>(), qrot.data_ptr<float>(), inv_freq.data_ptr<float>(), B, Tq, D, H, d);\n rope_kernel<<<(n_k + threads - 1) / threads, threads, 0, stream>>>(\n k.data_ptr<float>(), krot.data_ptr<float>(), inv_freq.data_ptr<float>(), B, Tk, D, H, d);\n dim3 grid(Tq, H, B);\n int attn_threads = (Tk <= 128 ? 128 : 256);\n attn_kernel<<<grid, attn_threads, Tk * sizeof(float), stream>>>(\n qrot.data_ptr<float>(), krot.data_ptr<float>(), v.data_ptr<float>(), out.data_ptr<float>(),\n B, Tq, Tk, D, H, d, rsqrtf((float)d));\n return out;\n}\n'''\n\n_ext = load_inline(name='cross_attn_rope_ext_v2', cpp_sources=_cpp_src, cuda_sources=_cuda_src,\n functions=['attn_rope_forward'], extra_cuda_cflags=['-O3', '--use_fast_math'], verbose=False)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim: int, num_heads: int, dropout: float = 0.0, rope_base: float = 10000.0):\n super().__init__()\n assert dim % num_heads == 0\n self.dim = dim\n self.num_heads = num_heads\n self.head_dim = dim // num_heads\n assert self.head_dim % 2 == 0\n self.q_proj = nn.Linear(dim, dim, bias=False)\n self.k_proj = nn.Linear(dim, dim, bias=False)\n self.v_proj = nn.Linear(dim, dim, bias=False)\n self.out_proj = nn.Linear(dim, dim, bias=False)\n self.dropout = nn.Dropout(dropout)\n self.scale = 1.0 / math.sqrt(self.head_dim)\n inv_freq = 1.0 / (rope_base ** (torch.arange(0, self.head_dim, 2).float() / self.head_dim))\n self.register_buffer('inv_freq', inv_freq, persistent=False)\n def forward(self, x_query: torch.Tensor, x_kv: torch.Tensor) -> torch.Tensor:\n q = self.q_proj(x_query).contiguous()\n k = self.k_proj(x_kv).contiguous()\n v = self.v_proj(x_kv).contiguous()\n out = _ext.attn_rope_forward(q, k, v, self.inv_freq, self.num_heads)\n return self.out_proj(out)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 3 · 1 tool call(s) · 15.3s turn · 15.3s LLM · 6,062 chars (click to expand)
turn 315.3s turn · 15.3s LLM · 0.0s tools
call → run_correctness
callrun_correctness
{
"kernel_code": "import math\nimport torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n_cpp_src = 'torch::Tensor attn_rope_forward(torch::Tensor q, torch::Tensor k, torch::Tensor v, torch::Tensor inv_freq, int64_t num_heads);'\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cfloat>\n\n__global__ void rope_kernel(const float* __restrict__ x, float* __restrict__ y,\n const float* __restrict__ inv_freq,\n int B, int T, int D, int H, int d) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * T * D;\n if (idx >= total) return;\n int j_global = idx % D;\n int j = j_global % d;\n int half = d >> 1;\n int pair = j < half ? j : j - half;\n int base = idx - j;\n float c, s;\n int t = (idx / D) % T;\n sincosf((float)t * inv_freq[pair], &s, &c);\n float a = x[idx];\n float b = (j < half) ? x[base + j + half] : x[base + j - half];\n y[idx] = (j < half) ? (a * c - b * s) : (a * c + b * s);\n}\n\n__global__ void attn_kernel(const float* __restrict__ q,\n const float* __restrict__ k,\n const float* __restrict__ v,\n float* __restrict__ out,\n int B, int Tq, int Tk, int D, int H, int d,\n float scale) {\n extern __shared__ float scores[];\n int tq = blockIdx.x;\n int h = blockIdx.y;\n int b = blockIdx.z;\n int tx = threadIdx.x;\n int qbase = (b * Tq + tq) * D + h * d;\n\n if (tx < Tk) {\n int kbase = (b * Tk + tx) * D + h * d;\n float sum = 0.0f;\n for (int j = 0; j < d; ++j) sum += q[qbase + j] * k[kbase + j];\n scores[tx] = sum * scale;\n }\n __syncthreads();\n\n if (tx == 0) {\n float m = -FLT_MAX;\n for (int i = 0; i < Tk; ++i) m = fmaxf(m, scores[i]);\n float denom = 0.0f;\n for (int i = 0; i < Tk; ++i) {\n float e = __expf(scores[i] - m);\n scores[i] = e;\n denom += e;\n }\n float inv = 1.0f / denom;\n for (int i = 0; i < Tk; ++i) scores[i] *= inv;\n }\n __syncthreads();\n\n if (tx < d) {\n float acc = 0.0f;\n for (int i = 0; i < Tk; ++i) {\n int vbase = (b * Tk + i) * D + h * d;\n acc += scores[i] * v[vbase + tx];\n }\n out[qbase + tx] = acc;\n }\n}\n\ntorch::Tensor attn_rope_forward(torch::Tensor q, torch::Tensor k, torch::Tensor v,\n torch::Tensor inv_freq, int64_t num_heads) {\n int B = q.size(0);\n int Tq = q.size(1);\n int D = q.size(2);\n int Tk = k.size(1);\n int H = (int)num_heads;\n int d = D / H;\n auto qrot = torch::empty_like(q);\n auto krot = torch::empty_like(k);\n auto out = torch::empty_like(q);\n cudaStream_t stream = at::cuda::getCurrentCUDAStream();\n int threads = 256;\n int n_q = B * Tq * D;\n int n_k = B * Tk * D;\n rope_kernel<<<(n_q + threads - 1) / threads, threads, 0, stream>>>(\n q.data_ptr<float>(), qrot.data_ptr<float>(), inv_freq.data_ptr<float>(), B, Tq, D, H, d);\n rope_kernel<<<(n_k + threads - 1) / threads, threads, 0, stream>>>(\n k.data_ptr<float>(), krot.data_ptr<float>(), inv_freq.data_ptr<float>(), B, Tk, D, H, d);\n dim3 grid(Tq, H, B);\n int attn_threads = (Tk <= 128 ? 128 : 256);\n attn_kernel<<<grid, attn_threads, Tk * sizeof(float), stream>>>(\n qrot.data_ptr<float>(), krot.data_ptr<float>(), v.data_ptr<float>(), out.data_ptr<float>(),\n B, Tq, Tk, D, H, d, rsqrtf((float)d));\n return out;\n}\n'''\n\n_ext = load_inline(name='cross_attn_rope_ext_v2', cpp_sources=_cpp_src, cuda_sources=_cuda_src,\n functions=['attn_rope_forward'], extra_cuda_cflags=['-O3', '--use_fast_math'], verbose=False)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim: int, num_heads: int, dropout: float = 0.0, rope_base: float = 10000.0):\n super().__init__()\n assert dim % num_heads == 0\n self.dim = dim\n self.num_heads = num_heads\n self.head_dim = dim // num_heads\n assert self.head_dim % 2 == 0\n self.q_proj = nn.Linear(dim, dim, bias=False)\n self.k_proj = nn.Linear(dim, dim, bias=False)\n self.v_proj = nn.Linear(dim, dim, bias=False)\n self.out_proj = nn.Linear(dim, dim, bias=False)\n self.dropout = nn.Dropout(dropout)\n self.scale = 1.0 / math.sqrt(self.head_dim)\n inv_freq = 1.0 / (rope_base ** (torch.arange(0, self.head_dim, 2).float() / self.head_dim))\n self.register_buffer('inv_freq', inv_freq, persistent=False)\n def forward(self, x_query: torch.Tensor, x_kv: torch.Tensor) -> torch.Tensor:\n q = self.q_proj(x_query).contiguous()\n k = self.k_proj(x_kv).contiguous()\n v = self.v_proj(x_kv).contiguous()\n out = _ext.attn_rope_forward(q, k, v, self.inv_freq, self.num_heads)\n return self.out_proj(out)\n"
}
result ← run_correctness (ok)
okrun_correctness
run_correctness PASSED: (5 / 5) trials all matched the reference.
Numerical precision: max_abs_err=1.28e-07 mean_abs_err=1.65e-08 max_rel_err=2.22e-01
turn 4 · FINAL · 1 tool call(s) · 19.0s turn · 19.0s LLM · 6,262 chars (click to expand)
turn 419.0s turn · 19.0s LLM · 0.0s tools · FINAL
call → submit_kernel
callsubmit_kernel
{
"kernel_code": "import math\nimport torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n_cpp_src = 'torch::Tensor attn_rope_forward(torch::Tensor q, torch::Tensor k, torch::Tensor v, torch::Tensor inv_freq, int64_t num_heads);'\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cfloat>\n\n__global__ void rope_kernel(const float* __restrict__ x, float* __restrict__ y,\n const float* __restrict__ inv_freq,\n int B, int T, int D, int H, int d) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * T * D;\n if (idx >= total) return;\n int j_global = idx % D;\n int j = j_global % d;\n int half = d >> 1;\n int pair = j < half ? j : j - half;\n int base = idx - j;\n float c, s;\n int t = (idx / D) % T;\n sincosf((float)t * inv_freq[pair], &s, &c);\n float a = x[idx];\n float b = (j < half) ? x[base + j + half] : x[base + j - half];\n y[idx] = (j < half) ? (a * c - b * s) : (a * c + b * s);\n}\n\n__global__ void attn_kernel(const float* __restrict__ q,\n const float* __restrict__ k,\n const float* __restrict__ v,\n float* __restrict__ out,\n int B, int Tq, int Tk, int D, int H, int d,\n float scale) {\n extern __shared__ float scores[];\n int tq = blockIdx.x;\n int h = blockIdx.y;\n int b = blockIdx.z;\n int tx = threadIdx.x;\n int qbase = (b * Tq + tq) * D + h * d;\n\n if (tx < Tk) {\n int kbase = (b * Tk + tx) * D + h * d;\n float sum = 0.0f;\n for (int j = 0; j < d; ++j) sum += q[qbase + j] * k[kbase + j];\n scores[tx] = sum * scale;\n }\n __syncthreads();\n\n if (tx == 0) {\n float m = -FLT_MAX;\n for (int i = 0; i < Tk; ++i) m = fmaxf(m, scores[i]);\n float denom = 0.0f;\n for (int i = 0; i < Tk; ++i) {\n float e = __expf(scores[i] - m);\n scores[i] = e;\n denom += e;\n }\n float inv = 1.0f / denom;\n for (int i = 0; i < Tk; ++i) scores[i] *= inv;\n }\n __syncthreads();\n\n if (tx < d) {\n float acc = 0.0f;\n for (int i = 0; i < Tk; ++i) {\n int vbase = (b * Tk + i) * D + h * d;\n acc += scores[i] * v[vbase + tx];\n }\n out[qbase + tx] = acc;\n }\n}\n\ntorch::Tensor attn_rope_forward(torch::Tensor q, torch::Tensor k, torch::Tensor v,\n torch::Tensor inv_freq, int64_t num_heads) {\n int B = q.size(0);\n int Tq = q.size(1);\n int D = q.size(2);\n int Tk = k.size(1);\n int H = (int)num_heads;\n int d = D / H;\n auto qrot = torch::empty_like(q);\n auto krot = torch::empty_like(k);\n auto out = torch::empty_like(q);\n cudaStream_t stream = at::cuda::getCurrentCUDAStream();\n int threads = 256;\n int n_q = B * Tq * D;\n int n_k = B * Tk * D;\n rope_kernel<<<(n_q + threads - 1) / threads, threads, 0, stream>>>(\n q.data_ptr<float>(), qrot.data_ptr<float>(), inv_freq.data_ptr<float>(), B, Tq, D, H, d);\n rope_kernel<<<(n_k + threads - 1) / threads, threads, 0, stream>>>(\n k.data_ptr<float>(), krot.data_ptr<float>(), inv_freq.data_ptr<float>(), B, Tk, D, H, d);\n dim3 grid(Tq, H, B);\n int attn_threads = (Tk <= 128 ? 128 : 256);\n attn_kernel<<<grid, attn_threads, Tk * sizeof(float), stream>>>(\n qrot.data_ptr<float>(), krot.data_ptr<float>(), v.data_ptr<float>(), out.data_ptr<float>(),\n B, Tq, Tk, D, H, d, rsqrtf((float)d));\n return out;\n}\n'''\n\n_ext = load_inline(name='cross_attn_rope_ext_v2', cpp_sources=_cpp_src, cuda_sources=_cuda_src,\n functions=['attn_rope_forward'], extra_cuda_cflags=['-O3', '--use_fast_math'], verbose=False)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim: int, num_heads: int, dropout: float = 0.0, rope_base: float = 10000.0):\n super().__init__()\n assert dim % num_heads == 0\n self.dim = dim\n self.num_heads = num_heads\n self.head_dim = dim // num_heads\n assert self.head_dim % 2 == 0\n self.q_proj = nn.Linear(dim, dim, bias=False)\n self.k_proj = nn.Linear(dim, dim, bias=False)\n self.v_proj = nn.Linear(dim, dim, bias=False)\n self.out_proj = nn.Linear(dim, dim, bias=False)\n self.dropout = nn.Dropout(dropout)\n self.scale = 1.0 / math.sqrt(self.head_dim)\n inv_freq = 1.0 / (rope_base ** (torch.arange(0, self.head_dim, 2).float() / self.head_dim))\n self.register_buffer('inv_freq', inv_freq, persistent=False)\n def forward(self, x_query: torch.Tensor, x_kv: torch.Tensor) -> torch.Tensor:\n q = self.q_proj(x_query).contiguous()\n k = self.k_proj(x_kv).contiguous()\n v = self.v_proj(x_kv).contiguous()\n out = _ext.attn_rope_forward(q, k, v, self.inv_freq, self.num_heads)\n return self.out_proj(out)\n"
}
result ← submit_kernel (ok)
oksubmit_kernel
submit_kernel PASSED: (5 / 5) correctness trials all passed.
Kernel runtime: 0.10 μs
Runtime stats: mean=0.10μs median=?μs std=0.02μs
Numerical precision: max_abs_err=1.28e-07 mean_abs_err=1.65e-08 max_rel_err=2.22e-01
Memory: 38.0 MB peak (ref: 39.0 MB, ratio: 0.9744x)
Kernel launches: 42 (ref: 145, fusion ratio: 3.4524)
SOL score: 0.9597
final submitted kernel (131 lines, 5,006 chars) — full source, untruncated
import math
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline
_cpp_src = 'torch::Tensor attn_rope_forward(torch::Tensor q, torch::Tensor k, torch::Tensor v, torch::Tensor inv_freq, int64_t num_heads);'
_cuda_src = r'''
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <ATen/cuda/CUDAContext.h>
#include <cfloat>
__global__ void rope_kernel(const float* __restrict__ x, float* __restrict__ y,
const float* __restrict__ inv_freq,
int B, int T, int D, int H, int d) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int total = B * T * D;
if (idx >= total) return;
int j_global = idx % D;
int j = j_global % d;
int half = d >> 1;
int pair = j < half ? j : j - half;
int base = idx - j;
float c, s;
int t = (idx / D) % T;
sincosf((float)t * inv_freq[pair], &s, &c);
float a = x[idx];
float b = (j < half) ? x[base + j + half] : x[base + j - half];
y[idx] = (j < half) ? (a * c - b * s) : (a * c + b * s);
}
__global__ void attn_kernel(const float* __restrict__ q,
const float* __restrict__ k,
const float* __restrict__ v,
float* __restrict__ out,
int B, int Tq, int Tk, int D, int H, int d,
float scale) {
extern __shared__ float scores[];
int tq = blockIdx.x;
int h = blockIdx.y;
int b = blockIdx.z;
int tx = threadIdx.x;
int qbase = (b * Tq + tq) * D + h * d;
if (tx < Tk) {
int kbase = (b * Tk + tx) * D + h * d;
float sum = 0.0f;
for (int j = 0; j < d; ++j) sum += q[qbase + j] * k[kbase + j];
scores[tx] = sum * scale;
}
__syncthreads();
if (tx == 0) {
float m = -FLT_MAX;
for (int i = 0; i < Tk; ++i) m = fmaxf(m, scores[i]);
float denom = 0.0f;
for (int i = 0; i < Tk; ++i) {
float e = __expf(scores[i] - m);
scores[i] = e;
denom += e;
}
float inv = 1.0f / denom;
for (int i = 0; i < Tk; ++i) scores[i] *= inv;
}
__syncthreads();
if (tx < d) {
float acc = 0.0f;
for (int i = 0; i < Tk; ++i) {
int vbase = (b * Tk + i) * D + h * d;
acc += scores[i] * v[vbase + tx];
}
out[qbase + tx] = acc;
}
}
torch::Tensor attn_rope_forward(torch::Tensor q, torch::Tensor k, torch::Tensor v,
torch::Tensor inv_freq, int64_t num_heads) {
int B = q.size(0);
int Tq = q.size(1);
int D = q.size(2);
int Tk = k.size(1);
int H = (int)num_heads;
int d = D / H;
auto qrot = torch::empty_like(q);
auto krot = torch::empty_like(k);
auto out = torch::empty_like(q);
cudaStream_t stream = at::cuda::getCurrentCUDAStream();
int threads = 256;
int n_q = B * Tq * D;
int n_k = B * Tk * D;
rope_kernel<<<(n_q + threads - 1) / threads, threads, 0, stream>>>(
q.data_ptr<float>(), qrot.data_ptr<float>(), inv_freq.data_ptr<float>(), B, Tq, D, H, d);
rope_kernel<<<(n_k + threads - 1) / threads, threads, 0, stream>>>(
k.data_ptr<float>(), krot.data_ptr<float>(), inv_freq.data_ptr<float>(), B, Tk, D, H, d);
dim3 grid(Tq, H, B);
int attn_threads = (Tk <= 128 ? 128 : 256);
attn_kernel<<<grid, attn_threads, Tk * sizeof(float), stream>>>(
qrot.data_ptr<float>(), krot.data_ptr<float>(), v.data_ptr<float>(), out.data_ptr<float>(),
B, Tq, Tk, D, H, d, rsqrtf((float)d));
return out;
}
'''
_ext = load_inline(name='cross_attn_rope_ext_v2', cpp_sources=_cpp_src, cuda_sources=_cuda_src,
functions=['attn_rope_forward'], extra_cuda_cflags=['-O3', '--use_fast_math'], verbose=False)
class ModelNew(nn.Module):
def __init__(self, dim: int, num_heads: int, dropout: float = 0.0, rope_base: float = 10000.0):
super().__init__()
assert dim % num_heads == 0
self.dim = dim
self.num_heads = num_heads
self.head_dim = dim // num_heads
assert self.head_dim % 2 == 0
self.q_proj = nn.Linear(dim, dim, bias=False)
self.k_proj = nn.Linear(dim, dim, bias=False)
self.v_proj = nn.Linear(dim, dim, bias=False)
self.out_proj = nn.Linear(dim, dim, bias=False)
self.dropout = nn.Dropout(dropout)
self.scale = 1.0 / math.sqrt(self.head_dim)
inv_freq = 1.0 / (rope_base ** (torch.arange(0, self.head_dim, 2).float() / self.head_dim))
self.register_buffer('inv_freq', inv_freq, persistent=False)
def forward(self, x_query: torch.Tensor, x_kv: torch.Tensor) -> torch.Tensor:
q = self.q_proj(x_query).contiguous()
k = self.k_proj(x_kv).contiguous()
v = self.v_proj(x_kv).contiguous()
out = _ext.attn_rope_forward(q, k, v, self.inv_freq, self.num_heads)
return self.out_proj(out)