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
import torch.nn.functional as F
import torch.utils.data
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_cat_elu_0(in_ptr0, out_ptr0, xnumel, XBLOCK: tl.constexpr
    ):
    xnumel = 512
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x1 = xindex // 16 % 8
    x0 = xindex % 16
    x2 = xindex // 128
    x3 = xindex
    tmp0 = x1
    tl.full([1], 0, tl.int64)
    tmp3 = tl.full([1], 4, tl.int64)
    tmp4 = tmp0 < tmp3
    tmp5 = tl.load(in_ptr0 + (x0 + 16 * x1 + 64 * x2), tmp4 & xmask, other=0.0)
    tmp6 = tmp0 >= tmp3
    tl.full([1], 8, tl.int64)
    tmp9 = tl.load(in_ptr0 + (x0 + 16 * (-4 + x1) + 64 * x2), tmp6 & xmask,
        other=0.0)
    tmp10 = -tmp9
    tmp11 = tl.full(tmp10.shape, 0.0, tmp10.dtype)
    tmp12 = tl.where(tmp6, tmp10, tmp11)
    tmp13 = tl.where(tmp4, tmp5, tmp12)
    tmp14 = 0.0
    tmp15 = tmp13 > tmp14
    tmp16 = 1.0
    tmp17 = tmp13 * tmp16
    tmp18 = libdevice.expm1(tmp17)
    tmp19 = tmp18 * tmp16
    tmp20 = tl.where(tmp15, tmp17, tmp19)
    tl.store(out_ptr0 + x3, tmp20, xmask)


@triton.jit
def triton_per_fused__weight_norm_interface_1(in_out_ptr0, in_ptr0, in_ptr1,
    out_ptr0, xnumel, rnumel, XBLOCK: tl.constexpr):
    xnumel = 4
    rnumel = 72
    RBLOCK: tl.constexpr = 128
    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 + 72 * 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 + 72 * x0), tmp9, rmask & xmask)


@triton.jit
def triton_poi_fused_cat_elu_2(in_ptr0, in_ptr1, out_ptr0, out_ptr1, xnumel,
    XBLOCK: tl.constexpr):
    xnumel = 512
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x1 = xindex // 16 % 8
    x0 = xindex % 16
    x2 = xindex // 128
    x3 = xindex
    tmp0 = x1
    tl.full([1], 0, tl.int64)
    tmp3 = tl.full([1], 4, tl.int64)
    tmp4 = tmp0 < tmp3
    tmp5 = tl.load(in_ptr0 + (x0 + 16 * x1 + 64 * x2), tmp4 & xmask, other=0.0)
    tmp6 = tl.load(in_ptr1 + x1, tmp4 & xmask, eviction_policy='evict_last',
        other=0.0)
    tmp7 = tmp5 + tmp6
    tmp8 = tl.full(tmp7.shape, 0.0, tmp7.dtype)
    tmp9 = tl.where(tmp4, tmp7, tmp8)
    tmp10 = tmp0 >= tmp3
    tl.full([1], 8, tl.int64)
    tmp13 = tl.load(in_ptr0 + (x0 + 16 * (-4 + x1) + 64 * x2), tmp10 &
        xmask, other=0.0)
    tmp14 = tl.load(in_ptr1 + (-4 + x1), tmp10 & xmask, eviction_policy=
        'evict_last', other=0.0)
    tmp15 = tmp13 + tmp14
    tmp16 = -tmp15
    tmp17 = tl.full(tmp16.shape, 0.0, tmp16.dtype)
    tmp18 = tl.where(tmp10, tmp16, tmp17)
    tmp19 = tl.where(tmp4, tmp9, tmp18)
    tmp20 = 0.0
    tmp21 = tmp19 > tmp20
    tmp22 = 1.0
    tmp23 = tmp19 * tmp22
    tmp24 = libdevice.expm1(tmp23)
    tmp25 = tmp24 * tmp22
    tmp26 = tl.where(tmp21, tmp23, tmp25)
    tl.store(out_ptr0 + x3, tmp19, xmask)
    tl.store(out_ptr1 + x3, tmp26, xmask)


@triton.jit
def triton_per_fused__weight_norm_interface_3(in_out_ptr0, in_ptr0, in_ptr1,
    out_ptr0, xnumel, rnumel, XBLOCK: tl.constexpr):
    xnumel = 8
    RBLOCK: tl.constexpr = 8
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    rindex = tl.arange(0, RBLOCK)[None, :]
    tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r1 = rindex
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + (r1 + 8 * x0), 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(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 + 8 * x0), tmp9, xmask)


@triton.jit
def triton_poi_fused_convolution_4(in_out_ptr0, in_ptr0, xnumel, XBLOCK: tl
    .constexpr):
    xnumel = 512
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x3 = xindex
    x1 = xindex // 16 % 8
    tmp0 = tl.load(in_out_ptr0 + x3, xmask)
    tmp1 = tl.load(in_ptr0 + x1, xmask, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tl.store(in_out_ptr0 + x3, tmp2, xmask)


@triton.jit
def triton_poi_fused_mul_sigmoid_5(in_ptr0, 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
    x0 = xindex % 64
    x1 = xindex // 64
    x2 = xindex
    tmp0 = tl.load(in_ptr0 + (64 + x0 + 128 * x1), xmask)
    tmp2 = tl.load(in_ptr0 + (x0 + 128 * x1), xmask)
    tmp1 = tl.sigmoid(tmp0)
    tmp3 = tmp2 * tmp1
    tl.store(out_ptr0 + x2, tmp1, xmask)
    tl.store(out_ptr1 + x2, tmp3, xmask)


def call(args):
    (primals_1, primals_2, primals_3, primals_4, primals_5, primals_6,
        primals_7) = args
    args.clear()
    assert_size_stride(primals_1, (4, 4, 4, 4), (64, 16, 4, 1))
    assert_size_stride(primals_2, (4, 1, 1, 1), (1, 1, 1, 1))
    assert_size_stride(primals_3, (4, 8, 3, 3), (72, 9, 3, 1))
    assert_size_stride(primals_4, (4,), (1,))
    assert_size_stride(primals_5, (8, 1, 1, 1), (1, 1, 1, 1))
    assert_size_stride(primals_6, (8, 8, 1, 1), (8, 1, 1, 1))
    assert_size_stride(primals_7, (8,), (1,))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((4, 8, 4, 4), (128, 16, 4, 1), torch.float32)
        get_raw_stream(0)
        triton_poi_fused_cat_elu_0[grid(512)](primals_1, buf0, 512, XBLOCK=
            128, num_warps=4, num_stages=1)
        del primals_1
        buf1 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf2 = reinterpret_tensor(buf1, (4, 1, 1, 1), (1, 1, 1, 1), 0)
        del buf1
        buf3 = empty_strided_cuda((4, 8, 3, 3), (72, 9, 3, 1), torch.float32)
        triton_per_fused__weight_norm_interface_1[grid(4)](buf2, primals_3,
            primals_2, buf3, 4, 72, XBLOCK=1, num_warps=2, num_stages=1)
        buf4 = extern_kernels.convolution(buf0, buf3, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf4, (4, 4, 4, 4), (64, 16, 4, 1))
        buf5 = empty_strided_cuda((4, 8, 4, 4), (128, 16, 4, 1), torch.float32)
        buf6 = empty_strided_cuda((4, 8, 4, 4), (128, 16, 4, 1), torch.float32)
        triton_poi_fused_cat_elu_2[grid(512)](buf4, primals_4, buf5, buf6, 
            512, XBLOCK=128, num_warps=4, num_stages=1)
        del primals_4
        buf7 = empty_strided_cuda((8, 1, 1, 1), (1, 8, 8, 8), torch.float32)
        buf8 = reinterpret_tensor(buf7, (8, 1, 1, 1), (1, 1, 1, 1), 0)
        del buf7
        buf9 = empty_strided_cuda((8, 8, 1, 1), (8, 1, 1, 1), torch.float32)
        triton_per_fused__weight_norm_interface_3[grid(8)](buf8, primals_6,
            primals_5, buf9, 8, 8, XBLOCK=8, num_warps=2, num_stages=1)
        buf10 = extern_kernels.convolution(buf6, buf9, stride=(1, 1),
            padding=(0, 0), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf10, (4, 8, 4, 4), (128, 16, 4, 1))
        buf11 = buf10
        del buf10
        triton_poi_fused_convolution_4[grid(512)](buf11, primals_7, 512,
            XBLOCK=128, num_warps=4, num_stages=1)
        del primals_7
        buf12 = buf4
        del buf4
        buf13 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.float32)
        triton_poi_fused_mul_sigmoid_5[grid(256)](buf11, buf12, buf13, 256,
            XBLOCK=128, num_warps=4, num_stages=1)
    return (buf13, buf3, buf9, primals_2, primals_3, primals_5, primals_6,
        buf0, buf2, buf3, buf5, buf6, buf8, buf9, reinterpret_tensor(buf11,
        (4, 4, 4, 4), (128, 16, 4, 1), 0), buf12)


def concat_elu(x):
    """Concatenated ReLU (http://arxiv.org/abs/1603.05201), but with ELU."""
    return F.elu(torch.cat((x, -x), dim=1))


class WNConv2d(nn.Module):
    """Weight-normalized 2d convolution.
	Args:
		in_channels (int): Number of channels in the input.
		out_channels (int): Number of channels in the output.
		kernel_size (int): Side length of each convolutional kernel.
		padding (int): Padding to add on edges of input.
		bias (bool): Use bias in the convolution operation.
	"""

    def __init__(self, in_channels, out_channels, kernel_size, padding,
        bias=True):
        super(WNConv2d, self).__init__()
        self.conv = nn.utils.weight_norm(nn.Conv2d(in_channels,
            out_channels, kernel_size, padding=padding, bias=bias))

    def forward(self, x):
        x = self.conv(x)
        return x


class GatedConvNew(nn.Module):
    """Gated Convolution Block

	Originally used by PixelCNN++ (https://arxiv.org/pdf/1701.05517).

	Args:
		num_channels (int): Number of channels in hidden activations.
		drop_prob (float): Dropout probability.
		aux_channels (int): Number of channels in optional auxiliary input.
	"""

    def __init__(self, num_channels, drop_prob=0.0, aux_channels=None):
        super(GatedConvNew, self).__init__()
        self.nlin = concat_elu
        self.conv = WNConv2d(2 * num_channels, num_channels, kernel_size=3,
            padding=1)
        self.drop = nn.Dropout2d(drop_prob)
        self.gate = WNConv2d(2 * num_channels, 2 * num_channels,
            kernel_size=1, padding=0)
        if aux_channels is not None:
            self.aux_conv = WNConv2d(2 * aux_channels, num_channels,
                kernel_size=1, padding=0)
        else:
            self.aux_conv = None

    def forward(self, input_0):
        primals_4 = self.conv.conv.bias
        primals_2 = self.conv.conv.weight_g
        primals_3 = self.conv.conv.weight_v
        primals_7 = self.gate.conv.bias
        primals_5 = self.gate.conv.weight_g
        primals_6 = self.gate.conv.weight_v
        primals_1 = input_0
        output = call([primals_1, primals_2, primals_3, primals_4,
            primals_5, primals_6, primals_7])
        return output[0]


Measured latency_ms: 0.015108