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 affine_relu_cuda(torch::Tensor a, torch::Tensor s, torch::Tensor b);\ntorch::Tensor affine_cuda(torch::Tensor a, torch::Tensor s, torch::Tensor b);\ntorch::Tensor affine_add_relu_cuda(torch::Tensor a, torch::Tensor s, torch::Tensor b, torch::Tensor id);\ntorch::Tensor affine_relu(torch::Tensor a, torch::Tensor s, torch::Tensor b) { return affine_relu_cuda(a,s,b); }\ntorch::Tensor affine(torch::Tensor a, torch::Tensor s, torch::Tensor b) { return affine_cuda(a,s,b); }\ntorch::Tensor affine_add_relu(torch::Tensor a, torch::Tensor s, torch::Tensor b, torch::Tensor id) { return affine_add_relu_cuda(a,s,b,id); }\n'''\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n__global__ void ar(const float* __restrict__ a,const float* __restrict__ s,const float* __restrict__ b,float* __restrict__ o,long n,int HW){long idx=blockIdx.x*blockDim.x+threadIdx.x;long st=(long)blockDim.x*gridDim.x;for(long i=idx;i<n;i+=st){int c=(int)((i/HW)%gridDim.y); float v=a[i]*s[c]+b[c]; o[i]=v>0?v:0;}}\n__global__ void af(const float* __restrict__ a,const float* __restrict__ s,const float* __restrict__ b,float* __restrict__ o,long n,int HW){long idx=blockIdx.x*blockDim.x+threadIdx.x;long st=(long)blockDim.x*gridDim.x;for(long i=idx;i<n;i+=st){int c=(int)((i/HW)%gridDim.y); o[i]=a[i]*s[c]+b[c];}}\n__global__ void aar(const float* __restrict__ a,const float* __restrict__ s,const float* __restrict__ b,const float* __restrict__ id,float* __restrict__ o,long n,int HW){long idx=blockIdx.x*blockDim.x+threadIdx.x;long st=(long)blockDim.x*gridDim.x;for(long i=idx;i<n;i+=st){int c=(int)((i/HW)%gridDim.y); float v=a[i]*s[c]+b[c]+id[i]; o[i]=v>0?v:0;}}\nstatic dim3 grid_for(long n,int C){long bx=(n+255)/256; if(bx>4096) bx=4096; return dim3((unsigned)bx,(unsigned)C,1);}\ntorch::Tensor affine_relu_cuda(torch::Tensor a, torch::Tensor s, torch::Tensor b){auto o=torch::empty_like(a); long n=a.numel(); int HW=(int)(a.size(2)*a.size(3)); ar<<<grid_for(n,(int)a.size(1)),256,0,at::cuda::getCurrentCUDAStream()>>>(a.data_ptr<float>(),s.data_ptr<float>(),b.data_ptr<float>(),o.data_ptr<float>(),n,HW); return o;}\ntorch::Tensor affine_cuda(torch::Tensor a, torch::Tensor s, torch::Tensor b){auto o=torch::empty_like(a); long n=a.numel(); int HW=(int)(a.size(2)*a.size(3)); af<<<grid_for(n,(int)a.size(1)),256,0,at::cuda::getCurrentCUDAStream()>>>(a.data_ptr<float>(),s.data_ptr<float>(),b.data_ptr<float>(),o.data_ptr<float>(),n,HW); return o;}\ntorch::Tensor affine_add_relu_cuda(torch::Tensor a, torch::Tensor s, torch::Tensor b, torch::Tensor id){auto o=torch::empty_like(a); long n=a.numel(); int HW=(int)(a.size(2)*a.size(3)); aar<<<grid_for(n,(int)a.size(1)),256,0,at::cuda::getCurrentCUDAStream()>>>(a.data_ptr<float>(),s.data_ptr<float>(),b.data_ptr<float>(),id.data_ptr<float>(),o.data_ptr<float>(),n,HW); return o;}\n'''\n_ext = load_inline(name='resnet18_fused_ext_v4', cpp_sources=[_cpp_src], cuda_sources=[_cuda_src], functions=['affine_relu','affine','affine_add_relu'], verbose=False, extra_cuda_cflags=['-O3'], extra_cflags=['-O3'])\nclass BasicBlockNew(nn.Module):\n expansion=1\n def __init__(self,in_channels,out_channels,stride=1,downsample=None):\n super().__init__(); self.conv1=nn.Conv2d(in_channels,out_channels,3,stride,1,bias=False); self.bn1=nn.BatchNorm2d(out_channels); self.relu=nn.ReLU(inplace=True); self.conv2=nn.Conv2d(out_channels,out_channels,3,1,1,bias=False); self.bn2=nn.BatchNorm2d(out_channels); self.downsample=downsample; self.stride=stride\nclass ModelNew(nn.Module):\n def __init__(self,num_classes=1000):\n super().__init__(); self.in_channels=64; self.conv1=nn.Conv2d(3,64,7,2,3,bias=False); self.bn1=nn.BatchNorm2d(64); self.relu=nn.ReLU(inplace=True); self.maxpool=nn.MaxPool2d(3,2,1); self.layer1=self._make_layer(BasicBlockNew,64,2,1); self.layer2=self._make_layer(BasicBlockNew,128,2,2); self.layer3=self._make_layer(BasicBlockNew,256,2,2); self.layer4=self._make_layer(BasicBlockNew,512,2,2); self.avgpool=nn.AdaptiveAvgPool2d((1,1)); self.fc=nn.Linear(512,num_classes); self._params=None\n def _make_layer(self,block,out_channels,blocks,stride=1):\n downsample=None\n if stride!=1 or self.in_channels!=out_channels:\n downsample=nn.Sequential(nn.Conv2d(self.in_channels,out_channels,1,stride,bias=False),nn.BatchNorm2d(out_channels))\n layers=[block(self.in_channels,out_channels,stride,downsample)]; self.in_channels=out_channels\n for _ in range(1,blocks): layers.append(block(self.in_channels,out_channels))\n return nn.Sequential(*layers)\n def _bn_params(self,bn):\n s=(bn.weight*torch.rsqrt(bn.running_var+bn.eps)).contiguous(); b=(bn.bias-bn.running_mean*s).contiguous(); return s,b\n def _build_params(self):\n p={}\n with torch.no_grad():\n p['bn1']=self._bn_params(self.bn1)\n for lname in ['layer1','layer2','layer3','layer4']:\n for i,blk in enumerate(getattr(self,lname)):\n k=f'{lname}.{i}'; p[k+'.bn1']=self._bn_params(blk.bn1); p[k+'.bn2']=self._bn_params(blk.bn2)\n if blk.downsample is not None: p[k+'.down']=self._bn_params(blk.downsample[1])\n self._params=p\n def _forward_ref(self,x):\n x=self.conv1(x); x=self.bn1(x); x=self.relu(x); x=self.maxpool(x)\n for layer in [self.layer1,self.layer2,self.layer3,self.layer4]:\n for b in layer:\n identity=x; out=b.conv1(x); out=b.bn1(out); out=self.relu(out); out=b.conv2(out); out=b.bn2(out)\n if b.downsample is not None: identity=b.downsample(x)\n x=self.relu(out+identity)\n return self.fc(torch.flatten(self.avgpool(x),1))\n def _block(self,x,blk,key):\n identity=x; out=F.conv2d(x,blk.conv1.weight,None,blk.stride,1); s,b=self._params[key+'.bn1']; out=_ext.affine_relu(out,s,b); out=F.conv2d(out,blk.conv2.weight,None,1,1)\n if blk.downsample is not None:\n identity=F.conv2d(x,blk.downsample[0].weight,None,blk.stride,0); s,b=self._params[key+'.down']; identity=_ext.affine(identity,s,b)\n s,b=self._params[key+'.bn2']; return _ext.affine_add_relu(out,s,b,identity)\n def forward(self,x):\n if self.training: return self._forward_ref(x)\n if self._params is None: self._build_params()\n x=x.contiguous(); x=F.conv2d(x,self.conv1.weight,None,2,3); s,b=self._params['bn1']; x=_ext.affine_relu(x,s,b); x=F.max_pool2d(x,3,2,1)\n for lname in ['layer1','layer2','layer3','layer4']:\n for i,blk in enumerate(getattr(self,lname)): x=self._block(x,blk,f'{lname}.{i}')\n return self.fc(torch.flatten(F.adaptive_avg_pool2d(x,(1,1)),1))\n"
}