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


class Model(nn.Module):
    """
    Stack of planar normalizing flow layers.  Each layer applies the
    invertible transform f(z) = z + u * tanh(w^T z + b), with a
    log-determinant-Jacobian correction.  Used to transform a simple
    base distribution (e.g. standard Gaussian) into a complex posterior.
    """

    def __init__(self, dim, num_flows):
        super().__init__()
        self.dim = dim
        self.num_flows = num_flows
        self.w = nn.ParameterList([nn.Parameter(torch.randn(dim)) for _ in range(num_flows)])
        self.u = nn.ParameterList([nn.Parameter(torch.randn(dim)) for _ in range(num_flows)])
        self.b = nn.ParameterList([nn.Parameter(torch.zeros(1)) for _ in range(num_flows)])

    def _constrain_u(self, w, u):
        """Ensure invertibility: w^T u >= -1."""
        wtu = (w * u).sum()
        m = -1 + F.softplus(wtu)
        u_hat = u + (m - wtu) * w / (w ** 2).sum()
        return u_hat

    def forward(self, z: torch.Tensor) -> tuple:
        """
        Args:
            z: (B, dim) – samples from base distribution
        Returns:
            (z_out, sum_log_det):
                z_out:       (B, dim) – transformed samples
                sum_log_det: (B,)     – total log |det df/dz|
        """
        sum_log_det = torch.zeros(z.shape[0], device=z.device)

        for i in range(self.num_flows):
            w = self.w[i]
            u = self._constrain_u(w, self.u[i])
            b = self.b[i]

            linear = z @ w + b  # (B,)
            h = torch.tanh(linear)
            h_prime = 1.0 - h ** 2

            z = z + u.unsqueeze(0) * h.unsqueeze(-1)
            psi = h_prime.unsqueeze(-1) * w.unsqueeze(0)  # (B, dim)
            log_det = torch.log(torch.abs(1.0 + psi @ u) + 1e-8)
            sum_log_det = sum_log_det + log_det

        return z, sum_log_det


dim = 32
num_flows = 8
batch_size = 1024


def get_inputs():
    return [torch.randn(popcorn_pri.jitter_int(batch_size), dim)]


def get_init_inputs():
    return [dim, num_flows]
turn 0 · FINAL · 1 tool call(s) · 210.9s turn · 34.2s LLM · tokens in=1,434 · out=2,292 · total=4,162 · 9,787 chars (click to expand)
turn 0210.9s turn · 34.2s LLM · 176.7s tools · tokens in=1,434 · out=2,292 · total=4,162 · FINAL
reasoning
The task is to replace the forward computation of this PyTorch model with a custom CUDA kernel. The model is a stack of planar normalizing flow layers. I need to implement a custom CUDA kernel for the forward pass of a stack of planar normalizing flow layers in PyTorch. The model uses PlanarFlowBuilder with parameters like num_flows=8, dim=32, batch=102 variable batch size up to 1024.
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 softplus(float x) { return log1pf(expf(x)); } __global__ void planar_flow_kernel( const float* __restrict__ z_in, float* __restrict__ z_out, float* __restrict__ sum_log_det, const float* __restrict__ w_all, const float* __restrict__ u_all, const float* __restrict__ b_all, int batch_size, int dim, int num_flows ) { int b = blockIdx.x * blockDim.x + threadIdx.x; if (b >= batch_size) return; float local_z[32]; // assume dim <= 32 for (int d = 0; d < dim; d++) { local_z[d] = z_in[b * dim + d]; } float log_det_acc = 0.0f; for (int i = 0; i < num_flows; i++) { const float* w = w_all + i * dim; const float* u = u_all + i * dim; float b = b_all[i]; // constrain u float wtu = 0.0f; float w_norm2 = 0.0f; for (int d = 0; d < dim; d++) { wtu += w[d] * u[d]; w_norm2 += w[d] * w[d]; } float m = -1.0f + softplus(wtu); float scale = (m - wtu) / (w_norm2 + 1e-12f); float linear = b; for (int d = 0; d < dim; d++) { linear += local_z[d] * w[d]; } float h = tanhf(linear); float h_prime = 1.0f - h * h; float psi_dot_u = 0.0f; for (int d = 0; d < dim; d++) { float u_hat_d = u[d] + scale * w[d]; float psi_d = h_prime * w[d]; local_z[d] += u_hat_d * h; psi_dot_u += psi_d * u_hat_d; } float log_det = logf(fabsf(1.0f + psi_dot_u) + 1e-8f); log_det_acc += log_det; } for (int d = 0; d < dim; d++) { z_out[b * dim + d] = local_z[d]; } sum_log_det[b] = log_det_acc; } torch::Tensor planar_flow_cuda( torch::Tensor z, torch::Tensor w_all, torch::Tensor u_all, torch::Tensor b_all, int num_flows ) { int batch_size = z.size(0); int dim = z.size(1); auto z_out = torch::empty_like(z); auto sum_log_det = torch::zeros({batch_size}, z.options()); const int threads = 256; const int blocks = (batch_size + threads - 1) / threads; planar_flow_kernel<<<blocks, threads>>>( z.data_ptr<float>(), z_out.data_ptr<float>(), sum_log_det.data_ptr<float>(), w_all.data_ptr<float>(), u_all.data_ptr<float>(), b_all.data_ptr<float>(), batch_size, dim, num_flows ); return z_out, sum_log_det; } ''' cpp_source = r''' #include <torch/extension.h> std::tuple<torch::Tensor, torch::Tensor> planar_flow_cuda( torch::Tensor z, torch::Tensor w_all, torch::Tensor u_all, torch::Tensor b_all, int num_flows ); ''' module = load_inline( name='planar_flow_ext', cpp_sources=[cpp_source], cuda_sources=[cuda_source], functions=['planar_flow_cuda'], verbose=False, extra_cflags=['-O3'] ) class ModelNew(nn.Module): def __init__(self, dim, num_flows): super().__init__() self.dim = dim self.num_flows = num_flows self.w = nn.ParameterList([nn.Parameter(torch.randn(dim)) for _ in range(num_flows)]) self.u = nn.ParameterList([nn.Parameter(torch.randn(dim)) for _ in range(num_flows)]) self.b = nn.ParameterList([nn.Parameter(torch.zeros(1)) for _ in range(num_flows)]) def forward(self, z: torch.Tensor) -> tuple: w_all = torch.stack([w.data for w in self.w]) u_all = torch.stack([u.data for u in self.u]) b_all = torch.stack([b.data for b in self.b]).view(-1) return module.planar_flow_cuda(z, w_all, u_all, b_all, self.num_flows) ``` The kernel uses register tiling for dim=32, fuses the entire flow stack per thread, and avoids Python overhead. It should be numerically equivalent and significantly faster.
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\ncuda_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 softplus(float x) {\n return log1pf(expf(x));\n}\n\n__global__ void planar_flow_kernel(\n const float* __restrict__ z_in,\n float* __restrict__ z_out,\n float* __restrict__ sum_log_det,\n const float* __restrict__ w_all,\n const float* __restrict__ u_all,\n const float* __restrict__ b_all,\n int batch_size,\n int dim,\n int num_flows\n) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= batch_size) return;\n\n float local_z[32]; // assume dim <= 32\n for (int d = 0; d < dim; d++) {\n local_z[d] = z_in[b * dim + d];\n }\n\n float log_det_acc = 0.0f;\n\n for (int i = 0; i < num_flows; i++) {\n const float* w = w_all + i * dim;\n const float* u = u_all + i * dim;\n float b = b_all[i];\n\n // constrain u\n float wtu = 0.0f;\n float w_norm2 = 0.0f;\n for (int d = 0; d < dim; d++) {\n wtu += w[d] * u[d];\n w_norm2 += w[d] * w[d];\n }\n float m = -1.0f + softplus(wtu);\n float scale = (m - wtu) / (w_norm2 + 1e-12f);\n\n float linear = b;\n for (int d = 0; d < dim; d++) {\n linear += local_z[d] * w[d];\n }\n float h = tanhf(linear);\n float h_prime = 1.0f - h * h;\n\n float psi_dot_u = 0.0f;\n for (int d = 0; d < dim; d++) {\n float u_hat_d = u[d] + scale * w[d];\n float psi_d = h_prime * w[d];\n local_z[d] += u_hat_d * h;\n psi_dot_u += psi_d * u_hat_d;\n }\n float log_det = logf(fabsf(1.0f + psi_dot_u) + 1e-8f);\n log_det_acc += log_det;\n }\n\n for (int d = 0; d < dim; d++) {\n z_out[b * dim + d] = local_z[d];\n }\n sum_log_det[b] = log_det_acc;\n}\n\nstd::tuple<torch::Tensor, torch::Tensor> planar_flow_cuda(\n torch::Tensor z,\n torch::Tensor w_all,\n torch::Tensor u_all,\n torch::Tensor b_all,\n int num_flows\n) {\n int batch_size = z.size(0);\n int dim = z.size(1);\n auto z_out = torch::empty_like(z);\n auto sum_log_det = torch::zeros({batch_size}, z.options());\n\n const int threads = 256;\n const int blocks = (batch_size + threads - 1) / threads;\n\n planar_flow_kernel<<<blocks, threads>>>(\n z.data_ptr<float>(),\n z_out.data_ptr<float>(),\n sum_log_det.data_ptr<float>(),\n w_all.data_ptr<float>(),\n u_all.data_ptr<float>(),\n b_all.data_ptr<float>(),\n batch_size,\n dim,\n num_flows\n );\n\n return std::make_tuple(z_out, sum_log_det);\n}\n'''\n\ncpp_source = r'''\n#include <torch/extension.h>\nstd::tuple<torch::Tensor, torch::Tensor> planar_flow_cuda(\n torch::Tensor z,\n torch::Tensor w_all,\n torch::Tensor u_all,\n torch::Tensor b_all,\n int num_flows\n);\n'''\n\nmodule = load_inline(\n name='planar_flow_ext',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['planar_flow_cuda'],\n verbose=False,\n extra_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim, num_flows):\n super().__init__()\n self.dim = dim\n self.num_flows = num_flows\n self.w = nn.ParameterList([nn.Parameter(torch.randn(dim)) for _ in range(num_flows)])\n self.u = nn.ParameterList([nn.Parameter(torch.randn(dim)) for _ in range(num_flows)])\n self.b = nn.ParameterList([nn.Parameter(torch.zeros(1)) for _ in range(num_flows)])\n\n def forward(self, z: torch.Tensor) -> tuple:\n w_all = torch.stack([w.data for w in self.w])\n u_all = torch.stack([u.data for u in self.u])\n b_all = torch.stack([b.data for b in self.b]).view(-1)\n return module.planar_flow_cuda(z, w_all, u_all, b_all, self.num_flows)\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 (141 lines, 3,921 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 softplus(float x) {
    return log1pf(expf(x));
}

__global__ void planar_flow_kernel(
    const float* __restrict__ z_in,
    float* __restrict__ z_out,
    float* __restrict__ sum_log_det,
    const float* __restrict__ w_all,
    const float* __restrict__ u_all,
    const float* __restrict__ b_all,
    int batch_size,
    int dim,
    int num_flows
) {
    int b = blockIdx.x * blockDim.x + threadIdx.x;
    if (b >= batch_size) return;

    float local_z[32];  // assume dim <= 32
    for (int d = 0; d < dim; d++) {
        local_z[d] = z_in[b * dim + d];
    }

    float log_det_acc = 0.0f;

    for (int i = 0; i < num_flows; i++) {
        const float* w = w_all + i * dim;
        const float* u = u_all + i * dim;
        float b = b_all[i];

        // constrain u
        float wtu = 0.0f;
        float w_norm2 = 0.0f;
        for (int d = 0; d < dim; d++) {
            wtu += w[d] * u[d];
            w_norm2 += w[d] * w[d];
        }
        float m = -1.0f + softplus(wtu);
        float scale = (m - wtu) / (w_norm2 + 1e-12f);

        float linear = b;
        for (int d = 0; d < dim; d++) {
            linear += local_z[d] * w[d];
        }
        float h = tanhf(linear);
        float h_prime = 1.0f - h * h;

        float psi_dot_u = 0.0f;
        for (int d = 0; d < dim; d++) {
            float u_hat_d = u[d] + scale * w[d];
            float psi_d = h_prime * w[d];
            local_z[d] += u_hat_d * h;
            psi_dot_u += psi_d * u_hat_d;
        }
        float log_det = logf(fabsf(1.0f + psi_dot_u) + 1e-8f);
        log_det_acc += log_det;
    }

    for (int d = 0; d < dim; d++) {
        z_out[b * dim + d] = local_z[d];
    }
    sum_log_det[b] = log_det_acc;
}

std::tuple<torch::Tensor, torch::Tensor> planar_flow_cuda(
    torch::Tensor z,
    torch::Tensor w_all,
    torch::Tensor u_all,
    torch::Tensor b_all,
    int num_flows
) {
    int batch_size = z.size(0);
    int dim = z.size(1);
    auto z_out = torch::empty_like(z);
    auto sum_log_det = torch::zeros({batch_size}, z.options());

    const int threads = 256;
    const int blocks = (batch_size + threads - 1) / threads;

    planar_flow_kernel<<<blocks, threads>>>(
        z.data_ptr<float>(),
        z_out.data_ptr<float>(),
        sum_log_det.data_ptr<float>(),
        w_all.data_ptr<float>(),
        u_all.data_ptr<float>(),
        b_all.data_ptr<float>(),
        batch_size,
        dim,
        num_flows
    );

    return std::make_tuple(z_out, sum_log_det);
}
'''

cpp_source = r'''
#include <torch/extension.h>
std::tuple<torch::Tensor, torch::Tensor> planar_flow_cuda(
    torch::Tensor z,
    torch::Tensor w_all,
    torch::Tensor u_all,
    torch::Tensor b_all,
    int num_flows
);
'''

module = load_inline(
    name='planar_flow_ext',
    cpp_sources=[cpp_source],
    cuda_sources=[cuda_source],
    functions=['planar_flow_cuda'],
    verbose=False,
    extra_cflags=['-O3']
)

class ModelNew(nn.Module):
    def __init__(self, dim, num_flows):
        super().__init__()
        self.dim = dim
        self.num_flows = num_flows
        self.w = nn.ParameterList([nn.Parameter(torch.randn(dim)) for _ in range(num_flows)])
        self.u = nn.ParameterList([nn.Parameter(torch.randn(dim)) for _ in range(num_flows)])
        self.b = nn.ParameterList([nn.Parameter(torch.zeros(1)) for _ in range(num_flows)])

    def forward(self, z: torch.Tensor) -> tuple:
        w_all = torch.stack([w.data for w in self.w])
        u_all = torch.stack([u.data for u in self.u])
        b_all = torch.stack([b.data for b in self.b]).view(-1)
        return module.planar_flow_cuda(z, w_all, u_all, b_all, self.num_flows)
```