grok-4.3 · popcorn · pop_l123_st_gpt
compile_fail 13_ConditionalVAELoss.py
model
grok-4.3
turns / max
1 / 1
tool calls
1
tokens (total)
4451
wall clock
71.3s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
—
ref runtime
-1.00 μs
speedup
—
started
2026-05-11T10:10:24Z
compile error: Error building extension 'cvae_cuda': [1/3] /opt/conda/bin/x86_64-conda-linux-gnu-c++ -MMD -MF main.o.d -DTORCH_EXTENSION_NAME=cvae_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/grok-4.3/level_3_problem_13_cache/k_c485bd631fb1/cvae_cuda/main.cpp -o main.o [2/3] /u
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):
    """
    Conditional Variational Autoencoder (CVAE) with full encoder–decoder
    forward pass and loss computation.  The encoder conditions on both
    input x and label y; the decoder conditions on z and y.  Returns
    reconstruction loss + KL divergence.
    """

    def __init__(self, input_dim, label_dim, hidden_dim, latent_dim):
        super().__init__()
        # Encoder: q(z | x, y)
        self.enc = nn.Sequential(
            nn.Linear(input_dim + label_dim, hidden_dim),
            nn.ReLU(),
        )
        self.enc_mu = nn.Linear(hidden_dim, latent_dim)
        self.enc_logvar = nn.Linear(hidden_dim, latent_dim)

        # Decoder: p(x | z, y)
        self.dec = nn.Sequential(
            nn.Linear(latent_dim + label_dim, hidden_dim),
            nn.ReLU(),
            nn.Linear(hidden_dim, input_dim),
        )

    def forward(self, x: torch.Tensor, y: torch.Tensor) -> tuple:
        """
        Args:
            x: (B, input_dim)  – input data
            y: (B, label_dim)  – conditioning label (one-hot)
        Returns:
            (loss, recon_loss, kl_loss):
                loss:       scalar – total CVAE loss
                recon_loss: scalar
                kl_loss:    scalar
        """
        h = self.enc(torch.cat([x, y], dim=-1))
        mu = self.enc_mu(h)
        logvar = self.enc_logvar(h)

        std = torch.exp(0.5 * logvar)
        z = mu + std * torch.randn_like(std)

        x_recon = self.dec(torch.cat([z, y], dim=-1))
        recon_loss = F.mse_loss(x_recon, x, reduction="sum") / x.shape[0]
        kl_loss = -0.5 * (1 + logvar - mu.pow(2) - logvar.exp()).sum() / x.shape[0]
        loss = recon_loss + kl_loss
        return loss, recon_loss, kl_loss


input_dim = 256
label_dim = 10
hidden_dim = 128
latent_dim = 32
batch_size = 128


def get_inputs():
    p = popcorn_pri
    x = torch.randn(p.jitter_int(batch_size), input_dim)
    y = torch.zeros(p.jitter_int(batch_size), label_dim)
    y.scatter_(1, torch.randint(0, label_dim, (p.jitter_int(batch_size), 1)), 1.0)
    return [x, y]


def get_init_inputs():
    return [input_dim, label_dim, hidden_dim, latent_dim]
turn 0 · FINAL · 1 tool call(s) · 71.3s turn · 42.3s LLM · tokens in=1,439 · out=2,256 · total=4,451 · 13,157 chars (click to expand)
turn 071.3s turn · 42.3s LLM · 29.0s tools · tokens in=1,439 · out=2,256 · total=4,451 · FINAL
reasoning
The task is to replace the forward computation of a PyTorch CVAE model with a custom CUDA kernel. I need to create a ModelNew class that does the same thing but faster using CUDA. I need to create a ModelNew class that replaces the forward computation of a PyTorch CVAE model with a custom CUDA kernel to make it faster. The original model is a Conditional VAE with an encoder and decoder, handling reconstruction loss and KL divergence. First, I concatenate the input x and condition y, pass ผ่าน a linear layer to a hidden size, apply ReLU, then two linear layers for mu and logvar.
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 <curand_kernel.h>\n\n__global__ void cvae_forward_kernel(\n const float* __restrict__ x,\n const float* __restrict__ y,\n const float* __restrict__ enc_weight,\n const float* __restrict__ enc_bias,\n const float* __restrict__ enc_mu_weight,\n const float* __restrict__ enc_mu_bias,\n const float* __restrict__ enc_logvar_weight,\n const float* __restrict__ enc_logvar_bias,\n const float* __restrict__ dec_weight1,\n const float* __restrict__ dec_bias1,\n const float* __restrict__ dec_weight2,\n const float* __restrict__ dec_bias2,\n float* __restrict__ loss_out,\n float* __recon_loss_out,\n float* __kl_loss_out,\n int batch_size,\n int input_dim,\n int label_dim,\n int hidden_dim,\n int latent_dim,\n unsigned long long seed\n) {\n int b = blockIdx.x;\n if (b >= batch_size) return;\n\n extern __shared__ float smem[];\n\n // Simple implementation: each block handles one sample, using shared mem for intermediates\n float* h = smem;\n float* mu = &smem[hidden_dim];\n float* logvar = &smem[hidden_dim + latent_dim];\n float* z = &smem[hidden_dim + 2*latent_dim];\n float* h2 = &smem[hidden_dim + 3*latent_dim];\n float* x_recon = &smem[hidden_dim + 3*latent_dim + hidden_dim];\n\n // Encoder first linear + ReLU\n for (int i = threadIdx.x; i < hidden_dim; i += blockDim.x) {\n float sum = enc_bias[i];\n for (int j = 0; j < input_dim + label_dim; j++) {\n float val = (j < input_dim) ? x[b * input_dim + j] : y[b * label_dim + (j - input_dim)];\n sum += val * enc_weight[i * (input_dim + label_dim) + j];\n }\n h[i] = fmaxf(sum, 0.0f);\n }\n __syncthreads();\n\n // mu and logvar\n for (int i = threadIdx.x; i < latent_dim; i += blockDim.x) {\n float sum_mu = enc_mu_bias[i];\n float sum_logvar = enc_logvar_bias[i];\n for (int j = 0; j < hidden_dim; j++) {\n sum_mu += h[j] * enc_mu_weight[i * hidden_dim + j];\n sum_logvar += h[j] * enc_logvar_weight[i * hidden_dim + j];\n }\n mu[i] = sum_mu;\n logvar[i] = sum_logvar;\n }\n __syncthreads();\n\n // Sample z\n for (int i = threadIdx.x; i < latent_dim; i += blockDim.x) {\n float std = expf(0.5f * logvar[i]);\n curandState state;\n curand_init(seed + b * latent_dim + i, 0, 0, &state);\n float eps = curand_normal(&state);\n z[i] = mu[i] + std * eps;\n }\n __syncthreads();\n\n // Decoder\n for (int i = threadIdx.x; i < hidden_dim; i += blockDim.x) {\n float sum = dec_bias1[i];\n for (int j = 0; j < latent_dim + label_dim; j++) {\n float val = (j < latent_dim) ? z[j] : y[b * label_dim + (j - latent_dim)];\n sum += val * dec_weight1[i * (latent_dim + label_dim) + j];\n }\n h2[i] = fmaxf(sum, 0.0f);\n }\n __syncthreads();\n\n for (int i = threadIdx.x; i < input_dim; i += blockDim.x) {\n float sum = dec_bias2[i];\n for (int j = 0; j < hidden_dim; j++) {\n sum += h2[j] * dec_weight2[i * hidden_dim + j];\n }\n x_recon[i] = sum;\n }\n __syncthreads();\n\n // Compute per-sample recon and kl (will reduce later)\n __shared__ float partial_recon;\n __shared__ float partial_kl;\n if (threadIdx.x == 0) {\n partial_recon = 0.0f;\n partial_kl = 0.0f;\n }\n __syncthreads();\n\n float local_recon = 0.0f;\n float local_kl = 0.0f;\n for (int i = threadIdx.x; i < input_dim; i += blockDim.x) {\n float diff = x_recon[i] - x[b * input_dim + i];\n local_recon += diff * diff;\n }\n for (int i = threadIdx.x; i < latent_dim; i += blockDim.x) {\n local_kl += -0.5f * (1.0f + logvar[i] - mu[i]*mu[i] - expf(logvar[i]));\n }\n atomicAdd(&partial_recon, local_recon);\n atomicAdd(&partial_kl, local_kl);\n __syncthreads();\n\n if (threadIdx.x == 0) {\n float recon_per_sample = partial_recon / batch_size;\n float kl_per_sample = partial_kl / batch_size;\n atomicAdd(loss_out, recon_per_sample + kl_per_sample);\n atomicAdd(recon_loss_out, recon_per_sample);\n atomicAdd(kl_loss_out, kl_per_sample);\n }\n}\n\nstd::vector<torch::Tensor> cvae_forward_cuda(\n torch::Tensor x,\n torch::Tensor y,\n torch::Tensor enc_weight,\n torch::Tensor enc_bias,\n torch::Tensor enc_mu_weight,\n torch::Tensor enc_mu_bias,\n torch::Tensor enc_logvar_weight,\n torch::Tensor enc_logvar_bias,\n torch::Tensor dec_weight1,\n torch::Tensor dec_bias1,\n torch::Tensor dec_weight2,\n torch::Tensor dec_bias2,\n int hidden_dim,\n int latent_dim\n) {\n int batch_size = x.size(0);\n int input_dim = x.size(1);\n int label_dim = y.size(1);\n\n auto options = torch::TensorOptions().device(x.device()).dtype(torch::kFloat32);\n auto loss_out = torch::zeros({}, options);\n auto recon_out = torch::zeros({}, options);\n auto kl_out = torch::zeros({}, options);\n\n int threads = 128;\n size_t smem_size = (hidden_dim + 3*latent_dim + hidden_dim + input_dim) * sizeof(float);\n\n unsigned long long seed = 123456789ULL;\n\n cvae_forward_kernel<<<batch_size, threads, smem_size>>>(\n x.data_ptr<float>(),\n y.data_ptr<float>(),\n enc_weight.data_ptr<float>(),\n enc_bias.data_ptr<float>(),\n enc_mu_weight.data_ptr<float>(),\n enc_mu_bias.data_ptr<float>(),\n enc_logvar_weight.data_ptr<float>(),\n enc_logvar_bias.data_ptr<float>(),\n dec_weight1.data_ptr<float>(),\n dec_bias1.data_ptr<float>(),\n dec_weight2.data_ptr<float>(),\n dec_bias2.data_ptr<float>(),\n loss_out.data_ptr<float>(),\n recon_out.data_ptr<float>(),\n kl_out.data_ptr<float>(),\n batch_size, input_dim, label_dim, hidden_dim, latent_dim, seed\n );\n\n return {loss_out, recon_out, kl_out};\n}\n'''\n\ncpp_source = r'''\nstd::vector<torch::Tensor> cvae_forward_cuda(\n torch::Tensor x, torch::Tensor y,\n torch::Tensor enc_weight, torch::Tensor enc_bias,\n torch::Tensor enc_mu_weight, torch::Tensor enc_mu_bias,\n torch::Tensor enc_logvar_weight, torch::Tensor enc_logvar_bias,\n torch::Tensor dec_weight1, torch::Tensor dec_bias1,\n torch::Tensor dec_weight2, torch::Tensor dec_bias2,\n int hidden_dim, int latent_dim\n);\n'''\n\nmodule = load_inline(\n name='cvae_cuda',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['cvae_forward_cuda'],\n verbose=False,\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, input_dim, label_dim, hidden_dim, latent_dim):\n super().__init__()\n self.enc = nn.Sequential(\n nn.Linear(input_dim + label_dim, hidden_dim),\n nn.ReLU(),\n )\n self.enc_mu = nn.Linear(hidden_dim, latent_dim)\n self.enc_logvar = nn.Linear(hidden_dim, latent_dim)\n self.dec = nn.Sequential(\n nn.Linear(latent_dim + label_dim, hidden_dim),\n nn.ReLU(),\n nn.Linear(hidden_dim, input_dim),\n )\n self.input_dim = input_dim\n self.label_dim = label_dim\n self.hidden_dim = hidden_dim\n self.latent_dim = latent_dim\n\n def forward(self, x: torch.Tensor, y: torch.Tensor) -> tuple:\n # Extract weights\n enc_w = self.enc[0].weight\n enc_b = self.enc[0].bias\n mu_w = self.enc_mu.weight\n mu_b = self.enc_mu.bias\n logvar_w = self.enc_logvar.weight\n logvar_b = self.enc_logvar.bias\n dec_w1 = self.dec[0].weight\n dec_b1 = self.dec[0].bias\n dec_w2 = self.dec[2].weight\n dec_b2 = self.dec[2].bias\n\n outs = module.cvae_forward_cuda(\n x, y, enc_w, enc_b, mu_w, mu_b, logvar_w, logvar_b,\n dec_w1, dec_b1, dec_w2, dec_b2, self.hidden_dim, self.latent_dim\n )\n loss = outs[0]\n recon = outs[1]\n kl = outs[2]\n return loss, recon, kl\n" }
result ← submit_kernel (fail)
failsubmit_kernel submit_kernel FAILED: kernel did not compile. Error building extension 'cvae_cuda': [1/3] /opt/conda/bin/x86_64-conda-linux-gnu-c++ -MMD -MF main.o.d -DTORCH_EXTENSION_NAME=cvae_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/grok-4.3/level_3_problem_13_cache/k_c485bd631fb1/cvae_cuda/main.cpp -o main.o [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=cvae_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 -std=c++17 -c /scratch/abaner/KernelBench/runs/pop_l123_st_gpt/popcorn/grok-4.3/level_3_problem_13_cache/k_c485bd631fb1/cvae_cuda/cuda.cu -o cuda.cuda.o FAILED: [code=2] cuda.cuda.o /usr/local/cuda/bin/nvcc -MD -MF cuda.cuda.o.d -ccbin /opt/conda/bin/x86_64-conda-linux-gnu-cc -DTORCH_EXTENSION_NAME=cvae_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 -std=c++17 -c /scratch/abaner/KernelBench/runs/pop_l123_st_gpt/popcorn/grok-4.3/level_3_problem_13_cache/k_c485bd631fb1/cvae_cuda/cuda.cu -o cuda.cuda.o nvcc warning : incompatible redefinition for option 'compiler-bindir', the last value of this option was used /scratch/abaner/KernelBench/runs/pop_l123_st_gpt/popcorn/grok-4.3/level_3_problem_13_cache/k_c485bd631fb1/cvae_cuda/cuda.cu(126): error: identifier "recon_loss_out" is undefined atomicAdd(recon_loss_out, recon_per_sample); ^ /scratch/abaner/KernelBench/runs/pop_l123_st_gpt/popcorn/grok-4.3/level_3_problem_13_cache/k_c485bd631fb1/cvae_cuda/cuda.cu(127): error: identifier "kl_loss_out" is undefined atomicAdd(kl_loss_out, kl_per_sample); ^ 2 errors detected in the compilation of "/scratch/abaner/KernelBench/runs/pop_l123_st_gpt/popcorn/grok-4.3/level_3_problem_13_cache/k_c485bd631fb1/cvae_cuda/cuda.cu". ninja: build stopped: subcommand failed.
final submitted kernel (247 lines, 8,094 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 <curand_kernel.h>

__global__ void cvae_forward_kernel(
    const float* __restrict__ x,
    const float* __restrict__ y,
    const float* __restrict__ enc_weight,
    const float* __restrict__ enc_bias,
    const float* __restrict__ enc_mu_weight,
    const float* __restrict__ enc_mu_bias,
    const float* __restrict__ enc_logvar_weight,
    const float* __restrict__ enc_logvar_bias,
    const float* __restrict__ dec_weight1,
    const float* __restrict__ dec_bias1,
    const float* __restrict__ dec_weight2,
    const float* __restrict__ dec_bias2,
    float* __restrict__ loss_out,
    float* __recon_loss_out,
    float* __kl_loss_out,
    int batch_size,
    int input_dim,
    int label_dim,
    int hidden_dim,
    int latent_dim,
    unsigned long long seed
) {
    int b = blockIdx.x;
    if (b >= batch_size) return;

    extern __shared__ float smem[];

    // Simple implementation: each block handles one sample, using shared mem for intermediates
    float* h = smem;
    float* mu = &smem[hidden_dim];
    float* logvar = &smem[hidden_dim + latent_dim];
    float* z = &smem[hidden_dim + 2*latent_dim];
    float* h2 = &smem[hidden_dim + 3*latent_dim];
    float* x_recon = &smem[hidden_dim + 3*latent_dim + hidden_dim];

    // Encoder first linear + ReLU
    for (int i = threadIdx.x; i < hidden_dim; i += blockDim.x) {
        float sum = enc_bias[i];
        for (int j = 0; j < input_dim + label_dim; j++) {
            float val = (j < input_dim) ? x[b * input_dim + j] : y[b * label_dim + (j - input_dim)];
            sum += val * enc_weight[i * (input_dim + label_dim) + j];
        }
        h[i] = fmaxf(sum, 0.0f);
    }
    __syncthreads();

    // mu and logvar
    for (int i = threadIdx.x; i < latent_dim; i += blockDim.x) {
        float sum_mu = enc_mu_bias[i];
        float sum_logvar = enc_logvar_bias[i];
        for (int j = 0; j < hidden_dim; j++) {
            sum_mu += h[j] * enc_mu_weight[i * hidden_dim + j];
            sum_logvar += h[j] * enc_logvar_weight[i * hidden_dim + j];
        }
        mu[i] = sum_mu;
        logvar[i] = sum_logvar;
    }
    __syncthreads();

    // Sample z
    for (int i = threadIdx.x; i < latent_dim; i += blockDim.x) {
        float std = expf(0.5f * logvar[i]);
        curandState state;
        curand_init(seed + b * latent_dim + i, 0, 0, &state);
        float eps = curand_normal(&state);
        z[i] = mu[i] + std * eps;
    }
    __syncthreads();

    // Decoder
    for (int i = threadIdx.x; i < hidden_dim; i += blockDim.x) {
        float sum = dec_bias1[i];
        for (int j = 0; j < latent_dim + label_dim; j++) {
            float val = (j < latent_dim) ? z[j] : y[b * label_dim + (j - latent_dim)];
            sum += val * dec_weight1[i * (latent_dim + label_dim) + j];
        }
        h2[i] = fmaxf(sum, 0.0f);
    }
    __syncthreads();

    for (int i = threadIdx.x; i < input_dim; i += blockDim.x) {
        float sum = dec_bias2[i];
        for (int j = 0; j < hidden_dim; j++) {
            sum += h2[j] * dec_weight2[i * hidden_dim + j];
        }
        x_recon[i] = sum;
    }
    __syncthreads();

    // Compute per-sample recon and kl (will reduce later)
    __shared__ float partial_recon;
    __shared__ float partial_kl;
    if (threadIdx.x == 0) {
        partial_recon = 0.0f;
        partial_kl = 0.0f;
    }
    __syncthreads();

    float local_recon = 0.0f;
    float local_kl = 0.0f;
    for (int i = threadIdx.x; i < input_dim; i += blockDim.x) {
        float diff = x_recon[i] - x[b * input_dim + i];
        local_recon += diff * diff;
    }
    for (int i = threadIdx.x; i < latent_dim; i += blockDim.x) {
        local_kl += -0.5f * (1.0f + logvar[i] - mu[i]*mu[i] - expf(logvar[i]));
    }
    atomicAdd(&partial_recon, local_recon);
    atomicAdd(&partial_kl, local_kl);
    __syncthreads();

    if (threadIdx.x == 0) {
        float recon_per_sample = partial_recon / batch_size;
        float kl_per_sample = partial_kl / batch_size;
        atomicAdd(loss_out, recon_per_sample + kl_per_sample);
        atomicAdd(recon_loss_out, recon_per_sample);
        atomicAdd(kl_loss_out, kl_per_sample);
    }
}

std::vector<torch::Tensor> cvae_forward_cuda(
    torch::Tensor x,
    torch::Tensor y,
    torch::Tensor enc_weight,
    torch::Tensor enc_bias,
    torch::Tensor enc_mu_weight,
    torch::Tensor enc_mu_bias,
    torch::Tensor enc_logvar_weight,
    torch::Tensor enc_logvar_bias,
    torch::Tensor dec_weight1,
    torch::Tensor dec_bias1,
    torch::Tensor dec_weight2,
    torch::Tensor dec_bias2,
    int hidden_dim,
    int latent_dim
) {
    int batch_size = x.size(0);
    int input_dim = x.size(1);
    int label_dim = y.size(1);

    auto options = torch::TensorOptions().device(x.device()).dtype(torch::kFloat32);
    auto loss_out = torch::zeros({}, options);
    auto recon_out = torch::zeros({}, options);
    auto kl_out = torch::zeros({}, options);

    int threads = 128;
    size_t smem_size = (hidden_dim + 3*latent_dim + hidden_dim + input_dim) * sizeof(float);

    unsigned long long seed = 123456789ULL;

    cvae_forward_kernel<<<batch_size, threads, smem_size>>>(
        x.data_ptr<float>(),
        y.data_ptr<float>(),
        enc_weight.data_ptr<float>(),
        enc_bias.data_ptr<float>(),
        enc_mu_weight.data_ptr<float>(),
        enc_mu_bias.data_ptr<float>(),
        enc_logvar_weight.data_ptr<float>(),
        enc_logvar_bias.data_ptr<float>(),
        dec_weight1.data_ptr<float>(),
        dec_bias1.data_ptr<float>(),
        dec_weight2.data_ptr<float>(),
        dec_bias2.data_ptr<float>(),
        loss_out.data_ptr<float>(),
        recon_out.data_ptr<float>(),
        kl_out.data_ptr<float>(),
        batch_size, input_dim, label_dim, hidden_dim, latent_dim, seed
    );

    return {loss_out, recon_out, kl_out};
}
'''

cpp_source = r'''
std::vector<torch::Tensor> cvae_forward_cuda(
    torch::Tensor x, torch::Tensor y,
    torch::Tensor enc_weight, torch::Tensor enc_bias,
    torch::Tensor enc_mu_weight, torch::Tensor enc_mu_bias,
    torch::Tensor enc_logvar_weight, torch::Tensor enc_logvar_bias,
    torch::Tensor dec_weight1, torch::Tensor dec_bias1,
    torch::Tensor dec_weight2, torch::Tensor dec_bias2,
    int hidden_dim, int latent_dim
);
'''

module = load_inline(
    name='cvae_cuda',
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=['cvae_forward_cuda'],
    verbose=False,
    extra_cuda_cflags=['-O3']
)

class ModelNew(nn.Module):
    def __init__(self, input_dim, label_dim, hidden_dim, latent_dim):
        super().__init__()
        self.enc = nn.Sequential(
            nn.Linear(input_dim + label_dim, hidden_dim),
            nn.ReLU(),
        )
        self.enc_mu = nn.Linear(hidden_dim, latent_dim)
        self.enc_logvar = nn.Linear(hidden_dim, latent_dim)
        self.dec = nn.Sequential(
            nn.Linear(latent_dim + label_dim, hidden_dim),
            nn.ReLU(),
            nn.Linear(hidden_dim, input_dim),
        )
        self.input_dim = input_dim
        self.label_dim = label_dim
        self.hidden_dim = hidden_dim
        self.latent_dim = latent_dim

    def forward(self, x: torch.Tensor, y: torch.Tensor) -> tuple:
        # Extract weights
        enc_w = self.enc[0].weight
        enc_b = self.enc[0].bias
        mu_w = self.enc_mu.weight
        mu_b = self.enc_mu.bias
        logvar_w = self.enc_logvar.weight
        logvar_b = self.enc_logvar.bias
        dec_w1 = self.dec[0].weight
        dec_b1 = self.dec[0].bias
        dec_w2 = self.dec[2].weight
        dec_b2 = self.dec[2].bias

        outs = module.cvae_forward_cuda(
            x, y, enc_w, enc_b, mu_w, mu_b, logvar_w, logvar_b,
            dec_w1, dec_b1, dec_w2, dec_b2, self.hidden_dim, self.latent_dim
        )
        loss = outs[0]
        recon = outs[1]
        kl = outs[2]
        return loss, recon, kl