FW-GLM-5-1 · popcorn · pop_l123_st_gpt
incorrect 44_ReparameterizationTrick.py
model
FW-GLM-5-1
turns / max
1 / 1
tool calls
1
tokens (total)
6315
wall clock
823.7s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
—
ref runtime
0.42 μs
speedup
—
started
2026-05-11T09:03:23Z
correctness issue:
timingkernelbench.timing
kernel runtime
—
ref runtime
0.42 μs
speedup
—
kernel mean / median / std
—
ref mean / median / std
0.42 / — / 0.03 μs
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 0 · FINAL · 1 tool call(s) · 823.7s turn · 47.8s LLM · tokens in=1,475 · out=4,840 · total=6,315 · 20,333 chars (click to expand)
turn 0823.7s turn · 47.8s LLM · 775.9s tools · tokens in=1,475 · out=4,840 · total=6,315 · FINAL
reasoning
Let me analyze the reference implementation to understand what needs to be computed: 1. The model does reparameterization trick for Gaussian sampling 2. It supports mixture of Gaussians with `num_components` components 3. When `num_components > 1` and `mu.dim() == 3`: - Compute weights via softmax of logits - Sample component using multinomial - Select mu_k and std_k based on sampled component - Sample z = mu_k + std_k * noise - Compute log q(z) under mixture using logsumexp 4. Otherwise: - If 3D, take first component - Sample z = mu + std * noise - Compute log q(z) as Gaussian log-density The key operations: 1. `std = exp(0.5 * log_var)` - element-wise 2. Softmax of logits 3. Multinomial sampling for component selection 4. Random normal sampling 5. Various arithmetic for log probability computation 6. Logsumexp reduction For a CUDA kernel, I need to: 1. Generate random numbers (or use torch.randn_like externally) 2. Handle the mixture case with sampling 3. Compute the log probabilities Actually, looking at this more carefully, some operations like multinomial sampling are complex to implement in CUDA efficiently. Let me think about what can be efficiently parallelized: 1. Computing `std = exp(0.5 * log_var)` - trivial element-wise 2. Sampling z after component selection - element-wise with random 3. Computing log_q_z - involves reduction (logsumexp) The logsumexp is the main reduction operation. Let me write a kernel that: 1. Uses PyTorch for the multinomial sampling (hard to beat) 2. Uses custom CUDA for the dense computation like std computation, z sampling, and log probability computation Actually, let me reconsider. The main computational bottleneck would be: - For the mixture case: the log_probs computation over all K components for all D dimensions, then logsumexp over K - For single component: element-wise operations and sum over D Let me write a fused kernel that handles the sampling and log probability computation together. Key insight: We can fuse multiple operations: 1. std computation 2. Component selection (can use torch.multinomial, or implement our own) 3. z sampling 4. Log probability computation and reduction Let me start with a simpler approach - fuse what we can and use PyTorch for complex sampling operations. Actually, for a real speedup, I should focus on: 1. Fusing the std computation, random sampling, and z computation 2. Fusing the log probability computation with the logsumexp The logsumexp reduction across components is interesting - we can use parallel reduction. Let me write a comprehensive CUDA kernel: ```cpp #include <torch/extension.h> #include <cuda.h> #include <cuda_runtime.h> #include <curand_kernel.h> #include <cmath> // For mixture case __global__ void mixture_reparam_kernel( const float* mu, const float* log_var, const float* weights, const int* components, float* z, float* log_q_z, int batch_size, int num_components, int latent_dim, uint64_t seed ) { int b = blockIdx.x * blockDim.x + threadIdx.x; if (b >= batch_size) return; // Initialize random state curandState state; curand_init(seed, b, 0, &state); // Get selected component int comp = components[b]; // Sample z and compute log_q float log_q_sum = 0.0f; for (int d = 0; d < latent_dim; d++) { int idx_k = b * num_components * latent_dim + comp * latent_dim + d; float mu_val = mu[idx_k]; float log_var_val = log_var[idx_k]; float std_val = expf(0.5f * log_var_val); float noise = curand_normal(&state); float z_val = mu_val + std_val * noise; z[b * latent_dim + d] = z_val; // For log_q we need to compute over all components... } // ... more complex log_q computation } ``` Actually this is getting complex. Let me think about a cleaner approach. The main bottleneck operations: 1. Element-wise: std = exp(0.5 * log_var) - easy to fuse 2. Sampling: z = mu + std * randn - can fuse with std 3. Log prob: log_q computation - can fuse 4. Logsumexp: reduction - can optimize For the non-mixture case, it's simple element-wise operations followed by a sum reduction. For the mixture case: - Need to sample component (multinomial) - Then similar to non-mixture - But log_q needs to consider all components Let me write a kernel that handles both cases efficiently. I'll create a design where: 1. Use a separate kernel for multinomial sampling (or use torch's) 2. Fuse std computation, sampling, and z computation 3. Fuse log probability computation with reduction Here's my approach: ```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 <curand_kernel.h> #include <cmath> // Constants constexpr float LOG_2PI = 1.8378770664093453f; // log(2*pi) // Kernel for single component case __global__ void single_component_kernel( const float* __restrict__ mu, const float* __restrict__ log_var, float* __restrict__ z, float* __restrict__ log_q_z, int batch_size, int latent_dim, uint64_t seed ) { int b = blockIdx.x; int d = threadIdx.x; if (b >= batch_size || d >= latent_dim) return; // Initialize random state curandState state; curand_init(seed, b * latent_dim + d, 0, &state); int idx = b * latent_dim + d; float mu_val = mu[idx]; float log_var_val = log_var[idx]; float std_val = expf(0.5f * log_var_val); float noise = curand_normal(&state); float z_val = mu_val + std_val * noise; z[idx] = z_val; // Compute contribution to log_q float log_p = -0.5f * noise * noise - 0.5f * log_var_val - 0.5f * LOG_2PI; // Block-level reduction for log_q_z __shared__ float sdata[1024]; sdata[d] = log_p; __syncthreads(); // Reduce for (int s = blockDim.x >> 1; s > 0; s >>= 1) { if (d < s) { sdata[d] += sdata[d + s]; } __syncthreads(); } if (d == 0) { log_q_z[b] = sdata[0]; } } // ... more kernels """ cpp_source = """ torch::Tensor single_component_forward( torch::Tensor mu, torch::Tensor log_var, int64_t latent_dim, uint64_t seed ); """ # Load extension... ``` Actually, let me simplify and think about what really matters for performance: 1. The element-wise operations (exp, multiply, add) are memory-bound 2. Random number generation is compute-heavy 3. Reductions need efficient parallel algorithms For practical speedup, I should: 1. Fuse memory operations to reduce memory traffic 2. Use efficient random number generation 3. Use warp-level primitives for reductions Let me write a complete implementation:
call → submit_kernel
callsubmit_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\nimport math\n\ncuda_source = r\"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <curand_kernel.h>\n#include <cmath>\n#include <limits>\n\n// Constants\nstatic constexpr float LOG_2PI = 1.8378770664093453f; // log(2*pi)\nstatic constexpr float PI = 3.14159265358979323846f;\n\n// Warp-level reduction for sum\n__inline__ __device__ float warp_reduce_sum(float val) {\n for (int offset = 16; offset > 0; offset >>= 1) {\n val += __shfl_down_sync(0xffffffff, val, offset);\n }\n return val;\n}\n\n// Warp-level reduction for max\n__inline__ __device__ float warp_reduce_max(float val) {\n for (int offset = 16; offset > 0; offset >>= 1) {\n val = fmaxf(val, __shfl_down_sync(0xffffffff, val, offset));\n }\n return val;\n}\n\n// Single component kernel (for 2D input)\n__global__ void single_component_2d_kernel(\n const float* __restrict__ mu,\n const float* __restrict__ log_var,\n float* __restrict__ z,\n float* __restrict__ log_q_z,\n const int batch_size,\n const int latent_dim,\n const uint64_t seed\n) {\n const int b = blockIdx.x;\n const int tid = threadIdx.x;\n const int lane_id = tid & 31;\n const int warp_id = tid >> 5;\n const int warps_per_block = blockDim.x >> 5;\n \n if (b >= batch_size) return;\n \n // Each thread handles multiple latent dimensions\n float local_sum = 0.0f;\n \n // Initialize random state per thread\n curandState state;\n curand_init(seed, b * blockDim.x + tid, 0, &state);\n \n for (int d = tid; d < latent_dim; d += blockDim.x) {\n const int idx = b * latent_dim + d;\n const float mu_val = mu[idx];\n const float log_var_val = log_var[idx];\n const float std_val = expf(0.5f * log_var_val);\n \n const float noise = curand_normal(&state);\n const float z_val = mu_val + std_val * noise;\n z[idx] = z_val;\n \n // log_p contribution: -0.5 * ((z-mu)/std)^2 - 0.5*log_var - 0.5*log(2*pi)\n // = -0.5 * noise^2 - 0.5*log_var - 0.5*log(2*pi)\n const float log_p = -0.5f * noise * noise - 0.5f * log_var_val - 0.5f * LOG_2PI;\n local_sum += log_p;\n }\n \n // Warp-level reduction\n local_sum = warp_reduce_sum(local_sum);\n \n // First thread in each warp writes to shared memory\n __shared__ float s_partial[32]; // Max 32 warps per block\n if (lane_id == 0) {\n s_partial[warp_id] = local_sum;\n }\n __syncthreads();\n \n // Final reduction by first warp\n if (warp_id == 0) {\n float sum = (tid < warps_per_block) ? s_partial[tid] : 0.0f;\n sum = warp_reduce_sum(sum);\n if (tid == 0) {\n log_q_z[b] = sum;\n }\n }\n}\n\n// Mixture of Gaussians kernel - step 1: sample z\n__global__ void mixture_sample_kernel(\n const float* __restrict__ mu,\n const float* __restrict__ log_var,\n const int* __restrict__ components,\n float* __restrict__ z,\n const int batch_size,\n const int num_components,\n const int latent_dim,\n const uint64_t seed\n) {\n const int b = blockIdx.x;\n const int tid = threadIdx.x;\n \n if (b >= batch_size) return;\n \n const int comp = components[b];\n \n // Initialize random state\n curandState state;\n curand_init(seed, b * blockDim.x + tid, 0, &state);\n \n for (int d = tid; d < latent_dim; d += blockDim.x) {\n const int idx_mu = b * num_components * latent_dim + comp * latent_dim + d;\n const int idx_z = b * latent_dim + d;\n \n const float mu_val = mu[idx_mu];\n const float log_var_val = log_var[idx_mu];\n const float std_val = expf(0.5f * log_var_val);\n \n const float noise = curand_normal(&state);\n z[idx_z] = mu_val + std_val * noise;\n }\n}\n\n// Mixture of Gaussians kernel - step 2: compute log_q_z with logsumexp\n__global__ void mixture_logprob_kernel(\n const float* __restrict__ mu,\n const float* __restrict__ log_var,\n const float* __restrict__ z,\n const float* __restrict__ log_weights,\n float* __restrict__ log_q_z,\n const int batch_size,\n const int num_components,\n const int latent_dim\n) {\n const int b = blockIdx.x;\n const int tid = threadIdx.x;\n \n if (b >= batch_size) return;\n \n // Each thread computes log probability for one component\n // Then we do logsumexp across components\n extern __shared__ float s_log_probs[]; // size: num_components\n \n float max_log = -INFINITY;\n \n for (int k = tid; k < num_components; k += blockDim.x) {\n float sum_log_p = log_weights[k]; // Start with log(weight)\n \n for (int d = 0; d < latent_dim; d++) {\n const int idx_mu = b * num_components * latent_dim + k * latent_dim + d;\n const int idx_z = b * latent_dim + d;\n \n const float z_val = z[idx_z];\n const float mu_val = mu[idx_mu];\n const float log_var_val = log_var[idx_mu];\n const float std_val = expf(0.5f * log_var_val);\n \n // log N(z | mu, std) = -0.5 * ((z-mu)/std)^2 - 0.5*log_var - 0.5*log(2*pi)\n const float diff = (z_val - mu_val) / std_val;\n sum_log_p += -0.5f * diff * diff - 0.5f * log_var_val - 0.5f * LOG_2PI;\n }\n \n s_log_probs[k] = sum_log_p;\n max_log = fmaxf(max_log, sum_log_p);\n }\n __syncthreads();\n \n // Reduce max_log across threads\n __shared__ float s_max;\n if (tid == 0) {\n s_max = s_log_probs[0];\n for (int k = 1; k < num_components; k++) {\n s_max = fmaxf(s_max, s_log_probs[k]);\n }\n }\n __syncthreads();\n \n // Compute logsumexp\n const float max_val = s_max;\n float sum_exp = 0.0f;\n \n for (int k = tid; k < num_components; k += blockDim.x) {\n sum_exp += expf(s_log_probs[k] - max_val);\n }\n \n // Warp-level reduction\n sum_exp = warp_reduce_sum(sum_exp);\n \n __shared__ float s_sum_exp;\n if (tid == 0) {\n s_sum_exp = sum_exp;\n }\n __syncthreads();\n \n if (tid == 0) {\n log_q_z[b] = max_val + logf(s_sum_exp);\n }\n}\n\n// Host functions\nstd::tuple<torch::Tensor, torch::Tensor> single_component_forward(\n torch::Tensor mu, torch::Tensor log_var,\n int64_t latent_dim, uint64_t seed\n) {\n const int batch_size = mu.size(0);\n \n auto z = torch::empty({batch_size, latent_dim}, mu.options());\n auto log_q_z = torch::empty({batch_size}, mu.options());\n \n const int threads = 256;\n const int blocks = batch_size;\n \n single_component_2d_kernel<<<blocks, threads>>>(\n mu.data_ptr<float>(),\n log_var.data_ptr<float>(),\n z.data_ptr<float>(),\n log_q_z.data_ptr<float>(),\n batch_size,\n latent_dim,\n seed\n );\n \n return std::make_tuple(z, log_q_z);\n}\n\nstd::tuple<torch::Tensor, torch::Tensor> mixture_forward(\n torch::Tensor mu, torch::Tensor log_var,\n torch::Tensor weights,\n int64_t latent_dim, int64_t num_components,\n uint64_t seed\n) {\n const int batch_size = mu.size(0);\n \n // Sample components using multinomial (PyTorch function)\n auto weights_expanded = weights.expand({batch_size, -1});\n auto components = torch::multinomial(weights_expanded, 1).squeeze(-1);\n \n auto z = torch::empty({batch_size, latent_dim}, mu.options());\n auto log_q_z = torch::empty({batch_size}, mu.options());\n auto log_weights = torch::log(weights);\n \n const int threads = 256;\n const int blocks = batch_size;\n \n // Step 1: Sample z\n mixture_sample_kernel<<<blocks, threads>>>(\n mu.data_ptr<float>(),\n log_var.data_ptr<float>(),\n components.data_ptr<int>(),\n z.data_ptr<float>(),\n batch_size,\n num_components,\n latent_dim,\n seed\n );\n \n // Step 2: Compute log_q_z\n size_t shared_mem = num_components * sizeof(float);\n mixture_logprob_kernel<<<blocks, threads, shared_mem>>>(\n mu.data_ptr<float>(),\n log_var.data_ptr<float>(),\n z.data_ptr<float>(),\n log_weights.data_ptr<float>(),\n log_q_z.data_ptr<float>(),\n batch_size,\n num_components,\n latent_dim\n );\n \n return std::make_tuple(z, log_q_z);\n}\n\nstd::tuple<torch::Tensor, torch::Tensor> forward_3d_single(\n torch::Tensor mu, torch::Tensor log_var,\n int64_t latent_dim, uint64_t seed\n) {\n // 3D input but single component - take first component\n const int batch_size = mu.size(0);\n \n auto mu_2d = mu.select(1, 0).contiguous();\n auto log_var_2d = log_var.select(1, 0).contiguous();\n \n return single_component_forward(mu_2d, log_var_2d, latent_dim, seed);\n}\n\"\"\"\n\ncpp_source = \"\"\"\nstd::tuple<torch::Tensor, torch::Tensor> single_component_forward(\n torch::Tensor mu, torch::Tensor log_var,\n int64_t latent_dim, uint64_t seed\n);\n\nstd::tuple<torch::Tensor, torch::Tensor> mixture_forward(\n torch::Tensor mu, torch::Tensor log_var,\n torch::Tensor weights,\n int64_t latent_dim, int64_t num_components,\n uint64_t seed\n);\n\nstd::tuple<torch::Tensor, torch::Tensor> forward_3d_single(\n torch::Tensor mu, torch::Tensor log_var,\n int64_t latent_dim, uint64_t seed\n);\n\"\"\"\n\n# Compile the extension\nmodule = load_inline(\n name='reparam_cuda',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['single_component_forward', 'mixture_forward', 'forward_3d_single'],\n extra_cuda_cflags=['-O3', '--use_fast_math', '-gencode=arch=compute_80,code=sm_80'],\n verbose=False\n)\n\nimport os\nimport time\n\nclass ModelNew(nn.Module):\n \"\"\"\n Reparameterization trick for differentiable sampling from a diagonal\n Gaussian. Uses custom CUDA kernels for faster computation.\n \"\"\"\n \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 # Seed for reproducibility (can be changed)\n self._seed_counter = 0\n \n def _get_seed(self):\n # Use a combination of time and counter for different seeds\n self._seed_counter += 1\n return int(time.time() * 1e6 + self._seed_counter) % (2**63)\n \n def forward(self, mu: torch.Tensor, log_var: torch.Tensor) -> tuple:\n \"\"\"\n Args:\n mu: (B, K, D) or (B, D) - means\n log_var: (B, K, D) or (B, D) - log-variances\n Returns:\n (z, log_q_z):\n z: (B, D) - sampled latents\n log_q_z: (B,) - log-density of sample under q\n \"\"\"\n # Ensure contiguous float32 tensors\n mu = mu.contiguous().float()\n log_var = log_var.contiguous().float()\n \n seed = self._get_seed()\n \n if self.num_components > 1 and mu.dim() == 3:\n # Mixture of Gaussians case\n weights = torch.softmax(self.logits, dim=0)\n z, log_q_z = module.mixture_forward(\n mu, log_var, weights,\n self.latent_dim, self.num_components,\n seed\n )\n else:\n # Single component case\n if mu.dim() == 3:\n # Take first component\n mu = mu[:, 0].contiguous()\n log_var = log_var[:, 0].contiguous()\n \n z, log_q_z = module.single_component_forward(\n mu, log_var, self.latent_dim, seed\n )\n \n return z, log_q_z" }
result ← submit_kernel (fail)
failsubmit_kernel submit_kernel FAILED: correctness check did not pass (? trials). runtime_error: expected scalar type Int but found Long
final submitted kernel (381 lines, 11,626 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline
import math

cuda_source = r"""
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <curand_kernel.h>
#include <cmath>
#include <limits>

// Constants
static constexpr float LOG_2PI = 1.8378770664093453f;  // log(2*pi)
static constexpr float PI = 3.14159265358979323846f;

// Warp-level reduction for sum
__inline__ __device__ float warp_reduce_sum(float val) {
    for (int offset = 16; offset > 0; offset >>= 1) {
        val += __shfl_down_sync(0xffffffff, val, offset);
    }
    return val;
}

// Warp-level reduction for max
__inline__ __device__ float warp_reduce_max(float val) {
    for (int offset = 16; offset > 0; offset >>= 1) {
        val = fmaxf(val, __shfl_down_sync(0xffffffff, val, offset));
    }
    return val;
}

// Single component kernel (for 2D input)
__global__ void single_component_2d_kernel(
    const float* __restrict__ mu,
    const float* __restrict__ log_var,
    float* __restrict__ z,
    float* __restrict__ log_q_z,
    const int batch_size,
    const int latent_dim,
    const uint64_t seed
) {
    const int b = blockIdx.x;
    const int tid = threadIdx.x;
    const int lane_id = tid & 31;
    const int warp_id = tid >> 5;
    const int warps_per_block = blockDim.x >> 5;
    
    if (b >= batch_size) return;
    
    // Each thread handles multiple latent dimensions
    float local_sum = 0.0f;
    
    // Initialize random state per thread
    curandState state;
    curand_init(seed, b * blockDim.x + tid, 0, &state);
    
    for (int d = tid; d < latent_dim; d += blockDim.x) {
        const int idx = b * latent_dim + d;
        const float mu_val = mu[idx];
        const float log_var_val = log_var[idx];
        const float std_val = expf(0.5f * log_var_val);
        
        const float noise = curand_normal(&state);
        const float z_val = mu_val + std_val * noise;
        z[idx] = z_val;
        
        // log_p contribution: -0.5 * ((z-mu)/std)^2 - 0.5*log_var - 0.5*log(2*pi)
        // = -0.5 * noise^2 - 0.5*log_var - 0.5*log(2*pi)
        const float log_p = -0.5f * noise * noise - 0.5f * log_var_val - 0.5f * LOG_2PI;
        local_sum += log_p;
    }
    
    // Warp-level reduction
    local_sum = warp_reduce_sum(local_sum);
    
    // First thread in each warp writes to shared memory
    __shared__ float s_partial[32];  // Max 32 warps per block
    if (lane_id == 0) {
        s_partial[warp_id] = local_sum;
    }
    __syncthreads();
    
    // Final reduction by first warp
    if (warp_id == 0) {
        float sum = (tid < warps_per_block) ? s_partial[tid] : 0.0f;
        sum = warp_reduce_sum(sum);
        if (tid == 0) {
            log_q_z[b] = sum;
        }
    }
}

// Mixture of Gaussians kernel - step 1: sample z
__global__ void mixture_sample_kernel(
    const float* __restrict__ mu,
    const float* __restrict__ log_var,
    const int* __restrict__ components,
    float* __restrict__ z,
    const int batch_size,
    const int num_components,
    const int latent_dim,
    const uint64_t seed
) {
    const int b = blockIdx.x;
    const int tid = threadIdx.x;
    
    if (b >= batch_size) return;
    
    const int comp = components[b];
    
    // Initialize random state
    curandState state;
    curand_init(seed, b * blockDim.x + tid, 0, &state);
    
    for (int d = tid; d < latent_dim; d += blockDim.x) {
        const int idx_mu = b * num_components * latent_dim + comp * latent_dim + d;
        const int idx_z = b * latent_dim + d;
        
        const float mu_val = mu[idx_mu];
        const float log_var_val = log_var[idx_mu];
        const float std_val = expf(0.5f * log_var_val);
        
        const float noise = curand_normal(&state);
        z[idx_z] = mu_val + std_val * noise;
    }
}

// Mixture of Gaussians kernel - step 2: compute log_q_z with logsumexp
__global__ void mixture_logprob_kernel(
    const float* __restrict__ mu,
    const float* __restrict__ log_var,
    const float* __restrict__ z,
    const float* __restrict__ log_weights,
    float* __restrict__ log_q_z,
    const int batch_size,
    const int num_components,
    const int latent_dim
) {
    const int b = blockIdx.x;
    const int tid = threadIdx.x;
    
    if (b >= batch_size) return;
    
    // Each thread computes log probability for one component
    // Then we do logsumexp across components
    extern __shared__ float s_log_probs[];  // size: num_components
    
    float max_log = -INFINITY;
    
    for (int k = tid; k < num_components; k += blockDim.x) {
        float sum_log_p = log_weights[k];  // Start with log(weight)
        
        for (int d = 0; d < latent_dim; d++) {
            const int idx_mu = b * num_components * latent_dim + k * latent_dim + d;
            const int idx_z = b * latent_dim + d;
            
            const float z_val = z[idx_z];
            const float mu_val = mu[idx_mu];
            const float log_var_val = log_var[idx_mu];
            const float std_val = expf(0.5f * log_var_val);
            
            // log N(z | mu, std) = -0.5 * ((z-mu)/std)^2 - 0.5*log_var - 0.5*log(2*pi)
            const float diff = (z_val - mu_val) / std_val;
            sum_log_p += -0.5f * diff * diff - 0.5f * log_var_val - 0.5f * LOG_2PI;
        }
        
        s_log_probs[k] = sum_log_p;
        max_log = fmaxf(max_log, sum_log_p);
    }
    __syncthreads();
    
    // Reduce max_log across threads
    __shared__ float s_max;
    if (tid == 0) {
        s_max = s_log_probs[0];
        for (int k = 1; k < num_components; k++) {
            s_max = fmaxf(s_max, s_log_probs[k]);
        }
    }
    __syncthreads();
    
    // Compute logsumexp
    const float max_val = s_max;
    float sum_exp = 0.0f;
    
    for (int k = tid; k < num_components; k += blockDim.x) {
        sum_exp += expf(s_log_probs[k] - max_val);
    }
    
    // Warp-level reduction
    sum_exp = warp_reduce_sum(sum_exp);
    
    __shared__ float s_sum_exp;
    if (tid == 0) {
        s_sum_exp = sum_exp;
    }
    __syncthreads();
    
    if (tid == 0) {
        log_q_z[b] = max_val + logf(s_sum_exp);
    }
}

// Host functions
std::tuple<torch::Tensor, torch::Tensor> single_component_forward(
    torch::Tensor mu, torch::Tensor log_var,
    int64_t latent_dim, uint64_t seed
) {
    const int batch_size = mu.size(0);
    
    auto z = torch::empty({batch_size, latent_dim}, mu.options());
    auto log_q_z = torch::empty({batch_size}, mu.options());
    
    const int threads = 256;
    const int blocks = batch_size;
    
    single_component_2d_kernel<<<blocks, threads>>>(
        mu.data_ptr<float>(),
        log_var.data_ptr<float>(),
        z.data_ptr<float>(),
        log_q_z.data_ptr<float>(),
        batch_size,
        latent_dim,
        seed
    );
    
    return std::make_tuple(z, log_q_z);
}

std::tuple<torch::Tensor, torch::Tensor> mixture_forward(
    torch::Tensor mu, torch::Tensor log_var,
    torch::Tensor weights,
    int64_t latent_dim, int64_t num_components,
    uint64_t seed
) {
    const int batch_size = mu.size(0);
    
    // Sample components using multinomial (PyTorch function)
    auto weights_expanded = weights.expand({batch_size, -1});
    auto components = torch::multinomial(weights_expanded, 1).squeeze(-1);
    
    auto z = torch::empty({batch_size, latent_dim}, mu.options());
    auto log_q_z = torch::empty({batch_size}, mu.options());
    auto log_weights = torch::log(weights);
    
    const int threads = 256;
    const int blocks = batch_size;
    
    // Step 1: Sample z
    mixture_sample_kernel<<<blocks, threads>>>(
        mu.data_ptr<float>(),
        log_var.data_ptr<float>(),
        components.data_ptr<int>(),
        z.data_ptr<float>(),
        batch_size,
        num_components,
        latent_dim,
        seed
    );
    
    // Step 2: Compute log_q_z
    size_t shared_mem = num_components * sizeof(float);
    mixture_logprob_kernel<<<blocks, threads, shared_mem>>>(
        mu.data_ptr<float>(),
        log_var.data_ptr<float>(),
        z.data_ptr<float>(),
        log_weights.data_ptr<float>(),
        log_q_z.data_ptr<float>(),
        batch_size,
        num_components,
        latent_dim
    );
    
    return std::make_tuple(z, log_q_z);
}

std::tuple<torch::Tensor, torch::Tensor> forward_3d_single(
    torch::Tensor mu, torch::Tensor log_var,
    int64_t latent_dim, uint64_t seed
) {
    // 3D input but single component - take first component
    const int batch_size = mu.size(0);
    
    auto mu_2d = mu.select(1, 0).contiguous();
    auto log_var_2d = log_var.select(1, 0).contiguous();
    
    return single_component_forward(mu_2d, log_var_2d, latent_dim, seed);
}
"""

cpp_source = """
std::tuple<torch::Tensor, torch::Tensor> single_component_forward(
    torch::Tensor mu, torch::Tensor log_var,
    int64_t latent_dim, uint64_t seed
);

std::tuple<torch::Tensor, torch::Tensor> mixture_forward(
    torch::Tensor mu, torch::Tensor log_var,
    torch::Tensor weights,
    int64_t latent_dim, int64_t num_components,
    uint64_t seed
);

std::tuple<torch::Tensor, torch::Tensor> forward_3d_single(
    torch::Tensor mu, torch::Tensor log_var,
    int64_t latent_dim, uint64_t seed
);
"""

# Compile the extension
module = load_inline(
    name='reparam_cuda',
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=['single_component_forward', 'mixture_forward', 'forward_3d_single'],
    extra_cuda_cflags=['-O3', '--use_fast_math', '-gencode=arch=compute_80,code=sm_80'],
    verbose=False
)

import os
import time

class ModelNew(nn.Module):
    """
    Reparameterization trick for differentiable sampling from a diagonal
    Gaussian. Uses custom CUDA kernels for faster computation.
    """
    
    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))
        
        # Seed for reproducibility (can be changed)
        self._seed_counter = 0
    
    def _get_seed(self):
        # Use a combination of time and counter for different seeds
        self._seed_counter += 1
        return int(time.time() * 1e6 + self._seed_counter) % (2**63)
    
    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
        """
        # Ensure contiguous float32 tensors
        mu = mu.contiguous().float()
        log_var = log_var.contiguous().float()
        
        seed = self._get_seed()
        
        if self.num_components > 1 and mu.dim() == 3:
            # Mixture of Gaussians case
            weights = torch.softmax(self.logits, dim=0)
            z, log_q_z = module.mixture_forward(
                mu, log_var, weights,
                self.latent_dim, self.num_components,
                seed
            )
        else:
            # Single component case
            if mu.dim() == 3:
                # Take first component
                mu = mu[:, 0].contiguous()
                log_var = log_var[:, 0].contiguous()
            
            z, log_q_z = module.single_component_forward(
                mu, log_var, self.latent_dim, seed
            )
        
        return z, log_q_z