import torch
from torch import device
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, math as tl_math
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_red_fused_convolution_native_group_norm_0(in_out_ptr0, in_ptr0,
    in_ptr1, in_ptr2, out_ptr0, out_ptr2, out_ptr3, xnumel, rnumel, XBLOCK:
    tl.constexpr, RBLOCK: tl.constexpr):
    xnumel = 4
    rnumel = 4096
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    rbase = tl.arange(0, RBLOCK)[None, :]
    x0 = xindex
    tmp6_mean = tl.zeros([XBLOCK, RBLOCK], tl.float32)
    tmp6_m2 = tl.zeros([XBLOCK, RBLOCK], tl.float32)
    tmp6_weight = tl.zeros([XBLOCK, RBLOCK], tl.float32)
    for roffset in range(0, rnumel, RBLOCK):
        rindex = roffset + rbase
        rmask = rindex < rnumel
        r3 = rindex
        r2 = rindex // 1024
        tmp0 = tl.load(in_out_ptr0 + (r3 + 4096 * x0), rmask & xmask,
            eviction_policy='evict_first', other=0.0)
        tmp1 = tl.load(in_ptr0 + r2, rmask, eviction_policy='evict_last',
            other=0.0)
        tmp2 = tmp0 + tmp1
        tmp3 = tl.full([1, 1], 0, tl.int32)
        tmp4 = triton_helpers.maximum(tmp3, tmp2)
        tmp5 = tl.broadcast_to(tmp4, [XBLOCK, RBLOCK])
        tmp6_mean_next, tmp6_m2_next, tmp6_weight_next = (triton_helpers.
            welford_reduce(tmp5, tmp6_mean, tmp6_m2, tmp6_weight, roffset == 0)
            )
        tmp6_mean = tl.where(rmask & xmask, tmp6_mean_next, tmp6_mean)
        tmp6_m2 = tl.where(rmask & xmask, tmp6_m2_next, tmp6_m2)
        tmp6_weight = tl.where(rmask & xmask, tmp6_weight_next, tmp6_weight)
        tl.store(in_out_ptr0 + (r3 + 4096 * x0), tmp2, rmask & xmask)
    tmp6_tmp, tmp7_tmp, tmp8_tmp = triton_helpers.welford(tmp6_mean,
        tmp6_m2, tmp6_weight, 1)
    tmp6 = tmp6_tmp[:, None]
    tmp7 = tmp7_tmp[:, None]
    tmp8_tmp[:, None]
    tl.store(out_ptr0 + x0, tmp6, xmask)
    for roffset in range(0, rnumel, RBLOCK):
        rindex = roffset + rbase
        rmask = rindex < rnumel
        r3 = rindex
        r2 = rindex // 1024
        tmp9 = tl.load(in_out_ptr0 + (r3 + 4096 * x0), rmask & xmask,
            eviction_policy='evict_first', other=0.0)
        tmp19 = tl.load(in_ptr1 + r2, rmask, eviction_policy='evict_last',
            other=0.0)
        tmp21 = tl.load(in_ptr2 + r2, rmask, eviction_policy='evict_last',
            other=0.0)
        tmp10 = tl.full([1, 1], 0, tl.int32)
        tmp11 = triton_helpers.maximum(tmp10, tmp9)
        tmp12 = tmp11 - tmp6
        tmp13 = 4096.0
        tmp14 = tmp7 / tmp13
        tmp15 = 1e-05
        tmp16 = tmp14 + tmp15
        tmp17 = libdevice.rsqrt(tmp16)
        tmp18 = tmp12 * tmp17
        tmp20 = tmp18 * tmp19
        tmp22 = tmp20 + tmp21
        tl.store(out_ptr2 + (r3 + 4096 * x0), tmp22, rmask & xmask)
    tmp23 = 4096.0
    tmp24 = tmp7 / tmp23
    tmp25 = 1e-05
    tmp26 = tmp24 + tmp25
    tmp27 = libdevice.rsqrt(tmp26)
    tl.store(out_ptr3 + x0, tmp27, xmask)


@triton.jit
def triton_red_fused_convolution_native_group_norm_1(in_out_ptr0, in_ptr0,
    in_ptr1, in_ptr2, out_ptr0, out_ptr2, out_ptr3, xnumel, rnumel, XBLOCK:
    tl.constexpr, RBLOCK: tl.constexpr):
    xnumel = 4
    rnumel = 4096
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    rbase = tl.arange(0, RBLOCK)[None, :]
    x0 = xindex
    tmp6_mean = tl.zeros([XBLOCK, RBLOCK], tl.float32)
    tmp6_m2 = tl.zeros([XBLOCK, RBLOCK], tl.float32)
    tmp6_weight = tl.zeros([XBLOCK, RBLOCK], tl.float32)
    for roffset in range(0, rnumel, RBLOCK):
        rindex = roffset + rbase
        rmask = rindex < rnumel
        r3 = rindex
        r2 = rindex // 256
        tmp0 = tl.load(in_out_ptr0 + (r3 + 4096 * x0), rmask & xmask,
            eviction_policy='evict_first', other=0.0)
        tmp1 = tl.load(in_ptr0 + r2, rmask, eviction_policy='evict_last',
            other=0.0)
        tmp2 = tmp0 + tmp1
        tmp3 = tl.full([1, 1], 0, tl.int32)
        tmp4 = triton_helpers.maximum(tmp3, tmp2)
        tmp5 = tl.broadcast_to(tmp4, [XBLOCK, RBLOCK])
        tmp6_mean_next, tmp6_m2_next, tmp6_weight_next = (triton_helpers.
            welford_reduce(tmp5, tmp6_mean, tmp6_m2, tmp6_weight, roffset == 0)
            )
        tmp6_mean = tl.where(rmask & xmask, tmp6_mean_next, tmp6_mean)
        tmp6_m2 = tl.where(rmask & xmask, tmp6_m2_next, tmp6_m2)
        tmp6_weight = tl.where(rmask & xmask, tmp6_weight_next, tmp6_weight)
        tl.store(in_out_ptr0 + (r3 + 4096 * x0), tmp2, rmask & xmask)
    tmp6_tmp, tmp7_tmp, tmp8_tmp = triton_helpers.welford(tmp6_mean,
        tmp6_m2, tmp6_weight, 1)
    tmp6 = tmp6_tmp[:, None]
    tmp7 = tmp7_tmp[:, None]
    tmp8_tmp[:, None]
    tl.store(out_ptr0 + x0, tmp6, xmask)
    for roffset in range(0, rnumel, RBLOCK):
        rindex = roffset + rbase
        rmask = rindex < rnumel
        r3 = rindex
        r2 = rindex // 256
        tmp9 = tl.load(in_out_ptr0 + (r3 + 4096 * x0), rmask & xmask,
            eviction_policy='evict_first', other=0.0)
        tmp19 = tl.load(in_ptr1 + r2, rmask, eviction_policy='evict_last',
            other=0.0)
        tmp21 = tl.load(in_ptr2 + r2, rmask, eviction_policy='evict_last',
            other=0.0)
        tmp10 = tl.full([1, 1], 0, tl.int32)
        tmp11 = triton_helpers.maximum(tmp10, tmp9)
        tmp12 = tmp11 - tmp6
        tmp13 = 4096.0
        tmp14 = tmp7 / tmp13
        tmp15 = 1e-05
        tmp16 = tmp14 + tmp15
        tmp17 = libdevice.rsqrt(tmp16)
        tmp18 = tmp12 * tmp17
        tmp20 = tmp18 * tmp19
        tmp22 = tmp20 + tmp21
        tl.store(out_ptr2 + (r3 + 4096 * x0), tmp22, rmask & xmask)
    tmp23 = 4096.0
    tmp24 = tmp7 / tmp23
    tmp25 = 1e-05
    tmp26 = tmp24 + tmp25
    tmp27 = libdevice.rsqrt(tmp26)
    tl.store(out_ptr3 + x0, tmp27, xmask)


@triton.jit
def triton_red_fused_convolution_native_group_norm_2(in_out_ptr0, in_ptr0,
    in_ptr1, in_ptr2, out_ptr0, out_ptr2, out_ptr3, xnumel, rnumel, XBLOCK:
    tl.constexpr, RBLOCK: tl.constexpr):
    xnumel = 4
    rnumel = 2048
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    rbase = tl.arange(0, RBLOCK)[None, :]
    x0 = xindex
    tmp6_mean = tl.zeros([XBLOCK, RBLOCK], tl.float32)
    tmp6_m2 = tl.zeros([XBLOCK, RBLOCK], tl.float32)
    tmp6_weight = tl.zeros([XBLOCK, RBLOCK], tl.float32)
    for roffset in range(0, rnumel, RBLOCK):
        rindex = roffset + rbase
        rmask = rindex < rnumel
        r3 = rindex
        r2 = rindex // 64
        tmp0 = tl.load(in_out_ptr0 + (r3 + 2048 * x0), rmask & xmask,
            eviction_policy='evict_first', other=0.0)
        tmp1 = tl.load(in_ptr0 + r2, rmask, eviction_policy='evict_last',
            other=0.0)
        tmp2 = tmp0 + tmp1
        tmp3 = tl.full([1, 1], 0, tl.int32)
        tmp4 = triton_helpers.maximum(tmp3, tmp2)
        tmp5 = tl.broadcast_to(tmp4, [XBLOCK, RBLOCK])
        tmp6_mean_next, tmp6_m2_next, tmp6_weight_next = (triton_helpers.
            welford_reduce(tmp5, tmp6_mean, tmp6_m2, tmp6_weight, roffset == 0)
            )
        tmp6_mean = tl.where(rmask & xmask, tmp6_mean_next, tmp6_mean)
        tmp6_m2 = tl.where(rmask & xmask, tmp6_m2_next, tmp6_m2)
        tmp6_weight = tl.where(rmask & xmask, tmp6_weight_next, tmp6_weight)
        tl.store(in_out_ptr0 + (r3 + 2048 * x0), tmp2, rmask & xmask)
    tmp6_tmp, tmp7_tmp, tmp8_tmp = triton_helpers.welford(tmp6_mean,
        tmp6_m2, tmp6_weight, 1)
    tmp6 = tmp6_tmp[:, None]
    tmp7 = tmp7_tmp[:, None]
    tmp8_tmp[:, None]
    tl.store(out_ptr0 + x0, tmp6, xmask)
    for roffset in range(0, rnumel, RBLOCK):
        rindex = roffset + rbase
        rmask = rindex < rnumel
        r3 = rindex
        r2 = rindex // 64
        tmp9 = tl.load(in_out_ptr0 + (r3 + 2048 * x0), rmask & xmask,
            eviction_policy='evict_first', other=0.0)
        tmp19 = tl.load(in_ptr1 + r2, rmask, eviction_policy='evict_last',
            other=0.0)
        tmp21 = tl.load(in_ptr2 + r2, rmask, eviction_policy='evict_last',
            other=0.0)
        tmp10 = tl.full([1, 1], 0, tl.int32)
        tmp11 = triton_helpers.maximum(tmp10, tmp9)
        tmp12 = tmp11 - tmp6
        tmp13 = 2048.0
        tmp14 = tmp7 / tmp13
        tmp15 = 1e-05
        tmp16 = tmp14 + tmp15
        tmp17 = libdevice.rsqrt(tmp16)
        tmp18 = tmp12 * tmp17
        tmp20 = tmp18 * tmp19
        tmp22 = tmp20 + tmp21
        tl.store(out_ptr2 + (r3 + 2048 * x0), tmp22, rmask & xmask)
    tmp23 = 2048.0
    tmp24 = tmp7 / tmp23
    tmp25 = 1e-05
    tmp26 = tmp24 + tmp25
    tmp27 = libdevice.rsqrt(tmp26)
    tl.store(out_ptr3 + x0, tmp27, xmask)


@triton.jit
def triton_poi_fused_add_convolution_exp_mul_3(in_out_ptr0, in_out_ptr1,
    in_ptr0, in_ptr1, in_ptr2, out_ptr0, xnumel, XBLOCK: tl.constexpr):
    xnumel = 1024
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x3 = xindex
    x1 = xindex // 16 % 16
    tmp0 = tl.load(in_out_ptr0 + x3, xmask)
    tmp1 = tl.load(in_ptr0 + x1, xmask, eviction_policy='evict_last')
    tmp3 = tl.load(in_out_ptr1 + x3, xmask)
    tmp4 = tl.load(in_ptr1 + x1, xmask, eviction_policy='evict_last')
    tmp6 = tl.load(in_ptr2 + x3, xmask)
    tmp2 = tmp0 + tmp1
    tmp5 = tmp3 + tmp4
    tmp7 = 0.5
    tmp8 = tmp5 * tmp7
    tmp9 = tl_math.exp(tmp8)
    tmp10 = tmp6 * tmp9
    tmp11 = tmp2 + tmp10
    tl.store(in_out_ptr0 + x3, tmp2, xmask)
    tl.store(in_out_ptr1 + x3, tmp5, xmask)
    tl.store(out_ptr0 + x3, tmp11, xmask)


@triton.jit
def triton_red_fused_convolution_native_group_norm_4(in_out_ptr0, in_ptr0,
    out_ptr0, out_ptr1, out_ptr2, xnumel, rnumel, XBLOCK: tl.constexpr,
    RBLOCK: tl.constexpr):
    xnumel = 8
    rnumel = 8192
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    rbase = tl.arange(0, RBLOCK)[None, :]
    x4 = xindex
    x0 = xindex % 2
    tmp6_mean = tl.zeros([XBLOCK, RBLOCK], tl.float32)
    tmp6_m2 = tl.zeros([XBLOCK, RBLOCK], tl.float32)
    tmp6_weight = tl.zeros([XBLOCK, RBLOCK], tl.float32)
    for roffset in range(0, rnumel, RBLOCK):
        rindex = roffset + rbase
        rmask = rindex < rnumel
        r5 = rindex
        r3 = rindex // 256
        tmp0 = tl.load(in_out_ptr0 + (r5 + 8192 * x4), rmask & xmask,
            eviction_policy='evict_first', other=0.0)
        tmp1 = tl.load(in_ptr0 + (r3 + 32 * x0), rmask & xmask,
            eviction_policy='evict_last', other=0.0)
        tmp2 = tmp0 + tmp1
        tmp3 = tl.full([1, 1], 0, tl.int32)
        tmp4 = triton_helpers.maximum(tmp3, tmp2)
        tmp5 = tl.broadcast_to(tmp4, [XBLOCK, RBLOCK])
        tmp6_mean_next, tmp6_m2_next, tmp6_weight_next = (triton_helpers.
            welford_reduce(tmp5, tmp6_mean, tmp6_m2, tmp6_weight, roffset == 0)
            )
        tmp6_mean = tl.where(rmask & xmask, tmp6_mean_next, tmp6_mean)
        tmp6_m2 = tl.where(rmask & xmask, tmp6_m2_next, tmp6_m2)
        tmp6_weight = tl.where(rmask & xmask, tmp6_weight_next, tmp6_weight)
        tl.store(in_out_ptr0 + (r5 + 8192 * x4), tmp2, rmask & xmask)
    tmp6_tmp, tmp7_tmp, tmp8_tmp = triton_helpers.welford(tmp6_mean,
        tmp6_m2, tmp6_weight, 1)
    tmp6 = tmp6_tmp[:, None]
    tmp7 = tmp7_tmp[:, None]
    tmp8 = tmp8_tmp[:, None]
    tl.store(out_ptr0 + x4, tmp6, xmask)
    tl.store(out_ptr1 + x4, tmp7, xmask)
    tl.store(out_ptr2 + x4, tmp8, xmask)


@triton.jit
def triton_per_fused_native_group_norm_5(in_ptr0, in_ptr1, in_ptr2,
    out_ptr0, out_ptr1, out_ptr2, xnumel, rnumel, XBLOCK: tl.constexpr):
    xnumel = 4
    RBLOCK: tl.constexpr = 2
    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 + 2 * x0), xmask, other=0.0)
    tmp1 = tl.load(in_ptr1 + (r1 + 2 * x0), xmask, other=0.0)
    tmp2 = tl.load(in_ptr2 + (r1 + 2 * x0), xmask, other=0.0)
    tmp3 = tl.broadcast_to(tmp0, [XBLOCK, RBLOCK])
    tmp4 = tl.broadcast_to(tmp1, [XBLOCK, RBLOCK])
    tmp5 = tl.broadcast_to(tmp2, [XBLOCK, RBLOCK])
    tmp7 = tl.where(xmask, tmp3, 0)
    tmp8 = tl.where(xmask, tmp4, 0)
    tmp9 = tl.where(xmask, tmp5, 0)
    tmp10, tmp11, tmp12 = triton_helpers.welford(tmp7, tmp8, tmp9, 1)
    tmp13 = tmp10[:, None]
    tmp14 = tmp11[:, None]
    tmp12[:, None]
    tmp16 = 16384.0
    tmp17 = tmp14 / tmp16
    tmp18 = 1e-05
    tmp19 = tmp17 + tmp18
    tmp20 = libdevice.rsqrt(tmp19)
    tl.store(out_ptr2 + x0, tmp20, xmask)
    tl.store(out_ptr0 + x0, tmp13, xmask)
    tl.store(out_ptr1 + x0, tmp14, xmask)


@triton.jit
def triton_poi_fused_native_group_norm_6(in_ptr0, in_ptr1, in_ptr2, in_ptr3,
    in_ptr4, out_ptr0, xnumel, XBLOCK: tl.constexpr):
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    tl.full([XBLOCK], True, tl.int1)
    x3 = xindex
    x2 = xindex // 16384
    x1 = xindex // 256 % 64
    tmp0 = tl.load(in_ptr0 + x3, None)
    tmp3 = tl.load(in_ptr1 + x2, None, eviction_policy='evict_last')
    tmp5 = tl.load(in_ptr2 + x2, None, eviction_policy='evict_last')
    tmp12 = tl.load(in_ptr3 + x1, None, eviction_policy='evict_last')
    tmp14 = tl.load(in_ptr4 + x1, None, eviction_policy='evict_last')
    tmp1 = tl.full([1], 0, tl.int32)
    tmp2 = triton_helpers.maximum(tmp1, tmp0)
    tmp4 = tmp2 - tmp3
    tmp6 = 16384.0
    tmp7 = tmp5 / tmp6
    tmp8 = 1e-05
    tmp9 = tmp7 + tmp8
    tmp10 = libdevice.rsqrt(tmp9)
    tmp11 = tmp4 * tmp10
    tmp13 = tmp11 * tmp12
    tmp15 = tmp13 + tmp14
    tl.store(out_ptr0 + x3, tmp15, None)


@triton.jit
def triton_red_fused_convolution_native_group_norm_7(in_out_ptr0, in_ptr0,
    out_ptr0, out_ptr1, out_ptr2, xnumel, rnumel, XBLOCK: tl.constexpr,
    RBLOCK: tl.constexpr):
    xnumel = 16
    rnumel = 8192
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    rbase = tl.arange(0, RBLOCK)[None, :]
    x4 = xindex
    x0 = xindex % 4
    tmp6_mean = tl.zeros([XBLOCK, RBLOCK], tl.float32)
    tmp6_m2 = tl.zeros([XBLOCK, RBLOCK], tl.float32)
    tmp6_weight = tl.zeros([XBLOCK, RBLOCK], tl.float32)
    for roffset in range(0, rnumel, RBLOCK):
        rindex = roffset + rbase
        rmask = rindex < rnumel
        r5 = rindex
        r3 = rindex // 1024
        tmp0 = tl.load(in_out_ptr0 + (r5 + 8192 * x4), rmask & xmask,
            eviction_policy='evict_first', other=0.0)
        tmp1 = tl.load(in_ptr0 + (r3 + 8 * x0), rmask & xmask,
            eviction_policy='evict_last', other=0.0)
        tmp2 = tmp0 + tmp1
        tmp3 = tl.full([1, 1], 0, tl.int32)
        tmp4 = triton_helpers.maximum(tmp3, tmp2)
        tmp5 = tl.broadcast_to(tmp4, [XBLOCK, RBLOCK])
        tmp6_mean_next, tmp6_m2_next, tmp6_weight_next = (triton_helpers.
            welford_reduce(tmp5, tmp6_mean, tmp6_m2, tmp6_weight, roffset == 0)
            )
        tmp6_mean = tl.where(rmask & xmask, tmp6_mean_next, tmp6_mean)
        tmp6_m2 = tl.where(rmask & xmask, tmp6_m2_next, tmp6_m2)
        tmp6_weight = tl.where(rmask & xmask, tmp6_weight_next, tmp6_weight)
        tl.store(in_out_ptr0 + (r5 + 8192 * x4), tmp2, rmask & xmask)
    tmp6_tmp, tmp7_tmp, tmp8_tmp = triton_helpers.welford(tmp6_mean,
        tmp6_m2, tmp6_weight, 1)
    tmp6 = tmp6_tmp[:, None]
    tmp7 = tmp7_tmp[:, None]
    tmp8 = tmp8_tmp[:, None]
    tl.store(out_ptr0 + x4, tmp6, xmask)
    tl.store(out_ptr1 + x4, tmp7, xmask)
    tl.store(out_ptr2 + x4, tmp8, xmask)


@triton.jit
def triton_per_fused_native_group_norm_8(in_ptr0, in_ptr1, in_ptr2,
    out_ptr0, out_ptr1, out_ptr2, xnumel, rnumel, XBLOCK: tl.constexpr):
    xnumel = 4
    RBLOCK: tl.constexpr = 4
    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 + 4 * x0), xmask, other=0.0)
    tmp1 = tl.load(in_ptr1 + (r1 + 4 * x0), xmask, other=0.0)
    tmp2 = tl.load(in_ptr2 + (r1 + 4 * x0), xmask, other=0.0)
    tmp3 = tl.broadcast_to(tmp0, [XBLOCK, RBLOCK])
    tmp4 = tl.broadcast_to(tmp1, [XBLOCK, RBLOCK])
    tmp5 = tl.broadcast_to(tmp2, [XBLOCK, RBLOCK])
    tmp7 = tl.where(xmask, tmp3, 0)
    tmp8 = tl.where(xmask, tmp4, 0)
    tmp9 = tl.where(xmask, tmp5, 0)
    tmp10, tmp11, tmp12 = triton_helpers.welford(tmp7, tmp8, tmp9, 1)
    tmp13 = tmp10[:, None]
    tmp14 = tmp11[:, None]
    tmp12[:, None]
    tmp16 = 32768.0
    tmp17 = tmp14 / tmp16
    tmp18 = 1e-05
    tmp19 = tmp17 + tmp18
    tmp20 = libdevice.rsqrt(tmp19)
    tl.store(out_ptr2 + x0, tmp20, xmask)
    tl.store(out_ptr0 + x0, tmp13, xmask)
    tl.store(out_ptr1 + x0, tmp14, xmask)


@triton.jit
def triton_poi_fused_native_group_norm_9(in_ptr0, in_ptr1, in_ptr2, in_ptr3,
    in_ptr4, out_ptr0, xnumel, XBLOCK: tl.constexpr):
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    tl.full([XBLOCK], True, tl.int1)
    x3 = xindex
    x2 = xindex // 32768
    x1 = xindex // 1024 % 32
    tmp0 = tl.load(in_ptr0 + x3, None)
    tmp3 = tl.load(in_ptr1 + x2, None, eviction_policy='evict_last')
    tmp5 = tl.load(in_ptr2 + x2, None, eviction_policy='evict_last')
    tmp12 = tl.load(in_ptr3 + x1, None, eviction_policy='evict_last')
    tmp14 = tl.load(in_ptr4 + x1, None, eviction_policy='evict_last')
    tmp1 = tl.full([1], 0, tl.int32)
    tmp2 = triton_helpers.maximum(tmp1, tmp0)
    tmp4 = tmp2 - tmp3
    tmp6 = 32768.0
    tmp7 = tmp5 / tmp6
    tmp8 = 1e-05
    tmp9 = tmp7 + tmp8
    tmp10 = libdevice.rsqrt(tmp9)
    tmp11 = tmp4 * tmp10
    tmp13 = tmp11 * tmp12
    tmp15 = tmp13 + tmp14
    tl.store(out_ptr0 + x3, tmp15, None)


@triton.jit
def triton_poi_fused_convolution_sigmoid_10(in_out_ptr0, in_ptr0, xnumel,
    XBLOCK: tl.constexpr):
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    tl.full([XBLOCK], True, tl.int1)
    x0 = xindex
    tmp0 = tl.load(in_out_ptr0 + x0, None)
    tmp1 = tl.load(in_ptr0 + 0)
    tmp2 = tl.broadcast_to(tmp1, [XBLOCK])
    tmp3 = tmp0 + tmp2
    tmp4 = tl.sigmoid(tmp3)
    tl.store(in_out_ptr0 + x0, tmp4, None)


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, primals_29, primals_30, primals_31) = args
    args.clear()
    assert_size_stride(primals_1, (4, 1, 4, 4), (16, 16, 4, 1))
    assert_size_stride(primals_2, (4,), (1,))
    assert_size_stride(primals_3, (4, 1, 64, 64), (4096, 4096, 64, 1))
    assert_size_stride(primals_4, (4,), (1,))
    assert_size_stride(primals_5, (4,), (1,))
    assert_size_stride(primals_6, (16, 4, 4, 4), (64, 16, 4, 1))
    assert_size_stride(primals_7, (16,), (1,))
    assert_size_stride(primals_8, (16,), (1,))
    assert_size_stride(primals_9, (16,), (1,))
    assert_size_stride(primals_10, (32, 16, 4, 4), (256, 16, 4, 1))
    assert_size_stride(primals_11, (32,), (1,))
    assert_size_stride(primals_12, (32,), (1,))
    assert_size_stride(primals_13, (32,), (1,))
    assert_size_stride(primals_14, (16, 32, 4, 4), (512, 16, 4, 1))
    assert_size_stride(primals_15, (16,), (1,))
    assert_size_stride(primals_16, (16, 32, 4, 4), (512, 16, 4, 1))
    assert_size_stride(primals_17, (16,), (1,))
    assert_size_stride(primals_18, (16, 32, 8, 8), (2048, 64, 8, 1))
    assert_size_stride(primals_19, (32,), (1,))
    assert_size_stride(primals_20, (32,), (1,))
    assert_size_stride(primals_21, (32,), (1,))
    assert_size_stride(primals_22, (32, 64, 8, 8), (4096, 64, 8, 1))
    assert_size_stride(primals_23, (64,), (1,))
    assert_size_stride(primals_24, (64,), (1,))
    assert_size_stride(primals_25, (64,), (1,))
    assert_size_stride(primals_26, (64, 32, 8, 8), (2048, 64, 8, 1))
    assert_size_stride(primals_27, (32,), (1,))
    assert_size_stride(primals_28, (32,), (1,))
    assert_size_stride(primals_29, (32,), (1,))
    assert_size_stride(primals_30, (32, 1, 8, 8), (64, 64, 8, 1))
    assert_size_stride(primals_31, (1,), (1,))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = extern_kernels.convolution(primals_3, primals_1, stride=(2, 
            2), padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf0, (4, 4, 32, 32), (4096, 1024, 32, 1))
        buf1 = buf0
        del buf0
        buf2 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf5 = empty_strided_cuda((4, 4, 32, 32), (4096, 1024, 32, 1),
            torch.float32)
        buf6 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        get_raw_stream(0)
        triton_red_fused_convolution_native_group_norm_0[grid(4)](buf1,
            primals_2, primals_4, primals_5, buf2, buf5, buf6, 4, 4096,
            XBLOCK=1, RBLOCK=2048, num_warps=16, num_stages=1)
        del primals_2
        del primals_5
        buf7 = extern_kernels.convolution(buf5, primals_6, stride=(2, 2),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf7, (4, 16, 16, 16), (4096, 256, 16, 1))
        buf8 = buf7
        del buf7
        buf9 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf12 = empty_strided_cuda((4, 16, 16, 16), (4096, 256, 16, 1),
            torch.float32)
        buf13 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        triton_red_fused_convolution_native_group_norm_1[grid(4)](buf8,
            primals_7, primals_8, primals_9, buf9, buf12, buf13, 4, 4096,
            XBLOCK=1, RBLOCK=2048, num_warps=16, num_stages=1)
        del primals_7
        del primals_9
        buf14 = extern_kernels.convolution(buf12, primals_10, stride=(2, 2),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf14, (4, 32, 8, 8), (2048, 64, 8, 1))
        buf15 = buf14
        del buf14
        buf16 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf19 = empty_strided_cuda((4, 32, 8, 8), (2048, 64, 8, 1), torch.
            float32)
        buf20 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        triton_red_fused_convolution_native_group_norm_2[grid(4)](buf15,
            primals_11, primals_12, primals_13, buf16, buf19, buf20, 4, 
            2048, XBLOCK=1, RBLOCK=2048, num_warps=16, num_stages=1)
        del primals_11
        del primals_13
        buf21 = extern_kernels.convolution(buf19, primals_14, stride=(2, 2),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf21, (4, 16, 4, 4), (256, 16, 4, 1))
        buf23 = extern_kernels.convolution(buf19, primals_16, stride=(2, 2),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf23, (4, 16, 4, 4), (256, 16, 4, 1))
        buf25 = torch.ops.aten.randn.default([4, 16, 4, 4], dtype=torch.
            float32, device=device(type='cuda', index=0), pin_memory=False)
        buf26 = buf25
        del buf25
        buf22 = buf21
        del buf21
        buf24 = buf23
        del buf23
        buf27 = empty_strided_cuda((4, 16, 4, 4), (256, 16, 4, 1), torch.
            float32)
        triton_poi_fused_add_convolution_exp_mul_3[grid(1024)](buf22, buf24,
            primals_15, primals_17, buf26, buf27, 1024, XBLOCK=128,
            num_warps=4, num_stages=1)
        del primals_15
        del primals_17
        buf28 = extern_kernels.convolution(buf27, primals_18, stride=(2, 2),
            padding=(3, 3), dilation=(1, 1), transposed=True,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf28, (4, 32, 8, 8), (2048, 64, 8, 1))
        buf29 = buf28
        del buf28
        buf30 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf33 = empty_strided_cuda((4, 32, 8, 8), (2048, 64, 8, 1), torch.
            float32)
        buf34 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        triton_red_fused_convolution_native_group_norm_2[grid(4)](buf29,
            primals_19, primals_20, primals_21, buf30, buf33, buf34, 4, 
            2048, XBLOCK=1, RBLOCK=2048, num_warps=16, num_stages=1)
        del primals_19
        del primals_21
        buf35 = extern_kernels.convolution(buf33, primals_22, stride=(2, 2),
            padding=(3, 3), dilation=(1, 1), transposed=True,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf35, (4, 64, 16, 16), (16384, 256, 16, 1))
        buf36 = buf35
        del buf35
        buf37 = empty_strided_cuda((4, 1, 1, 1, 2), (2, 8, 8, 8, 1), torch.
            float32)
        buf38 = empty_strided_cuda((4, 1, 1, 1, 2), (2, 8, 8, 8, 1), torch.
            float32)
        buf39 = empty_strided_cuda((4, 1, 1, 1, 2), (2, 8, 8, 8, 1), torch.
            float32)
        triton_red_fused_convolution_native_group_norm_4[grid(8)](buf36,
            primals_23, buf37, buf38, buf39, 8, 8192, XBLOCK=1, RBLOCK=2048,
            num_warps=16, num_stages=1)
        del primals_23
        buf40 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf41 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf44 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        triton_per_fused_native_group_norm_5[grid(4)](buf37, buf38, buf39,
            buf40, buf41, buf44, 4, 2, XBLOCK=1, num_warps=2, num_stages=1)
        del buf37
        del buf38
        del buf39
        buf43 = empty_strided_cuda((4, 64, 16, 16), (16384, 256, 16, 1),
            torch.float32)
        triton_poi_fused_native_group_norm_6[grid(65536)](buf36, buf40,
            buf41, primals_24, primals_25, buf43, 65536, XBLOCK=512,
            num_warps=4, num_stages=1)
        del primals_25
        buf45 = extern_kernels.convolution(buf43, primals_26, stride=(2, 2),
            padding=(3, 3), dilation=(1, 1), transposed=True,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf45, (4, 32, 32, 32), (32768, 1024, 32, 1))
        buf46 = buf45
        del buf45
        buf47 = empty_strided_cuda((4, 1, 1, 1, 4), (4, 16, 16, 16, 1),
            torch.float32)
        buf48 = empty_strided_cuda((4, 1, 1, 1, 4), (4, 16, 16, 16, 1),
            torch.float32)
        buf49 = empty_strided_cuda((4, 1, 1, 1, 4), (4, 16, 16, 16, 1),
            torch.float32)
        triton_red_fused_convolution_native_group_norm_7[grid(16)](buf46,
            primals_27, buf47, buf48, buf49, 16, 8192, XBLOCK=1, RBLOCK=
            2048, num_warps=16, num_stages=1)
        del primals_27
        buf50 = buf41
        del buf41
        buf51 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf54 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        triton_per_fused_native_group_norm_8[grid(4)](buf47, buf48, buf49,
            buf50, buf51, buf54, 4, 4, XBLOCK=1, num_warps=2, num_stages=1)
        del buf47
        del buf48
        del buf49
        buf53 = empty_strided_cuda((4, 32, 32, 32), (32768, 1024, 32, 1),
            torch.float32)
        triton_poi_fused_native_group_norm_9[grid(131072)](buf46, buf50,
            buf51, primals_28, primals_29, buf53, 131072, XBLOCK=512,
            num_warps=8, num_stages=1)
        del buf51
        del primals_29
        buf55 = extern_kernels.convolution(buf53, primals_30, stride=(2, 2),
            padding=(3, 3), dilation=(1, 1), transposed=True,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf55, (4, 1, 64, 64), (4096, 4096, 64, 1))
        buf56 = buf55
        del buf55
        triton_poi_fused_convolution_sigmoid_10[grid(16384)](buf56,
            primals_31, 16384, XBLOCK=256, num_warps=4, num_stages=1)
        del primals_31
    return (buf56, buf22, buf24, buf27, buf12, buf43, primals_1, primals_3,
        primals_4, primals_6, primals_8, primals_10, primals_12, primals_14,
        primals_16, primals_18, primals_20, primals_22, primals_24,
        primals_26, primals_28, primals_30, buf1, buf5, reinterpret_tensor(
        buf2, (4, 1), (1, 1), 0), reinterpret_tensor(buf6, (4, 1), (1, 1), 
        0), buf8, buf12, reinterpret_tensor(buf9, (4, 1), (1, 1), 0),
        reinterpret_tensor(buf13, (4, 1), (1, 1), 0), buf15, buf19,
        reinterpret_tensor(buf16, (4, 1), (1, 1), 0), reinterpret_tensor(
        buf20, (4, 1), (1, 1), 0), buf24, buf26, buf27, buf29, buf33,
        reinterpret_tensor(buf30, (4, 1), (1, 1), 0), reinterpret_tensor(
        buf34, (4, 1), (1, 1), 0), buf36, buf43, reinterpret_tensor(buf40,
        (4, 1), (1, 1), 0), reinterpret_tensor(buf44, (4, 1), (1, 1), 0),
        buf46, buf53, reinterpret_tensor(buf50, (4, 1), (1, 1), 0),
        reinterpret_tensor(buf54, (4, 1), (1, 1), 0), buf56)


class TextureFinderNew(nn.Module):

    def __init__(self):
        super(TextureFinderNew, self).__init__()
        self.encoder_conv1 = nn.Conv2d(in_channels=1, out_channels=4,
            kernel_size=4, stride=2, padding=1, dilation=1, groups=1, bias=True
            )
        self.encoder_conv1.bias.data.zero_()
        self.encoder_conv1.weight.data[:, :, :, :] = 1 / 0.32 + torch.normal(
            mean=torch.tensor(0.0), std=torch.tensor(0.0001))
        self.encoder_normalization1 = nn.GroupNorm(1, 4, eps=1e-05, affine=True
            )
        self.encoder_conv2 = nn.Conv2d(in_channels=4, out_channels=16,
            kernel_size=4, stride=2, padding=1, dilation=1, groups=1, bias=True
            )
        self.encoder_conv2.bias.data.zero_()
        self.encoder_conv2.weight.data[:, :, :, :] = 1 / (8 * 16
            ) + torch.normal(mean=torch.tensor(0.0), std=torch.tensor(0.0001))
        self.encoder_normalization2 = nn.GroupNorm(1, 16, eps=1e-05, affine
            =True)
        self.encoder_conv3 = nn.Conv2d(in_channels=16, out_channels=32,
            kernel_size=4, stride=2, padding=1, dilation=1, groups=1, bias=True
            )
        self.encoder_conv3.bias.data.zero_()
        self.encoder_conv3.weight.data[:, :, :, :] = 1 / (8 * 16
            ) + torch.normal(mean=torch.tensor(0.0), std=torch.tensor(0.001))
        self.encoder_normalization3 = nn.GroupNorm(1, 32, eps=1e-05, affine
            =True)
        self.encoder_mu = nn.Conv2d(in_channels=32, out_channels=16,
            kernel_size=4, stride=2, padding=1, dilation=1, groups=1, bias=True
            )
        self.encoder_mu.bias.data.zero_()
        self.encoder_mu.weight.data[:, :, :, :] = 1 / (8 * 16) + torch.normal(
            mean=torch.tensor(0.0), std=torch.tensor(0.0001))
        self.encoder_log_var = nn.Conv2d(in_channels=32, out_channels=16,
            kernel_size=4, stride=2, padding=1, dilation=1, groups=1, bias=True
            )
        self.encoder_log_var.bias.data[:] = -2.3
        self.encoder_log_var.weight.data.zero_()
        self.decoder_conv1 = nn.ConvTranspose2d(16, 32, kernel_size=8,
            stride=2, padding=3)
        self.decoder_conv1.bias.data.zero_()
        self.decoder_conv1.weight.data[:, :, :, :] = 1 / (8 * 8 * 8
            ) + torch.normal(mean=torch.tensor(0.0), std=torch.tensor(0.0001))
        self.decoder_normalization1 = nn.GroupNorm(1, 32, eps=1e-05, affine
            =True)
        self.decoder_conv2 = nn.ConvTranspose2d(32, 64, kernel_size=8,
            stride=2, padding=3, output_padding=0, groups=1, bias=True,
            dilation=1)
        self.decoder_conv2.bias.data.zero_()
        self.decoder_conv2.weight.data[:, :, :, :] = 1 / (8 * 8 * 8
            ) + torch.normal(mean=torch.tensor(0.0), std=torch.tensor(0.0001))
        self.decoder_normalization2 = nn.GroupNorm(1, 64, eps=1e-05, affine
            =True)
        self.decoder_conv3 = nn.ConvTranspose2d(64, 32, kernel_size=8,
            stride=2, padding=3, output_padding=0, groups=1, bias=True,
            dilation=1)
        self.decoder_conv3.bias.data.zero_()
        self.decoder_conv3.weight.data[:, :, :, :] = 1 / (4 * 8 * 8
            ) + torch.normal(mean=torch.tensor(0.0), std=torch.tensor(0.001))
        self.decoder_normalization3 = nn.GroupNorm(1, 32, eps=1e-05, affine
            =True)
        self.decoder_conv5 = nn.ConvTranspose2d(32, 1, kernel_size=8,
            stride=2, padding=3, output_padding=0, groups=1, bias=True,
            dilation=1)
        self.decoder_conv5.bias.data[:] = -(0.5 / 0.24)
        self.decoder_conv5.weight.data[:, :, :, :] = 1 / (32 * 8 * 8 * 0.24
            ) + torch.normal(mean=torch.tensor(0.0), std=torch.tensor(0.001))

    def sample_from_mu_log_var(self, mu, log_var):
        std = torch.exp(0.5 * log_var)
        eps = torch.randn_like(std)
        sample = mu + eps * std
        return sample

    def forward(self, input_0):
        primals_1 = self.encoder_conv1.weight
        primals_2 = self.encoder_conv1.bias
        primals_4 = self.encoder_normalization1.weight
        primals_5 = self.encoder_normalization1.bias
        primals_6 = self.encoder_conv2.weight
        primals_7 = self.encoder_conv2.bias
        primals_8 = self.encoder_normalization2.weight
        primals_9 = self.encoder_normalization2.bias
        primals_10 = self.encoder_conv3.weight
        primals_11 = self.encoder_conv3.bias
        primals_12 = self.encoder_normalization3.weight
        primals_13 = self.encoder_normalization3.bias
        primals_14 = self.encoder_mu.weight
        primals_15 = self.encoder_mu.bias
        primals_16 = self.encoder_log_var.weight
        primals_17 = self.encoder_log_var.bias
        primals_18 = self.decoder_conv1.weight
        primals_19 = self.decoder_conv1.bias
        primals_20 = self.decoder_normalization1.weight
        primals_21 = self.decoder_normalization1.bias
        primals_22 = self.decoder_conv2.weight
        primals_23 = self.decoder_conv2.bias
        primals_24 = self.decoder_normalization2.weight
        primals_25 = self.decoder_normalization2.bias
        primals_26 = self.decoder_conv3.weight
        primals_27 = self.decoder_conv3.bias
        primals_28 = self.decoder_normalization3.weight
        primals_29 = self.decoder_normalization3.bias
        primals_30 = self.decoder_conv5.weight
        primals_31 = self.decoder_conv5.bias
        primals_3 = 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, primals_29,
            primals_30, primals_31])
        return output[0], output[1], output[2], output[3], output[4], output[5]


Measured latency_ms: 0.11818