grok-4.3 · popcorn · pop_l123_all_gpt
compile_fail 2_FusedMultiStepDiffusionSolver.py
model
grok-4.3
turns / max
3 / 15
tool calls
3
tokens (total)
18790
wall clock
1057.6s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
—
ref runtime
-1.00 μs
speedup
—
started
2026-05-11T18:16:18Z
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 · 1 tool call(s) · 429.5s turn · 32.4s LLM · tokens in=2,843 · out=1,522 · total=4,682 · 6,866 chars (click to expand)
turn 0429.5s turn · 32.4s LLM · 397.0s tools · tokens in=2,843 · out=1,522 · total=4,682
reasoning
The task is to replace the forward computation of the PyTorch model with a custom CUDA kernel. The model performs fused multi-step diffusion denoising. I need to replace the forward computation of a PyTorch model with a custom CUDA kernel for a fused multi-step diffusion denoising process. The model, called ModelNew, must take the same inputs as the original Model, implement forward with the一樣 signature, and produce the same outputs numerically.
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 = '''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cmath>\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,\n const long* __restrict__ t_start,\n float* __restrict__ output,\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 * blockDim.x + threadIdx.x;\n if (b >= B) return;\n\n float x_cur[64]; // assuming dim <= 64\n for (int i = 0; i < dim; i++) {\n x_cur[i] = x[b * dim + i];\n }\n\n long t = t_start[b];\n\n for (int k = 0; k < num_fused_steps; k++) {\n long t_idx = t - k;\n float t_norm = (float)t_idx / (float)total_timesteps;\n\n // MLP forward\n float h1[256]; // hidden = dim*4 = 256\n for (int i = 0; i < hidden; i++) {\n float sum = b1[i];\n for (int j = 0; j < dim; j++) {\n sum += w1[i * (dim + 1) + j] * x_cur[j];\n }\n sum += w1[i * (dim + 1) + dim] * t_norm;\n h1[i] = silu(sum);\n }\n\n float h2[256];\n for (int i = 0; i < hidden; i++) {\n float sum = b2[i];\n for (int j = 0; j < hidden; j++) {\n sum += w2[i * hidden + j] * h1[j];\n }\n h2[i] = silu(sum);\n }\n\n float eps[64];\n for (int i = 0; i < dim; i++) {\n float sum = b3[i];\n for (int j = 0; j < hidden; j++) {\n sum += w3[i * hidden + j] * h2[j];\n }\n eps[i] = sum;\n }\n\n float c1 = sqrt_recip_alphas[t_idx];\n float c2 = coef_eps[t_idx];\n for (int i = 0; i < dim; i++) {\n x_cur[i] = c1 * (x_cur[i] - c2 * eps[i]);\n }\n }\n\n for (int i = 0; i < dim; i++) {\n output[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 = dim * 4;\n\n auto output = torch::zeros_like(x);\n\n dim3 blocks((B + 255) / 256);\n dim3 threads(256);\n\n fused_denoise_kernel<<<blocks, threads>>>(\n x.data_ptr<float>(),\n t_start.data_ptr<long>(),\n output.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 output;\n}\n'''\n\ncpp_source = '''\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\nmodule = load_inline(\n name='fused_denoise_ext',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['fused_denoise'],\n verbose=True\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(\n \"coef_eps\",\n betas / torch.sqrt(1.0 - alphas_cumprod),\n )\n\n def forward(self, x: torch.Tensor, t_start: torch.Tensor) -> torch.Tensor:\n # Get weights\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 module.fused_denoise(\n x, t_start.long(),\n w1, b1, w2, b2, w3, b3,\n self.sqrt_recip_alphas,\n self.coef_eps,\n self.num_fused_steps,\n self.total_timesteps\n )\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 1 · 1 tool call(s) · 342.3s turn · 14.9s LLM · tokens in=4,699 · out=1,523 · total=6,230 · 6,562 chars (click to expand)
turn 1342.3s turn · 14.9s LLM · 327.4s tools · tokens in=4,699 · out=1,523 · total=6,230
reasoning
The kernel compiled. Now check correctness.
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 = '''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cmath>\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,\n const long* __restrict__ t_start,\n float* __restrict__ output,\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 * blockDim.x + threadIdx.x;\n if (b >= B) return;\n\n float x_cur[64]; // assuming dim <= 64\n for (int i = 0; i < dim; i++) {\n x_cur[i] = x[b * dim + i];\n }\n\n long t = t_start[b];\n\n for (int k = 0; k < num_fused_steps; k++) {\n long t_idx = t - k;\n float t_norm = (float)t_idx / (float)total_timesteps;\n\n // MLP forward\n float h1[256]; // hidden = dim*4 = 256\n for (int i = 0; i < hidden; i++) {\n float sum = b1[i];\n for (int j = 0; j < dim; j++) {\n sum += w1[i * (dim + 1) + j] * x_cur[j];\n }\n sum += w1[i * (dim + 1) + dim] * t_norm;\n h1[i] = silu(sum);\n }\n\n float h2[256];\n for (int i = 0; i < hidden; i++) {\n float sum = b2[i];\n for (int j = 0; j < hidden; j++) {\n sum += w2[i * hidden + j] * h1[j];\n }\n h2[i] = silu(sum);\n }\n\n float eps[64];\n for (int i = 0; i < dim; i++) {\n float sum = b3[i];\n for (int j = 0; j < hidden; j++) {\n sum += w3[i * hidden + j] * h2[j];\n }\n eps[i] = sum;\n }\n\n float c1 = sqrt_recip_alphas[t_idx];\n float c2 = coef_eps[t_idx];\n for (int i = 0; i < dim; i++) {\n x_cur[i] = c1 * (x_cur[i] - c2 * eps[i]);\n }\n }\n\n for (int i = 0; i < dim; i++) {\n output[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 = dim * 4;\n\n auto output = torch::zeros_like(x);\n\n dim3 blocks((B + 255) / 256);\n dim3 threads(256);\n\n fused_denoise_kernel<<<blocks, threads>>>(\n x.data_ptr<float>(),\n t_start.data_ptr<long>(),\n output.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 output;\n}\n'''\n\ncpp_source = '''\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\nmodule = load_inline(\n name='fused_denoise_ext',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['fused_denoise'],\n verbose=True\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(\n \"coef_eps\",\n betas / torch.sqrt(1.0 - alphas_cumprod),\n )\n\n def forward(self, x: torch.Tensor, t_start: torch.Tensor) -> torch.Tensor:\n # Get weights\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 module.fused_denoise(\n x, t_start.long(),\n w1, b1, w2, b2, w3, b3,\n self.sqrt_recip_alphas,\n self.coef_eps,\n self.num_fused_steps,\n self.total_timesteps\n )\n" }
result ← run_correctness (ok)
okrun_correctness run_correctness PASSED: (3 / 3) trials all matched the reference. Numerical precision: max_abs_err=2.38e-07 mean_abs_err=4.69e-09 max_rel_err=1.42e-06
turn 2 · FINAL · 1 tool call(s) · 285.9s turn · 17.7s LLM · tokens in=6,290 · out=1,522 · total=7,878 · 6,702 chars (click to expand)
turn 2285.9s turn · 17.7s LLM · 268.2s tools · tokens in=6,290 · out=1,522 · total=7,878 · FINAL
reasoning
Correctness passed with very small error. Now, to check if it's faster, I should submit it, but since the goal is faster, and this is a naive kernel, it might already be faster due to fusion, or not. ...
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 = '''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cmath>\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,\n const long* __restrict__ t_start,\n float* __restrict__ output,\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 * blockDim.x + threadIdx.x;\n if (b >= B) return;\n\n float x_cur[64]; // assuming dim <= 64\n for (int i = 0; i < dim; i++) {\n x_cur[i] = x[b * dim + i];\n }\n\n long t = t_start[b];\n\n for (int k = 0; k < num_fused_steps; k++) {\n long t_idx = t - k;\n float t_norm = (float)t_idx / (float)total_timesteps;\n\n // MLP forward\n float h1[256]; // hidden = dim*4 = 256\n for (int i = 0; i < hidden; i++) {\n float sum = b1[i];\n for (int j = 0; j < dim; j++) {\n sum += w1[i * (dim + 1) + j] * x_cur[j];\n }\n sum += w1[i * (dim + 1) + dim] * t_norm;\n h1[i] = silu(sum);\n }\n\n float h2[256];\n for (int i = 0; i < hidden; i++) {\n float sum = b2[i];\n for (int j = 0; j < hidden; j++) {\n sum += w2[i * hidden + j] * h1[j];\n }\n h2[i] = silu(sum);\n }\n\n float eps[64];\n for (int i = 0; i < dim; i++) {\n float sum = b3[i];\n for (int j = 0; j < hidden; j++) {\n sum += w3[i * hidden + j] * h2[j];\n }\n eps[i] = sum;\n }\n\n float c1 = sqrt_recip_alphas[t_idx];\n float c2 = coef_eps[t_idx];\n for (int i = 0; i < dim; i++) {\n x_cur[i] = c1 * (x_cur[i] - c2 * eps[i]);\n }\n }\n\n for (int i = 0; i < dim; i++) {\n output[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 = dim * 4;\n\n auto output = torch::zeros_like(x);\n\n dim3 blocks((B + 255) / 256);\n dim3 threads(256);\n\n fused_denoise_kernel<<<blocks, threads>>>(\n x.data_ptr<float>(),\n t_start.data_ptr<long>(),\n output.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 output;\n}\n'''\n\ncpp_source = '''\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\nmodule = load_inline(\n name='fused_denoise_ext',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['fused_denoise'],\n verbose=True\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(\n \"coef_eps\",\n betas / torch.sqrt(1.0 - alphas_cumprod),\n )\n\n def forward(self, x: torch.Tensor, t_start: torch.Tensor) -> torch.Tensor:\n // Get weights\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 module.fused_denoise(\n x, t_start.long(),\n w1, b1, w2, b2, w3, b3,\n self.sqrt_recip_alphas,\n self.coef_eps,\n self.num_fused_steps,\n 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 (181 lines, 5,246 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

cuda_source = '''
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <cmath>

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

__global__ void fused_denoise_kernel(
    const float* __restrict__ x,
    const long* __restrict__ t_start,
    float* __restrict__ output,
    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 * blockDim.x + threadIdx.x;
    if (b >= B) return;

    float x_cur[64];  // assuming dim <= 64
    for (int i = 0; i < dim; i++) {
        x_cur[i] = x[b * dim + i];
    }

    long t = t_start[b];

    for (int k = 0; k < num_fused_steps; k++) {
        long t_idx = t - k;
        float t_norm = (float)t_idx / (float)total_timesteps;

        // MLP forward
        float h1[256];  // hidden = dim*4 = 256
        for (int i = 0; i < hidden; i++) {
            float sum = b1[i];
            for (int j = 0; j < dim; j++) {
                sum += w1[i * (dim + 1) + j] * x_cur[j];
            }
            sum += w1[i * (dim + 1) + dim] * t_norm;
            h1[i] = silu(sum);
        }

        float h2[256];
        for (int i = 0; i < hidden; i++) {
            float sum = b2[i];
            for (int j = 0; j < hidden; j++) {
                sum += w2[i * hidden + j] * h1[j];
            }
            h2[i] = silu(sum);
        }

        float eps[64];
        for (int i = 0; i < dim; i++) {
            float sum = b3[i];
            for (int j = 0; j < hidden; j++) {
                sum += w3[i * hidden + j] * h2[j];
            }
            eps[i] = sum;
        }

        float c1 = sqrt_recip_alphas[t_idx];
        float c2 = coef_eps[t_idx];
        for (int i = 0; i < dim; i++) {
            x_cur[i] = c1 * (x_cur[i] - c2 * eps[i]);
        }
    }

    for (int i = 0; i < dim; i++) {
        output[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 = dim * 4;

    auto output = torch::zeros_like(x);

    dim3 blocks((B + 255) / 256);
    dim3 threads(256);

    fused_denoise_kernel<<<blocks, threads>>>(
        x.data_ptr<float>(),
        t_start.data_ptr<long>(),
        output.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 output;
}
'''

cpp_source = '''
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
);
'''

module = load_inline(
    name='fused_denoise_ext',
    cpp_sources=[cpp_source],
    cuda_sources=[cuda_source],
    functions=['fused_denoise'],
    verbose=True
)

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:
        // Get weights
        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 module.fused_denoise(
            x, t_start.long(),
            w1, b1, w2, b2, w3, b3,
            self.sqrt_recip_alphas,
            self.coef_eps,
            self.num_fused_steps,
            self.total_timesteps
        )