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
from torch._inductor.runtime.triton_helpers import libdevice
import torch.nn as nn
from abc import ABC
from torch.nn.utils.weight_norm import weight_norm
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_per_fused__weight_norm_interface_0(in_out_ptr0, in_ptr0, in_ptr1,
    out_ptr0, xnumel, rnumel, XBLOCK: tl.constexpr):
    xnumel = 4
    rnumel = 36
    RBLOCK: tl.constexpr = 64
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    rindex = tl.arange(0, RBLOCK)[None, :]
    rmask = rindex < rnumel
    r1 = rindex
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + (r1 + 36 * x0), rmask & xmask, other=0.0)
    tmp7 = tl.load(in_ptr1 + x0, xmask, eviction_policy='evict_last')
    tmp1 = tmp0 * tmp0
    tmp2 = tl.broadcast_to(tmp1, [XBLOCK, RBLOCK])
    tmp4 = tl.where(rmask & xmask, tmp2, 0)
    tmp5 = tl.sum(tmp4, 1)[:, None]
    tmp6 = libdevice.sqrt(tmp5)
    tmp8 = tmp7 / tmp6
    tmp9 = tmp0 * tmp8
    tl.debug_barrier()
    tl.store(in_out_ptr0 + x0, tmp6, xmask)
    tl.store(out_ptr0 + (r1 + 36 * x0), tmp9, rmask & xmask)


@triton.jit
def triton_poi_fused_relu_1(in_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_out_ptr0 + x0, xmask)
    tmp1 = tl.full([1], 0, tl.int32)
    tmp2 = triton_helpers.maximum(tmp1, tmp0)
    tl.store(in_out_ptr0 + x0, tmp2, xmask)


@triton.jit
def triton_poi_fused_add_relu_threshold_backward_2(in_out_ptr0, in_ptr0,
    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_out_ptr0 + x0, xmask)
    tmp1 = tl.load(in_ptr0 + x0, xmask)
    tmp2 = tmp0 + tmp1
    tmp3 = tl.full([1], 0, tl.int32)
    tmp4 = triton_helpers.maximum(tmp3, tmp2)
    tmp5 = 0.0
    tmp6 = tmp4 <= tmp5
    tl.store(in_out_ptr0 + x0, tmp4, xmask)
    tl.store(out_ptr0 + x0, tmp6, xmask)


def call(args):
    primals_1, primals_2, primals_3, primals_4, primals_5 = args
    args.clear()
    assert_size_stride(primals_1, (4, 1, 1, 1), (1, 1, 1, 1))
    assert_size_stride(primals_2, (4, 4, 3, 3), (36, 9, 3, 1))
    assert_size_stride(primals_3, (4, 4, 4, 4), (64, 16, 4, 1))
    assert_size_stride(primals_4, (4, 1, 1, 1), (1, 1, 1, 1))
    assert_size_stride(primals_5, (4, 4, 3, 3), (36, 9, 3, 1))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf1 = reinterpret_tensor(buf0, (4, 1, 1, 1), (1, 1, 1, 1), 0)
        del buf0
        buf2 = empty_strided_cuda((4, 4, 3, 3), (36, 9, 3, 1), torch.float32)
        get_raw_stream(0)
        triton_per_fused__weight_norm_interface_0[grid(4)](buf1, primals_2,
            primals_1, buf2, 4, 36, XBLOCK=1, num_warps=2, num_stages=1)
        buf3 = extern_kernels.convolution(primals_3, buf2, 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
        triton_poi_fused_relu_1[grid(256)](buf4, 256, XBLOCK=128, num_warps
            =4, num_stages=1)
        buf5 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf6 = reinterpret_tensor(buf5, (4, 1, 1, 1), (1, 1, 1, 1), 0)
        del buf5
        buf7 = empty_strided_cuda((4, 4, 3, 3), (36, 9, 3, 1), torch.float32)
        triton_per_fused__weight_norm_interface_0[grid(4)](buf6, primals_5,
            primals_4, buf7, 4, 36, XBLOCK=1, num_warps=2, num_stages=1)
        buf8 = extern_kernels.convolution(buf4, buf7, 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 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.bool)
        triton_poi_fused_add_relu_threshold_backward_2[grid(256)](buf9,
            primals_3, buf10, 256, XBLOCK=128, num_warps=4, num_stages=1)
    return (buf9, buf2, buf7, primals_1, primals_2, primals_3, primals_4,
        primals_5, buf1, buf2, buf4, buf6, buf7, buf10)


def conv1x1(in_planes, out_planes, stride=1):
    """
    Create a 1x1 2d convolution block
    """
    return nn.Conv2d(in_planes, out_planes, kernel_size=1, stride=stride,
        bias=False)


def conv3x3(in_planes, out_planes, stride=1):
    """
    Create a 3x3 2d convolution block
    """
    return nn.Conv2d(in_planes, out_planes, kernel_size=3, stride=stride,
        padding=1, bias=False)


class NeuralNetworkModule(nn.Module, ABC):
    """
    Note: input device and output device are determined by module parameters,
          your input module / output submodule should not store parameters on
          more than one device, and you also should not move your output to
          other devices other than your parameter storage device in forward().
    """

    def __init__(self):
        super().__init__()
        self.input_module = None
        self.output_module = None

    def set_input_module(self, input_module: 'nn.Module'):
        """
        Set the input submodule of current module.
        """
        self.input_module = input_module
        if not isinstance(input_module, NeuralNetworkModule):
            if isinstance(input_module, nn.Sequential):
                input_module = self.find_child(input_module, True)
            if len({p.device for p in input_module.parameters()}) > 1:
                raise RuntimeError(
                    'Input module must be another NeuralNetworkModule or locate on one single device.'
                    )

    def set_output_module(self, output_module: 'nn.Module'):
        """
        Set the output submodule of current module.
        """
        self.output_module = output_module
        if not isinstance(output_module, NeuralNetworkModule):
            if isinstance(output_module, nn.Sequential):
                output_module = self.find_child(output_module, False)
            if len({p.device for p in output_module.parameters()}) > 1:
                raise RuntimeError(
                    'Output module must be another NeuralNetworkModule or locate on one single device.'
                    )

    @property
    def input_device(self):
        if self.input_module is None:
            raise RuntimeError('Input module not set.')
        elif not isinstance(self.input_module, NeuralNetworkModule):
            dev_set = {p.device for p in self.input_module.parameters()}
            if len(dev_set) != 1:
                raise RuntimeError(
                    'This input module contains parameters on different devices, please consider about splitting it.'
                    )
            else:
                return list(dev_set)[0]
        else:
            return self.input_module.input_device

    @property
    def output_device(self):
        if self.output_module is None and self.input_module is None:
            raise RuntimeError('Output module not set.')
        elif self.output_module is not None:
            if not isinstance(self.output_module, NeuralNetworkModule):
                dev_set = {p.device for p in self.output_module.parameters()}
                if len(dev_set) != 1:
                    raise RuntimeError(
                        'This output module contains parameters on different devices, please consider about splitting it.'
                        )
                else:
                    return list(dev_set)[0]
            else:
                return self.output_module.output_device
        else:
            return self.input_device

    @staticmethod
    def find_child(seq, is_first=True):
        """
        Find the first / last leaf child module.
        """
        if isinstance(seq, nn.Sequential):
            if is_first:
                return NeuralNetworkModule.find_child(seq[0], is_first)
            else:
                return NeuralNetworkModule.find_child(seq[-1], is_first)
        else:
            return seq

    def forward(self, *_):
        pass


class BasicBlockWNNew(NeuralNetworkModule):
    """
    Basic block with weight normalization
    """
    expansion = 1

    def __init__(self, in_planes, out_planes, stride=1, **__):
        """
        Create a basic block of resnet.

        Args:
            in_planes:  Number of input planes.
            out_planes: Number of output planes.
            stride:     Stride of convolution.
        """
        super().__init__()
        self.conv1 = weight_norm(conv3x3(in_planes, out_planes, stride))
        self.conv2 = weight_norm(conv3x3(out_planes, self.expansion *
            out_planes))
        self.shortcut = nn.Sequential()
        self.set_input_module(self.conv1)
        if stride != 1 or in_planes != self.expansion * out_planes:
            self.shortcut = nn.Sequential(weight_norm(conv1x1(in_planes, 
                self.expansion * out_planes, stride)))

    def forward(self, input_0):
        primals_1 = self.conv1.weight_g
        primals_2 = self.conv1.weight_v
        primals_4 = self.conv2.weight_g
        primals_5 = self.conv2.weight_v
        primals_3 = input_0
        output = call([primals_1, primals_2, primals_3, primals_4, primals_5])
        return output[0]


Measured latency_ms: 0.06338