compile_fail
9_MetropolisHastingsStep.py
started
2026-05-11T19:37:15Z
compile error: Syntax error in custom generated code or ModelNew not found
timingkernelbench.timing
kernel 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 039.4s turn · 39.4s LLM · 0.0s tools · tokens in=2,156 · out=677 · total=4,192
reasoning
The task is to replace the forward computation of a PyTorch model with a custom CUDA kernel. The model performs a batched Metropolis-Hastings MCMC step.
The task requires replacing the forward pass of a PyTorch model that implements a batched Metropolis-Hastings MCMC step with a custom CUDA kernel. The model takes a current state tensor x of shape (batch_size, dim), proposes a new state by adding noise scaled by step_size, computes the log acceptance ratio, decides Accepted acceptance based on a uniform random variable, and selects between the proposal and the current state element-wise.
To achieve this with a CUDA kernel, I need to fuse the computations for speed, performing them directly on the GPU. Since the computation is stochastic due to random number generation, achieving exact numerical equivalence is challenging because Py PyTorch uses a default random number generator, while a custom kernel would use something like cuRAND, potentially generating different values. The tolerance for floating-point comparisons is set to 1e-4 for fp32, which accommodates small differences, likely accounting for the randomness. The signature of forward must remain identical: def forward forward(self, x): ...
To ensure outputs match closely, I'll generate the proposal and random numbers in PyTorch on the CPU or GPU side before entering the kernel, passing these precomputed values into the kernel. Inside forward, I'll compute proposal = x + self.step_size * torch.
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 <math.h>
__global__ void mh_kernel(
const float* __restrict__ x,
const float* __restrict__ proposal,
const float* __restrict__ log_u,
float* __restrict__ x_new,
bool* __restrict__ accepted,
int B,
int D
) {
int b = blockIdx.x * blockDim.x + threadIdx.x;
if (b >= B) return;
float sum_x2 = 0.0f;
float sum_p2 = 0.0f;
for (int i = 0; i < D; ++i) {
float xv = x[b * D + i];
float pv = proposal[b * D + i];
sum_x2 += xv * xv;
sum_p2 += pv * pv;
}
float log_alpha = -0.5f * (sum_p2 - sum_x2);
bool acc = log_u[b] < log_alpha;
accepted[b] = acc;
for (int i = 0; i < D; ++i) {
x_new[b * D + i] = acc ? proposal[b * D + i] : x[b * D + i];
}
}
std::vector<torch::Tensor> mh_step_cuda(
torch::Tensor x,
torch::Tensor proposal,
torch::Tensor log_u
) {
int B = x.size(0);
int D = x.size(1);
auto x_new = torch::empty_like(x);
auto accepted = torch::empty({B}, x.options().dtype(torch::kBool));
int threads = 256;
int blocks = (B + threads - 1) / threads;
mh_kernel<<<blocks, threads>>>(
x.data_ptr<float>(),
proposal.data_ptr<float>(),
log_u.data_ptr<float>(),
x_new.data_ptr<float>(),
accepted.data_ptr<bool>(),
B,
D
);
return {x_new, accepted};
}
'''
cpp_source = r'''
std::vector<torch::Tensor> mh_step_cuda(torch::Tensor x, torch::Tensor proposal, torch::Tensor log_u);
'''
module = load_inline(
name='mh_cuda',
cpp_sources=[cpp_source],
cuda_sources=[cuda_source],
functions=['mh_step_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:
proposal = x + self.step_size * torch.randn_like(x)
log_u = torch.log(torch.rand(x.shape[0], device=x.device).clamp(min=1e-10))
x_new, accepted = module.mh_step_cuda(x, proposal, log_u)
return x_new, accepted
```
turn 1124.4s turn · 9.3s LLM · 115.1s tools · tokens in=4,260 · out=684 · total=4,992
reasoning
The system requires me to use tools via function calls. My previous response was the kernel code, but now I need to actually use the tools to compile and test it.
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.h>\n#include <cuda_runtime.h>\n#include <math.h>\n\n__global__ void mh_kernel(\n const float* __restrict__ x,\n const float* __restrict__ proposal,\n const float* __restrict__ log_u,\n float* __restrict__ x_new,\n bool* __restrict__ accepted,\n int B,\n int D\n) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= B) return;\n\n float sum_x2 = 0.0f;\n float sum_p2 = 0.0f;\n for (int i = 0; i < D; ++i) {\n float xv = x[b * D + i];\n float pv = proposal[b * D + i];\n sum_x2 += xv * xv;\n sum_p2 += pv * pv;\n }\n\n float log_alpha = -0.5f * (sum_p2 - sum_x2);\n bool acc = log_u[b] < log_alpha;\n\n accepted[b] = acc;\n for (int i = 0; i < D; ++i) {\n x_new[b * D + i] = acc ? proposal[b * D + i] : x[b * D + i];\n }\n}\n\nstd::vector<torch::Tensor> mh_step_cuda(\n torch::Tensor x,\n torch::Tensor proposal,\n torch::Tensor log_u\n) {\n int B = x.size(0);\n 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 int threads = 256;\n int blocks = (B + threads - 1) / threads;\n\n mh_kernel<<<blocks, threads>>>(\n x.data_ptr<float>(),\n proposal.data_ptr<float>(),\n log_u.data_ptr<float>(),\n x_new.data_ptr<float>(),\n accepted.data_ptr<bool>(),\n B,\n D\n );\n\n return {x_new, accepted};\n}\n'''\n\ncpp_source = r'''\nstd::vector<torch::Tensor> mh_step_cuda(torch::Tensor x, torch::Tensor proposal, torch::Tensor log_u);\n'''\n\nmodule = load_inline(\n name='mh_cuda',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['mh_step_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 proposal = x + self.step_size * torch.randn_like(x)\n log_u = torch.log(torch.rand(x.shape[0], device=x.device).clamp(min=1e-10))\n x_new, accepted = module.mh_step_cuda(x, proposal, log_u)\n return x_new, accepted\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 276.5s turn · 10.1s LLM · 66.4s tools · tokens in=5,009 · out=685 · total=5,707
reasoning
The kernel compiled successfully. Now I should run the correctness check.
call → run_correctness
callrun_correctness
{
"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 <math.h>\n\n__global__ void mh_kernel(\n const float* __restrict__ x,\n const float* __restrict__ proposal,\n const float* __restrict__ log_u,\n float* __restrict__ x_new,\n bool* __restrict__ accepted,\n int B,\n int D\n) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= B) return;\n\n float sum_x2 = 0.0f;\n float sum_p2 = 0.0f;\n for (int i = 0; i < D; ++i) {\n float xv = x[b * D + i];\n float pv = proposal[b * D + i];\n sum_x2 += xv * xv;\n sum_p2 += pv * pv;\n }\n\n float log_alpha = -0.5f * (sum_p2 - sum_x2);\n bool acc = log_u[b] < log_alpha;\n\n accepted[b] = acc;\n for (int i = 0; i < D; ++i) {\n x_new[b * D + i] = acc ? proposal[b * D + i] : x[b * D + i];\n }\n}\n\nstd::vector<torch::Tensor> mh_step_cuda(\n torch::Tensor x,\n torch::Tensor proposal,\n torch::Tensor log_u\n) {\n int B = x.size(0);\n 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 int threads = 256;\n int blocks = (B + threads - 1) / threads;\n\n mh_kernel<<<blocks, threads>>>(\n x.data_ptr<float>(),\n proposal.data_ptr<float>(),\n log_u.data_ptr<float>(),\n x_new.data_ptr<float>(),\n accepted.data_ptr<bool>(),\n B,\n D\n );\n\n return {x_new, accepted};\n}\n'''\n\ncpp_source = r'''\nstd::vector<torch::Tensor> mh_step_cuda(torch::Tensor x, torch::Tensor proposal, torch::Tensor log_u);\n'''\n\nmodule = load_inline(\n name='mh_cuda',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['mh_step_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 proposal = x + self.step_size * torch.randn_like(x)\n log_u = torch.log(torch.rand(x.shape[0], device=x.device).clamp(min=1e-10))\n x_new, accepted = module.mh_step_cuda(x, proposal, log_u)\n return x_new, accepted\n"
}
result ← run_correctness (ok)
okrun_correctness
run_correctness PASSED: (3 / 3) trials all matched the reference.
Numerical precision: max_abs_err=0.00e+00 mean_abs_err=0.00e+00 max_rel_err=0.00e+00
turn 38.3s turn · 8.3s LLM · 0.0s tools · tokens in=5,767 · out=685 · total=6,466 · FINAL
reasoning
Correctness passed with exact match. Now I can submit the kernel.
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 <math.h>\n\n__global__ void mh_kernel(\n const float* __restrict__ x,\n const float* __restrict__ proposal,\n const float* __restrict__ log_u,\n float* __restrict__ x_new,\n bool* __restrict__ accepted,\n int B,\n int D\n) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= B) return;\n\n float sum_x2 = 0.0f;\n float sum_p2 = 0.0f;\n for (int i = 0; i < D; ++i) {\n float xv = x[b * D + i];\n float pv = proposal[b * D + i];\n sum_x2 += xv * xv;\n sum_p2 += pv * pv;\n }\n\n float log_alpha = -0.5f * (sum_p2 - sum_x2);\n bool acc = log_u[b] < log_alpha;\n\n accepted[b] = acc;\n for (int i = 0; i < D; ++i) {\n x_new[b * D + i] = acc ? proposal[b * D + i] : x[b * D + i];\n }\n}\n\nstd::vector<torch::Tensor> mh_step_cuda(\n torch::Tensor x,\n torch::Tensor proposal,\n torch::Tensor log_u\n) {\n int B = x.size(0);\n 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 int threads = 256;\n int blocks = (B + threads - 1) / threads;\n\n mh_kernel<<<blocks, threads>>>(\n x.data_ptr<float>(),\n proposal.data_ptr<float>(),\n log_u.data_ptr<float>(),\n x_new.data_ptr<float>(),\n accepted.data_ptr<bool>(),\n B,\n D\n );\n\n return {x_new, accepted};\n}\n'''\n\ncpp_source = r'''\nstd::vector<torch::Tensor> mh_step_cuda(torch::Tensor x, torch::Tensor proposal, torch::Tensor log_u);\n'''\n\nmodule = load_inline(\n name='mh_cuda',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['mh_step_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 proposal = x + self.step_size * torch.randn_like(x)\n log_u = torch.log(torch.rand(x.shape[0], device=x.device).clamp(min=1e-10))\n x_new, accepted = module.mh_step_cuda(x, proposal, log_u)\n return x_new, accepted\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 (93 lines, 2,394 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 <math.h>
__global__ void mh_kernel(
const float* __restrict__ x,
const float* __restrict__ proposal,
const float* __restrict__ log_u,
float* __restrict__ x_new,
bool* __restrict__ accepted,
int B,
int D
) {
int b = blockIdx.x * blockDim.x + threadIdx.x;
if (b >= B) return;
float sum_x2 = 0.0f;
float sum_p2 = 0.0f;
for (int i = 0; i < D; ++i) {
float xv = x[b * D + i];
float pv = proposal[b * D + i];
sum_x2 += xv * xv;
sum_p2 += pv * pv;
}
float log_alpha = -0.5f * (sum_p2 - sum_x2);
bool acc = log_u[b] < log_alpha;
accepted[b] = acc;
for (int i = 0; i < D; ++i) {
x_new[b * D + i] = acc ? proposal[b * D + i] : x[b * D + i];
}
}
std::vector<torch::Tensor> mh_step_cuda(
torch::Tensor x,
torch::Tensor proposal,
torch::Tensor log_u
) {
int B = x.size(0);
int D = x.size(1);
auto x_new = torch::empty_like(x);
auto accepted = torch::empty({B}, x.options().dtype(torch::kBool));
int threads = 256;
int blocks = (B + threads - 1) / threads;
mh_kernel<<<blocks, threads>>>(
x.data_ptr<float>(),
proposal.data_ptr<float>(),
log_u.data_ptr<float>(),
x_new.data_ptr<float>(),
accepted.data_ptr<bool>(),
B,
D
);
return {x_new, accepted};
}
'''
cpp_source = r'''
std::vector<torch::Tensor> mh_step_cuda(torch::Tensor x, torch::Tensor proposal, torch::Tensor log_u);
'''
module = load_inline(
name='mh_cuda',
cpp_sources=[cpp_source],
cuda_sources=[cuda_source],
functions=['mh_step_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:
proposal = x + self.step_size * torch.randn_like(x)
log_u = torch.log(torch.rand(x.shape[0], device=x.device).clamp(min=1e-10))
x_new, accepted = module.mh_step_cuda(x, proposal, log_u)
return x_new, accepted
```