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.triton_helpers import libdevice
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_constant_pad_nd_0(in_ptr0, out_ptr0, xnumel, XBLOCK:
    tl.constexpr):
    xnumel = 1024
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x1 = xindex // 8 % 8
    x0 = xindex % 8
    x2 = xindex // 64
    x4 = xindex
    tmp0 = -2 + x1
    tmp1 = tl.full([1], 0, tl.int64)
    tmp2 = tmp0 >= tmp1
    tmp3 = tl.full([1], 4, tl.int64)
    tmp4 = tmp0 < tmp3
    tmp5 = -2 + x0
    tmp6 = tmp5 >= tmp1
    tmp7 = tmp5 < tmp3
    tmp8 = tmp2 & tmp4
    tmp9 = tmp8 & tmp6
    tmp10 = tmp9 & tmp7
    tmp11 = tl.load(in_ptr0 + (-10 + x0 + 4 * x1 + 16 * x2), tmp10 & xmask,
        other=0.0)
    tl.store(out_ptr0 + x4, tmp11, xmask)


@triton.jit
def triton_per_fused_convolution_native_group_norm_pixel_shuffle_1(in_out_ptr0,
    in_out_ptr1, in_ptr0, in_ptr1, in_ptr2, out_ptr0, out_ptr2, xnumel,
    rnumel, XBLOCK: tl.constexpr):
    xnumel = 64
    rnumel = 25
    RBLOCK: tl.constexpr = 32
    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
    r2 = rindex
    x3 = xindex
    x0 = xindex % 16
    r7 = rindex % 5
    r8 = rindex // 5
    x4 = xindex % 2
    x5 = xindex // 2 % 2
    x6 = xindex // 4
    tmp0 = tl.load(in_out_ptr0 + (r2 + 25 * x3), rmask & xmask, other=0.0)
    tmp1 = tl.load(in_ptr0 + x0, xmask, eviction_policy='evict_last')
    tmp26 = tl.load(in_ptr1 + x0, xmask, eviction_policy='evict_last')
    tmp28 = tl.load(in_ptr2 + x0, xmask, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tmp3 = tl.broadcast_to(tmp2, [XBLOCK, RBLOCK])
    tl.where(rmask & xmask, tmp3, 0)
    tmp6 = tl.broadcast_to(tmp3, [XBLOCK, RBLOCK])
    tmp8 = tl.where(rmask & xmask, tmp6, 0)
    tmp9 = tl.sum(tmp8, 1)[:, None]
    tmp10 = tl.full([XBLOCK, 1], 25, tl.int32)
    tmp11 = tmp10.to(tl.float32)
    tmp12 = tmp9 / tmp11
    tmp13 = tmp3 - tmp12
    tmp14 = tmp13 * tmp13
    tmp15 = tl.broadcast_to(tmp14, [XBLOCK, RBLOCK])
    tmp17 = tl.where(rmask & xmask, tmp15, 0)
    tmp18 = tl.sum(tmp17, 1)[:, None]
    tmp19 = 25.0
    tmp20 = tmp18 / tmp19
    tmp21 = 1e-05
    tmp22 = tmp20 + tmp21
    tmp23 = libdevice.rsqrt(tmp22)
    tmp24 = tmp2 - tmp12
    tmp25 = tmp24 * tmp23
    tmp27 = tmp25 * tmp26
    tmp29 = tmp27 + tmp28
    tmp30 = 0.0
    tmp31 = tmp29 > tmp30
    tmp32 = 1.0
    tmp33 = tmp29 * tmp32
    tmp34 = libdevice.expm1(tmp33)
    tmp35 = tmp34 * tmp32
    tmp36 = tl.where(tmp31, tmp33, tmp35)
    tl.store(in_out_ptr0 + (r2 + 25 * x3), tmp2, rmask & xmask)
    tl.debug_barrier()
    tl.store(in_out_ptr1 + x3, tmp23, xmask)
    tl.store(out_ptr2 + (x4 + 2 * r7 + 10 * x5 + 20 * r8 + 100 * x6), tmp36,
        rmask & xmask)
    tl.store(out_ptr0 + x3, tmp12, xmask)


def call(args):
    primals_1, primals_2, primals_3, primals_4, primals_5 = args
    args.clear()
    assert_size_stride(primals_1, (4, 4, 4, 4), (64, 16, 4, 1))
    assert_size_stride(primals_2, (16, 4, 4, 4), (64, 16, 4, 1))
    assert_size_stride(primals_3, (16,), (1,))
    assert_size_stride(primals_4, (16,), (1,))
    assert_size_stride(primals_5, (16,), (1,))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((4, 4, 8, 8), (256, 64, 8, 1), torch.float32)
        get_raw_stream(0)
        triton_poi_fused_constant_pad_nd_0[grid(1024)](primals_1, buf0, 
            1024, XBLOCK=128, num_warps=4, num_stages=1)
        del primals_1
        buf1 = extern_kernels.convolution(buf0, primals_2, stride=(1, 1),
            padding=(0, 0), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf1, (4, 16, 5, 5), (400, 25, 5, 1))
        buf2 = buf1
        del buf1
        buf3 = empty_strided_cuda((4, 16, 1, 1), (16, 1, 1, 1), torch.float32)
        buf4 = empty_strided_cuda((4, 16, 1, 1), (16, 1, 64, 64), torch.float32
            )
        buf6 = reinterpret_tensor(buf4, (4, 16, 1, 1), (16, 1, 1, 1), 0)
        del buf4
        buf8 = empty_strided_cuda((4, 4, 5, 2, 5, 2), (400, 100, 20, 10, 2,
            1), torch.float32)
        triton_per_fused_convolution_native_group_norm_pixel_shuffle_1[grid(64)
            ](buf2, buf6, primals_3, primals_4, primals_5, buf3, buf8, 64, 
            25, XBLOCK=8, num_warps=2, num_stages=1)
        del primals_3
    return reinterpret_tensor(buf8, (4, 4, 10, 10), (400, 100, 10, 1), 0
        ), primals_2, primals_4, primals_5, buf0, buf2, buf3, buf6


class Conv2D(nn.Module):
    """
    2D convolution with GroupNorm and ELU

    Parameters
    ----------
    in_channels : int
        Number of input channels
    out_channels : int
        Number of output channels
    kernel_size : int
        Kernel size
    stride : int
        Stride
    """

    def __init__(self, in_channels, out_channels, kernel_size, stride):
        super().__init__()
        self.kernel_size = kernel_size
        self.conv_base = nn.Conv2d(in_channels, out_channels, kernel_size=
            kernel_size, stride=stride)
        self.pad = nn.ConstantPad2d([kernel_size // 2] * 4, value=0)
        self.normalize = torch.nn.GroupNorm(16, out_channels)
        self.activ = nn.ELU(inplace=True)

    def forward(self, x):
        """Runs the Conv2D layer."""
        x = self.conv_base(self.pad(x))
        return self.activ(self.normalize(x))


class UnpackLayerConv2dNew(nn.Module):
    """
    Unpacking layer with 2d convolutions. Takes a [B,C,H,W] tensor, convolves it
    to produce [B,(r^2)C,H,W] and then unpacks it to produce [B,C,rH,rW].
    """

    def __init__(self, in_channels, out_channels, kernel_size, r=2):
        """
        Initializes a UnpackLayerConv2d object.

        Parameters
        ----------
        in_channels : int
            Number of input channels
        out_channels : int
            Number of output channels
        kernel_size : int
            Kernel size
        r : int
            Packing ratio
        """
        super().__init__()
        self.conv = Conv2D(in_channels, out_channels * r ** 2, kernel_size, 1)
        self.unpack = nn.PixelShuffle(r)

    def forward(self, input_0):
        primals_2 = self.conv.conv_base.weight
        primals_3 = self.conv.conv_base.bias
        primals_4 = self.conv.normalize.weight
        primals_5 = self.conv.normalize.bias
        primals_1 = input_0
        output = call([primals_1, primals_2, primals_3, primals_4, primals_5])
        return output[0]


Measured latency_ms: 0.043832