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"
}