incorrect
29_MSARowAttention.py
started
2026-05-11T14:43:38Z
correctness issue: Output mismatch
timingkernelbench.timing
kernel mean / median / std
—
ref mean / median / std
0.54 / — / 0.11 μs
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
dram bandwidth
724.52 GB/s
warp stalls (top 5)nsight
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn
import torch.nn.functional as F
import math
class Model(nn.Module):
"""
Row-wise gated self-attention over Multiple Sequence Alignment (MSA)
representations. Each row of the MSA is an independent sequence; for
each row, standard multi-head attention is performed across residue
positions, with an additive bias from the pair representation.
"""
def __init__(self, msa_dim, pair_dim, num_heads, dropout=0.0):
super().__init__()
self.num_heads = num_heads
self.head_dim = msa_dim // num_heads
assert msa_dim % num_heads == 0
self.layer_norm = nn.LayerNorm(msa_dim)
self.query = nn.Linear(msa_dim, msa_dim, bias=False)
self.key = nn.Linear(msa_dim, msa_dim, bias=False)
self.value = nn.Linear(msa_dim, msa_dim, bias=False)
self.gate = nn.Linear(msa_dim, msa_dim)
self.pair_bias = nn.Linear(pair_dim, num_heads, bias=False)
self.out_proj = nn.Linear(msa_dim, msa_dim)
self.dropout = nn.Dropout(dropout)
def forward(
self,
msa_repr: torch.Tensor,
pair_repr: torch.Tensor,
msa_mask: torch.Tensor,
) -> torch.Tensor:
"""
Args:
msa_repr: (B, S, N, msa_dim)
pair_repr: (B, N, N, pair_dim)
msa_mask: (B, S, N)
Returns:
updated MSA (B, S, N, msa_dim)
"""
B, S, N, D = msa_repr.shape
h = self.num_heads
d = self.head_dim
x = self.layer_norm(msa_repr)
q = self.query(x).view(B, S, N, h, d).permute(0, 1, 3, 2, 4)
k = self.key(x).view(B, S, N, h, d).permute(0, 1, 3, 2, 4)
v = self.value(x).view(B, S, N, h, d).permute(0, 1, 3, 2, 4)
g = torch.sigmoid(self.gate(x)).view(B, S, N, h, d).permute(0, 1, 3, 2, 4)
# Pair bias: shared across MSA rows
bias = self.pair_bias(pair_repr).permute(0, 3, 1, 2) # (B, h, N, N)
attn = torch.einsum("bshid,bshjd->bshij", q, k) / math.sqrt(d)
attn = attn + bias.unsqueeze(1)
mask = msa_mask.view(B, S, 1, 1, N).expand_as(attn)
attn = attn.masked_fill(mask == 0, float("-inf"))
attn = F.softmax(attn, dim=-1)
attn = self.dropout(attn)
out = torch.einsum("bshij,bshjd->bshid", attn, v)
out = g * out
out = out.permute(0, 1, 3, 2, 4).reshape(B, S, N, D)
return self.out_proj(out)
msa_dim = 64
pair_dim = 64
num_heads = 4
seq_len = 32
msa_depth = 8
batch_size = 2
def get_inputs():
p = popcorn_pri
msa_repr = torch.randn(p.jitter_int(batch_size), p.jitter_int(msa_depth), p.jitter_int(seq_len), msa_dim)
pair_repr = torch.randn(p.jitter_int(batch_size), p.jitter_int(seq_len), p.jitter_int(seq_len), pair_dim)
msa_mask = torch.ones(p.jitter_int(batch_size), p.jitter_int(msa_depth), p.jitter_int(seq_len))
return [msa_repr, pair_repr, msa_mask]
def get_init_inputs():
return [msa_dim, pair_dim, num_heads]
turn 0248.4s turn · 3.8s LLM · 244.6s tools · tokens in=2,116 · out=6 · total=2,460
reasoning
The task is to replace the forward computation of this PyTorch model with a custom CUDA kernel. The model is a row-wise gated self-attention over MSA representations.
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 1332.9s turn · 11.7s LLM · 321.2s tools · tokens in=2,693 · out=284 · total=3,392
reasoning
It's an H100, high bandwidth, lots of TFLOPS.
I need to replace the forward computation of a PyTorch model with a custom CUDA kernel. The new ModelNew class must take the same inputs as the original Model, have the softmax over the attention heads, and produce the same outputs: msa_repr, pair_repr, and m的核心mask.
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 dummy_kernel(float* out, const float* in, int n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n if (idx < n) out[idx] = in[idx];\n}\n\ntorch::Tensor dummy_forward(torch::Tensor input) {\n auto output = torch::empty_like(input);\n int n = input.numel();\n dummy_kernel<<<(n+255)/256, 256>>>(output.data_ptr<float>(), input.data_ptr<float>(), n);\n return output;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor dummy_forward(torch::Tensor input);\n\"\"\"\n\nmodule = load_inline(\n name=\"dummy_module\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"dummy_forward\"],\n verbose=True\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, msa_dim, pair_dim, num_heads, dropout=0.0):\n super().__init__()\n self.ref = nn.Module() # placeholder\n def forward(self, msa_repr, pair_repr, msa_mask):\n return module.dummy_forward(msa_repr)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 2480.5s turn · 8.1s LLM · 472.4s tools · tokens in=3,409 · out=385 · total=4,078
reasoning
Good, compilation works. Now I need to implement the real thing. This will be a substantial effort.
I need to replace the forward computation of a PyTorch model with a custom CUDA kernel. The new ModelNew class must take the same inputs as the original Model, have the forward method with the same signature, and produce the same outputs numerically.
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_runtime.h>\n#include <math.h>\n\n__global__ void gated_msa_attention(\n const float* __restrict__ msa, // (B, S, N, D)\n const float* __restrict__ pair, // (B, N, N, P)\n const float* __restrict__ msa_mask, // (B, S, N)\n float* __restrict__ out, // (B, S, N, D)\n // weights omitted for brevity in skeleton\n int B, int S, int N, int D, int h, int P, float scale\n) {\n // Placeholder: copy input to output for now\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n if (idx < B*S*N*D) out[idx] = msa[idx];\n}\n'''\n\ncpp_source = r'''\ntorch::Tensor gated_attention_cuda(torch::Tensor msa_repr, torch::Tensor pair_repr, torch::Tensor msa_mask);\n'''\n\nmodule = load_inline(\n name='gated_msa',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['gated_attention_cuda'],\n verbose=True\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, msa_dim, pair_dim, num_heads, dropout=0.0):\n super().__init__()\n self.msa_dim = msa_dim\n self.pair_dim = pair_dim\n self.num_heads = num_heads\n def forward(self, msa_repr, pair_repr, msa_mask):\n return module.gated_attention_cuda(msa_repr, pair_repr, msa_mask)\n"
}
result ← compile_kernel (fail)
failcompile_kernel
compile_kernel FAILED: ImportError.
ImportError: /scratch/abaner/KernelBench/runs/pop_l123_default_gpt/popcorn/grok-4.3/level_2_problem_29_cache/k_09e6787cdee7/gated_msa/gated_msa.so: undefined symbol: _Z20gated_attention_cudaN2at6TensorES0_S0_
turn 3249.5s turn · 5.0s LLM · 244.5s tools · tokens in=4,165 · out=283 · total=4,476
reasoning
Error because I didn't define the function in CUDA source properly.
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_runtime.h>\n\ntorch::Tensor gated_attention_cuda(torch::Tensor msa_repr, torch::Tensor pair_repr, torch::Tensor msa_mask) {\n auto output = torch::empty_like(msa_repr);\n // TODO: implement kernel\n // For now, just return zeros or copy\n output.copy_(msa_repr);\n return output;\n}\n'''\n\ncpp_source = r'''\ntorch::Tensor gated_attention_cuda(torch::Tensor msa_repr, torch::Tensor pair_repr, torch::Tensor msa_mask);\n'''\n\nmodule = load_inline(\n name='gated_msa',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['gated_attention_cuda'],\n verbose=True\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, msa_dim, pair_dim, num_heads, dropout=0.0):\n super().__init__()\n self.msa_dim = msa_dim\n self.pair_dim = pair_dim\n self.num_heads = num_heads\n def forward(self, msa_repr, pair_repr, msa_mask):\n return module.gated_attention_cuda(msa_repr, pair_repr, msa_mask)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 4 · FINAL · 1 tool call(s) · 460.9s turn · 23.9s LLM · tokens in=4,493 · out=1,339 · total=6,160 · 6,021 chars (click to expand)
turn 4460.9s turn · 23.9s LLM · 437.0s tools · tokens in=4,493 · out=1,339 · total=6,160 · FINAL
reasoning
Good. Now to implement the real thing, I need to write a full CUDA kernel that does the computation.
I need to replace the forward computation of a PyTorch model with a custom CUDA kernel. The new ModelNew class must accept the same inputs as the original Model, have the same forward signature, and produce equivalent outputs. It should include all necessary parameters like Embedder, Attention, etc.
call → submit_kernel
callsubmit_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nimport torch.nn.functional as F\nimport math\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = r'''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n#include <math.h>\n\n__global__ void attention_kernel(\n const float* q, const float* k, const float* v, const float* bias, const float* mask,\n float* out, float* gate,\n int B, int S, int h, int N, int d, float scale\n) {\n // Very naive implementation for small N=32\n int bsh = blockIdx.x;\n int b = bsh / (S * h);\n int sh = bsh % (S * h);\n int s = sh / h;\n int head = sh % h;\n \n extern __shared__ float smem[];\n \n // Compute attn scores\n for(int i = threadIdx.x; i < N; i += blockDim.x) {\n float sum = 0.0f;\n for(int j=0; j<N; j++) {\n float score = 0.0f;\n for(int dd=0; dd<d; dd++) {\n score += q[((b*S + s)*h + head)*N*d + i*d + dd] * k[((b*S + s)*h + head)*N*d + j*d + dd];\n }\n score = score * scale + bias[(b*h + head)*N*N + i*N + j];\n if (mask[(b*S + s)*N + j] == 0) score = -INFINITY;\n smem[j] = score;\n }\n // softmax\n float maxv = -INFINITY;\n for(int j=0; j<N; j++) maxv = fmaxf(maxv, smem[j]);\n float sumexp = 0.0f;\n for(int j=0; j<N; j++) {\n smem[j] = expf(smem[j] - maxv);\n sumexp += smem[j];\n }\n for(int j=0; j<N; j++) smem[j] /= sumexp;\n \n // matmul with v and gate\n for(int dd=0; dd<d; dd++) {\n float res = 0.0f;\n for(int j=0; j<N; j++) {\n res += smem[j] * v[((b*S + s)*h + head)*N*d + j*d + dd];\n }\n int out_idx = ((b*S + s)*h + head)*N*d + i*d + dd;\n out[out_idx] = res * gate[out_idx];\n }\n }\n}\n\ntorch::Tensor custom_attention(torch::Tensor q, torch::Tensor k, torch::Tensor v, torch::Tensor bias, torch::Tensor g, torch::Tensor mask, float scale) {\n auto options = torch::TensorOptions().dtype(torch::kFloat32).device(q.device());\n auto out = torch::empty_like(q);\n int B = q.size(0);\n int S = q.size(1);\n int h = q.size(2);\n int N = q.size(3);\n int d = q.size(4);\n dim3 grid(B*S*h);\n int threads = 32;\n attention_kernel<<<grid, threads, N*sizeof(float)>>>(q.data_ptr<float>(), k.data_ptr<float>(), v.data_ptr<float>(), bias.data_ptr<float>(), mask.data_ptr<float>(), out.data_ptr<float>(), g.data_ptr<float>(), B, S, h, N, d, scale);\n return out;\n}\n'''\n\ncpp_source = r'''\ntorch::Tensor custom_attention(torch::Tensor q, torch::Tensor k, torch::Tensor v, torch::Tensor bias, torch::Tensor g, torch::Tensor mask, float scale);\n'''\n\next = load_inline(\n name='custom_attn',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['custom_attention'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, msa_dim, pair_dim, num_heads, dropout=0.0):\n super().__init__()\n self.num_heads = num_heads\n self.head_dim = msa_dim // num_heads\n self.layer_norm = nn.LayerNorm(msa_dim)\n self.query = nn.Linear(msa_dim, msa_dim, bias=False)\n self.key = nn.Linear(msa_dim, msa_dim, bias=False)\n self.value = nn.Linear(msa_dim, msa_dim, bias=False)\n self.gate = nn.Linear(msa_dim, msa_dim)\n self.pair_bias = nn.Linear(pair_dim, num_heads, bias=False)\n self.out_proj = nn.Linear(msa_dim, msa_dim)\n self.dropout = nn.Dropout(dropout)\n\n def forward(self, msa_repr, pair_repr, msa_mask):\n B, S, N, D = msa_repr.shape\n h = self.num_heads\n d = self.head_dim\n x = self.layer_norm(msa_repr)\n q = self.query(x).view(B, S, N, h, d).permute(0, 1, 3, 2, 4).contiguous()\n k = self.key(x).view(B, S, N, h, d).permute(0, 1, 3, 2, 4).contiguous()\n v = self.value(x).view(B, S, N, h, d).permute(0, 1, 3, 2, 4).contiguous()\n g = torch.sigmoid(self.gate(x)).view(B, S, N, h, d).permute(0, 1, 3, 2, 4).contiguous()\n bias = self.pair_bias(pair_repr).permute(0, 3, 1, 2).contiguous()\n scale = 1.0 / math.sqrt(d)\n out = ext.custom_attention(q, k, v, bias, g, msa_mask, scale)\n out = out.permute(0, 1, 3, 2, 4).reshape(B, S, N, D)\n return self.out_proj(out)\n"
}
result ← submit_kernel (fail)
failsubmit_kernel
submit_kernel FAILED: correctness check did not pass ((0 / 5) trials).
correctness_issue: Output mismatch
max_difference: ['0.157683', '0.144166', '0.170426', '0.170848', '0.154564']
avg_difference: ['0.026097', '0.025011', '0.024860', '0.025226', '0.025312']
final submitted kernel (115 lines, 4,289 chars) — full source, untruncated
import torch
import torch.nn as nn
import torch.nn.functional as F
import math
from torch.utils.cpp_extension import load_inline
cuda_source = r'''
#include <torch/extension.h>
#include <cuda_runtime.h>
#include <math.h>
__global__ void attention_kernel(
const float* q, const float* k, const float* v, const float* bias, const float* mask,
float* out, float* gate,
int B, int S, int h, int N, int d, float scale
) {
// Very naive implementation for small N=32
int bsh = blockIdx.x;
int b = bsh / (S * h);
int sh = bsh % (S * h);
int s = sh / h;
int head = sh % h;
extern __shared__ float smem[];
// Compute attn scores
for(int i = threadIdx.x; i < N; i += blockDim.x) {
float sum = 0.0f;
for(int j=0; j<N; j++) {
float score = 0.0f;
for(int dd=0; dd<d; dd++) {
score += q[((b*S + s)*h + head)*N*d + i*d + dd] * k[((b*S + s)*h + head)*N*d + j*d + dd];
}
score = score * scale + bias[(b*h + head)*N*N + i*N + j];
if (mask[(b*S + s)*N + j] == 0) score = -INFINITY;
smem[j] = score;
}
// softmax
float maxv = -INFINITY;
for(int j=0; j<N; j++) maxv = fmaxf(maxv, smem[j]);
float sumexp = 0.0f;
for(int j=0; j<N; j++) {
smem[j] = expf(smem[j] - maxv);
sumexp += smem[j];
}
for(int j=0; j<N; j++) smem[j] /= sumexp;
// matmul with v and gate
for(int dd=0; dd<d; dd++) {
float res = 0.0f;
for(int j=0; j<N; j++) {
res += smem[j] * v[((b*S + s)*h + head)*N*d + j*d + dd];
}
int out_idx = ((b*S + s)*h + head)*N*d + i*d + dd;
out[out_idx] = res * gate[out_idx];
}
}
}
torch::Tensor custom_attention(torch::Tensor q, torch::Tensor k, torch::Tensor v, torch::Tensor bias, torch::Tensor g, torch::Tensor mask, float scale) {
auto options = torch::TensorOptions().dtype(torch::kFloat32).device(q.device());
auto out = torch::empty_like(q);
int B = q.size(0);
int S = q.size(1);
int h = q.size(2);
int N = q.size(3);
int d = q.size(4);
dim3 grid(B*S*h);
int threads = 32;
attention_kernel<<<grid, threads, N*sizeof(float)>>>(q.data_ptr<float>(), k.data_ptr<float>(), v.data_ptr<float>(), bias.data_ptr<float>(), mask.data_ptr<float>(), out.data_ptr<float>(), g.data_ptr<float>(), B, S, h, N, d, scale);
return out;
}
'''
cpp_source = r'''
torch::Tensor custom_attention(torch::Tensor q, torch::Tensor k, torch::Tensor v, torch::Tensor bias, torch::Tensor g, torch::Tensor mask, float scale);
'''
ext = load_inline(
name='custom_attn',
cpp_sources=cpp_source,
cuda_sources=cuda_source,
functions=['custom_attention'],
verbose=False
)
class ModelNew(nn.Module):
def __init__(self, msa_dim, pair_dim, num_heads, dropout=0.0):
super().__init__()
self.num_heads = num_heads
self.head_dim = msa_dim // num_heads
self.layer_norm = nn.LayerNorm(msa_dim)
self.query = nn.Linear(msa_dim, msa_dim, bias=False)
self.key = nn.Linear(msa_dim, msa_dim, bias=False)
self.value = nn.Linear(msa_dim, msa_dim, bias=False)
self.gate = nn.Linear(msa_dim, msa_dim)
self.pair_bias = nn.Linear(pair_dim, num_heads, bias=False)
self.out_proj = nn.Linear(msa_dim, msa_dim)
self.dropout = nn.Dropout(dropout)
def forward(self, msa_repr, pair_repr, msa_mask):
B, S, N, D = msa_repr.shape
h = self.num_heads
d = self.head_dim
x = self.layer_norm(msa_repr)
q = self.query(x).view(B, S, N, h, d).permute(0, 1, 3, 2, 4).contiguous()
k = self.key(x).view(B, S, N, h, d).permute(0, 1, 3, 2, 4).contiguous()
v = self.value(x).view(B, S, N, h, d).permute(0, 1, 3, 2, 4).contiguous()
g = torch.sigmoid(self.gate(x)).view(B, S, N, h, d).permute(0, 1, 3, 2, 4).contiguous()
bias = self.pair_bias(pair_repr).permute(0, 3, 1, 2).contiguous()
scale = 1.0 / math.sqrt(d)
out = ext.custom_attention(q, k, v, bias, g, msa_mask, scale)
out = out.permute(0, 1, 3, 2, 4).reshape(B, S, N, D)
return self.out_proj(out)