call → run_correctness
callrun_correctness
{
"kernel_code": "import torch\nimport torch.nn as nn\nimport torch.nn.functional as F\nfrom torch.utils.cpp_extension import load_inline\n\n_cpp_src = r\"\"\"\n#include <torch/extension.h>\ntorch::Tensor bn_relu_forward(torch::Tensor x, torch::Tensor weight, torch::Tensor bias, torch::Tensor running_mean, torch::Tensor running_var, double eps);\ntorch::Tensor final_bn_relu_avg_forward(torch::Tensor x, torch::Tensor weight, torch::Tensor bias, torch::Tensor running_mean, torch::Tensor running_var, double eps);\ntorch::Tensor avgpool2x2_forward(torch::Tensor x);\ntorch::Tensor relu_avg_forward(torch::Tensor x);\n\"\"\"\n_cuda_src = r\"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <math.h>\n__global__ void bn_relu_kernel(const float* __restrict__ x, float* __restrict__ y,const float* __restrict__ w,const float* __restrict__ b,const float* __restrict__ mean,const float* __restrict__ var,int total,int C,int HW,float eps){int idx=blockIdx.x*blockDim.x+threadIdx.x; if(idx<total){int c=(idx/HW)%C; float v=(x[idx]-mean[c])*rsqrtf(var[c]+eps)*w[c]+b[c]; y[idx]=v>0.0f?v:0.0f;}}\n__global__ void avgpool2x2_kernel(const float* __restrict__ x,float* __restrict__ y,int N,int C,int H,int W,int OH,int OW){int idx=blockIdx.x*blockDim.x+threadIdx.x; int total=N*C*OH*OW; if(idx<total){int ow=idx%OW; int t=idx/OW; int oh=t%OH; t/=OH; int c=t%C; int n=t/C; int base=((n*C+c)*H+(oh<<1))*W+(ow<<1); float s=x[base]+x[base+1]+x[base+W]+x[base+W+1]; y[idx]=s*0.25f;}}\n__global__ void final_bn_relu_avg_kernel(const float* __restrict__ x,float* __restrict__ y,const float* __restrict__ w,const float* __restrict__ b,const float* __restrict__ mean,const float* __restrict__ var,int N,int C,int HW,float eps){int nc=blockIdx.x; int c=nc%C; int n=nc/C; const float* xp=x+(n*C+c)*HW; float scale=rsqrtf(var[c]+eps)*w[c]; float shift=b[c]-mean[c]*scale; float sum=0.0f; for(int i=threadIdx.x;i<HW;i+=blockDim.x){float v=xp[i]*scale+shift; sum+=v>0.0f?v:0.0f;} __shared__ float smem[256]; smem[threadIdx.x]=sum; __syncthreads(); for(int stride=blockDim.x>>1;stride>0;stride>>=1){if(threadIdx.x<stride) smem[threadIdx.x]+=smem[threadIdx.x+stride]; __syncthreads();} if(threadIdx.x==0) y[n*C+c]=smem[0]/(float)HW;}\n__global__ void relu_avg_kernel(const float* __restrict__ x,float* __restrict__ y,int N,int C,int HW){int nc=blockIdx.x; int c=nc%C; int n=nc/C; const float* xp=x+(n*C+c)*HW; float sum=0.0f; for(int i=threadIdx.x;i<HW;i+=blockDim.x){float v=xp[i]; sum+=v>0.0f?v:0.0f;} __shared__ float smem[256]; smem[threadIdx.x]=sum; __syncthreads(); for(int stride=blockDim.x>>1;stride>0;stride>>=1){if(threadIdx.x<stride) smem[threadIdx.x]+=smem[threadIdx.x+stride]; __syncthreads();} if(threadIdx.x==0) y[n*C+c]=smem[0]/(float)HW;}\n\ntorch::Tensor bn_relu_forward(torch::Tensor x,torch::Tensor weight,torch::Tensor bias,torch::Tensor running_mean,torch::Tensor running_var,double eps){x=x.contiguous(); auto y=torch::empty_like(x); int64_t N=x.size(0),C=x.size(1),H=x.size(2),W=x.size(3); int total=(int)(N*C*H*W), HW=(int)(H*W); int threads=256, blocks=(total+threads-1)/threads; bn_relu_kernel<<<blocks,threads>>>(x.data_ptr<float>(),y.data_ptr<float>(),weight.data_ptr<float>(),bias.data_ptr<float>(),running_mean.data_ptr<float>(),running_var.data_ptr<float>(),total,(int)C,HW,(float)eps); return y;}\ntorch::Tensor avgpool2x2_forward(torch::Tensor x){x=x.contiguous(); int64_t N=x.size(0),C=x.size(1),H=x.size(2),W=x.size(3),OH=H/2,OW=W/2; auto y=torch::empty({N,C,OH,OW},x.options()); int total=(int)(N*C*OH*OW); int threads=256, blocks=(total+threads-1)/threads; avgpool2x2_kernel<<<blocks,threads>>>(x.data_ptr<float>(),y.data_ptr<float>(),(int)N,(int)C,(int)H,(int)W,(int)OH,(int)OW); return y;}\ntorch::Tensor final_bn_relu_avg_forward(torch::Tensor x,torch::Tensor weight,torch::Tensor bias,torch::Tensor running_mean,torch::Tensor running_var,double eps){x=x.contiguous(); int64_t N=x.size(0),C=x.size(1),H=x.size(2),W=x.size(3); auto y=torch::empty({N,C},x.options()); final_bn_relu_avg_kernel<<<(int)(N*C),256>>>(x.data_ptr<float>(),y.data_ptr<float>(),weight.data_ptr<float>(),bias.data_ptr<float>(),running_mean.data_ptr<float>(),running_var.data_ptr<float>(),(int)N,(int)C,(int)(H*W),(float)eps); return y;}\ntorch::Tensor relu_avg_forward(torch::Tensor x){x=x.contiguous(); int64_t N=x.size(0),C=x.size(1),H=x.size(2),W=x.size(3); auto y=torch::empty({N,C},x.options()); relu_avg_kernel<<<(int)(N*C),256>>>(x.data_ptr<float>(),y.data_ptr<float>(),(int)N,(int)C,(int)(H*W)); return y;}\n\"\"\"\n_ext=load_inline(name=\"densenet_fused_ext_v2\", cpp_sources=_cpp_src, cuda_sources=_cuda_src, functions=[\"bn_relu_forward\",\"final_bn_relu_avg_forward\",\"avgpool2x2_forward\",\"relu_avg_forward\"], extra_cflags=[\"-O3\"], extra_cuda_cflags=[\"-O3\",\"--use_fast_math\"], verbose=False)\n\nclass DenseBlock(nn.Module):\n def __init__(self,num_layers:int,num_input_features:int,growth_rate:int):\n super().__init__(); self.layers=nn.ModuleList([self._make_layer(num_input_features+i*growth_rate,growth_rate) for i in range(num_layers)])\n def _make_layer(self,in_features:int,growth_rate:int):\n return nn.Sequential(nn.BatchNorm2d(in_features), nn.ReLU(inplace=True), nn.Conv2d(in_features,growth_rate,kernel_size=3,padding=1,bias=False), nn.Dropout(0.0))\n def forward(self,x):\n features=[x]\n for layer in self.layers:\n if x.is_cuda and x.dtype is torch.float32 and (not self.training):\n y=_ext.bn_relu_forward(x,layer[0].weight,layer[0].bias,layer[0].running_mean,layer[0].running_var,layer[0].eps)\n else:\n y=layer[0](x); y=F.relu(y,inplace=True)\n new_feature=layer[2](y)\n features.append(new_feature); x=torch.cat(features,1)\n return x\nclass TransitionLayer(nn.Module):\n def __init__(self,num_input_features:int,num_output_features:int):\n super().__init__(); self.transition=nn.Sequential(nn.BatchNorm2d(num_input_features), nn.ReLU(inplace=True), nn.Conv2d(num_input_features,num_output_features,kernel_size=1,bias=False), nn.AvgPool2d(kernel_size=2,stride=2))\n def forward(self,x):\n t=self.transition\n if x.is_cuda and x.dtype is torch.float32 and (not self.training):\n y=_ext.bn_relu_forward(x,t[0].weight,t[0].bias,t[0].running_mean,t[0].running_var,t[0].eps)\n y=_ext.avgpool2x2_forward(y); return t[2](y)\n # training: keep exact BN, but move avgpool before linear 1x1 conv\n y=t[0](x); y=F.relu(y,inplace=True)\n if y.is_cuda and y.dtype is torch.float32:\n y=_ext.avgpool2x2_forward(y)\n else:\n y=F.avg_pool2d(y,2,2)\n return t[2](y)\nclass ModelNew(nn.Module):\n def __init__(self,growth_rate:int=32,num_classes:int=1000):\n super().__init__(); self.features=nn.Sequential(nn.Conv2d(3,64,kernel_size=7,stride=2,padding=3,bias=False), nn.BatchNorm2d(64), nn.ReLU(inplace=True), nn.MaxPool2d(kernel_size=3,stride=2,padding=1)); num_features=64; block_layers=[6,12,24,16]; self.dense_blocks=nn.ModuleList(); self.transition_layers=nn.ModuleList()\n for i,num_layers in enumerate(block_layers):\n block=DenseBlock(num_layers,num_features,growth_rate); self.dense_blocks.append(block); num_features += num_layers*growth_rate\n if i != len(block_layers)-1:\n self.transition_layers.append(TransitionLayer(num_features,num_features//2)); num_features//=2\n self.final_bn=nn.BatchNorm2d(num_features); self.classifier=nn.Linear(num_features,num_classes)\n def forward(self,x:torch.Tensor)->torch.Tensor:\n if x.is_cuda and x.dtype is torch.float32 and (not self.training):\n x=self.features[0](x); x=_ext.bn_relu_forward(x,self.features[1].weight,self.features[1].bias,self.features[1].running_mean,self.features[1].running_var,self.features[1].eps); x=self.features[3](x)\n else:\n x=self.features[0](x); x=self.features[1](x); x=F.relu(x,inplace=True); x=self.features[3](x)\n for i,block in enumerate(self.dense_blocks):\n x=block(x)\n if i != len(self.dense_blocks)-1: x=self.transition_layers[i](x)\n if x.is_cuda and x.dtype is torch.float32 and (not self.training):\n x=_ext.final_bn_relu_avg_forward(x,self.final_bn.weight,self.final_bn.bias,self.final_bn.running_mean,self.final_bn.running_var,self.final_bn.eps)\n else:\n x=self.final_bn(x)\n if x.is_cuda and x.dtype is torch.float32: x=_ext.relu_avg_forward(x)\n else: x=F.adaptive_avg_pool2d(F.relu(x,inplace=True),(1,1)).view(x.size(0),-1)\n return self.classifier(x)\n"
}