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
from torch._inductor.runtime.triton_helpers import math as tl_math
from torch.nn import Module
from typing import Union
from typing import Tuple
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_abs_add_div_max_pool2d_with_indices_0(in_out_ptr0,
    in_ptr0, xnumel, XBLOCK: tl.constexpr):
    xnumel = 256
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x1 = xindex // 4 % 4
    x0 = xindex % 4
    x3 = xindex
    tmp191 = tl.load(in_ptr0 + x3, xmask)
    tmp0 = -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 = -2 + x0
    tmp7 = tmp6 >= tmp1
    tmp8 = tmp6 < tmp3
    tmp9 = tmp7 & tmp8
    tmp10 = tmp5 & tmp9
    tmp11 = tl.load(in_ptr0 + (-10 + x3), tmp10 & xmask, other=0.0)
    tmp12 = tl_math.abs(tmp11)
    tmp13 = tl.full(tmp12.shape, float('-inf'), tmp12.dtype)
    tmp14 = tl.where(tmp10, tmp12, tmp13)
    tmp15 = -1 + x0
    tmp16 = tmp15 >= tmp1
    tmp17 = tmp15 < tmp3
    tmp18 = tmp16 & tmp17
    tmp19 = tmp5 & tmp18
    tmp20 = tl.load(in_ptr0 + (-9 + x3), tmp19 & xmask, other=0.0)
    tmp21 = tl_math.abs(tmp20)
    tmp22 = tl.full(tmp21.shape, float('-inf'), tmp21.dtype)
    tmp23 = tl.where(tmp19, tmp21, tmp22)
    tmp24 = triton_helpers.maximum(tmp23, tmp14)
    tmp25 = x0
    tmp26 = tmp25 >= tmp1
    tmp27 = tmp25 < tmp3
    tmp28 = tmp26 & tmp27
    tmp29 = tmp5 & tmp28
    tmp30 = tl.load(in_ptr0 + (-8 + x3), tmp29 & xmask, other=0.0)
    tmp31 = tl_math.abs(tmp30)
    tmp32 = tl.full(tmp31.shape, float('-inf'), tmp31.dtype)
    tmp33 = tl.where(tmp29, tmp31, tmp32)
    tmp34 = triton_helpers.maximum(tmp33, tmp24)
    tmp35 = 1 + x0
    tmp36 = tmp35 >= tmp1
    tmp37 = tmp35 < tmp3
    tmp38 = tmp36 & tmp37
    tmp39 = tmp5 & tmp38
    tmp40 = tl.load(in_ptr0 + (-7 + x3), tmp39 & xmask, other=0.0)
    tmp41 = tl_math.abs(tmp40)
    tmp42 = tl.full(tmp41.shape, float('-inf'), tmp41.dtype)
    tmp43 = tl.where(tmp39, tmp41, tmp42)
    tmp44 = triton_helpers.maximum(tmp43, tmp34)
    tmp45 = 2 + x0
    tmp46 = tmp45 >= tmp1
    tmp47 = tmp45 < tmp3
    tmp48 = tmp46 & tmp47
    tmp49 = tmp5 & tmp48
    tmp50 = tl.load(in_ptr0 + (-6 + x3), tmp49 & xmask, other=0.0)
    tmp51 = tl_math.abs(tmp50)
    tmp52 = tl.full(tmp51.shape, float('-inf'), tmp51.dtype)
    tmp53 = tl.where(tmp49, tmp51, tmp52)
    tmp54 = triton_helpers.maximum(tmp53, tmp44)
    tmp55 = -1 + x1
    tmp56 = tmp55 >= tmp1
    tmp57 = tmp55 < tmp3
    tmp58 = tmp56 & tmp57
    tmp59 = tmp58 & tmp9
    tmp60 = tl.load(in_ptr0 + (-6 + x3), tmp59 & xmask, other=0.0)
    tmp61 = tl_math.abs(tmp60)
    tmp62 = tl.full(tmp61.shape, float('-inf'), tmp61.dtype)
    tmp63 = tl.where(tmp59, tmp61, tmp62)
    tmp64 = triton_helpers.maximum(tmp63, tmp54)
    tmp65 = tmp58 & tmp18
    tmp66 = tl.load(in_ptr0 + (-5 + x3), tmp65 & xmask, other=0.0)
    tmp67 = tl_math.abs(tmp66)
    tmp68 = tl.full(tmp67.shape, float('-inf'), tmp67.dtype)
    tmp69 = tl.where(tmp65, tmp67, tmp68)
    tmp70 = triton_helpers.maximum(tmp69, tmp64)
    tmp71 = tmp58 & tmp28
    tmp72 = tl.load(in_ptr0 + (-4 + x3), tmp71 & xmask, other=0.0)
    tmp73 = tl_math.abs(tmp72)
    tmp74 = tl.full(tmp73.shape, float('-inf'), tmp73.dtype)
    tmp75 = tl.where(tmp71, tmp73, tmp74)
    tmp76 = triton_helpers.maximum(tmp75, tmp70)
    tmp77 = tmp58 & tmp38
    tmp78 = tl.load(in_ptr0 + (-3 + x3), tmp77 & xmask, other=0.0)
    tmp79 = tl_math.abs(tmp78)
    tmp80 = tl.full(tmp79.shape, float('-inf'), tmp79.dtype)
    tmp81 = tl.where(tmp77, tmp79, tmp80)
    tmp82 = triton_helpers.maximum(tmp81, tmp76)
    tmp83 = tmp58 & tmp48
    tmp84 = tl.load(in_ptr0 + (-2 + x3), tmp83 & xmask, other=0.0)
    tmp85 = tl_math.abs(tmp84)
    tmp86 = tl.full(tmp85.shape, float('-inf'), tmp85.dtype)
    tmp87 = tl.where(tmp83, tmp85, tmp86)
    tmp88 = triton_helpers.maximum(tmp87, tmp82)
    tmp89 = x1
    tmp90 = tmp89 >= tmp1
    tmp91 = tmp89 < tmp3
    tmp92 = tmp90 & tmp91
    tmp93 = tmp92 & tmp9
    tmp94 = tl.load(in_ptr0 + (-2 + x3), tmp93 & xmask, other=0.0)
    tmp95 = tl_math.abs(tmp94)
    tmp96 = tl.full(tmp95.shape, float('-inf'), tmp95.dtype)
    tmp97 = tl.where(tmp93, tmp95, tmp96)
    tmp98 = triton_helpers.maximum(tmp97, tmp88)
    tmp99 = tmp92 & tmp18
    tmp100 = tl.load(in_ptr0 + (-1 + x3), tmp99 & xmask, other=0.0)
    tmp101 = tl_math.abs(tmp100)
    tmp102 = tl.full(tmp101.shape, float('-inf'), tmp101.dtype)
    tmp103 = tl.where(tmp99, tmp101, tmp102)
    tmp104 = triton_helpers.maximum(tmp103, tmp98)
    tmp105 = tmp92 & tmp28
    tmp106 = tl.load(in_ptr0 + x3, tmp105 & xmask, other=0.0)
    tmp107 = tl_math.abs(tmp106)
    tmp108 = tl.full(tmp107.shape, float('-inf'), tmp107.dtype)
    tmp109 = tl.where(tmp105, tmp107, tmp108)
    tmp110 = triton_helpers.maximum(tmp109, tmp104)
    tmp111 = tmp92 & tmp38
    tmp112 = tl.load(in_ptr0 + (1 + x3), tmp111 & xmask, other=0.0)
    tmp113 = tl_math.abs(tmp112)
    tmp114 = tl.full(tmp113.shape, float('-inf'), tmp113.dtype)
    tmp115 = tl.where(tmp111, tmp113, tmp114)
    tmp116 = triton_helpers.maximum(tmp115, tmp110)
    tmp117 = tmp92 & tmp48
    tmp118 = tl.load(in_ptr0 + (2 + x3), tmp117 & xmask, other=0.0)
    tmp119 = tl_math.abs(tmp118)
    tmp120 = tl.full(tmp119.shape, float('-inf'), tmp119.dtype)
    tmp121 = tl.where(tmp117, tmp119, tmp120)
    tmp122 = triton_helpers.maximum(tmp121, tmp116)
    tmp123 = 1 + x1
    tmp124 = tmp123 >= tmp1
    tmp125 = tmp123 < tmp3
    tmp126 = tmp124 & tmp125
    tmp127 = tmp126 & tmp9
    tmp128 = tl.load(in_ptr0 + (2 + x3), tmp127 & xmask, other=0.0)
    tmp129 = tl_math.abs(tmp128)
    tmp130 = tl.full(tmp129.shape, float('-inf'), tmp129.dtype)
    tmp131 = tl.where(tmp127, tmp129, tmp130)
    tmp132 = triton_helpers.maximum(tmp131, tmp122)
    tmp133 = tmp126 & tmp18
    tmp134 = tl.load(in_ptr0 + (3 + x3), tmp133 & xmask, other=0.0)
    tmp135 = tl_math.abs(tmp134)
    tmp136 = tl.full(tmp135.shape, float('-inf'), tmp135.dtype)
    tmp137 = tl.where(tmp133, tmp135, tmp136)
    tmp138 = triton_helpers.maximum(tmp137, tmp132)
    tmp139 = tmp126 & tmp28
    tmp140 = tl.load(in_ptr0 + (4 + x3), tmp139 & xmask, other=0.0)
    tmp141 = tl_math.abs(tmp140)
    tmp142 = tl.full(tmp141.shape, float('-inf'), tmp141.dtype)
    tmp143 = tl.where(tmp139, tmp141, tmp142)
    tmp144 = triton_helpers.maximum(tmp143, tmp138)
    tmp145 = tmp126 & tmp38
    tmp146 = tl.load(in_ptr0 + (5 + x3), tmp145 & xmask, other=0.0)
    tmp147 = tl_math.abs(tmp146)
    tmp148 = tl.full(tmp147.shape, float('-inf'), tmp147.dtype)
    tmp149 = tl.where(tmp145, tmp147, tmp148)
    tmp150 = triton_helpers.maximum(tmp149, tmp144)
    tmp151 = tmp126 & tmp48
    tmp152 = tl.load(in_ptr0 + (6 + x3), tmp151 & xmask, other=0.0)
    tmp153 = tl_math.abs(tmp152)
    tmp154 = tl.full(tmp153.shape, float('-inf'), tmp153.dtype)
    tmp155 = tl.where(tmp151, tmp153, tmp154)
    tmp156 = triton_helpers.maximum(tmp155, tmp150)
    tmp157 = 2 + x1
    tmp158 = tmp157 >= tmp1
    tmp159 = tmp157 < tmp3
    tmp160 = tmp158 & tmp159
    tmp161 = tmp160 & tmp9
    tmp162 = tl.load(in_ptr0 + (6 + x3), tmp161 & xmask, other=0.0)
    tmp163 = tl_math.abs(tmp162)
    tmp164 = tl.full(tmp163.shape, float('-inf'), tmp163.dtype)
    tmp165 = tl.where(tmp161, tmp163, tmp164)
    tmp166 = triton_helpers.maximum(tmp165, tmp156)
    tmp167 = tmp160 & tmp18
    tmp168 = tl.load(in_ptr0 + (7 + x3), tmp167 & xmask, other=0.0)
    tmp169 = tl_math.abs(tmp168)
    tmp170 = tl.full(tmp169.shape, float('-inf'), tmp169.dtype)
    tmp171 = tl.where(tmp167, tmp169, tmp170)
    tmp172 = triton_helpers.maximum(tmp171, tmp166)
    tmp173 = tmp160 & tmp28
    tmp174 = tl.load(in_ptr0 + (8 + x3), tmp173 & xmask, other=0.0)
    tmp175 = tl_math.abs(tmp174)
    tmp176 = tl.full(tmp175.shape, float('-inf'), tmp175.dtype)
    tmp177 = tl.where(tmp173, tmp175, tmp176)
    tmp178 = triton_helpers.maximum(tmp177, tmp172)
    tmp179 = tmp160 & tmp38
    tmp180 = tl.load(in_ptr0 + (9 + x3), tmp179 & xmask, other=0.0)
    tmp181 = tl_math.abs(tmp180)
    tmp182 = tl.full(tmp181.shape, float('-inf'), tmp181.dtype)
    tmp183 = tl.where(tmp179, tmp181, tmp182)
    tmp184 = triton_helpers.maximum(tmp183, tmp178)
    tmp185 = tmp160 & tmp48
    tmp186 = tl.load(in_ptr0 + (10 + x3), tmp185 & xmask, other=0.0)
    tmp187 = tl_math.abs(tmp186)
    tmp188 = tl.full(tmp187.shape, float('-inf'), tmp187.dtype)
    tmp189 = tl.where(tmp185, tmp187, tmp188)
    tmp190 = triton_helpers.maximum(tmp189, tmp184)
    tmp192 = 1.0
    tmp193 = tmp190 + tmp192
    tmp194 = tmp191 / tmp193
    tl.store(in_out_ptr0 + x3, tmp194, 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, 4, 4), (64, 16, 4, 1), torch.float32)
        buf1 = buf0
        del buf0
        get_raw_stream(0)
        triton_poi_fused_abs_add_div_max_pool2d_with_indices_0[grid(256)](buf1,
            arg0_1, 256, XBLOCK=128, num_warps=4, num_stages=1)
        del arg0_1
    return buf1,


class DivisiveNormalization2dNew(Module):
    """Applies a 2D divisive normalization over an input signal composed of several input
    planes.

    In the simplest case, the output value of the layer with input size :math:`(N, C, H, W)`
    and output :math:`(N, C, H, W)`.

    Args:
        b_type: Type of suppressin field, must be one of (`linf`, `l1`, `l2`).
        b_size: The size of the suppression field, must be > 0.
        sigma: Constant added to suppression field, must be > 0.

    Shape:
        - Input: :math:`(N, C, H, W)`
        - Output: :math:`(N, C, H, W)`

    Examples::

        >>> # suppression of size=3, sigma=1
        >>> d = DivisiveNormalization2d(b_size=3, sigma=1)
        >>> input = torch.randn(20, 16, 50, 50)
        >>> output = d(input)
    """

    def __init__(self, b_type: 'str'='linf', b_size:
        'Union[int, Tuple[int, int]]'=(5, 5), sigma: 'float'=1.0) ->None:
        super(DivisiveNormalization2dNew, self).__init__()
        self.sigma = sigma
        if isinstance(b_size, int):
            self.b_size = b_size, b_size
        else:
            self.b_size = b_size
        self.padding = self.b_size[0] // 2, self.b_size[1] // 2
        self.b_type = b_type

    def __repr__(self) ->str:
        s = 'DivisiveNormalization2d('
        s += f'b_type={self.b_type}, b_size={self.b_size}, sigma={self.sigma}'
        s += ')'
        return s.format(**self.__dict__)

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


Measured latency_ms: 0.014342