grok-4.3 · popcorn · pop_l123_st_gpt
compile_fail 43_ContactMapPrediction.py
model
grok-4.3
turns / max
1 / 1
tool calls
1
tokens (total)
4645
wall clock
557.2s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
—
ref runtime
-1.00 μs
speedup
—
started
2026-05-11T09:54:16Z
compile error: Syntax error in custom generated code or ModelNew not found
timingkernelbench.timing
kernel runtime
—
ref runtime
—
speedup
—
kernel mean / median / std
—
ref mean / median / std
—
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn
import torch.nn.functional as F


class Model(nn.Module):
    """
    Contact map prediction head.  Takes residue-pair embeddings and
    predicts a symmetric binary contact probability matrix (residue i
    within 8 Å of residue j).  Uses a small MLP on the symmetrised pair
    representation followed by a sigmoid.  Common final stage in protein
    structure prediction pipelines.
    """

    def __init__(self, pair_dim, hidden_dim):
        super().__init__()
        self.layer_norm = nn.LayerNorm(pair_dim)
        self.mlp = nn.Sequential(
            nn.Linear(pair_dim, hidden_dim),
            nn.ReLU(),
            nn.Linear(hidden_dim, hidden_dim),
            nn.ReLU(),
            nn.Linear(hidden_dim, 1),
        )

    def forward(self, pair_repr: torch.Tensor) -> torch.Tensor:
        """
        Args:
            pair_repr: (B, N, N, pair_dim)
        Returns:
            contact_probs: (B, N, N) – symmetric contact probability
        """
        x = self.layer_norm(pair_repr)
        # Symmetrise
        x = (x + x.permute(0, 2, 1, 3)) / 2.0
        logits = self.mlp(x).squeeze(-1)  # (B, N, N)
        return torch.sigmoid(logits)


pair_dim = 64
hidden_dim = 32
seq_len = 64
batch_size = 4


def get_inputs():
    return [torch.randn(popcorn_pri.jitter_int(batch_size), popcorn_pri.jitter_int(seq_len), popcorn_pri.jitter_int(seq_len), pair_dim)]


def get_init_inputs():
    return [pair_dim, hidden_dim]
turn 0 · FINAL · 1 tool call(s) · 557.2s turn · 17.9s LLM · tokens in=1,221 · out=3,038 · total=4,645 · 11,642 chars (click to expand)
turn 0557.2s turn · 17.9s LLM · 539.3s tools · tokens in=1,221 · out=3,038 · total=4,645 · FINAL
reasoning
The task is to replace the forward computation of the PyTorch model with a custom CUDA kernel. The model is a contact map prediction head with LayerNorm, symmetrization, MLP with ReLU and sigmoid.
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 <cmath> __global__ void fused_contact_kernel( const float* __restrict__ pair_repr, const float* __restrict__ ln_weight, const float* __restrict__ ln_bias, const float* __restrict__ w1, const float* __restrict__ b1, const float* __restrict__ w2, const float* __restrict__ b2, const float* __restrict__ w3, const float* __restrict__ b3, float* __restrict__ output, int B, int N, int D, int H ) { int idx = blockIdx.x * blockDim.x + threadIdx.x; int total = B * N * N; if (idx >= total) return; int b = idx / (N * N); int rem = idx % (N * N); int i = rem / N; int j = rem % N; // LayerNorm for position (b,i,j) over D const float* x_ptr = pair_repr + (b * N * N + i * N + j) * D; float mean = 0.0f; for (int d = 0; d < D; d++) mean += x_ptr[d]; mean /= D; float var = 0.0f; for (int d = 0; d < D; d++) { float diff = x_ptr[d] - mean; var += diff * diff; } var /= D; float inv_std = rsqrtf(var + 1e-5f); // Symmetrized LN: average with (j,i) const float* x_t_ptr = pair_repr + (b * N * N + j * N + i) * D; float mean_t = 0.0f; for (int d = 0; d < D; d++) mean_t += x_t_ptr[d]; mean_t /= D; float var_t = 0.0f; for (int d = 0; d < D; d++) { float diff = x_t_ptr[d] - mean_t; var_t += diff * diff; } var_t /= D; float inv_std_t = rsqrtf(var_t + 1e-5f); // Compute averaged normalized features float norm_feat[64]; for (int d = 0; d < D; d++) { float n1 = (x_ptr[d] - mean) * inv_std * ln_weight[d] + ln_bias[d]; float n2 = (x_t_ptr[d] - mean_t) * inv_std_t * ln_weight[d] + ln_bias[d]; norm_feat[d] = (n1 + n2) * 0.5f; } // MLP: D -> H -> H -> 1 float h1[32]; for (int h = 0; h < H; h++) { float sum = b1[h]; for (int d = 0; d < D; d++) { sum += norm_feat[d] * w1[h * D + d]; } h1[h] = fmaxf(sum, 0.0f); // ReLU } float h2[32]; for (int h = 0; h < H; h++) { float sum = b2[h]; for (int hh = 0; hh < H; hh++) { sum += h1[hh] * w2[h * H + hh]; } h2[h] = fmaxf(sum, 0.0f); // ReLU } float logit = b3[0]; for (int h = 0; h < H; h++) { logit += h2[h] * w3[h]; } output[idx] = 1.0f / (1.0f + expf(-logit)); // sigmoid } torch::Tensor fused_contact_forward( torch::Tensor pair_repr, torch::Tensor ln_weight, torch::Tensor ln_bias, torch::Tensor w1, torch::Tensor b1, torch::Tensor w2, torch::Tensor b2, torch::Tensor w3, torch::Tensor b3 ) { int B = pair_repr.size(0); int N = pair_repr.size(1); int D = pair_repr.size(3); int H = w1.size(0); auto output = torch::empty({B, N, N}, pair_repr.options()); int total = B * N * N; int threads = 256; int blocks = (total + threads - 1) / threads; fused_contact_kernel<<<blocks, threads>>>( pair_repr.data_ptr<float>(), ln_weight.data_ptr<float>(), ln_bias.data_ptr<float>(), w1.data_ptr<float>(), b1.data_ptr<float>(), w2.data_ptr<float>(), b2.data_ptr<float>(), w3.data_ptr<float>(), b3.data_ptr<float>(), output.data_ptr<float>(), B, N, D, H ); return output; } ''' cpp_source = ''' torch::Tensor fused_contact_forward( torch::Tensor pair_repr, torch::Tensor ln_weight, torch::Tensor ln_bias, torch::Tensor w1, torch::Tensor b1, torch::Tensor w2, torch::Tensor b2, torch::Tensor w3, torch::Tensor b3 ); ''' module = load_inline( name='fused_contact', cpp_sources=[cpp_source], cuda_sources=[cuda_source], functions=['fused_contact_forward'], verbose=False, extra_cuda_cflags=['-O3'] ) class ModelNew(nn.Module): def __init__(self, pair_dim, hidden_dim): super().__init__() self.layer_norm = nn.LayerNorm(pair_dim) self.mlp = nn.Sequential( nn.Linear(pair_dim, hidden_dim), nn.ReLU(), nn.Linear(hidden_dim, hidden_dim), nn.ReLU(), nn.Linear(hidden_dim, 1), ) self.pair_dim = pair_dim self.hidden_dim = hidden_dim def forward(self, pair_repr: torch.Tensor) -> torch.Tensor: ln = self.layer_norm mlp = self.mlp return module.fused_contact_forward( pair_repr.contiguous(), ln.weight, ln.bias, mlp[0].weight, mlp[0].bias, mlp[2].weight, mlp[2].bias, mlp[4].weight, mlp[4].bias ) ''' # noqa: E501 submit_kernel(kernel_code=open(__file__).read() if False else globals().get('__kernel_code__', __import__('inspect').getsource(__import__('__main__')))) # placeholder, actual submit below ``` Wait, the final submission must be correct. I am fixing the submit.
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 <cmath>\n\n__global__ void fused_contact_kernel(\n const float* __restrict__ pair_repr,\n const float* __restrict__ ln_weight,\n const float* __restrict__ ln_bias,\n const float* __restrict__ w1,\n const float* __restrict__ b1,\n const float* __restrict__ w2,\n const float* __restrict__ b2,\n const float* __restrict__ w3,\n const float* __restrict__ b3,\n float* __restrict__ output,\n int B, int N, int D, int H\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * N * N;\n if (idx >= total) return;\n\n int b = idx / (N * N);\n int rem = idx % (N * N);\n int i = rem / N;\n int j = rem % N;\n\n // LayerNorm for position (b,i,j) over D\n const float* x_ptr = pair_repr + (b * N * N + i * N + j) * D;\n float mean = 0.0f;\n for (int d = 0; d < D; d++) mean += x_ptr[d];\n mean /= D;\n\n float var = 0.0f;\n for (int d = 0; d < D; d++) {\n float diff = x_ptr[d] - mean;\n var += diff * diff;\n }\n var /= D;\n float inv_std = rsqrtf(var + 1e-5f);\n\n // Symmetrized LN: average with (j,i)\n const float* x_t_ptr = pair_repr + (b * N * N + j * N + i) * D;\n float mean_t = 0.0f;\n for (int d = 0; d < D; d++) mean_t += x_t_ptr[d];\n mean_t /= D;\n\n float var_t = 0.0f;\n for (int d = 0; d < D; d++) {\n float diff = x_t_ptr[d] - mean_t;\n var_t += diff * diff;\n }\n var_t /= D;\n float inv_std_t = rsqrtf(var_t + 1e-5f);\n\n // Compute averaged normalized features\n float norm_feat[64];\n for (int d = 0; d < D; d++) {\n float n1 = (x_ptr[d] - mean) * inv_std * ln_weight[d] + ln_bias[d];\n float n2 = (x_t_ptr[d] - mean_t) * inv_std_t * ln_weight[d] + ln_bias[d];\n norm_feat[d] = (n1 + n2) * 0.5f;\n }\n\n // MLP: D -> H -> H -> 1\n float h1[32];\n for (int h = 0; h < H; h++) {\n float sum = b1[h];\n for (int d = 0; d < D; d++) {\n sum += norm_feat[d] * w1[h * D + d];\n }\n h1[h] = fmaxf(sum, 0.0f); // ReLU\n }\n\n float h2[32];\n for (int h = 0; h < H; h++) {\n float sum = b2[h];\n for (int hh = 0; hh < H; hh++) {\n sum += h1[hh] * w2[h * H + hh];\n }\n h2[h] = fmaxf(sum, 0.0f); // ReLU\n }\n\n float logit = b3[0];\n for (int h = 0; h < H; h++) {\n logit += h2[h] * w3[h];\n }\n\n output[idx] = 1.0f / (1.0f + expf(-logit)); // sigmoid\n}\n\ntorch::Tensor fused_contact_forward(\n torch::Tensor pair_repr,\n torch::Tensor ln_weight,\n torch::Tensor ln_bias,\n torch::Tensor w1, torch::Tensor b1,\n torch::Tensor w2, torch::Tensor b2,\n torch::Tensor w3, torch::Tensor b3\n) {\n int B = pair_repr.size(0);\n int N = pair_repr.size(1);\n int D = pair_repr.size(3);\n int H = w1.size(0);\n\n auto output = torch::empty({B, N, N}, pair_repr.options());\n\n int total = B * N * N;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n\n fused_contact_kernel<<<blocks, threads>>>(\n pair_repr.data_ptr<float>(),\n ln_weight.data_ptr<float>(),\n ln_bias.data_ptr<float>(),\n w1.data_ptr<float>(),\n b1.data_ptr<float>(),\n w2.data_ptr<float>(),\n b2.data_ptr<float>(),\n w3.data_ptr<float>(),\n b3.data_ptr<float>(),\n output.data_ptr<float>(),\n B, N, D, H\n );\n\n return output;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor fused_contact_forward(\n torch::Tensor pair_repr,\n torch::Tensor ln_weight,\n torch::Tensor ln_bias,\n torch::Tensor w1, torch::Tensor b1,\n torch::Tensor w2, torch::Tensor b2,\n torch::Tensor w3, torch::Tensor b3\n);\n'''\n\nmodule = load_inline(\n name='fused_contact',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['fused_contact_forward'],\n verbose=False,\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, pair_dim, hidden_dim):\n super().__init__()\n self.layer_norm = nn.LayerNorm(pair_dim)\n self.mlp = nn.Sequential(\n nn.Linear(pair_dim, hidden_dim),\n nn.ReLU(),\n nn.Linear(hidden_dim, hidden_dim),\n nn.ReLU(),\n nn.Linear(hidden_dim, 1),\n )\n self.pair_dim = pair_dim\n self.hidden_dim = hidden_dim\n\n def forward(self, pair_repr: torch.Tensor) -> torch.Tensor:\n ln = self.layer_norm\n mlp = self.mlp\n return module.fused_contact_forward(\n pair_repr.contiguous(),\n ln.weight, ln.bias,\n mlp[0].weight, mlp[0].bias,\n mlp[2].weight, mlp[2].bias,\n mlp[4].weight, mlp[4].bias\n )\n'''" }
result ← submit_kernel (fail)
failsubmit_kernel submit_kernel FAILED: kernel did not compile. Syntax error in custom generated code or ModelNew not found
final submitted kernel (177 lines, 4,796 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 <cmath>

__global__ void fused_contact_kernel(
    const float* __restrict__ pair_repr,
    const float* __restrict__ ln_weight,
    const float* __restrict__ ln_bias,
    const float* __restrict__ w1,
    const float* __restrict__ b1,
    const float* __restrict__ w2,
    const float* __restrict__ b2,
    const float* __restrict__ w3,
    const float* __restrict__ b3,
    float* __restrict__ output,
    int B, int N, int D, int H
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int total = B * N * N;
    if (idx >= total) return;

    int b = idx / (N * N);
    int rem = idx % (N * N);
    int i = rem / N;
    int j = rem % N;

    // LayerNorm for position (b,i,j) over D
    const float* x_ptr = pair_repr + (b * N * N + i * N + j) * D;
    float mean = 0.0f;
    for (int d = 0; d < D; d++) mean += x_ptr[d];
    mean /= D;

    float var = 0.0f;
    for (int d = 0; d < D; d++) {
        float diff = x_ptr[d] - mean;
        var += diff * diff;
    }
    var /= D;
    float inv_std = rsqrtf(var + 1e-5f);

    // Symmetrized LN: average with (j,i)
    const float* x_t_ptr = pair_repr + (b * N * N + j * N + i) * D;
    float mean_t = 0.0f;
    for (int d = 0; d < D; d++) mean_t += x_t_ptr[d];
    mean_t /= D;

    float var_t = 0.0f;
    for (int d = 0; d < D; d++) {
        float diff = x_t_ptr[d] - mean_t;
        var_t += diff * diff;
    }
    var_t /= D;
    float inv_std_t = rsqrtf(var_t + 1e-5f);

    // Compute averaged normalized features
    float norm_feat[64];
    for (int d = 0; d < D; d++) {
        float n1 = (x_ptr[d] - mean) * inv_std * ln_weight[d] + ln_bias[d];
        float n2 = (x_t_ptr[d] - mean_t) * inv_std_t * ln_weight[d] + ln_bias[d];
        norm_feat[d] = (n1 + n2) * 0.5f;
    }

    // MLP: D -> H -> H -> 1
    float h1[32];
    for (int h = 0; h < H; h++) {
        float sum = b1[h];
        for (int d = 0; d < D; d++) {
            sum += norm_feat[d] * w1[h * D + d];
        }
        h1[h] = fmaxf(sum, 0.0f);  // ReLU
    }

    float h2[32];
    for (int h = 0; h < H; h++) {
        float sum = b2[h];
        for (int hh = 0; hh < H; hh++) {
            sum += h1[hh] * w2[h * H + hh];
        }
        h2[h] = fmaxf(sum, 0.0f);  // ReLU
    }

    float logit = b3[0];
    for (int h = 0; h < H; h++) {
        logit += h2[h] * w3[h];
    }

    output[idx] = 1.0f / (1.0f + expf(-logit));  // sigmoid
}

torch::Tensor fused_contact_forward(
    torch::Tensor pair_repr,
    torch::Tensor ln_weight,
    torch::Tensor ln_bias,
    torch::Tensor w1, torch::Tensor b1,
    torch::Tensor w2, torch::Tensor b2,
    torch::Tensor w3, torch::Tensor b3
) {
    int B = pair_repr.size(0);
    int N = pair_repr.size(1);
    int D = pair_repr.size(3);
    int H = w1.size(0);

    auto output = torch::empty({B, N, N}, pair_repr.options());

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

    fused_contact_kernel<<<blocks, threads>>>(
        pair_repr.data_ptr<float>(),
        ln_weight.data_ptr<float>(),
        ln_bias.data_ptr<float>(),
        w1.data_ptr<float>(),
        b1.data_ptr<float>(),
        w2.data_ptr<float>(),
        b2.data_ptr<float>(),
        w3.data_ptr<float>(),
        b3.data_ptr<float>(),
        output.data_ptr<float>(),
        B, N, D, H
    );

    return output;
}
'''

cpp_source = '''
torch::Tensor fused_contact_forward(
    torch::Tensor pair_repr,
    torch::Tensor ln_weight,
    torch::Tensor ln_bias,
    torch::Tensor w1, torch::Tensor b1,
    torch::Tensor w2, torch::Tensor b2,
    torch::Tensor w3, torch::Tensor b3
);
'''

module = load_inline(
    name='fused_contact',
    cpp_sources=[cpp_source],
    cuda_sources=[cuda_source],
    functions=['fused_contact_forward'],
    verbose=False,
    extra_cuda_cflags=['-O3']
)

class ModelNew(nn.Module):
    def __init__(self, pair_dim, hidden_dim):
        super().__init__()
        self.layer_norm = nn.LayerNorm(pair_dim)
        self.mlp = nn.Sequential(
            nn.Linear(pair_dim, hidden_dim),
            nn.ReLU(),
            nn.Linear(hidden_dim, hidden_dim),
            nn.ReLU(),
            nn.Linear(hidden_dim, 1),
        )
        self.pair_dim = pair_dim
        self.hidden_dim = hidden_dim

    def forward(self, pair_repr: torch.Tensor) -> torch.Tensor:
        ln = self.layer_norm
        mlp = self.mlp
        return module.fused_contact_forward(
            pair_repr.contiguous(),
            ln.weight, ln.bias,
            mlp[0].weight, mlp[0].bias,
            mlp[2].weight, mlp[2].bias,
            mlp[4].weight, mlp[4].bias
        )
'''