correct
16_HamiltonianMonteCarloStep.py
started
2026-05-05T00:17:02Z
timingkernelbench.timing
kernel mean / median / std
0.01 / — / 1.55e-04 μs
ref mean / median / std
0.54 / — / 0.02 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
11.60x
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
dram bandwidth
214.84 GB/s
warp stalls (top 5)nsight
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn
class Model(nn.Module):
"""
Hamiltonian Monte Carlo (HMC) proposal using leapfrog integration.
Simulates Hamiltonian dynamics on a batched set of particles to
produce proposals with high acceptance probability. The target
distribution is a standard Gaussian for benchmarking.
"""
def __init__(self, dim, num_leapfrog_steps=10, step_size=0.1):
super().__init__()
self.dim = dim
self.num_leapfrog_steps = num_leapfrog_steps
self.step_size = step_size
self.register_buffer("mu", torch.zeros(dim))
def _grad_log_prob(self, x: torch.Tensor) -> torch.Tensor:
"""Gradient of log p(x) for N(0, I)."""
return -(x - self.mu)
def _log_prob(self, x: torch.Tensor) -> torch.Tensor:
return -0.5 * (x ** 2).sum(dim=-1)
def forward(self, q: torch.Tensor) -> tuple:
"""
Args:
q: (B, dim) – current positions
Returns:
(q_new, accepted):
q_new: (B, dim)
accepted: (B,)
"""
eps = self.step_size
p = torch.randn_like(q)
current_p = p.clone()
current_q = q.clone()
# Half step for momentum
p = p + 0.5 * eps * self._grad_log_prob(q)
for _ in range(self.num_leapfrog_steps - 1):
q = q + eps * p
p = p + eps * self._grad_log_prob(q)
q = q + eps * p
p = p + 0.5 * eps * self._grad_log_prob(q)
p = -p # negate for reversibility
current_H = -self._log_prob(current_q) + 0.5 * (current_p ** 2).sum(-1)
proposed_H = -self._log_prob(q) + 0.5 * (p ** 2).sum(-1)
log_alpha = current_H - proposed_H
log_u = torch.log(torch.rand(q.shape[0], device=q.device).clamp(min=1e-10))
accepted = log_u < log_alpha
q_new = torch.where(accepted.unsqueeze(-1), q, current_q)
return q_new, accepted
dim = 64
batch_size = 2048
def get_inputs():
return [torch.randn(popcorn_pri.jitter_int(batch_size), dim)]
def get_init_inputs():
return [dim]
turn 042.8s turn · 42.8s LLM · 0.0s tools · FINAL
reasoning
**Considering computation efficiency**
I’m thinking about kernel operations per row that can compute local quantities for strided dimensions and sum energies while reducing within a block. After that, it sounds like each thread will write its elements of q_new, which seems straightforward. Now, should I consider using dynamic shared memory? Maybe I could use 128 floats and an accepted boolean. By the way, generating random numbers might add overhead, similar to the reference implementation.
call → submit_kernel
callsubmit_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n\n_cpp_src = r\"\"\"\n#include <torch/extension.h>\n#include <vector>\n\nstd::vector<torch::Tensor> hmc_forward(torch::Tensor q,\n torch::Tensor p,\n torch::Tensor u,\n int64_t steps,\n double eps);\n\"\"\"\n\n_cuda_src = r\"\"\"\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <vector>\n\n__global__ void hmc_kernel(const float* __restrict__ q0,\n const float* __restrict__ p0,\n const float* __restrict__ u,\n float* __restrict__ qout,\n bool* __restrict__ accepted,\n int B, int D, int steps, float eps) {\n int b = blockIdx.x;\n int tid = threadIdx.x;\n if (b >= B) return;\n\n const int base = b * D;\n float local = 0.0f;\n\n for (int d = tid; d < D; d += blockDim.x) {\n int idx = base + d;\n float q = q0[idx];\n float p = p0[idx];\n float cq = q;\n float cp = p;\n\n // Half momentum step: grad log p(q) = -q for N(0,I)\n p = p - (0.5f * eps) * q;\n\n // Full leapfrog steps except the last position update.\n for (int s = 0; s < steps - 1; ++s) {\n q = q + eps * p;\n p = p - eps * q;\n }\n\n q = q + eps * p;\n p = p - (0.5f * eps) * q;\n p = -p;\n\n qout[idx] = q; // kept if accepted, overwritten below if rejected\n local += 0.5f * (cq * cq + cp * cp - q * q - p * p);\n }\n\n __shared__ float smem[256];\n smem[tid] = local;\n __syncthreads();\n\n // blockDim is 128 in the wrapper, but keep this generic for powers of two <=256.\n for (int stride = blockDim.x >> 1; stride > 0; stride >>= 1) {\n if (tid < stride) smem[tid] += smem[tid + stride];\n __syncthreads();\n }\n\n __shared__ int acc_s;\n if (tid == 0) {\n float uu = u[b];\n uu = uu < 1.0e-10f ? 1.0e-10f : uu;\n int acc = (logf(uu) < smem[0]);\n acc_s = acc;\n accepted[b] = (bool)acc;\n }\n __syncthreads();\n\n if (!acc_s) {\n for (int d = tid; d < D; d += blockDim.x) {\n int idx = base + d;\n qout[idx] = q0[idx];\n }\n }\n}\n\nstd::vector<torch::Tensor> hmc_forward(torch::Tensor q,\n torch::Tensor p,\n torch::Tensor u,\n int64_t steps,\n double eps) {\n const int B = (int)q.size(0);\n const int D = (int)q.size(1);\n auto qout = torch::empty_like(q);\n auto acc = torch::empty({B}, q.options().dtype(torch::kBool));\n\n const int threads = 128;\n dim3 grid(B);\n cudaStream_t stream = at::cuda::getCurrentCUDAStream();\n hmc_kernel<<<grid, threads, 0, stream>>>(\n q.data_ptr<float>(), p.data_ptr<float>(), u.data_ptr<float>(),\n qout.data_ptr<float>(), acc.data_ptr<bool>(), B, D, (int)steps, (float)eps);\n return {qout, acc};\n}\n\"\"\"\n\n_hmc_ext = load_inline(\n name=\"hmc_leapfrog_ext_v1\",\n cpp_sources=_cpp_src,\n cuda_sources=_cuda_src,\n functions=[\"hmc_forward\"],\n extra_cflags=[\"-O3\"],\n extra_cuda_cflags=[\"-O3\", \"--use_fast_math\"],\n verbose=False,\n)\n\n\nclass ModelNew(nn.Module):\n def __init__(self, dim, num_leapfrog_steps=10, step_size=0.1):\n super().__init__()\n self.dim = dim\n self.num_leapfrog_steps = num_leapfrog_steps\n self.step_size = step_size\n self.register_buffer(\"mu\", torch.zeros(dim))\n\n def forward(self, q: torch.Tensor) -> tuple:\n p = torch.randn_like(q)\n u = torch.rand((q.shape[0],), device=q.device)\n q_new, accepted = _hmc_ext.hmc_forward(q, p, u, int(self.num_leapfrog_steps), float(self.step_size))\n return q_new, accepted\n"
}
result ← submit_kernel (ok)
oksubmit_kernel
submit_kernel PASSED: (5 / 5) correctness trials all passed.
Kernel runtime: 0.01 μs
Runtime stats: mean=0.01μs median=?μs std=0.00μs
Numerical precision: max_abs_err=7.15e-07 mean_abs_err=2.46e-08 max_rel_err=9.40e-03
Memory: 34.51 MB peak (ref: 36.5 MB, ratio: 0.9455x)
Kernel launches: 10 (ref: 116, fusion ratio: 11.6)
SOL score: 0.4884
Flagged for excessive speedup — this submission may have been rejected by automated review.
final submitted kernel (131 lines, 3,975 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline
_cpp_src = r"""
#include <torch/extension.h>
#include <vector>
std::vector<torch::Tensor> hmc_forward(torch::Tensor q,
torch::Tensor p,
torch::Tensor u,
int64_t steps,
double eps);
"""
_cuda_src = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <vector>
__global__ void hmc_kernel(const float* __restrict__ q0,
const float* __restrict__ p0,
const float* __restrict__ u,
float* __restrict__ qout,
bool* __restrict__ accepted,
int B, int D, int steps, float eps) {
int b = blockIdx.x;
int tid = threadIdx.x;
if (b >= B) return;
const int base = b * D;
float local = 0.0f;
for (int d = tid; d < D; d += blockDim.x) {
int idx = base + d;
float q = q0[idx];
float p = p0[idx];
float cq = q;
float cp = p;
// Half momentum step: grad log p(q) = -q for N(0,I)
p = p - (0.5f * eps) * q;
// Full leapfrog steps except the last position update.
for (int s = 0; s < steps - 1; ++s) {
q = q + eps * p;
p = p - eps * q;
}
q = q + eps * p;
p = p - (0.5f * eps) * q;
p = -p;
qout[idx] = q; // kept if accepted, overwritten below if rejected
local += 0.5f * (cq * cq + cp * cp - q * q - p * p);
}
__shared__ float smem[256];
smem[tid] = local;
__syncthreads();
// blockDim is 128 in the wrapper, but keep this generic for powers of two <=256.
for (int stride = blockDim.x >> 1; stride > 0; stride >>= 1) {
if (tid < stride) smem[tid] += smem[tid + stride];
__syncthreads();
}
__shared__ int acc_s;
if (tid == 0) {
float uu = u[b];
uu = uu < 1.0e-10f ? 1.0e-10f : uu;
int acc = (logf(uu) < smem[0]);
acc_s = acc;
accepted[b] = (bool)acc;
}
__syncthreads();
if (!acc_s) {
for (int d = tid; d < D; d += blockDim.x) {
int idx = base + d;
qout[idx] = q0[idx];
}
}
}
std::vector<torch::Tensor> hmc_forward(torch::Tensor q,
torch::Tensor p,
torch::Tensor u,
int64_t steps,
double eps) {
const int B = (int)q.size(0);
const int D = (int)q.size(1);
auto qout = torch::empty_like(q);
auto acc = torch::empty({B}, q.options().dtype(torch::kBool));
const int threads = 128;
dim3 grid(B);
cudaStream_t stream = at::cuda::getCurrentCUDAStream();
hmc_kernel<<<grid, threads, 0, stream>>>(
q.data_ptr<float>(), p.data_ptr<float>(), u.data_ptr<float>(),
qout.data_ptr<float>(), acc.data_ptr<bool>(), B, D, (int)steps, (float)eps);
return {qout, acc};
}
"""
_hmc_ext = load_inline(
name="hmc_leapfrog_ext_v1",
cpp_sources=_cpp_src,
cuda_sources=_cuda_src,
functions=["hmc_forward"],
extra_cflags=["-O3"],
extra_cuda_cflags=["-O3", "--use_fast_math"],
verbose=False,
)
class ModelNew(nn.Module):
def __init__(self, dim, num_leapfrog_steps=10, step_size=0.1):
super().__init__()
self.dim = dim
self.num_leapfrog_steps = num_leapfrog_steps
self.step_size = step_size
self.register_buffer("mu", torch.zeros(dim))
def forward(self, q: torch.Tensor) -> tuple:
p = torch.randn_like(q)
u = torch.rand((q.shape[0],), device=q.device)
q_new, accepted = _hmc_ext.hmc_forward(q, p, u, int(self.num_leapfrog_steps), float(self.step_size))
return q_new, accepted