import torch
from torch._inductor.select_algorithm import extern_kernels
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 numpy as np
import torch as th
import torch.utils.data
import torch.nn as nn
from collections import OrderedDict
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_add_0(in_ptr0, in_ptr1, out_ptr0, xnumel, XBLOCK: tl.
    constexpr):
    xnumel = 256
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + x0, xmask)
    tmp1 = tl.load(in_ptr1 + 0)
    tmp2 = tl.broadcast_to(tmp1, [XBLOCK])
    tmp3 = tmp0 + tmp2
    tl.store(out_ptr0 + x0, tmp3, xmask)


@triton.jit
def triton_poi_fused_add_convolution_relu_threshold_backward_1(in_ptr0,
    in_ptr1, in_ptr2, in_ptr3, out_ptr0, out_ptr1, xnumel, XBLOCK: tl.constexpr
    ):
    xnumel = 256
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x3 = xindex
    x1 = xindex // 16 % 4
    tmp0 = tl.load(in_ptr0 + x3, xmask)
    tmp1 = tl.load(in_ptr1 + x1, xmask, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr2 + 0)
    tmp4 = tl.broadcast_to(tmp3, [XBLOCK])
    tmp8 = tl.load(in_ptr3 + 0)
    tmp9 = tl.broadcast_to(tmp8, [XBLOCK])
    tmp2 = tmp0 + tmp1
    tmp5 = tmp2 + tmp4
    tmp6 = tl.full([1], 0, tl.int32)
    tmp7 = triton_helpers.maximum(tmp6, tmp5)
    tmp10 = tmp7 + tmp9
    tmp11 = 0.0
    tmp12 = tmp7 <= tmp11
    tl.store(out_ptr0 + x3, tmp10, xmask)
    tl.store(out_ptr1 + x3, tmp12, xmask)


@triton.jit
def triton_poi_fused_add_convolution_mul_relu_2(in_out_ptr0, in_ptr0,
    in_ptr1, in_ptr2, in_ptr3, in_ptr4, out_ptr0, xnumel, XBLOCK: tl.constexpr
    ):
    xnumel = 256
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x3 = xindex
    x1 = xindex // 16 % 4
    tmp0 = tl.load(in_out_ptr0 + x3, xmask)
    tmp1 = tl.load(in_ptr0 + x1, xmask, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr1 + 0)
    tmp4 = tl.broadcast_to(tmp3, [XBLOCK])
    tmp6 = tl.load(in_ptr2 + 0)
    tmp7 = tl.broadcast_to(tmp6, [XBLOCK])
    tmp9 = tl.load(in_ptr3 + x3, xmask)
    tmp13 = tl.load(in_ptr4 + 0)
    tmp14 = tl.broadcast_to(tmp13, [XBLOCK])
    tmp2 = tmp0 + tmp1
    tmp5 = tmp2 * tmp4
    tmp8 = tmp5 + tmp7
    tmp10 = tmp8 + tmp9
    tmp11 = tl.full([1], 0, tl.int32)
    tmp12 = triton_helpers.maximum(tmp11, tmp10)
    tmp15 = tmp12 + tmp14
    tl.store(in_out_ptr0 + x3, tmp2, xmask)
    tl.store(out_ptr0 + x3, tmp15, xmask)


@triton.jit
def triton_poi_fused_add_convolution_mul_relu_threshold_backward_3(in_out_ptr0,
    in_ptr0, in_ptr1, in_ptr2, in_ptr3, in_ptr4, in_ptr5, in_ptr6, in_ptr7,
    out_ptr0, out_ptr1, out_ptr2, xnumel, XBLOCK: tl.constexpr):
    xnumel = 256
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x3 = xindex
    x1 = xindex // 16 % 4
    tmp0 = tl.load(in_out_ptr0 + x3, xmask)
    tmp1 = tl.load(in_ptr0 + x1, xmask, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr1 + 0)
    tmp4 = tl.broadcast_to(tmp3, [XBLOCK])
    tmp6 = tl.load(in_ptr2 + 0)
    tmp7 = tl.broadcast_to(tmp6, [XBLOCK])
    tmp9 = tl.load(in_ptr3 + x3, xmask)
    tmp10 = tl.load(in_ptr4 + 0)
    tmp11 = tl.broadcast_to(tmp10, [XBLOCK])
    tmp13 = tl.load(in_ptr5 + 0)
    tmp14 = tl.broadcast_to(tmp13, [XBLOCK])
    tmp16 = tl.load(in_ptr6 + x3, xmask)
    tmp24 = tl.load(in_ptr7 + 0)
    tmp25 = tl.broadcast_to(tmp24, [XBLOCK])
    tmp2 = tmp0 + tmp1
    tmp5 = tmp2 * tmp4
    tmp8 = tmp5 + tmp7
    tmp12 = tmp9 * tmp11
    tmp15 = tmp12 + tmp14
    tmp17 = tmp15 + tmp16
    tmp18 = tl.full([1], 0, tl.int32)
    tmp19 = triton_helpers.maximum(tmp18, tmp17)
    tmp20 = tmp8 + tmp19
    tmp21 = triton_helpers.maximum(tmp18, tmp20)
    tmp22 = 0.0
    tmp23 = tmp19 <= tmp22
    tmp26 = tmp21 + tmp25
    tl.store(in_out_ptr0 + x3, tmp2, xmask)
    tl.store(out_ptr0 + x3, tmp21, xmask)
    tl.store(out_ptr1 + x3, tmp23, xmask)
    tl.store(out_ptr2 + x3, tmp26, xmask)


@triton.jit
def triton_poi_fused_add_convolution_mul_relu_threshold_backward_4(in_out_ptr0,
    in_ptr0, in_ptr1, in_ptr2, in_ptr3, out_ptr0, out_ptr1, out_ptr2,
    xnumel, XBLOCK: tl.constexpr):
    xnumel = 256
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x3 = xindex
    x1 = xindex // 16 % 4
    tmp0 = tl.load(in_out_ptr0 + x3, xmask)
    tmp1 = tl.load(in_ptr0 + x1, xmask, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr1 + 0)
    tmp4 = tl.broadcast_to(tmp3, [XBLOCK])
    tmp6 = tl.load(in_ptr2 + 0)
    tmp7 = tl.broadcast_to(tmp6, [XBLOCK])
    tmp9 = tl.load(in_ptr3 + x3, xmask)
    tmp2 = tmp0 + tmp1
    tmp5 = tmp2 * tmp4
    tmp8 = tmp5 + tmp7
    tmp10 = tmp8 + tmp9
    tmp11 = tl.full([1], 0, tl.int32)
    tmp12 = triton_helpers.maximum(tmp11, tmp10)
    tmp13 = 0.0
    tmp14 = tmp12 <= tmp13
    tmp15 = tmp9 <= tmp13
    tl.store(in_out_ptr0 + x3, tmp2, xmask)
    tl.store(out_ptr0 + x3, tmp12, xmask)
    tl.store(out_ptr1 + x3, tmp14, xmask)
    tl.store(out_ptr2 + x3, tmp15, xmask)


def call(args):
    (primals_1, primals_2, primals_3, primals_4, primals_5, primals_6,
        primals_7, primals_8, primals_9, primals_10, primals_11, primals_12,
        primals_13, primals_14, primals_15, primals_16, primals_17,
        primals_18, primals_19, primals_20, primals_21, primals_22,
        primals_23, primals_24, primals_25, primals_26, primals_27, primals_28
        ) = args
    args.clear()
    assert_size_stride(primals_1, (4, 4, 4, 4), (64, 16, 4, 1))
    assert_size_stride(primals_2, (1,), (1,))
    assert_size_stride(primals_3, (4, 4, 3, 3), (36, 9, 3, 1))
    assert_size_stride(primals_4, (4,), (1,))
    assert_size_stride(primals_5, (1,), (1,))
    assert_size_stride(primals_6, (1,), (1,))
    assert_size_stride(primals_7, (4, 4, 3, 3), (36, 9, 3, 1))
    assert_size_stride(primals_8, (4,), (1,))
    assert_size_stride(primals_9, (1,), (1,))
    assert_size_stride(primals_10, (1,), (1,))
    assert_size_stride(primals_11, (1,), (1,))
    assert_size_stride(primals_12, (4, 4, 3, 3), (36, 9, 3, 1))
    assert_size_stride(primals_13, (4,), (1,))
    assert_size_stride(primals_14, (1,), (1,))
    assert_size_stride(primals_15, (1,), (1,))
    assert_size_stride(primals_16, (4, 4, 3, 3), (36, 9, 3, 1))
    assert_size_stride(primals_17, (4,), (1,))
    assert_size_stride(primals_18, (1,), (1,))
    assert_size_stride(primals_19, (1,), (1,))
    assert_size_stride(primals_20, (1,), (1,))
    assert_size_stride(primals_21, (4, 4, 3, 3), (36, 9, 3, 1))
    assert_size_stride(primals_22, (4,), (1,))
    assert_size_stride(primals_23, (1,), (1,))
    assert_size_stride(primals_24, (1,), (1,))
    assert_size_stride(primals_25, (4, 4, 3, 3), (36, 9, 3, 1))
    assert_size_stride(primals_26, (4,), (1,))
    assert_size_stride(primals_27, (1,), (1,))
    assert_size_stride(primals_28, (1,), (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)
        get_raw_stream(0)
        triton_poi_fused_add_0[grid(256)](primals_1, primals_2, buf0, 256,
            XBLOCK=256, num_warps=4, num_stages=1)
        del primals_2
        buf1 = extern_kernels.convolution(buf0, primals_3, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf1, (4, 4, 4, 4), (64, 16, 4, 1))
        buf2 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.float32)
        buf22 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.bool)
        triton_poi_fused_add_convolution_relu_threshold_backward_1[grid(256)](
            buf1, primals_4, primals_5, primals_6, buf2, buf22, 256, XBLOCK
            =256, num_warps=4, num_stages=1)
        del primals_4
        del primals_5
        del primals_6
        buf3 = extern_kernels.convolution(buf2, primals_7, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf3, (4, 4, 4, 4), (64, 16, 4, 1))
        buf4 = buf3
        del buf3
        buf5 = buf1
        del buf1
        triton_poi_fused_add_convolution_mul_relu_2[grid(256)](buf4,
            primals_8, primals_9, primals_10, primals_1, primals_11, buf5, 
            256, XBLOCK=256, num_warps=4, num_stages=1)
        del primals_11
        del primals_8
        buf6 = extern_kernels.convolution(buf5, primals_12, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf6, (4, 4, 4, 4), (64, 16, 4, 1))
        buf7 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.float32)
        buf20 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.bool)
        triton_poi_fused_add_convolution_relu_threshold_backward_1[grid(256)](
            buf6, primals_13, primals_14, primals_15, buf7, buf20, 256,
            XBLOCK=256, num_warps=4, num_stages=1)
        del primals_13
        del primals_14
        del primals_15
        buf8 = extern_kernels.convolution(buf7, primals_16, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf8, (4, 4, 4, 4), (64, 16, 4, 1))
        buf9 = buf8
        del buf8
        buf10 = buf6
        del buf6
        buf21 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.bool)
        buf11 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.float32)
        triton_poi_fused_add_convolution_mul_relu_threshold_backward_3[grid
            (256)](buf9, primals_17, primals_18, primals_19, buf4,
            primals_9, primals_10, primals_1, primals_20, buf10, buf21,
            buf11, 256, XBLOCK=256, num_warps=4, num_stages=1)
        del primals_1
        del primals_10
        del primals_17
        del primals_19
        del primals_20
        buf12 = extern_kernels.convolution(buf11, primals_21, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf12, (4, 4, 4, 4), (64, 16, 4, 1))
        buf13 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.float32)
        buf18 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.bool)
        triton_poi_fused_add_convolution_relu_threshold_backward_1[grid(256)](
            buf12, primals_22, primals_23, primals_24, buf13, buf18, 256,
            XBLOCK=256, num_warps=4, num_stages=1)
        del primals_22
        del primals_23
        del primals_24
        buf14 = extern_kernels.convolution(buf13, primals_25, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf14, (4, 4, 4, 4), (64, 16, 4, 1))
        buf15 = buf14
        del buf14
        buf16 = buf12
        del buf12
        buf17 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.bool)
        buf19 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.bool)
        triton_poi_fused_add_convolution_mul_relu_threshold_backward_4[grid
            (256)](buf15, primals_26, primals_27, primals_28, buf10, buf16,
            buf17, buf19, 256, XBLOCK=256, num_warps=4, num_stages=1)
        del buf10
        del primals_26
        del primals_28
    return (buf16, primals_3, primals_7, primals_9, primals_12, primals_16,
        primals_18, primals_21, primals_25, primals_27, buf0, buf2, buf4,
        buf5, buf7, buf9, buf11, buf13, buf15, buf17, buf18, buf19, buf20,
        buf21, buf22)


def _get_activation(activation):
    valid = ['relu', 'leaky_relu', 'lrelu', 'tanh', 'sigmoid']
    assert activation in valid, 'activation should be one of {}'.format(valid)
    if activation == 'relu':
        return nn.ReLU(inplace=True)
    if activation == 'leaky_relu' or activation == 'lrelu':
        return nn.LeakyReLU(inplace=True)
    if activation == 'sigmoid':
        return nn.Sigmoid()
    if activation == 'tanh':
        return nn.Tanh()
    return None


def _init_fc_or_conv(fc_conv, activation):
    gain = 1.0
    if activation is not None:
        gain = nn.init.calculate_gain(activation)
    nn.init.xavier_uniform_(fc_conv.weight, gain)
    if fc_conv.bias is not None:
        nn.init.constant_(fc_conv.bias, 0.0)


def _get_norm_layer(norm_layer, channels):
    valid = ['instance', 'batch']
    assert norm_layer in valid, 'norm_layer should be one of {}'.format(valid)
    if norm_layer == 'instance':
        layer = nn.InstanceNorm2d(channels, affine=True)
    elif norm_layer == 'batch':
        layer = nn.BatchNorm2d(channels, affine=True)
    nn.init.constant_(layer.bias, 0.0)
    nn.init.constant_(layer.weight, 1.0)
    return layer


class ConvModule(nn.Module):
    """Basic convolution module with conv + norm(optional) + activation(optional).

    Args:
      n_in(int): number of input channels.
      n_out(int): number of output channels.
      ksize(int): size of the convolution kernel (square).
      stride(int): downsampling factor
      pad(bool): if True, zero pad the convolutions to maintain a constant size.
      activation(str): nonlinear activation function between convolutions.
      norm_layer(str): normalization to apply between the convolution modules.
    """

    def __init__(self, n_in, n_out, ksize=3, stride=1, pad=True, activation
        =None, norm_layer=None):
        super(ConvModule, self).__init__()
        assert isinstance(n_in, int
            ) and n_in > 0, 'Input channels should be a positive integer got {}'.format(
            n_in)
        assert isinstance(n_out, int
            ) and n_out > 0, 'Output channels should be a positive integer got {}'.format(
            n_out)
        assert isinstance(ksize, int
            ) and ksize > 0, 'Kernel size should be a positive integer got {}'.format(
            ksize)
        padding = (ksize - 1) // 2 if pad else 0
        use_bias_in_conv = norm_layer is None
        self.add_module('conv', nn.Conv2d(n_in, n_out, ksize, stride=stride,
            padding=padding, bias=use_bias_in_conv))
        if norm_layer is not None:
            self.add_module('norm', _get_norm_layer(norm_layer, n_out))
        if activation is not None:
            self.add_module('activation', _get_activation(activation))
        _init_fc_or_conv(self.conv, activation)

    def forward(self, x):
        for c in self.children():
            x = c(x)
        return x


class FixupBasicBlock(nn.Module):
    expansion = 1

    def __init__(self, n_features, ksize=3, pad=True, activation='relu'):
        super(FixupBasicBlock, self).__init__()
        self.bias1a = nn.Parameter(th.zeros(1))
        self.conv1 = ConvModule(n_features, n_features, ksize=ksize, stride
            =1, pad=pad, activation=None, norm_layer=None)
        self.bias1b = nn.Parameter(th.zeros(1))
        self.activation = _get_activation(activation)
        self.bias2a = nn.Parameter(th.zeros(1))
        self.conv2 = ConvModule(n_features, n_features, ksize=ksize, stride
            =1, pad=pad, activation=None, norm_layer=None)
        self.scale = nn.Parameter(th.ones(1))
        self.bias2b = nn.Parameter(th.zeros(1))
        self.activation2 = _get_activation(activation)
        self.ksize = 3
        self.pad = pad

    def forward(self, x):
        identity = x
        out = self.conv1(x + self.bias1a)
        out = self.activation(out + self.bias1b)
        out = self.conv2(out + self.bias2a)
        out = out * self.scale + self.bias2b
        crop = (self.ksize - 1) // 2 * 2
        if crop > 0 and not self.pad:
            identity = identity[:, :, crop:-crop, crop:-crop]
        out += identity
        out = self.activation2(out)
        return out


class FixupResidualChainNew(nn.Module):
    """Linear chain of residual blocks.

    Args:
      n_features(int): number of input channels.
      depth(int): number of residual blocks
      ksize(int): size of the convolution kernel (square).
      activation(str): nonlinear activation function between convolutions.
      norm_layer(str): normalization to apply between the convolution modules.
      pad(bool): if True, zero pad the convs to maintain a constant size.
    """

    def __init__(self, n_features, depth=3, ksize=3, activation='relu',
        norm_layer=None, pad=True):
        super(FixupResidualChainNew, self).__init__()
        assert isinstance(n_features, int
            ) and n_features > 0, 'Number of feature channels should be a positive integer'
        assert isinstance(ksize, int) and ksize > 0 or isinstance(ksize, list
            ), 'Kernel size should be a positive integer or a list of integers'
        assert isinstance(depth, int
            ) and depth > 0 and depth < 16, 'Depth should be a positive integer lower than 16'
        self.depth = depth
        layers = OrderedDict()
        for lvl in range(depth):
            blockname = 'resblock{}'.format(lvl)
            layers[blockname] = FixupBasicBlock(n_features, ksize=ksize,
                activation=activation, pad=pad)
        self.net = nn.Sequential(layers)
        self._reset_weights()

    def _reset_weights(self):
        for m in self.net.modules():
            if isinstance(m, FixupBasicBlock):
                nn.init.normal_(m.conv1.conv.weight, mean=0, std=np.sqrt(2 /
                    (m.conv1.conv.weight.shape[0] * np.prod(m.conv1.conv.
                    weight.shape[2:]))) * self.depth ** -0.5)
                nn.init.constant_(m.conv2.conv.weight, 0)

    def forward(self, input_0):
        primals_2 = self.net.resblock0.bias1a
        primals_5 = self.net.resblock0.bias1b
        primals_6 = self.net.resblock0.bias2a
        primals_9 = self.net.resblock0.scale
        primals_10 = self.net.resblock0.bias2b
        primals_3 = self.net.resblock0.conv1.conv.weight
        primals_4 = self.net.resblock0.conv1.conv.bias
        primals_7 = self.net.resblock0.conv2.conv.weight
        primals_8 = self.net.resblock0.conv2.conv.bias
        primals_11 = self.net.resblock1.bias1a
        primals_14 = self.net.resblock1.bias1b
        primals_15 = self.net.resblock1.bias2a
        primals_18 = self.net.resblock1.scale
        primals_19 = self.net.resblock1.bias2b
        primals_12 = self.net.resblock1.conv1.conv.weight
        primals_13 = self.net.resblock1.conv1.conv.bias
        primals_16 = self.net.resblock1.conv2.conv.weight
        primals_17 = self.net.resblock1.conv2.conv.bias
        primals_20 = self.net.resblock2.bias1a
        primals_23 = self.net.resblock2.bias1b
        primals_24 = self.net.resblock2.bias2a
        primals_27 = self.net.resblock2.scale
        primals_28 = self.net.resblock2.bias2b
        primals_21 = self.net.resblock2.conv1.conv.weight
        primals_22 = self.net.resblock2.conv1.conv.bias
        primals_25 = self.net.resblock2.conv2.conv.weight
        primals_26 = self.net.resblock2.conv2.conv.bias
        primals_1 = input_0
        output = call([primals_1, primals_2, primals_3, primals_4,
            primals_5, primals_6, primals_7, primals_8, primals_9,
            primals_10, primals_11, primals_12, primals_13, primals_14,
            primals_15, primals_16, primals_17, primals_18, primals_19,
            primals_20, primals_21, primals_22, primals_23, primals_24,
            primals_25, primals_26, primals_27, primals_28])
        return output[0]


Measured latency_ms: 0.015053