grok-4.3 · popcorn · pop_l123_st_gpt
compile_fail 2_FusedMultiStepDiffusionSolver.py
model
grok-4.3
turns / max
1 / 1
tool calls
1
tokens (total)
5655
wall clock
128.8s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
—
ref runtime
-1.00 μs
speedup
—
started
2026-05-11T09:02:28Z
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

# Default ModelNew is PyTorch-only. run_and_check: use check_kernel=False with backend=cuda
# until the candidate file contains a real CUDA kernel (see kernel_static_checker).


def _linear_beta_schedule(num_timesteps: int, beta_start: float, beta_end: float) -> torch.Tensor:
    return torch.linspace(beta_start, beta_end, num_timesteps)


class Model(nn.Module):
    """
    Fused **multi-step** diffusion denoising (VP-DDPM-style) in a single ``forward``.

    Instead of one scheduler step per Python call, this reference unrolls
    ``num_fused_steps`` consecutive **deterministic** reverse updates (posterior mean,
    no sampler noise)—a typical fusion target for robotics / world-model diffusion
    heads where latency matters.

    Each inner step: predict ``ε_θ(x, t)`` with a small MLP, then apply the standard
    DDPM mean update from discrete time ``t`` to ``t-1``. Batch elements may start
    from different ``t_start``; inputs must satisfy ``t_start >= num_fused_steps`` so
    indices stay valid.
    """

    def __init__(self, dim: int, num_fused_steps: int, total_timesteps: int = 1000):
        super().__init__()
        assert num_fused_steps >= 1
        assert total_timesteps >= num_fused_steps + 1
        self.dim = dim
        self.num_fused_steps = num_fused_steps
        self.total_timesteps = total_timesteps

        self.eps_net = nn.Sequential(
            nn.Linear(dim + 1, dim * 4),
            nn.SiLU(),
            nn.Linear(dim * 4, dim * 4),
            nn.SiLU(),
            nn.Linear(dim * 4, dim),
        )

        betas = _linear_beta_schedule(total_timesteps, 1e-4, 0.02)
        alphas = 1.0 - betas
        alphas_cumprod = torch.cumprod(alphas, dim=0)
        self.register_buffer("betas", betas)
        self.register_buffer("alphas", alphas)
        self.register_buffer("alphas_cumprod", alphas_cumprod)
        # x_{t-1} = (1/sqrt(alpha_t)) * (x_t - (beta_t / sqrt(1-alpha_bar_t)) * eps)
        self.register_buffer("sqrt_recip_alphas", torch.sqrt(1.0 / alphas))
        self.register_buffer(
            "coef_eps",
            betas / torch.sqrt(1.0 - alphas_cumprod),
        )

    def forward(self, x: torch.Tensor, t_start: torch.Tensor) -> torch.Tensor:
        """
        Args:
            x: (B, dim) noisy latent at discrete time ``t_start`` (same semantics as DDPM).
            t_start: (B,) int64 timesteps in ``[num_fused_steps, total_timesteps - 1]``.
        Returns:
            (B, dim) latent after ``num_fused_steps`` fused denoise steps toward 0.
        """
        B, d = x.shape
        assert d == self.dim
        assert t_start.shape == (B,)
        # Harness paths (measure_ref_program_time / eval) cast activations to fp32;
        # discrete timesteps must survive as whole-valued floats → long indices.
        t_start = t_start.long()

        x_cur = x
        for k in range(self.num_fused_steps):
            t_idx = t_start - k
            # ``get_inputs`` samples t_start >= num_fused_steps so t_idx >= 1 always.
            t_norm = t_idx.float() / float(self.total_timesteps)
            eps = self.eps_net(torch.cat([x_cur, t_norm.unsqueeze(-1)], dim=-1))

            c1 = self.sqrt_recip_alphas[t_idx].view(B, 1)
            c2 = self.coef_eps[t_idx].view(B, 1)
            x_cur = c1 * (x_cur - c2 * eps)

        return x_cur


class ModelNew(nn.Module):
    """KernelBench candidate entry point; default delegates to ``Model``."""

    def __init__(self, dim: int, num_fused_steps: int, total_timesteps: int = 1000):
        super().__init__()
        self._impl = Model(dim, num_fused_steps, total_timesteps=total_timesteps)

    def forward(self, x: torch.Tensor, t_start: torch.Tensor) -> torch.Tensor:
        return self._impl(x, t_start)


batch_size = 8
dim = 64
num_fused_steps = 16
total_timesteps = 1000


def get_inputs():
    p = popcorn_pri
    x = torch.randn(p.jitter_int(batch_size), dim)
    low = num_fused_steps
    high = total_timesteps
    t_start = torch.randint(low, high, (p.jitter_int(batch_size),), dtype=torch.long)
    return [x, t_start]


def get_init_inputs():
    return [dim, num_fused_steps, total_timesteps]
turn 0 · FINAL · 1 tool call(s) · 128.8s turn · 73.1s LLM · tokens in=1,945 · out=3,257 · total=5,655 · 13,299 chars (click to expand)
turn 0128.8s turn · 73.1s LLM · 55.6s tools · tokens in=1,945 · out=3,257 · total=5,655 · FINAL
reasoning
The task is to replace the forward computation of a PyTorch model with a custom CUDA kernel. The model is a fused multi-step diffusion denoising process. The task is to replace the forward computation of a PyTorch model with a potential CUDA kernel for a fused multi-step diffusion denoising process. The reference model performs several denoising steps, calculating indices, normalizing times, concatenating inputs, passing them through an MLP to predict noise, and updating the latent variable using precomputed buffers.
assistant```python import torch import torch.nn as nn import torch.nn.functional as F 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> __device__ __forceinline__ float silu(float x) { return x / (1.0f + expf(-x)); } __global__ void fused_denoise_kernel( const float* __restrict__ x_in, const long* __restrict__ t_start, float* __restrict__ x_out, const float* __restrict__ w1, const float* __restrict__ b1, const float* __restrict__ w2, const float* __restrict__ b2, const float* __restrict__ w3, const float* __restrict__ b3, const float* __restrict__ sqrt_recip_alphas, const float* __restrict__ coef_eps, int B, int dim, int hidden, int num_fused_steps, int total_timesteps ) { int b = blockIdx.x; if (b >= B) return; extern __shared__ float smem[]; float* x_cur = smem; // dim floats per block float* h1 = smem + dim; float* h2 = h1 + hidden; // load initial x for (int i = threadIdx.x; i < dim; i += blockDim.x) { x_cur[i] = x_in[b * dim + i]; } __syncthreads(); long t0 = t_start[b]; for (int k = 0; k < num_fused_steps; ++k) { long t_idx = t0 - k; float t_norm = (float)t_idx / (float)total_timesteps; // MLP forward: concat [x, t_norm] -> h1 (dim+1 -> hidden) for (int o = threadIdx.x; o < hidden; o += blockDim.x) { float sum = b1[o]; for (int i = 0; i < dim; ++i) { sum += x_cur[i] * w1[o * (dim + 1) + i]; } sum += t_norm * w1[o * (dim + 1) + dim]; h1[o] = silu(sum); } __syncthreads(); // h1 -> h2 (hidden -> hidden) for (int o = threadIdx.x; o < hidden; o += blockDim.x) { float sum = b2[o]; for (int i = 0; i < hidden; ++i) { sum += h1[i] * w2[o * hidden + i]; } h2[o] = silu(sum); } __syncthreads(); // h2 -> eps (hidden -> dim) for (int o = threadIdx.x; o < dim; o += blockDim.x) { float sum = b3[o]; for (int i = 0; i < hidden; ++i) { sum += h2[i] * w3[o * hidden + i]; } float eps_val = sum; float c1 = sqrt_recip_alphas[t_idx]; float c2 = coef_eps[t_idx]; x_cur[o] = c1 * (x_cur[o] - c2 * eps_val); } __syncthreads(); } // write out for (int i = threadIdx.x; i < dim; i += blockDim.x) { x_out[b * dim + i] = x_cur[i]; } } torch::Tensor fused_denoise( torch::Tensor x, torch::Tensor t_start, torch::Tensor w1, torch::Tensor b1, torch::Tensor w2, torch::Tensor b2, torch::Tensor w3, torch::Tensor b3, torch::Tensor sqrt_recip_alphas, torch::Tensor coef_eps, int num_fused_steps, int total_timesteps ) { int B = x.size(0); int dim = x.size(1); int hidden = w1.size(0); auto x_out = torch::empty_like(x); dim3 grid(B); dim3 block(256); int smem_size = (dim + 2 * hidden) * sizeof(float); fused_denoise_kernel<<<grid, block, smem_size>>>( x.data_ptr<float>(), t_start.data_ptr<long>(), x_out.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>(), sqrt_recip_alphas.data_ptr<float>(), coef_eps.data_ptr<float>(), B, dim, hidden, num_fused_steps, total_timesteps ); return x_out; } ''' _cpp_source = r''' torch::Tensor fused_denoise( torch::Tensor x, torch::Tensor t_start, torch::Tensor w1, torch::Tensor b1, torch::Tensor w2, torch::Tensor b2, torch::Tensor w3, torch::Tensor b3, torch::Tensor sqrt_recip_alphas, torch::Tensor coef_eps, int num_fused_steps, int total_timesteps ); ''' _fused = load_inline( name='fused_denoise_ext', cpp_sources=_cpp_source, cuda_sources=_cuda_source, functions=['fused_denoise'], verbose=False ) class ModelNew(nn.Module): def __init__(self, dim: int, num_fused_steps: int, total_timesteps: int = 1000): super().__init__() self.dim = dim self.num_fused_steps = num_fused_steps self.total_timesteps = total_timesteps self.eps_net = nn.Sequential( nn.Linear(dim + 1, dim * 4), nn.SiLU(), nn.Linear(dim * 4, dim * 4), nn.SiLU(), nn.Linear(dim * 4, dim), ) betas = torch.linspace(1e-4, 0.02, total_timesteps) alphas = 1.0 - betas alphas_cumprod = torch.cumprod(alphas, dim=0) self.register_buffer("betas", betas) self.register_buffer("alphas", alphas) self.register_buffer("alphas_cumprod", alphas_cumprod) self.register_buffer("sqrt_recip_alphas", torch.sqrt(1.0 / alphas)) self.register_buffer("coef_eps", betas / torch.sqrt(1.0 - alphas_cumprod)) def forward(self, x: torch.Tensor, t_start: torch.Tensor) -> torch.Tensor: w1 = self.eps_net[0].weight b1 = self.eps_net[0].bias w2 = self.eps_net[2].weight b2 = self.eps_net[2].bias w3 = self.eps_net[4].weight b3 = self.eps_net[4].bias return _fused.fused_denoise( x, t_start, w1, b1, w2, b2, w3, b3, self.sqrt_recip_alphas, self.coef_eps, self.num_fused_steps, self.total_timesteps ) ``` submit_kernel with above code.
call → submit_kernel
callsubmit_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nimport torch.nn.functional as F\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 <math.h>\n\n__device__ __forceinline__ float silu(float x) {\n return x / (1.0f + expf(-x));\n}\n\n__global__ void fused_denoise_kernel(\n const float* __restrict__ x_in,\n const long* __restrict__ t_start,\n float* __restrict__ x_out,\n const float* __restrict__ w1, const float* __restrict__ b1,\n const float* __restrict__ w2, const float* __restrict__ b2,\n const float* __restrict__ w3, const float* __restrict__ b3,\n const float* __restrict__ sqrt_recip_alphas,\n const float* __restrict__ coef_eps,\n int B, int dim, int hidden, int num_fused_steps, int total_timesteps\n) {\n int b = blockIdx.x;\n if (b >= B) return;\n\n extern __shared__ float smem[];\n float* x_cur = smem; // dim floats per block\n float* h1 = smem + dim;\n float* h2 = h1 + hidden;\n\n // load initial x\n for (int i = threadIdx.x; i < dim; i += blockDim.x) {\n x_cur[i] = x_in[b * dim + i];\n }\n __syncthreads();\n\n long t0 = t_start[b];\n for (int k = 0; k < num_fused_steps; ++k) {\n long t_idx = t0 - k;\n float t_norm = (float)t_idx / (float)total_timesteps;\n\n // MLP forward: concat [x, t_norm] -> h1 (dim+1 -> hidden)\n for (int o = threadIdx.x; o < hidden; o += blockDim.x) {\n float sum = b1[o];\n for (int i = 0; i < dim; ++i) {\n sum += x_cur[i] * w1[o * (dim + 1) + i];\n }\n sum += t_norm * w1[o * (dim + 1) + dim];\n h1[o] = silu(sum);\n }\n __syncthreads();\n\n // h1 -> h2 (hidden -> hidden)\n for (int o = threadIdx.x; o < hidden; o += blockDim.x) {\n float sum = b2[o];\n for (int i = 0; i < hidden; ++i) {\n sum += h1[i] * w2[o * hidden + i];\n }\n h2[o] = silu(sum);\n }\n __syncthreads();\n\n // h2 -> eps (hidden -> dim)\n for (int o = threadIdx.x; o < dim; o += blockDim.x) {\n float sum = b3[o];\n for (int i = 0; i < hidden; ++i) {\n sum += h2[i] * w3[o * hidden + i];\n }\n float eps_val = sum;\n float c1 = sqrt_recip_alphas[t_idx];\n float c2 = coef_eps[t_idx];\n x_cur[o] = c1 * (x_cur[o] - c2 * eps_val);\n }\n __syncthreads();\n }\n\n // write out\n for (int i = threadIdx.x; i < dim; i += blockDim.x) {\n x_out[b * dim + i] = x_cur[i];\n }\n}\n\ntorch::Tensor fused_denoise(\n torch::Tensor x,\n torch::Tensor t_start,\n torch::Tensor w1, torch::Tensor b1,\n torch::Tensor w2, torch::Tensor b2,\n torch::Tensor w3, torch::Tensor b3,\n torch::Tensor sqrt_recip_alphas,\n torch::Tensor coef_eps,\n int num_fused_steps, int total_timesteps\n) {\n int B = x.size(0);\n int dim = x.size(1);\n int hidden = w1.size(0);\n\n auto x_out = torch::empty_like(x);\n\n dim3 grid(B);\n dim3 block(256);\n int smem_size = (dim + 2 * hidden) * sizeof(float);\n\n fused_denoise_kernel<<<grid, block, smem_size>>>(\n x.data_ptr<float>(),\n t_start.data_ptr<long>(),\n x_out.data_ptr<float>(),\n w1.data_ptr<float>(), b1.data_ptr<float>(),\n w2.data_ptr<float>(), b2.data_ptr<float>(),\n w3.data_ptr<float>(), b3.data_ptr<float>(),\n sqrt_recip_alphas.data_ptr<float>(),\n coef_eps.data_ptr<float>(),\n B, dim, hidden, num_fused_steps, total_timesteps\n );\n\n return x_out;\n}\n'''\n\n_cpp_source = r'''\ntorch::Tensor fused_denoise(\n torch::Tensor x,\n torch::Tensor t_start,\n torch::Tensor w1, torch::Tensor b1,\n torch::Tensor w2, torch::Tensor b2,\n torch::Tensor w3, torch::Tensor b3,\n torch::Tensor sqrt_recip_alphas,\n torch::Tensor coef_eps,\n int num_fused_steps, int total_timesteps\n);\n'''\n\n_fused = load_inline(\n name='fused_denoise_ext',\n cpp_sources=_cpp_source,\n cuda_sources=_cuda_source,\n functions=['fused_denoise'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim: int, num_fused_steps: int, total_timesteps: int = 1000):\n super().__init__()\n self.dim = dim\n self.num_fused_steps = num_fused_steps\n self.total_timesteps = total_timesteps\n\n self.eps_net = nn.Sequential(\n nn.Linear(dim + 1, dim * 4),\n nn.SiLU(),\n nn.Linear(dim * 4, dim * 4),\n nn.SiLU(),\n nn.Linear(dim * 4, dim),\n )\n\n betas = torch.linspace(1e-4, 0.02, total_timesteps)\n alphas = 1.0 - betas\n alphas_cumprod = torch.cumprod(alphas, dim=0)\n self.register_buffer(\"betas\", betas)\n self.register_buffer(\"alphas\", alphas)\n self.register_buffer(\"alphas_cumprod\", alphas_cumprod)\n self.register_buffer(\"sqrt_recip_alphas\", torch.sqrt(1.0 / alphas))\n self.register_buffer(\"coef_eps\", betas / torch.sqrt(1.0 - alphas_cumprod))\n\n def forward(self, x: torch.Tensor, t_start: torch.Tensor) -> torch.Tensor:\n w1 = self.eps_net[0].weight\n b1 = self.eps_net[0].bias\n w2 = self.eps_net[2].weight\n b2 = self.eps_net[2].bias\n w3 = self.eps_net[4].weight\n b3 = self.eps_net[4].bias\n\n return _fused.fused_denoise(\n x, t_start,\n w1, b1, w2, b2, w3, b3,\n self.sqrt_recip_alphas, self.coef_eps,\n self.num_fused_steps, self.total_timesteps\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 (182 lines, 5,555 chars) — full source, untruncated
import torch
import torch.nn as nn
import torch.nn.functional as F
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>

__device__ __forceinline__ float silu(float x) {
    return x / (1.0f + expf(-x));
}

__global__ void fused_denoise_kernel(
    const float* __restrict__ x_in,
    const long* __restrict__ t_start,
    float* __restrict__ x_out,
    const float* __restrict__ w1, const float* __restrict__ b1,
    const float* __restrict__ w2, const float* __restrict__ b2,
    const float* __restrict__ w3, const float* __restrict__ b3,
    const float* __restrict__ sqrt_recip_alphas,
    const float* __restrict__ coef_eps,
    int B, int dim, int hidden, int num_fused_steps, int total_timesteps
) {
    int b = blockIdx.x;
    if (b >= B) return;

    extern __shared__ float smem[];
    float* x_cur = smem;  // dim floats per block
    float* h1 = smem + dim;
    float* h2 = h1 + hidden;

    // load initial x
    for (int i = threadIdx.x; i < dim; i += blockDim.x) {
        x_cur[i] = x_in[b * dim + i];
    }
    __syncthreads();

    long t0 = t_start[b];
    for (int k = 0; k < num_fused_steps; ++k) {
        long t_idx = t0 - k;
        float t_norm = (float)t_idx / (float)total_timesteps;

        // MLP forward: concat [x, t_norm] -> h1 (dim+1 -> hidden)
        for (int o = threadIdx.x; o < hidden; o += blockDim.x) {
            float sum = b1[o];
            for (int i = 0; i < dim; ++i) {
                sum += x_cur[i] * w1[o * (dim + 1) + i];
            }
            sum += t_norm * w1[o * (dim + 1) + dim];
            h1[o] = silu(sum);
        }
        __syncthreads();

        // h1 -> h2 (hidden -> hidden)
        for (int o = threadIdx.x; o < hidden; o += blockDim.x) {
            float sum = b2[o];
            for (int i = 0; i < hidden; ++i) {
                sum += h1[i] * w2[o * hidden + i];
            }
            h2[o] = silu(sum);
        }
        __syncthreads();

        // h2 -> eps (hidden -> dim)
        for (int o = threadIdx.x; o < dim; o += blockDim.x) {
            float sum = b3[o];
            for (int i = 0; i < hidden; ++i) {
                sum += h2[i] * w3[o * hidden + i];
            }
            float eps_val = sum;
            float c1 = sqrt_recip_alphas[t_idx];
            float c2 = coef_eps[t_idx];
            x_cur[o] = c1 * (x_cur[o] - c2 * eps_val);
        }
        __syncthreads();
    }

    // write out
    for (int i = threadIdx.x; i < dim; i += blockDim.x) {
        x_out[b * dim + i] = x_cur[i];
    }
}

torch::Tensor fused_denoise(
    torch::Tensor x,
    torch::Tensor t_start,
    torch::Tensor w1, torch::Tensor b1,
    torch::Tensor w2, torch::Tensor b2,
    torch::Tensor w3, torch::Tensor b3,
    torch::Tensor sqrt_recip_alphas,
    torch::Tensor coef_eps,
    int num_fused_steps, int total_timesteps
) {
    int B = x.size(0);
    int dim = x.size(1);
    int hidden = w1.size(0);

    auto x_out = torch::empty_like(x);

    dim3 grid(B);
    dim3 block(256);
    int smem_size = (dim + 2 * hidden) * sizeof(float);

    fused_denoise_kernel<<<grid, block, smem_size>>>(
        x.data_ptr<float>(),
        t_start.data_ptr<long>(),
        x_out.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>(),
        sqrt_recip_alphas.data_ptr<float>(),
        coef_eps.data_ptr<float>(),
        B, dim, hidden, num_fused_steps, total_timesteps
    );

    return x_out;
}
'''

_cpp_source = r'''
torch::Tensor fused_denoise(
    torch::Tensor x,
    torch::Tensor t_start,
    torch::Tensor w1, torch::Tensor b1,
    torch::Tensor w2, torch::Tensor b2,
    torch::Tensor w3, torch::Tensor b3,
    torch::Tensor sqrt_recip_alphas,
    torch::Tensor coef_eps,
    int num_fused_steps, int total_timesteps
);
'''

_fused = load_inline(
    name='fused_denoise_ext',
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=['fused_denoise'],
    verbose=False
)

class ModelNew(nn.Module):
    def __init__(self, dim: int, num_fused_steps: int, total_timesteps: int = 1000):
        super().__init__()
        self.dim = dim
        self.num_fused_steps = num_fused_steps
        self.total_timesteps = total_timesteps

        self.eps_net = nn.Sequential(
            nn.Linear(dim + 1, dim * 4),
            nn.SiLU(),
            nn.Linear(dim * 4, dim * 4),
            nn.SiLU(),
            nn.Linear(dim * 4, dim),
        )

        betas = torch.linspace(1e-4, 0.02, total_timesteps)
        alphas = 1.0 - betas
        alphas_cumprod = torch.cumprod(alphas, dim=0)
        self.register_buffer("betas", betas)
        self.register_buffer("alphas", alphas)
        self.register_buffer("alphas_cumprod", alphas_cumprod)
        self.register_buffer("sqrt_recip_alphas", torch.sqrt(1.0 / alphas))
        self.register_buffer("coef_eps", betas / torch.sqrt(1.0 - alphas_cumprod))

    def forward(self, x: torch.Tensor, t_start: torch.Tensor) -> torch.Tensor:
        w1 = self.eps_net[0].weight
        b1 = self.eps_net[0].bias
        w2 = self.eps_net[2].weight
        b2 = self.eps_net[2].bias
        w3 = self.eps_net[4].weight
        b3 = self.eps_net[4].bias

        return _fused.fused_denoise(
            x, t_start,
            w1, b1, w2, b2, w3, b3,
            self.sqrt_recip_alphas, self.coef_eps,
            self.num_fused_steps, self.total_timesteps
        )
```