grok-4.3 · popcorn · pop_l123_st_gpt
correct 25_OuterProductMean.py
model
grok-4.3
turns / max
1 / 1
tool calls
1
tokens (total)
3860
wall clock
425.0s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
0.12 μs
ref runtime
0.12 μs
speedup
0.97x
started
2026-05-11T08:53:39Z
timingkernelbench.timing
kernel runtime
0.12 μs
ref runtime
0.12 μs
speedup
0.97x
kernel mean / median / std
0.12 / — / 6.78e-04 μs
ref mean / median / std
0.12 / — / 8.16e-04 μs
numerical precisioneval_kernel_against_ref
max abs error
0.00
mean abs error
0.00
max rel error
0.00
mean rel error
0.00
gpu memoryextended_metrics.measure_memory
kernel peak
42.41 MB
ref peak
49.91 MB
kernel / ref
0.85x
kernel launches / fusiontorch.profiler
kernel launches
38
ref launches
75
fusion ratio (ref / kernel)
1.97x
total cuda time
555 μs
energyNVML (pynvml)
kernel mJ / run
229.36
ref mJ / run
0.00
ratio (ref / kernel)
—
avg power
648.88 W
ref avg power
0.00 W
SOL (speed-of-light)heuristic_from_speedup
SOL score (0–1)
0.098
dram util %
—
compute util %
—
bottleneck
unknown
dominant pipe
—
arithmetic intensity
—
ridge point
—
roofline / occupancyheuristic
dram bandwidth
—
dram util %
—
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
—
bottleneck
—
dominant pipe %
—
arithmetic intensity
—
ridge point
—
L1 hit %
—
L2 hit %
—
ld sectors / req
—
st sectors / req
—
regs / thread
—
smem / block
—
block size
—
peak bw
—
peak fp32 tflops
—
peak fp16 tflops
—
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn
import math


class Model(nn.Module):
    """
    Outer-product mean: converts MSA representation to a pair
    representation update.  Used in AlphaFold2 to communicate information
    from the MSA stack to the pair stack.  For each pair (i,j), the
    outer product of projected MSA features at positions i and j is
    averaged over the MSA depth dimension.
    """

    def __init__(self, msa_dim, pair_dim, hidden_dim=32):
        super().__init__()
        self.layer_norm = nn.LayerNorm(msa_dim)
        self.proj_a = nn.Linear(msa_dim, hidden_dim)
        self.proj_b = nn.Linear(msa_dim, hidden_dim)
        self.out_proj = nn.Linear(hidden_dim * hidden_dim, pair_dim)

    def forward(self, msa_repr: torch.Tensor, msa_mask: torch.Tensor) -> torch.Tensor:
        """
        Args:
            msa_repr: (B, S, N, msa_dim) – MSA representation
            msa_mask: (B, S, N)          – 1 for valid rows/positions
        Returns:
            pair update (B, N, N, pair_dim)
        """
        x = self.layer_norm(msa_repr)
        a = self.proj_a(x)  # (B, S, N, H)
        b = self.proj_b(x)

        mask = msa_mask.unsqueeze(-1)  # (B, S, N, 1)
        a = a * mask
        b = b * mask

        outer = torch.einsum("bsih,bsjk->bijhk", a, b)  # (B, N, N, H, H)
        denom = msa_mask.sum(dim=1).unsqueeze(-1).unsqueeze(-1).unsqueeze(-1).clamp(min=1)
        outer = outer / denom  # mean over S

        B, N1, N2, H1, H2 = outer.shape
        outer = outer.reshape(B, N1, N2, H1 * H2)
        return self.out_proj(outer)


msa_dim = 64
pair_dim = 64
hidden_dim = 32
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)
    msa_mask = torch.ones(p.jitter_int(batch_size), p.jitter_int(msa_depth), p.jitter_int(seq_len))
    return [msa_repr, msa_mask]


def get_init_inputs():
    return [msa_dim, pair_dim, hidden_dim]
turn 0 · FINAL · 1 tool call(s) · 425.0s turn · 33.2s LLM · tokens in=1,429 · out=1,585 · total=3,860 · 7,467 chars (click to expand)
turn 0425.0s turn · 33.2s LLM · 391.9s tools · tokens in=1,429 · out=1,585 · total=3,860 · FINAL
reasoning
The task is to replace the forward computation of the PyTorch model with a custom CUDA kernel. The model is an outer-product mean from AlphaFold2. I need to implement a custom CUDA kernel for the outer-product mean computation in the AlphaFold2 model. The `ModelNew` class should match the original constructor and forward signature, replacing the forward computation with a CUDA kernel for speed, while ensuring equivalent outputs. The computation involves: - Applying layer norm to `msa (B, S, N, msa_dim)` - Two linear projections to `hidden_dim`: `a` and `b` - Multiplying by an unsqueezed `msa_mask` - Performing an einsum: `outer = torch.
assistant```python 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 outer_mean_kernel( const float* __restrict__ a, const float* __restrict__ b, const float* __restrict__ denom, float* __restrict__ out, int B, int S, int N, int H ) { int idx = blockIdx.x * blockDim.x + threadIdx.x; int total = B * N * N * H * H; if (idx >= total) return; int h2 = idx % H; int h1 = (idx / H) % H; int j = (idx / (H*H)) % N; int i = (idx / (H*H*N)) % N; int batch = idx / (H*H*N*N); float sum = 0.0f; float d = denom[batch * N + i]; if (d < 1.0f) d = 1.0f; for (int s = 0; s < S; s++) { float aval = a[((batch * S + s) * N + i) * H + h1]; float bval = b[((batch * S + s) * N + j) * H + h2]; sum += aval * bval; } out[idx] = sum / d; } torch::Tensor outer_mean_cuda(torch::Tensor a, torch::Tensor b, torch::Tensor msa_mask) { auto B = a.size(0); auto S = a.size(1); auto N = a.size(2); auto H = a.size(3); auto denom = msa_mask.sum(1); // (B, N) auto options = torch::TensorOptions().dtype(torch::kFloat32).device(a.device()); auto out = torch::zeros({B, N, N, H*H}, options); int total = B * N * N * H * H; int threads = 256; int blocks = (total + threads - 1) / threads; outer_mean_kernel<<<blocks, threads>>>( a.data_ptr<float>(), b.data_ptr<float>(), denom.data_ptr<float>(), out.data_ptr<float>(), B, S, N, H ); return out; } ''' cpp_source = ''' torch::Tensor outer_mean_cuda(torch::Tensor a, torch::Tensor b, torch::Tensor msa_mask); ''' module = load_inline( name='outer_mean_ext', cpp_sources=[cpp_source], cuda_sources=[cuda_source], functions=['outer_mean_cuda'], verbose=False ) class ModelNew(nn.Module): def __init__(self, msa_dim, pair_dim, hidden_dim=32): super().__init__() self.layer_norm = nn.LayerNorm(msa_dim) self.proj_a = nn.Linear(msa_dim, hidden_dim) self.proj_b = nn.Linear(msa_dim, hidden_dim) self.out_proj = nn.Linear(hidden_dim * hidden_dim, pair_dim) self.ext = module def forward(self, msa_repr: torch.Tensor, msa_mask: torch.Tensor) -> torch.Tensor: x = self.layer_norm(msa_repr) a = self.proj_a(x) b = self.proj_b(x) mask = msa_mask.unsqueeze(-1) a = a * mask b = b * mask outer = self.ext.outer_mean_cuda(a, b, msa_mask) return self.out_proj(outer) ```
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 outer_mean_kernel(\n const float* __restrict__ a,\n const float* __restrict__ b,\n const float* __restrict__ denom,\n float* __restrict__ out,\n int B, int S, int N, int H\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * N * N * H * H;\n if (idx >= total) return;\n\n int h2 = idx % H;\n int h1 = (idx / H) % H;\n int j = (idx / (H*H)) % N;\n int i = (idx / (H*H*N)) % N;\n int batch = idx / (H*H*N*N);\n\n float sum = 0.0f;\n float d = denom[batch * N + i];\n if (d < 1.0f) d = 1.0f;\n\n for (int s = 0; s < S; s++) {\n float aval = a[((batch * S + s) * N + i) * H + h1];\n float bval = b[((batch * S + s) * N + j) * H + h2];\n sum += aval * bval;\n }\n out[idx] = sum / d;\n}\n\ntorch::Tensor outer_mean_cuda(torch::Tensor a, torch::Tensor b, torch::Tensor msa_mask) {\n auto B = a.size(0);\n auto S = a.size(1);\n auto N = a.size(2);\n auto H = a.size(3);\n\n auto denom = msa_mask.sum(1); // (B, N)\n auto options = torch::TensorOptions().dtype(torch::kFloat32).device(a.device());\n auto out = torch::zeros({B, N, N, H*H}, options);\n\n int total = B * N * N * H * H;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n\n outer_mean_kernel<<<blocks, threads>>>(\n a.data_ptr<float>(),\n b.data_ptr<float>(),\n denom.data_ptr<float>(),\n out.data_ptr<float>(),\n B, S, N, H\n );\n\n return out;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor outer_mean_cuda(torch::Tensor a, torch::Tensor b, torch::Tensor msa_mask);\n'''\n\nmodule = load_inline(\n name='outer_mean_ext',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['outer_mean_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, msa_dim, pair_dim, hidden_dim=32):\n super().__init__()\n self.layer_norm = nn.LayerNorm(msa_dim)\n self.proj_a = nn.Linear(msa_dim, hidden_dim)\n self.proj_b = nn.Linear(msa_dim, hidden_dim)\n self.out_proj = nn.Linear(hidden_dim * hidden_dim, pair_dim)\n self.ext = module\n\n def forward(self, msa_repr: torch.Tensor, msa_mask: torch.Tensor) -> torch.Tensor:\n x = self.layer_norm(msa_repr)\n a = self.proj_a(x)\n b = self.proj_b(x)\n\n mask = msa_mask.unsqueeze(-1)\n a = a * mask\n b = b * mask\n\n outer = self.ext.outer_mean_cuda(a, b, msa_mask)\n return self.out_proj(outer)\n" }
result ← submit_kernel (ok)
oksubmit_kernel submit_kernel PASSED: (5 / 5) correctness trials all passed. Kernel runtime: 0.12 μs Runtime stats: mean=0.12μs median=?μs std=0.00μs Numerical precision: max_abs_err=0.00e+00 mean_abs_err=0.00e+00 max_rel_err=0.00e+00 Memory: 42.41 MB peak (ref: 49.91 MB, ratio: 0.8497x) Kernel launches: 38 (ref: 75, fusion ratio: 1.9737) Energy: 229.36 mJ/run (ref: 0.0 mJ, ratio: -1x) SOL score: 0.0975
final submitted kernel (97 lines, 2,639 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 outer_mean_kernel(
    const float* __restrict__ a,
    const float* __restrict__ b,
    const float* __restrict__ denom,
    float* __restrict__ out,
    int B, int S, int N, int H
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int total = B * N * N * H * H;
    if (idx >= total) return;

    int h2 = idx % H;
    int h1 = (idx / H) % H;
    int j = (idx / (H*H)) % N;
    int i = (idx / (H*H*N)) % N;
    int batch = idx / (H*H*N*N);

    float sum = 0.0f;
    float d = denom[batch * N + i];
    if (d < 1.0f) d = 1.0f;

    for (int s = 0; s < S; s++) {
        float aval = a[((batch * S + s) * N + i) * H + h1];
        float bval = b[((batch * S + s) * N + j) * H + h2];
        sum += aval * bval;
    }
    out[idx] = sum / d;
}

torch::Tensor outer_mean_cuda(torch::Tensor a, torch::Tensor b, torch::Tensor msa_mask) {
    auto B = a.size(0);
    auto S = a.size(1);
    auto N = a.size(2);
    auto H = a.size(3);

    auto denom = msa_mask.sum(1); // (B, N)
    auto options = torch::TensorOptions().dtype(torch::kFloat32).device(a.device());
    auto out = torch::zeros({B, N, N, H*H}, options);

    int total = B * N * N * H * H;
    int threads = 256;
    int blocks = (total + threads - 1) / threads;

    outer_mean_kernel<<<blocks, threads>>>(
        a.data_ptr<float>(),
        b.data_ptr<float>(),
        denom.data_ptr<float>(),
        out.data_ptr<float>(),
        B, S, N, H
    );

    return out;
}
'''

cpp_source = '''
torch::Tensor outer_mean_cuda(torch::Tensor a, torch::Tensor b, torch::Tensor msa_mask);
'''

module = load_inline(
    name='outer_mean_ext',
    cpp_sources=[cpp_source],
    cuda_sources=[cuda_source],
    functions=['outer_mean_cuda'],
    verbose=False
)

class ModelNew(nn.Module):
    def __init__(self, msa_dim, pair_dim, hidden_dim=32):
        super().__init__()
        self.layer_norm = nn.LayerNorm(msa_dim)
        self.proj_a = nn.Linear(msa_dim, hidden_dim)
        self.proj_b = nn.Linear(msa_dim, hidden_dim)
        self.out_proj = nn.Linear(hidden_dim * hidden_dim, pair_dim)
        self.ext = module

    def forward(self, msa_repr: torch.Tensor, msa_mask: torch.Tensor) -> torch.Tensor:
        x = self.layer_norm(msa_repr)
        a = self.proj_a(x)
        b = self.proj_b(x)

        mask = msa_mask.unsqueeze(-1)
        a = a * mask
        b = b * mask

        outer = self.ext.outer_mean_cuda(a, b, msa_mask)
        return self.out_proj(outer)