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\nimport numpy as np, collections\nfrom itertools import repeat\n_cpp_src='''#include <torch/extension.h>\\ntorch::Tensor layer_norm_cuda(torch::Tensor x, torch::Tensor weight, torch::Tensor bias, double eps);\\ntorch::Tensor add_layer_norm_cuda(torch::Tensor residual, torch::Tensor x, torch::Tensor weight, torch::Tensor bias, double eps);\\ntorch::Tensor layer_norm(torch::Tensor x, torch::Tensor weight, torch::Tensor bias, double eps){return layer_norm_cuda(x,weight,bias,eps);}\\ntorch::Tensor add_layer_norm(torch::Tensor residual, torch::Tensor x, torch::Tensor weight, torch::Tensor bias, double eps){return add_layer_norm_cuda(residual,x,weight,bias,eps);}\\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__inline__ __device__ float warp_sum(float v){for(int o=16;o>0;o/=2)v+=__shfl_down_sync(0xffffffff,v,o);return v;}\n__inline__ __device__ float block_sum(float v){__shared__ float s[32];int lane=threadIdx.x&31,wid=threadIdx.x>>5;v=warp_sum(v);if(lane==0)s[wid]=v;__syncthreads();v=(threadIdx.x<(blockDim.x>>5))?s[lane]:0.0f;if(wid==0)v=warp_sum(v);return v;}\ntemplate<bool ADD> __global__ void ln_kernel(const float* __restrict__ r,const float* __restrict__ x,const float* __restrict__ w,const float* __restrict__ b,float* __restrict__ out,int rows,int C,float eps){int row=blockIdx.x;if(row>=rows)return;const float* xr=x+(long long)row*C;float sum=0,ss=0;for(int c=threadIdx.x;c<C;c+=blockDim.x){float v=xr[c];sum+=v;ss+=v*v;}sum=block_sum(sum);ss=block_sum(ss);float mean=sum/(float)C;float var=ss/(float)C-mean*mean;if(var<0)var=0;float inv=rsqrtf(var+eps);float* yr=out+(long long)row*C;const float* rr=r+(long long)row*C;for(int c=threadIdx.x;c<C;c+=blockDim.x){float y=(xr[c]-mean)*inv*w[c]+b[c];if(ADD)y+=rr[c];yr[c]=y;}}\ntorch::Tensor layer_norm_cuda(torch::Tensor x,torch::Tensor weight,torch::Tensor bias,double eps){auto xc=x.contiguous();auto out=torch::empty_like(xc);int C=(int)xc.size(-1),rows=(int)(xc.numel()/C);ln_kernel<false><<<rows,256,0,at::cuda::getCurrentCUDAStream()>>>(xc.data_ptr<float>(),xc.data_ptr<float>(),weight.data_ptr<float>(),bias.data_ptr<float>(),out.data_ptr<float>(),rows,C,(float)eps);return out;}\ntorch::Tensor add_layer_norm_cuda(torch::Tensor residual,torch::Tensor x,torch::Tensor weight,torch::Tensor bias,double eps){auto rc=residual.contiguous();auto xc=x.contiguous();auto out=torch::empty_like(xc);int C=(int)xc.size(-1),rows=(int)(xc.numel()/C);ln_kernel<true><<<rows,256,0,at::cuda::getCurrentCUDAStream()>>>(rc.data_ptr<float>(),xc.data_ptr<float>(),weight.data_ptr<float>(),bias.data_ptr<float>(),out.data_ptr<float>(),rows,C,(float)eps);return out;}\n'''\n_ext=load_inline(name='swinv2_fast_ln_ext2',cpp_sources=[_cpp_src],cuda_sources=[_cuda_src],functions=['layer_norm','add_layer_norm'],extra_cuda_cflags=['-O3','--use_fast_math'],extra_cflags=['-O3'],verbose=False)\ndef _ntuple(n):\n def parse(x): return tuple(x) if isinstance(x,collections.abc.Iterable) and not isinstance(x,str) else tuple(repeat(x,n))\n return parse\nto_2tuple=_ntuple(2)\nclass FastLayerNorm(nn.Module):\n def __init__(self,normalized_shape,eps=1e-5,elementwise_affine=True,bias=True,device=None,dtype=None):\n super().__init__(); self.normalized_shape=(normalized_shape,) if isinstance(normalized_shape,int) else tuple(normalized_shape); self.eps=eps; self.elementwise_affine=elementwise_affine\n if elementwise_affine:\n self.weight=nn.Parameter(torch.empty(self.normalized_shape,device=device,dtype=dtype)); self.bias=nn.Parameter(torch.empty(self.normalized_shape,device=device,dtype=dtype))\n else: self.register_parameter('weight',None); self.register_parameter('bias',None)\n self.reset_parameters()\n def reset_parameters(self):\n if self.elementwise_affine: nn.init.ones_(self.weight); nn.init.zeros_(self.bias)\n def forward(self,x):\n if x.is_cuda and x.dtype==torch.float32 and self.elementwise_affine and len(self.normalized_shape)==1: return _ext.layer_norm(x,self.weight,self.bias,float(self.eps))\n return F.layer_norm(x,self.normalized_shape,self.weight,self.bias,self.eps)\ndef fast_add_norm(residual,x,norm):\n if x.is_cuda and x.dtype==torch.float32 and isinstance(norm,FastLayerNorm): return _ext.add_layer_norm(residual,x,norm.weight,norm.bias,float(norm.eps))\n return residual+norm(x)\nclass Mlp(nn.Module):\n def __init__(self,in_features,hidden_features=None,out_features=None,act_layer=nn.GELU,drop=0.):\n super().__init__(); out_features=out_features or in_features; hidden_features=hidden_features or in_features; self.fc1=nn.Linear(in_features,hidden_features); self.act=act_layer(); self.fc2=nn.Linear(hidden_features,out_features); self.drop=nn.Dropout(drop)\n def forward(self,x): x=self.fc1(x); x=self.act(x); x=self.drop(x); x=self.fc2(x); x=self.drop(x); return x\ndef window_partition(x,window_size):\n B,H,W,C=x.shape; x=x.view(B,H//window_size,window_size,W//window_size,window_size,C); return x.permute(0,1,3,2,4,5).contiguous().view(-1,window_size,window_size,C)\ndef window_reverse(windows,window_size,H,W):\n B=int(windows.shape[0]/(H*W/window_size/window_size)); x=windows.view(B,H//window_size,W//window_size,window_size,window_size,-1); return x.permute(0,1,3,2,4,5).contiguous().view(B,H,W,-1)\nclass WindowAttention(nn.Module):\n def __init__(self,dim,window_size,num_heads,qkv_bias=True,attn_drop=0.,proj_drop=0.,pretrained_window_size=[0,0]):\n super().__init__(); self.dim=dim; self.window_size=window_size; self.pretrained_window_size=pretrained_window_size; self.num_heads=num_heads; self.logit_scale=nn.Parameter(torch.log(10*torch.ones((num_heads,1,1))),requires_grad=True); self.cpb_mlp=nn.Sequential(nn.Linear(2,512,bias=True),nn.ReLU(inplace=True),nn.Linear(512,num_heads,bias=False))\n rh=torch.arange(-(window_size[0]-1),window_size[0],dtype=torch.float32); rw=torch.arange(-(window_size[1]-1),window_size[1],dtype=torch.float32); table=torch.stack(torch.meshgrid([rh,rw])).permute(1,2,0).contiguous().unsqueeze(0)\n if pretrained_window_size[0]>0: table[:,:,:,0]/=(pretrained_window_size[0]-1); table[:,:,:,1]/=(pretrained_window_size[1]-1)\n else: table[:,:,:,0]/=(window_size[0]-1); table[:,:,:,1]/=(window_size[1]-1)\n table*=8; table=torch.sign(table)*torch.log2(torch.abs(table)+1.0)/np.log2(8); self.register_buffer('relative_coords_table',table)\n ch=torch.arange(window_size[0]); cw=torch.arange(window_size[1]); coords=torch.stack(torch.meshgrid([ch,cw])); cf=torch.flatten(coords,1); rc=cf[:,:,None]-cf[:,None,:]; rc=rc.permute(1,2,0).contiguous(); rc[:,:,0]+=window_size[0]-1; rc[:,:,1]+=window_size[1]-1; rc[:,:,0]*=2*window_size[1]-1; self.register_buffer('relative_position_index',rc.sum(-1))\n self.qkv=nn.Linear(dim,dim*3,bias=False); self.q_bias=nn.Parameter(torch.zeros(dim)) if qkv_bias else None; self.v_bias=nn.Parameter(torch.zeros(dim)) if qkv_bias else None; self.attn_drop=nn.Dropout(attn_drop); self.proj=nn.Linear(dim,dim); self.proj_drop=nn.Dropout(proj_drop); self.softmax=nn.Softmax(dim=-1)\n def forward(self,x,mask=None):\n B_,N,C=x.shape; bias=torch.cat((self.q_bias,torch.zeros_like(self.v_bias,requires_grad=False),self.v_bias)) if self.q_bias is not None else None; qkv=F.linear(x,self.qkv.weight,bias); qkv=qkv.reshape(B_,N,3,self.num_heads,-1).permute(2,0,3,1,4); q,k,v=qkv[0],qkv[1],qkv[2]\n attn=(F.normalize(q,dim=-1)@F.normalize(k,dim=-1).transpose(-2,-1)); logit_scale=torch.clamp(self.logit_scale.to(x.device),max=torch.log(torch.tensor(100.0,device=x.device))).exp(); attn=attn*logit_scale\n tbl=self.cpb_mlp(self.relative_coords_table).view(-1,self.num_heads); rb=tbl[self.relative_position_index.view(-1)].view(N,N,-1).permute(2,0,1).contiguous(); attn=attn+(16*torch.sigmoid(rb)).unsqueeze(0)\n if mask is not None:\n nW=mask.shape[0]; attn=attn.view(B_//nW,nW,self.num_heads,N,N)+mask.unsqueeze(1).unsqueeze(0); attn=attn.view(-1,self.num_heads,N,N)\n attn=self.softmax(attn); attn=self.attn_drop(attn); x=(attn@v).transpose(1,2).reshape(B_,N,C); x=self.proj(x); x=self.proj_drop(x); return x\nclass SwinTransformerBlock(nn.Module):\n def __init__(self,dim,input_resolution,num_heads,window_size=7,shift_size=0,mlp_ratio=4.,qkv_bias=True,drop=0.,attn_drop=0.,drop_path=0.,act_layer=nn.GELU,norm_layer=nn.LayerNorm,pretrained_window_size=0):\n super().__init__(); self.dim=dim; self.input_resolution=input_resolution; self.num_heads=num_heads; self.window_size=window_size; self.shift_size=shift_size; self.mlp_ratio=mlp_ratio\n if min(input_resolution)<=self.window_size: self.shift_size=0; self.window_size=min(input_resolution)\n self.norm1=norm_layer(dim); self.attn=WindowAttention(dim,to_2tuple(self.window_size),num_heads,qkv_bias,attn_drop,drop,to_2tuple(pretrained_window_size)); self.drop_path=nn.Identity(); self.norm2=norm_layer(dim); self.mlp=Mlp(dim,int(dim*mlp_ratio),act_layer=act_layer,drop=drop)\n if self.shift_size>0:\n H,W=input_resolution; img_mask=torch.zeros((1,H,W,1)); hs=(slice(0,-self.window_size),slice(-self.window_size,-self.shift_size),slice(-self.shift_size,None)); ws=hs; cnt=0\n for h in hs:\n for w in ws: img_mask[:,h,w,:]=cnt; cnt+=1\n mw=window_partition(img_mask,self.window_size).view(-1,self.window_size*self.window_size); am=mw.unsqueeze(1)-mw.unsqueeze(2); attn_mask=am.masked_fill(am!=0,float(-100.0)).masked_fill(am==0,float(0.0))\n else: attn_mask=None\n self.register_buffer('attn_mask',attn_mask)\n def forward(self,x):\n H,W=self.input_resolution; B,L,C=x.shape; shortcut=x; x=x.view(B,H,W,C); sx=torch.roll(x,shifts=(-self.shift_size,-self.shift_size),dims=(1,2)) if self.shift_size>0 else x; xw=window_partition(sx,self.window_size).view(-1,self.window_size*self.window_size,C); aw=self.attn(xw,mask=self.attn_mask); sx=window_reverse(aw.view(-1,self.window_size,self.window_size,C),self.window_size,H,W); x=torch.roll(sx,shifts=(self.shift_size,self.shift_size),dims=(1,2)) if self.shift_size>0 else sx; x=x.view(B,H*W,C); x=fast_add_norm(shortcut,x,self.norm1); x=fast_add_norm(x,self.mlp(x),self.norm2); return x\nclass PatchMerging(nn.Module):\n def __init__(self,input_resolution,dim,norm_layer=nn.LayerNorm): super().__init__(); self.input_resolution=input_resolution; self.dim=dim; self.reduction=nn.Linear(4*dim,2*dim,bias=False); self.norm=norm_layer(2*dim)\n def forward(self,x):\n H,W=self.input_resolution; B,L,C=x.shape; x=x.view(B,H,W,C); x=torch.cat([x[:,0::2,0::2,:],x[:,1::2,0::2,:],x[:,0::2,1::2,:],x[:,1::2,1::2,:]],-1).view(B,-1,4*C); x=self.reduction(x); x=self.norm(x); return x\nclass BasicLayer(nn.Module):\n def __init__(self,dim,input_resolution,depth,num_heads,window_size,mlp_ratio=4.,qkv_bias=True,drop=0.,attn_drop=0.,drop_path=0.,norm_layer=nn.LayerNorm,downsample=None,use_checkpoint=False,pretrained_window_size=0):\n super().__init__(); self.dim=dim; self.input_resolution=input_resolution; self.depth=depth; self.use_checkpoint=use_checkpoint; self.blocks=nn.ModuleList([SwinTransformerBlock(dim,input_resolution,num_heads,window_size,0 if i%2==0 else window_size//2,mlp_ratio,qkv_bias,drop,attn_drop,drop_path[i] if isinstance(drop_path,list) else drop_path,norm_layer=norm_layer,pretrained_window_size=pretrained_window_size) for i in range(depth)]); self.downsample=downsample(input_resolution,dim=dim,norm_layer=norm_layer) if downsample is not None else None\n def forward(self,x):\n for blk in self.blocks: x=blk(x)\n if self.downsample is not None: x=self.downsample(x)\n return x\nclass PatchEmbed(nn.Module):\n def __init__(self,img_size=224,patch_size=4,in_chans=3,embed_dim=96,norm_layer=None):\n super().__init__(); img_size=to_2tuple(img_size); patch_size=to_2tuple(patch_size); self.img_size=img_size; self.patch_size=patch_size; self.patches_resolution=[img_size[0]//patch_size[0],img_size[1]//patch_size[1]]; self.num_patches=self.patches_resolution[0]*self.patches_resolution[1]; self.in_chans=in_chans; self.embed_dim=embed_dim; self.proj=nn.Conv2d(in_chans,embed_dim,kernel_size=patch_size,stride=patch_size); self.norm=norm_layer(embed_dim) if norm_layer is not None else None\n def forward(self,x):\n x=self.proj(x).flatten(2).transpose(1,2)\n if self.norm is not None: x=self.norm(x)\n return x\nclass ModelNew(nn.Module):\n def __init__(self,img_size=224,patch_size=4,in_chans=3,num_classes=1000,embed_dim=96,depths=[2,2,6,2],num_heads=[3,6,12,24],window_size=7,mlp_ratio=4.,qkv_bias=True,drop_rate=0.,attn_drop_rate=0.,drop_path_rate=0.1,norm_layer=nn.LayerNorm,patch_norm=True,use_checkpoint=False,pretrained_window_sizes=[0,0,0,0],**kwargs):\n super().__init__(); self.num_classes=num_classes; self.num_layers=len(depths); self.embed_dim=embed_dim; self.patch_norm=patch_norm; self.num_features=int(embed_dim*2**(self.num_layers-1)); self.mlp_ratio=mlp_ratio; nl=FastLayerNorm if norm_layer is nn.LayerNorm else norm_layer; self.patch_embed=PatchEmbed(img_size,patch_size,in_chans,embed_dim,norm_layer=nl if patch_norm else None); pr=self.patch_embed.patches_resolution; self.patches_resolution=pr; self.pos_drop=nn.Dropout(p=drop_rate); dpr=[x.item() for x in torch.linspace(0,drop_path_rate,sum(depths))]; self.layers=nn.ModuleList()\n for i in range(self.num_layers): self.layers.append(BasicLayer(dim=int(embed_dim*2**i),input_resolution=(pr[0]//(2**i),pr[1]//(2**i)),depth=depths[i],num_heads=num_heads[i],window_size=window_size,mlp_ratio=mlp_ratio,qkv_bias=qkv_bias,drop=drop_rate,attn_drop=attn_drop_rate,drop_path=dpr[sum(depths[:i]):sum(depths[:i+1])],norm_layer=nl,downsample=PatchMerging if i<self.num_layers-1 else None,use_checkpoint=use_checkpoint,pretrained_window_size=pretrained_window_sizes[i]))\n self.norm=nl(self.num_features); self.avgpool=nn.AdaptiveAvgPool1d(1); self.head=nn.Linear(self.num_features,num_classes) if num_classes>0 else nn.Identity()\n def forward_features(self,x):\n x=self.patch_embed(x); x=self.pos_drop(x)\n for layer in self.layers: x=layer(x)\n x=self.norm(x); x=self.avgpool(x.transpose(1,2)); x=torch.flatten(x,1); return x\n def forward(self,x): return self.head(self.forward_features(x))\n"
}