FW-GLM-5-1 · popcorn · pop_l123_st_gpt
compile_fail 16_HamiltonianMonteCarloStep.py
model
FW-GLM-5-1
turns / max
1 / 1
tool calls
1
tokens (total)
5603
wall clock
777.9s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
—
ref runtime
-1.00 μs
speedup
—
started
2026-05-11T09:35:50Z
compile error: Error building extension 'hmc_cuda': [1/3] /opt/conda/bin/x86_64-conda-linux-gnu-c++ -MMD -MF main.o.d -DTORCH_EXTENSION_NAME=hmc_cuda -DTORCH_API_INCLUDE_EXTENSION_H -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /usr/local/cuda/include -isystem /opt/conda/include/python3.10 -fPIC -std=c++17 -c /scratch/abaner/KernelBench/runs/pop_l123_st_gpt/popcorn/FW-GLM-5-1/level_2_problem_16_cache/k_fd9dfbc5aaee/hmc_cuda/main.cpp -o main.o FAILED: [
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


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 0 · FINAL · 1 tool call(s) · 777.9s turn · 48.3s LLM · tokens in=1,381 · out=4,222 · total=5,603 · 22,201 chars (click to expand)
turn 0777.9s turn · 48.3s LLM · 729.7s tools · tokens in=1,381 · out=4,222 · total=5,603 · FINAL
reasoning
Let me analyze the task. I need to implement a custom CUDA kernel for Hamiltonian Monte Carlo (HMC) proposal using leapfrog integration. The key operations are: 1. Generate random momentum `p` from Gaussian distribution 2. Perform leapfrog integration steps: - Half step for momentum: `p = p + 0.5 * eps * grad_log_prob(q)` - Multiple full steps: `q = q + eps * p; p = p + eps * grad_log_prob(q)` - Final half step for momentum - Negate momentum for reversibility 3. Compute Hamiltonians for current and proposed states 4. Metropolis-Hastings acceptance step The gradient of log probability for standard Gaussian is: `grad_log_prob(q) = -(q - mu) = -q` (since mu=0) So the operations are: - `p = randn_like(q)` - `p = p + 0.5 * eps * (-q) = p - 0.5 * eps * q` - Loop: `q = q + eps * p; p = p - eps * q` - Final: `q = q + eps * p; p = p - 0.5 * eps * q; p = -p` - `current_H = -log_prob(current_q) + 0.5 * sum(current_p^2) = 0.5 * sum(current_q^2) + 0.5 * sum(current_p^2)` - `proposed_H = 0.5 * sum(q^2) + 0.5 * sum(p^2)` - `log_alpha = current_H - proposed_H` - Accept if `log(rand()) < log_alpha` Let me design a CUDA kernel for this: 1. I need to generate random numbers - this is tricky in CUDA. I could use cuRAND or pass pre-generated random numbers from PyTorch. 2. The leapfrog integration involves multiple position and momentum updates that can be parallelized across dimensions. 3. Computing Hamiltonians requires summing squares across dimensions for each batch element. 4. The acceptance decision needs to be made per batch element. The main challenge is handling the random number generation. I could pre-generate random numbers in PyTorch and pass them to the kernel, or use cuRAND within the kernel for better performance. Since there are multiple random number draws (initial momentum and uniform random for acceptance), I need to account for the total number of random values needed. Looking at the actual requirements: the momentum `p` is drawn once at the start, and a uniform random `u` is drawn once per batch element for the acceptance test. I'll generate both in PyTorch beforehand and pass them as arguments to keep the kernel simpler and more efficient. The Hamiltonian computation involves summing across dimensions, which can be done efficiently with parallel reduction, followed by an element-wise acceptance check. However, the loop structure reveals a complexity—the leapfrog integration alternates between updating momentum with the gradient and updating position with the momentum, and these operations depend on each other's results. This dependency chain makes straightforward parallelization challenging, since each step requires the output of the previous step. The gradient computation itself is element-wise, which is promising, but the real bottleneck is that I need to compute `grad_log_prob(q)` at multiple different `q` values throughout the trajectory. This means I can't precompute all gradients ahead of time—they must be computed at each step as `q` evolves. The dependencies create a sequential structure: updating position requires the current momentum, and updating momentum requires the current position. Since all dimensions of position and momentum get updated simultaneously at each leapfrog step, I could parallelize within each dimension at a given step, but the steps themselves must proceed in order. Each thread could handle one dimension, updating `q[d]` and `p[d]` in lockstep across the trajectory. For the Hamiltonian calculation, I'd need to sum across dimensions using parallel reduction. I could pre-generate random numbers in PyTorch rather than relying on cuRAND to keep things simpler. The kernel would take the position `q`, random momentum `p_rand`, uniform samples `u`, and parameters like step size and leapfrog steps, then have each thread process one dimension through the leapfrog sequence. The momentum and position updates interleave in a specific sequence during leapfrog integration. After the initial half-step momentum update, we alternate between full position and momentum updates for the middle steps, then finish with a final position update and another half-step momentum update before negation. This structure ensures proper time-reversibility in the Hamiltonian dynamics. Since we can't easily parallelize the leapfrog steps themselves due to their sequential nature, I should focus on optimizing memory access and leveraging PyTorch's existing operations where possible. The most computationally intensive parts—the random number generation and basic arithmetic—are already well-optimized in PyTorch, so a custom CUDA kernel would need to be carefully designed to provide real speedup over just using `torch.randn` for momentum initialization and standard tensor operations for the integration steps. A fusion approach could work well here: pre-generate the random momentum and uniform samples on the GPU, then use a dedicated kernel that handles the entire leapfrog trajectory and acceptance logic in one pass. Each thread would process one batch element, computing all dimension updates for each integration step while maintaining the sequential dependency of the leapfrog scheme. This way we avoid repeated kernel launches and memory transfers between steps. Now I'll implement the CUDA kernel. Each thread processes one batch element, storing position and momentum in registers for fast access. The leapfrog steps update q and p alternately, starting with a half-step on momentum, then alternating full steps on position and momentum. After the specified number of steps, I need to complete the final half-step on momentum and then compute the Hamiltonian to determine acceptance. For the Hamiltonian calculation, I sum the squared values of the proposed position and momentum across dimensions. I also need to compute the current Hamiltonian using the original position and momentum values that were saved at the beginning. The acceptance probability is then determined by comparing these two energy values. I compute the log acceptance ratio as the difference between current and proposed Hamiltonians, then compare it against a log uniform sample. If the proposal is accepted, I copy the new position into the output; otherwise, I preserve the original position and record the acceptance status. The kernel launch configuration uses a block size of 256 threads, with enough blocks to cover all batch elements. I'm passing the device pointers for positions, momenta, uniform samples, outputs, and acceptance flags along with the step parameters. The forward method accepts the position tensor and returns both the updated positions and a boolean tensor indicating which proposals were accepted. To implement this, I generate random momenta from a standard normal distribution and uniform samples for the acceptance test, then invoke the kernel with the appropriate dimension and leapfrog parameters. The positions tensor is already on the correct device, and I ensure the acceptance tensor matches its device placement. I need to reconsider the random number handling. The reference generates momentum via `torch.randn_like(q)` and uniform samples for acceptance, which I should preserve. The key constraint is that each batch element needs one uniform random number for the acceptance step. I can generate these tensors in PyTorch upfront and pass them to the kernel, keeping the kernel focused on the leapfrog integration and acceptance logic rather than random number generation. The leapfrog step itself is where the optimization matters most—it's the computational core where I can fuse operations for better performance. The Hamiltonian computation requires a sum across dimensions, which presents a challenge for parallelization. I'm considering having each thread handle one batch element across all dimensions—if the dimensionality is moderate (like 64), the thread can efficiently manage all data, and I can leverage shuffle operations for reduction if needed. For correctness, I need to verify that `grad_log_prob(q) = -(q - mu) = -q` when mu is zero. I should also examine how `log_prob` computes the full log-likelihood, as this affects the Hamiltonian calculation. Looking at the reference implementation, `log_prob(x)` returns -0.5 * ||x||², which means the Hamiltonian H = -log_prob(q) + 0.5 * ||p||² = 0.5 * ||q||² + 0.5 * ||p||². In the leapfrog integration loop, I notice there's a final position update and half-step momentum update outside the main iteration cycle, followed by negation of momentum for reversibility. After negating momentum, the proposed Hamiltonian uses the final position and the negated momentum value, which maintains the correct symplectic structure needed for the Metropolis-Hastings acceptance step. For the CUDA kernel, I'll keep the structure straightforward: each thread handles one batch element across all dimensions, pre-generating random numbers for the momentum initialization and acceptance decisions before passing them to the kernel. I need to complete the Hamiltonian computation for the proposed state and then handle the acceptance logic. The log probability sums over all dimensions, so I accumulate the squared momentum and position values. Then I compute the Hamiltonian as the negative log probability plus the kinetic energy term. For acceptance, I compare the current Hamiltonian against the proposed one—if the proposed state has lower energy, it's more likely to be accepted. The acceptance decision uses a uniform random number and the log of that number compared against the energy difference. If accepted, I copy the proposed state; otherwise I keep the current state. The kernel needs to be launched with appropriate block and grid dimensions based on the batch size. In the forward pass, I generate random momentum and uniform samples, then prepare the output tensors before launching the kernel with the position data, random samples, and leapfrog parameters like step size and number of steps. # Final half step for momentum for (int d = 0; d < dim; d++) { T q_val = q[d]; p[d] = p[d] - half_eps * q_val; } // Negate momentum for (int d = 0; d < dim; d++) { p[d] = -p[d]; } // Compute Hamiltonians T H_current = 0, H_proposed = 0; for (int d = 0; d < dim; d++) { H_current += 0.5f * (current_q[d] * current_q[d] + current_p[d] * current_p[d]); H_proposed += 0.5f * (q[d] * q[d] + p[d] * p[d]); } // Metropolis acceptance T log_alpha = H_current - H_proposed; The kernel completes the leapfrog integration with a final momentum half-step, then negates momentum for reversibility. Hamiltonians for both current and proposed states are computed by summing kinetic and potential energy across dimensions. The Metropolis-Hastings acceptance probability follows from the energy difference. The loop unrolling should work well given that dim is fixed at 64—compilers typically handle small constant iteration counts effectively. Register usage depends on whether the compiler keeps `q`, `p`, `current_q`, and `current_p` in registers across the leapfrog steps, which could pressure register availability but should remain manageable for this dimensionality. Looking back at the reference implementation, I notice the initial momentum update is a half-step (`p = p + 0.5 * eps * grad_log_prob(q)`), then the loop runs `num_leapfrog_steps - 1` iterations with full steps, and there's a final position update followed by another half-step on momentum. This matches what I have, so the structure is correct. Computing the gradients inline eliminates the intermediate arrays, which should improve memory efficiency. For handling different precision types, I can leverage CUDA's function overloading to define type-specific absolute value operations rather than relying on a potentially problematic macro approach.
call → submit_kernel
callsubmit_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\nhmc_cuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\ntemplate <typename T>\n__device__ T my_fabs(T x);\n\ntemplate <> __device__ float my_fabs(float x) { return fabsf(x); }\ntemplate <> __device__ double my_fabs(double x) { return fabs(x); }\n\ntemplate <typename T>\n__global__ void hmc_leapfrog_kernel(\n const T* __restrict__ q_in,\n const T* __restrict__ p_rand,\n const T* __restrict__ u_rand,\n T* __restrict__ q_out,\n bool* __restrict__ accepted_out,\n int batch_size,\n int dim,\n T eps,\n int num_steps\n) {\n int batch_idx = blockIdx.x * blockDim.x + threadIdx.x;\n if (batch_idx >= batch_size) return;\n \n const T* q_ptr = q_in + batch_idx * dim;\n const T* p_rand_ptr = p_rand + batch_idx * dim;\n T* q_out_ptr = q_out + batch_idx * dim;\n \n T eps_sq = eps * eps;\n T eps_half = eps * static_cast<T>(0.5);\n \n // Local storage for position and momentum\n // For small dim (64), this fits in registers\n T q[64], p[64], current_q[64], current_p[64];\n \n #pragma unroll\n for (int d = 0; d < dim; d++) {\n current_q[d] = q_ptr[d];\n current_p[d] = p_rand_ptr[d];\n q[d] = current_q[d];\n p[d] = current_p[d];\n }\n \n // Initial half step for momentum: p = p - 0.5 * eps * q\n #pragma unroll\n for (int d = 0; d < dim; d++) {\n p[d] = p[d] - eps_half * q[d];\n }\n \n // Main leapfrog loop: (num_steps - 1) full steps\n for (int step = 0; step < num_steps - 1; step++) {\n // Full position update: q = q + eps * p\n #pragma unroll\n for (int d = 0; d < dim; d++) {\n q[d] = q[d] + eps * p[d];\n }\n // Full momentum update: p = p - eps * q\n #pragma unroll\n for (int d = 0; d < dim; d++) {\n p[d] = p[d] - eps * q[d];\n }\n }\n \n // Final full position update\n #pragma unroll\n for (int d = 0; d < dim; d++) {\n q[d] = q[d] + eps * p[d];\n }\n \n // Final half step for momentum\n #pragma unroll\n for (int d = 0; d < dim; d++) {\n p[d] = p[d] - eps_half * q[d];\n }\n \n // Negate momentum for reversibility\n #pragma unroll\n for (int d = 0; d < dim; d++) {\n p[d] = -p[d];\n }\n \n // Compute Hamiltonians: H = 0.5 * (||q||^2 + ||p||^2)\n T H_current = static_cast<T>(0);\n T H_proposed = static_cast<T>(0);\n \n #pragma unroll\n for (int d = 0; d < dim; d++) {\n H_current += current_q[d] * current_q[d] + current_p[d] * current_p[d];\n H_proposed += q[d] * q[d] + p[d] * p[d];\n }\n H_current *= static_cast<T>(0.5);\n H_proposed *= static_cast<T>(0.5);\n \n // Metropolis-Hastings acceptance\n T log_alpha = H_current - H_proposed;\n T log_u = log(u_rand[batch_idx]);\n bool accept = log_u < log_alpha;\n \n accepted_out[batch_idx] = accept;\n \n if (accept) {\n #pragma unroll\n for (int d = 0; d < dim; d++) {\n q_out_ptr[d] = q[d];\n }\n } else {\n #pragma unroll\n for (int d = 0; d < dim; d++) {\n q_out_ptr[d] = current_q[d];\n }\n }\n}\n\n// Define types for template instantiation\n#define INSTANTIATE_HMC_KERNEL(T) \\\ntemplate __global__ void hmc_leapfrog_kernel<T>( \\\n const T* __restrict__ q_in, \\\n const T* __restrict__ p_rand, \\\n const T* __restrict__ u_rand, \\\n T* __restrict__ q_out, \\\n bool* __restrict__ accepted_out, \\\n int batch_size, \\\n int dim, \\\n T eps, \\\n int num_steps \\\n);\n\nINSTANTIATE_HMC_KERNEL(float)\nINSTANTIATE_HMC_KERNEL(double)\n\nstd::tuple<torch::Tensor, torch::Tensor> hmc_forward(\n torch::Tensor q,\n torch::Tensor p_rand,\n torch::Tensor u_rand,\n double step_size,\n int64_t num_leapfrog_steps\n) {\n TORCH_CHECK(q.dim() == 2, \"q must be 2D\");\n int batch_size = q.size(0);\n int dim = q.size(1);\n \n auto q_out = torch::empty_like(q);\n auto accepted = torch::empty({batch_size}, torch::dtype(torch::kBool).device(q.device()));\n \n const int block_size = 256;\n const int num_blocks = (batch_size + block_size - 1) / block_size;\n \n AT_DISPATCH_FLOATING_TYPES(q.scalar_type(), \"hmc_leapfrog_kernel\", ([&] {\n hmc_leapfrog_kernel<scalar_t><<<num_blocks, block_size>>>(\n q.data_ptr<scalar_t>(),\n p_rand.data_ptr<scalar_t>(),\n u_rand.data_ptr<scalar_t>(),\n q_out.data_ptr<scalar_t>(),\n accepted.data_ptr<bool>(),\n batch_size,\n dim,\n static_cast<scalar_t>(step_size),\n static_cast<int>(num_leapfrog_steps)\n );\n }));\n \n return std::make_tuple(q_out, accepted);\n}\n\"\"\"\n\nhmc_module = load_inline(\n name='hmc_cuda',\n cpp_sources='',\n cuda_sources=hmc_cuda_source,\n functions=['hmc_forward'],\n verbose=False,\n extra_cuda_cflags=['-O3', '--use_fast_math']\n)\n\n\nclass ModelNew(nn.Module):\n \"\"\"\n Hamiltonian Monte Carlo (HMC) proposal using leapfrog integration.\n Optimized with a custom CUDA kernel that fuses the leapfrog steps\n and acceptance computation.\n \"\"\"\n\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\n def forward(self, q: torch.Tensor) -> tuple:\n \"\"\"\n Args:\n q: (B, dim) - current positions\n Returns:\n (q_new, accepted):\n q_new: (B, dim)\n accepted: (B,)\n \"\"\"\n # Generate random numbers for momentum and acceptance\n p_rand = torch.randn_like(q)\n u_rand = torch.rand(q.shape[0], device=q.device, dtype=q.dtype).clamp(min=1e-10)\n \n # Call the fused CUDA kernel\n q_new, accepted = hmc_module.hmc_forward(\n q, p_rand, u_rand, self.step_size, self.num_leapfrog_steps\n )\n \n return q_new, accepted" }
result ← submit_kernel (fail)
failsubmit_kernel submit_kernel FAILED: kernel did not compile. Error building extension 'hmc_cuda': [1/3] /opt/conda/bin/x86_64-conda-linux-gnu-c++ -MMD -MF main.o.d -DTORCH_EXTENSION_NAME=hmc_cuda -DTORCH_API_INCLUDE_EXTENSION_H -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /usr/local/cuda/include -isystem /opt/conda/include/python3.10 -fPIC -std=c++17 -c /scratch/abaner/KernelBench/runs/pop_l123_st_gpt/popcorn/FW-GLM-5-1/level_2_problem_16_cache/k_fd9dfbc5aaee/hmc_cuda/main.cpp -o main.o FAILED: [code=1] main.o /opt/conda/bin/x86_64-conda-linux-gnu-c++ -MMD -MF main.o.d -DTORCH_EXTENSION_NAME=hmc_cuda -DTORCH_API_INCLUDE_EXTENSION_H -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /usr/local/cuda/include -isystem /opt/conda/include/python3.10 -fPIC -std=c++17 -c /scratch/abaner/KernelBench/runs/pop_l123_st_gpt/popcorn/FW-GLM-5-1/level_2_problem_16_cache/k_fd9dfbc5aaee/hmc_cuda/main.cpp -o main.o /scratch/abaner/KernelBench/runs/pop_l123_st_gpt/popcorn/FW-GLM-5-1/level_2_problem_16_cache/k_fd9dfbc5aaee/hmc_cuda/main.cpp: In function 'void pybind11_init_hmc_cuda(pybind11::module_&)': /scratch/abaner/KernelBench/runs/pop_l123_st_gpt/popcorn/FW-GLM-5-1/level_2_problem_16_cache/k_fd9dfbc5aaee/hmc_cuda/main.cpp:4:50: error: 'hmc_forward' was not declared in this scope 4 | m.def("hmc_forward", torch::wrap_pybind_function(hmc_forward), "hmc_forward"); | ^~~~~~~~~~~ [2/3] /usr/local/cuda/bin/nvcc -MD -MF cuda.cuda.o.d -ccbin /opt/conda/bin/x86_64-conda-linux-gnu-cc -DTORCH_EXTENSION_NAME=hmc_cuda -DTORCH_API_INCLUDE_EXTENSION_H -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /usr/local/cuda/include -isystem /opt/conda/include/python3.10 -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr -gencode=arch=compute_90,code=compute_90 -gencode=arch=compute_90,code=sm_90 --compiler-options '-fPIC' -O3 --use_fast_math -std=c++17 -c /scratch/abaner/KernelBench/runs/pop_l123_st_gpt/popcorn/FW-GLM-5-1/level_2_problem_16_cache/k_fd9dfbc5aaee/hmc_cuda/cuda.cu -o cuda.cuda.o nvcc warning : incompatible redefinition for option 'compiler-bindir', the last value of this option was used ninja: build stopped: subcommand failed.
final submitted kernel (213 lines, 6,025 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

hmc_cuda_source = """
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>

template <typename T>
__device__ T my_fabs(T x);

template <> __device__ float my_fabs(float x) { return fabsf(x); }
template <> __device__ double my_fabs(double x) { return fabs(x); }

template <typename T>
__global__ void hmc_leapfrog_kernel(
    const T* __restrict__ q_in,
    const T* __restrict__ p_rand,
    const T* __restrict__ u_rand,
    T* __restrict__ q_out,
    bool* __restrict__ accepted_out,
    int batch_size,
    int dim,
    T eps,
    int num_steps
) {
    int batch_idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (batch_idx >= batch_size) return;
    
    const T* q_ptr = q_in + batch_idx * dim;
    const T* p_rand_ptr = p_rand + batch_idx * dim;
    T* q_out_ptr = q_out + batch_idx * dim;
    
    T eps_sq = eps * eps;
    T eps_half = eps * static_cast<T>(0.5);
    
    // Local storage for position and momentum
    // For small dim (64), this fits in registers
    T q[64], p[64], current_q[64], current_p[64];
    
    #pragma unroll
    for (int d = 0; d < dim; d++) {
        current_q[d] = q_ptr[d];
        current_p[d] = p_rand_ptr[d];
        q[d] = current_q[d];
        p[d] = current_p[d];
    }
    
    // Initial half step for momentum: p = p - 0.5 * eps * q
    #pragma unroll
    for (int d = 0; d < dim; d++) {
        p[d] = p[d] - eps_half * q[d];
    }
    
    // Main leapfrog loop: (num_steps - 1) full steps
    for (int step = 0; step < num_steps - 1; step++) {
        // Full position update: q = q + eps * p
        #pragma unroll
        for (int d = 0; d < dim; d++) {
            q[d] = q[d] + eps * p[d];
        }
        // Full momentum update: p = p - eps * q
        #pragma unroll
        for (int d = 0; d < dim; d++) {
            p[d] = p[d] - eps * q[d];
        }
    }
    
    // Final full position update
    #pragma unroll
    for (int d = 0; d < dim; d++) {
        q[d] = q[d] + eps * p[d];
    }
    
    // Final half step for momentum
    #pragma unroll
    for (int d = 0; d < dim; d++) {
        p[d] = p[d] - eps_half * q[d];
    }
    
    // Negate momentum for reversibility
    #pragma unroll
    for (int d = 0; d < dim; d++) {
        p[d] = -p[d];
    }
    
    // Compute Hamiltonians: H = 0.5 * (||q||^2 + ||p||^2)
    T H_current = static_cast<T>(0);
    T H_proposed = static_cast<T>(0);
    
    #pragma unroll
    for (int d = 0; d < dim; d++) {
        H_current += current_q[d] * current_q[d] + current_p[d] * current_p[d];
        H_proposed += q[d] * q[d] + p[d] * p[d];
    }
    H_current *= static_cast<T>(0.5);
    H_proposed *= static_cast<T>(0.5);
    
    // Metropolis-Hastings acceptance
    T log_alpha = H_current - H_proposed;
    T log_u = log(u_rand[batch_idx]);
    bool accept = log_u < log_alpha;
    
    accepted_out[batch_idx] = accept;
    
    if (accept) {
        #pragma unroll
        for (int d = 0; d < dim; d++) {
            q_out_ptr[d] = q[d];
        }
    } else {
        #pragma unroll
        for (int d = 0; d < dim; d++) {
            q_out_ptr[d] = current_q[d];
        }
    }
}

// Define types for template instantiation
#define INSTANTIATE_HMC_KERNEL(T) \
template __global__ void hmc_leapfrog_kernel<T>( \
    const T* __restrict__ q_in, \
    const T* __restrict__ p_rand, \
    const T* __restrict__ u_rand, \
    T* __restrict__ q_out, \
    bool* __restrict__ accepted_out, \
    int batch_size, \
    int dim, \
    T eps, \
    int num_steps \
);

INSTANTIATE_HMC_KERNEL(float)
INSTANTIATE_HMC_KERNEL(double)

std::tuple<torch::Tensor, torch::Tensor> hmc_forward(
    torch::Tensor q,
    torch::Tensor p_rand,
    torch::Tensor u_rand,
    double step_size,
    int64_t num_leapfrog_steps
) {
    TORCH_CHECK(q.dim() == 2, "q must be 2D");
    int batch_size = q.size(0);
    int dim = q.size(1);
    
    auto q_out = torch::empty_like(q);
    auto accepted = torch::empty({batch_size}, torch::dtype(torch::kBool).device(q.device()));
    
    const int block_size = 256;
    const int num_blocks = (batch_size + block_size - 1) / block_size;
    
    AT_DISPATCH_FLOATING_TYPES(q.scalar_type(), "hmc_leapfrog_kernel", ([&] {
        hmc_leapfrog_kernel<scalar_t><<<num_blocks, block_size>>>(
            q.data_ptr<scalar_t>(),
            p_rand.data_ptr<scalar_t>(),
            u_rand.data_ptr<scalar_t>(),
            q_out.data_ptr<scalar_t>(),
            accepted.data_ptr<bool>(),
            batch_size,
            dim,
            static_cast<scalar_t>(step_size),
            static_cast<int>(num_leapfrog_steps)
        );
    }));
    
    return std::make_tuple(q_out, accepted);
}
"""

hmc_module = load_inline(
    name='hmc_cuda',
    cpp_sources='',
    cuda_sources=hmc_cuda_source,
    functions=['hmc_forward'],
    verbose=False,
    extra_cuda_cflags=['-O3', '--use_fast_math']
)


class ModelNew(nn.Module):
    """
    Hamiltonian Monte Carlo (HMC) proposal using leapfrog integration.
    Optimized with a custom CUDA kernel that fuses the leapfrog steps
    and acceptance computation.
    """

    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

    def forward(self, q: torch.Tensor) -> tuple:
        """
        Args:
            q: (B, dim) - current positions
        Returns:
            (q_new, accepted):
                q_new:    (B, dim)
                accepted: (B,)
        """
        # Generate random numbers for momentum and acceptance
        p_rand = torch.randn_like(q)
        u_rand = torch.rand(q.shape[0], device=q.device, dtype=q.dtype).clamp(min=1e-10)
        
        # Call the fused CUDA kernel
        q_new, accepted = hmc_module.hmc_forward(
            q, p_rand, u_rand, self.step_size, self.num_leapfrog_steps
        )
        
        return q_new, accepted