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
reinterpret_tensor = torch._C._dynamo.guards._reinterpret_tensor


@triton.jit
def triton_poi_fused_gt_index_max_pool2d_with_indices_0(in_ptr0, in_ptr1,
    out_ptr2, out_ptr3, 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)
    tmp31 = tmp1 > tmp0
    tmp32 = tl.full([1], 1, tl.int8)
    tmp33 = tl.full([1], 0, tl.int8)
    tmp34 = tl.where(tmp31, tmp32, tmp33)
    tmp35 = tmp3 > tmp2
    tmp36 = tl.full([1], 2, tl.int8)
    tmp37 = tl.where(tmp35, tmp36, tmp34)
    tmp38 = tmp5 > tmp4
    tmp39 = tl.full([1], 3, tl.int8)
    tmp40 = tl.where(tmp38, tmp39, tmp37)
    tmp41 = tmp7 > tmp6
    tmp42 = tl.full([1], 4, tl.int8)
    tmp43 = tl.where(tmp41, tmp42, tmp40)
    tmp44 = tmp9 > tmp8
    tmp45 = tl.full([1], 5, tl.int8)
    tmp46 = tl.where(tmp44, tmp45, tmp43)
    tmp47 = tmp11 > tmp10
    tmp48 = tl.full([1], 6, tl.int8)
    tmp49 = tl.where(tmp47, tmp48, tmp46)
    tmp50 = tmp13 > tmp12
    tmp51 = tl.full([1], 7, tl.int8)
    tmp52 = tl.where(tmp50, tmp51, tmp49)
    tmp53 = tmp15 > tmp14
    tmp54 = tl.full([1], 8, tl.int8)
    tmp55 = tl.where(tmp53, tmp54, tmp52)
    tmp56 = tmp17 > tmp16
    tmp57 = tl.full([1], 9, tl.int8)
    tmp58 = tl.where(tmp56, tmp57, tmp55)
    tmp59 = tmp19 > tmp18
    tmp60 = tl.full([1], 10, tl.int8)
    tmp61 = tl.where(tmp59, tmp60, tmp58)
    tmp62 = tmp21 > tmp20
    tmp63 = tl.full([1], 11, tl.int8)
    tmp64 = tl.where(tmp62, tmp63, tmp61)
    tmp65 = tmp23 > tmp22
    tmp66 = tl.full([1], 12, tl.int8)
    tmp67 = tl.where(tmp65, tmp66, tmp64)
    tmp68 = tmp25 > tmp24
    tmp69 = tl.full([1], 13, tl.int8)
    tmp70 = tl.where(tmp68, tmp69, tmp67)
    tmp71 = tmp27 > tmp26
    tmp72 = tl.full([1], 14, tl.int8)
    tmp73 = tl.where(tmp71, tmp72, tmp70)
    tmp74 = tmp29 > tmp28
    tmp75 = tl.full([1], 15, tl.int8)
    tmp76 = tl.where(tmp74, tmp75, tmp73)
    tmp77 = 0.0
    tmp78 = tmp30 > tmp77
    tmp79 = tl.full([1], 4, tl.int32)
    tmp80 = tl.where((tmp76 < 0) != (tmp79 < 0), tl.where(tmp76 % tmp79 != 
        0, tmp76 // tmp79 - 1, tmp76 // tmp79), tmp76 // tmp79)
    tmp81 = tmp80 * tmp79
    tmp82 = tmp76 - tmp81
    tmp83 = tl.full([1], 0, tl.int64)
    tmp84 = tmp83 + tmp80
    tmp85 = tmp83 + tmp82
    tmp86 = tl.full([1], 4, tl.int64)
    tmp87 = tmp84 * tmp86
    tmp88 = tmp87 + tmp85
    tmp89 = tl.full([XBLOCK], 256, tl.int32)
    tmp90 = tmp88 + tmp89
    tmp91 = tmp88 < 0
    tmp92 = tl.where(tmp91, tmp90, tmp88)
    tl.device_assert((0 <= tmp92) & (tmp92 < 256) | ~xmask,
        'index out of bounds: 0 <= tmp92 < 256')
    tmp94 = tl.load(in_ptr1 + tmp92, xmask, eviction_policy='evict_last')
    tl.store(out_ptr2 + x0, tmp78, xmask)
    tl.store(out_ptr3 + x0, tmp94, xmask)


def call(args):
    arg0_1, arg1_1 = args
    args.clear()
    assert_size_stride(arg0_1, (4, 4, 4, 4), (64, 16, 4, 1))
    assert_size_stride(arg1_1, (4, 4, 4, 4), (64, 16, 4, 1))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf2 = empty_strided_cuda((4, 4, 1, 1), (4, 1, 1, 1), torch.bool)
        buf3 = empty_strided_cuda((16,), (1,), torch.float32)
        get_raw_stream(0)
        triton_poi_fused_gt_index_max_pool2d_with_indices_0[grid(16)](arg0_1,
            arg1_1, buf2, buf3, 16, XBLOCK=16, num_warps=1, num_stages=1)
        del arg0_1
        del arg1_1
    return buf2, reinterpret_tensor(buf3, (4, 4, 1, 1), (4, 1, 1, 1), 0)


class _SpikeMaxPoolNd(nn.Module):

    def __init__(self, kernel_size, stride=None, padding=0, dilation=1,
        ceil_mode=False):
        super(_SpikeMaxPoolNd, self).__init__()
        self.kernel_size = kernel_size
        self.stride = stride or kernel_size
        self.padding = padding
        self.dilation = dilation
        self.ceil_mode = ceil_mode
        self.return_indices = True

    def reset_state(self):
        pass


class MaxPool2dNew(_SpikeMaxPoolNd):
    """Simple port of PyTorch MaxPool2d with small adjustment for spiking operations.
    
    Currently pooling only supports operations on floating point numbers, thus it casts the uint8 spikes to floats back and forth.
    The trace of the 'maximum' spike is also returned. In case of multiple spikes within pooling window, returns first spike of 
    the window (top left corner).
    """

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


Measured latency_ms: 0.027298