grok-4.3 · popcorn · pop_l123_st_gpt
compile_fail 9_MetropolisHastingsStep.py
model
grok-4.3
turns / max
1 / 1
tool calls
1
tokens (total)
3258
wall clock
463.6s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
—
ref runtime
-1.00 μs
speedup
—
started
2026-05-11T09:25:30Z
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 math


class Model(nn.Module):
    """
    Batched Metropolis-Hastings MCMC step.  Given a current state and a
    log-probability function (here a Gaussian target for benchmarking),
    proposes a new state via a Gaussian random walk, computes the
    acceptance ratio, and stochastically accepts or rejects.
    """

    def __init__(self, dim, step_size=0.1):
        super().__init__()
        self.dim = dim
        self.step_size = step_size
        # Target: N(mu, sigma^2 I)
        self.register_buffer("mu", torch.zeros(dim))
        self.register_buffer("sigma", torch.ones(dim))

    def _log_prob(self, x: torch.Tensor) -> torch.Tensor:
        return -0.5 * (((x - self.mu) / self.sigma) ** 2).sum(dim=-1)

    def forward(self, x: torch.Tensor) -> tuple:
        """
        Args:
            x: (B, dim) – current chain states
        Returns:
            (x_new, accepted):
                x_new:    (B, dim) – states after accept/reject
                accepted: (B,)    – boolean mask of accepted proposals
        """
        proposal = x + self.step_size * torch.randn_like(x)
        log_alpha = self._log_prob(proposal) - self._log_prob(x)
        log_u = torch.log(torch.rand(x.shape[0], device=x.device).clamp(min=1e-10))
        accepted = log_u < log_alpha
        x_new = torch.where(accepted.unsqueeze(-1), proposal, x)
        return x_new, accepted


dim = 64
batch_size = 4096


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


def get_init_inputs():
    return [dim]
turn 0 · FINAL · 1 tool call(s) · 463.6s turn · 16.7s LLM · tokens in=1,258 · out=1,490 · total=3,258 · 7,027 chars (click to expand)
turn 0463.6s turn · 16.7s LLM · 447.0s tools · tokens in=1,258 · out=1,490 · total=3,258 · FINAL
reasoning
The task is to replace the forward computation of a PyTorch model with a custom CUDA kernel. The model is a batched Metropolis-Hastings MCMC step. I need to create a `ModelNew` class that replaces the forward computation of a batched Metropolis-Hastings MCMC step with a custom CUDA kernel for speed. The output must be a complete Python file defining `ModelNew`, compiled using `torch.utils.cpp_extension.load_inline` with CUDA code.
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 <curand_kernel.h> __global__ void mh_kernel( const float* __restrict__ x, float* __restrict__ x_new, bool* __restrict__ accepted, const float step_size, const int B, const int D, unsigned long long seed ) { int b = blockIdx.x * blockDim.x + threadIdx.x; if (b >= B) return; curandStatePhilox4_32_10_t state; curand_init(seed, b, 0, &state); float log_prob_x = 0.0f; float log_prob_prop = 0.0f; for (int d = 0; d < D; d++) { float val = x[b * D + d]; log_prob_x += -0.5f * val * val; } for (int d = 0; d < D; d++) { float noise = curand_normal(&state); float prop = x[b * D + d] + step_size * noise; x_new[b * D + d] = prop; // temporarily store proposal log_prob_prop += -0.5f * prop * prop; } float log_alpha = log_prob_prop - log_prob_x; float u = curand_uniform(&state); float log_u = logf(fmaxf(u, 1e-10f)); bool acc = log_u < log_alpha; accepted[b] = acc; if (!acc) { for (int d = 0; d < D; d++) { x_new[b * D + d] = x[b * D + d]; } } } std::vector<torch::Tensor> mh_cuda(torch::Tensor x, float step_size) { const int B = x.size(0); const int D = x.size(1); auto x_new = torch::empty_like(x); auto accepted = torch::empty({B}, x.options().dtype(torch::kBool)); const int threads = 256; const int blocks = (B + threads - 1) / threads; unsigned long long seed = 123456789ULL; mh_kernel<<<blocks, threads>>>( x.data_ptr<float>(), x_new.data_ptr<float>(), accepted.data_ptr<bool>(), step_size, B, D, seed ); return {x_new, accepted}; } ''' _cpp_source = r''' std::vector<torch::Tensor> mh_cuda(torch::Tensor x, float step_size); ''' _mh_module = load_inline( name='mh_cuda', cpp_sources=_cpp_source, cuda_sources=_cuda_source, functions=['mh_cuda'], verbose=False ) class ModelNew(nn.Module): def __init__(self, dim, step_size=0.1): super().__init__() self.dim = dim self.step_size = step_size self.register_buffer("mu", torch.zeros(dim)) self.register_buffer("sigma", torch.ones(dim)) def forward(self, x: torch.Tensor) -> tuple: result = _mh_module.mh_cuda(x.contiguous(), self.step_size) return result[0], result[1] ``` The above is the complete submission. Call the tool.
call → submit_kernel
callsubmit_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n_cuda_source = r'''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <curand_kernel.h>\n\n__global__ void mh_kernel(\n const float* __restrict__ x,\n float* __restrict__ x_new,\n bool* __restrict__ accepted,\n const float step_size,\n const int B,\n const int D,\n unsigned long long seed\n) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= B) return;\n\n curandStatePhilox4_32_10_t state;\n curand_init(seed, b, 0, &state);\n\n float log_prob_x = 0.0f;\n float log_prob_prop = 0.0f;\n\n for (int d = 0; d < D; d++) {\n float val = x[b * D + d];\n log_prob_x += -0.5f * val * val;\n }\n\n for (int d = 0; d < D; d++) {\n float noise = curand_normal(&state);\n float prop = x[b * D + d] + step_size * noise;\n x_new[b * D + d] = prop; // temporarily store proposal\n log_prob_prop += -0.5f * prop * prop;\n }\n\n float log_alpha = log_prob_prop - log_prob_x;\n float u = curand_uniform(&state);\n float log_u = logf(fmaxf(u, 1e-10f));\n bool acc = log_u < log_alpha;\n\n accepted[b] = acc;\n\n if (!acc) {\n for (int d = 0; d < D; d++) {\n x_new[b * D + d] = x[b * D + d];\n }\n }\n}\n\nstd::vector<torch::Tensor> mh_cuda(torch::Tensor x, float step_size) {\n const int B = x.size(0);\n const int D = x.size(1);\n auto x_new = torch::empty_like(x);\n auto accepted = torch::empty({B}, x.options().dtype(torch::kBool));\n\n const int threads = 256;\n const int blocks = (B + threads - 1) / threads;\n\n unsigned long long seed = 123456789ULL;\n\n mh_kernel<<<blocks, threads>>>(\n x.data_ptr<float>(),\n x_new.data_ptr<float>(),\n accepted.data_ptr<bool>(),\n step_size,\n B,\n D,\n seed\n );\n\n return {x_new, accepted};\n}\n'''\n\n_cpp_source = r'''\nstd::vector<torch::Tensor> mh_cuda(torch::Tensor x, float step_size);\n'''\n\n_mh_module = load_inline(\n name='mh_cuda',\n cpp_sources=_cpp_source,\n cuda_sources=_cuda_source,\n functions=['mh_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim, step_size=0.1):\n super().__init__()\n self.dim = dim\n self.step_size = step_size\n self.register_buffer(\"mu\", torch.zeros(dim))\n self.register_buffer(\"sigma\", torch.ones(dim))\n\n def forward(self, x: torch.Tensor) -> tuple:\n result = _mh_module.mh_cuda(x.contiguous(), self.step_size)\n return result[0], result[1]\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 (103 lines, 2,581 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 <curand_kernel.h>

__global__ void mh_kernel(
    const float* __restrict__ x,
    float* __restrict__ x_new,
    bool* __restrict__ accepted,
    const float step_size,
    const int B,
    const int D,
    unsigned long long seed
) {
    int b = blockIdx.x * blockDim.x + threadIdx.x;
    if (b >= B) return;

    curandStatePhilox4_32_10_t state;
    curand_init(seed, b, 0, &state);

    float log_prob_x = 0.0f;
    float log_prob_prop = 0.0f;

    for (int d = 0; d < D; d++) {
        float val = x[b * D + d];
        log_prob_x += -0.5f * val * val;
    }

    for (int d = 0; d < D; d++) {
        float noise = curand_normal(&state);
        float prop = x[b * D + d] + step_size * noise;
        x_new[b * D + d] = prop;  // temporarily store proposal
        log_prob_prop += -0.5f * prop * prop;
    }

    float log_alpha = log_prob_prop - log_prob_x;
    float u = curand_uniform(&state);
    float log_u = logf(fmaxf(u, 1e-10f));
    bool acc = log_u < log_alpha;

    accepted[b] = acc;

    if (!acc) {
        for (int d = 0; d < D; d++) {
            x_new[b * D + d] = x[b * D + d];
        }
    }
}

std::vector<torch::Tensor> mh_cuda(torch::Tensor x, float step_size) {
    const int B = x.size(0);
    const int D = x.size(1);
    auto x_new = torch::empty_like(x);
    auto accepted = torch::empty({B}, x.options().dtype(torch::kBool));

    const int threads = 256;
    const int blocks = (B + threads - 1) / threads;

    unsigned long long seed = 123456789ULL;

    mh_kernel<<<blocks, threads>>>(
        x.data_ptr<float>(),
        x_new.data_ptr<float>(),
        accepted.data_ptr<bool>(),
        step_size,
        B,
        D,
        seed
    );

    return {x_new, accepted};
}
'''

_cpp_source = r'''
std::vector<torch::Tensor> mh_cuda(torch::Tensor x, float step_size);
'''

_mh_module = load_inline(
    name='mh_cuda',
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=['mh_cuda'],
    verbose=False
)

class ModelNew(nn.Module):
    def __init__(self, dim, step_size=0.1):
        super().__init__()
        self.dim = dim
        self.step_size = step_size
        self.register_buffer("mu", torch.zeros(dim))
        self.register_buffer("sigma", torch.ones(dim))

    def forward(self, x: torch.Tensor) -> tuple:
        result = _mh_module.mh_cuda(x.contiguous(), self.step_size)
        return result[0], result[1]
```