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