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
from torch import 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_per_fused_elu_native_group_norm_0(in_out_ptr0, in_ptr0, in_ptr1,
    in_ptr2, out_ptr0, out_ptr2, xnumel, rnumel, XBLOCK: tl.constexpr):
    xnumel = 4
    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, :]
    tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r1 = rindex
    x0 = xindex
    r3 = rindex // 16
    tmp0 = tl.load(in_ptr0 + (r1 + 64 * x0), xmask, other=0.0)
    tmp24 = tl.load(in_ptr1 + r3, None, eviction_policy='evict_last')
    tmp26 = tl.load(in_ptr2 + r3, None, eviction_policy='evict_last')
    tmp1 = tl.broadcast_to(tmp0, [XBLOCK, RBLOCK])
    tl.where(xmask, tmp1, 0)
    tmp4 = tl.broadcast_to(tmp1, [XBLOCK, RBLOCK])
    tmp6 = tl.where(xmask, tmp4, 0)
    tmp7 = tl.sum(tmp6, 1)[:, None]
    tmp8 = tl.full([XBLOCK, 1], 64, tl.int32)
    tmp9 = tmp8.to(tl.float32)
    tmp10 = tmp7 / tmp9
    tmp11 = tmp1 - tmp10
    tmp12 = tmp11 * tmp11
    tmp13 = tl.broadcast_to(tmp12, [XBLOCK, RBLOCK])
    tmp15 = tl.where(xmask, tmp13, 0)
    tmp16 = tl.sum(tmp15, 1)[:, None]
    tmp17 = tmp0 - tmp10
    tmp18 = 64.0
    tmp19 = tmp16 / tmp18
    tmp20 = 1e-05
    tmp21 = tmp19 + tmp20
    tmp22 = libdevice.rsqrt(tmp21)
    tmp23 = tmp17 * tmp22
    tmp25 = tmp23 * tmp24
    tmp27 = tmp25 + tmp26
    tmp28 = 0.0
    tmp29 = tmp27 > tmp28
    tmp30 = 1.0
    tmp31 = tmp27 * tmp30
    tmp32 = libdevice.expm1(tmp31)
    tmp33 = tmp32 * tmp30
    tmp34 = tl.where(tmp29, tmp31, tmp33)
    tl.store(in_out_ptr0 + (r1 + 64 * x0), tmp34, xmask)
    tl.store(out_ptr2 + x0, tmp22, xmask)
    tl.store(out_ptr0 + x0, tmp10, xmask)


@triton.jit
def triton_per_fused_add_elu_native_group_norm_1(in_out_ptr0, in_ptr0,
    in_ptr1, in_ptr2, in_ptr3, out_ptr0, out_ptr2, xnumel, rnumel, XBLOCK:
    tl.constexpr):
    xnumel = 4
    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, :]
    tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r1 = rindex
    x0 = xindex
    r3 = rindex // 16
    tmp0 = tl.load(in_ptr0 + (r1 + 64 * x0), xmask, other=0.0)
    tmp24 = tl.load(in_ptr1 + r3, None, eviction_policy='evict_last')
    tmp26 = tl.load(in_ptr2 + r3, None, eviction_policy='evict_last')
    tmp28 = tl.load(in_ptr3 + (r1 + 64 * x0), xmask, other=0.0)
    tmp1 = tl.broadcast_to(tmp0, [XBLOCK, RBLOCK])
    tl.where(xmask, tmp1, 0)
    tmp4 = tl.broadcast_to(tmp1, [XBLOCK, RBLOCK])
    tmp6 = tl.where(xmask, tmp4, 0)
    tmp7 = tl.sum(tmp6, 1)[:, None]
    tmp8 = tl.full([XBLOCK, 1], 64, tl.int32)
    tmp9 = tmp8.to(tl.float32)
    tmp10 = tmp7 / tmp9
    tmp11 = tmp1 - tmp10
    tmp12 = tmp11 * tmp11
    tmp13 = tl.broadcast_to(tmp12, [XBLOCK, RBLOCK])
    tmp15 = tl.where(xmask, tmp13, 0)
    tmp16 = tl.sum(tmp15, 1)[:, None]
    tmp17 = tmp0 - tmp10
    tmp18 = 64.0
    tmp19 = tmp16 / tmp18
    tmp20 = 1e-05
    tmp21 = tmp19 + tmp20
    tmp22 = libdevice.rsqrt(tmp21)
    tmp23 = tmp17 * tmp22
    tmp25 = tmp23 * tmp24
    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 + (r1 + 64 * x0), tmp36, xmask)
    tl.store(out_ptr2 + x0, tmp22, xmask)
    tl.store(out_ptr0 + x0, tmp10, xmask)


def call(args):
    (primals_1, primals_2, primals_3, primals_4, primals_5, primals_6,
        primals_7, primals_8, primals_9, primals_10) = args
    args.clear()
    assert_size_stride(primals_1, (4, 4, 3, 3, 3), (108, 27, 9, 3, 1))
    assert_size_stride(primals_2, (4, 4, 4, 4), (64, 16, 4, 1))
    assert_size_stride(primals_3, (4,), (1,))
    assert_size_stride(primals_4, (4,), (1,))
    assert_size_stride(primals_5, (4, 4, 3, 3, 3), (108, 27, 9, 3, 1))
    assert_size_stride(primals_6, (4,), (1,))
    assert_size_stride(primals_7, (4,), (1,))
    assert_size_stride(primals_8, (4, 4, 3, 3, 3), (108, 27, 9, 3, 1))
    assert_size_stride(primals_9, (4,), (1,))
    assert_size_stride(primals_10, (4,), (1,))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = extern_kernels.convolution(reinterpret_tensor(primals_2, (1,
            4, 4, 4, 4), (256, 64, 16, 4, 1), 0), primals_1, stride=(1, 1, 
            1), padding=(1, 1, 1), dilation=(1, 1, 1), transposed=False,
            output_padding=(0, 0, 0), groups=1, bias=None)
        assert_size_stride(buf0, (1, 4, 4, 4, 4), (256, 64, 16, 4, 1))
        buf1 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf4 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.float32)
        buf6 = buf4
        del buf4
        buf5 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        get_raw_stream(0)
        triton_per_fused_elu_native_group_norm_0[grid(4)](buf6, buf0,
            primals_3, primals_4, buf1, buf5, 4, 64, XBLOCK=1, num_warps=2,
            num_stages=1)
        del primals_4
        buf7 = extern_kernels.convolution(reinterpret_tensor(buf6, (1, 4, 4,
            4, 4), (256, 64, 16, 4, 1), 0), primals_5, stride=(1, 1, 1),
            padding=(1, 1, 1), dilation=(1, 1, 1), transposed=False,
            output_padding=(0, 0, 0), groups=1, bias=None)
        assert_size_stride(buf7, (1, 4, 4, 4, 4), (256, 64, 16, 4, 1))
        buf8 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf11 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.float32)
        buf13 = buf11
        del buf11
        buf12 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        triton_per_fused_elu_native_group_norm_0[grid(4)](buf13, buf7,
            primals_6, primals_7, buf8, buf12, 4, 64, XBLOCK=1, num_warps=2,
            num_stages=1)
        del primals_7
        buf14 = extern_kernels.convolution(reinterpret_tensor(buf13, (1, 4,
            4, 4, 4), (256, 64, 16, 4, 1), 0), primals_8, stride=(1, 1, 1),
            padding=(1, 1, 1), dilation=(1, 1, 1), transposed=False,
            output_padding=(0, 0, 0), groups=1, bias=None)
        assert_size_stride(buf14, (1, 4, 4, 4, 4), (256, 64, 16, 4, 1))
        buf15 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf19 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.float32)
        buf20 = buf19
        del buf19
        buf18 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        triton_per_fused_add_elu_native_group_norm_1[grid(4)](buf20, buf14,
            primals_9, primals_10, buf6, buf15, buf18, 4, 64, XBLOCK=1,
            num_warps=2, num_stages=1)
        del primals_10
    return (buf20, primals_1, primals_3, primals_5, primals_6, primals_8,
        primals_9, reinterpret_tensor(primals_2, (1, 4, 4, 4, 4), (256, 64,
        16, 4, 1), 0), buf0, reinterpret_tensor(buf1, (4, 1), (1, 1), 0),
        reinterpret_tensor(buf5, (4, 1), (1, 1), 0), buf6, buf7,
        reinterpret_tensor(buf8, (4, 1), (1, 1), 0), reinterpret_tensor(
        buf12, (4, 1), (1, 1), 0), buf13, buf14, reinterpret_tensor(buf15,
        (4, 1), (1, 1), 0), reinterpret_tensor(buf18, (4, 1), (1, 1), 0), buf20
        )


def conv3d(in_channels, out_channels, kernel_size, bias, padding):
    return nn.Conv3d(in_channels, out_channels, kernel_size, padding=
        padding, bias=bias)


def create_conv(in_channels, out_channels, kernel_size, order, num_groups,
    padding):
    """
    Create a list of modules with together constitute a single conv layer with non-linearity
    and optional batchnorm/groupnorm.

    Args:
        in_channels (int): number of input channels
        out_channels (int): number of output channels
        kernel_size(int or tuple): size of the convolving kernel
        order (string): order of things, e.g.
            'cr' -> conv + ReLU
            'gcr' -> groupnorm + conv + ReLU
            'cl' -> conv + LeakyReLU
            'ce' -> conv + ELU
            'bcr' -> batchnorm + conv + ReLU
        num_groups (int): number of groups for the GroupNorm
        padding (int or tuple): add zero-padding added to all three sides of the input

    Return:
        list of tuple (name, module)
    """
    assert 'c' in order, 'Conv layer MUST be present'
    assert order[0
        ] not in 'rle', 'Non-linearity cannot be the first operation in the layer'
    modules = []
    for i, char in enumerate(order):
        if char == 'r':
            modules.append(('ReLU', nn.ReLU(inplace=True)))
        elif char == 'l':
            modules.append(('LeakyReLU', nn.LeakyReLU(negative_slope=0.1,
                inplace=True)))
        elif char == 'e':
            modules.append(('ELU', nn.ELU(inplace=True)))
        elif char == 'c':
            bias = not ('g' in order or 'b' in order)
            modules.append(('conv', conv3d(in_channels, out_channels,
                kernel_size, bias, padding=padding)))
        elif char == 'g':
            is_before_conv = i < order.index('c')
            if is_before_conv:
                num_channels = in_channels
            else:
                num_channels = out_channels
            if num_channels < num_groups:
                num_groups = 1
            assert num_channels % num_groups == 0, f'Expected number of channels in input to be divisible by num_groups. num_channels={num_channels}, num_groups={num_groups}'
            modules.append(('groupnorm', nn.GroupNorm(num_groups=num_groups,
                num_channels=num_channels)))
        elif char == 'b':
            is_before_conv = i < order.index('c')
            if is_before_conv:
                modules.append(('batchnorm', nn.BatchNorm3d(in_channels)))
            else:
                modules.append(('batchnorm', nn.BatchNorm3d(out_channels)))
        else:
            raise ValueError(
                f"Unsupported layer type '{char}'. MUST be one of ['b', 'g', 'r', 'l', 'e', 'c']"
                )
    return modules


class SingleConv(nn.Sequential):
    """
    Basic convolutional module consisting of a Conv3d, non-linearity and optional batchnorm/groupnorm. The order
    of operations can be specified via the `order` parameter

    Args:
        in_channels (int): number of input channels
        out_channels (int): number of output channels
        kernel_size (int or tuple): size of the convolving kernel
        order (string): determines the order of layers, e.g.
            'cr' -> conv + ReLU
            'crg' -> conv + ReLU + groupnorm
            'cl' -> conv + LeakyReLU
            'ce' -> conv + ELU
        num_groups (int): number of groups for the GroupNorm
        padding (int or tuple):
    """

    def __init__(self, in_channels, out_channels, kernel_size=3, order=
        'gcr', num_groups=8, padding=1):
        super(SingleConv, self).__init__()
        for name, module in create_conv(in_channels, out_channels,
            kernel_size, order, num_groups, padding=padding):
            self.add_module(name, module)


class ExtResNetBlockNew(nn.Module):
    """
    Basic UNet block consisting of a SingleConv followed by the residual block.
    The SingleConv takes care of increasing/decreasing the number of channels and also ensures that the number
    of output channels is compatible with the residual block that follows.
    This block can be used instead of standard DoubleConv in the Encoder module.
    Motivated by: https://arxiv.org/pdf/1706.00120.pdf

    Notice we use ELU instead of ReLU (order='cge') and put non-linearity after the groupnorm.
    """

    def __init__(self, in_channels, out_channels, kernel_size=3, order=
        'cge', num_groups=8, **kwargs):
        super(ExtResNetBlockNew, self).__init__()
        self.conv1 = SingleConv(in_channels, out_channels, kernel_size=
            kernel_size, order=order, num_groups=num_groups)
        self.conv2 = SingleConv(out_channels, out_channels, kernel_size=
            kernel_size, order=order, num_groups=num_groups)
        n_order = order
        for c in 'rel':
            n_order = n_order.replace(c, '')
        self.conv3 = SingleConv(out_channels, out_channels, kernel_size=
            kernel_size, order=n_order, num_groups=num_groups)
        if 'l' in order:
            self.non_linearity = nn.LeakyReLU(negative_slope=0.1, inplace=True)
        elif 'e' in order:
            self.non_linearity = nn.ELU(inplace=True)
        else:
            self.non_linearity = nn.ReLU(inplace=True)

    def forward(self, input_0):
        primals_1 = self.conv1.conv.weight
        primals_3 = self.conv1.groupnorm.weight
        primals_4 = self.conv1.groupnorm.bias
        primals_5 = self.conv2.conv.weight
        primals_6 = self.conv2.groupnorm.weight
        primals_7 = self.conv2.groupnorm.bias
        primals_8 = self.conv3.conv.weight
        primals_9 = self.conv3.groupnorm.weight
        primals_10 = self.conv3.groupnorm.bias
        primals_2 = input_0
        output = call([primals_1, primals_2, primals_3, primals_4,
            primals_5, primals_6, primals_7, primals_8, primals_9, primals_10])
        return output[0]


Measured latency_ms: 0.115499