call → compile_kernel
callcompile_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\ncpp_src = r'''\n#include <torch/extension.h>\n#include <vector>\n\nstd::vector<torch::Tensor> mu_logvar_forward(torch::Tensor h, torch::Tensor wmu, torch::Tensor bmu, torch::Tensor wlv, torch::Tensor blv);\ntorch::Tensor z_kl_forward(torch::Tensor mu, torch::Tensor logvar, torch::Tensor eps, torch::Tensor z);\ntorch::Tensor latent_hidden_relu_forward(torch::Tensor z, torch::Tensor w, torch::Tensor bias);\nstd::vector<torch::Tensor> recon_elbo_forward(torch::Tensor logits, torch::Tensor x, torch::Tensor kl);\n'''\n\ncuda_src = r'''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <vector>\n\n__global__ void mu_logvar_kernel(const float* __restrict__ h,\n const float* __restrict__ wmu, const float* __restrict__ bmu,\n const float* __restrict__ wlv, const float* __restrict__ blv,\n float* __restrict__ mu, float* __restrict__ lv,\n int B, int H, int L) {\n int b = blockIdx.x;\n int t = threadIdx.x;\n if (t < 64) {\n int j = t & 31;\n bool is_lv = (t >= 32);\n const float* w = is_lv ? wlv : wmu;\n const float* bias = is_lv ? blv : bmu;\n float sum = bias[j];\n const float* hp = h + b * H;\n const float* wp = w + j * H;\n #pragma unroll 4\n for (int k = 0; k < 256; ++k) sum = fmaf(hp[k], wp[k], sum);\n if (is_lv) lv[b * L + j] = sum; else mu[b * L + j] = sum;\n }\n}\n\n__global__ void z_kl_kernel(const float* __restrict__ mu, const float* __restrict__ logvar,\n const float* __restrict__ eps, float* __restrict__ z,\n float* __restrict__ kl, int L) {\n int b = blockIdx.x;\n int tid = threadIdx.x;\n float sum = 0.0f;\n if (tid < 32) {\n int idx = b * L + tid;\n float m = mu[idx];\n float l = logvar[idx];\n float e = expf(l);\n sum = 0.5f * (m * m + e - 1.0f - l);\n z[idx] = m + expf(0.5f * l) * eps[idx];\n }\n __shared__ float smem[32];\n if (tid < 32) smem[tid] = sum;\n __syncthreads();\n if (tid < 16) smem[tid] += smem[tid + 16];\n __syncthreads();\n if (tid < 8) smem[tid] += smem[tid + 8];\n __syncthreads();\n if (tid < 4) smem[tid] += smem[tid + 4];\n __syncthreads();\n if (tid < 2) smem[tid] += smem[tid + 2];\n __syncthreads();\n if (tid == 0) kl[b] = smem[0] + smem[1];\n}\n\n__global__ void latent_hidden_relu_kernel(const float* __restrict__ z, const float* __restrict__ w,\n const float* __restrict__ bias, float* __restrict__ out,\n int L, int H) {\n int b = blockIdx.x;\n int j = threadIdx.x;\n if (j < H) {\n float sum = bias[j];\n const float* zp = z + b * L;\n const float* wp = w + j * L;\n #pragma unroll\n for (int k = 0; k < 32; ++k) sum = fmaf(zp[k], wp[k], sum);\n out[b * H + j] = sum > 0.0f ? sum : 0.0f;\n }\n}\n\n__global__ void recon_elbo_kernel(const float* __restrict__ logits, const float* __restrict__ x,\n const float* __restrict__ kl, float* __restrict__ recon,\n float* __restrict__ elbo, int D) {\n int b = blockIdx.x;\n int tid = threadIdx.x;\n float sum = 0.0f;\n const int base = b * D;\n for (int j = tid; j < D; j += 256) {\n float y = logits[base + j];\n float xv = x[base + j];\n sum += fmaxf(y, 0.0f) - xv * y + log1pf(expf(-fabsf(y)));\n }\n __shared__ float smem[256];\n smem[tid] = sum;\n __syncthreads();\n for (int stride = 128; stride > 0; stride >>= 1) {\n if (tid < stride) smem[tid] += smem[tid + stride];\n __syncthreads();\n }\n if (tid == 0) {\n float r = smem[0];\n recon[b] = r;\n elbo[b] = -(r + kl[b]);\n }\n}\n\nstd::vector<torch::Tensor> mu_logvar_forward(torch::Tensor h, torch::Tensor wmu, torch::Tensor bmu, torch::Tensor wlv, torch::Tensor blv) {\n int B = h.size(0); int L = wmu.size(0); int H = h.size(1);\n auto mu = torch::empty({B, L}, h.options());\n auto lv = torch::empty({B, L}, h.options());\n mu_logvar_kernel<<<B, 64>>>(h.data_ptr<float>(), wmu.data_ptr<float>(), bmu.data_ptr<float>(),\n wlv.data_ptr<float>(), blv.data_ptr<float>(), mu.data_ptr<float>(), lv.data_ptr<float>(), B, H, L);\n return {mu, lv};\n}\n\ntorch::Tensor z_kl_forward(torch::Tensor mu, torch::Tensor logvar, torch::Tensor eps, torch::Tensor z) {\n int B = mu.size(0); int L = mu.size(1);\n auto kl = torch::empty({B}, mu.options());\n z_kl_kernel<<<B, 32>>>(mu.data_ptr<float>(), logvar.data_ptr<float>(), eps.data_ptr<float>(), z.data_ptr<float>(), kl.data_ptr<float>(), L);\n return kl;\n}\n\ntorch::Tensor latent_hidden_relu_forward(torch::Tensor z, torch::Tensor w, torch::Tensor bias) {\n int B = z.size(0); int H = w.size(0); int L = z.size(1);\n auto out = torch::empty({B, H}, z.options());\n latent_hidden_relu_kernel<<<B, 256>>>(z.data_ptr<float>(), w.data_ptr<float>(), bias.data_ptr<float>(), out.data_ptr<float>(), L, H);\n return out;\n}\n\nstd::vector<torch::Tensor> recon_elbo_forward(torch::Tensor logits, torch::Tensor x, torch::Tensor kl) {\n int B = logits.size(0); int D = logits.size(1);\n auto recon = torch::empty({B}, logits.options());\n auto elbo = torch::empty({B}, logits.options());\n recon_elbo_kernel<<<B, 256>>>(logits.data_ptr<float>(), x.data_ptr<float>(), kl.data_ptr<float>(), recon.data_ptr<float>(), elbo.data_ptr<float>(), D);\n return {elbo, recon};\n}\n'''\n\n_ext = load_inline(name='vae_elbo_ext_v2', cpp_sources=cpp_src, cuda_sources=cuda_src,\n functions=['mu_logvar_forward','z_kl_forward','latent_hidden_relu_forward','recon_elbo_forward'], verbose=False,\n extra_cuda_cflags=['-O3', '--use_fast_math'], extra_cflags=['-O3'])\n\nclass ModelNew(nn.Module):\n def __init__(self, input_dim, hidden_dim, latent_dim):\n super().__init__()\n self.encoder = nn.Sequential(nn.Linear(input_dim, hidden_dim), nn.ReLU())\n self.fc_mu = nn.Linear(hidden_dim, latent_dim)\n self.fc_logvar = nn.Linear(hidden_dim, latent_dim)\n self.decoder = nn.Sequential(nn.Linear(latent_dim, hidden_dim), nn.ReLU(), nn.Linear(hidden_dim, input_dim), nn.Sigmoid())\n\n def forward(self, x: torch.Tensor) -> tuple:\n h = F.linear(x, self.encoder[0].weight, self.encoder[0].bias)\n h = torch.relu(h)\n mu, logvar = _ext.mu_logvar_forward(h, self.fc_mu.weight, self.fc_mu.bias, self.fc_logvar.weight, self.fc_logvar.bias)\n eps = torch.randn_like(logvar)\n z = torch.empty_like(mu)\n kl = _ext.z_kl_forward(mu, logvar, eps, z)\n dh = _ext.latent_hidden_relu_forward(z, self.decoder[0].weight, self.decoder[0].bias)\n logits = F.linear(dh, self.decoder[2].weight, self.decoder[2].bias)\n elbo, recon = _ext.recon_elbo_forward(logits, x, kl)\n return elbo, recon, kl\n"
}