diff --git a/.gitignore b/.gitignore index 99cdf5a..294c423 100644 --- a/.gitignore +++ b/.gitignore @@ -6,4 +6,27 @@ checkpoints/ .idea/ wandb/ *.pth -debug* \ No newline at end of file + +# C extensions +*.so + +# Distribution / packaging +.Python +build/ +develop-eggs/ +dist/ +downloads/ +eggs/ +.eggs/ +lib/ +lib64/ +parts/ +sdist/ +var/ +wheels/ +pip-wheel-metadata/ +share/python-wheels/ +*.egg-info/ +.installed.cfg +*.egg +MANIFEST \ No newline at end of file diff --git a/datasets_builders.py b/datasets_builders.py index 03f49c2..e6e9cbe 100644 --- a/datasets_builders.py +++ b/datasets_builders.py @@ -8,7 +8,6 @@ from timm.data.constants import IMAGENET_DEFAULT_MEAN, IMAGENET_DEFAULT_STD from timm.data import create_transform -from datasets import load_dataset from debug import DEBUG class INatDataset(ImageFolder): @@ -65,6 +64,7 @@ def build_dataset(is_train, data_path, args): dataset = datasets.ImageFolder(root, transform=transform) nb_classes = 1000 elif args.data_set == 'HUGGINGFACE': + from datasets import load_dataset def huggingface_transform(examples): examples["image"] = [transform(x.convert(mode="RGB")) for x in examples["image"]] return examples diff --git a/debug.py b/debug.py new file mode 100644 index 0000000..9f72275 --- /dev/null +++ b/debug.py @@ -0,0 +1 @@ +DEBUG=False \ No newline at end of file diff --git a/main.py b/main.py index 7447943..44511a8 100644 --- a/main.py +++ b/main.py @@ -228,10 +228,15 @@ def get_args_parser(): help='use scale-aware embeds') parser.add_argument('--grid-to-random-ratio', default=0.7, type=float, help='hybrid sampler grid to random ratio') + parser.add_argument('--model-type', default='deit', type=str, choices=["deit", "swin", "pvt"]) + + return parser def main(args): + if args.model_type != "deit": + assert args.eval, "For non-deit models, only evaluation is implemented." utils.init_distributed_mode(args) print(args) @@ -356,19 +361,39 @@ def log_to_wandb(log_dict, step): mixup_fn = None print(f"Creating model: {args.model}") - model = create_model( - args.model, - pretrained=False, - num_classes=args.nb_classes, - drop_rate=args.drop, - drop_path_rate=args.drop_path, - drop_block_rate=None, - img_size=args.input_size, - - Patch_layer=PatchEmbedHybrid, - use_learned_pos_embed = args.use_learned_pos_embed, - quantize_pos_embed = args.quantize_pos_embed - ) + if args.model_type == "deit": + model = create_model( + args.model, + pretrained=False, + num_classes=args.nb_classes, + drop_rate=args.drop, + drop_path_rate=args.drop_path, + drop_block_rate=None, + img_size=args.input_size, + + Patch_layer=PatchEmbedHybrid, + use_learned_pos_embed = args.use_learned_pos_embed, + quantize_pos_embed = args.quantize_pos_embed + ) + elif args.model_type == "swin": + from swin.config import get_config + from swin.models import build_model + class _Args: + pass + _args = _Args() + cfg_path = "swin/swin_base_patch4_window7_224.yaml" + setattr(_args, "cfg", cfg_path) + setattr(_args, "opts", []) + setattr(_args, "local_rank", 0) + setattr(_args, "data_path", args.data_path) + config = get_config(_args) + model = build_model(config) + elif args.model_type == "pvt": + assert args.model == 'pvt_v2_b5' + from pvt.pvt_v2 import pvt_v2_b5 + model = pvt_v2_b5() + else: + assert False if args.finetune: if args.finetune.startswith('https'): @@ -490,6 +515,8 @@ def log_to_wandb(log_dict, step): else: checkpoint = torch.load(args.resume, map_location='cpu') #del checkpoint['model']['pos_embed'] + if args.model_type == "pvt": + checkpoint = {"model": checkpoint} print(model_without_ddp.load_state_dict(checkpoint['model'], strict=False)) if not args.eval and 'optimizer' in checkpoint and 'lr_scheduler' in checkpoint and 'epoch' in checkpoint: optimizer.load_state_dict(checkpoint['optimizer']) diff --git a/pvt/pvt_v2.py b/pvt/pvt_v2.py new file mode 100644 index 0000000..eb7a08c --- /dev/null +++ b/pvt/pvt_v2.py @@ -0,0 +1,403 @@ +import torch +import torch.nn as nn +import torch.nn.functional as F +from functools import partial + +from timm.models.layers import DropPath, to_2tuple, trunc_normal_ +from timm.models.registry import register_model +from timm.models.vision_transformer import _cfg +import math + + +class Mlp(nn.Module): + def __init__(self, in_features, hidden_features=None, out_features=None, act_layer=nn.GELU, drop=0., linear=False): + 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.dwconv = DWConv(hidden_features) + self.act = act_layer() + self.fc2 = nn.Linear(hidden_features, out_features) + self.drop = nn.Dropout(drop) + self.linear = linear + if self.linear: + self.relu = nn.ReLU(inplace=True) + self.apply(self._init_weights) + + def _init_weights(self, m): + if isinstance(m, nn.Linear): + trunc_normal_(m.weight, std=.02) + if isinstance(m, nn.Linear) and m.bias is not None: + nn.init.constant_(m.bias, 0) + elif isinstance(m, nn.LayerNorm): + nn.init.constant_(m.bias, 0) + nn.init.constant_(m.weight, 1.0) + elif isinstance(m, nn.Conv2d): + fan_out = m.kernel_size[0] * m.kernel_size[1] * m.out_channels + fan_out //= m.groups + m.weight.data.normal_(0, math.sqrt(2.0 / fan_out)) + if m.bias is not None: + m.bias.data.zero_() + + def forward(self, x, H, W): + x = self.fc1(x) + if self.linear: + x = self.relu(x) + x = self.dwconv(x, H, W) + x = self.act(x) + x = self.drop(x) + x = self.fc2(x) + x = self.drop(x) + return x + + +class Attention(nn.Module): + def __init__(self, dim, num_heads=8, qkv_bias=False, qk_scale=None, attn_drop=0., proj_drop=0., sr_ratio=1, linear=False): + super().__init__() + assert dim % num_heads == 0, f"dim {dim} should be divided by num_heads {num_heads}." + + self.dim = dim + self.num_heads = num_heads + head_dim = dim // num_heads + self.scale = qk_scale or head_dim ** -0.5 + + self.q = nn.Linear(dim, dim, bias=qkv_bias) + self.kv = nn.Linear(dim, dim * 2, bias=qkv_bias) + self.attn_drop = nn.Dropout(attn_drop) + self.proj = nn.Linear(dim, dim) + self.proj_drop = nn.Dropout(proj_drop) + + self.linear = linear + self.sr_ratio = sr_ratio + if not linear: + if sr_ratio > 1: + self.sr = nn.Conv2d(dim, dim, kernel_size=sr_ratio, stride=sr_ratio) + self.norm = nn.LayerNorm(dim) + else: + self.pool = nn.AdaptiveAvgPool2d(7) + self.sr = nn.Conv2d(dim, dim, kernel_size=1, stride=1) + self.norm = nn.LayerNorm(dim) + self.act = nn.GELU() + self.apply(self._init_weights) + + def _init_weights(self, m): + if isinstance(m, nn.Linear): + trunc_normal_(m.weight, std=.02) + if isinstance(m, nn.Linear) and m.bias is not None: + nn.init.constant_(m.bias, 0) + elif isinstance(m, nn.LayerNorm): + nn.init.constant_(m.bias, 0) + nn.init.constant_(m.weight, 1.0) + elif isinstance(m, nn.Conv2d): + fan_out = m.kernel_size[0] * m.kernel_size[1] * m.out_channels + fan_out //= m.groups + m.weight.data.normal_(0, math.sqrt(2.0 / fan_out)) + if m.bias is not None: + m.bias.data.zero_() + + def forward(self, x, H, W): + B, N, C = x.shape + q = self.q(x).reshape(B, N, self.num_heads, C // self.num_heads).permute(0, 2, 1, 3) + + if not self.linear: + if self.sr_ratio > 1: + x_ = x.permute(0, 2, 1).reshape(B, C, H, W) + x_ = self.sr(x_).reshape(B, C, -1).permute(0, 2, 1) + x_ = self.norm(x_) + kv = self.kv(x_).reshape(B, -1, 2, self.num_heads, C // self.num_heads).permute(2, 0, 3, 1, 4) + else: + kv = self.kv(x).reshape(B, -1, 2, self.num_heads, C // self.num_heads).permute(2, 0, 3, 1, 4) + else: + x_ = x.permute(0, 2, 1).reshape(B, C, H, W) + x_ = self.sr(self.pool(x_)).reshape(B, C, -1).permute(0, 2, 1) + x_ = self.norm(x_) + x_ = self.act(x_) + kv = self.kv(x_).reshape(B, -1, 2, self.num_heads, C // self.num_heads).permute(2, 0, 3, 1, 4) + k, v = kv[0], kv[1] + + attn = (q @ k.transpose(-2, -1)) * self.scale + attn = attn.softmax(dim=-1) + 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 Block(nn.Module): + + def __init__(self, dim, num_heads, mlp_ratio=4., qkv_bias=False, qk_scale=None, drop=0., attn_drop=0., + drop_path=0., act_layer=nn.GELU, norm_layer=nn.LayerNorm, sr_ratio=1, linear=False): + super().__init__() + self.norm1 = norm_layer(dim) + self.attn = Attention( + dim, + num_heads=num_heads, qkv_bias=qkv_bias, qk_scale=qk_scale, + attn_drop=attn_drop, proj_drop=drop, sr_ratio=sr_ratio, linear=linear) + # NOTE: drop path for stochastic depth, we shall see if this is better than dropout here + self.drop_path = DropPath(drop_path) if drop_path > 0. else 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, linear=linear) + + self.apply(self._init_weights) + + def _init_weights(self, m): + if isinstance(m, nn.Linear): + trunc_normal_(m.weight, std=.02) + if isinstance(m, nn.Linear) and m.bias is not None: + nn.init.constant_(m.bias, 0) + elif isinstance(m, nn.LayerNorm): + nn.init.constant_(m.bias, 0) + nn.init.constant_(m.weight, 1.0) + elif isinstance(m, nn.Conv2d): + fan_out = m.kernel_size[0] * m.kernel_size[1] * m.out_channels + fan_out //= m.groups + m.weight.data.normal_(0, math.sqrt(2.0 / fan_out)) + if m.bias is not None: + m.bias.data.zero_() + + def forward(self, x, H, W): + x = x + self.drop_path(self.attn(self.norm1(x), H, W)) + x = x + self.drop_path(self.mlp(self.norm2(x), H, W)) + + return x + + +class OverlapPatchEmbed(nn.Module): + """ Image to Patch Embedding + """ + + def __init__(self, img_size=224, patch_size=7, stride=4, in_chans=3, embed_dim=768): + super().__init__() + + img_size = to_2tuple(img_size) + patch_size = to_2tuple(patch_size) + + assert max(patch_size) > stride, "Set larger patch_size than stride" + + self.img_size = img_size + self.patch_size = patch_size + self.H, self.W = img_size[0] // stride, img_size[1] // stride + self.num_patches = self.H * self.W + self.proj = nn.Conv2d(in_chans, embed_dim, kernel_size=patch_size, stride=stride, + padding=(patch_size[0] // 2, patch_size[1] // 2)) + self.norm = nn.LayerNorm(embed_dim) + + self.apply(self._init_weights) + + def _init_weights(self, m): + if isinstance(m, nn.Linear): + trunc_normal_(m.weight, std=.02) + if isinstance(m, nn.Linear) and m.bias is not None: + nn.init.constant_(m.bias, 0) + elif isinstance(m, nn.LayerNorm): + nn.init.constant_(m.bias, 0) + nn.init.constant_(m.weight, 1.0) + elif isinstance(m, nn.Conv2d): + fan_out = m.kernel_size[0] * m.kernel_size[1] * m.out_channels + fan_out //= m.groups + m.weight.data.normal_(0, math.sqrt(2.0 / fan_out)) + if m.bias is not None: + m.bias.data.zero_() + + def forward(self, x): + x = self.proj(x) + _, _, H, W = x.shape + x = x.flatten(2).transpose(1, 2) + x = self.norm(x) + + return x, H, W + + +class PyramidVisionTransformerV2(nn.Module): + def __init__(self, img_size=224, patch_size=16, in_chans=3, num_classes=1000, embed_dims=[64, 128, 256, 512], + num_heads=[1, 2, 4, 8], mlp_ratios=[4, 4, 4, 4], qkv_bias=False, qk_scale=None, drop_rate=0., + attn_drop_rate=0., drop_path_rate=0., norm_layer=nn.LayerNorm, + depths=[3, 4, 6, 3], sr_ratios=[8, 4, 2, 1], num_stages=4, linear=False): + super().__init__() + self.num_classes = num_classes + self.depths = depths + self.num_stages = num_stages + + dpr = [x.item() for x in torch.linspace(0, drop_path_rate, sum(depths))] # stochastic depth decay rule + cur = 0 + + for i in range(num_stages): + patch_embed = OverlapPatchEmbed(img_size=img_size if i == 0 else img_size // (2 ** (i + 1)), + patch_size=7 if i == 0 else 3, + stride=4 if i == 0 else 2, + in_chans=in_chans if i == 0 else embed_dims[i - 1], + embed_dim=embed_dims[i]) + + block = nn.ModuleList([Block( + dim=embed_dims[i], num_heads=num_heads[i], mlp_ratio=mlp_ratios[i], qkv_bias=qkv_bias, qk_scale=qk_scale, + drop=drop_rate, attn_drop=attn_drop_rate, drop_path=dpr[cur + j], norm_layer=norm_layer, + sr_ratio=sr_ratios[i], linear=linear) + for j in range(depths[i])]) + norm = norm_layer(embed_dims[i]) + cur += depths[i] + + setattr(self, f"patch_embed{i + 1}", patch_embed) + setattr(self, f"block{i + 1}", block) + setattr(self, f"norm{i + 1}", norm) + + # classification head + self.head = nn.Linear(embed_dims[3], num_classes) if num_classes > 0 else nn.Identity() + + self.apply(self._init_weights) + + def _init_weights(self, m): + if isinstance(m, nn.Linear): + trunc_normal_(m.weight, std=.02) + if isinstance(m, nn.Linear) and m.bias is not None: + nn.init.constant_(m.bias, 0) + elif isinstance(m, nn.LayerNorm): + nn.init.constant_(m.bias, 0) + nn.init.constant_(m.weight, 1.0) + elif isinstance(m, nn.Conv2d): + fan_out = m.kernel_size[0] * m.kernel_size[1] * m.out_channels + fan_out //= m.groups + m.weight.data.normal_(0, math.sqrt(2.0 / fan_out)) + if m.bias is not None: + m.bias.data.zero_() + + def freeze_patch_emb(self): + self.patch_embed1.requires_grad = False + + @torch.jit.ignore + def no_weight_decay(self): + return {'pos_embed1', 'pos_embed2', 'pos_embed3', 'pos_embed4', 'cls_token'} # has pos_embed may be better + + def get_classifier(self): + return self.head + + def reset_classifier(self, num_classes, global_pool=''): + self.num_classes = num_classes + self.head = nn.Linear(self.embed_dim, num_classes) if num_classes > 0 else nn.Identity() + + def forward_features(self, x, coords=None, keeps=None): + # Note that pyramid vit doesn't have explicit position embeddings, hence 'coords' argument is ignored. + assert keeps is None, "PVT transformer works only on merged patches with soft dropout. Set args --hard_dropout=0, --merge_patches=1 to proceed." + B = x.shape[0] + + for i in range(self.num_stages): + patch_embed = getattr(self, f"patch_embed{i + 1}") + block = getattr(self, f"block{i + 1}") + norm = getattr(self, f"norm{i + 1}") + x, H, W = patch_embed(x) + for blk in block: + x = blk(x, H, W) + x = norm(x) + if i != self.num_stages - 1: + x = x.reshape(B, H, W, -1).permute(0, 3, 1, 2).contiguous() + + return x.mean(dim=1) + + def forward(self, x, coords=None, keeps=None): + x = self.forward_features(x, coords, keeps) + x = self.head(x) + + return x + + +class DWConv(nn.Module): + def __init__(self, dim=768): + super(DWConv, self).__init__() + self.dwconv = nn.Conv2d(dim, dim, 3, 1, 1, bias=True, groups=dim) + + def forward(self, x, H, W): + B, N, C = x.shape + x = x.transpose(1, 2).view(B, C, H, W) + x = self.dwconv(x) + x = x.flatten(2).transpose(1, 2) + + return x + + +def _conv_filter(state_dict, patch_size=16): + """ convert patch embedding weight from manual patchify + linear proj to conv""" + out_dict = {} + for k, v in state_dict.items(): + if 'patch_embed.proj.weight' in k: + v = v.reshape((v.shape[0], 3, patch_size, patch_size)) + out_dict[k] = v + + return out_dict + + +@register_model +def pvt_v2_b0(pretrained=False, **kwargs): + model = PyramidVisionTransformerV2( + patch_size=4, embed_dims=[32, 64, 160, 256], num_heads=[1, 2, 5, 8], mlp_ratios=[8, 8, 4, 4], qkv_bias=True, + norm_layer=partial(nn.LayerNorm, eps=1e-6), depths=[2, 2, 2, 2], sr_ratios=[8, 4, 2, 1], + **kwargs) + model.default_cfg = _cfg() + + return model + + +@register_model +def pvt_v2_b1(pretrained=False, **kwargs): + model = PyramidVisionTransformerV2( + patch_size=4, embed_dims=[64, 128, 320, 512], num_heads=[1, 2, 5, 8], mlp_ratios=[8, 8, 4, 4], qkv_bias=True, + norm_layer=partial(nn.LayerNorm, eps=1e-6), depths=[2, 2, 2, 2], sr_ratios=[8, 4, 2, 1], + **kwargs) + model.default_cfg = _cfg() + + return model + + +@register_model +def pvt_v2_b2(pretrained=False, **kwargs): + model = PyramidVisionTransformerV2( + patch_size=4, embed_dims=[64, 128, 320, 512], num_heads=[1, 2, 5, 8], mlp_ratios=[8, 8, 4, 4], qkv_bias=True, + norm_layer=partial(nn.LayerNorm, eps=1e-6), depths=[3, 4, 6, 3], sr_ratios=[8, 4, 2, 1], **kwargs) + model.default_cfg = _cfg() + + return model + + +@register_model +def pvt_v2_b3(pretrained=False, **kwargs): + model = PyramidVisionTransformerV2( + patch_size=4, embed_dims=[64, 128, 320, 512], num_heads=[1, 2, 5, 8], mlp_ratios=[8, 8, 4, 4], qkv_bias=True, + norm_layer=partial(nn.LayerNorm, eps=1e-6), depths=[3, 4, 18, 3], sr_ratios=[8, 4, 2, 1], + **kwargs) + model.default_cfg = _cfg() + + return model + + +@register_model +def pvt_v2_b4(pretrained=False, **kwargs): + model = PyramidVisionTransformerV2( + patch_size=4, embed_dims=[64, 128, 320, 512], num_heads=[1, 2, 5, 8], mlp_ratios=[8, 8, 4, 4], qkv_bias=True, + norm_layer=partial(nn.LayerNorm, eps=1e-6), depths=[3, 8, 27, 3], sr_ratios=[8, 4, 2, 1], + **kwargs) + model.default_cfg = _cfg() + + return model + + +@register_model +def pvt_v2_b5(pretrained=False, **kwargs): + model = PyramidVisionTransformerV2( + patch_size=4, embed_dims=[64, 128, 320, 512], num_heads=[1, 2, 5, 8], mlp_ratios=[4, 4, 4, 4], qkv_bias=True, + norm_layer=partial(nn.LayerNorm, eps=1e-6), depths=[3, 6, 40, 3], sr_ratios=[8, 4, 2, 1], + **kwargs) + model.default_cfg = _cfg() + + return model + + +@register_model +def pvt_v2_b2_li(pretrained=False, **kwargs): + model = PyramidVisionTransformerV2( + patch_size=4, embed_dims=[64, 128, 320, 512], num_heads=[1, 2, 5, 8], mlp_ratios=[8, 8, 4, 4], qkv_bias=True, + norm_layer=partial(nn.LayerNorm, eps=1e-6), depths=[3, 4, 6, 3], sr_ratios=[8, 4, 2, 1], linear=True, **kwargs) + model.default_cfg = _cfg() + + return model diff --git a/run_pvt.sh b/run_pvt.sh new file mode 100755 index 0000000..e44b522 --- /dev/null +++ b/run_pvt.sh @@ -0,0 +1,57 @@ +#!/bin/bash + +model_name=STD_PVT5B +model_path=models/pvt_v2_b5.pth +input_size=224 +patch_size=16 +patch_shake=0 +patch_dropout=0 +patch_select=0 +patch_select_mode=drop +patch_zoom=0 + +model_type=pvt +hard_dropout=0 +merge_patches=1 + +#model_name=$1 +#model_path=$2 +#input_size=$3 +#patch_size=$4 +#patch_shake=$5 +#patch_dropout=$6 +#patch_select=$7 +#patch_select_mode=$8 +#patch_zoom=$9 + +run_name="${model_name}_GRID${input_size}x${patch_size}_SHAKE${patch_shake}_DROP${patch_dropout}_SELECT${patch_select}_MODE${patch_select_mode}_ZOOM${patch_zoom}" +echo $run_name +if [[ $model_name == *"STD"* ]]; then + quantize_pos_embed=1 + use_learned_pos_embed=1 +else + quantize_pos_embed=0 + use_learned_pos_embed=0 +fi + +python main.py \ + --model-type $model_type \ + --hard_dropout $hard_dropout \ + --merge_patches $merge_patches \ + --no-log-to-wandb \ + --model pvt_v2_b5 \ + --data-path ~/datasets/imagenet \ + --batch 16 --input-size $input_size --eval-crop-ratio 1.0 --seed 0 \ + --grid-patch-size $patch_size \ + --num_workers 10 \ + --eval --resume=$model_path \ + --patch_shake=$patch_shake \ + --patch_dropout=$patch_dropout \ + --patch_select=$patch_select \ + --patch_select_mode=$patch_select_mode \ + --quantize_pos_embed=$quantize_pos_embed \ + --use_learned_pos_embed=$use_learned_pos_embed \ + --patch_zoom=$patch_zoom \ + --modify_patches + + #--log-to-wandb --wandb_run_name=$run_name --wandb_model_name=$model_name \ diff --git a/run_swin.sh b/run_swin.sh new file mode 100755 index 0000000..b047c10 --- /dev/null +++ b/run_swin.sh @@ -0,0 +1,59 @@ +#!/bin/bash + +model_name=STD_SWINB +model_path=models/swin_base_patch4_window7_224.pth + +input_size=224 +patch_size=16 +patch_shake=0 +patch_dropout=0 +patch_select=0 +patch_select_mode=drop +patch_zoom=0 + +model_type=swin +hard_dropout=0 +merge_patches=1 + + +#model_name=$1 +#model_path=$2 +#input_size=$3 +#patch_size=$4 +#patch_shake=$5 +#patch_dropout=$6 +#patch_select=$7 +#patch_select_mode=$8 +#patch_zoom=$9 + +run_name="${model_name}_GRID${input_size}x${patch_size}_SHAKE${patch_shake}_DROP${patch_dropout}_SELECT${patch_select}_MODE${patch_select_mode}_ZOOM${patch_zoom}" +echo $run_name +if [[ $model_name == *"STD"* ]]; then + quantize_pos_embed=1 + use_learned_pos_embed=1 +else + quantize_pos_embed=0 + use_learned_pos_embed=0 +fi + +python main.py \ + --model-type $model_type \ + --hard_dropout $hard_dropout \ + --merge_patches $merge_patches \ + --no-log-to-wandb \ + --model _ \ + --data-path ~/datasets/imagenet/ILSVRC/Data/CLS-LOC/ \ + --batch 16 --input-size $input_size --eval-crop-ratio 1.0 --seed 0 \ + --grid-patch-size $patch_size \ + --num_workers 10 \ + --eval --resume=$model_path \ + --patch_shake=$patch_shake \ + --patch_dropout=$patch_dropout \ + --patch_select=$patch_select \ + --patch_select_mode=$patch_select_mode \ + --quantize_pos_embed=$quantize_pos_embed \ + --use_learned_pos_embed=$use_learned_pos_embed \ + --patch_zoom=$patch_zoom \ + --modify_patches + + #--log-to-wandb --wandb_run_name=$run_name --wandb_model_name=$model_name \ \ No newline at end of file diff --git a/scripts/run_all.sh b/scripts/run_all.sh index aa4187c..d96f14f 100755 --- a/scripts/run_all.sh +++ b/scripts/run_all.sh @@ -1,18 +1,15 @@ #!/bin/bash -source scripts/sweep_grid_scale.sh +set -e -#source scripts/sweep_zoom_exp.sh -#source scripts/sweep_dropout_exp.sh -#source scripts/sweep_shake_exp.sh +#source scripts/sweep_grid_scale.sh + +source scripts/sweep_zoom_exp.sh +source scripts/sweep_dropout_exp.sh +source scripts/sweep_shake_exp.sh source scripts/sweep_zoom_shake_exp.sh source scripts/sweep_zoom_dropout_exp.sh -#source scripts/sweep_dropout_shake_exp.sh -#source random - -#source edge counting -#source scripts/sweep_central_sampling_exp.sh +source scripts/sweep_dropout_shake_exp.sh -#source 30/70, 70/30 vits -#source scripts/sweep_dropout_vs_rescale_exp.sh \ No newline at end of file +source scripts/sweep_dropout_vs_rescale_exp.sh diff --git a/scripts/single_grid_mod_exp.sh b/scripts/single_grid_mod_exp.sh index 8c6f150..21f71ad 100755 --- a/scripts/single_grid_mod_exp.sh +++ b/scripts/single_grid_mod_exp.sh @@ -10,7 +10,27 @@ patch_select=$7 patch_select_mode=$8 patch_zoom=$9 -run_name="${model_name}_GRID${input_size}x${patch_size}_SHAKE${patch_shake}_DROP${patch_dropout}_SELECT${patch_select}_MODE${patch_select_mode}_ZOOM${patch_zoom}" + +# change if needed +hard_dropout=0 +merge_patches=1 +wandb="--log-to-wandb" +data_set="IMNET" +data_path="~/datasets/imagenet" +# change if needed + +if [[ $model_name == *"PVT"* ]]; then + model_type="pvt" + model="pvt_v2_b5" +elif [[ $model_name == *"SWIN"* ]]; then + model_type="swin" + model="swinb_p4_w7" +else + model_type="deit" + model="deit_base_patch16_LS" +fi + +run_name="${model_type}_${model_name}_HARD_DROP${hard_dropout}_MERGE_PATCHES_${merge_patches}_GRID${input_size}x${patch_size}_SHAKE${patch_shake}_DROP${patch_dropout}_SELECT${patch_select}_MODE${patch_select_mode}_ZOOM${patch_zoom}" echo $run_name if [[ $model_name == *"STD"* ]]; then quantize_pos_embed=1 @@ -20,14 +40,19 @@ else use_learned_pos_embed=0 fi + python main.py \ - --model deit_base_patch16_LS \ - --data-set HUGGINGFACE \ - --data-path /home/jan.olszewski/datasets/imagenet-1k/default/1.0.0/a1e9bfc56c3a7350165007d1176b15e9128fcaf9ab972147840529aed3ae52bc \ + --model-type $model_type \ + --hard_dropout $hard_dropout \ + --merge_patches $merge_patches \ + --model $model \ + --data-set $data_set \ + --data-path $data_path \ --batch 16 --input-size $input_size --eval-crop-ratio 1.0 --seed 0 \ --grid-patch-size $patch_size \ --num_workers 10 \ - --log-to-wandb --wandb_run_name=$run_name --wandb_model_name=$model_name \ + $wandb \ + --wandb_run_name=$run_name --wandb_model_name=$model_name \ --eval --resume=$model_path \ --patch_shake=$patch_shake \ --patch_dropout=$patch_dropout \ @@ -36,4 +61,4 @@ python main.py \ --quantize_pos_embed=$quantize_pos_embed \ --use_learned_pos_embed=$use_learned_pos_embed \ --patch_zoom=$patch_zoom \ - --modify_patches \ No newline at end of file + --modify_patches diff --git a/scripts/sweep_dropout_exp.sh b/scripts/sweep_dropout_exp.sh index 212250f..7addb62 100755 --- a/scripts/sweep_dropout_exp.sh +++ b/scripts/sweep_dropout_exp.sh @@ -1,22 +1,9 @@ #!/bin/bash +set -e -model_name=ELASTIC_VITB -model_path=models/base_448.pth -input_size=448 -patch_size=32 -patch_shake=0 -patch_dropout=0 -patch_select=0 -patch_select_mode=drop -patch_zoom=0 - -for ((patch_dropout=0; patch_dropout<=160; patch_dropout+=20)); do - source scripts/single_grid_mod_exp.sh $model_name $model_path $input_size $patch_size $patch_shake $patch_dropout $patch_select $patch_select_mode $patch_zoom -done - -model_name=STD_VITB -model_path=https://dl.fbaipublicfiles.com/deit/deit_3_base_224_1k.pth +model_name=STD_PVT5B +model_path="models/pvt_v2_b5.pth" input_size=224 patch_size=16 patch_shake=0 diff --git a/scripts/sweep_dropout_shake_exp.sh b/scripts/sweep_dropout_shake_exp.sh index 2325a04..7fac610 100755 --- a/scripts/sweep_dropout_shake_exp.sh +++ b/scripts/sweep_dropout_shake_exp.sh @@ -1,34 +1,19 @@ #!/bin/bash +set -e -model_name=ELASTIC_VITB -model_path=models/base_448.pth -input_size=448 -patch_size=32 -patch_shake=0 -patch_dropout=0 -patch_select=0 -patch_select_mode=drop -patch_zoom=0 - -for ((i=0; i<=8; i+=1)); do - patch_shake=$((2*i)) - patch_dropout=$((20*i)) - source scripts/single_grid_mod_exp.sh $model_name $model_path $input_size $patch_size $patch_shake $patch_dropout $patch_select $patch_select_mode $patch_zoom -done - -model_name=STD_VITB -model_path=https://dl.fbaipublicfiles.com/deit/deit_3_base_224_1k.pth +model_name=STD_PVT5B +model_path="models/pvt_v2_b5.pth" input_size=224 patch_size=16 patch_shake=0 patch_dropout=0 patch_select=0 -patch_select_mode=drop +patch_select_mode="drop" patch_zoom=0 for ((i=0; i<=8; i+=1)); do patch_shake=$i patch_dropout=$((20*i)) source scripts/single_grid_mod_exp.sh $model_name $model_path $input_size $patch_size $patch_shake $patch_dropout $patch_select $patch_select_mode $patch_zoom -done \ No newline at end of file +done diff --git a/scripts/sweep_dropout_vs_rescale_exp.sh b/scripts/sweep_dropout_vs_rescale_exp.sh index 7240dd2..c611f88 100755 --- a/scripts/sweep_dropout_vs_rescale_exp.sh +++ b/scripts/sweep_dropout_vs_rescale_exp.sh @@ -1,54 +1,25 @@ #!/bin/bash +set -e -model_name=ELASTIC_VITB -model_path=models/base_448.pth -input_size=448 -patch_size=32 -patch_shake=0 -patch_dropout=0 -patch_select=0 -patch_select_mode=drop -patch_selects=(30 40 49) -patch_zoom=0 - -for patch_select in ${patch_selects[@]}; do - source scripts/single_grid_mod_exp.sh $model_name $model_path $input_size $patch_size $patch_shake $patch_dropout $patch_select $patch_select_mode $patch_zoom -done - - -patch_selects=(0 10 20 30 40 49) -model_name=ELASTIC_VITB -model_path=models/base_448.pth -input_size=448 -patch_size=32 -patch_shake=0 -patch_dropout=0 -patch_select=0 -patch_select_mode=resize -patch_zoom=0 - -for patch_select in ${patch_selects[@]}; do - source scripts/single_grid_mod_exp.sh $model_name $model_path $input_size $patch_size $patch_shake $patch_dropout $patch_select $patch_select_mode $patch_zoom -done - -model_name=STD_VITB -model_path=https://dl.fbaipublicfiles.com/deit/deit_3_base_224_1k.pth +model_name=STD_PVT5B +model_path="models/pvt_v2_b5.pth" input_size=224 patch_size=16 patch_shake=0 patch_dropout=0 patch_select=0 -patch_select_mode=drop +patch_selecqt_mode="drop" patch_zoom=0 +patch_selects=(0 10 20 30 40 49) + for patch_select in ${patch_selects[@]}; do source scripts/single_grid_mod_exp.sh $model_name $model_path $input_size $patch_size $patch_shake $patch_dropout $patch_select $patch_select_mode $patch_zoom done - -model_name=STD_VITB -model_path=https://dl.fbaipublicfiles.com/deit/deit_3_base_224_1k.pth +model_name=STD_PVT5B +model_path="models/pvt_v2_b5.pth" input_size=224 patch_size=16 patch_shake=0 @@ -59,4 +30,4 @@ patch_zoom=0 for patch_select in ${patch_selects[@]}; do source scripts/single_grid_mod_exp.sh $model_name $model_path $input_size $patch_size $patch_shake $patch_dropout $patch_select $patch_select_mode $patch_zoom -done \ No newline at end of file +done diff --git a/scripts/sweep_grid_scale.sh b/scripts/sweep_grid_scale.sh index d4b5a21..6a89ce9 100755 --- a/scripts/sweep_grid_scale.sh +++ b/scripts/sweep_grid_scale.sh @@ -1,31 +1,15 @@ #!/bin/bash +set -e -model_name=ELASTIC_VITB -model_path=models/base_448.pth -input_size=448 -patch_size=32 -patch_shake=0 -patch_dropout=0 -patch_select=0 -patch_select_mode=drop -patch_zoom=0 - -for ((patch_size=32; patch_size<=64; patch_size+=4)); do - source scripts/single_grid_mod_exp.sh $model_name $model_path $input_size $patch_size $patch_shake $patch_dropout $patch_select $patch_select_mode $patch_zoom -done -for ((patch_size=30; patch_size>=16; patch_size-=2)); do - source scripts/single_grid_mod_exp.sh $model_name $model_path $input_size $patch_size $patch_shake $patch_dropout $patch_select $patch_select_mode $patch_zoom -done - -model_name=STD_VITB -model_path=https://dl.fbaipublicfiles.com/deit/deit_3_base_224_1k.pth +model_name=STD_PVT5B +model_path="models/pvt_v2_b5.pth" input_size=224 patch_size=16 patch_shake=0 patch_dropout=0 patch_select=0 -patch_select_mode=drop +patch_select_mode="drop" patch_zoom=0 for ((patch_size=16; patch_size<=32; patch_size+=2)); do @@ -33,4 +17,4 @@ for ((patch_size=16; patch_size<=32; patch_size+=2)); do done for ((patch_size=16; patch_size>=8; patch_size-=1)); do source scripts/single_grid_mod_exp.sh $model_name $model_path $input_size $patch_size $patch_shake $patch_dropout $patch_select $patch_select_mode $patch_zoom -done \ No newline at end of file +done diff --git a/scripts/sweep_scale_exp.sh b/scripts/sweep_scale_exp.sh index 048bd3f..2202b26 100755 --- a/scripts/sweep_scale_exp.sh +++ b/scripts/sweep_scale_exp.sh @@ -1,28 +1,15 @@ #!/bin/bash +set -e -model_name=ELASTIC_VITB -model_path=models/base_448.pth -input_size=448 -patch_size=32 -patch_shake=0 -patch_dropout=0 -patch_select=0 -patch_select_mode=drop -patch_zoom=0 - -for ((patch_shake=0; patch_shake<=16; patch_shake+=2)); do - source scripts/single_grid_mod_exp.sh $model_name $model_path $input_size $patch_size $patch_shake $patch_dropout $patch_select $patch_select_mode $patch_zoom -done - -model_name=STD_VITB -model_path=https://dl.fbaipublicfiles.com/deit/deit_3_base_224_1k.pth +model_name=STD_PVT5B +model_path="models/pvt_v2_b5.pth" input_size=224 patch_size=16 patch_shake=0 patch_dropout=0 patch_select=0 -patch_select_mode=drop +patch_select_mode="drop" patch_zoom=0 for ((patch_shake=0; patch_shake<=8; patch_shake+=1)); do diff --git a/scripts/sweep_shake_exp.sh b/scripts/sweep_shake_exp.sh index 048bd3f..2202b26 100755 --- a/scripts/sweep_shake_exp.sh +++ b/scripts/sweep_shake_exp.sh @@ -1,28 +1,15 @@ #!/bin/bash +set -e -model_name=ELASTIC_VITB -model_path=models/base_448.pth -input_size=448 -patch_size=32 -patch_shake=0 -patch_dropout=0 -patch_select=0 -patch_select_mode=drop -patch_zoom=0 - -for ((patch_shake=0; patch_shake<=16; patch_shake+=2)); do - source scripts/single_grid_mod_exp.sh $model_name $model_path $input_size $patch_size $patch_shake $patch_dropout $patch_select $patch_select_mode $patch_zoom -done - -model_name=STD_VITB -model_path=https://dl.fbaipublicfiles.com/deit/deit_3_base_224_1k.pth +model_name=STD_PVT5B +model_path="models/pvt_v2_b5.pth" input_size=224 patch_size=16 patch_shake=0 patch_dropout=0 patch_select=0 -patch_select_mode=drop +patch_select_mode="drop" patch_zoom=0 for ((patch_shake=0; patch_shake<=8; patch_shake+=1)); do diff --git a/scripts/sweep_zoom_dropout_exp.sh b/scripts/sweep_zoom_dropout_exp.sh index b7f7056..7ea1c32 100755 --- a/scripts/sweep_zoom_dropout_exp.sh +++ b/scripts/sweep_zoom_dropout_exp.sh @@ -1,35 +1,15 @@ #!/bin/bash +set -e -model_name=ELASTIC_VITB -model_path=models/base_448.pth -input_size=448 -patch_size=32 -patch_shake=0 -patch_dropout=0 -patch_select=0 -patch_select_mode=drop -patch_zoom=0 - -for ((i=0; i<=8; i+=1)); do - patch_dropout=$((20*i)) - patch_zoom=$((2*i)) - source scripts/single_grid_mod_exp.sh $model_name $model_path $input_size $patch_size $patch_shake $patch_dropout $patch_select $patch_select_mode $patch_zoom -done -for ((i=1; i<=8; i+=1)); do - patch_dropout=$((20*i)) - patch_zoom=$((-4*i)) - source scripts/single_grid_mod_exp.sh $model_name $model_path $input_size $patch_size $patch_shake $patch_dropout $patch_select $patch_select_mode $patch_zoom -done - -model_name=STD_VITB -model_path=https://dl.fbaipublicfiles.com/deit/deit_3_base_224_1k.pth +model_name=STD_PVT5B +model_path="models/pvt_v2_b5.pth" input_size=224 patch_size=16 patch_shake=0 patch_dropout=0 patch_select=0 -patch_select_mode=drop +patch_select_mode="drop" patch_zoom=0 for ((i=0; i<=8; i+=1)); do @@ -41,4 +21,4 @@ for ((i=1; i<=8; i+=1)); do patch_dropout=$((20*i)) patch_zoom=$((-2*i)) source scripts/single_grid_mod_exp.sh $model_name $model_path $input_size $patch_size $patch_shake $patch_dropout $patch_select $patch_select_mode $patch_zoom -done \ No newline at end of file +done diff --git a/scripts/sweep_zoom_exp.sh b/scripts/sweep_zoom_exp.sh index 394b7aa..ff1e61e 100755 --- a/scripts/sweep_zoom_exp.sh +++ b/scripts/sweep_zoom_exp.sh @@ -1,31 +1,15 @@ #!/bin/bash +set -e -model_name=ELASTIC_VITB -model_path=models/base_448.pth -input_size=448 -patch_size=32 -patch_shake=0 -patch_dropout=0 -patch_select=0 -patch_select_mode=drop -patch_zoom=0 - -for ((patch_zoom=0; patch_zoom<=16; patch_zoom+=2)); do - source scripts/single_grid_mod_exp.sh $model_name $model_path $input_size $patch_size $patch_shake $patch_dropout $patch_select $patch_select_mode $patch_zoom -done -for ((patch_zoom=-4; patch_zoom>=-32; patch_zoom-=4)); do - source scripts/single_grid_mod_exp.sh $model_name $model_path $input_size $patch_size $patch_shake $patch_dropout $patch_select $patch_select_mode $patch_zoom -done - -model_name=STD_VITB -model_path=https://dl.fbaipublicfiles.com/deit/deit_3_base_224_1k.pth +model_name=STD_PVT5B +model_path="models/pvt_v2_b5.pth" input_size=224 patch_size=16 patch_shake=0 patch_dropout=0 patch_select=0 -patch_select_mode=drop +patch_select_mode="drop" patch_zoom=0 @@ -34,4 +18,4 @@ for ((patch_zoom=0; patch_zoom<=8; patch_zoom+=1)); do done for ((patch_zoom=-2; patch_zoom>=-16; patch_zoom-=2)); do source scripts/single_grid_mod_exp.sh $model_name $model_path $input_size $patch_size $patch_shake $patch_dropout $patch_select $patch_select_mode $patch_zoom -done \ No newline at end of file +done diff --git a/scripts/sweep_zoom_shake_exp.sh b/scripts/sweep_zoom_shake_exp.sh index bddd0c5..35d06d1 100755 --- a/scripts/sweep_zoom_shake_exp.sh +++ b/scripts/sweep_zoom_shake_exp.sh @@ -1,35 +1,15 @@ #!/bin/bash +set -e -model_name=ELASTIC_VITB -model_path=models/base_448.pth -input_size=448 -patch_size=32 -patch_shake=0 -patch_dropout=0 -patch_select=0 -patch_select_mode=drop -patch_zoom=0 - -for ((i=0; i<=8; i+=1)); do - patch_shake=$((2*i)) - patch_zoom=$((2*i)) - source scripts/single_grid_mod_exp.sh $model_name $model_path $input_size $patch_size $patch_shake $patch_dropout $patch_select $patch_select_mode $patch_zoom -done -for ((i=1; i<=8; i+=1)); do - patch_shake=$((2*i)) - patch_zoom=$((-4*i)) - source scripts/single_grid_mod_exp.sh $model_name $model_path $input_size $patch_size $patch_shake $patch_dropout $patch_select $patch_select_mode $patch_zoom -done - -model_name=STD_VITB -model_path=https://dl.fbaipublicfiles.com/deit/deit_3_base_224_1k.pth +model_name=STD_PVT5B +model_path="models/pvt_v2_b5.pth" input_size=224 patch_size=16 patch_shake=0 patch_dropout=0 patch_select=0 -patch_select_mode=drop +patch_select_mode="drop" patch_zoom=0 for ((i=0; i<=8; i+=1)); do @@ -41,4 +21,4 @@ for ((i=1; i<=8; i+=1)); do patch_shake=$i patch_zoom=$((-2*i)) source scripts/single_grid_mod_exp.sh $model_name $model_path $input_size $patch_size $patch_shake $patch_dropout $patch_select $patch_select_mode $patch_zoom -done \ No newline at end of file +done diff --git a/swin/config.py b/swin/config.py new file mode 100644 index 0000000..1671ec3 --- /dev/null +++ b/swin/config.py @@ -0,0 +1,352 @@ +# -------------------------------------------------------- +# Swin Transformer +# Copyright (c) 2021 Microsoft +# Licensed under The MIT License [see LICENSE for details] +# Written by Ze Liu +# --------------------------------------------------------' + +import os +import yaml +from yacs.config import CfgNode as CN + +_C = CN() + +# Base config files +_C.BASE = [''] + +# ----------------------------------------------------------------------------- +# Data settings +# ----------------------------------------------------------------------------- +_C.DATA = CN() +# Batch size for a single GPU, could be overwritten by command line argument +_C.DATA.BATCH_SIZE = 128 +# Path to dataset, could be overwritten by command line argument +_C.DATA.DATA_PATH = '' +# Dataset name +_C.DATA.DATASET = 'imagenet' +# Input image size +_C.DATA.IMG_SIZE = 224 +# Interpolation to resize image (random, bilinear, bicubic) +_C.DATA.INTERPOLATION = 'bicubic' +# Use zipped dataset instead of folder dataset +# could be overwritten by command line argument +_C.DATA.ZIP_MODE = False +# Cache Data in Memory, could be overwritten by command line argument +_C.DATA.CACHE_MODE = 'part' +# Pin CPU memory in DataLoader for more efficient (sometimes) transfer to GPU. +_C.DATA.PIN_MEMORY = True +# Number of data loading threads +_C.DATA.NUM_WORKERS = 8 + +# [SimMIM] Mask patch size for MaskGenerator +_C.DATA.MASK_PATCH_SIZE = 32 +# [SimMIM] Mask ratio for MaskGenerator +_C.DATA.MASK_RATIO = 0.6 + +# ----------------------------------------------------------------------------- +# Model settings +# ----------------------------------------------------------------------------- +_C.MODEL = CN() +# Model type +_C.MODEL.TYPE = 'swin' +# Model name +_C.MODEL.NAME = 'swin_tiny_patch4_window7_224' +# Pretrained weight from checkpoint, could be imagenet22k pretrained weight +# could be overwritten by command line argument +_C.MODEL.PRETRAINED = '' +# Checkpoint to resume, could be overwritten by command line argument +_C.MODEL.RESUME = '' +# Number of classes, overwritten in data preparation +_C.MODEL.NUM_CLASSES = 1000 +# Dropout rate +_C.MODEL.DROP_RATE = 0.0 +# Drop path rate +_C.MODEL.DROP_PATH_RATE = 0.1 +# Label Smoothing +_C.MODEL.LABEL_SMOOTHING = 0.1 + +# Swin Transformer parameters +_C.MODEL.SWIN = CN() +_C.MODEL.SWIN.PATCH_SIZE = 4 +_C.MODEL.SWIN.IN_CHANS = 3 +_C.MODEL.SWIN.EMBED_DIM = 96 +_C.MODEL.SWIN.DEPTHS = [2, 2, 6, 2] +_C.MODEL.SWIN.NUM_HEADS = [3, 6, 12, 24] +_C.MODEL.SWIN.WINDOW_SIZE = 7 +_C.MODEL.SWIN.MLP_RATIO = 4. +_C.MODEL.SWIN.QKV_BIAS = True +_C.MODEL.SWIN.QK_SCALE = None +_C.MODEL.SWIN.APE = False +_C.MODEL.SWIN.PATCH_NORM = True + +# Swin Transformer V2 parameters +_C.MODEL.SWINV2 = CN() +_C.MODEL.SWINV2.PATCH_SIZE = 4 +_C.MODEL.SWINV2.IN_CHANS = 3 +_C.MODEL.SWINV2.EMBED_DIM = 96 +_C.MODEL.SWINV2.DEPTHS = [2, 2, 6, 2] +_C.MODEL.SWINV2.NUM_HEADS = [3, 6, 12, 24] +_C.MODEL.SWINV2.WINDOW_SIZE = 7 +_C.MODEL.SWINV2.MLP_RATIO = 4. +_C.MODEL.SWINV2.QKV_BIAS = True +_C.MODEL.SWINV2.APE = False +_C.MODEL.SWINV2.PATCH_NORM = True +_C.MODEL.SWINV2.PRETRAINED_WINDOW_SIZES = [0, 0, 0, 0] + +# Swin Transformer MoE parameters +_C.MODEL.SWIN_MOE = CN() +_C.MODEL.SWIN_MOE.PATCH_SIZE = 4 +_C.MODEL.SWIN_MOE.IN_CHANS = 3 +_C.MODEL.SWIN_MOE.EMBED_DIM = 96 +_C.MODEL.SWIN_MOE.DEPTHS = [2, 2, 6, 2] +_C.MODEL.SWIN_MOE.NUM_HEADS = [3, 6, 12, 24] +_C.MODEL.SWIN_MOE.WINDOW_SIZE = 7 +_C.MODEL.SWIN_MOE.MLP_RATIO = 4. +_C.MODEL.SWIN_MOE.QKV_BIAS = True +_C.MODEL.SWIN_MOE.QK_SCALE = None +_C.MODEL.SWIN_MOE.APE = False +_C.MODEL.SWIN_MOE.PATCH_NORM = True +_C.MODEL.SWIN_MOE.MLP_FC2_BIAS = True +_C.MODEL.SWIN_MOE.INIT_STD = 0.02 +_C.MODEL.SWIN_MOE.PRETRAINED_WINDOW_SIZES = [0, 0, 0, 0] +_C.MODEL.SWIN_MOE.MOE_BLOCKS = [[-1], [-1], [-1], [-1]] +_C.MODEL.SWIN_MOE.NUM_LOCAL_EXPERTS = 1 +_C.MODEL.SWIN_MOE.TOP_VALUE = 1 +_C.MODEL.SWIN_MOE.CAPACITY_FACTOR = 1.25 +_C.MODEL.SWIN_MOE.COSINE_ROUTER = False +_C.MODEL.SWIN_MOE.NORMALIZE_GATE = False +_C.MODEL.SWIN_MOE.USE_BPR = True +_C.MODEL.SWIN_MOE.IS_GSHARD_LOSS = False +_C.MODEL.SWIN_MOE.GATE_NOISE = 1.0 +_C.MODEL.SWIN_MOE.COSINE_ROUTER_DIM = 256 +_C.MODEL.SWIN_MOE.COSINE_ROUTER_INIT_T = 0.5 +_C.MODEL.SWIN_MOE.MOE_DROP = 0.0 +_C.MODEL.SWIN_MOE.AUX_LOSS_WEIGHT = 0.01 + +# Swin MLP parameters +_C.MODEL.SWIN_MLP = CN() +_C.MODEL.SWIN_MLP.PATCH_SIZE = 4 +_C.MODEL.SWIN_MLP.IN_CHANS = 3 +_C.MODEL.SWIN_MLP.EMBED_DIM = 96 +_C.MODEL.SWIN_MLP.DEPTHS = [2, 2, 6, 2] +_C.MODEL.SWIN_MLP.NUM_HEADS = [3, 6, 12, 24] +_C.MODEL.SWIN_MLP.WINDOW_SIZE = 7 +_C.MODEL.SWIN_MLP.MLP_RATIO = 4. +_C.MODEL.SWIN_MLP.APE = False +_C.MODEL.SWIN_MLP.PATCH_NORM = True + +# [SimMIM] Norm target during training +_C.MODEL.SIMMIM = CN() +_C.MODEL.SIMMIM.NORM_TARGET = CN() +_C.MODEL.SIMMIM.NORM_TARGET.ENABLE = False +_C.MODEL.SIMMIM.NORM_TARGET.PATCH_SIZE = 47 + +# ----------------------------------------------------------------------------- +# Training settings +# ----------------------------------------------------------------------------- +_C.TRAIN = CN() +_C.TRAIN.START_EPOCH = 0 +_C.TRAIN.EPOCHS = 300 +_C.TRAIN.WARMUP_EPOCHS = 20 +_C.TRAIN.WEIGHT_DECAY = 0.05 +_C.TRAIN.BASE_LR = 5e-4 +_C.TRAIN.WARMUP_LR = 5e-7 +_C.TRAIN.MIN_LR = 5e-6 +# Clip gradient norm +_C.TRAIN.CLIP_GRAD = 5.0 +# Auto resume from latest checkpoint +_C.TRAIN.AUTO_RESUME = True +# Gradient accumulation steps +# could be overwritten by command line argument +_C.TRAIN.ACCUMULATION_STEPS = 1 +# Whether to use gradient checkpointing to save memory +# could be overwritten by command line argument +_C.TRAIN.USE_CHECKPOINT = False + +# LR scheduler +_C.TRAIN.LR_SCHEDULER = CN() +_C.TRAIN.LR_SCHEDULER.NAME = 'cosine' +# Epoch interval to decay LR, used in StepLRScheduler +_C.TRAIN.LR_SCHEDULER.DECAY_EPOCHS = 30 +# LR decay rate, used in StepLRScheduler +_C.TRAIN.LR_SCHEDULER.DECAY_RATE = 0.1 +# warmup_prefix used in CosineLRScheduler +_C.TRAIN.LR_SCHEDULER.WARMUP_PREFIX = True +# [SimMIM] Gamma / Multi steps value, used in MultiStepLRScheduler +_C.TRAIN.LR_SCHEDULER.GAMMA = 0.1 +_C.TRAIN.LR_SCHEDULER.MULTISTEPS = [] + +# Optimizer +_C.TRAIN.OPTIMIZER = CN() +_C.TRAIN.OPTIMIZER.NAME = 'adamw' +# Optimizer Epsilon +_C.TRAIN.OPTIMIZER.EPS = 1e-8 +# Optimizer Betas +_C.TRAIN.OPTIMIZER.BETAS = (0.9, 0.999) +# SGD momentum +_C.TRAIN.OPTIMIZER.MOMENTUM = 0.9 + +# [SimMIM] Layer decay for fine-tuning +_C.TRAIN.LAYER_DECAY = 1.0 + +# MoE +_C.TRAIN.MOE = CN() +# Only save model on master device +_C.TRAIN.MOE.SAVE_MASTER = False +# ----------------------------------------------------------------------------- +# Augmentation settings +# ----------------------------------------------------------------------------- +_C.AUG = CN() +# Color jitter factor +_C.AUG.COLOR_JITTER = 0.4 +# Use AutoAugment policy. "v0" or "original" +_C.AUG.AUTO_AUGMENT = 'rand-m9-mstd0.5-inc1' +# Random erase prob +_C.AUG.REPROB = 0.25 +# Random erase mode +_C.AUG.REMODE = 'pixel' +# Random erase count +_C.AUG.RECOUNT = 1 +# Mixup alpha, mixup enabled if > 0 +_C.AUG.MIXUP = 0.8 +# Cutmix alpha, cutmix enabled if > 0 +_C.AUG.CUTMIX = 1.0 +# Cutmix min/max ratio, overrides alpha and enables cutmix if set +_C.AUG.CUTMIX_MINMAX = None +# Probability of performing mixup or cutmix when either/both is enabled +_C.AUG.MIXUP_PROB = 1.0 +# Probability of switching to cutmix when both mixup and cutmix enabled +_C.AUG.MIXUP_SWITCH_PROB = 0.5 +# How to apply mixup/cutmix params. Per "batch", "pair", or "elem" +_C.AUG.MIXUP_MODE = 'batch' + +# ----------------------------------------------------------------------------- +# Testing settings +# ----------------------------------------------------------------------------- +_C.TEST = CN() +# Whether to use center crop when testing +_C.TEST.CROP = True +# Whether to use SequentialSampler as validation sampler +_C.TEST.SEQUENTIAL = False +_C.TEST.SHUFFLE = False + +# ----------------------------------------------------------------------------- +# Misc +# ----------------------------------------------------------------------------- +# [SimMIM] Whether to enable pytorch amp, overwritten by command line argument +_C.ENABLE_AMP = False + +# Enable Pytorch automatic mixed precision (amp). +_C.AMP_ENABLE = True +# [Deprecated] Mixed precision opt level of apex, if O0, no apex amp is used ('O0', 'O1', 'O2') +_C.AMP_OPT_LEVEL = '' +# Path to output folder, overwritten by command line argument +_C.OUTPUT = '' +# Tag of experiment, overwritten by command line argument +_C.TAG = 'default' +# Frequency to save checkpoint +_C.SAVE_FREQ = 1 +# Frequency to logging info +_C.PRINT_FREQ = 10 +# Fixed random seed +_C.SEED = 0 +# Perform evaluation only, overwritten by command line argument +_C.EVAL_MODE = False +# Test throughput only, overwritten by command line argument +_C.THROUGHPUT_MODE = False +# local rank for DistributedDataParallel, given by command line argument +_C.LOCAL_RANK = 0 +# for acceleration +_C.FUSED_WINDOW_PROCESS = False +_C.FUSED_LAYERNORM = False + + +def _update_config_from_file(config, cfg_file): + config.defrost() + with open(cfg_file, 'r') as f: + yaml_cfg = yaml.load(f, Loader=yaml.FullLoader) + + for cfg in yaml_cfg.setdefault('BASE', ['']): + if cfg: + _update_config_from_file( + config, os.path.join(os.path.dirname(cfg_file), cfg) + ) + print('=> merge config from {}'.format(cfg_file)) + config.merge_from_file(cfg_file) + config.freeze() + + +def update_config(config, args): + _update_config_from_file(config, args.cfg) + + config.defrost() + if args.opts: + config.merge_from_list(args.opts) + + def _check_args(name): + if hasattr(args, name) and eval(f'args.{name}'): + return True + return False + + # merge from specific arguments + if _check_args('batch_size'): + config.DATA.BATCH_SIZE = args.batch_size + if _check_args('data_path'): + config.DATA.DATA_PATH = args.data_path + if _check_args('zip'): + config.DATA.ZIP_MODE = True + if _check_args('cache_mode'): + config.DATA.CACHE_MODE = args.cache_mode + if _check_args('pretrained'): + config.MODEL.PRETRAINED = args.pretrained + if _check_args('resume'): + config.MODEL.RESUME = args.resume + if _check_args('accumulation_steps'): + config.TRAIN.ACCUMULATION_STEPS = args.accumulation_steps + if _check_args('use_checkpoint'): + config.TRAIN.USE_CHECKPOINT = True + if _check_args('amp_opt_level'): + print("[warning] Apex amp has been deprecated, please use pytorch amp instead!") + if args.amp_opt_level == 'O0': + config.AMP_ENABLE = False + if _check_args('disable_amp'): + config.AMP_ENABLE = False + if _check_args('output'): + config.OUTPUT = args.output + if _check_args('tag'): + config.TAG = args.tag + if _check_args('eval'): + config.EVAL_MODE = True + if _check_args('throughput'): + config.THROUGHPUT_MODE = True + + # [SimMIM] + if _check_args('enable_amp'): + config.ENABLE_AMP = args.enable_amp + + # for acceleration + if _check_args('fused_window_process'): + config.FUSED_WINDOW_PROCESS = True + if _check_args('fused_layernorm'): + config.FUSED_LAYERNORM = True + ## Overwrite optimizer if not None, currently we use it for [fused_adam, fused_lamb] + if _check_args('optim'): + config.TRAIN.OPTIMIZER.NAME = args.optim + + # set local rank for distributed training + config.LOCAL_RANK = args.local_rank + + # output folder + config.OUTPUT = os.path.join(config.OUTPUT, config.MODEL.NAME, config.TAG) + + config.freeze() + + +def get_config(args): + """Get a yacs CfgNode object with default values.""" + # Return a clone so that the defaults will not be altered + # This is for the "local variable" use pattern + config = _C.clone() + update_config(config, args) + + return config diff --git a/swin/kernels/window_process/setup.py b/swin/kernels/window_process/setup.py new file mode 100644 index 0000000..c78526d --- /dev/null +++ b/swin/kernels/window_process/setup.py @@ -0,0 +1,12 @@ +from setuptools import setup +from torch.utils.cpp_extension import BuildExtension, CUDAExtension + + +setup(name='swin_window_process', + ext_modules=[ + CUDAExtension('swin_window_process', [ + 'swin_window_process.cpp', + 'swin_window_process_kernel.cu', + ]) + ], + cmdclass={'build_ext': BuildExtension}) \ No newline at end of file diff --git a/swin/kernels/window_process/swin_window_process.cpp b/swin/kernels/window_process/swin_window_process.cpp new file mode 100644 index 0000000..a7f15b0 --- /dev/null +++ b/swin/kernels/window_process/swin_window_process.cpp @@ -0,0 +1,132 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. All rights reserved. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include + + +at::Tensor roll_and_window_partition_forward_cuda( + at::Tensor & input, + //at::Tensor & output, + const int B, + const int H, + const int W, + const int C, + const int shift_size, + const int window_size); + + +at::Tensor roll_and_window_partition_backward_cuda( + at::Tensor & grad_in, + //at::Tensor & grad_out, + const int B, + const int H, + const int W, + const int C, + const int shift_size, + const int window_size); + + +at::Tensor window_merge_and_roll_forward_cuda( + at::Tensor & input, + //at::Tensor & output, + const int B, + const int H, + const int W, + const int C, + const int shift_size, + const int window_size); + +at::Tensor window_merge_and_roll_backward_cuda( + at::Tensor & grad_in, + //at::Tensor & grad_out, + const int B, + const int H, + const int W, + const int C, + const int shift_size, + const int window_size); + + +#define CHECK_CUDA(x) AT_ASSERTM(x.type().is_cuda(), #x " must be a CUDA tensor") +#define CHECK_CONTIGUOUS(x) AT_ASSERTM(x.is_contiguous(), #x " must be contiguous") +#define CHECK_INPUT(x) CHECK_CUDA(x); CHECK_CONTIGUOUS(x) + + + +at::Tensor roll_and_window_partition_forward( + at::Tensor & input, + //at::Tensor & output, + const int B, + const int H, + const int W, + const int C, + const int shift_size, + const int window_size){ + CHECK_INPUT(input); + return roll_and_window_partition_forward_cuda(input, B, H, W, C, shift_size, window_size); +} + + +at::Tensor roll_and_window_partition_backward( + at::Tensor & grad_in, + //at::Tensor & grad_out, + const int B, + const int H, + const int W, + const int C, + const int shift_size, + const int window_size){ + CHECK_INPUT(grad_in); + return roll_and_window_partition_backward_cuda(grad_in, B, H, W, C, shift_size, window_size); +} + + +at::Tensor window_merge_and_roll_forward( + at::Tensor & input, + //at::Tensor & output, + const int B, + const int H, + const int W, + const int C, + const int shift_size, + const int window_size){ + CHECK_INPUT(input); + return window_merge_and_roll_forward_cuda(input, B, H, W, C, shift_size, window_size); +} + + +at::Tensor window_merge_and_roll_backward( + at::Tensor & grad_in, + //at::Tensor & grad_out, + const int B, + const int H, + const int W, + const int C, + const int shift_size, + const int window_size){ + CHECK_INPUT(grad_in); + return window_merge_and_roll_backward_cuda(grad_in, B, H, W, C, shift_size, window_size); +} + + + +PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) { + m.def("roll_and_window_partition_forward", &roll_and_window_partition_forward, "torch.roll and window_partition."); + m.def("roll_and_window_partition_backward", &roll_and_window_partition_backward, "torch.roll and window_partition."); + m.def("window_merge_and_roll_forward", &window_merge_and_roll_forward, "window merge and torch.roll."); + m.def("window_merge_and_roll_backward", &window_merge_and_roll_backward, "window merge and torch.roll."); +} \ No newline at end of file diff --git a/swin/kernels/window_process/swin_window_process_kernel.cu b/swin/kernels/window_process/swin_window_process_kernel.cu new file mode 100644 index 0000000..4976949 --- /dev/null +++ b/swin/kernels/window_process/swin_window_process_kernel.cu @@ -0,0 +1,323 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. All rights reserved. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include +#include +#include +#include +#include + +int best_block_dim(int feat_dim){ + int best_dim; + if (feat_dim < 384){ + best_dim = 64; + } + else{ + if (feat_dim < 1024){ + best_dim = 128; + } + else{ + best_dim = 256; + } + } + return best_dim; +} + + +template +__global__ void roll_and_window_partition_forward_cuda_kernel( + T* input, + T* output, + const int B, + const int H, + const int W, + const int C, + const int shift_size, + const int window_size, + const int nH, + const int nW){ + // start + //bool qual = threadIdx.x < C; + int index = threadIdx.x; + int offset; + for (int i = index; i < C; i += blockDim.x) { + offset = ((blockIdx.z * gridDim.y + blockIdx.y) * gridDim.x + blockIdx.x) * C + i; // C = blocksize + int input_offset = blockIdx.z / (nH * nW) * H * W * C + + (blockIdx.z % (nH * nW) / nW * window_size + blockIdx.y - shift_size + H) % H * W * C + + (blockIdx.z % nW * window_size + blockIdx.x - shift_size + W) % W * C + + i; + output[offset] = (T)(__ldg(input + input_offset)); + } +} + + +template +__global__ void roll_and_window_partition_backward_cuda_kernel( + T* grad_in, + T* grad_out, + const int B, + const int H, + const int W, + const int C, + const int shift_size, + const int window_size, + const int nH, + const int nW){ + // start + int index = threadIdx.x; + int offset; + for (int i = index; i < C; i += blockDim.x) { + offset = ((blockIdx.z * gridDim.y + blockIdx.y) * gridDim.x + blockIdx.x) * C + i; // C = blocksize + int input_offset = + (blockIdx.z * nH * nW + (blockIdx.y + shift_size + H) % H / window_size * nW + (blockIdx.x + shift_size + W) % W / window_size) * window_size * window_size * C + + (blockIdx.y + shift_size + H ) % H % window_size * window_size * C + + (blockIdx.x + shift_size + W ) % W % window_size * C + + i; + grad_out[offset] = (T)(__ldg(grad_in + input_offset)); + } +} + + +template +__global__ void window_merge_and_roll_forward_cuda_kernel( + T* input, + T* output, + const int B, + const int H, + const int W, + const int C, + const int shift_size, + const int window_size, + const int nH, + const int nW){ + // start + int index = threadIdx.x; + int offset; + for (int i = index; i < C; i += blockDim.x) { + offset = ((blockIdx.z * gridDim.y + blockIdx.y) * gridDim.x + blockIdx.x) * C + i; // C = blocksize + int input_offset = + (blockIdx.z * nH * nW + (blockIdx.y - shift_size + H) % H / window_size * nH + (blockIdx.x - shift_size + W) % W / window_size) * window_size * window_size * C + + (blockIdx.y - shift_size + H) % window_size * window_size * C + + (blockIdx.x - shift_size + W) % window_size * C + + i; + output[offset] = (T)(__ldg(input + input_offset)); + } +} + + + +template +__global__ void window_merge_and_roll_backward_cuda_kernel( + T* grad_in, + T* grad_out, + const int B, + const int H, + const int W, + const int C, + const int shift_size, + const int window_size, + const int nH, + const int nW){ + // start + int index = threadIdx.x; + int offset; + for (int i = index; i < C; i += blockDim.x) { + offset = ((blockIdx.z * gridDim.y + blockIdx.y) * gridDim.x + blockIdx.x) * C + i; // C = blocksize + int input_offset = + (blockIdx.z / (nH * nW)) * H * W * C + + (blockIdx.z % (nH * nW) / nW * window_size + blockIdx.y + shift_size + H) % H * W * C + + (blockIdx.z % nW * window_size + blockIdx.x + shift_size + W) % W * C + + i; + grad_out[offset] = (T)(__ldg(grad_in + input_offset)); + } +} + +// input: [B, H, W, C] +// output: [B*nH*nW, window_size, window_size, C] +at::Tensor roll_and_window_partition_forward_cuda( + at::Tensor & input, + //at::Tensor & output, + const int B, + const int H, + const int W, + const int C, + const int shift_size, + const int window_size){ + + int nH = H / window_size; + int nW = W / window_size; + + dim3 grid(window_size, window_size, B * nH * nW); + //dim3 block((C + 31) / 32 * 32); + int blocknum = best_block_dim(C); + dim3 block(blocknum); + + at::Tensor output; + if (input.scalar_type() == torch::kFloat16){ + output = torch::empty({B*nH*nW, window_size, window_size, C}, torch::dtype(torch::kFloat16).device(torch::kCUDA).requires_grad(true)); + } + else{ + output = torch::empty({B*nH*nW, window_size, window_size, C}, torch::dtype(torch::kFloat32).device(torch::kCUDA).requires_grad(true)); + } + + AT_DISPATCH_FLOATING_TYPES_AND_HALF(input.type(), "roll_and_window_partition_forward_cuda_kernel", ([&] { + roll_and_window_partition_forward_cuda_kernel<<>>( + input.data(), + output.data(), + B, + H, + W, + C, + shift_size, + window_size, + nH, + nW); + })); + return output; +} + + +// grad_in: [B*nH*nW, window_size, window_size, C] +// grad_out: [B, H, W, C] +at::Tensor roll_and_window_partition_backward_cuda( + at::Tensor & grad_in, + const int B, + const int H, + const int W, + const int C, + const int shift_size, + const int window_size){ + + int nH = H / window_size; + int nW = W / window_size; + + dim3 grid(W, H, B); + //dim3 block((C + 31) / 32 * 32); + int blocknum = best_block_dim(C); + dim3 block(blocknum); + + at::Tensor grad_out; + if (grad_in.scalar_type() == torch::kFloat16){ + grad_out = torch::empty({B, H, W, C}, torch::dtype(torch::kFloat16).device(torch::kCUDA).requires_grad(false)); + } + else{ + grad_out = torch::empty({B, H, W, C}, torch::dtype(torch::kFloat32).device(torch::kCUDA).requires_grad(false)); + } + + AT_DISPATCH_FLOATING_TYPES_AND_HALF(grad_in.type(), "roll_and_window_partition_backward_cuda_kernel", ([&] { + roll_and_window_partition_backward_cuda_kernel<<>>( + grad_in.data(), + grad_out.data(), + B, + H, + W, + C, + shift_size, + window_size, + nH, + nW); + })); + return grad_out; +} + + +// input: [B*nH*nW, window_size, window_size, C] +// output: [B, H, W, C] +at::Tensor window_merge_and_roll_forward_cuda( + at::Tensor & input, + //at::Tensor & output, + const int B, + const int H, + const int W, + const int C, + const int shift_size, + const int window_size){ + + int nH = H / window_size; + int nW = W / window_size; + + dim3 grid(W, H, B); + //dim3 block((C + 31) / 32 * 32); + int blocknum = best_block_dim(C); + dim3 block(blocknum); + + //generate output tensor inside + at::Tensor output; + if (input.scalar_type() == torch::kFloat16){ + output = torch::empty({B, H, W, C}, torch::dtype(torch::kFloat16).device(torch::kCUDA).requires_grad(true)); + } + else{ + output = torch::empty({B, H, W, C}, torch::dtype(torch::kFloat32).device(torch::kCUDA).requires_grad(true)); + } + + AT_DISPATCH_FLOATING_TYPES_AND_HALF(input.type(), "window_merge_and_roll_forward_cuda_kernel", ([&] { + window_merge_and_roll_forward_cuda_kernel<<>>( + input.data(), + output.data(), + B, + H, + W, + C, + shift_size, + window_size, + nH, + nW); + })); + return output; +} + + +at::Tensor window_merge_and_roll_backward_cuda( + at::Tensor & grad_in, + const int B, + const int H, + const int W, + const int C, + const int shift_size, + const int window_size){ + + int nH = H / window_size; + int nW = W / window_size; + + dim3 grid(window_size, window_size, B * nH * nW); + //dim3 block((C + 31) / 32 * 32); + int blocknum = best_block_dim(C); + dim3 block(blocknum); + + at::Tensor grad_out; + if (grad_in.scalar_type() == torch::kFloat16){ + grad_out = torch::empty({B*nH*nW, window_size, window_size, C}, torch::dtype(torch::kFloat16).device(torch::kCUDA).requires_grad(false)); + } + else{ + grad_out = torch::empty({B*nH*nW, window_size, window_size, C}, torch::dtype(torch::kFloat32).device(torch::kCUDA).requires_grad(false)); + } + + AT_DISPATCH_FLOATING_TYPES_AND_HALF(grad_in.type(), "window_merge_and_roll_backward_cuda_kernel", ([&] { + window_merge_and_roll_backward_cuda_kernel<<>>( + grad_in.data(), + grad_out.data(), + B, + H, + W, + C, + shift_size, + window_size, + nH, + nW); + })); + return grad_out; +} \ No newline at end of file diff --git a/swin/kernels/window_process/unit_test.py b/swin/kernels/window_process/unit_test.py new file mode 100644 index 0000000..65dee56 --- /dev/null +++ b/swin/kernels/window_process/unit_test.py @@ -0,0 +1,250 @@ +# -------------------------------------------------------- +# Fused kernel for window process for SwinTransformer +# Copyright (c) 2022 Nvidia +# Licensed under The MIT License [see LICENSE for details] +# -------------------------------------------------------- + +import torch +import swin_window_process +import random +import time +import unittest + + +class WindowProcess(torch.autograd.Function): + @staticmethod + def forward(ctx, input, B, H, W, C, shift_size, window_size): + output = swin_window_process.roll_and_window_partition_forward(input, B, H, W, C, shift_size, window_size) + + ctx.B = B + ctx.H = H + ctx.W = W + ctx.C = C + ctx.shift_size = shift_size + ctx.window_size = window_size + return output + + @staticmethod + def backward(ctx, grad_in): + B = ctx.B + H = ctx.H + W = ctx.W + C = ctx.C + shift_size = ctx.shift_size + window_size = ctx.window_size + + grad_out = swin_window_process.roll_and_window_partition_backward(grad_in, B, H, W, C, shift_size, window_size) + return grad_out, None, None, None, None, None, None, None + + +class WindowProcessReverse(torch.autograd.Function): + @staticmethod + def forward(ctx, input, B, H, W, C, shift_size, window_size): + output = swin_window_process.window_merge_and_roll_forward(input, B, H, W, C, shift_size, window_size) + + ctx.B = B + ctx.H = H + ctx.W = W + ctx.C = C + ctx.shift_size = shift_size + ctx.window_size = window_size + + return output + + @staticmethod + def backward(ctx, grad_in): + B = ctx.B + H = ctx.H + W = ctx.W + C = ctx.C + shift_size = ctx.shift_size + window_size = ctx.window_size + + grad_out = swin_window_process.window_merge_and_roll_backward(grad_in, B, H, W, C, shift_size, window_size) + return grad_out, None, None, None, None, None, None, None + + +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 + + +def pyt_forward(x, shift_size, window_size): + # x in shape(B, H, W, C) + # cyclic shift + if shift_size > 0: + shifted_x = torch.roll(x, shifts=(-shift_size, -shift_size), dims=(1, 2)) + else: + shifted_x = x + # partition windows + x_windows = window_partition(shifted_x, window_size) + return x_windows + + +def reverse_pyt_forward(attn_windows, shift_size, window_size, H, W): + # x in shape(B*nH*nW, window_size, window_size, C) + shifted_x = window_reverse(attn_windows, window_size, H, W) + if shift_size > 0: + x = torch.roll(shifted_x, shifts=(shift_size, shift_size), dims=(1, 2)) + else: + x = shifted_x + return x + + +def copy_one_tensor(input, requires_grad=True): + input1 = input.clone().detach().requires_grad_(requires_grad).cuda() + return input1 + +class Test_WindowProcess(unittest.TestCase): + def setUp(self): + self.B = 192 + self.H = 56 + self.W = 56 + self.C = 96 + self.shift_size = 2 + self.window_size = 7 + self.nH = self.H // self.window_size + self.nW = self.W // self.window_size + + def test_roll_and_window_partition_forward(self, dtype=torch.float32): + input = torch.randn((self.B, self.H, self.W, self.C), dtype=dtype, requires_grad=True).cuda() + + input1 = copy_one_tensor(input, True) + input2 = copy_one_tensor(input, True) + + with torch.no_grad(): + # ori + expected = pyt_forward(input1, self.shift_size, self.window_size) + # fused kernel + fused_output = WindowProcess.apply(input2, self.B, self.H, self.W, self.C, -self.shift_size, self.window_size) + + self.assertTrue(torch.equal(expected, fused_output)) + #self.assertTrue(torch.allclose(expected, fused_output, rtol=1e-05, atol=1e-08)) + + def test_roll_and_window_partition_backward(self, dtype=torch.float32): + input = torch.randn((self.B, self.H, self.W, self.C), dtype=dtype, requires_grad=True).cuda() + d_loss_tensor = torch.randn((self.B*self.nW*self.nH, self.window_size, self.window_size, self.C), dtype=dtype).cuda() + + input1 = copy_one_tensor(input, True) + input2 = copy_one_tensor(input, True) + + # ori + expected = pyt_forward(input1, self.shift_size, self.window_size) + expected.backward(d_loss_tensor) + # fused kernel + fused_output = WindowProcess.apply(input2, self.B, self.H, self.W, self.C, -self.shift_size, self.window_size) + fused_output.backward(d_loss_tensor) + + self.assertTrue(torch.equal(expected, fused_output)) + #self.assertTrue(torch.allclose(expected, fused_output, rtol=1e-05, atol=1e-08)) + + def test_window_merge_and_roll_forward(self, dtype=torch.float32): + input = torch.randn((self.B*self.nH*self.nW, self.window_size, self.window_size, self.C), dtype=dtype, requires_grad=True).cuda() + + input1 = copy_one_tensor(input, True) + input2 = copy_one_tensor(input, True) + + with torch.no_grad(): + # ori + expected = reverse_pyt_forward(input1, self.shift_size, self.window_size, self.H, self.W) + # fused kernel + fused_output = WindowProcessReverse.apply(input2, self.B, self.H, self.W, self.C, self.shift_size, self.window_size) + + self.assertTrue(torch.equal(expected, fused_output)) + #self.assertTrue(torch.allclose(expected, fused_output, rtol=1e-05, atol=1e-08)) + + + def test_window_merge_and_roll_backward(self, dtype=torch.float32): + input = torch.randn((self.B*self.nH*self.nW, self.window_size, self.window_size, self.C), dtype=dtype, requires_grad=True).cuda() + d_loss_tensor = torch.randn((self.B, self.H, self.W, self.C), dtype=dtype, requires_grad=True).cuda() + + input1 = copy_one_tensor(input, True) + input2 = copy_one_tensor(input, True) + + # ori + expected = reverse_pyt_forward(input1, self.shift_size, self.window_size, self.H, self.W) + expected.backward(d_loss_tensor) + # fused kernel + fused_output = WindowProcessReverse.apply(input2, self.B, self.H, self.W, self.C, self.shift_size, self.window_size) + fused_output.backward(d_loss_tensor) + + self.assertTrue(torch.equal(expected, fused_output)) + #self.assertTrue(torch.allclose(expected, fused_output, rtol=1e-05, atol=1e-08)) + + def test_forward_backward_speed(self, dtype=torch.float32, times=1000): + input = torch.randn((self.B*self.nH*self.nW, self.window_size, self.window_size, self.C), dtype=dtype, requires_grad=True).cuda() + d_loss_tensor = torch.randn((self.B, self.H, self.W, self.C), dtype=dtype, requires_grad=True).cuda() + + input1 = copy_one_tensor(input, True) + input2 = copy_one_tensor(input, True) + + # SwinTransformer official + def run_pyt(t=1000): + for _ in range(t): + expected = reverse_pyt_forward(input1, self.shift_size, self.window_size, self.H, self.W) + expected.backward(d_loss_tensor) + + # my op + def run_fusedop(t=1000): + for _ in range(t): + fused_output = WindowProcessReverse.apply(input2, self.B, self.H, self.W, self.C, self.shift_size, self.window_size) + fused_output.backward(d_loss_tensor) + + torch.cuda.synchronize() + t1 = time.time() + run_pyt(t=times) + torch.cuda.synchronize() + t2 = time.time() + run_fusedop(t=times) + torch.cuda.synchronize() + t3 = time.time() + self.assertTrue((t3 - t2) < (t2 - t1)) + + print('Run {} times'.format(times)) + print('Original time cost: {}'.format(t2 - t1)) + print('Fused op time cost: {}'.format(t3 - t2)) + + def test_roll_and_window_partition_forward_fp16(self, dtype=torch.float16): + self.test_roll_and_window_partition_forward(dtype=dtype) + + def test_roll_and_window_partition_backward_fp16(self, dtype=torch.float16): + self.test_roll_and_window_partition_backward(dtype=dtype) + + def test_window_merge_and_roll_forward_fp16(self, dtype=torch.float16): + self.test_window_merge_and_roll_forward(dtype=dtype) + + def test_window_merge_and_roll_backward_fp16(self, dtype=torch.float16): + self.test_window_merge_and_roll_backward(dtype=dtype) + + def test_forward_backward_speed_fp16(self, dtype=torch.float16, times=1000): + self.test_forward_backward_speed(dtype=dtype, times=times) + + +if __name__ == '__main__': + print('Pass only two tensors are exactly the same (using torch.equal).\n') + torch.manual_seed(0) + unittest.main(verbosity=2) diff --git a/swin/kernels/window_process/window_process.py b/swin/kernels/window_process/window_process.py new file mode 100644 index 0000000..ee43e9e --- /dev/null +++ b/swin/kernels/window_process/window_process.py @@ -0,0 +1,63 @@ +# -------------------------------------------------------- +# Fused kernel for window process for SwinTransformer +# Copyright (c) 2022 Nvidia +# Licensed under The MIT License [see LICENSE for details] +# -------------------------------------------------------- + +import torch +import swin_window_process + + +class WindowProcess(torch.autograd.Function): + @staticmethod + def forward(ctx, input, B, H, W, C, shift_size, window_size): + output = swin_window_process.roll_and_window_partition_forward(input, B, H, W, C, shift_size, window_size) + + ctx.B = B + ctx.H = H + ctx.W = W + ctx.C = C + ctx.shift_size = shift_size + ctx.window_size = window_size + return output + + @staticmethod + def backward(ctx, grad_in): + B = ctx.B + H = ctx.H + W = ctx.W + C = ctx.C + shift_size = ctx.shift_size + window_size = ctx.window_size + + grad_out = swin_window_process.roll_and_window_partition_backward(grad_in, B, H, W, C, shift_size, window_size) + return grad_out, None, None, None, None, None, None, None + + +class WindowProcessReverse(torch.autograd.Function): + @staticmethod + def forward(ctx, input, B, H, W, C, shift_size, window_size): + output = swin_window_process.window_merge_and_roll_forward(input, B, H, W, C, shift_size, window_size) + + ctx.B = B + ctx.H = H + ctx.W = W + ctx.C = C + ctx.shift_size = shift_size + ctx.window_size = window_size + + return output + + @staticmethod + def backward(ctx, grad_in): + B = ctx.B + H = ctx.H + W = ctx.W + C = ctx.C + shift_size = ctx.shift_size + window_size = ctx.window_size + + #grad_out = ctx.saved_tensors[0] + #grad_out = torch.zeros((B, H, W, C), dtype=dtype).cuda() + grad_out = swin_window_process.window_merge_and_roll_backward(grad_in, B, H, W, C, shift_size, window_size) + return grad_out, None, None, None, None, None, None, None diff --git a/swin/models/__init__.py b/swin/models/__init__.py new file mode 100644 index 0000000..2d9c65e --- /dev/null +++ b/swin/models/__init__.py @@ -0,0 +1 @@ +from .build import build_model \ No newline at end of file diff --git a/swin/models/build.py b/swin/models/build.py new file mode 100644 index 0000000..ac7459b --- /dev/null +++ b/swin/models/build.py @@ -0,0 +1,24 @@ +from .swin_transformer import SwinTransformer +import torch.nn as nn +layernorm = nn.LayerNorm + + +def build_model(config): + return SwinTransformer(img_size=config.DATA.IMG_SIZE, + patch_size=config.MODEL.SWIN.PATCH_SIZE, + in_chans=config.MODEL.SWIN.IN_CHANS, + num_classes=config.MODEL.NUM_CLASSES, + embed_dim=config.MODEL.SWIN.EMBED_DIM, + depths=config.MODEL.SWIN.DEPTHS, + num_heads=config.MODEL.SWIN.NUM_HEADS, + window_size=config.MODEL.SWIN.WINDOW_SIZE, + mlp_ratio=config.MODEL.SWIN.MLP_RATIO, + qkv_bias=config.MODEL.SWIN.QKV_BIAS, + qk_scale=config.MODEL.SWIN.QK_SCALE, + drop_rate=config.MODEL.DROP_RATE, + drop_path_rate=config.MODEL.DROP_PATH_RATE, + ape=config.MODEL.SWIN.APE, + norm_layer=layernorm, + patch_norm=config.MODEL.SWIN.PATCH_NORM, + use_checkpoint=config.TRAIN.USE_CHECKPOINT, + fused_window_process=config.FUSED_WINDOW_PROCESS) \ No newline at end of file diff --git a/swin/models/swin_transformer.py b/swin/models/swin_transformer.py new file mode 100644 index 0000000..91cbe59 --- /dev/null +++ b/swin/models/swin_transformer.py @@ -0,0 +1,629 @@ +# -------------------------------------------------------- +# Swin Transformer +# Copyright (c) 2021 Microsoft +# Licensed under The MIT License [see LICENSE for details] +# Written by Ze Liu +# -------------------------------------------------------- + +import torch +import torch.nn as nn +import torch.utils.checkpoint as checkpoint +from timm.models.layers import DropPath, to_2tuple, trunc_normal_ + +try: + import os, sys + + kernel_path = os.path.abspath(os.path.join('..')) + sys.path.append(kernel_path) + from kernels.window_process.window_process import WindowProcess, WindowProcessReverse + +except: + WindowProcess = None + WindowProcessReverse = None + print("[Warning] Fused window process have not been installed. Please refer to get_started.md for installation.") + + +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 + qk_scale (float | None, optional): Override default qk scale of head_dim ** -0.5 if set + attn_drop (float, optional): Dropout ratio of attention weight. Default: 0.0 + proj_drop (float, optional): Dropout ratio of output. Default: 0.0 + """ + + def __init__(self, dim, window_size, num_heads, qkv_bias=True, qk_scale=None, attn_drop=0., proj_drop=0.): + + super().__init__() + self.dim = dim + self.window_size = window_size # Wh, Ww + self.num_heads = num_heads + head_dim = dim // num_heads + self.scale = qk_scale or head_dim ** -0.5 + + # define a parameter table of relative position bias + self.relative_position_bias_table = nn.Parameter( + torch.zeros((2 * window_size[0] - 1) * (2 * window_size[1] - 1), num_heads)) # 2*Wh-1 * 2*Ww-1, nH + + # 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=qkv_bias) + self.attn_drop = nn.Dropout(attn_drop) + self.proj = nn.Linear(dim, dim) + self.proj_drop = nn.Dropout(proj_drop) + + trunc_normal_(self.relative_position_bias_table, std=.02) + 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 = self.qkv(x).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] # make torchscript happy (cannot use tensor as tuple) + + q = q * self.scale + attn = (q @ k.transpose(-2, -1)) + + relative_position_bias = self.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 + 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 + + def extra_repr(self) -> str: + return f'dim={self.dim}, window_size={self.window_size}, num_heads={self.num_heads}' + + def flops(self, N): + # calculate flops for 1 window with token length of N + flops = 0 + # qkv = self.qkv(x) + flops += N * self.dim * 3 * self.dim + # attn = (q @ k.transpose(-2, -1)) + flops += self.num_heads * N * (self.dim // self.num_heads) * N + # x = (attn @ v) + flops += self.num_heads * N * N * (self.dim // self.num_heads) + # x = self.proj(x) + flops += N * self.dim * self.dim + return flops + + +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 + qk_scale (float | None, optional): Override default qk scale of head_dim ** -0.5 if set. + 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 + fused_window_process (bool, optional): If True, use one kernel to fused window shift & window partition for acceleration, similar for the reversed part. Default: False + """ + + def __init__(self, dim, input_resolution, num_heads, window_size=7, shift_size=0, + mlp_ratio=4., qkv_bias=True, qk_scale=None, drop=0., attn_drop=0., drop_path=0., + act_layer=nn.GELU, norm_layer=nn.LayerNorm, + fused_window_process=False): + 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, qk_scale=qk_scale, attn_drop=attn_drop, proj_drop=drop) + + self.drop_path = DropPath(drop_path) if drop_path > 0. else 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) + self.fused_window_process = fused_window_process + + 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 = self.norm1(x) + x = x.view(B, H, W, C) + + # cyclic shift + if self.shift_size > 0: + if not self.fused_window_process: + shifted_x = torch.roll(x, shifts=(-self.shift_size, -self.shift_size), dims=(1, 2)) + # partition windows + x_windows = window_partition(shifted_x, self.window_size) # nW*B, window_size, window_size, C + else: + x_windows = WindowProcess.apply(x, B, H, W, C, -self.shift_size, self.window_size) + 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) + + # reverse cyclic shift + if self.shift_size > 0: + if not self.fused_window_process: + shifted_x = window_reverse(attn_windows, self.window_size, H, W) # B H' W' C + x = torch.roll(shifted_x, shifts=(self.shift_size, self.shift_size), dims=(1, 2)) + else: + x = WindowProcessReverse.apply(attn_windows, B, H, W, C, self.shift_size, self.window_size) + else: + shifted_x = window_reverse(attn_windows, self.window_size, H, W) # B H' W' C + x = shifted_x + x = x.view(B, H * W, C) + x = shortcut + self.drop_path(x) + + # FFN + x = x + self.drop_path(self.mlp(self.norm2(x))) + + return x + + def extra_repr(self) -> str: + return f"dim={self.dim}, input_resolution={self.input_resolution}, num_heads={self.num_heads}, " \ + f"window_size={self.window_size}, shift_size={self.shift_size}, mlp_ratio={self.mlp_ratio}" + + def flops(self): + flops = 0 + H, W = self.input_resolution + # norm1 + flops += self.dim * H * W + # W-MSA/SW-MSA + nW = H * W / self.window_size / self.window_size + flops += nW * self.attn.flops(self.window_size * self.window_size) + # mlp + flops += 2 * H * W * self.dim * self.dim * self.mlp_ratio + # norm2 + flops += self.dim * H * W + return flops + + +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(4 * 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.norm(x) + x = self.reduction(x) + + return x + + def extra_repr(self) -> str: + return f"input_resolution={self.input_resolution}, dim={self.dim}" + + def flops(self): + H, W = self.input_resolution + flops = H * W * self.dim + flops += (H // 2) * (W // 2) * 4 * self.dim * 2 * self.dim + return flops + + +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 + qk_scale (float | None, optional): Override default qk scale of head_dim ** -0.5 if set. + 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. + fused_window_process (bool, optional): If True, use one kernel to fused window shift & window partition for acceleration, similar for the reversed part. Default: False + """ + + def __init__(self, dim, input_resolution, depth, num_heads, window_size, + mlp_ratio=4., qkv_bias=True, qk_scale=None, drop=0., attn_drop=0., + drop_path=0., norm_layer=nn.LayerNorm, downsample=None, use_checkpoint=False, + fused_window_process=False): + + 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, qk_scale=qk_scale, + drop=drop, attn_drop=attn_drop, + drop_path=drop_path[i] if isinstance(drop_path, list) else drop_path, + norm_layer=norm_layer, + fused_window_process=fused_window_process) + 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 + + def extra_repr(self) -> str: + return f"dim={self.dim}, input_resolution={self.input_resolution}, depth={self.depth}" + + def flops(self): + flops = 0 + for blk in self.blocks: + flops += blk.flops() + if self.downsample is not None: + flops += self.downsample.flops() + return flops + +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 + + def flops(self): + Ho, Wo = self.patches_resolution + flops = Ho * Wo * self.embed_dim * self.in_chans * (self.patch_size[0] * self.patch_size[1]) + if self.norm is not None: + flops += Ho * Wo * self.embed_dim + return flops + + +class SwinTransformer(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 + qk_scale (float): Override default qk scale of head_dim ** -0.5 if set. Default: None + 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. + ape (bool): If True, add absolute position embedding to the patch embedding. Default: False + patch_norm (bool): If True, add normalization after patch embedding. Default: True + use_checkpoint (bool): Whether to use checkpointing to save memory. Default: False + fused_window_process (bool, optional): If True, use one kernel to fused window shift & window partition for acceleration, similar for the reversed part. Default: False + """ + + 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, qk_scale=None, + drop_rate=0., attn_drop_rate=0., drop_path_rate=0.1, + norm_layer=nn.LayerNorm, ape=False, patch_norm=True, + use_checkpoint=False, fused_window_process=False, **kwargs): + super().__init__() + + self.num_classes = num_classes + self.num_layers = len(depths) + self.embed_dim = embed_dim + self.ape = ape + self.patch_norm = patch_norm + self.num_features = int(embed_dim * 2 ** (self.num_layers - 1)) + self.mlp_ratio = mlp_ratio + + self.img_size = img_size + self.patch_size = patch_size + + # 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 + + # absolute position embedding + if self.ape: + self.absolute_pos_embed = nn.Parameter(torch.zeros(1, num_patches, embed_dim)) + trunc_normal_(self.absolute_pos_embed, std=.02) + + 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, qk_scale=qk_scale, + 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, + fused_window_process=fused_window_process) + 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() + + self.apply(self._init_weights) + + def _init_weights(self, m): + if isinstance(m, nn.Linear): + trunc_normal_(m.weight, std=.02) + if isinstance(m, nn.Linear) and m.bias is not None: + nn.init.constant_(m.bias, 0) + elif isinstance(m, nn.LayerNorm): + nn.init.constant_(m.bias, 0) + nn.init.constant_(m.weight, 1.0) + + @torch.jit.ignore + def no_weight_decay(self): + return {'absolute_pos_embed'} + + @torch.jit.ignore + def no_weight_decay_keywords(self): + return {'relative_position_bias_table'} + + def quantize_coords(self, coords): + grid_size = self.img_size // self.patch_size + return ((coords[...,0]/self.patch_size).round()*grid_size + (coords[..., 1]/self.patch_size).round()).to(dtype=int).clip(0, grid_size**2-1) + + def forward_features(self, x, coords=None, keeps=None): + assert keeps is None, "Swin transformer works only on merged patches with soft dropout. Set args --hard_dropout=0, --merge_patches=1 to proceed." + x = self.patch_embed(x) + if self.ape: + assert False, "Although implemented, the function is not tested because absolute patch embedding in swin is turned off by default." + if coords is not None: + quantized = self.quantize_coords(coords) + x = x + self.absolute_pos_embed[[0] ,quantized, :] + else: + x = x + self.absolute_pos_embed + # ========================================== + + + 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, coords=None, keeps=None): + x = self.forward_features(x, coords, keeps) + x = self.head(x) + return x + + def flops(self): + flops = 0 + flops += self.patch_embed.flops() + for i, layer in enumerate(self.layers): + flops += layer.flops() + flops += self.num_features * self.patches_resolution[0] * self.patches_resolution[1] // (2 ** self.num_layers) + flops += self.num_features * self.num_classes + return flops diff --git a/swin/swin_base_patch4_window7_224.yaml b/swin/swin_base_patch4_window7_224.yaml new file mode 100644 index 0000000..b296128 --- /dev/null +++ b/swin/swin_base_patch4_window7_224.yaml @@ -0,0 +1,9 @@ +MODEL: + TYPE: swin + NAME: swin_base_patch4_window7_224 + DROP_PATH_RATE: 0.5 + SWIN: + EMBED_DIM: 128 + DEPTHS: [ 2, 2, 18, 2 ] + NUM_HEADS: [ 4, 8, 16, 32 ] + WINDOW_SIZE: 7 \ No newline at end of file