# --------------------------------------------------------
# Swin Transformer V2
# Copyright (c) 2022 Microsoft
# Licensed under The MIT License [see LICENSE for details]
# Written by Ze Liu
# --------------------------------------------------------

import torch
import torch.nn as nn
import torch.nn.functional as F
import numpy as np
import collections
from itertools import repeat

def _ntuple(n):
    def parse(x):
        if isinstance(x, collections.abc.Iterable) and not isinstance(x, str):
            return tuple(x)
        return tuple(repeat(x, n))
    return parse
to_2tuple = _ntuple(2)

class Mlp(nn.Module):
    def __init__(self, in_features, hidden_features=None, out_features=None, act_layer=nn.GELU, drop=0.):
        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)

    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


def window_partition(x, window_size):
    """
    Args:
        x: (B, H, W, C)
        window_size (int): window size

    Returns:
        windows: (num_windows*B, window_size, window_size, C)
    """
    B, H, W, C = x.shape
    x = x.view(B, H // window_size, window_size, W // window_size, window_size, C)
    windows = x.permute(0, 1, 3, 2, 4, 5).contiguous().view(-1, window_size, window_size, C)
    return windows


def window_reverse(windows, window_size, H, W):
    """
    Args:
        windows: (num_windows*B, window_size, window_size, C)
        window_size (int): Window size
        H (int): Height of image
        W (int): Width of image

    Returns:
        x: (B, H, W, C)
    """
    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)
    x = x.permute(0, 1, 3, 2, 4, 5).contiguous().view(B, H, W, -1)
    return x


class WindowAttention(nn.Module):
    r""" Window based multi-head self attention (W-MSA) module with relative position bias.
    It supports both of shifted and non-shifted window.

    Args:
        dim (int): Number of input channels.
        window_size (tuple[int]): The height and width of the window.
        num_heads (int): Number of attention heads.
        qkv_bias (bool, optional):  If True, add a learnable bias to query, key, value. Default: True
        attn_drop (float, optional): Dropout ratio of attention weight. Default: 0.0
        proj_drop (float, optional): Dropout ratio of output. Default: 0.0
        pretrained_window_size (tuple[int]): The height and width of the window in pre-training.
    """

    def __init__(self, dim, window_size, num_heads, qkv_bias=True, attn_drop=0., proj_drop=0.,
                 pretrained_window_size=[0, 0]):

        super().__init__()
        self.dim = dim
        self.window_size = window_size  # Wh, Ww
        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)

        # mlp to generate continuous relative position bias
        self.cpb_mlp = nn.Sequential(nn.Linear(2, 512, bias=True),
                                     nn.ReLU(inplace=True),
                                     nn.Linear(512, num_heads, bias=False))

        # get relative_coords_table
        relative_coords_h = torch.arange(-(self.window_size[0] - 1), self.window_size[0], dtype=torch.float32)
        relative_coords_w = torch.arange(-(self.window_size[1] - 1), self.window_size[1], dtype=torch.float32)
        relative_coords_table = torch.stack(
            torch.meshgrid([relative_coords_h,
                            relative_coords_w])).permute(1, 2, 0).contiguous().unsqueeze(0)  # 1, 2*Wh-1, 2*Ww-1, 2
        if pretrained_window_size[0] > 0:
            relative_coords_table[:, :, :, 0] /= (pretrained_window_size[0] - 1)
            relative_coords_table[:, :, :, 1] /= (pretrained_window_size[1] - 1)
        else:
            relative_coords_table[:, :, :, 0] /= (self.window_size[0] - 1)
            relative_coords_table[:, :, :, 1] /= (self.window_size[1] - 1)
        relative_coords_table *= 8  # normalize to -8, 8
        relative_coords_table = torch.sign(relative_coords_table) * torch.log2(
            torch.abs(relative_coords_table) + 1.0) / np.log2(8)

        self.register_buffer("relative_coords_table", relative_coords_table)

        # get pair-wise relative position index for each token inside the window
        coords_h = torch.arange(self.window_size[0])
        coords_w = torch.arange(self.window_size[1])
        coords = torch.stack(torch.meshgrid([coords_h, coords_w]))  # 2, Wh, Ww
        coords_flatten = torch.flatten(coords, 1)  # 2, Wh*Ww
        relative_coords = coords_flatten[:, :, None] - coords_flatten[:, None, :]  # 2, Wh*Ww, Wh*Ww
        relative_coords = relative_coords.permute(1, 2, 0).contiguous()  # Wh*Ww, Wh*Ww, 2
        relative_coords[:, :, 0] += self.window_size[0] - 1  # shift to start from 0
        relative_coords[:, :, 1] += self.window_size[1] - 1
        relative_coords[:, :, 0] *= 2 * self.window_size[1] - 1
        relative_position_index = relative_coords.sum(-1)  # Wh*Ww, Wh*Ww
        self.register_buffer("relative_position_index", relative_position_index)

        self.qkv = nn.Linear(dim, dim * 3, bias=False)
        if qkv_bias:
            self.q_bias = nn.Parameter(torch.zeros(dim))
            self.v_bias = nn.Parameter(torch.zeros(dim))
        else:
            self.q_bias = None
            self.v_bias = 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)

    def forward(self, x, mask=None):
        """
        Args:
            x: input features with shape of (num_windows*B, N, C)
            mask: (0/-inf) mask with shape of (num_windows, Wh*Ww, Wh*Ww) or None
        """
        B_, N, C = x.shape
        qkv_bias = None
        if self.q_bias is not None:
            qkv_bias = torch.cat((self.q_bias, torch.zeros_like(self.v_bias, requires_grad=False), self.v_bias))
        qkv = F.linear(input=x, weight=self.qkv.weight, bias=qkv_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]  # make torchscript happy (cannot use tensor as tuple)

        # cosine attention
        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(1. / 0.01, device=x.device))).exp()
        attn = attn * logit_scale

        relative_position_bias_table = self.cpb_mlp(self.relative_coords_table).view(-1, self.num_heads)
        relative_position_bias = relative_position_bias_table[self.relative_position_index.view(-1)].view(
            self.window_size[0] * self.window_size[1], self.window_size[0] * self.window_size[1], -1)  # Wh*Ww,Wh*Ww,nH
        relative_position_bias = relative_position_bias.permute(2, 0, 1).contiguous()  # nH, Wh*Ww, Wh*Ww
        relative_position_bias = 16 * torch.sigmoid(relative_position_bias)
        attn = attn + relative_position_bias.unsqueeze(0)

        if mask is not None:
            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)
            attn = self.softmax(attn)
        else:
            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

class SwinTransformerBlock(nn.Module):
    r""" Swin Transformer Block.

    Args:
        dim (int): Number of input channels.
        input_resolution (tuple[int]): Input resulotion.
        num_heads (int): Number of attention heads.
        window_size (int): Window size.
        shift_size (int): Shift size for SW-MSA.
        mlp_ratio (float): Ratio of mlp hidden dim to embedding dim.
        qkv_bias (bool, optional): If True, add a learnable bias to query, key, value. Default: True
        drop (float, optional): Dropout rate. Default: 0.0
        attn_drop (float, optional): Attention dropout rate. Default: 0.0
        drop_path (float, optional): Stochastic depth rate. Default: 0.0
        act_layer (nn.Module, optional): Activation layer. Default: nn.GELU
        norm_layer (nn.Module, optional): Normalization layer.  Default: nn.LayerNorm
        pretrained_window_size (int): Window size in pre-training.
    """

    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):
        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
        if min(self.input_resolution) <= self.window_size:
            # if window size is larger than input resolution, we don't partition windows
            self.shift_size = 0
            self.window_size = min(self.input_resolution)
        assert 0 <= self.shift_size < self.window_size, "shift_size must in 0-window_size"

        self.norm1 = norm_layer(dim)
        self.attn = WindowAttention(
            dim, window_size=to_2tuple(self.window_size), num_heads=num_heads,
            qkv_bias=qkv_bias, attn_drop=attn_drop, proj_drop=drop,
            pretrained_window_size=to_2tuple(pretrained_window_size))

        self.drop_path = nn.Identity()
        self.norm2 = norm_layer(dim)
        mlp_hidden_dim = int(dim * mlp_ratio)
        self.mlp = Mlp(in_features=dim, hidden_features=mlp_hidden_dim, act_layer=act_layer, drop=drop)

        if self.shift_size > 0:
            # calculate attention mask for SW-MSA
            H, W = self.input_resolution
            img_mask = torch.zeros((1, H, W, 1))  # 1 H W 1
            h_slices = (slice(0, -self.window_size),
                        slice(-self.window_size, -self.shift_size),
                        slice(-self.shift_size, None))
            w_slices = (slice(0, -self.window_size),
                        slice(-self.window_size, -self.shift_size),
                        slice(-self.shift_size, None))
            cnt = 0
            for h in h_slices:
                for w in w_slices:
                    img_mask[:, h, w, :] = cnt
                    cnt += 1

            mask_windows = window_partition(img_mask, self.window_size)  # nW, window_size, window_size, 1
            mask_windows = mask_windows.view(-1, self.window_size * self.window_size)
            attn_mask = mask_windows.unsqueeze(1) - mask_windows.unsqueeze(2)
            attn_mask = attn_mask.masked_fill(attn_mask != 0, float(-100.0)).masked_fill(attn_mask == 0, float(0.0))
        else:
            attn_mask = None

        self.register_buffer("attn_mask", attn_mask)

    def forward(self, x):
        H, W = self.input_resolution
        B, L, C = x.shape
        assert L == H * W, "input feature has wrong size"

        shortcut = x
        x = x.view(B, H, W, C)

        # cyclic shift
        if self.shift_size > 0:
            shifted_x = torch.roll(x, shifts=(-self.shift_size, -self.shift_size), dims=(1, 2))
        else:
            shifted_x = x

        # partition windows
        x_windows = window_partition(shifted_x, self.window_size)  # nW*B, window_size, window_size, C
        x_windows = x_windows.view(-1, self.window_size * self.window_size, C)  # nW*B, window_size*window_size, C

        # W-MSA/SW-MSA
        attn_windows = self.attn(x_windows, mask=self.attn_mask)  # nW*B, window_size*window_size, C

        # merge windows
        attn_windows = attn_windows.view(-1, self.window_size, self.window_size, C)
        shifted_x = window_reverse(attn_windows, self.window_size, H, W)  # B H' W' C

        # reverse cyclic shift
        if self.shift_size > 0:
            x = torch.roll(shifted_x, shifts=(self.shift_size, self.shift_size), dims=(1, 2))
        else:
            x = shifted_x
        x = x.view(B, H * W, C)
        x = shortcut + self.drop_path(self.norm1(x))

        # FFN
        x = x + self.drop_path(self.norm2(self.mlp(x)))

        return x


class PatchMerging(nn.Module):
    r""" Patch Merging Layer.

    Args:
        input_resolution (tuple[int]): Resolution of input feature.
        dim (int): Number of input channels.
        norm_layer (nn.Module, optional): Normalization layer.  Default: nn.LayerNorm
    """

    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)

    def forward(self, x):
        """
        x: B, H*W, C
        """
        H, W = self.input_resolution
        B, L, C = x.shape
        assert L == H * W, "input feature has wrong size"
        assert H % 2 == 0 and W % 2 == 0, f"x size ({H}*{W}) are not even."

        x = x.view(B, H, W, C)

        x0 = x[:, 0::2, 0::2, :]  # B H/2 W/2 C
        x1 = x[:, 1::2, 0::2, :]  # B H/2 W/2 C
        x2 = x[:, 0::2, 1::2, :]  # B H/2 W/2 C
        x3 = x[:, 1::2, 1::2, :]  # B H/2 W/2 C
        x = torch.cat([x0, x1, x2, x3], -1)  # B H/2 W/2 4*C
        x = x.view(B, -1, 4 * C)  # B H/2*W/2 4*C

        x = self.reduction(x)
        x = self.norm(x)

        return x

class BasicLayer(nn.Module):
    """ A basic Swin Transformer layer for one stage.

    Args:
        dim (int): Number of input channels.
        input_resolution (tuple[int]): Input resolution.
        depth (int): Number of blocks.
        num_heads (int): Number of attention heads.
        window_size (int): Local window size.
        mlp_ratio (float): Ratio of mlp hidden dim to embedding dim.
        qkv_bias (bool, optional): If True, add a learnable bias to query, key, value. Default: True
        drop (float, optional): Dropout rate. Default: 0.0
        attn_drop (float, optional): Attention dropout rate. Default: 0.0
        drop_path (float | tuple[float], optional): Stochastic depth rate. Default: 0.0
        norm_layer (nn.Module, optional): Normalization layer. Default: nn.LayerNorm
        downsample (nn.Module | None, optional): Downsample layer at the end of the layer. Default: None
        use_checkpoint (bool): Whether to use checkpointing to save memory. Default: False.
        pretrained_window_size (int): Local window size in pre-training.
    """

    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):

        super().__init__()
        self.dim = dim
        self.input_resolution = input_resolution
        self.depth = depth
        self.use_checkpoint = use_checkpoint

        # build blocks
        self.blocks = nn.ModuleList([
            SwinTransformerBlock(dim=dim, input_resolution=input_resolution,
                                 num_heads=num_heads, window_size=window_size,
                                 shift_size=0 if (i % 2 == 0) else window_size // 2,
                                 mlp_ratio=mlp_ratio,
                                 qkv_bias=qkv_bias,
                                 drop=drop, attn_drop=attn_drop,
                                 drop_path=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)])

        # patch merging layer
        if downsample is not None:
            self.downsample = downsample(input_resolution, dim=dim, norm_layer=norm_layer)
        else:
            self.downsample = None

    def forward(self, x):
        for blk in self.blocks:
            if self.use_checkpoint:
                x = checkpoint.checkpoint(blk, x)
            else:
                x = blk(x)
        if self.downsample is not None:
            x = self.downsample(x)
        return x
    

class PatchEmbed(nn.Module):
    r""" Image to Patch Embedding

    Args:
        img_size (int): Image size.  Default: 224.
        patch_size (int): Patch token size. Default: 4.
        in_chans (int): Number of input image channels. Default: 3.
        embed_dim (int): Number of linear projection output channels. Default: 96.
        norm_layer (nn.Module, optional): Normalization layer. Default: None
    """

    def __init__(self, img_size=224, patch_size=4, in_chans=3, embed_dim=96, norm_layer=None):
        super().__init__()
        img_size = to_2tuple(img_size)
        patch_size = to_2tuple(patch_size)
        patches_resolution = [img_size[0] // patch_size[0], img_size[1] // patch_size[1]]
        self.img_size = img_size
        self.patch_size = patch_size
        self.patches_resolution = patches_resolution
        self.num_patches = patches_resolution[0] * 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)
        if norm_layer is not None:
            self.norm = norm_layer(embed_dim)
        else:
            self.norm = None

    def forward(self, x):
        B, C, H, W = x.shape
        # FIXME look at relaxing size constraints
        assert H == self.img_size[0] and W == self.img_size[1], \
            f"Input image size ({H}*{W}) doesn't match model ({self.img_size[0]}*{self.img_size[1]})."
        x = self.proj(x).flatten(2).transpose(1, 2)  # B Ph*Pw C
        if self.norm is not None:
            x = self.norm(x)
        return x


class Model(nn.Module):
    r""" Swin Transformer
        A PyTorch impl of : `Swin Transformer: Hierarchical Vision Transformer using Shifted Windows`  -
          https://arxiv.org/pdf/2103.14030

    Args:
        img_size (int | tuple(int)): Input image size. Default 224
        patch_size (int | tuple(int)): Patch size. Default: 4
        in_chans (int): Number of input image channels. Default: 3
        num_classes (int): Number of classes for classification head. Default: 1000
        embed_dim (int): Patch embedding dimension. Default: 96
        depths (tuple(int)): Depth of each Swin Transformer layer.
        num_heads (tuple(int)): Number of attention heads in different layers.
        window_size (int): Window size. Default: 7
        mlp_ratio (float): Ratio of mlp hidden dim to embedding dim. Default: 4
        qkv_bias (bool): If True, add a learnable bias to query, key, value. Default: True
        drop_rate (float): Dropout rate. Default: 0
        attn_drop_rate (float): Attention dropout rate. Default: 0
        drop_path_rate (float): Stochastic depth rate. Default: 0.1
        norm_layer (nn.Module): Normalization layer. Default: nn.LayerNorm.
        patch_norm (bool): If True, add normalization after patch embedding. Default: True
        use_checkpoint (bool): Whether to use checkpointing to save memory. Default: False
        pretrained_window_sizes (tuple(int)): Pretrained window sizes of each layer.
    """

    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):
        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

        # split image into non-overlapping patches
        self.patch_embed = PatchEmbed(
            img_size=img_size, patch_size=patch_size, in_chans=in_chans, embed_dim=embed_dim,
            norm_layer=norm_layer if self.patch_norm else None)
        num_patches = self.patch_embed.num_patches
        patches_resolution = self.patch_embed.patches_resolution
        self.patches_resolution = patches_resolution

        self.pos_drop = nn.Dropout(p=drop_rate)

        # stochastic depth
        dpr = [x.item() for x in torch.linspace(0, drop_path_rate, sum(depths))]  # stochastic depth decay rule

        # build layers
        self.layers = nn.ModuleList()
        for i_layer in range(self.num_layers):
            layer = BasicLayer(dim=int(embed_dim * 2 ** i_layer),
                               input_resolution=(patches_resolution[0] // (2 ** i_layer),
                                                 patches_resolution[1] // (2 ** i_layer)),
                               depth=depths[i_layer],
                               num_heads=num_heads[i_layer],
                               window_size=window_size,
                               mlp_ratio=self.mlp_ratio,
                               qkv_bias=qkv_bias,
                               drop=drop_rate, attn_drop=attn_drop_rate,
                               drop_path=dpr[sum(depths[:i_layer]):sum(depths[:i_layer + 1])],
                               norm_layer=norm_layer,
                               downsample=PatchMerging if (i_layer < self.num_layers - 1) else None,
                               use_checkpoint=use_checkpoint,
                               pretrained_window_size=pretrained_window_sizes[i_layer])
            self.layers.append(layer)

        self.norm = norm_layer(self.num_features)
        self.avgpool = nn.AdaptiveAvgPool1d(1)
        self.head = nn.Linear(self.num_features, num_classes) if num_classes > 0 else nn.Identity()

    def forward_features(self, x):
        x = self.patch_embed(x)
        x = self.pos_drop(x)

        for layer in self.layers:
            x = layer(x)

        x = self.norm(x)  # B L C
        x = self.avgpool(x.transpose(1, 2))  # B C 1
        x = torch.flatten(x, 1)
        return x

    def forward(self, x):
        x = self.forward_features(x)
        x = self.head(x)
        return x

batch_size = 10
image_size = 224

def get_inputs():
    return [torch.rand(batch_size, 3, image_size, image_size)]

def get_init_inputs():
    return []

import torch
import torch.nn as nn
import torch.nn.functional as F
import numpy as np
import collections
from itertools import repeat
from torch.utils.cpp_extension import load_inline
_cpp = r'''
#include <torch/extension.h>
torch::Tensor swin_partition_cuda(torch::Tensor x, int64_t H, int64_t W, int64_t window_size, int64_t shift);
torch::Tensor swin_reverse_cuda(torch::Tensor windows, int64_t B, int64_t H, int64_t W, int64_t C, int64_t window_size, int64_t shift);
torch::Tensor swin_partition(torch::Tensor x, int64_t H, int64_t W, int64_t window_size, int64_t shift) { return swin_partition_cuda(x, H, W, window_size, shift); }
torch::Tensor swin_reverse(torch::Tensor windows, int64_t B, int64_t H, int64_t W, int64_t C, int64_t window_size, int64_t shift) { return swin_reverse_cuda(windows, B, H, W, C, window_size, shift); }
'''
_cuda = r'''
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
template <typename scalar_t> __global__ void partition_kernel(const scalar_t* __restrict__ x, scalar_t* __restrict__ out,long total,int B,int H,int W,int C,int ws,int shift,int nWh,int nWw,int N){long idx=blockIdx.x*blockDim.x+threadIdx.x;if(idx>=total)return;int c=idx%C;long t=idx/C;int token=t%N;int bw=t/N;int win=bw%(nWh*nWw);int b=bw/(nWh*nWw);int th=token/ws;int tw=token-th*ws;int wh=win/nWw;int ww=win-wh*nWw;int ih=wh*ws+th;int iw=ww*ws+tw;if(shift>0){ih+=shift;if(ih>=H)ih-=H;iw+=shift;if(iw>=W)iw-=W;}out[idx]=x[((long)b*H*W+ih*W+iw)*C+c];}
template <typename scalar_t> __global__ void reverse_kernel(const scalar_t* __restrict__ windows, scalar_t* __restrict__ out,long total,int B,int H,int W,int C,int ws,int shift,int nWh,int nWw,int N){long idx=blockIdx.x*blockDim.x+threadIdx.x;if(idx>=total)return;int c=idx%C;long t=idx/C;int pos=t%(H*W);int b=t/(H*W);int h=pos/W;int w=pos-h*W;int sh=h;int sw=w;if(shift>0){sh-=shift;if(sh<0)sh+=H;sw-=shift;if(sw<0)sw+=W;}int wh=sh/ws;int ww=sw/ws;int th=sh-wh*ws;int tw=sw-ww*ws;int win=wh*nWw+ww;int token=th*ws+tw;long in_idx=(((long)b*(nWh*nWw)+win)*N+token)*C+c;out[idx]=windows[in_idx];}
torch::Tensor swin_partition_cuda(torch::Tensor x,int64_t H64,int64_t W64,int64_t ws64,int64_t shift64){int B=(int)x.size(0),H=(int)H64,W=(int)W64,C=(int)x.size(2),ws=(int)ws64,shift=(int)shift64;int nWh=H/ws,nWw=W/ws,N=ws*ws;auto out=torch::empty({(long)B*nWh*nWw,N,C},x.options());long total=out.numel();int threads=256;int blocks=(total+threads-1)/threads;AT_DISPATCH_FLOATING_TYPES(x.scalar_type(),"swin_partition_cuda",([&]{partition_kernel<scalar_t><<<blocks,threads,0,at::cuda::getCurrentCUDAStream()>>>(x.data_ptr<scalar_t>(),out.data_ptr<scalar_t>(),total,B,H,W,C,ws,shift,nWh,nWw,N);}));return out;}
torch::Tensor swin_reverse_cuda(torch::Tensor windows,int64_t B64,int64_t H64,int64_t W64,int64_t C64,int64_t ws64,int64_t shift64){int B=(int)B64,H=(int)H64,W=(int)W64,C=(int)C64,ws=(int)ws64,shift=(int)shift64;int nWh=H/ws,nWw=W/ws,N=ws*ws;auto out=torch::empty({B,H*W,C},windows.options());long total=out.numel();int threads=256;int blocks=(total+threads-1)/threads;AT_DISPATCH_FLOATING_TYPES(windows.scalar_type(),"swin_reverse_cuda",([&]{reverse_kernel<scalar_t><<<blocks,threads,0,at::cuda::getCurrentCUDAStream()>>>(windows.data_ptr<scalar_t>(),out.data_ptr<scalar_t>(),total,B,H,W,C,ws,shift,nWh,nWw,N);}));return out;}
'''
_ext=load_inline(name='swin_opt_ext_v1',cpp_sources=_cpp,cuda_sources=_cuda,functions=['swin_partition','swin_reverse'],extra_cuda_cflags=['-O3'],extra_cflags=['-O3'],verbose=False)
def _ntuple(n):
    def parse(x):
        if isinstance(x,collections.abc.Iterable) and not isinstance(x,str):return tuple(x)
        return tuple(repeat(x,n))
    return parse
to_2tuple=_ntuple(2)
class Mlp(nn.Module):
    def __init__(self,in_features,hidden_features=None,out_features=None,act_layer=nn.GELU,drop=0.):
        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)
    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
def window_partition(x,window_size):
    B,H,W,C=x.shape;return x.view(B,H//window_size,window_size,W//window_size,window_size,C).permute(0,1,3,2,4,5).contiguous().view(-1,window_size,window_size,C)
class WindowAttention(nn.Module):
    def __init__(self,dim,window_size,num_heads,qkv_bias=True,attn_drop=0.,proj_drop=0.,pretrained_window_size=[0,0]):
        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));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);rct=torch.stack(torch.meshgrid([rh,rw])).permute(1,2,0).contiguous().unsqueeze(0)
        if pretrained_window_size[0]>0:rct[:,:,:,0]/=(pretrained_window_size[0]-1);rct[:,:,:,1]/=(pretrained_window_size[1]-1)
        else:rct[:,:,:,0]/=(window_size[0]-1);rct[:,:,:,1]/=(window_size[1]-1)
        rct*=8;rct=torch.sign(rct)*torch.log2(torch.abs(rct)+1.0)/np.log2(8);self.register_buffer('relative_coords_table',rct);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));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);self._bias_cache=None;self._bias_cache_key=None
    def _rel_bias(self,device):
        key=(device,self.cpb_mlp[0].weight._version,self.cpb_mlp[0].bias._version,self.cpb_mlp[2].weight._version)
        if (not self.training) and self._bias_cache is not None and self._bias_cache_key==key:return self._bias_cache
        table=self.cpb_mlp(self.relative_coords_table.to(device)).view(-1,self.num_heads);N=self.window_size[0]*self.window_size[1];bias=table[self.relative_position_index.view(-1).to(device)].view(N,N,-1).permute(2,0,1).contiguous();bias=16*torch.sigmoid(bias)
        if not self.training:self._bias_cache=bias;self._bias_cache_key=key
        return bias
    def forward(self,x,mask=None):
        B_,N,C=x.shape;qkv_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,qkv_bias).reshape(B_,N,3,self.num_heads,C//self.num_heads).permute(2,0,3,1,4);q,k,v=qkv[0],qkv[1],qkv[2];q=F.normalize(q,dim=-1);k=F.normalize(k,dim=-1);logit_scale=torch.clamp(self.logit_scale.to(x.device),max=torch.log(torch.tensor(100.0,device=x.device))).exp();q=q*logit_scale.unsqueeze(0);rel=self._rel_bias(x.device);attn_mask=(rel.unsqueeze(0)+mask.to(x.device).repeat(B_//mask.shape[0],1,1).unsqueeze(1)) if mask is not None else rel.unsqueeze(0);y=F.scaled_dot_product_attention(q,k,v,attn_mask=attn_mask,dropout_p=(self.attn_drop.p if self.training else 0.0),scale=1.0);x=y.transpose(1,2).reshape(B_,N,C);x=self.proj(x);x=self.proj_drop(x);return x
class SwinTransformerBlock(nn.Module):
    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):
        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
        if min(input_resolution)<=self.window_size:self.shift_size=0;self.window_size=min(input_resolution)
        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)
        if self.shift_size>0:
            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));cnt=0
            for h in hs:
                for w in hs:img_mask[:,h,w,:]=cnt;cnt+=1
            mw=window_partition(img_mask,self.window_size).view(-1,self.window_size*self.window_size);am=mw.unsqueeze(1)-mw.unsqueeze(2);am=am.masked_fill(am!=0,float(-100.0)).masked_fill(am==0,float(0.0))
        else:am=None
        self.register_buffer('attn_mask',am)
    def forward(self,x):
        H,W=self.input_resolution;B,L,C=x.shape;shortcut=x;xw=_ext.swin_partition(x.contiguous(),H,W,self.window_size,self.shift_size);aw=self.attn(xw,mask=self.attn_mask);x=_ext.swin_reverse(aw.contiguous(),B,H,W,C,self.window_size,self.shift_size);x=shortcut+self.drop_path(self.norm1(x));x=x+self.drop_path(self.norm2(self.mlp(x)));return x
class PatchMerging(nn.Module):
    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)
    def forward(self,x):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
class BasicLayer(nn.Module):
    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):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
    def forward(self,x):
        for blk in self.blocks:x=blk(x)
        if self.downsample is not None:x=self.downsample(x)
        return x
class PatchEmbed(nn.Module):
    def __init__(self,img_size=224,patch_size=4,in_chans=3,embed_dim=96,norm_layer=None):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
    def forward(self,x):x=self.proj(x).permute(0,2,3,1).reshape(x.shape[0],-1,self.embed_dim);x=self.norm(x) if self.norm is not None else x;return x
class ModelNew(nn.Module):
    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):
        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;self.patch_embed=PatchEmbed(img_size,patch_size,in_chans,embed_dim,norm_layer if self.patch_norm else None);self.patches_resolution=self.patch_embed.patches_resolution;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()
        for i_layer in range(self.num_layers):self.layers.append(BasicLayer(int(embed_dim*2**i_layer),(self.patches_resolution[0]//(2**i_layer),self.patches_resolution[1]//(2**i_layer)),depths[i_layer],num_heads[i_layer],window_size,self.mlp_ratio,qkv_bias,drop_rate,attn_drop_rate,dpr[sum(depths[:i_layer]):sum(depths[:i_layer+1])],norm_layer,PatchMerging if i_layer<self.num_layers-1 else None,use_checkpoint,pretrained_window_sizes[i_layer]))
        self.norm=norm_layer(self.num_features);self.avgpool=nn.AdaptiveAvgPool1d(1);self.head=nn.Linear(self.num_features,num_classes) if num_classes>0 else nn.Identity()
    def forward_features(self,x):
        x=self.patch_embed(x);x=self.pos_drop(x)
        for layer in self.layers:x=layer(x)
        x=self.norm(x);x=self.avgpool(x.transpose(1,2));return torch.flatten(x,1)
    def forward(self,x):return self.head(self.forward_features(x))