grok-4.3 · popcorn · pop_l123_default_gpt
correct 44_ReparameterizationTrick.py
model
grok-4.3
turns / max
4 / 12
tool calls
3
tokens (total)
24510
wall clock
476.5s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
0.46 μs
ref runtime
0.34 μs
speedup
0.74x
started
2026-05-11T14:09:13Z
timingkernelbench.timing
kernel runtime
0.46 μs
ref runtime
0.34 μs
speedup
0.74x
kernel mean / median / std
0.46 / — / 0.03 μs
ref mean / median / std
0.34 / — / 0.06 μs
numerical precisioneval_kernel_against_ref
max abs error
3.81e-05
mean abs error
1.38e-07
max rel error
4.38e-07
mean rel error
1.51e-09
gpu memoryextended_metrics.measure_memory
kernel peak
34.01 MB
ref peak
34.69 MB
kernel / ref
0.98x
kernel launches / fusiontorch.profiler
kernel launches
67
ref launches
113
fusion ratio (ref / kernel)
1.69x
total cuda time
1236 μs
energyNVML (pynvml)
kernel mJ / run
274.94
ref mJ / run
0.00
ratio (ref / kernel)
—
avg power
744.88 W
ref avg power
0.00 W
SOL (speed-of-light)heuristic_from_speedup
SOL score (0–1)
0.074
dram util %
—
compute util %
—
bottleneck
unknown
dominant pipe
—
arithmetic intensity
—
ridge point
—
roofline / occupancyheuristic
dram bandwidth
—
dram util %
—
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
—
bottleneck
—
dominant pipe %
—
arithmetic intensity
—
ridge point
—
L1 hit %
—
L2 hit %
—
ld sectors / req
—
st sectors / req
—
regs / thread
—
smem / block
—
block size
—
peak bw
—
peak fp32 tflops
—
peak fp16 tflops
—
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn


class Model(nn.Module):
    """
    Reparameterization trick for differentiable sampling from a diagonal
    Gaussian.  Separates stochasticity from parameters so that gradients
    flow through mu and log_var.  Also demonstrates a "mixture of
    Gaussians" reparameterization with learnable component weights.
    """

    def __init__(self, latent_dim, num_components=1):
        super().__init__()
        self.latent_dim = latent_dim
        self.num_components = num_components
        if num_components > 1:
            self.logits = nn.Parameter(torch.zeros(num_components))

    def forward(self, mu: torch.Tensor, log_var: torch.Tensor) -> tuple:
        """
        Args:
            mu:      (B, K, D) or (B, D) – means
            log_var: (B, K, D) or (B, D) – log-variances
        Returns:
            (z, log_q_z):
                z:       (B, D) – sampled latents
                log_q_z: (B,)   – log-density of sample under q
        """
        std = torch.exp(0.5 * log_var)

        if self.num_components > 1 and mu.dim() == 3:
            weights = torch.softmax(self.logits, dim=0)  # (K,)
            # Sample component
            comp = torch.multinomial(weights.expand(mu.shape[0], -1), 1).squeeze(-1)
            batch_idx = torch.arange(mu.shape[0], device=mu.device)
            mu_k = mu[batch_idx, comp]
            std_k = std[batch_idx, comp]
            z = mu_k + std_k * torch.randn_like(mu_k)
            # log q(z) under mixture
            log_probs = -0.5 * ((z.unsqueeze(1) - mu) / std) ** 2 - log_var * 0.5 - 0.5 * torch.log(torch.tensor(2 * 3.14159265))
            log_probs = log_probs.sum(-1) + weights.log().unsqueeze(0)
            log_q_z = torch.logsumexp(log_probs, dim=1)
        else:
            if mu.dim() == 3:
                mu = mu[:, 0]
                std = std[:, 0]
                log_var = log_var[:, 0]
            z = mu + std * torch.randn_like(mu)
            log_q_z = (-0.5 * ((z - mu) / std) ** 2 - 0.5 * log_var - 0.5 * torch.log(torch.tensor(2 * 3.14159265))).sum(-1)

        return z, log_q_z


latent_dim = 64
num_components = 4
batch_size = 256


def get_inputs():
    p = popcorn_pri
    mu = torch.randn(p.jitter_int(batch_size), num_components, latent_dim)
    log_var = torch.randn(p.jitter_int(batch_size), num_components, latent_dim) * 0.5
    return [mu, log_var]


def get_init_inputs():
    return [latent_dim, num_components]
turn 036.8s turn · 36.8s LLM · 0.0s tools · tokens in=1,891 · out=1,121 · total=4,307
reasoning
The task is to replace the forward computation of this PyTorch model with a custom CUDA kernel. The model is a reparameterization trick for sampling from a diagonal Gaussian, with a mixture of Gaussia...I need to replace the forward computation of a PyTorch model with a custom CUDA kernel. The model uses a reparameterization trick for sampling from a diagonal Gaussian, possibly with a mixture of Gaussians when there are multiple components. My ModelNew class will take the same inputs as the original: latent_dim and num_components in its constructor. In the forward method, it receives mu and log_var, shaped as (batch_size, num_components, latent_dim) when there are multiple components. It should output z and log_q_z. For the mixture case, I first compute weights as softmax of some logits, but actually, the internal reasoning mentions softmax(logits), but logits aren't mentioned in the question, wait, the question is truncated, but assuming it's a standard VAE with mixture. To implement efficiently in CUDA. First, understand the math. For single component, it's simple: std = exp(0.5 * log_var), z = mu + std * randn, log_q_z = -0 (truncated) For mixture, sample component k ~ multinomial(weights), then use mu_k, log_var_k for that component. Then z = mu_k + std_k * randn For log_q_z, it's the log prob under the mixture.
assistant```python 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 <math.h> __global__ void mixture_logprob_kernel( const float* __restrict__ z, const float* __restrict__ mu, const float* __restrict__ std, const float* __restrict__ log_var, const float* __restrict__ log_weights, float* __restrict__ log_q_z, int B, int K, int D, float log_2pi ) { int b = blockIdx.x * blockDim.x + threadIdx.x; if (b >= B) return; float max_log = -1e30f; float sum_exp = 0.0f; for (int k = 0; k < K; k++) { float sum = 0.0f; for (int d = 0; d < D; d++) { float diff = (z[b * D + d] - mu[(b * K + k) * D + d]) / std[(b * K + k) * D + d]; sum += -0.5f * diff * diff - 0.5f * log_var[(b * K + k) * D + d] - 0.5f * log_2pi; } float lp = sum + log_weights[k]; if (lp > max_log) max_log = lp; } for (int k = 0; k < K; k++) { float sum = 0.0f; for (int d = 0; d < D; d++) { float diff = (z[b * D + d] - mu[(b * K + k) * D + d]) / std[(b * K + k) * D + d]; sum += -0.5f * diff * diff - 0.5f * log_var[(b * K + k) * D + d] - 0.5f * log_2pi; } float lp = sum + log_weights[k]; sum_exp += expf(lp - max_log); } log_q_z[b] = max_log + logf(sum_exp); } torch::Tensor compute_log_q_z(torch::Tensor z, torch::Tensor mu, torch::Tensor std, torch::Tensor log_var, torch::Tensor log_weights) { int B = z.size(0); int K = mu.size(1); int D = z.size(1); auto log_q_z = torch::empty({B}, mu.options()); float log_2pi = logf(2.0f * 3.14159265f); const int threads = 256; const int blocks = (B + threads - 1) / threads; mixture_logprob_kernel<<<blocks, threads>>>( z.data_ptr<float>(), mu.data_ptr<float>(), std.data_ptr<float>(), log_var.data_ptr<float>(), log_weights.data_ptr<float>(), log_q_z.data_ptr<float>(), B, K, D, log_2pi ); return log_q_z; } ''' cpp_source = ''' torch::Tensor compute_log_q_z(torch::Tensor z, torch::Tensor mu, torch::Tensor std, torch::Tensor log_var, torch::Tensor log_weights); ''' module = load_inline( name='mixture_reparam', cpp_sources=[cpp_source], cuda_sources=[cuda_source], functions=['compute_log_q_z'], verbose=False ) class ModelNew(nn.Module): def __init__(self, latent_dim, num_components=1): super().__init__() self.latent_dim = latent_dim self.num_components = num_components if num_components > 1: self.logits = nn.Parameter(torch.zeros(num_components)) def forward(self, mu: torch.Tensor, log_var: torch.Tensor) -> tuple: std = torch.exp(0.5 * log_var) if self.num_components > 1 and mu.dim() == 3: weights = torch.softmax(self.logits, dim=0) comp = torch.multinomial(weights.expand(mu.shape[0], -1), 1).squeeze(-1) batch_idx = torch.arange(mu.shape[0], device=mu.device) mu_k = mu[batch_idx, comp] std_k = std[batch_idx, comp] z = mu_k + std_k * torch.randn_like(mu_k) log_weights = weights.log() log_q_z = module.compute_log_q_z(z, mu, std, log_var, log_weights) else: if mu.dim() == 3: mu = mu[:, 0] std = std[:, 0] log_var = log_var[:, 0] z = mu + std * torch.randn_like(mu) log_q_z = (-0.5 * ((z - mu) / std) ** 2 - 0.5 * log_var - 0.5 * torch.log(torch.tensor(2 * 3.14159265, device=mu.device))).sum(-1) return z, log_q_z ```
turn 1108.4s turn · 14.3s LLM · 94.1s tools · tokens in=4,375 · out=1,100 · total=5,577
reasoning
The instruction says: "You have 12 turns and 40 tool calls. Each turn is one response from you, optionally with tool calls."
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 <math.h>\n\n__global__ void mixture_logprob_kernel(\n const float* __restrict__ z,\n const float* __restrict__ mu,\n const float* __restrict__ std,\n const float* __restrict__ log_var,\n const float* __restrict__ log_weights,\n float* __restrict__ log_q_z,\n int B, int K, int D, float log_2pi\n) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= B) return;\n\n float max_log = -1e30f;\n for (int k = 0; k < K; k++) {\n float sum = 0.0f;\n for (int d = 0; d < D; d++) {\n int idx = (b * K + k) * D + d;\n float diff = (z[b * D + d] - mu[idx]) / std[idx];\n sum += -0.5f * diff * diff - 0.5f * log_var[idx] - 0.5f * log_2pi;\n }\n float lp = sum + log_weights[k];\n if (lp > max_log) max_log = lp;\n }\n float sum_exp = 0.0f;\n for (int k = 0; k < K; k++) {\n float sum = 0.0f;\n for (int d = 0; d < D; d++) {\n int idx = (b * K + k) * D + d;\n float diff = (z[b * D + d] - mu[idx]) / std[idx];\n sum += -0.5f * diff * diff - 0.5f * log_var[idx] - 0.5f * log_2pi;\n }\n float lp = sum + log_weights[k];\n sum_exp += expf(lp - max_log);\n }\n log_q_z[b] = max_log + logf(sum_exp);\n}\n\ntorch::Tensor compute_log_q_z(torch::Tensor z, torch::Tensor mu, torch::Tensor std, torch::Tensor log_var, torch::Tensor log_weights) {\n int B = z.size(0);\n int K = mu.size(1);\n int D = z.size(1);\n auto log_q_z = torch::empty({B}, mu.options());\n float log_2pi = logf(2.0f * 3.14159265f);\n\n const int threads = 256;\n const int blocks = (B + threads - 1) / threads;\n\n mixture_logprob_kernel<<<blocks, threads>>>(\n z.data_ptr<float>(),\n mu.data_ptr<float>(),\n std.data_ptr<float>(),\n log_var.data_ptr<float>(),\n log_weights.data_ptr<float>(),\n log_q_z.data_ptr<float>(),\n B, K, D, log_2pi\n );\n return log_q_z;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor compute_log_q_z(torch::Tensor z, torch::Tensor mu, torch::Tensor std, torch::Tensor log_var, torch::Tensor log_weights);\n'''\n\nmodule = load_inline(\n name='mixture_reparam',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['compute_log_q_z'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, latent_dim, num_components=1):\n super().__init__()\n self.latent_dim = latent_dim\n self.num_components = num_components\n if num_components > 1:\n self.logits = nn.Parameter(torch.zeros(num_components))\n\n def forward(self, mu: torch.Tensor, log_var: torch.Tensor) -> tuple:\n std = torch.exp(0.5 * log_var)\n\n if self.num_components > 1 and mu.dim() == 3:\n weights = torch.softmax(self.logits, dim=0)\n comp = torch.multinomial(weights.expand(mu.shape[0], -1), 1).squeeze(-1)\n batch_idx = torch.arange(mu.shape[0], device=mu.device)\n mu_k = mu[batch_idx, comp]\n std_k = std[batch_idx, comp]\n z = mu_k + std_k * torch.randn_like(mu_k)\n log_weights = weights.log()\n log_q_z = module.compute_log_q_z(z, mu, std, log_var, log_weights)\n else:\n if mu.dim() == 3:\n mu = mu[:, 0]\n std = std[:, 0]\n log_var = log_var[:, 0]\n z = mu + std * torch.randn_like(mu)\n log_q_z = (-0.5 * ((z - mu) / std) ** 2 - 0.5 * log_var - 0.5 * torch.log(torch.tensor(2 * 3.14159265, device=mu.device))).sum(-1)\n\n return z, log_q_z\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 215.3s turn · 14.2s LLM · 1.1s tools · tokens in=5,594 · out=1,101 · total=6,705
reasoning
Compilation succeeded. Now I should run correctness check.
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 <math.h>\n\n__global__ void mixture_logprob_kernel(\n const float* __restrict__ z,\n const float* __restrict__ mu,\n const float* __restrict__ std,\n const float* __restrict__ log_var,\n const float* __restrict__ log_weights,\n float* __restrict__ log_q_z,\n int B, int K, int D, float log_2pi\n) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= B) return;\n\n float max_log = -1e30f;\n for (int k = 0; k < K; k++) {\n float sum = 0.0f;\n for (int d = 0; d < D; d++) {\n int idx = (b * K + k) * D + d;\n float diff = (z[b * D + d] - mu[idx]) / std[idx];\n sum += -0.5f * diff * diff - 0.5f * log_var[idx] - 0.5f * log_2pi;\n }\n float lp = sum + log_weights[k];\n if (lp > max_log) max_log = lp;\n }\n float sum_exp = 0.0f;\n for (int k = 0; k < K; k++) {\n float sum = 0.0f;\n for (int d = 0; d < D; d++) {\n int idx = (b * K + k) * D + d;\n float diff = (z[b * D + d] - mu[idx]) / std[idx];\n sum += -0.5f * diff * diff - 0.5f * log_var[idx] - 0.5f * log_2pi;\n }\n float lp = sum + log_weights[k];\n sum_exp += expf(lp - max_log);\n }\n log_q_z[b] = max_log + logf(sum_exp);\n}\n\ntorch::Tensor compute_log_q_z(torch::Tensor z, torch::Tensor mu, torch::Tensor std, torch::Tensor log_var, torch::Tensor log_weights) {\n int B = z.size(0);\n int K = mu.size(1);\n int D = z.size(1);\n auto log_q_z = torch::empty({B}, mu.options());\n float log_2pi = logf(2.0f * 3.14159265f);\n\n const int threads = 256;\n const int blocks = (B + threads - 1) / threads;\n\n mixture_logprob_kernel<<<blocks, threads>>>(\n z.data_ptr<float>(),\n mu.data_ptr<float>(),\n std.data_ptr<float>(),\n log_var.data_ptr<float>(),\n log_weights.data_ptr<float>(),\n log_q_z.data_ptr<float>(),\n B, K, D, log_2pi\n );\n return log_q_z;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor compute_log_q_z(torch::Tensor z, torch::Tensor mu, torch::Tensor std, torch::Tensor log_var, torch::Tensor log_weights);\n'''\n\nmodule = load_inline(\n name='mixture_reparam',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['compute_log_q_z'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, latent_dim, num_components=1):\n super().__init__()\n self.latent_dim = latent_dim\n self.num_components = num_components\n if num_components > 1:\n self.logits = nn.Parameter(torch.zeros(num_components))\n\n def forward(self, mu: torch.Tensor, log_var: torch.Tensor) -> tuple:\n std = torch.exp(0.5 * log_var)\n\n if self.num_components > 1 and mu.dim() == 3:\n weights = torch.softmax(self.logits, dim=0)\n comp = torch.multinomial(weights.expand(mu.shape[0], -1), 1).squeeze(-1)\n batch_idx = torch.arange(mu.shape[0], device=mu.device)\n mu_k = mu[batch_idx, comp]\n std_k = std[batch_idx, comp]\n z = mu_k + std_k * torch.randn_like(mu_k)\n log_weights = weights.log()\n log_q_z = module.compute_log_q_z(z, mu, std, log_var, log_weights)\n else:\n if mu.dim() == 3:\n mu = mu[:, 0]\n std = std[:, 0]\n log_var = log_var[:, 0]\n z = mu + std * torch.randn_like(mu)\n log_q_z = (-0.5 * ((z - mu) / std) ** 2 - 0.5 * log_var - 0.5 * torch.log(torch.tensor(2 * 3.14159265, device=mu.device))).sum(-1)\n\n return z, log_q_z\n" }
result ← run_correctness (ok)
okrun_correctness run_correctness PASSED: (5 / 5) trials all matched the reference. Numerical precision: max_abs_err=3.81e-05 mean_abs_err=1.38e-07 max_rel_err=4.38e-07
turn 3316.0s turn · 15.0s LLM · 301.0s tools · tokens in=6,765 · out=1,100 · total=7,921 · FINAL
reasoning
Correctness passed with good tolerance. Now, since it's correct and should be faster (the kernel computes logsumexp efficiently, though tiny data), I can submit.
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 <math.h>\n\n__global__ void mixture_logprob_kernel(\n const float* __restrict__ z,\n const float* __restrict__ mu,\n const float* __restrict__ std,\n const float* __restrict__ log_var,\n const float* __restrict__ log_weights,\n float* __restrict__ log_q_z,\n int B, int K, int D, float log_2pi\n) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= B) return;\n\n float max_log = -1e30f;\n for (int k = 0; k < K; k++) {\n float sum = 0.0f;\n for (int d = 0; d < D; d++) {\n int idx = (b * K + k) * D + d;\n float diff = (z[b * D + d] - mu[idx]) / std[idx];\n sum += -0.5f * diff * diff - 0.5f * log_var[idx] - 0.5f * log_2pi;\n }\n float lp = sum + log_weights[k];\n if (lp > max_log) max_log = lp;\n }\n float sum_exp = 0.0f;\n for (int k = 0; k < K; k++) {\n float sum = 0.0f;\n for (int d = 0; d < D; d++) {\n int idx = (b * K + k) * D + d;\n float diff = (z[b * D + d] - mu[idx]) / std[idx];\n sum += -0.5f * diff * diff - 0.5f * log_var[idx] - 0.5f * log_2pi;\n }\n float lp = sum + log_weights[k];\n sum_exp += expf(lp - max_log);\n }\n log_q_z[b] = max_log + logf(sum_exp);\n}\n\ntorch::Tensor compute_log_q_z(torch::Tensor z, torch::Tensor mu, torch::Tensor std, torch::Tensor log_var, torch::Tensor log_weights) {\n int B = z.size(0);\n int K = mu.size(1);\n int D = z.size(1);\n auto log_q_z = torch::empty({B}, mu.options());\n float log_2pi = logf(2.0f * 3.14159265f);\n\n const int threads = 256;\n const int blocks = (B + threads - 1) / threads;\n\n mixture_logprob_kernel<<<blocks, threads>>>(\n z.data_ptr<float>(),\n mu.data_ptr<float>(),\n std.data_ptr<float>(),\n log_var.data_ptr<float>(),\n log_weights.data_ptr<float>(),\n log_q_z.data_ptr<float>(),\n B, K, D, log_2pi\n );\n return log_q_z;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor compute_log_q_z(torch::Tensor z, torch::Tensor mu, torch::Tensor std, torch::Tensor log_var, torch::Tensor log_weights);\n'''\n\nmodule = load_inline(\n name='mixture_reparam',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['compute_log_q_z'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, latent_dim, num_components=1):\n super().__init__()\n self.latent_dim = latent_dim\n self.num_components = num_components\n if num_components > 1:\n self.logits = nn.Parameter(torch.zeros(num_components))\n\n def forward(self, mu: torch.Tensor, log_var: torch.Tensor) -> tuple:\n std = torch.exp(0.5 * log_var)\n\n if self.num_components > 1 and mu.dim() == 3:\n weights = torch.softmax(self.logits, dim=0)\n comp = torch.multinomial(weights.expand(mu.shape[0], -1), 1).squeeze(-1)\n batch_idx = torch.arange(mu.shape[0], device=mu.device)\n mu_k = mu[batch_idx, comp]\n std_k = std[batch_idx, comp]\n z = mu_k + std_k * torch.randn_like(mu_k)\n log_weights = weights.log()\n log_q_z = module.compute_log_q_z(z, mu, std, log_var, log_weights)\n else:\n if mu.dim() == 3:\n mu = mu[:, 0]\n std = std[:, 0]\n log_var = log_var[:, 0]\n z = mu + std * torch.randn_like(mu)\n log_q_z = (-0.5 * ((z - mu) / std) ** 2 - 0.5 * log_var - 0.5 * torch.log(torch.tensor(2 * 3.14159265, device=mu.device))).sum(-1)\n\n return z, log_q_z\n" }
result ← submit_kernel (ok)
oksubmit_kernel submit_kernel PASSED: (5 / 5) correctness trials all passed. Kernel runtime: 0.46 μs Runtime stats: mean=0.46μs median=?μs std=0.03μs Numerical precision: max_abs_err=3.81e-05 mean_abs_err=1.38e-07 max_rel_err=4.38e-07 Memory: 34.01 MB peak (ref: 34.69 MB, ratio: 0.9802x) Kernel launches: 67 (ref: 113, fusion ratio: 1.6866) Energy: 274.94 mJ/run (ref: 0.0 mJ, ratio: -1x) SOL score: 0.0742
final submitted kernel (112 lines, 3,730 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 <math.h>

__global__ void mixture_logprob_kernel(
    const float* __restrict__ z,
    const float* __restrict__ mu,
    const float* __restrict__ std,
    const float* __restrict__ log_var,
    const float* __restrict__ log_weights,
    float* __restrict__ log_q_z,
    int B, int K, int D, float log_2pi
) {
    int b = blockIdx.x * blockDim.x + threadIdx.x;
    if (b >= B) return;

    float max_log = -1e30f;
    for (int k = 0; k < K; k++) {
        float sum = 0.0f;
        for (int d = 0; d < D; d++) {
            int idx = (b * K + k) * D + d;
            float diff = (z[b * D + d] - mu[idx]) / std[idx];
            sum += -0.5f * diff * diff - 0.5f * log_var[idx] - 0.5f * log_2pi;
        }
        float lp = sum + log_weights[k];
        if (lp > max_log) max_log = lp;
    }
    float sum_exp = 0.0f;
    for (int k = 0; k < K; k++) {
        float sum = 0.0f;
        for (int d = 0; d < D; d++) {
            int idx = (b * K + k) * D + d;
            float diff = (z[b * D + d] - mu[idx]) / std[idx];
            sum += -0.5f * diff * diff - 0.5f * log_var[idx] - 0.5f * log_2pi;
        }
        float lp = sum + log_weights[k];
        sum_exp += expf(lp - max_log);
    }
    log_q_z[b] = max_log + logf(sum_exp);
}

torch::Tensor compute_log_q_z(torch::Tensor z, torch::Tensor mu, torch::Tensor std, torch::Tensor log_var, torch::Tensor log_weights) {
    int B = z.size(0);
    int K = mu.size(1);
    int D = z.size(1);
    auto log_q_z = torch::empty({B}, mu.options());
    float log_2pi = logf(2.0f * 3.14159265f);

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

    mixture_logprob_kernel<<<blocks, threads>>>(
        z.data_ptr<float>(),
        mu.data_ptr<float>(),
        std.data_ptr<float>(),
        log_var.data_ptr<float>(),
        log_weights.data_ptr<float>(),
        log_q_z.data_ptr<float>(),
        B, K, D, log_2pi
    );
    return log_q_z;
}
'''

cpp_source = '''
torch::Tensor compute_log_q_z(torch::Tensor z, torch::Tensor mu, torch::Tensor std, torch::Tensor log_var, torch::Tensor log_weights);
'''

module = load_inline(
    name='mixture_reparam',
    cpp_sources=[cpp_source],
    cuda_sources=[cuda_source],
    functions=['compute_log_q_z'],
    verbose=False
)

class ModelNew(nn.Module):
    def __init__(self, latent_dim, num_components=1):
        super().__init__()
        self.latent_dim = latent_dim
        self.num_components = num_components
        if num_components > 1:
            self.logits = nn.Parameter(torch.zeros(num_components))

    def forward(self, mu: torch.Tensor, log_var: torch.Tensor) -> tuple:
        std = torch.exp(0.5 * log_var)

        if self.num_components > 1 and mu.dim() == 3:
            weights = torch.softmax(self.logits, dim=0)
            comp = torch.multinomial(weights.expand(mu.shape[0], -1), 1).squeeze(-1)
            batch_idx = torch.arange(mu.shape[0], device=mu.device)
            mu_k = mu[batch_idx, comp]
            std_k = std[batch_idx, comp]
            z = mu_k + std_k * torch.randn_like(mu_k)
            log_weights = weights.log()
            log_q_z = module.compute_log_q_z(z, mu, std, log_var, log_weights)
        else:
            if mu.dim() == 3:
                mu = mu[:, 0]
                std = std[:, 0]
                log_var = log_var[:, 0]
            z = mu + std * torch.randn_like(mu)
            log_q_z = (-0.5 * ((z - mu) / std) ** 2 - 0.5 * log_var - 0.5 * torch.log(torch.tensor(2 * 3.14159265, device=mu.device))).sum(-1)

        return z, log_q_z