gpt-5.5-priority · original · l123_single_gpt
compile_fail 29_SwinMLP.py
model
gpt-5.5-priority
turns / max
1 / 1
tool calls
1
tokens (total)
0
wall clock
—
truncated
no
backend
cuda fp32
variant
original
runtime
—
ref runtime
-1.00 μs
speedup
—
started
2026-05-05T02:53:43Z
compile error: /lib/x86_64-linux-gnu/libstdc++.so.6: version `CXXABI_1.3.15' not found (required by /scratch/abaner/KernelBench/runs/l123_single_gpt/original/gpt-5.5-priority/level_3_problem_29_cache/k_a4849eb17f94/swin_mlp_spatial_ext/swin_mlp_spatial_ext.so)
timingkernelbench.timing
kernel runtime
—
ref runtime
—
speedup
—
kernel mean / median / std
—
ref mean / median / std
—
reference PyTorch source (the prompt the agent sees)
# --------------------------------------------------------
# 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.nn.functional as F
from itertools import repeat
import collections.abc

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 SwinMLPBlock(nn.Module):
    r""" Swin MLP 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.
        drop (float, optional): 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
    """

    def __init__(self, dim, input_resolution, num_heads, window_size=7, shift_size=0,
                 mlp_ratio=4., drop=0., drop_path=0.,
                 act_layer=nn.GELU, norm_layer=nn.LayerNorm):
        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.padding = [self.window_size - self.shift_size, self.shift_size,
                        self.window_size - self.shift_size, self.shift_size]  # P_l,P_r,P_t,P_b

        self.norm1 = norm_layer(dim)
        # use group convolution to implement multi-head MLP
        self.spatial_mlp = nn.Conv1d(self.num_heads * self.window_size ** 2,
                                     self.num_heads * self.window_size ** 2,
                                     kernel_size=1,
                                     groups=self.num_heads)

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

    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)

        # shift
        if self.shift_size > 0:
            P_l, P_r, P_t, P_b = self.padding
            shifted_x = F.pad(x, [0, 0, P_l, P_r, P_t, P_b], "constant", 0)
        else:
            shifted_x = x
        _, _H, _W, _ = shifted_x.shape

        # 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

        # Window/Shifted-Window Spatial MLP
        x_windows_heads = x_windows.view(-1, self.window_size * self.window_size, self.num_heads, C // self.num_heads)
        x_windows_heads = x_windows_heads.transpose(1, 2)  # nW*B, nH, window_size*window_size, C//nH
        x_windows_heads = x_windows_heads.reshape(-1, self.num_heads * self.window_size * self.window_size,
                                                  C // self.num_heads)
        spatial_mlp_windows = self.spatial_mlp(x_windows_heads)  # nW*B, nH*window_size*window_size, C//nH
        spatial_mlp_windows = spatial_mlp_windows.view(-1, self.num_heads, self.window_size * self.window_size,
                                                       C // self.num_heads).transpose(1, 2)
        spatial_mlp_windows = spatial_mlp_windows.reshape(-1, self.window_size * self.window_size, C)

        # merge windows
        spatial_mlp_windows = spatial_mlp_windows.reshape(-1, self.window_size, self.window_size, C)
        shifted_x = window_reverse(spatial_mlp_windows, self.window_size, _H, _W)  # B H' W' C

        # reverse shift
        if self.shift_size > 0:
            P_l, P_r, P_t, P_b = self.padding
            x = shifted_x[:, P_t:-P_b, P_l:-P_r, :].contiguous()
        else:
            x = shifted_x
        x = x.view(B, H * W, C)

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

        return x


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

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

    def __init__(self, input_resolution, dim, norm_layer=nn.LayerNorm):
        super().__init__()
        self.input_resolution = input_resolution
        self.dim = dim
        self.reduction = nn.Linear(4 * dim, 2 * dim, bias=False)
        self.norm = norm_layer(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

class BasicLayer(nn.Module):
    """ A basic Swin MLP 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.
        drop (float, optional): 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.
    """

    def __init__(self, dim, input_resolution, depth, num_heads, window_size,
                 mlp_ratio=4., drop=0., drop_path=0.,
                 norm_layer=nn.LayerNorm, downsample=None, use_checkpoint=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([
            SwinMLPBlock(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,
                         drop=drop,
                         drop_path=drop_path[i] if isinstance(drop_path, list) else drop_path,
                         norm_layer=norm_layer)
            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 _ntuple(n):
    def parse(x):
        if isinstance(x, collections.abc.Iterable) and not isinstance(x, str):
            return tuple(x)
        return tuple(repeat(x, n))
    return parse
to_2tuple = _ntuple(2)

class 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 Model(nn.Module):
    r""" Swin MLP

    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 MLP 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
        drop_rate (float): Dropout rate. Default: 0
        drop_path_rate (float): Stochastic depth rate. Default: 0.1
        norm_layer (nn.Module): Normalization layer. Default: nn.LayerNorm.
        patch_norm (bool): If True, add normalization after patch embedding. Default: True
        use_checkpoint (bool): Whether to use checkpointing to save memory. Default: False
    """

    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., drop_rate=0., drop_path_rate=0.1,
                 norm_layer=nn.LayerNorm, patch_norm=True,
                 use_checkpoint=False, **kwargs):
        super().__init__()

        self.num_classes = num_classes
        self.num_layers = len(depths)
        self.embed_dim = embed_dim
        self.patch_norm = patch_norm
        self.num_features = int(embed_dim * 2 ** (self.num_layers - 1))
        self.mlp_ratio = mlp_ratio

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

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

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

        # build layers
        self.layers = nn.ModuleList()
        for i_layer in range(self.num_layers):
            layer = BasicLayer(dim=int(embed_dim * 2 ** i_layer),
                               input_resolution=(patches_resolution[0] // (2 ** i_layer),
                                                 patches_resolution[1] // (2 ** i_layer)),
                               depth=depths[i_layer],
                               num_heads=num_heads[i_layer],
                               window_size=window_size,
                               mlp_ratio=self.mlp_ratio,
                               drop=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)
            self.layers.append(layer)

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

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

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

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

    def forward(self, x):
        x = self.forward_features(x)
        x = self.head(x)
        return x
    
batch_size = 10
image_size = 224

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

def get_init_inputs():
    return []
turn 0 · FINAL · 1 tool call(s) · 52.4s turn · 52.4s LLM · 12,149 chars (click to expand)
turn 052.4s turn · 52.4s LLM · 0.0s tools · FINAL
call → submit_kernel
callsubmit_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nimport torch.nn.functional as F\nfrom torch.utils.cpp_extension import load_inline\nfrom itertools import repeat\nimport collections.abc\n\n_cpp_src = r'''\n#include <torch/extension.h>\n\ntorch::Tensor spatial_mlp_forward_cuda(torch::Tensor x, torch::Tensor weight, torch::Tensor bias, int H, int W, int window_size, int shift_size, int num_heads);\n\ntorch::Tensor spatial_mlp_forward(torch::Tensor x, torch::Tensor weight, torch::Tensor bias, int H, int W, int window_size, int shift_size, int num_heads) {\n return spatial_mlp_forward_cuda(x, weight, bias, H, W, window_size, shift_size, num_heads);\n}\n'''\n\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\n__global__ void spatial_mlp_kernel(const float* __restrict__ x,\n const float* __restrict__ w,\n const float* __restrict__ b,\n float* __restrict__ y,\n int B, int H, int W, int C,\n int ws, int shift, int heads, int chd,\n long total) {\n long idx = (long)blockIdx.x * blockDim.x + threadIdx.x;\n if (idx >= total) return;\n int c = idx % C;\n long t = idx / C;\n int ww = t % W;\n t /= W;\n int hh = t % H;\n int bb = t / H;\n\n int pad_l, pad_t, Hp, Wp;\n if (shift > 0) {\n pad_l = ws - shift;\n pad_t = ws - shift;\n Hp = H + ws;\n Wp = W + ws;\n } else {\n pad_l = 0;\n pad_t = 0;\n Hp = H;\n Wp = W;\n }\n\n int ph = hh + pad_t;\n int pw = ww + pad_l;\n int local_h = ph % ws;\n int local_w = pw % ws;\n int p = local_h * ws + local_w;\n int base_h = ph - local_h;\n int base_w = pw - local_w;\n int hidx = c / chd;\n const float* wp = w + (hidx * ws * ws + p) * (ws * ws);\n float acc = b[hidx * ws * ws + p];\n\n #pragma unroll\n for (int q = 0; q < 64; ++q) {\n if (q >= ws * ws) break;\n int qh = q / ws;\n int qw = q - qh * ws;\n int sph = base_h + qh;\n int spw = base_w + qw;\n int oh = sph - pad_t;\n int ow = spw - pad_l;\n float xv = 0.0f;\n if ((unsigned)oh < (unsigned)H && (unsigned)ow < (unsigned)W) {\n xv = x[((bb * H + oh) * W + ow) * C + c];\n }\n acc = fmaf(wp[q], xv, acc);\n }\n y[idx] = acc;\n}\n\ntorch::Tensor spatial_mlp_forward_cuda(torch::Tensor x, torch::Tensor weight, torch::Tensor bias, int H, int W, int window_size, int shift_size, int num_heads) {\n auto y = torch::empty_like(x);\n int B = x.size(0);\n int C = x.size(2);\n long total = (long)B * H * W * C;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n int chd = C / num_heads;\n spatial_mlp_kernel<<<blocks, threads, 0, at::cuda::getCurrentCUDAStream()>>>(\n x.data_ptr<float>(), weight.data_ptr<float>(), bias.data_ptr<float>(), y.data_ptr<float>(),\n B, H, W, C, window_size, shift_size, num_heads, chd, total);\n return y;\n}\n'''\n\n_ext = load_inline(name='swin_mlp_spatial_ext', cpp_sources=_cpp_src, cuda_sources=_cuda_src,\n functions=['spatial_mlp_forward'], extra_cuda_cflags=['-O3'], verbose=False)\n\ndef _ntuple(n):\n def parse(x):\n if isinstance(x, collections.abc.Iterable) and not isinstance(x, str):\n return tuple(x)\n return tuple(repeat(x, n))\n return parse\nto_2tuple = _ntuple(2)\n\nclass Mlp(nn.Module):\n def __init__(self, in_features, hidden_features=None, out_features=None, act_layer=nn.GELU, drop=0.):\n super().__init__()\n out_features = out_features or in_features\n hidden_features = hidden_features or in_features\n self.fc1 = nn.Linear(in_features, hidden_features)\n self.act = act_layer()\n self.fc2 = nn.Linear(hidden_features, out_features)\n self.drop = nn.Dropout(drop)\n def forward(self, x):\n x = self.fc1(x)\n x = self.act(x)\n x = self.drop(x)\n x = self.fc2(x)\n x = self.drop(x)\n return x\n\nclass SwinMLPBlock(nn.Module):\n def __init__(self, dim, input_resolution, num_heads, window_size=7, shift_size=0,\n mlp_ratio=4., drop=0., drop_path=0., act_layer=nn.GELU, norm_layer=nn.LayerNorm):\n super().__init__()\n self.dim = dim; self.input_resolution = input_resolution; self.num_heads = num_heads\n self.window_size = window_size; self.shift_size = shift_size; self.mlp_ratio = mlp_ratio\n if min(self.input_resolution) <= self.window_size:\n self.shift_size = 0; self.window_size = min(self.input_resolution)\n self.padding = [self.window_size - self.shift_size, self.shift_size,\n self.window_size - self.shift_size, self.shift_size]\n self.norm1 = norm_layer(dim)\n self.spatial_mlp = nn.Conv1d(self.num_heads * self.window_size ** 2,\n self.num_heads * self.window_size ** 2,\n kernel_size=1, groups=self.num_heads)\n self.drop_path = nn.Identity()\n self.norm2 = norm_layer(dim)\n mlp_hidden_dim = int(dim * mlp_ratio)\n self.mlp = Mlp(in_features=dim, hidden_features=mlp_hidden_dim, act_layer=act_layer, drop=drop)\n def forward(self, x):\n H, W = self.input_resolution\n shortcut = x\n xn = self.norm1(x).contiguous()\n y = _ext.spatial_mlp_forward(xn, self.spatial_mlp.weight, self.spatial_mlp.bias,\n H, W, self.window_size, self.shift_size, self.num_heads)\n x = shortcut + y\n x = x + self.mlp(self.norm2(x))\n return x\n\nclass PatchMerging(nn.Module):\n def __init__(self, input_resolution, dim, norm_layer=nn.LayerNorm):\n super().__init__()\n self.input_resolution = input_resolution; self.dim = dim\n self.reduction = nn.Linear(4 * dim, 2 * dim, bias=False)\n self.norm = norm_layer(4 * dim)\n def forward(self, x):\n H, W = self.input_resolution\n B, L, C = x.shape\n x = x.view(B, H, W, C)\n x0 = x[:, 0::2, 0::2, :]\n x1 = x[:, 1::2, 0::2, :]\n x2 = x[:, 0::2, 1::2, :]\n x3 = x[:, 1::2, 1::2, :]\n x = torch.cat([x0, x1, x2, x3], -1).view(B, -1, 4 * C)\n x = self.norm(x)\n x = self.reduction(x)\n return x\n\nclass BasicLayer(nn.Module):\n def __init__(self, dim, input_resolution, depth, num_heads, window_size,\n mlp_ratio=4., drop=0., drop_path=0., norm_layer=nn.LayerNorm,\n downsample=None, use_checkpoint=False):\n super().__init__()\n self.dim = dim; self.input_resolution = input_resolution; self.depth = depth; self.use_checkpoint = use_checkpoint\n self.blocks = nn.ModuleList([\n SwinMLPBlock(dim=dim, input_resolution=input_resolution, num_heads=num_heads, window_size=window_size,\n shift_size=0 if (i % 2 == 0) else window_size // 2, mlp_ratio=mlp_ratio,\n drop=drop, drop_path=drop_path[i] if isinstance(drop_path, list) else drop_path,\n norm_layer=norm_layer) for i in range(depth)])\n self.downsample = downsample(input_resolution, dim=dim, norm_layer=norm_layer) if downsample is not None else None\n def forward(self, x):\n for blk in self.blocks:\n x = blk(x)\n if self.downsample is not None:\n x = self.downsample(x)\n return x\n\nclass PatchEmbed(nn.Module):\n def __init__(self, img_size=224, patch_size=4, in_chans=3, embed_dim=96, norm_layer=None):\n super().__init__()\n img_size = to_2tuple(img_size); patch_size = to_2tuple(patch_size)\n patches_resolution = [img_size[0] // patch_size[0], img_size[1] // patch_size[1]]\n self.img_size = img_size; self.patch_size = patch_size; self.patches_resolution = patches_resolution\n self.num_patches = patches_resolution[0] * patches_resolution[1]\n self.in_chans = in_chans; self.embed_dim = embed_dim\n self.proj = nn.Conv2d(in_chans, embed_dim, kernel_size=patch_size, stride=patch_size)\n self.norm = norm_layer(embed_dim) if norm_layer is not None else None\n def forward(self, x):\n x = self.proj(x).flatten(2).transpose(1, 2)\n if self.norm is not None:\n x = self.norm(x)\n return x\n\nclass ModelNew(nn.Module):\n def __init__(self, img_size=224, patch_size=4, in_chans=3, num_classes=1000,\n embed_dim=96, depths=[2, 2, 6, 2], num_heads=[3, 6, 12, 24],\n window_size=7, mlp_ratio=4., drop_rate=0., drop_path_rate=0.1,\n norm_layer=nn.LayerNorm, patch_norm=True, use_checkpoint=False, **kwargs):\n super().__init__()\n self.num_classes = num_classes; self.num_layers = len(depths); self.embed_dim = embed_dim\n self.patch_norm = patch_norm; self.num_features = int(embed_dim * 2 ** (self.num_layers - 1)); self.mlp_ratio = mlp_ratio\n self.patch_embed = PatchEmbed(img_size=img_size, patch_size=patch_size, in_chans=in_chans, embed_dim=embed_dim,\n norm_layer=norm_layer if self.patch_norm else None)\n patches_resolution = self.patch_embed.patches_resolution\n self.patches_resolution = patches_resolution\n self.pos_drop = nn.Dropout(p=drop_rate)\n dpr = [x.item() for x in torch.linspace(0, drop_path_rate, sum(depths))]\n self.layers = nn.ModuleList()\n for i_layer in range(self.num_layers):\n layer = BasicLayer(dim=int(embed_dim * 2 ** i_layer),\n input_resolution=(patches_resolution[0] // (2 ** i_layer), patches_resolution[1] // (2 ** i_layer)),\n depth=depths[i_layer], num_heads=num_heads[i_layer], window_size=window_size,\n mlp_ratio=self.mlp_ratio, drop=drop_rate,\n drop_path=dpr[sum(depths[:i_layer]):sum(depths[:i_layer + 1])],\n norm_layer=norm_layer,\n downsample=PatchMerging if (i_layer < self.num_layers - 1) else None,\n use_checkpoint=use_checkpoint)\n self.layers.append(layer)\n self.norm = norm_layer(self.num_features)\n self.avgpool = nn.AdaptiveAvgPool1d(1)\n self.head = nn.Linear(self.num_features, num_classes) if num_classes > 0 else nn.Identity()\n def forward_features(self, x):\n x = self.patch_embed(x)\n x = self.pos_drop(x)\n for layer in self.layers:\n x = layer(x)\n x = self.norm(x)\n x = self.avgpool(x.transpose(1, 2))\n x = torch.flatten(x, 1)\n return x\n def forward(self, x):\n x = self.forward_features(x)\n x = self.head(x)\n return x\n" }
result ← submit_kernel (fail)
failsubmit_kernel submit_kernel FAILED: kernel did not compile. /lib/x86_64-linux-gnu/libstdc++.so.6: version `CXXABI_1.3.15' not found (required by /scratch/abaner/KernelBench/runs/l123_single_gpt/original/gpt-5.5-priority/level_3_problem_29_cache/k_a4849eb17f94/swin_mlp_spatial_ext/swin_mlp_spatial_ext.so)
final submitted kernel (248 lines, 10,901 chars) — full source, untruncated
import torch
import torch.nn as nn
import torch.nn.functional as F
from torch.utils.cpp_extension import load_inline
from itertools import repeat
import collections.abc

_cpp_src = r'''
#include <torch/extension.h>

torch::Tensor spatial_mlp_forward_cuda(torch::Tensor x, torch::Tensor weight, torch::Tensor bias, int H, int W, int window_size, int shift_size, int num_heads);

torch::Tensor spatial_mlp_forward(torch::Tensor x, torch::Tensor weight, torch::Tensor bias, int H, int W, int window_size, int shift_size, int num_heads) {
    return spatial_mlp_forward_cuda(x, weight, bias, H, W, window_size, shift_size, num_heads);
}
'''

_cuda_src = r'''
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda.h>
#include <cuda_runtime.h>

__global__ void spatial_mlp_kernel(const float* __restrict__ x,
                                   const float* __restrict__ w,
                                   const float* __restrict__ b,
                                   float* __restrict__ y,
                                   int B, int H, int W, int C,
                                   int ws, int shift, int heads, int chd,
                                   long total) {
    long idx = (long)blockIdx.x * blockDim.x + threadIdx.x;
    if (idx >= total) return;
    int c = idx % C;
    long t = idx / C;
    int ww = t % W;
    t /= W;
    int hh = t % H;
    int bb = t / H;

    int pad_l, pad_t, Hp, Wp;
    if (shift > 0) {
        pad_l = ws - shift;
        pad_t = ws - shift;
        Hp = H + ws;
        Wp = W + ws;
    } else {
        pad_l = 0;
        pad_t = 0;
        Hp = H;
        Wp = W;
    }

    int ph = hh + pad_t;
    int pw = ww + pad_l;
    int local_h = ph % ws;
    int local_w = pw % ws;
    int p = local_h * ws + local_w;
    int base_h = ph - local_h;
    int base_w = pw - local_w;
    int hidx = c / chd;
    const float* wp = w + (hidx * ws * ws + p) * (ws * ws);
    float acc = b[hidx * ws * ws + p];

    #pragma unroll
    for (int q = 0; q < 64; ++q) {
        if (q >= ws * ws) break;
        int qh = q / ws;
        int qw = q - qh * ws;
        int sph = base_h + qh;
        int spw = base_w + qw;
        int oh = sph - pad_t;
        int ow = spw - pad_l;
        float xv = 0.0f;
        if ((unsigned)oh < (unsigned)H && (unsigned)ow < (unsigned)W) {
            xv = x[((bb * H + oh) * W + ow) * C + c];
        }
        acc = fmaf(wp[q], xv, acc);
    }
    y[idx] = acc;
}

torch::Tensor spatial_mlp_forward_cuda(torch::Tensor x, torch::Tensor weight, torch::Tensor bias, int H, int W, int window_size, int shift_size, int num_heads) {
    auto y = torch::empty_like(x);
    int B = x.size(0);
    int C = x.size(2);
    long total = (long)B * H * W * C;
    int threads = 256;
    int blocks = (total + threads - 1) / threads;
    int chd = C / num_heads;
    spatial_mlp_kernel<<<blocks, threads, 0, at::cuda::getCurrentCUDAStream()>>>(
        x.data_ptr<float>(), weight.data_ptr<float>(), bias.data_ptr<float>(), y.data_ptr<float>(),
        B, H, W, C, window_size, shift_size, num_heads, chd, total);
    return y;
}
'''

_ext = load_inline(name='swin_mlp_spatial_ext', cpp_sources=_cpp_src, cuda_sources=_cuda_src,
                   functions=['spatial_mlp_forward'], extra_cuda_cflags=['-O3'], verbose=False)

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

class Mlp(nn.Module):
    def __init__(self, in_features, hidden_features=None, out_features=None, act_layer=nn.GELU, drop=0.):
        super().__init__()
        out_features = out_features or in_features
        hidden_features = hidden_features or in_features
        self.fc1 = nn.Linear(in_features, hidden_features)
        self.act = act_layer()
        self.fc2 = nn.Linear(hidden_features, out_features)
        self.drop = nn.Dropout(drop)
    def forward(self, x):
        x = self.fc1(x)
        x = self.act(x)
        x = self.drop(x)
        x = self.fc2(x)
        x = self.drop(x)
        return x

class SwinMLPBlock(nn.Module):
    def __init__(self, dim, input_resolution, num_heads, window_size=7, shift_size=0,
                 mlp_ratio=4., drop=0., drop_path=0., act_layer=nn.GELU, norm_layer=nn.LayerNorm):
        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:
            self.shift_size = 0; self.window_size = min(self.input_resolution)
        self.padding = [self.window_size - self.shift_size, self.shift_size,
                        self.window_size - self.shift_size, self.shift_size]
        self.norm1 = norm_layer(dim)
        self.spatial_mlp = nn.Conv1d(self.num_heads * self.window_size ** 2,
                                     self.num_heads * self.window_size ** 2,
                                     kernel_size=1, groups=self.num_heads)
        self.drop_path = nn.Identity()
        self.norm2 = norm_layer(dim)
        mlp_hidden_dim = int(dim * mlp_ratio)
        self.mlp = Mlp(in_features=dim, hidden_features=mlp_hidden_dim, act_layer=act_layer, drop=drop)
    def forward(self, x):
        H, W = self.input_resolution
        shortcut = x
        xn = self.norm1(x).contiguous()
        y = _ext.spatial_mlp_forward(xn, self.spatial_mlp.weight, self.spatial_mlp.bias,
                                     H, W, self.window_size, self.shift_size, self.num_heads)
        x = shortcut + y
        x = x + self.mlp(self.norm2(x))
        return x

class PatchMerging(nn.Module):
    def __init__(self, input_resolution, dim, norm_layer=nn.LayerNorm):
        super().__init__()
        self.input_resolution = input_resolution; self.dim = dim
        self.reduction = nn.Linear(4 * dim, 2 * dim, bias=False)
        self.norm = norm_layer(4 * dim)
    def forward(self, x):
        H, W = self.input_resolution
        B, L, C = x.shape
        x = x.view(B, H, W, C)
        x0 = x[:, 0::2, 0::2, :]
        x1 = x[:, 1::2, 0::2, :]
        x2 = x[:, 0::2, 1::2, :]
        x3 = x[:, 1::2, 1::2, :]
        x = torch.cat([x0, x1, x2, x3], -1).view(B, -1, 4 * C)
        x = self.norm(x)
        x = self.reduction(x)
        return x

class BasicLayer(nn.Module):
    def __init__(self, dim, input_resolution, depth, num_heads, window_size,
                 mlp_ratio=4., drop=0., drop_path=0., norm_layer=nn.LayerNorm,
                 downsample=None, use_checkpoint=False):
        super().__init__()
        self.dim = dim; self.input_resolution = input_resolution; self.depth = depth; self.use_checkpoint = use_checkpoint
        self.blocks = nn.ModuleList([
            SwinMLPBlock(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,
                         drop=drop, drop_path=drop_path[i] if isinstance(drop_path, list) else drop_path,
                         norm_layer=norm_layer) for i in range(depth)])
        self.downsample = downsample(input_resolution, dim=dim, norm_layer=norm_layer) if downsample is not None else None
    def forward(self, x):
        for blk in self.blocks:
            x = blk(x)
        if self.downsample is not None:
            x = self.downsample(x)
        return x

class PatchEmbed(nn.Module):
    def __init__(self, img_size=224, patch_size=4, in_chans=3, embed_dim=96, norm_layer=None):
        super().__init__()
        img_size = to_2tuple(img_size); patch_size = to_2tuple(patch_size)
        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)
        self.norm = norm_layer(embed_dim) if norm_layer is not None else None
    def forward(self, x):
        x = self.proj(x).flatten(2).transpose(1, 2)
        if self.norm is not None:
            x = self.norm(x)
        return x

class ModelNew(nn.Module):
    def __init__(self, img_size=224, patch_size=4, in_chans=3, num_classes=1000,
                 embed_dim=96, depths=[2, 2, 6, 2], num_heads=[3, 6, 12, 24],
                 window_size=7, mlp_ratio=4., drop_rate=0., drop_path_rate=0.1,
                 norm_layer=nn.LayerNorm, patch_norm=True, use_checkpoint=False, **kwargs):
        super().__init__()
        self.num_classes = num_classes; self.num_layers = len(depths); self.embed_dim = embed_dim
        self.patch_norm = patch_norm; self.num_features = int(embed_dim * 2 ** (self.num_layers - 1)); self.mlp_ratio = mlp_ratio
        self.patch_embed = PatchEmbed(img_size=img_size, patch_size=patch_size, in_chans=in_chans, embed_dim=embed_dim,
                                      norm_layer=norm_layer if self.patch_norm else None)
        patches_resolution = self.patch_embed.patches_resolution
        self.patches_resolution = patches_resolution
        self.pos_drop = nn.Dropout(p=drop_rate)
        dpr = [x.item() for x in torch.linspace(0, drop_path_rate, sum(depths))]
        self.layers = nn.ModuleList()
        for i_layer in range(self.num_layers):
            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, drop=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)
            self.layers.append(layer)
        self.norm = norm_layer(self.num_features)
        self.avgpool = nn.AdaptiveAvgPool1d(1)
        self.head = nn.Linear(self.num_features, num_classes) if num_classes > 0 else nn.Identity()
    def forward_features(self, x):
        x = self.patch_embed(x)
        x = self.pos_drop(x)
        for layer in self.layers:
            x = layer(x)
        x = self.norm(x)
        x = self.avgpool(x.transpose(1, 2))
        x = torch.flatten(x, 1)
        return x
    def forward(self, x):
        x = self.forward_features(x)
        x = self.head(x)
        return x