call → run_correctness
callrun_correctness
{
"kernel_code": "import torch\nimport torch.nn as nn\nimport math\nfrom torch.utils.cpp_extension import load_inline\ncpp_src = r\"\"\"\n#include <torch/extension.h>\n#include <vector>\nstd::vector<torch::Tensor> reparam_forward(torch::Tensor mu, torch::Tensor log_var, torch::Tensor comp, torch::Tensor eps, torch::Tensor logw);\n\"\"\"\ncuda_src = r\"\"\"\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda_runtime.h>\n#include <vector>\n#include <math.h>\n#define LOG_SQRT_2PI 0.9189385175704956f\n__inline__ __device__ float warp_sum(float v){\n v += __shfl_down_sync(0xffffffff, v, 16);\n v += __shfl_down_sync(0xffffffff, v, 8);\n v += __shfl_down_sync(0xffffffff, v, 4);\n v += __shfl_down_sync(0xffffffff, v, 2);\n v += __shfl_down_sync(0xffffffff, v, 1);\n return v;\n}\n__global__ void mix4_kernel(const float* __restrict__ mu,const float* __restrict__ lv,const long long* __restrict__ comp,const float* __restrict__ eps,const float* __restrict__ logw,float* __restrict__ z,float* __restrict__ logq){\n int b=blockIdx.x, t=threadIdx.x, c=(int)comp[b]; int base=b*256+t; int sel=b*256+c*64+t;\n float lvc=lv[sel]; float zv=mu[sel]+expf(0.5f*lvc)*eps[b*64+t]; z[b*64+t]=zv;\n float m0=mu[base], l0=lv[base], d0=zv-m0; float a0=-0.5f*d0*d0*expf(-l0)-0.5f*l0-LOG_SQRT_2PI;\n float m1=mu[base+64], l1=lv[base+64], d1=zv-m1; float a1=-0.5f*d1*d1*expf(-l1)-0.5f*l1-LOG_SQRT_2PI;\n float m2=mu[base+128], l2=lv[base+128], d2=zv-m2; float a2=-0.5f*d2*d2*expf(-l2)-0.5f*l2-LOG_SQRT_2PI;\n float m3=mu[base+192], l3=lv[base+192], d3=zv-m3; float a3=-0.5f*d3*d3*expf(-l3)-0.5f*l3-LOG_SQRT_2PI;\n a0=warp_sum(a0); a1=warp_sum(a1); a2=warp_sum(a2); a3=warp_sum(a3);\n __shared__ float part[8]; int lane=t&31, w=t>>5;\n if(lane==0){part[w*4]=a0; part[w*4+1]=a1; part[w*4+2]=a2; part[w*4+3]=a3;}\n __syncthreads();\n if(t==0){\n float v0=part[0]+part[4]+logw[0]; float v1=part[1]+part[5]+logw[1]; float v2=part[2]+part[6]+logw[2]; float v3=part[3]+part[7]+logw[3];\n float mx=fmaxf(fmaxf(v0,v1),fmaxf(v2,v3));\n logq[b]=mx+logf(expf(v0-mx)+expf(v1-mx)+expf(v2-mx)+expf(v3-mx));\n }\n}\n__global__ void generic_kernel(const float* __restrict__ mu,const float* __restrict__ lv,const long long* __restrict__ comp,const float* __restrict__ eps,const float* __restrict__ logw,float* __restrict__ z,float* __restrict__ logq,int K,int D){\n int b=blockIdx.x,t=threadIdx.x,c=(int)comp[b]; float acc[8]; #pragma unroll for(int k=0;k<8;k++) acc[k]=0.f;\n for(int d=t;d<D;d+=blockDim.x){int s=(b*K+c)*D+d; float lvc=lv[s]; float zv=mu[s]+expf(.5f*lvc)*eps[b*D+d]; z[b*D+d]=zv; for(int k=0;k<K;k++){int i=(b*K+k)*D+d; float diff=zv-mu[i]; float l=lv[i]; acc[k]+=-.5f*diff*diff*expf(-l)-.5f*l-LOG_SQRT_2PI;}}\n extern __shared__ float sm[]; for(int k=0;k<K;k++) sm[k*blockDim.x+t]=acc[k]; __syncthreads(); for(int s=blockDim.x/2;s;s/=2){if(t<s){for(int k=0;k<K;k++) sm[k*blockDim.x+t]+=sm[k*blockDim.x+t+s];} __syncthreads();} if(t==0){float mx=-3.4e38f,v[8]; for(int k=0;k<K;k++){v[k]=sm[k*blockDim.x]+logw[k]; mx=fmaxf(mx,v[k]);} float se=0; for(int k=0;k<K;k++) se+=expf(v[k]-mx); logq[b]=mx+logf(se);} }\nstd::vector<torch::Tensor> reparam_forward(torch::Tensor mu,torch::Tensor log_var,torch::Tensor comp,torch::Tensor eps,torch::Tensor logw){int B=mu.size(0),D=eps.size(1),K=mu.dim()==3?mu.size(1):1; auto z=torch::empty({B,D},mu.options()); auto q=torch::empty({B},mu.options()); auto st=at::cuda::getCurrentCUDAStream(); if(K==4 && D==64){mix4_kernel<<<B,64,0,st>>>(mu.data_ptr<float>(),log_var.data_ptr<float>(),(const long long*)comp.data_ptr<int64_t>(),eps.data_ptr<float>(),logw.data_ptr<float>(),z.data_ptr<float>(),q.data_ptr<float>());} else {generic_kernel<<<B,64,(size_t)K*64*sizeof(float),st>>>(mu.data_ptr<float>(),log_var.data_ptr<float>(),(const long long*)comp.data_ptr<int64_t>(),eps.data_ptr<float>(),logw.data_ptr<float>(),z.data_ptr<float>(),q.data_ptr<float>(),K,D);} return {z,q};}\n\"\"\"\n_ext=load_inline(name='reparam_ext_v4warp',cpp_sources=cpp_src,cuda_sources=cuda_src,functions=['reparam_forward'],extra_cuda_cflags=['-O3','--use_fast_math'],extra_cflags=['-O3'],verbose=False)\nclass ModelNew(nn.Module):\n def __init__(self, latent_dim, num_components=1):\n super().__init__(); self.latent_dim=latent_dim; self.num_components=num_components\n if num_components>1:\n self.logits=nn.Parameter(torch.zeros(num_components)); self.register_buffer('_weights',torch.full((num_components,),1.0/float(num_components))); self.register_buffer('_logw',torch.full((num_components,),-math.log(float(num_components))))\n def forward(self, mu, log_var):\n B=mu.shape[0]\n if self.num_components>1 and mu.dim()==3:\n comp=torch.multinomial(self._weights.expand(B,-1),1).squeeze(-1); eps=torch.randn((B,mu.shape[2]),device=mu.device,dtype=mu.dtype); return tuple(_ext.reparam_forward(mu,log_var,comp,eps,self._logw))\n else:\n D=mu.shape[-1]; comp=torch.zeros((B,),device=mu.device,dtype=torch.long); eps=torch.randn((B,D),device=mu.device,dtype=mu.dtype); lw=torch.zeros((1,),device=mu.device,dtype=mu.dtype); return tuple(_ext.reparam_forward(mu,log_var,comp,eps,lw))\n"
}