import torch
import triton
import triton.language as tl
from torch._inductor.runtime.triton_heuristics import grid
from torch._C import _cuda_getCurrentRawStream as get_raw_stream
from torch._inductor.runtime import triton_helpers
import torch.nn as nn
assert_size_stride = torch._C._dynamo.guards.assert_size_stride
empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda


@triton.jit
def triton_poi_fused_max_pool2d_with_indices_0(in_ptr0, out_ptr0, xnumel,
    XBLOCK: tl.constexpr):
    xnumel = 16
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + 16 * x0, xmask, eviction_policy='evict_last')
    tmp1 = tl.load(in_ptr0 + (1 + 16 * x0), xmask, eviction_policy='evict_last'
        )
    tmp3 = tl.load(in_ptr0 + (2 + 16 * x0), xmask, eviction_policy='evict_last'
        )
    tmp5 = tl.load(in_ptr0 + (3 + 16 * x0), xmask, eviction_policy='evict_last'
        )
    tmp7 = tl.load(in_ptr0 + (4 + 16 * x0), xmask, eviction_policy='evict_last'
        )
    tmp9 = tl.load(in_ptr0 + (5 + 16 * x0), xmask, eviction_policy='evict_last'
        )
    tmp11 = tl.load(in_ptr0 + (6 + 16 * x0), xmask, eviction_policy=
        'evict_last')
    tmp13 = tl.load(in_ptr0 + (7 + 16 * x0), xmask, eviction_policy=
        'evict_last')
    tmp15 = tl.load(in_ptr0 + (8 + 16 * x0), xmask, eviction_policy=
        'evict_last')
    tmp17 = tl.load(in_ptr0 + (9 + 16 * x0), xmask, eviction_policy=
        'evict_last')
    tmp19 = tl.load(in_ptr0 + (10 + 16 * x0), xmask, eviction_policy=
        'evict_last')
    tmp21 = tl.load(in_ptr0 + (11 + 16 * x0), xmask, eviction_policy=
        'evict_last')
    tmp23 = tl.load(in_ptr0 + (12 + 16 * x0), xmask, eviction_policy=
        'evict_last')
    tmp25 = tl.load(in_ptr0 + (13 + 16 * x0), xmask, eviction_policy=
        'evict_last')
    tmp27 = tl.load(in_ptr0 + (14 + 16 * x0), xmask, eviction_policy=
        'evict_last')
    tmp29 = tl.load(in_ptr0 + (15 + 16 * x0), xmask, eviction_policy=
        'evict_last')
    tmp2 = triton_helpers.maximum(tmp1, tmp0)
    tmp4 = triton_helpers.maximum(tmp3, tmp2)
    tmp6 = triton_helpers.maximum(tmp5, tmp4)
    tmp8 = triton_helpers.maximum(tmp7, tmp6)
    tmp10 = triton_helpers.maximum(tmp9, tmp8)
    tmp12 = triton_helpers.maximum(tmp11, tmp10)
    tmp14 = triton_helpers.maximum(tmp13, tmp12)
    tmp16 = triton_helpers.maximum(tmp15, tmp14)
    tmp18 = triton_helpers.maximum(tmp17, tmp16)
    tmp20 = triton_helpers.maximum(tmp19, tmp18)
    tmp22 = triton_helpers.maximum(tmp21, tmp20)
    tmp24 = triton_helpers.maximum(tmp23, tmp22)
    tmp26 = triton_helpers.maximum(tmp25, tmp24)
    tmp28 = triton_helpers.maximum(tmp27, tmp26)
    tmp30 = triton_helpers.maximum(tmp29, tmp28)
    tl.store(out_ptr0 + x0, tmp30, xmask)


@triton.jit
def triton_poi_fused_max_pool2d_with_indices_1(in_ptr0, out_ptr0, xnumel,
    XBLOCK: tl.constexpr):
    xnumel = 144
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x1 = xindex // 3 % 3
    x0 = xindex % 3
    x2 = xindex // 9
    x4 = xindex
    tmp0 = -1 + 2 * x1
    tmp1 = tl.full([1], 0, tl.int64)
    tmp2 = tmp0 >= tmp1
    tmp3 = tl.full([1], 4, tl.int64)
    tmp4 = tmp0 < tmp3
    tmp5 = tmp2 & tmp4
    tmp6 = -1 + 2 * x0
    tmp7 = tmp6 >= tmp1
    tmp8 = tmp6 < tmp3
    tmp9 = tmp7 & tmp8
    tmp10 = tmp5 & tmp9
    tmp11 = tl.load(in_ptr0 + (-5 + 2 * x0 + 8 * x1 + 16 * x2), tmp10 &
        xmask, eviction_policy='evict_last', other=float('-inf'))
    tmp12 = 2 * x0
    tmp13 = tmp12 >= tmp1
    tmp14 = tmp12 < tmp3
    tmp15 = tmp13 & tmp14
    tmp16 = tmp5 & tmp15
    tmp17 = tl.load(in_ptr0 + (-4 + 2 * x0 + 8 * x1 + 16 * x2), tmp16 &
        xmask, eviction_policy='evict_last', other=float('-inf'))
    tmp18 = triton_helpers.maximum(tmp17, tmp11)
    tmp19 = 2 * x1
    tmp20 = tmp19 >= tmp1
    tmp21 = tmp19 < tmp3
    tmp22 = tmp20 & tmp21
    tmp23 = tmp22 & tmp9
    tmp24 = tl.load(in_ptr0 + (-1 + 2 * x0 + 8 * x1 + 16 * x2), tmp23 &
        xmask, eviction_policy='evict_last', other=float('-inf'))
    tmp25 = triton_helpers.maximum(tmp24, tmp18)
    tmp26 = tmp22 & tmp15
    tmp27 = tl.load(in_ptr0 + (2 * x0 + 8 * x1 + 16 * x2), tmp26 & xmask,
        eviction_policy='evict_last', other=float('-inf'))
    tmp28 = triton_helpers.maximum(tmp27, tmp25)
    tl.store(out_ptr0 + x4, tmp28, xmask)


@triton.jit
def triton_poi_fused_cat_2(in_ptr0, in_ptr1, in_ptr2, out_ptr0, xnumel,
    XBLOCK: tl.constexpr):
    xnumel = 480
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex % 120
    x1 = xindex // 120
    x2 = xindex
    tmp0 = x0
    tl.full([1], 0, tl.int64)
    tmp3 = tl.full([1], 56, tl.int64)
    tmp4 = tmp0 < tmp3
    tmp5 = tl.full([1], 20, tl.int64)
    tmp6 = tmp0 < tmp5
    tmp7 = tmp6 & tmp4
    tmp8 = tl.full([1], 4, tl.int64)
    tmp9 = tmp0 < tmp8
    tmp10 = tmp9 & tmp7
    tmp11 = tl.load(in_ptr0 + (4 * x1 + x0), tmp10 & xmask, eviction_policy
        ='evict_last', other=0.0)
    tmp12 = tmp0 >= tmp8
    tmp13 = tmp12 & tmp7
    tmp14 = tl.load(in_ptr1 + (2 * ((-4 + x0) % 2) + 8 * ((-4 + x0) // 2 % 
        8) + 64 * x1), tmp13 & xmask, eviction_policy='evict_last', other=0.0)
    tmp15 = tl.load(in_ptr1 + (1 + 2 * ((-4 + x0) % 2) + 8 * ((-4 + x0) // 
        2 % 8) + 64 * x1), tmp13 & xmask, eviction_policy='evict_last',
        other=0.0)
    tmp16 = triton_helpers.maximum(tmp15, tmp14)
    tmp17 = tl.load(in_ptr1 + (4 + 2 * ((-4 + x0) % 2) + 8 * ((-4 + x0) // 
        2 % 8) + 64 * x1), tmp13 & xmask, eviction_policy='evict_last',
        other=0.0)
    tmp18 = triton_helpers.maximum(tmp17, tmp16)
    tmp19 = tl.load(in_ptr1 + (5 + 2 * ((-4 + x0) % 2) + 8 * ((-4 + x0) // 
        2 % 8) + 64 * x1), tmp13 & xmask, eviction_policy='evict_last',
        other=0.0)
    tmp20 = triton_helpers.maximum(tmp19, tmp18)
    tmp21 = tl.full(tmp20.shape, 0.0, tmp20.dtype)
    tmp22 = tl.where(tmp13, tmp20, tmp21)
    tmp23 = tl.where(tmp9, tmp11, tmp22)
    tmp24 = tl.full(tmp23.shape, 0.0, tmp23.dtype)
    tmp25 = tl.where(tmp7, tmp23, tmp24)
    tmp26 = tmp0 >= tmp5
    tmp27 = tmp26 & tmp4
    tmp28 = tl.load(in_ptr2 + (36 * x1 + (-20 + x0)), tmp27 & xmask,
        eviction_policy='evict_last', other=0.0)
    tmp29 = tl.where(tmp6, tmp25, tmp28)
    tmp30 = tl.full(tmp29.shape, 0.0, tmp29.dtype)
    tmp31 = tl.where(tmp4, tmp29, tmp30)
    tmp32 = tmp0 >= tmp3
    tl.full([1], 120, tl.int64)
    tmp35 = tl.load(in_ptr1 + (64 * x1 + (-56 + x0) % 64), tmp32 & xmask,
        eviction_policy='evict_last', other=0.0)
    tmp36 = tl.where(tmp4, tmp31, tmp35)
    tl.store(out_ptr0 + x2, tmp36, xmask)


def call(args):
    arg0_1, = args
    args.clear()
    assert_size_stride(arg0_1, (4, 4, 4, 4), (64, 16, 4, 1))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((4, 4, 1, 1), (4, 1, 1, 1), torch.float32)
        get_raw_stream(0)
        triton_poi_fused_max_pool2d_with_indices_0[grid(16)](arg0_1, buf0, 
            16, XBLOCK=16, num_warps=1, num_stages=1)
        buf1 = empty_strided_cuda((4, 4, 3, 3), (36, 9, 3, 1), torch.float32)
        triton_poi_fused_max_pool2d_with_indices_1[grid(144)](arg0_1, buf1,
            144, XBLOCK=256, num_warps=4, num_stages=1)
        buf2 = empty_strided_cuda((4, 120), (120, 1), torch.float32)
        triton_poi_fused_cat_2[grid(480)](buf0, arg0_1, buf1, buf2, 480,
            XBLOCK=256, num_warps=4, num_stages=1)
        del arg0_1
        del buf0
        del buf1
    return buf2,


class SpatialPyramidPooling2dNew(nn.Module):
    """apply spatial pyramid pooling over a 4d input(a mini-batch of 2d inputs
    with additional channel dimension) as described in the paper
    'Spatial Pyramid Pooling in deep convolutional Networks for visual recognition'
    Args:
        num_level:
        pool_type: max_pool, avg_pool, Default:max_pool
    By the way, the target output size is num_grid:
        num_grid = 0
        for i in range num_level:
            num_grid += (i + 1) * (i + 1)
        num_grid = num_grid * channels # channels is the channel dimension of input data
    examples:
        >>> input = torch.randn((1,3,32,32), dtype=torch.float32)
        >>> net = torch.nn.Sequential(nn.Conv2d(in_channels=3,out_channels=32,kernel_size=3,stride=1),\\
                                      nn.ReLU(),\\
                                      SpatialPyramidPooling2d(num_level=2,pool_type='avg_pool'),\\
                                      nn.Linear(32 * (1*1 + 2*2), 10))
        >>> output = net(input)
    """

    def __init__(self, num_level, pool_type='max_pool'):
        super(SpatialPyramidPooling2dNew, self).__init__()
        self.num_level = num_level
        self.pool_type = pool_type

    def __repr__(self):
        return self.__class__.__name__ + '(' + 'num_level = ' + str(self.
            num_level) + ', pool_type = ' + str(self.pool_type) + ')'

    def forward(self, input_0):
        arg0_1 = input_0
        output = call([arg0_1])
        return output[0]


Measured latency_ms: 0.039417