# AOT ID: ['0_forward']
from ctypes import c_void_p, c_long, c_int
import torch
import math
import random
import os
import tempfile
from math import inf, nan
from torch._inductor.hooks import run_intermediate_hooks
from torch._inductor.utils import maybe_profile
from torch._inductor.codegen.memory_planning import _align as align
from torch import device, empty_strided
from torch._inductor.async_compile import AsyncCompile
from torch._inductor.select_algorithm import extern_kernels
from torch._inductor.codegen.multi_kernel import MultiKernelCall
import triton
import triton.language as tl
from torch._inductor.runtime.triton_heuristics import grid, split_scan_grid, grid_combo_kernels, start_graph, end_graph
from torch._C import _cuda_getCurrentRawStream as get_raw_stream

aten = torch.ops.aten
inductor_ops = torch.ops.inductor
_quantized = torch.ops._quantized
assert_size_stride = torch._C._dynamo.guards.assert_size_stride
empty_strided_cpu = torch._C._dynamo.guards._empty_strided_cpu
empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda
empty_strided_xpu = torch._C._dynamo.guards._empty_strided_xpu
reinterpret_tensor = torch._C._dynamo.guards._reinterpret_tensor
alloc_from_pool = torch.ops.inductor._alloc_from_pool
async_compile = AsyncCompile()


# kernel path: runs/run_shard_1/inductor_cache/xq/cxqxvwhuevcb5oe7gsgfej3rmce7mdc6vqg3mkviccgugis2c7ro.py
# Unsorted Source Nodes: [], Original ATen: []
# Source node to ATen node mapping:
triton_poi_fused_0 = async_compile.triton('triton_poi_fused_0', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[256, 16], tile_hint=TileHint.SQUARE,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: 'i32', 3: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_0', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 1, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_0(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    ynumel = 192
    xnumel = 9
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = yindex < ynumel
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y3 = yindex
    y0 = yindex % 3
    y1 = (yindex // 3)
    tmp0 = tl.load(in_ptr0 + (x2 + (9*y3)), xmask & ymask, eviction_policy='evict_last')
    tl.store(out_ptr0 + (y0 + (3*x2) + (27*y1)), tmp0, xmask & ymask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/eb/cebfmp3xsydvhof2vuiuzrtwr7fwapeufpm52glmntqds2lsygkv.py
# Unsorted Source Nodes: [], Original ATen: []
# Source node to ATen node mapping:
triton_poi_fused_1 = async_compile.triton('triton_poi_fused_1', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[4096, 16], tile_hint=TileHint.SQUARE,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: 'i32', 3: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_1', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 1, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_1(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    ynumel = 4096
    xnumel = 9
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y3 = yindex
    y0 = yindex % 64
    y1 = (yindex // 64)
    tmp0 = tl.load(in_ptr0 + (x2 + (9*y3)), xmask, eviction_policy='evict_last')
    tl.store(out_ptr0 + (y0 + (64*x2) + (576*y1)), tmp0, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/xs/cxsarnw2wm2gid2judloczqftyialh3etpmvbejw7tuglcm5m2ir.py
# Unsorted Source Nodes: [], Original ATen: []
# Source node to ATen node mapping:
triton_poi_fused_2 = async_compile.triton('triton_poi_fused_2', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[8192, 16], tile_hint=TileHint.SQUARE,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: 'i32', 3: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_2', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 1, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_2(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    ynumel = 8192
    xnumel = 9
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y3 = yindex
    y0 = yindex % 64
    y1 = (yindex // 64)
    tmp0 = tl.load(in_ptr0 + (x2 + (9*y3)), xmask, eviction_policy='evict_last')
    tl.store(out_ptr0 + (y0 + (64*x2) + (576*y1)), tmp0, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/fh/cfhbha4jjott434uisf5mi7hyc6aabnyin37w2nsewlcpyvjmf25.py
# Unsorted Source Nodes: [], Original ATen: []
# Source node to ATen node mapping:
triton_poi_fused_3 = async_compile.triton('triton_poi_fused_3', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[16384, 16], tile_hint=TileHint.SQUARE,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: 'i32', 3: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_3', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 1, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_3(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    ynumel = 16384
    xnumel = 9
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y3 = yindex
    y0 = yindex % 128
    y1 = (yindex // 128)
    tmp0 = tl.load(in_ptr0 + (x2 + (9*y3)), xmask, eviction_policy='evict_last')
    tl.store(out_ptr0 + (y0 + (128*x2) + (1152*y1)), tmp0, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/tb/ctbdw725wibroaudcyozlixnt22th2hpv7lozyazwday5aem5n4w.py
# Unsorted Source Nodes: [], Original ATen: []
# Source node to ATen node mapping:
triton_poi_fused_4 = async_compile.triton('triton_poi_fused_4', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[32768, 16], tile_hint=TileHint.SQUARE,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: 'i32', 3: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_4', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 1, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_4(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    ynumel = 32768
    xnumel = 9
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y3 = yindex
    y0 = yindex % 128
    y1 = (yindex // 128)
    tmp0 = tl.load(in_ptr0 + (x2 + (9*y3)), xmask, eviction_policy='evict_last')
    tl.store(out_ptr0 + (y0 + (128*x2) + (1152*y1)), tmp0, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/wu/cwu76v5k3qykuedq2kwod6ewkatwicgnt222bneru5mrfolaws3b.py
# Unsorted Source Nodes: [], Original ATen: []
# Source node to ATen node mapping:
triton_poi_fused_5 = async_compile.triton('triton_poi_fused_5', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[65536, 16], tile_hint=TileHint.SQUARE,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: 'i32', 3: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_5', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 1, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_5(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    ynumel = 65536
    xnumel = 9
    yoffset = (tl.program_id(1) + tl.program_id(2) * tl.num_programs(1)) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y3 = yindex
    y0 = yindex % 256
    y1 = (yindex // 256)
    tmp0 = tl.load(in_ptr0 + (x2 + (9*y3)), xmask, eviction_policy='evict_last')
    tl.store(out_ptr0 + (y0 + (256*x2) + (2304*y1)), tmp0, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/qc/cqcnwdigorioolv3gd37hby6ssaofoagcgegs6zyzz2tzdydz7r3.py
# Unsorted Source Nodes: [], Original ATen: []
# Source node to ATen node mapping:
triton_poi_fused_6 = async_compile.triton('triton_poi_fused_6', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[131072, 16], tile_hint=TileHint.SQUARE,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: 'i32', 3: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_6', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 1, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_6(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    ynumel = 131072
    xnumel = 9
    yoffset = (tl.program_id(1) + tl.program_id(2) * tl.num_programs(1)) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y3 = yindex
    y0 = yindex % 256
    y1 = (yindex // 256)
    tmp0 = tl.load(in_ptr0 + (x2 + (9*y3)), xmask, eviction_policy='evict_last')
    tl.store(out_ptr0 + (y0 + (256*x2) + (2304*y1)), tmp0, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/mv/cmvl76eg7bes5atj6xgol54kxqhtjnd2g6nvulp4v4n7jkcg7kt6.py
# Unsorted Source Nodes: [], Original ATen: []
# Source node to ATen node mapping:
triton_poi_fused_7 = async_compile.triton('triton_poi_fused_7', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[262144, 16], tile_hint=TileHint.SQUARE,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: 'i32', 3: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_7', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 1, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_7(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    ynumel = 262144
    xnumel = 9
    yoffset = (tl.program_id(1) + tl.program_id(2) * tl.num_programs(1)) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y3 = yindex
    y0 = yindex % 512
    y1 = (yindex // 512)
    tmp0 = tl.load(in_ptr0 + (x2 + (9*y3)), xmask, eviction_policy='evict_last')
    tl.store(out_ptr0 + (y0 + (512*x2) + (4608*y1)), tmp0, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/ve/cveuvmjivs2j3hozpkoamfacfv3pk35fk44c2uhmxr36r5jmxnwa.py
# Topologically Sorted Source Nodes: [output_1], Original ATen: [aten.view]
# Source node to ATen node mapping:
#   output_1 => view
# Graph fragment:
#   %view : [num_users=2] = call_function[target=torch.ops.aten.reshape.default](args = (%primals_1, [-1, 3, 64, 64]), kwargs = {})
triton_poi_fused_view_8 = async_compile.triton('triton_poi_fused_view_8', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[16, 4096], tile_hint=TileHint.SQUARE,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: 'i32', 3: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 3), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_view_8', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 1, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_view_8(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    ynumel = 12
    xnumel = 4096
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = yindex < ynumel
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = tl.full([XBLOCK, YBLOCK], True, tl.int1)
    x2 = xindex
    y3 = yindex
    y0 = yindex % 3
    y1 = (yindex // 3)
    tmp0 = tl.load(in_ptr0 + (x2 + (4096*y3)), ymask, eviction_policy='evict_last')
    tl.store(out_ptr0 + (y0 + (3*x2) + (12288*y1)), tmp0, ymask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/rb/crbcpoenob43a6eznxwv5whebropvznt67ku5dtspf4ukokcg5ti.py
# Topologically Sorted Source Nodes: [output_2], Original ATen: [aten.convolution]
# Source node to ATen node mapping:
#   output_2 => convolution
# Graph fragment:
#   %convolution : [num_users=5] = call_function[target=torch.ops.aten.convolution.default](args = (%view, %primals_2, %primals_3, [1, 1], [1, 1], [1, 1], False, [0, 0], 1), kwargs = {})
triton_poi_fused_convolution_9 = async_compile.triton('triton_poi_fused_convolution_9', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[1048576], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_convolution_9', 'mutated_arg_names': ['in_out_ptr0'], 'no_x_dim': False, 'num_load': 2, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_convolution_9(in_out_ptr0, in_ptr0, xnumel, XBLOCK : tl.constexpr):
    xnumel = 1048576
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = tl.full([XBLOCK], True, tl.int1)
    x2 = xindex
    x0 = xindex % 64
    tmp0 = tl.load(in_out_ptr0 + (x2), None)
    tmp1 = tl.load(in_ptr0 + (x0), None, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tl.store(in_out_ptr0 + (x2), tmp2, None)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/ty/ctycparcvrzotu7vcwq3aou6btdg3ik53e3ulfl34hmlnyadiggk.py
# Topologically Sorted Source Nodes: [add, add_1, add_2, output_3], Original ATen: [aten.add, aten.div]
# Source node to ATen node mapping:
#   add => add
#   add_1 => add_1
#   add_2 => add_2
#   output_3 => div
# Graph fragment:
#   %add : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%slice_4, %slice_8), kwargs = {})
#   %add_1 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%add, %slice_12), kwargs = {})
#   %add_2 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%add_1, %slice_16), kwargs = {})
#   %div : [num_users=2] = call_function[target=torch.ops.aten.div.Tensor](args = (%add_2, 4), kwargs = {})
triton_poi_fused_add_div_10 = async_compile.triton('triton_poi_fused_add_div_10', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[262144], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_add_div_10', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 4, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_add_div_10(in_ptr0, out_ptr0, xnumel, XBLOCK : tl.constexpr):
    xnumel = 262144
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = tl.full([XBLOCK], True, tl.int1)
    x0 = xindex % 64
    x1 = (xindex // 64) % 32
    x2 = (xindex // 2048)
    x3 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + (128*x1) + (8192*x2)), None)
    tmp1 = tl.load(in_ptr0 + (4096 + x0 + (128*x1) + (8192*x2)), None)
    tmp3 = tl.load(in_ptr0 + (64 + x0 + (128*x1) + (8192*x2)), None)
    tmp5 = tl.load(in_ptr0 + (4160 + x0 + (128*x1) + (8192*x2)), None)
    tmp2 = tmp0 + tmp1
    tmp4 = tmp2 + tmp3
    tmp6 = tmp4 + tmp5
    tmp7 = 0.25
    tmp8 = tmp6 * tmp7
    tl.store(out_ptr0 + (x3), tmp8, None)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/4g/c4gclmoqb6epmybyisrgafec4pjthbjsaiifk6uhni4dfjgdebva.py
# Topologically Sorted Source Nodes: [output_5], Original ATen: [aten.native_layer_norm]
# Source node to ATen node mapping:
#   output_5 => var_mean
# Graph fragment:
#   %var_mean : [num_users=2] = call_function[target=torch.ops.aten.var_mean.correction](args = (%convolution, [1, 2, 3]), kwargs = {correction: 0, keepdim: True})
triton_per_fused_native_layer_norm_11 = async_compile.triton('triton_per_fused_native_layer_norm_11', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.persistent_reduction(
    size_hints=[8192, 128],
    reduction_hint=ReductionHint.INNER,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: 'i32', 5: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_per_fused_native_layer_norm_11', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 1, 'num_reduction': 5, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False}
)
@triton.jit
def triton_per_fused_native_layer_norm_11(in_ptr0, out_ptr0, out_ptr1, out_ptr2, xnumel, rnumel, XBLOCK : tl.constexpr):
    xnumel = 8192
    rnumel = 128
    RBLOCK: tl.constexpr = 128
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = tl.full([XBLOCK, RBLOCK], True, tl.int1)
    rindex = tl.arange(0, RBLOCK)[None, :]
    roffset = 0
    rmask = tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r3 = rindex
    x0 = xindex % 32
    x1 = (xindex // 32) % 64
    x2 = (xindex // 2048)
    x4 = xindex
    tmp0 = tl.load(in_ptr0 + ((2*x0) + (64*(r3 % 64)) + (4096*(((r3 + (128*x1)) // 64) % 64)) + (262144*x2) + ((r3 + (128*x1)) // 4096)), None, eviction_policy='evict_last')
    tmp1 = tl.broadcast_to(tmp0, [XBLOCK, RBLOCK])
    tmp3 = tl.broadcast_to(tmp1, [XBLOCK, RBLOCK])
    tmp5 = tl.sum(tmp3, 1)[:, None]
    tmp6 = tl.full([XBLOCK, 1], 128, tl.int32)
    tmp7 = tmp6.to(tl.float32)
    tmp8 = tmp5 / tmp7
    tmp9 = tmp1 - tmp8
    tmp10 = tmp9 * tmp9
    tmp11 = tl.broadcast_to(tmp10, [XBLOCK, RBLOCK])
    tmp13 = tl.sum(tmp11, 1)[:, None]
    tl.store(out_ptr0 + (x4), tmp8, None)
    tl.store(out_ptr1 + (x4), tmp13, None)
    tl.store(out_ptr2 + (x4), tmp7, None)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/ek/cekhypacuiewh3wcjubkxn3ldjt3zbe2f2bulbuapps35f6ffk4f.py
# Topologically Sorted Source Nodes: [output_5], Original ATen: [aten.native_layer_norm]
# Source node to ATen node mapping:
#   output_5 => var_mean
# Graph fragment:
#   %var_mean : [num_users=2] = call_function[target=torch.ops.aten.var_mean.correction](args = (%convolution, [1, 2, 3]), kwargs = {correction: 0, keepdim: True})
triton_per_fused_native_layer_norm_12 = async_compile.triton('triton_per_fused_native_layer_norm_12', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.persistent_reduction(
    size_hints=[128, 64],
    reduction_hint=ReductionHint.INNER,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: '*fp32', 5: '*fp32', 6: 'i32', 7: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5, 6, 7), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_per_fused_native_layer_norm_12', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 3, 'num_reduction': 3, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False}
)
@triton.jit
def triton_per_fused_native_layer_norm_12(in_ptr0, in_ptr1, in_ptr2, out_ptr0, out_ptr1, out_ptr2, xnumel, rnumel, XBLOCK : tl.constexpr):
    xnumel = 128
    rnumel = 64
    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, :]
    roffset = 0
    rmask = tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r2 = rindex
    x0 = xindex % 32
    x1 = (xindex // 32)
    x3 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + (32*r2) + (2048*x1)), xmask, other=0.0)
    tmp1 = tl.load(in_ptr1 + (x0 + (32*r2) + (2048*x1)), xmask, other=0.0)
    tmp2 = tl.load(in_ptr2 + (x0 + (32*r2) + (2048*x1)), 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]
    tmp15 = tmp12[:, None]
    tl.store(out_ptr0 + (x3), tmp13, xmask)
    tl.store(out_ptr1 + (x3), tmp14, xmask)
    tl.store(out_ptr2 + (x3), tmp15, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/ls/clsuuh4bjbkuusigtbskl3bqvmae5njtd4xgvyg6pfnk3jcbjxsg.py
# Topologically Sorted Source Nodes: [output_5], Original ATen: [aten.native_layer_norm]
# Source node to ATen node mapping:
#   output_5 => add_3, rsqrt, var_mean
# Graph fragment:
#   %var_mean : [num_users=2] = call_function[target=torch.ops.aten.var_mean.correction](args = (%convolution, [1, 2, 3]), kwargs = {correction: 0, keepdim: True})
#   %add_3 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%getitem, 1e-05), kwargs = {})
#   %rsqrt : [num_users=2] = call_function[target=torch.ops.aten.rsqrt.default](args = (%add_3,), kwargs = {})
triton_per_fused_native_layer_norm_13 = async_compile.triton('triton_per_fused_native_layer_norm_13', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.persistent_reduction(
    size_hints=[4, 32],
    reduction_hint=ReductionHint.INNER,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: '*fp32', 5: 'i32', 6: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 6), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_per_fused_native_layer_norm_13', 'mutated_arg_names': ['in_out_ptr0'], 'no_x_dim': False, 'num_load': 3, 'num_reduction': 2, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False}
)
@triton.jit
def triton_per_fused_native_layer_norm_13(in_out_ptr0, in_ptr0, in_ptr1, in_ptr2, out_ptr0, xnumel, rnumel, XBLOCK : tl.constexpr):
    xnumel = 4
    rnumel = 32
    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, :]
    roffset = 0
    rmask = tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r1 = rindex
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + (r1 + (32*x0)), xmask, other=0.0)
    tmp1 = tl.load(in_ptr1 + (r1 + (32*x0)), xmask, other=0.0)
    tmp2 = tl.load(in_ptr2 + (r1 + (32*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]
    tmp15 = tmp12[:, None]
    tmp16 = 262144.0
    tmp17 = tmp14 / tmp16
    tmp18 = 1e-05
    tmp19 = tmp17 + tmp18
    tmp20 = libdevice.rsqrt(tmp19)
    tl.debug_barrier()
    tl.store(in_out_ptr0 + (x0), tmp20, xmask)
    tl.store(out_ptr0 + (x0), tmp13, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/c5/cc5y5r4worxkwpf5psagtt4y3sm4xik7ldm3kdhkhz3yzkehusqn.py
# Topologically Sorted Source Nodes: [output_5, output_6], Original ATen: [aten.native_layer_norm, aten.relu]
# Source node to ATen node mapping:
#   output_5 => add_4, mul, mul_1, sub
#   output_6 => relu
# Graph fragment:
#   %sub : [num_users=1] = call_function[target=torch.ops.aten.sub.Tensor](args = (%convolution, %getitem_1), kwargs = {})
#   %mul : [num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%sub, %rsqrt), kwargs = {})
#   %mul_1 : [num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%mul, %primals_6), kwargs = {})
#   %add_4 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%mul_1, %primals_7), kwargs = {})
#   %relu : [num_users=2] = call_function[target=torch.ops.aten.relu.default](args = (%add_4,), kwargs = {})
triton_poi_fused_native_layer_norm_relu_14 = async_compile.triton('triton_poi_fused_native_layer_norm_relu_14', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[256, 4096], tile_hint=TileHint.DEFAULT,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: '*fp32', 5: '*fp32', 6: 'i32', 7: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5, 6, 7), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_native_layer_norm_relu_14', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 5, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_native_layer_norm_relu_14(in_ptr0, in_ptr1, in_ptr2, in_ptr3, in_ptr4, out_ptr0, ynumel, xnumel, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    ynumel = 256
    xnumel = 4096
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = yindex < ynumel
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = tl.full([XBLOCK, YBLOCK], True, tl.int1)
    x2 = xindex
    y0 = yindex % 64
    y1 = (yindex // 64)
    tmp0 = tl.load(in_ptr0 + (y0 + (64*x2) + (262144*y1)), ymask, eviction_policy='evict_last')
    tmp1 = tl.load(in_ptr1 + (y1), ymask, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr2 + (y1), ymask, eviction_policy='evict_last')
    tmp5 = tl.load(in_ptr3 + (x2 + (4096*y0)), ymask, eviction_policy='evict_last')
    tmp7 = tl.load(in_ptr4 + (x2 + (4096*y0)), ymask, eviction_policy='evict_last')
    tmp2 = tmp0 - tmp1
    tmp4 = tmp2 * tmp3
    tmp6 = tmp4 * tmp5
    tmp8 = tmp6 + tmp7
    tmp9 = tl.full([1, 1], 0, tl.int32)
    tmp10 = triton_helpers.maximum(tmp9, tmp8)
    tl.store(out_ptr0 + (y0 + (64*x2) + (262144*y1)), tmp10, ymask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/vb/cvbwszludzdze53oohssaok6zdrj2fuaroqpeelolqqkdtwkryhp.py
# Topologically Sorted Source Nodes: [output_4, add_3, add_4, add_5, output_11, output_12], Original ATen: [aten.convolution, aten.add, aten.div]
# Source node to ATen node mapping:
#   add_3 => add_7
#   add_4 => add_8
#   add_5 => add_9
#   output_11 => div_1
#   output_12 => add_10
#   output_4 => convolution_1
# Graph fragment:
#   %convolution_1 : [num_users=1] = call_function[target=torch.ops.aten.convolution.default](args = (%div, %primals_4, %primals_5, [1, 1], [0, 0], [1, 1], False, [0, 0], 1), kwargs = {})
#   %add_7 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%slice_20, %slice_24), kwargs = {})
#   %add_8 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%add_7, %slice_28), kwargs = {})
#   %add_9 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%add_8, %slice_32), kwargs = {})
#   %div_1 : [num_users=1] = call_function[target=torch.ops.aten.div.Tensor](args = (%add_9, 4), kwargs = {})
#   %add_10 : [num_users=4] = call_function[target=torch.ops.aten.add.Tensor](args = (%convolution_1, %div_1), kwargs = {})
triton_poi_fused_add_convolution_div_15 = async_compile.triton('triton_poi_fused_add_convolution_div_15', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[524288], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_add_convolution_div_15', 'mutated_arg_names': ['in_out_ptr0'], 'no_x_dim': False, 'num_load': 7, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_add_convolution_div_15(in_out_ptr0, in_ptr0, in_ptr1, in_ptr2, xnumel, XBLOCK : tl.constexpr):
    xnumel = 524288
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = tl.full([XBLOCK], True, tl.int1)
    x3 = xindex
    x0 = xindex % 128
    x1 = (xindex // 128) % 32
    x2 = (xindex // 4096)
    tmp0 = tl.load(in_out_ptr0 + (x3), None)
    tmp1 = tl.load(in_ptr0 + (x0), None, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr1 + (x0 + (256*x1) + (16384*x2)), None)
    tmp4 = tl.load(in_ptr2 + (x0), None, eviction_policy='evict_last')
    tmp6 = tl.load(in_ptr1 + (8192 + x0 + (256*x1) + (16384*x2)), None)
    tmp9 = tl.load(in_ptr1 + (128 + x0 + (256*x1) + (16384*x2)), None)
    tmp12 = tl.load(in_ptr1 + (8320 + x0 + (256*x1) + (16384*x2)), None)
    tmp2 = tmp0 + tmp1
    tmp5 = tmp3 + tmp4
    tmp7 = tmp6 + tmp4
    tmp8 = tmp5 + tmp7
    tmp10 = tmp9 + tmp4
    tmp11 = tmp8 + tmp10
    tmp13 = tmp12 + tmp4
    tmp14 = tmp11 + tmp13
    tmp15 = 0.25
    tmp16 = tmp14 * tmp15
    tmp17 = tmp2 + tmp16
    tl.store(in_out_ptr0 + (x3), tmp17, None)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/xs/cxsp4rxlz5wry2p4nhkugd7cglt4u3cqfe2impeaj4u6upjyvfae.py
# Topologically Sorted Source Nodes: [add_7, add_8, add_9, output_13], Original ATen: [aten.add, aten.div]
# Source node to ATen node mapping:
#   add_7 => add_11
#   add_8 => add_12
#   add_9 => add_13
#   output_13 => div_2
# Graph fragment:
#   %add_11 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%slice_36, %slice_40), kwargs = {})
#   %add_12 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%add_11, %slice_44), kwargs = {})
#   %add_13 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%add_12, %slice_48), kwargs = {})
#   %div_2 : [num_users=2] = call_function[target=torch.ops.aten.div.Tensor](args = (%add_13, 4), kwargs = {})
triton_poi_fused_add_div_16 = async_compile.triton('triton_poi_fused_add_div_16', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[131072], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_add_div_16', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 4, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_add_div_16(in_ptr0, out_ptr0, xnumel, XBLOCK : tl.constexpr):
    xnumel = 131072
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = tl.full([XBLOCK], True, tl.int1)
    x0 = xindex % 128
    x1 = (xindex // 128) % 16
    x2 = (xindex // 2048)
    x3 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + (256*x1) + (8192*x2)), None)
    tmp1 = tl.load(in_ptr0 + (4096 + x0 + (256*x1) + (8192*x2)), None)
    tmp3 = tl.load(in_ptr0 + (128 + x0 + (256*x1) + (8192*x2)), None)
    tmp5 = tl.load(in_ptr0 + (4224 + x0 + (256*x1) + (8192*x2)), None)
    tmp2 = tmp0 + tmp1
    tmp4 = tmp2 + tmp3
    tmp6 = tmp4 + tmp5
    tmp7 = 0.25
    tmp8 = tmp6 * tmp7
    tl.store(out_ptr0 + (x3), tmp8, None)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/5h/c5hhfuf5llp3hp6ihn4gotkpsvne3uans7ynm3ntkzr4wvdwpc5h.py
# Topologically Sorted Source Nodes: [output_15], Original ATen: [aten.native_layer_norm]
# Source node to ATen node mapping:
#   output_15 => var_mean_2
# Graph fragment:
#   %var_mean_2 : [num_users=2] = call_function[target=torch.ops.aten.var_mean.correction](args = (%add_10, [1, 2, 3]), kwargs = {correction: 0, keepdim: True})
triton_per_fused_native_layer_norm_17 = async_compile.triton('triton_per_fused_native_layer_norm_17', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.persistent_reduction(
    size_hints=[4096, 128],
    reduction_hint=ReductionHint.INNER,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: 'i32', 5: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_per_fused_native_layer_norm_17', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 1, 'num_reduction': 5, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False}
)
@triton.jit
def triton_per_fused_native_layer_norm_17(in_ptr0, out_ptr0, out_ptr1, out_ptr2, xnumel, rnumel, XBLOCK : tl.constexpr):
    xnumel = 4096
    rnumel = 128
    RBLOCK: tl.constexpr = 128
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = tl.full([XBLOCK, RBLOCK], True, tl.int1)
    rindex = tl.arange(0, RBLOCK)[None, :]
    roffset = 0
    rmask = tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r3 = rindex
    x0 = xindex % 16
    x1 = (xindex // 16) % 64
    x2 = (xindex // 1024)
    x4 = xindex
    tmp0 = tl.load(in_ptr0 + ((8*x0) + (128*(r3 % 32)) + (4096*(((r3 + (128*x1)) // 32) % 32)) + (131072*x2) + ((r3 + (128*x1)) // 1024)), None, eviction_policy='evict_last')
    tmp1 = tl.broadcast_to(tmp0, [XBLOCK, RBLOCK])
    tmp3 = tl.broadcast_to(tmp1, [XBLOCK, RBLOCK])
    tmp5 = tl.sum(tmp3, 1)[:, None]
    tmp6 = tl.full([XBLOCK, 1], 128, tl.int32)
    tmp7 = tmp6.to(tl.float32)
    tmp8 = tmp5 / tmp7
    tmp9 = tmp1 - tmp8
    tmp10 = tmp9 * tmp9
    tmp11 = tl.broadcast_to(tmp10, [XBLOCK, RBLOCK])
    tmp13 = tl.sum(tmp11, 1)[:, None]
    tl.store(out_ptr0 + (x4), tmp8, None)
    tl.store(out_ptr1 + (x4), tmp13, None)
    tl.store(out_ptr2 + (x4), tmp7, None)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/2g/c2gbre7yucwpv2on2cuyzjbhajtnkevcfmzexd6yp3i6cu7gmghz.py
# Topologically Sorted Source Nodes: [output_15], Original ATen: [aten.native_layer_norm]
# Source node to ATen node mapping:
#   output_15 => var_mean_2
# Graph fragment:
#   %var_mean_2 : [num_users=2] = call_function[target=torch.ops.aten.var_mean.correction](args = (%add_10, [1, 2, 3]), kwargs = {correction: 0, keepdim: True})
triton_per_fused_native_layer_norm_18 = async_compile.triton('triton_per_fused_native_layer_norm_18', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.persistent_reduction(
    size_hints=[64, 64],
    reduction_hint=ReductionHint.INNER,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: '*fp32', 5: '*fp32', 6: 'i32', 7: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5, 6, 7), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_per_fused_native_layer_norm_18', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 3, 'num_reduction': 3, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False}
)
@triton.jit
def triton_per_fused_native_layer_norm_18(in_ptr0, in_ptr1, in_ptr2, out_ptr0, out_ptr1, out_ptr2, xnumel, rnumel, XBLOCK : tl.constexpr):
    xnumel = 64
    rnumel = 64
    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, :]
    roffset = 0
    rmask = tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r2 = rindex
    x0 = xindex % 16
    x1 = (xindex // 16)
    x3 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + (16*r2) + (1024*x1)), xmask, other=0.0)
    tmp1 = tl.load(in_ptr1 + (x0 + (16*r2) + (1024*x1)), xmask, other=0.0)
    tmp2 = tl.load(in_ptr2 + (x0 + (16*r2) + (1024*x1)), 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]
    tmp15 = tmp12[:, None]
    tl.store(out_ptr0 + (x3), tmp13, xmask)
    tl.store(out_ptr1 + (x3), tmp14, xmask)
    tl.store(out_ptr2 + (x3), tmp15, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/vv/cvvmbq3chmn3xa3e75fiba7uazpfuow2u3js3vkjb7tnao5w6anu.py
# Topologically Sorted Source Nodes: [output_15], Original ATen: [aten.native_layer_norm, aten.native_layer_norm_backward]
# Source node to ATen node mapping:
#   output_15 => add_14, rsqrt_2, var_mean_2
# Graph fragment:
#   %var_mean_2 : [num_users=2] = call_function[target=torch.ops.aten.var_mean.correction](args = (%add_10, [1, 2, 3]), kwargs = {correction: 0, keepdim: True})
#   %add_14 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%getitem_4, 1e-05), kwargs = {})
#   %rsqrt_2 : [num_users=2] = call_function[target=torch.ops.aten.rsqrt.default](args = (%add_14,), kwargs = {})
#   %div_18 : [num_users=1] = call_function[target=torch.ops.aten.div.Tensor](args = (%rsqrt_2, 131072), kwargs = {})
triton_per_fused_native_layer_norm_native_layer_norm_backward_19 = async_compile.triton('triton_per_fused_native_layer_norm_native_layer_norm_backward_19', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.persistent_reduction(
    size_hints=[4, 16],
    reduction_hint=ReductionHint.INNER,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: '*fp32', 5: '*fp32', 6: 'i32', 7: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5, 7), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_per_fused_native_layer_norm_native_layer_norm_backward_19', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 3, 'num_reduction': 2, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False}
)
@triton.jit
def triton_per_fused_native_layer_norm_native_layer_norm_backward_19(in_ptr0, in_ptr1, in_ptr2, out_ptr0, out_ptr1, out_ptr2, xnumel, rnumel, XBLOCK : tl.constexpr):
    xnumel = 4
    rnumel = 16
    RBLOCK: tl.constexpr = 16
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    rindex = tl.arange(0, RBLOCK)[None, :]
    roffset = 0
    rmask = tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r1 = rindex
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + (r1 + (16*x0)), xmask, other=0.0)
    tmp1 = tl.load(in_ptr1 + (r1 + (16*x0)), xmask, other=0.0)
    tmp2 = tl.load(in_ptr2 + (r1 + (16*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]
    tmp15 = tmp12[:, None]
    tmp16 = 131072.0
    tmp17 = tmp14 / tmp16
    tmp18 = 1e-05
    tmp19 = tmp17 + tmp18
    tmp20 = libdevice.rsqrt(tmp19)
    tmp21 = 7.62939453125e-06
    tmp22 = tmp20 * tmp21
    tl.store(out_ptr2 + (x0), tmp22, xmask)
    tl.store(out_ptr0 + (x0), tmp13, xmask)
    tl.store(out_ptr1 + (x0), tmp14, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/zb/czbu6vhwq66osjuqkwlhcfjhywqltx7u2ozqhhge5fv5z4fdfops.py
# Topologically Sorted Source Nodes: [output_15], Original ATen: [aten.native_layer_norm]
# Source node to ATen node mapping:
#   output_15 => add_14, mul_4, rsqrt_2, sub_2, var_mean_2
# Graph fragment:
#   %var_mean_2 : [num_users=2] = call_function[target=torch.ops.aten.var_mean.correction](args = (%add_10, [1, 2, 3]), kwargs = {correction: 0, keepdim: True})
#   %add_14 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%getitem_4, 1e-05), kwargs = {})
#   %rsqrt_2 : [num_users=2] = call_function[target=torch.ops.aten.rsqrt.default](args = (%add_14,), kwargs = {})
#   %sub_2 : [num_users=1] = call_function[target=torch.ops.aten.sub.Tensor](args = (%add_10, %getitem_5), kwargs = {})
#   %mul_4 : [num_users=2] = call_function[target=torch.ops.aten.mul.Tensor](args = (%sub_2, %rsqrt_2), kwargs = {})
triton_poi_fused_native_layer_norm_20 = async_compile.triton('triton_poi_fused_native_layer_norm_20', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[524288], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_native_layer_norm_20', 'mutated_arg_names': ['in_out_ptr0'], 'no_x_dim': False, 'num_load': 3, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_native_layer_norm_20(in_out_ptr0, in_ptr0, in_ptr1, xnumel, XBLOCK : tl.constexpr):
    xnumel = 524288
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = tl.full([XBLOCK], True, tl.int1)
    x2 = xindex
    x1 = (xindex // 131072)
    tmp0 = tl.load(in_out_ptr0 + (x2), None)
    tmp1 = tl.load(in_ptr0 + (x1), None, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr1 + (x1), None, eviction_policy='evict_last')
    tmp2 = tmp0 - tmp1
    tmp4 = 131072.0
    tmp5 = tmp3 / tmp4
    tmp6 = 1e-05
    tmp7 = tmp5 + tmp6
    tmp8 = libdevice.rsqrt(tmp7)
    tmp9 = tmp2 * tmp8
    tl.store(in_out_ptr0 + (x2), tmp9, None)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/r4/cr4f3ds4ycxdmppitnzzaehb6q76tci7ek34wp6wlo5rur3vby4f.py
# Topologically Sorted Source Nodes: [output_15, output_16], Original ATen: [aten.native_layer_norm, aten.relu]
# Source node to ATen node mapping:
#   output_15 => add_15, mul_5
#   output_16 => relu_2
# Graph fragment:
#   %mul_5 : [num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%mul_4, %primals_15), kwargs = {})
#   %add_15 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%mul_5, %primals_16), kwargs = {})
#   %relu_2 : [num_users=2] = call_function[target=torch.ops.aten.relu.default](args = (%add_15,), kwargs = {})
triton_poi_fused_native_layer_norm_relu_21 = async_compile.triton('triton_poi_fused_native_layer_norm_relu_21', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[512, 1024], tile_hint=TileHint.DEFAULT,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: 'i32', 5: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_native_layer_norm_relu_21', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 3, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_native_layer_norm_relu_21(in_ptr0, in_ptr1, in_ptr2, out_ptr0, ynumel, xnumel, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    ynumel = 512
    xnumel = 1024
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = yindex < ynumel
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y0 = yindex % 128
    y1 = (yindex // 128)
    tmp0 = tl.load(in_ptr0 + (y0 + (128*x2) + (131072*y1)), xmask & ymask, eviction_policy='evict_last')
    tmp1 = tl.load(in_ptr1 + (x2 + (1024*y0)), xmask & ymask, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr2 + (x2 + (1024*y0)), xmask & ymask, eviction_policy='evict_last')
    tmp2 = tmp0 * tmp1
    tmp4 = tmp2 + tmp3
    tmp5 = tl.full([1, 1], 0, tl.int32)
    tmp6 = triton_helpers.maximum(tmp5, tmp4)
    tl.store(out_ptr0 + (y0 + (128*x2) + (131072*y1)), tmp6, xmask & ymask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/if/ciftkdf7qfoumdtbvbupl6vawwgedvo55e2eeb6bdjcogrvxaca5.py
# Topologically Sorted Source Nodes: [output_18], Original ATen: [aten.native_layer_norm]
# Source node to ATen node mapping:
#   output_18 => add_16, rsqrt_3, var_mean_3
# Graph fragment:
#   %var_mean_3 : [num_users=2] = call_function[target=torch.ops.aten.var_mean.correction](args = (%convolution_5, [1, 2, 3]), kwargs = {correction: 0, keepdim: True})
#   %add_16 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%getitem_6, 1e-05), kwargs = {})
#   %rsqrt_3 : [num_users=2] = call_function[target=torch.ops.aten.rsqrt.default](args = (%add_16,), kwargs = {})
triton_per_fused_native_layer_norm_22 = async_compile.triton('triton_per_fused_native_layer_norm_22', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.persistent_reduction(
    size_hints=[4, 16],
    reduction_hint=ReductionHint.INNER,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: '*fp32', 5: 'i32', 6: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 6), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_per_fused_native_layer_norm_22', 'mutated_arg_names': ['in_out_ptr0'], 'no_x_dim': False, 'num_load': 3, 'num_reduction': 2, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False}
)
@triton.jit
def triton_per_fused_native_layer_norm_22(in_out_ptr0, in_ptr0, in_ptr1, in_ptr2, out_ptr0, xnumel, rnumel, XBLOCK : tl.constexpr):
    xnumel = 4
    rnumel = 16
    RBLOCK: tl.constexpr = 16
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    rindex = tl.arange(0, RBLOCK)[None, :]
    roffset = 0
    rmask = tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r1 = rindex
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + (r1 + (16*x0)), xmask, other=0.0)
    tmp1 = tl.load(in_ptr1 + (r1 + (16*x0)), xmask, other=0.0)
    tmp2 = tl.load(in_ptr2 + (r1 + (16*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]
    tmp15 = tmp12[:, None]
    tmp16 = 131072.0
    tmp17 = tmp14 / tmp16
    tmp18 = 1e-05
    tmp19 = tmp17 + tmp18
    tmp20 = libdevice.rsqrt(tmp19)
    tl.debug_barrier()
    tl.store(in_out_ptr0 + (x0), tmp20, xmask)
    tl.store(out_ptr0 + (x0), tmp13, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/ye/cyep4gzpczc7wycskgflj7u5mrzigyuowlpl55yyb6r4jkmh52ir.py
# Topologically Sorted Source Nodes: [output_18, output_19], Original ATen: [aten.native_layer_norm, aten.relu]
# Source node to ATen node mapping:
#   output_18 => add_17, mul_6, mul_7, sub_3
#   output_19 => relu_3
# Graph fragment:
#   %sub_3 : [num_users=1] = call_function[target=torch.ops.aten.sub.Tensor](args = (%convolution_5, %getitem_7), kwargs = {})
#   %mul_6 : [num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%sub_3, %rsqrt_3), kwargs = {})
#   %mul_7 : [num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%mul_6, %primals_18), kwargs = {})
#   %add_17 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%mul_7, %primals_19), kwargs = {})
#   %relu_3 : [num_users=2] = call_function[target=torch.ops.aten.relu.default](args = (%add_17,), kwargs = {})
triton_poi_fused_native_layer_norm_relu_23 = async_compile.triton('triton_poi_fused_native_layer_norm_relu_23', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[512, 1024], tile_hint=TileHint.DEFAULT,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: '*fp32', 5: '*fp32', 6: 'i32', 7: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5, 6, 7), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_native_layer_norm_relu_23', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 5, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_native_layer_norm_relu_23(in_ptr0, in_ptr1, in_ptr2, in_ptr3, in_ptr4, out_ptr0, ynumel, xnumel, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    ynumel = 512
    xnumel = 1024
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = yindex < ynumel
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y0 = yindex % 128
    y1 = (yindex // 128)
    tmp0 = tl.load(in_ptr0 + (y0 + (128*x2) + (131072*y1)), xmask & ymask, eviction_policy='evict_last')
    tmp1 = tl.load(in_ptr1 + (y1), ymask, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr2 + (y1), ymask, eviction_policy='evict_last')
    tmp5 = tl.load(in_ptr3 + (x2 + (1024*y0)), xmask & ymask, eviction_policy='evict_last')
    tmp7 = tl.load(in_ptr4 + (x2 + (1024*y0)), xmask & ymask, eviction_policy='evict_last')
    tmp2 = tmp0 - tmp1
    tmp4 = tmp2 * tmp3
    tmp6 = tmp4 * tmp5
    tmp8 = tmp6 + tmp7
    tmp9 = tl.full([1, 1], 0, tl.int32)
    tmp10 = triton_helpers.maximum(tmp9, tmp8)
    tl.store(out_ptr0 + (y0 + (128*x2) + (131072*y1)), tmp10, xmask & ymask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/lh/clhmrmx7jbepxzntvg7r37dkp74kyidudrpiswbfnvyrhaozhxis.py
# Topologically Sorted Source Nodes: [output_14, add_10, add_11, add_12, output_21, output_22], Original ATen: [aten.convolution, aten.add, aten.div]
# Source node to ATen node mapping:
#   add_10 => add_18
#   add_11 => add_19
#   add_12 => add_20
#   output_14 => convolution_4
#   output_21 => div_3
#   output_22 => add_21
# Graph fragment:
#   %convolution_4 : [num_users=1] = call_function[target=torch.ops.aten.convolution.default](args = (%div_2, %primals_13, %primals_14, [1, 1], [0, 0], [1, 1], False, [0, 0], 1), kwargs = {})
#   %add_18 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%slice_52, %slice_56), kwargs = {})
#   %add_19 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%add_18, %slice_60), kwargs = {})
#   %add_20 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%add_19, %slice_64), kwargs = {})
#   %div_3 : [num_users=1] = call_function[target=torch.ops.aten.div.Tensor](args = (%add_20, 4), kwargs = {})
#   %add_21 : [num_users=4] = call_function[target=torch.ops.aten.add.Tensor](args = (%convolution_4, %div_3), kwargs = {})
triton_poi_fused_add_convolution_div_24 = async_compile.triton('triton_poi_fused_add_convolution_div_24', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[262144], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_add_convolution_div_24', 'mutated_arg_names': ['in_out_ptr0'], 'no_x_dim': False, 'num_load': 7, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_add_convolution_div_24(in_out_ptr0, in_ptr0, in_ptr1, in_ptr2, xnumel, XBLOCK : tl.constexpr):
    xnumel = 262144
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = tl.full([XBLOCK], True, tl.int1)
    x3 = xindex
    x0 = xindex % 256
    x1 = (xindex // 256) % 16
    x2 = (xindex // 4096)
    tmp0 = tl.load(in_out_ptr0 + (x3), None)
    tmp1 = tl.load(in_ptr0 + (x0), None, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr1 + (x0 + (512*x1) + (16384*x2)), None)
    tmp4 = tl.load(in_ptr2 + (x0), None, eviction_policy='evict_last')
    tmp6 = tl.load(in_ptr1 + (8192 + x0 + (512*x1) + (16384*x2)), None)
    tmp9 = tl.load(in_ptr1 + (256 + x0 + (512*x1) + (16384*x2)), None)
    tmp12 = tl.load(in_ptr1 + (8448 + x0 + (512*x1) + (16384*x2)), None)
    tmp2 = tmp0 + tmp1
    tmp5 = tmp3 + tmp4
    tmp7 = tmp6 + tmp4
    tmp8 = tmp5 + tmp7
    tmp10 = tmp9 + tmp4
    tmp11 = tmp8 + tmp10
    tmp13 = tmp12 + tmp4
    tmp14 = tmp11 + tmp13
    tmp15 = 0.25
    tmp16 = tmp14 * tmp15
    tmp17 = tmp2 + tmp16
    tl.store(in_out_ptr0 + (x3), tmp17, None)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/il/cilusw2anp3yj64iukiuryikrd5mjpkpmc63xeqvvclpdfynm4fp.py
# Topologically Sorted Source Nodes: [add_14, add_15, add_16, output_23], Original ATen: [aten.add, aten.div]
# Source node to ATen node mapping:
#   add_14 => add_22
#   add_15 => add_23
#   add_16 => add_24
#   output_23 => div_4
# Graph fragment:
#   %add_22 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%slice_68, %slice_72), kwargs = {})
#   %add_23 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%add_22, %slice_76), kwargs = {})
#   %add_24 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%add_23, %slice_80), kwargs = {})
#   %div_4 : [num_users=2] = call_function[target=torch.ops.aten.div.Tensor](args = (%add_24, 4), kwargs = {})
triton_poi_fused_add_div_25 = async_compile.triton('triton_poi_fused_add_div_25', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[65536], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_add_div_25', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 4, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_add_div_25(in_ptr0, out_ptr0, xnumel, XBLOCK : tl.constexpr):
    xnumel = 65536
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = tl.full([XBLOCK], True, tl.int1)
    x0 = xindex % 256
    x1 = (xindex // 256) % 8
    x2 = (xindex // 2048)
    x3 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + (512*x1) + (8192*x2)), None)
    tmp1 = tl.load(in_ptr0 + (4096 + x0 + (512*x1) + (8192*x2)), None)
    tmp3 = tl.load(in_ptr0 + (256 + x0 + (512*x1) + (8192*x2)), None)
    tmp5 = tl.load(in_ptr0 + (4352 + x0 + (512*x1) + (8192*x2)), None)
    tmp2 = tmp0 + tmp1
    tmp4 = tmp2 + tmp3
    tmp6 = tmp4 + tmp5
    tmp7 = 0.25
    tmp8 = tmp6 * tmp7
    tl.store(out_ptr0 + (x3), tmp8, None)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/bz/cbzdz3tuarnaiygtfp74mr76ojrpdcz5xckdknn74njonejzonkf.py
# Topologically Sorted Source Nodes: [output_25], Original ATen: [aten.native_layer_norm]
# Source node to ATen node mapping:
#   output_25 => var_mean_4
# Graph fragment:
#   %var_mean_4 : [num_users=2] = call_function[target=torch.ops.aten.var_mean.correction](args = (%add_21, [1, 2, 3]), kwargs = {correction: 0, keepdim: True})
triton_per_fused_native_layer_norm_26 = async_compile.triton('triton_per_fused_native_layer_norm_26', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.persistent_reduction(
    size_hints=[2048, 128],
    reduction_hint=ReductionHint.INNER,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: 'i32', 5: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_per_fused_native_layer_norm_26', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 1, 'num_reduction': 5, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False}
)
@triton.jit
def triton_per_fused_native_layer_norm_26(in_ptr0, out_ptr0, out_ptr1, out_ptr2, xnumel, rnumel, XBLOCK : tl.constexpr):
    xnumel = 2048
    rnumel = 128
    RBLOCK: tl.constexpr = 128
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = tl.full([XBLOCK, RBLOCK], True, tl.int1)
    rindex = tl.arange(0, RBLOCK)[None, :]
    roffset = 0
    rmask = tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r3 = rindex
    x0 = xindex % 8
    x1 = (xindex // 8) % 64
    x2 = (xindex // 512)
    x4 = xindex
    tmp0 = tl.load(in_ptr0 + ((32*x0) + (256*(r3 % 16)) + (4096*(((r3 + (128*x1)) // 16) % 16)) + (65536*x2) + ((r3 + (128*x1)) // 256)), None, eviction_policy='evict_last')
    tmp1 = tl.broadcast_to(tmp0, [XBLOCK, RBLOCK])
    tmp3 = tl.broadcast_to(tmp1, [XBLOCK, RBLOCK])
    tmp5 = tl.sum(tmp3, 1)[:, None]
    tmp6 = tl.full([XBLOCK, 1], 128, tl.int32)
    tmp7 = tmp6.to(tl.float32)
    tmp8 = tmp5 / tmp7
    tmp9 = tmp1 - tmp8
    tmp10 = tmp9 * tmp9
    tmp11 = tl.broadcast_to(tmp10, [XBLOCK, RBLOCK])
    tmp13 = tl.sum(tmp11, 1)[:, None]
    tl.store(out_ptr0 + (x4), tmp8, None)
    tl.store(out_ptr1 + (x4), tmp13, None)
    tl.store(out_ptr2 + (x4), tmp7, None)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/li/clibis357s7sopbpor2nbk7azfabwgbrwpqwyiwe6khhhafxyi62.py
# Topologically Sorted Source Nodes: [output_25], Original ATen: [aten.native_layer_norm]
# Source node to ATen node mapping:
#   output_25 => var_mean_4
# Graph fragment:
#   %var_mean_4 : [num_users=2] = call_function[target=torch.ops.aten.var_mean.correction](args = (%add_21, [1, 2, 3]), kwargs = {correction: 0, keepdim: True})
triton_per_fused_native_layer_norm_27 = async_compile.triton('triton_per_fused_native_layer_norm_27', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.persistent_reduction(
    size_hints=[32, 64],
    reduction_hint=ReductionHint.INNER,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: '*fp32', 5: '*fp32', 6: 'i32', 7: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5, 6, 7), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_per_fused_native_layer_norm_27', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 3, 'num_reduction': 3, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False}
)
@triton.jit
def triton_per_fused_native_layer_norm_27(in_ptr0, in_ptr1, in_ptr2, out_ptr0, out_ptr1, out_ptr2, xnumel, rnumel, XBLOCK : tl.constexpr):
    xnumel = 32
    rnumel = 64
    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, :]
    roffset = 0
    rmask = tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r2 = rindex
    x0 = xindex % 8
    x1 = (xindex // 8)
    x3 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + (8*r2) + (512*x1)), xmask, other=0.0)
    tmp1 = tl.load(in_ptr1 + (x0 + (8*r2) + (512*x1)), xmask, other=0.0)
    tmp2 = tl.load(in_ptr2 + (x0 + (8*r2) + (512*x1)), 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]
    tmp15 = tmp12[:, None]
    tl.store(out_ptr0 + (x3), tmp13, xmask)
    tl.store(out_ptr1 + (x3), tmp14, xmask)
    tl.store(out_ptr2 + (x3), tmp15, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/kb/ckb7hkzz5gbdm7sqdka7jrk62mkno3dot6batylguwoc2aw2ey4o.py
# Topologically Sorted Source Nodes: [output_25], Original ATen: [aten.native_layer_norm, aten.native_layer_norm_backward]
# Source node to ATen node mapping:
#   output_25 => add_25, rsqrt_4, var_mean_4
# Graph fragment:
#   %var_mean_4 : [num_users=2] = call_function[target=torch.ops.aten.var_mean.correction](args = (%add_21, [1, 2, 3]), kwargs = {correction: 0, keepdim: True})
#   %add_25 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%getitem_8, 1e-05), kwargs = {})
#   %rsqrt_4 : [num_users=2] = call_function[target=torch.ops.aten.rsqrt.default](args = (%add_25,), kwargs = {})
#   %div_14 : [num_users=1] = call_function[target=torch.ops.aten.div.Tensor](args = (%rsqrt_4, 65536), kwargs = {})
triton_per_fused_native_layer_norm_native_layer_norm_backward_28 = async_compile.triton('triton_per_fused_native_layer_norm_native_layer_norm_backward_28', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.persistent_reduction(
    size_hints=[4, 8],
    reduction_hint=ReductionHint.INNER,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: '*fp32', 5: '*fp32', 6: 'i32', 7: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_per_fused_native_layer_norm_native_layer_norm_backward_28', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 3, 'num_reduction': 2, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False}
)
@triton.jit
def triton_per_fused_native_layer_norm_native_layer_norm_backward_28(in_ptr0, in_ptr1, in_ptr2, out_ptr0, out_ptr1, out_ptr2, xnumel, rnumel, XBLOCK : tl.constexpr):
    xnumel = 4
    rnumel = 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, :]
    roffset = 0
    rmask = tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r1 = rindex
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + (r1 + (8*x0)), xmask, other=0.0)
    tmp1 = tl.load(in_ptr1 + (r1 + (8*x0)), xmask, other=0.0)
    tmp2 = tl.load(in_ptr2 + (r1 + (8*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]
    tmp15 = tmp12[:, None]
    tmp16 = 65536.0
    tmp17 = tmp14 / tmp16
    tmp18 = 1e-05
    tmp19 = tmp17 + tmp18
    tmp20 = libdevice.rsqrt(tmp19)
    tmp21 = 1.52587890625e-05
    tmp22 = tmp20 * tmp21
    tl.store(out_ptr2 + (x0), tmp22, xmask)
    tl.store(out_ptr0 + (x0), tmp13, xmask)
    tl.store(out_ptr1 + (x0), tmp14, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/kg/ckgdvaxpopdb3fe35qadraus5smfssbvpcxkxl4hbc3kbjyhefo6.py
# Topologically Sorted Source Nodes: [output_25], Original ATen: [aten.native_layer_norm]
# Source node to ATen node mapping:
#   output_25 => add_25, mul_8, rsqrt_4, sub_4, var_mean_4
# Graph fragment:
#   %var_mean_4 : [num_users=2] = call_function[target=torch.ops.aten.var_mean.correction](args = (%add_21, [1, 2, 3]), kwargs = {correction: 0, keepdim: True})
#   %add_25 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%getitem_8, 1e-05), kwargs = {})
#   %rsqrt_4 : [num_users=2] = call_function[target=torch.ops.aten.rsqrt.default](args = (%add_25,), kwargs = {})
#   %sub_4 : [num_users=1] = call_function[target=torch.ops.aten.sub.Tensor](args = (%add_21, %getitem_9), kwargs = {})
#   %mul_8 : [num_users=2] = call_function[target=torch.ops.aten.mul.Tensor](args = (%sub_4, %rsqrt_4), kwargs = {})
triton_poi_fused_native_layer_norm_29 = async_compile.triton('triton_poi_fused_native_layer_norm_29', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[262144], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_native_layer_norm_29', 'mutated_arg_names': ['in_out_ptr0'], 'no_x_dim': False, 'num_load': 3, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_native_layer_norm_29(in_out_ptr0, in_ptr0, in_ptr1, xnumel, XBLOCK : tl.constexpr):
    xnumel = 262144
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = tl.full([XBLOCK], True, tl.int1)
    x2 = xindex
    x1 = (xindex // 65536)
    tmp0 = tl.load(in_out_ptr0 + (x2), None)
    tmp1 = tl.load(in_ptr0 + (x1), None, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr1 + (x1), None, eviction_policy='evict_last')
    tmp2 = tmp0 - tmp1
    tmp4 = 65536.0
    tmp5 = tmp3 / tmp4
    tmp6 = 1e-05
    tmp7 = tmp5 + tmp6
    tmp8 = libdevice.rsqrt(tmp7)
    tmp9 = tmp2 * tmp8
    tl.store(in_out_ptr0 + (x2), tmp9, None)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/tj/ctjclajsdp3de7el3lchd4ukevwslhdet2xkfhg3gzvkil76z5dp.py
# Topologically Sorted Source Nodes: [output_25, output_26], Original ATen: [aten.native_layer_norm, aten.relu]
# Source node to ATen node mapping:
#   output_25 => add_26, mul_9
#   output_26 => relu_4
# Graph fragment:
#   %mul_9 : [num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%mul_8, %primals_24), kwargs = {})
#   %add_26 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%mul_9, %primals_25), kwargs = {})
#   %relu_4 : [num_users=2] = call_function[target=torch.ops.aten.relu.default](args = (%add_26,), kwargs = {})
triton_poi_fused_native_layer_norm_relu_30 = async_compile.triton('triton_poi_fused_native_layer_norm_relu_30', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[1024, 256], tile_hint=TileHint.DEFAULT,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: 'i32', 5: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_native_layer_norm_relu_30', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 3, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_native_layer_norm_relu_30(in_ptr0, in_ptr1, in_ptr2, out_ptr0, ynumel, xnumel, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    ynumel = 1024
    xnumel = 256
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y0 = yindex % 256
    y1 = (yindex // 256)
    tmp0 = tl.load(in_ptr0 + (y0 + (256*x2) + (65536*y1)), xmask, eviction_policy='evict_last')
    tmp1 = tl.load(in_ptr1 + (x2 + (256*y0)), xmask, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr2 + (x2 + (256*y0)), xmask, eviction_policy='evict_last')
    tmp2 = tmp0 * tmp1
    tmp4 = tmp2 + tmp3
    tmp5 = tl.full([1, 1], 0, tl.int32)
    tmp6 = triton_helpers.maximum(tmp5, tmp4)
    tl.store(out_ptr0 + (y0 + (256*x2) + (65536*y1)), tmp6, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/54/c5444b5nerk747lh5yx7vmux7r5n32zgot7c2bcrkxrr2pypb7fz.py
# Topologically Sorted Source Nodes: [output_28], Original ATen: [aten.native_layer_norm]
# Source node to ATen node mapping:
#   output_28 => add_27, rsqrt_5, var_mean_5
# Graph fragment:
#   %var_mean_5 : [num_users=2] = call_function[target=torch.ops.aten.var_mean.correction](args = (%convolution_8, [1, 2, 3]), kwargs = {correction: 0, keepdim: True})
#   %add_27 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%getitem_10, 1e-05), kwargs = {})
#   %rsqrt_5 : [num_users=2] = call_function[target=torch.ops.aten.rsqrt.default](args = (%add_27,), kwargs = {})
triton_per_fused_native_layer_norm_31 = async_compile.triton('triton_per_fused_native_layer_norm_31', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.persistent_reduction(
    size_hints=[4, 8],
    reduction_hint=ReductionHint.INNER,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: '*fp32', 5: 'i32', 6: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_per_fused_native_layer_norm_31', 'mutated_arg_names': ['in_out_ptr0'], 'no_x_dim': False, 'num_load': 3, 'num_reduction': 2, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False}
)
@triton.jit
def triton_per_fused_native_layer_norm_31(in_out_ptr0, in_ptr0, in_ptr1, in_ptr2, out_ptr0, xnumel, rnumel, XBLOCK : tl.constexpr):
    xnumel = 4
    rnumel = 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, :]
    roffset = 0
    rmask = tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r1 = rindex
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + (r1 + (8*x0)), xmask, other=0.0)
    tmp1 = tl.load(in_ptr1 + (r1 + (8*x0)), xmask, other=0.0)
    tmp2 = tl.load(in_ptr2 + (r1 + (8*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]
    tmp15 = tmp12[:, None]
    tmp16 = 65536.0
    tmp17 = tmp14 / tmp16
    tmp18 = 1e-05
    tmp19 = tmp17 + tmp18
    tmp20 = libdevice.rsqrt(tmp19)
    tl.debug_barrier()
    tl.store(in_out_ptr0 + (x0), tmp20, xmask)
    tl.store(out_ptr0 + (x0), tmp13, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/y5/cy5rbi2j7f5oxhldvtztdbbxvbhae3hv55frx5v4gnhlvv7ovqqf.py
# Topologically Sorted Source Nodes: [output_28, output_29], Original ATen: [aten.native_layer_norm, aten.relu]
# Source node to ATen node mapping:
#   output_28 => add_28, mul_10, mul_11, sub_5
#   output_29 => relu_5
# Graph fragment:
#   %sub_5 : [num_users=1] = call_function[target=torch.ops.aten.sub.Tensor](args = (%convolution_8, %getitem_11), kwargs = {})
#   %mul_10 : [num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%sub_5, %rsqrt_5), kwargs = {})
#   %mul_11 : [num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%mul_10, %primals_27), kwargs = {})
#   %add_28 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%mul_11, %primals_28), kwargs = {})
#   %relu_5 : [num_users=2] = call_function[target=torch.ops.aten.relu.default](args = (%add_28,), kwargs = {})
triton_poi_fused_native_layer_norm_relu_32 = async_compile.triton('triton_poi_fused_native_layer_norm_relu_32', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[1024, 256], tile_hint=TileHint.DEFAULT,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: '*fp32', 5: '*fp32', 6: 'i32', 7: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5, 6, 7), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_native_layer_norm_relu_32', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 5, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_native_layer_norm_relu_32(in_ptr0, in_ptr1, in_ptr2, in_ptr3, in_ptr4, out_ptr0, ynumel, xnumel, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    ynumel = 1024
    xnumel = 256
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y0 = yindex % 256
    y1 = (yindex // 256)
    tmp0 = tl.load(in_ptr0 + (y0 + (256*x2) + (65536*y1)), xmask, eviction_policy='evict_last')
    tmp1 = tl.load(in_ptr1 + (y1), None, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr2 + (y1), None, eviction_policy='evict_last')
    tmp5 = tl.load(in_ptr3 + (x2 + (256*y0)), xmask, eviction_policy='evict_last')
    tmp7 = tl.load(in_ptr4 + (x2 + (256*y0)), xmask, eviction_policy='evict_last')
    tmp2 = tmp0 - tmp1
    tmp4 = tmp2 * tmp3
    tmp6 = tmp4 * tmp5
    tmp8 = tmp6 + tmp7
    tmp9 = tl.full([1, 1], 0, tl.int32)
    tmp10 = triton_helpers.maximum(tmp9, tmp8)
    tl.store(out_ptr0 + (y0 + (256*x2) + (65536*y1)), tmp10, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/5b/c5bjs23zf242vy4grfagtkpvkltveo34gwvwzxfh3qehycpzx76i.py
# Topologically Sorted Source Nodes: [output_24, add_17, add_18, add_19, output_31, output_32], Original ATen: [aten.convolution, aten.add, aten.div]
# Source node to ATen node mapping:
#   add_17 => add_29
#   add_18 => add_30
#   add_19 => add_31
#   output_24 => convolution_7
#   output_31 => div_5
#   output_32 => add_32
# Graph fragment:
#   %convolution_7 : [num_users=1] = call_function[target=torch.ops.aten.convolution.default](args = (%div_4, %primals_22, %primals_23, [1, 1], [0, 0], [1, 1], False, [0, 0], 1), kwargs = {})
#   %add_29 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%slice_84, %slice_88), kwargs = {})
#   %add_30 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%add_29, %slice_92), kwargs = {})
#   %add_31 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%add_30, %slice_96), kwargs = {})
#   %div_5 : [num_users=1] = call_function[target=torch.ops.aten.div.Tensor](args = (%add_31, 4), kwargs = {})
#   %add_32 : [num_users=4] = call_function[target=torch.ops.aten.add.Tensor](args = (%convolution_7, %div_5), kwargs = {})
triton_poi_fused_add_convolution_div_33 = async_compile.triton('triton_poi_fused_add_convolution_div_33', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[131072], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_add_convolution_div_33', 'mutated_arg_names': ['in_out_ptr0'], 'no_x_dim': False, 'num_load': 7, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_add_convolution_div_33(in_out_ptr0, in_ptr0, in_ptr1, in_ptr2, xnumel, XBLOCK : tl.constexpr):
    xnumel = 131072
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = tl.full([XBLOCK], True, tl.int1)
    x3 = xindex
    x0 = xindex % 512
    x1 = (xindex // 512) % 8
    x2 = (xindex // 4096)
    tmp0 = tl.load(in_out_ptr0 + (x3), None)
    tmp1 = tl.load(in_ptr0 + (x0), None, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr1 + (x0 + (1024*x1) + (16384*x2)), None)
    tmp4 = tl.load(in_ptr2 + (x0), None, eviction_policy='evict_last')
    tmp6 = tl.load(in_ptr1 + (8192 + x0 + (1024*x1) + (16384*x2)), None)
    tmp9 = tl.load(in_ptr1 + (512 + x0 + (1024*x1) + (16384*x2)), None)
    tmp12 = tl.load(in_ptr1 + (8704 + x0 + (1024*x1) + (16384*x2)), None)
    tmp2 = tmp0 + tmp1
    tmp5 = tmp3 + tmp4
    tmp7 = tmp6 + tmp4
    tmp8 = tmp5 + tmp7
    tmp10 = tmp9 + tmp4
    tmp11 = tmp8 + tmp10
    tmp13 = tmp12 + tmp4
    tmp14 = tmp11 + tmp13
    tmp15 = 0.25
    tmp16 = tmp14 * tmp15
    tmp17 = tmp2 + tmp16
    tl.store(in_out_ptr0 + (x3), tmp17, None)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/x4/cx452hmilix63dphhq66is5cre2qakcycukghzslonzf7d25juda.py
# Topologically Sorted Source Nodes: [add_21, add_22, add_23, output_33], Original ATen: [aten.add, aten.div]
# Source node to ATen node mapping:
#   add_21 => add_33
#   add_22 => add_34
#   add_23 => add_35
#   output_33 => div_6
# Graph fragment:
#   %add_33 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%slice_100, %slice_104), kwargs = {})
#   %add_34 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%add_33, %slice_108), kwargs = {})
#   %add_35 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%add_34, %slice_112), kwargs = {})
#   %div_6 : [num_users=2] = call_function[target=torch.ops.aten.div.Tensor](args = (%add_35, 4), kwargs = {})
triton_poi_fused_add_div_34 = async_compile.triton('triton_poi_fused_add_div_34', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[32768], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_add_div_34', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 4, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_add_div_34(in_ptr0, out_ptr0, xnumel, XBLOCK : tl.constexpr):
    xnumel = 32768
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = tl.full([XBLOCK], True, tl.int1)
    x0 = xindex % 512
    x1 = (xindex // 512) % 4
    x2 = (xindex // 2048)
    x3 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + (1024*x1) + (8192*x2)), None)
    tmp1 = tl.load(in_ptr0 + (4096 + x0 + (1024*x1) + (8192*x2)), None)
    tmp3 = tl.load(in_ptr0 + (512 + x0 + (1024*x1) + (8192*x2)), None)
    tmp5 = tl.load(in_ptr0 + (4608 + x0 + (1024*x1) + (8192*x2)), None)
    tmp2 = tmp0 + tmp1
    tmp4 = tmp2 + tmp3
    tmp6 = tmp4 + tmp5
    tmp7 = 0.25
    tmp8 = tmp6 * tmp7
    tl.store(out_ptr0 + (x3), tmp8, None)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/2y/c2yvtpggflzc26vsl2ze47vkze2d4smrgmyvcgasvgu4bn4toiqe.py
# Topologically Sorted Source Nodes: [output_35], Original ATen: [aten.native_layer_norm]
# Source node to ATen node mapping:
#   output_35 => var_mean_6
# Graph fragment:
#   %var_mean_6 : [num_users=2] = call_function[target=torch.ops.aten.var_mean.correction](args = (%add_32, [1, 2, 3]), kwargs = {correction: 0, keepdim: True})
triton_per_fused_native_layer_norm_35 = async_compile.triton('triton_per_fused_native_layer_norm_35', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.persistent_reduction(
    size_hints=[1024, 128],
    reduction_hint=ReductionHint.INNER,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: 'i32', 5: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_per_fused_native_layer_norm_35', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 1, 'num_reduction': 5, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False}
)
@triton.jit
def triton_per_fused_native_layer_norm_35(in_ptr0, out_ptr0, out_ptr1, out_ptr2, xnumel, rnumel, XBLOCK : tl.constexpr):
    xnumel = 1024
    rnumel = 128
    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, :]
    roffset = 0
    rmask = tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r2 = rindex
    x0 = xindex % 256
    x1 = (xindex // 256)
    x3 = xindex
    tmp0 = tl.load(in_ptr0 + ((2*x0) + (512*(r2 % 64)) + (32768*x1) + (r2 // 64)), xmask, eviction_policy='evict_last', other=0.0)
    tmp1 = tl.broadcast_to(tmp0, [XBLOCK, RBLOCK])
    tmp3 = 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], 128, 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]
    tl.store(out_ptr0 + (x3), tmp10, xmask)
    tl.store(out_ptr1 + (x3), tmp16, xmask)
    tl.store(out_ptr2 + (x3), tmp9, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/dj/cdjsp2qz2ulj4fjoal7hcy4thngy677m6e5pkz67brwn7hx7ldcx.py
# Topologically Sorted Source Nodes: [output_35], Original ATen: [aten.native_layer_norm]
# Source node to ATen node mapping:
#   output_35 => var_mean_6
# Graph fragment:
#   %var_mean_6 : [num_users=2] = call_function[target=torch.ops.aten.var_mean.correction](args = (%add_32, [1, 2, 3]), kwargs = {correction: 0, keepdim: True})
triton_per_fused_native_layer_norm_36 = async_compile.triton('triton_per_fused_native_layer_norm_36', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.persistent_reduction(
    size_hints=[16, 64],
    reduction_hint=ReductionHint.INNER,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: '*fp32', 5: '*fp32', 6: 'i32', 7: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5, 6, 7), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_per_fused_native_layer_norm_36', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 3, 'num_reduction': 3, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False}
)
@triton.jit
def triton_per_fused_native_layer_norm_36(in_ptr0, in_ptr1, in_ptr2, out_ptr0, out_ptr1, out_ptr2, xnumel, rnumel, XBLOCK : tl.constexpr):
    xnumel = 16
    rnumel = 64
    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, :]
    roffset = 0
    rmask = tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r1 = rindex
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + (r1 + (64*x0)), xmask, other=0.0)
    tmp1 = tl.load(in_ptr1 + (r1 + (64*x0)), xmask, other=0.0)
    tmp2 = tl.load(in_ptr2 + (r1 + (64*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]
    tmp15 = tmp12[:, None]
    tl.store(out_ptr0 + (x0), tmp13, xmask)
    tl.store(out_ptr1 + (x0), tmp14, xmask)
    tl.store(out_ptr2 + (x0), tmp15, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/vd/cvd7ajuy77b7jihf6tgodxqccvvuo6tafktb65qfokdsr7i3ahpt.py
# Topologically Sorted Source Nodes: [output_35], Original ATen: [aten.native_layer_norm, aten.native_layer_norm_backward]
# Source node to ATen node mapping:
#   output_35 => add_36, rsqrt_6, var_mean_6
# Graph fragment:
#   %var_mean_6 : [num_users=2] = call_function[target=torch.ops.aten.var_mean.correction](args = (%add_32, [1, 2, 3]), kwargs = {correction: 0, keepdim: True})
#   %add_36 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%getitem_12, 1e-05), kwargs = {})
#   %rsqrt_6 : [num_users=2] = call_function[target=torch.ops.aten.rsqrt.default](args = (%add_36,), kwargs = {})
#   %div_10 : [num_users=1] = call_function[target=torch.ops.aten.div.Tensor](args = (%rsqrt_6, 32768), kwargs = {})
triton_per_fused_native_layer_norm_native_layer_norm_backward_37 = async_compile.triton('triton_per_fused_native_layer_norm_native_layer_norm_backward_37', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.persistent_reduction(
    size_hints=[4, 4],
    reduction_hint=ReductionHint.INNER,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: '*fp32', 5: '*fp32', 6: 'i32', 7: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_per_fused_native_layer_norm_native_layer_norm_backward_37', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 3, 'num_reduction': 2, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False}
)
@triton.jit
def triton_per_fused_native_layer_norm_native_layer_norm_backward_37(in_ptr0, in_ptr1, in_ptr2, out_ptr0, out_ptr1, out_ptr2, xnumel, rnumel, XBLOCK : tl.constexpr):
    xnumel = 4
    rnumel = 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, :]
    roffset = 0
    rmask = 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]
    tmp15 = tmp12[:, None]
    tmp16 = 32768.0
    tmp17 = tmp14 / tmp16
    tmp18 = 1e-05
    tmp19 = tmp17 + tmp18
    tmp20 = libdevice.rsqrt(tmp19)
    tmp21 = 3.0517578125e-05
    tmp22 = tmp20 * tmp21
    tl.store(out_ptr2 + (x0), tmp22, xmask)
    tl.store(out_ptr0 + (x0), tmp13, xmask)
    tl.store(out_ptr1 + (x0), tmp14, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/f6/cf6wbbpc7lp5olc2yrx2524wnjgkepvyxyb6tedcjlmvoafaxtel.py
# Topologically Sorted Source Nodes: [output_35], Original ATen: [aten.native_layer_norm]
# Source node to ATen node mapping:
#   output_35 => add_36, mul_12, rsqrt_6, sub_6, var_mean_6
# Graph fragment:
#   %var_mean_6 : [num_users=2] = call_function[target=torch.ops.aten.var_mean.correction](args = (%add_32, [1, 2, 3]), kwargs = {correction: 0, keepdim: True})
#   %add_36 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%getitem_12, 1e-05), kwargs = {})
#   %rsqrt_6 : [num_users=2] = call_function[target=torch.ops.aten.rsqrt.default](args = (%add_36,), kwargs = {})
#   %sub_6 : [num_users=1] = call_function[target=torch.ops.aten.sub.Tensor](args = (%add_32, %getitem_13), kwargs = {})
#   %mul_12 : [num_users=2] = call_function[target=torch.ops.aten.mul.Tensor](args = (%sub_6, %rsqrt_6), kwargs = {})
triton_poi_fused_native_layer_norm_38 = async_compile.triton('triton_poi_fused_native_layer_norm_38', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[131072], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_native_layer_norm_38', 'mutated_arg_names': ['in_out_ptr0'], 'no_x_dim': False, 'num_load': 3, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_native_layer_norm_38(in_out_ptr0, in_ptr0, in_ptr1, xnumel, XBLOCK : tl.constexpr):
    xnumel = 131072
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = tl.full([XBLOCK], True, tl.int1)
    x2 = xindex
    x1 = (xindex // 32768)
    tmp0 = tl.load(in_out_ptr0 + (x2), None)
    tmp1 = tl.load(in_ptr0 + (x1), None, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr1 + (x1), None, eviction_policy='evict_last')
    tmp2 = tmp0 - tmp1
    tmp4 = 32768.0
    tmp5 = tmp3 / tmp4
    tmp6 = 1e-05
    tmp7 = tmp5 + tmp6
    tmp8 = libdevice.rsqrt(tmp7)
    tmp9 = tmp2 * tmp8
    tl.store(in_out_ptr0 + (x2), tmp9, None)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/pz/cpzfjfx3egsjclk5nt6vkyv5snw46qkemttiqrkdcznwciez5l56.py
# Topologically Sorted Source Nodes: [output_35, output_36], Original ATen: [aten.native_layer_norm, aten.relu]
# Source node to ATen node mapping:
#   output_35 => add_37, mul_13
#   output_36 => relu_6
# Graph fragment:
#   %mul_13 : [num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%mul_12, %primals_33), kwargs = {})
#   %add_37 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%mul_13, %primals_34), kwargs = {})
#   %relu_6 : [num_users=2] = call_function[target=torch.ops.aten.relu.default](args = (%add_37,), kwargs = {})
triton_poi_fused_native_layer_norm_relu_39 = async_compile.triton('triton_poi_fused_native_layer_norm_relu_39', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[2048, 64], tile_hint=TileHint.DEFAULT,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: 'i32', 5: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_native_layer_norm_relu_39', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 3, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_native_layer_norm_relu_39(in_ptr0, in_ptr1, in_ptr2, out_ptr0, ynumel, xnumel, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    ynumel = 2048
    xnumel = 64
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y0 = yindex % 512
    y1 = (yindex // 512)
    tmp0 = tl.load(in_ptr0 + (y0 + (512*x2) + (32768*y1)), xmask, eviction_policy='evict_last')
    tmp1 = tl.load(in_ptr1 + (x2 + (64*y0)), xmask, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr2 + (x2 + (64*y0)), xmask, eviction_policy='evict_last')
    tmp2 = tmp0 * tmp1
    tmp4 = tmp2 + tmp3
    tmp5 = tl.full([1, 1], 0, tl.int32)
    tmp6 = triton_helpers.maximum(tmp5, tmp4)
    tl.store(out_ptr0 + (y0 + (512*x2) + (32768*y1)), tmp6, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/aq/caqcuedjp4pnbu5aquyzwwrm7jjokxioy2cp72wxancl47k5woxh.py
# Topologically Sorted Source Nodes: [output_38], Original ATen: [aten.native_layer_norm]
# Source node to ATen node mapping:
#   output_38 => add_38, rsqrt_7, var_mean_7
# Graph fragment:
#   %var_mean_7 : [num_users=2] = call_function[target=torch.ops.aten.var_mean.correction](args = (%convolution_11, [1, 2, 3]), kwargs = {correction: 0, keepdim: True})
#   %add_38 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%getitem_14, 1e-05), kwargs = {})
#   %rsqrt_7 : [num_users=2] = call_function[target=torch.ops.aten.rsqrt.default](args = (%add_38,), kwargs = {})
triton_per_fused_native_layer_norm_40 = async_compile.triton('triton_per_fused_native_layer_norm_40', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.persistent_reduction(
    size_hints=[4, 4],
    reduction_hint=ReductionHint.INNER,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: '*fp32', 5: 'i32', 6: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_per_fused_native_layer_norm_40', 'mutated_arg_names': ['in_out_ptr0'], 'no_x_dim': False, 'num_load': 3, 'num_reduction': 2, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False}
)
@triton.jit
def triton_per_fused_native_layer_norm_40(in_out_ptr0, in_ptr0, in_ptr1, in_ptr2, out_ptr0, xnumel, rnumel, XBLOCK : tl.constexpr):
    xnumel = 4
    rnumel = 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, :]
    roffset = 0
    rmask = 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]
    tmp15 = tmp12[:, None]
    tmp16 = 32768.0
    tmp17 = tmp14 / tmp16
    tmp18 = 1e-05
    tmp19 = tmp17 + tmp18
    tmp20 = libdevice.rsqrt(tmp19)
    tl.debug_barrier()
    tl.store(in_out_ptr0 + (x0), tmp20, xmask)
    tl.store(out_ptr0 + (x0), tmp13, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/bf/cbfb6bzdoso2looxw3nury2nigxo7q3nputk5vc2az4vk4qpotqs.py
# Topologically Sorted Source Nodes: [output_38, output_39], Original ATen: [aten.native_layer_norm, aten.relu]
# Source node to ATen node mapping:
#   output_38 => add_39, mul_14, mul_15, sub_7
#   output_39 => relu_7
# Graph fragment:
#   %sub_7 : [num_users=1] = call_function[target=torch.ops.aten.sub.Tensor](args = (%convolution_11, %getitem_15), kwargs = {})
#   %mul_14 : [num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%sub_7, %rsqrt_7), kwargs = {})
#   %mul_15 : [num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%mul_14, %primals_36), kwargs = {})
#   %add_39 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%mul_15, %primals_37), kwargs = {})
#   %relu_7 : [num_users=2] = call_function[target=torch.ops.aten.relu.default](args = (%add_39,), kwargs = {})
triton_poi_fused_native_layer_norm_relu_41 = async_compile.triton('triton_poi_fused_native_layer_norm_relu_41', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[2048, 64], tile_hint=TileHint.DEFAULT,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: '*fp32', 5: '*fp32', 6: 'i32', 7: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5, 6, 7), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_native_layer_norm_relu_41', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 5, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_native_layer_norm_relu_41(in_ptr0, in_ptr1, in_ptr2, in_ptr3, in_ptr4, out_ptr0, ynumel, xnumel, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    ynumel = 2048
    xnumel = 64
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y0 = yindex % 512
    y1 = (yindex // 512)
    tmp0 = tl.load(in_ptr0 + (y0 + (512*x2) + (32768*y1)), xmask, eviction_policy='evict_last')
    tmp1 = tl.load(in_ptr1 + (y1), None, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr2 + (y1), None, eviction_policy='evict_last')
    tmp5 = tl.load(in_ptr3 + (x2 + (64*y0)), xmask, eviction_policy='evict_last')
    tmp7 = tl.load(in_ptr4 + (x2 + (64*y0)), xmask, eviction_policy='evict_last')
    tmp2 = tmp0 - tmp1
    tmp4 = tmp2 * tmp3
    tmp6 = tmp4 * tmp5
    tmp8 = tmp6 + tmp7
    tmp9 = tl.full([1, 1], 0, tl.int32)
    tmp10 = triton_helpers.maximum(tmp9, tmp8)
    tl.store(out_ptr0 + (y0 + (512*x2) + (32768*y1)), tmp10, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/m3/cm3bujlziwh7ohtxync7v6dnnqofm3d2y4bdvzj5ny2royc3pzlu.py
# Topologically Sorted Source Nodes: [output_34, add_24, add_25, add_26, output_41, output_42], Original ATen: [aten.convolution, aten.add, aten.div]
# Source node to ATen node mapping:
#   add_24 => add_40
#   add_25 => add_41
#   add_26 => add_42
#   output_34 => convolution_10
#   output_41 => div_7
#   output_42 => add_43
# Graph fragment:
#   %convolution_10 : [num_users=1] = call_function[target=torch.ops.aten.convolution.default](args = (%div_6, %primals_31, %primals_32, [1, 1], [0, 0], [1, 1], False, [0, 0], 1), kwargs = {})
#   %add_40 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%slice_116, %slice_120), kwargs = {})
#   %add_41 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%add_40, %slice_124), kwargs = {})
#   %add_42 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%add_41, %slice_128), kwargs = {})
#   %div_7 : [num_users=1] = call_function[target=torch.ops.aten.div.Tensor](args = (%add_42, 4), kwargs = {})
#   %add_43 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%convolution_10, %div_7), kwargs = {})
triton_poi_fused_add_convolution_div_42 = async_compile.triton('triton_poi_fused_add_convolution_div_42', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[64, 512], tile_hint=TileHint.DEFAULT,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: '*fp32', 5: 'i32', 6: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5, 6), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_add_convolution_div_42', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 7, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_add_convolution_div_42(in_ptr0, in_ptr1, in_ptr2, in_ptr3, out_ptr0, ynumel, xnumel, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    ynumel = 64
    xnumel = 512
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = yindex < ynumel
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x3 = xindex
    y4 = yindex
    y0 = yindex % 4
    y5 = (yindex // 4)
    y2 = (yindex // 16)
    y6 = yindex % 16
    tmp0 = tl.load(in_ptr0 + (x3 + (512*y4)), xmask & ymask, eviction_policy='evict_last')
    tmp1 = tl.load(in_ptr1 + (x3), xmask, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr2 + (x3 + (1024*y0) + (8192*y5)), xmask & ymask, eviction_policy='evict_last')
    tmp4 = tl.load(in_ptr3 + (x3), xmask, eviction_policy='evict_last')
    tmp6 = tl.load(in_ptr2 + (4096 + x3 + (1024*y0) + (8192*y5)), xmask & ymask, eviction_policy='evict_last')
    tmp9 = tl.load(in_ptr2 + (512 + x3 + (1024*y0) + (8192*y5)), xmask & ymask, eviction_policy='evict_last')
    tmp12 = tl.load(in_ptr2 + (4608 + x3 + (1024*y0) + (8192*y5)), xmask & ymask, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tmp5 = tmp3 + tmp4
    tmp7 = tmp6 + tmp4
    tmp8 = tmp5 + tmp7
    tmp10 = tmp9 + tmp4
    tmp11 = tmp8 + tmp10
    tmp13 = tmp12 + tmp4
    tmp14 = tmp11 + tmp13
    tmp15 = 0.25
    tmp16 = tmp14 * tmp15
    tmp17 = tmp2 + tmp16
    tl.store(out_ptr0 + (y6 + (16*x3) + (8192*y2)), tmp17, xmask & ymask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/rb/crb3cioqlsxtuhqenrqknr63jotj5whj373mclq56rcu3eg7y5gi.py
# Topologically Sorted Source Nodes: [output_45], Original ATen: [aten.tanh]
# Source node to ATen node mapping:
#   output_45 => tanh
# Graph fragment:
#   %add_tensor : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%mm_default, %primals_41), kwargs = {})
#   %tanh : [num_users=1] = call_function[target=torch.ops.aten.tanh.default](args = (%add_tensor,), kwargs = {})
triton_poi_fused_tanh_43 = async_compile.triton('triton_poi_fused_tanh_43', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[512], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_tanh_43', 'mutated_arg_names': ['in_out_ptr0'], 'no_x_dim': False, 'num_load': 2, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_tanh_43(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
    x2 = xindex
    x0 = xindex % 128
    tmp0 = tl.load(in_out_ptr0 + (x2), xmask)
    tmp1 = tl.load(in_ptr0 + (x0), xmask, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tmp3 = libdevice.tanh(tmp2)
    tl.store(in_out_ptr0 + (x2), tmp3, xmask)
''', device_str='cuda')


async_compile.wait(globals())
del async_compile

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, primals_32, primals_33, primals_34, primals_35, primals_36, primals_37, primals_38, primals_39, primals_40, primals_41 = args
    args.clear()
    assert_size_stride(primals_1, (4, 3, 64, 64), (12288, 4096, 64, 1))
    assert_size_stride(primals_2, (64, 3, 3, 3), (27, 9, 3, 1))
    assert_size_stride(primals_3, (64, ), (1, ))
    assert_size_stride(primals_4, (128, 64, 1, 1), (64, 1, 1, 1))
    assert_size_stride(primals_5, (128, ), (1, ))
    assert_size_stride(primals_6, (64, 64, 64), (4096, 64, 1))
    assert_size_stride(primals_7, (64, 64, 64), (4096, 64, 1))
    assert_size_stride(primals_8, (64, 64, 3, 3), (576, 9, 3, 1))
    assert_size_stride(primals_9, (64, 64, 64), (4096, 64, 1))
    assert_size_stride(primals_10, (64, 64, 64), (4096, 64, 1))
    assert_size_stride(primals_11, (128, 64, 3, 3), (576, 9, 3, 1))
    assert_size_stride(primals_12, (128, ), (1, ))
    assert_size_stride(primals_13, (256, 128, 1, 1), (128, 1, 1, 1))
    assert_size_stride(primals_14, (256, ), (1, ))
    assert_size_stride(primals_15, (128, 32, 32), (1024, 32, 1))
    assert_size_stride(primals_16, (128, 32, 32), (1024, 32, 1))
    assert_size_stride(primals_17, (128, 128, 3, 3), (1152, 9, 3, 1))
    assert_size_stride(primals_18, (128, 32, 32), (1024, 32, 1))
    assert_size_stride(primals_19, (128, 32, 32), (1024, 32, 1))
    assert_size_stride(primals_20, (256, 128, 3, 3), (1152, 9, 3, 1))
    assert_size_stride(primals_21, (256, ), (1, ))
    assert_size_stride(primals_22, (512, 256, 1, 1), (256, 1, 1, 1))
    assert_size_stride(primals_23, (512, ), (1, ))
    assert_size_stride(primals_24, (256, 16, 16), (256, 16, 1))
    assert_size_stride(primals_25, (256, 16, 16), (256, 16, 1))
    assert_size_stride(primals_26, (256, 256, 3, 3), (2304, 9, 3, 1))
    assert_size_stride(primals_27, (256, 16, 16), (256, 16, 1))
    assert_size_stride(primals_28, (256, 16, 16), (256, 16, 1))
    assert_size_stride(primals_29, (512, 256, 3, 3), (2304, 9, 3, 1))
    assert_size_stride(primals_30, (512, ), (1, ))
    assert_size_stride(primals_31, (512, 512, 1, 1), (512, 1, 1, 1))
    assert_size_stride(primals_32, (512, ), (1, ))
    assert_size_stride(primals_33, (512, 8, 8), (64, 8, 1))
    assert_size_stride(primals_34, (512, 8, 8), (64, 8, 1))
    assert_size_stride(primals_35, (512, 512, 3, 3), (4608, 9, 3, 1))
    assert_size_stride(primals_36, (512, 8, 8), (64, 8, 1))
    assert_size_stride(primals_37, (512, 8, 8), (64, 8, 1))
    assert_size_stride(primals_38, (512, 512, 3, 3), (4608, 9, 3, 1))
    assert_size_stride(primals_39, (512, ), (1, ))
    assert_size_stride(primals_40, (128, 8192), (8192, 1))
    assert_size_stride(primals_41, (128, ), (1, ))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((64, 3, 3, 3), (27, 1, 9, 3), torch.float32)
        # Unsorted Source Nodes: [], Original ATen: []
        stream0 = get_raw_stream(0)
        triton_poi_fused_0.run(primals_2, buf0, 192, 9, grid=grid(192, 9), stream=stream0)
        del primals_2
        buf1 = empty_strided_cuda((64, 64, 3, 3), (576, 1, 192, 64), torch.float32)
        # Unsorted Source Nodes: [], Original ATen: []
        triton_poi_fused_1.run(primals_8, buf1, 4096, 9, grid=grid(4096, 9), stream=stream0)
        del primals_8
        buf2 = empty_strided_cuda((128, 64, 3, 3), (576, 1, 192, 64), torch.float32)
        # Unsorted Source Nodes: [], Original ATen: []
        triton_poi_fused_2.run(primals_11, buf2, 8192, 9, grid=grid(8192, 9), stream=stream0)
        del primals_11
        buf3 = empty_strided_cuda((128, 128, 3, 3), (1152, 1, 384, 128), torch.float32)
        # Unsorted Source Nodes: [], Original ATen: []
        triton_poi_fused_3.run(primals_17, buf3, 16384, 9, grid=grid(16384, 9), stream=stream0)
        del primals_17
        buf4 = empty_strided_cuda((256, 128, 3, 3), (1152, 1, 384, 128), torch.float32)
        # Unsorted Source Nodes: [], Original ATen: []
        triton_poi_fused_4.run(primals_20, buf4, 32768, 9, grid=grid(32768, 9), stream=stream0)
        del primals_20
        buf5 = empty_strided_cuda((256, 256, 3, 3), (2304, 1, 768, 256), torch.float32)
        # Unsorted Source Nodes: [], Original ATen: []
        triton_poi_fused_5.run(primals_26, buf5, 65536, 9, grid=grid(65536, 9), stream=stream0)
        del primals_26
        buf6 = empty_strided_cuda((512, 256, 3, 3), (2304, 1, 768, 256), torch.float32)
        # Unsorted Source Nodes: [], Original ATen: []
        triton_poi_fused_6.run(primals_29, buf6, 131072, 9, grid=grid(131072, 9), stream=stream0)
        del primals_29
        buf7 = empty_strided_cuda((512, 512, 3, 3), (4608, 1, 1536, 512), torch.float32)
        # Unsorted Source Nodes: [], Original ATen: []
        triton_poi_fused_7.run(primals_35, buf7, 262144, 9, grid=grid(262144, 9), stream=stream0)
        del primals_35
        buf8 = empty_strided_cuda((512, 512, 3, 3), (4608, 1, 1536, 512), torch.float32)
        # Unsorted Source Nodes: [], Original ATen: []
        triton_poi_fused_7.run(primals_38, buf8, 262144, 9, grid=grid(262144, 9), stream=stream0)
        del primals_38
        buf9 = empty_strided_cuda((4, 3, 64, 64), (12288, 1, 192, 3), torch.float32)
        # Topologically Sorted Source Nodes: [output_1], Original ATen: [aten.view]
        triton_poi_fused_view_8.run(primals_1, buf9, 12, 4096, grid=grid(12, 4096), stream=stream0)
        del primals_1
        # Topologically Sorted Source Nodes: [output_2], Original ATen: [aten.convolution]
        buf10 = extern_kernels.convolution(buf9, buf0, stride=(1, 1), padding=(1, 1), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf10, (4, 64, 64, 64), (262144, 1, 4096, 64))
        buf11 = buf10; del buf10  # reuse
        # Topologically Sorted Source Nodes: [output_2], Original ATen: [aten.convolution]
        triton_poi_fused_convolution_9.run(buf11, primals_3, 1048576, grid=grid(1048576), stream=stream0)
        del primals_3
        buf12 = empty_strided_cuda((4, 64, 32, 32), (65536, 1, 2048, 64), torch.float32)
        # Topologically Sorted Source Nodes: [add, add_1, add_2, output_3], Original ATen: [aten.add, aten.div]
        triton_poi_fused_add_div_10.run(buf11, buf12, 262144, grid=grid(262144), stream=stream0)
        # Topologically Sorted Source Nodes: [output_4], Original ATen: [aten.convolution]
        buf13 = extern_kernels.convolution(buf12, primals_4, stride=(1, 1), padding=(0, 0), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf13, (4, 128, 32, 32), (131072, 1, 4096, 128))
        buf14 = empty_strided_cuda((4, 1, 1, 1, 32, 64), (2048, 8192, 8192, 8192, 1, 32), torch.float32)
        buf15 = empty_strided_cuda((4, 1, 1, 1, 32, 64), (2048, 8192, 8192, 8192, 1, 32), torch.float32)
        buf16 = empty_strided_cuda((4, 1, 1, 1, 32, 64), (2048, 8192, 8192, 8192, 1, 32), torch.float32)
        # Topologically Sorted Source Nodes: [output_5], Original ATen: [aten.native_layer_norm]
        triton_per_fused_native_layer_norm_11.run(buf11, buf14, buf15, buf16, 8192, 128, grid=grid(8192), stream=stream0)
        buf17 = empty_strided_cuda((4, 1, 1, 1, 32), (32, 128, 128, 128, 1), torch.float32)
        buf18 = empty_strided_cuda((4, 1, 1, 1, 32), (32, 128, 128, 128, 1), torch.float32)
        buf19 = empty_strided_cuda((4, 1, 1, 1, 32), (32, 128, 128, 128, 1), torch.float32)
        # Topologically Sorted Source Nodes: [output_5], Original ATen: [aten.native_layer_norm]
        triton_per_fused_native_layer_norm_12.run(buf14, buf15, buf16, buf17, buf18, buf19, 128, 64, grid=grid(128), stream=stream0)
        buf20 = empty_strided_cuda((4, 1, 1, 1), (1, 1, 1, 1), torch.float32)
        buf21 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf23 = reinterpret_tensor(buf21, (4, 1, 1, 1), (1, 1, 1, 1), 0); del buf21  # reuse
        # Topologically Sorted Source Nodes: [output_5], Original ATen: [aten.native_layer_norm]
        triton_per_fused_native_layer_norm_13.run(buf23, buf17, buf18, buf19, buf20, 4, 32, grid=grid(4), stream=stream0)
        buf24 = empty_strided_cuda((4, 64, 64, 64), (262144, 1, 4096, 64), torch.float32)
        # Topologically Sorted Source Nodes: [output_5, output_6], Original ATen: [aten.native_layer_norm, aten.relu]
        triton_poi_fused_native_layer_norm_relu_14.run(buf11, buf20, buf23, primals_6, primals_7, buf24, 256, 4096, grid=grid(256, 4096), stream=stream0)
        del primals_7
        # Topologically Sorted Source Nodes: [output_7], Original ATen: [aten.convolution]
        buf25 = extern_kernels.convolution(buf24, buf1, stride=(1, 1), padding=(1, 1), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf25, (4, 64, 64, 64), (262144, 1, 4096, 64))
        buf26 = buf16; del buf16  # reuse
        buf27 = buf15; del buf15  # reuse
        buf28 = buf14; del buf14  # reuse
        # Topologically Sorted Source Nodes: [output_8], Original ATen: [aten.native_layer_norm]
        triton_per_fused_native_layer_norm_11.run(buf25, buf26, buf27, buf28, 8192, 128, grid=grid(8192), stream=stream0)
        buf29 = buf19; del buf19  # reuse
        buf30 = buf18; del buf18  # reuse
        buf31 = buf17; del buf17  # reuse
        # Topologically Sorted Source Nodes: [output_8], Original ATen: [aten.native_layer_norm]
        triton_per_fused_native_layer_norm_12.run(buf26, buf27, buf28, buf29, buf30, buf31, 128, 64, grid=grid(128), stream=stream0)
        del buf26
        del buf27
        del buf28
        buf32 = empty_strided_cuda((4, 1, 1, 1), (1, 1, 1, 1), torch.float32)
        buf33 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf35 = reinterpret_tensor(buf33, (4, 1, 1, 1), (1, 1, 1, 1), 0); del buf33  # reuse
        # Topologically Sorted Source Nodes: [output_8], Original ATen: [aten.native_layer_norm]
        triton_per_fused_native_layer_norm_13.run(buf35, buf29, buf30, buf31, buf32, 4, 32, grid=grid(4), stream=stream0)
        del buf29
        del buf30
        del buf31
        buf36 = empty_strided_cuda((4, 64, 64, 64), (262144, 1, 4096, 64), torch.float32)
        # Topologically Sorted Source Nodes: [output_8, output_9], Original ATen: [aten.native_layer_norm, aten.relu]
        triton_poi_fused_native_layer_norm_relu_14.run(buf25, buf32, buf35, primals_9, primals_10, buf36, 256, 4096, grid=grid(256, 4096), stream=stream0)
        del primals_10
        # Topologically Sorted Source Nodes: [output_10], Original ATen: [aten.convolution]
        buf37 = extern_kernels.convolution(buf36, buf2, stride=(1, 1), padding=(1, 1), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf37, (4, 128, 64, 64), (524288, 1, 8192, 128))
        buf38 = buf13; del buf13  # reuse
        # Topologically Sorted Source Nodes: [output_4, add_3, add_4, add_5, output_11, output_12], Original ATen: [aten.convolution, aten.add, aten.div]
        triton_poi_fused_add_convolution_div_15.run(buf38, primals_5, buf37, primals_12, 524288, grid=grid(524288), stream=stream0)
        del buf37
        del primals_12
        del primals_5
        buf39 = empty_strided_cuda((4, 128, 16, 16), (32768, 1, 2048, 128), torch.float32)
        # Topologically Sorted Source Nodes: [add_7, add_8, add_9, output_13], Original ATen: [aten.add, aten.div]
        triton_poi_fused_add_div_16.run(buf38, buf39, 131072, grid=grid(131072), stream=stream0)
        # Topologically Sorted Source Nodes: [output_14], Original ATen: [aten.convolution]
        buf40 = extern_kernels.convolution(buf39, primals_13, stride=(1, 1), padding=(0, 0), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf40, (4, 256, 16, 16), (65536, 1, 4096, 256))
        buf41 = empty_strided_cuda((4, 1, 1, 1, 16, 64), (1024, 4096, 4096, 4096, 1, 16), torch.float32)
        buf42 = empty_strided_cuda((4, 1, 1, 1, 16, 64), (1024, 4096, 4096, 4096, 1, 16), torch.float32)
        buf43 = empty_strided_cuda((4, 1, 1, 1, 16, 64), (1024, 4096, 4096, 4096, 1, 16), torch.float32)
        # Topologically Sorted Source Nodes: [output_15], Original ATen: [aten.native_layer_norm]
        triton_per_fused_native_layer_norm_17.run(buf38, buf41, buf42, buf43, 4096, 128, grid=grid(4096), stream=stream0)
        buf44 = empty_strided_cuda((4, 1, 1, 1, 16), (16, 64, 64, 64, 1), torch.float32)
        buf45 = empty_strided_cuda((4, 1, 1, 1, 16), (16, 64, 64, 64, 1), torch.float32)
        buf46 = empty_strided_cuda((4, 1, 1, 1, 16), (16, 64, 64, 64, 1), torch.float32)
        # Topologically Sorted Source Nodes: [output_15], Original ATen: [aten.native_layer_norm]
        triton_per_fused_native_layer_norm_18.run(buf41, buf42, buf43, buf44, buf45, buf46, 64, 64, grid=grid(64), stream=stream0)
        buf47 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf48 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf124 = empty_strided_cuda((4, 1, 1, 1), (1, 1, 1, 1), torch.float32)
        # Topologically Sorted Source Nodes: [output_15], Original ATen: [aten.native_layer_norm, aten.native_layer_norm_backward]
        triton_per_fused_native_layer_norm_native_layer_norm_backward_19.run(buf44, buf45, buf46, buf47, buf48, buf124, 4, 16, grid=grid(4), stream=stream0)
        buf50 = buf38; del buf38  # reuse
        # Topologically Sorted Source Nodes: [output_15], Original ATen: [aten.native_layer_norm]
        triton_poi_fused_native_layer_norm_20.run(buf50, buf47, buf48, 524288, grid=grid(524288), stream=stream0)
        buf51 = empty_strided_cuda((4, 128, 32, 32), (131072, 1, 4096, 128), torch.float32)
        # Topologically Sorted Source Nodes: [output_15, output_16], Original ATen: [aten.native_layer_norm, aten.relu]
        triton_poi_fused_native_layer_norm_relu_21.run(buf50, primals_15, primals_16, buf51, 512, 1024, grid=grid(512, 1024), stream=stream0)
        del primals_16
        # Topologically Sorted Source Nodes: [output_17], Original ATen: [aten.convolution]
        buf52 = extern_kernels.convolution(buf51, buf3, stride=(1, 1), padding=(1, 1), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf52, (4, 128, 32, 32), (131072, 1, 4096, 128))
        buf53 = buf43; del buf43  # reuse
        buf54 = buf42; del buf42  # reuse
        buf55 = buf41; del buf41  # reuse
        # Topologically Sorted Source Nodes: [output_18], Original ATen: [aten.native_layer_norm]
        triton_per_fused_native_layer_norm_17.run(buf52, buf53, buf54, buf55, 4096, 128, grid=grid(4096), stream=stream0)
        buf56 = buf46; del buf46  # reuse
        buf57 = buf45; del buf45  # reuse
        buf58 = buf44; del buf44  # reuse
        # Topologically Sorted Source Nodes: [output_18], Original ATen: [aten.native_layer_norm]
        triton_per_fused_native_layer_norm_18.run(buf53, buf54, buf55, buf56, buf57, buf58, 64, 64, grid=grid(64), stream=stream0)
        del buf53
        del buf54
        del buf55
        buf59 = reinterpret_tensor(buf48, (4, 1, 1, 1), (1, 1, 1, 1), 0); del buf48  # reuse
        buf60 = buf47; del buf47  # reuse
        buf62 = reinterpret_tensor(buf60, (4, 1, 1, 1), (1, 1, 1, 1), 0); del buf60  # reuse
        # Topologically Sorted Source Nodes: [output_18], Original ATen: [aten.native_layer_norm]
        triton_per_fused_native_layer_norm_22.run(buf62, buf56, buf57, buf58, buf59, 4, 16, grid=grid(4), stream=stream0)
        del buf56
        del buf57
        del buf58
        buf63 = empty_strided_cuda((4, 128, 32, 32), (131072, 1, 4096, 128), torch.float32)
        # Topologically Sorted Source Nodes: [output_18, output_19], Original ATen: [aten.native_layer_norm, aten.relu]
        triton_poi_fused_native_layer_norm_relu_23.run(buf52, buf59, buf62, primals_18, primals_19, buf63, 512, 1024, grid=grid(512, 1024), stream=stream0)
        del primals_19
        # Topologically Sorted Source Nodes: [output_20], Original ATen: [aten.convolution]
        buf64 = extern_kernels.convolution(buf63, buf4, stride=(1, 1), padding=(1, 1), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf64, (4, 256, 32, 32), (262144, 1, 8192, 256))
        buf65 = buf40; del buf40  # reuse
        # Topologically Sorted Source Nodes: [output_14, add_10, add_11, add_12, output_21, output_22], Original ATen: [aten.convolution, aten.add, aten.div]
        triton_poi_fused_add_convolution_div_24.run(buf65, primals_14, buf64, primals_21, 262144, grid=grid(262144), stream=stream0)
        del buf64
        del primals_14
        del primals_21
        buf66 = empty_strided_cuda((4, 256, 8, 8), (16384, 1, 2048, 256), torch.float32)
        # Topologically Sorted Source Nodes: [add_14, add_15, add_16, output_23], Original ATen: [aten.add, aten.div]
        triton_poi_fused_add_div_25.run(buf65, buf66, 65536, grid=grid(65536), stream=stream0)
        # Topologically Sorted Source Nodes: [output_24], Original ATen: [aten.convolution]
        buf67 = extern_kernels.convolution(buf66, primals_22, stride=(1, 1), padding=(0, 0), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf67, (4, 512, 8, 8), (32768, 1, 4096, 512))
        buf68 = empty_strided_cuda((4, 1, 1, 1, 8, 64), (512, 2048, 2048, 2048, 1, 8), torch.float32)
        buf69 = empty_strided_cuda((4, 1, 1, 1, 8, 64), (512, 2048, 2048, 2048, 1, 8), torch.float32)
        buf70 = empty_strided_cuda((4, 1, 1, 1, 8, 64), (512, 2048, 2048, 2048, 1, 8), torch.float32)
        # Topologically Sorted Source Nodes: [output_25], Original ATen: [aten.native_layer_norm]
        triton_per_fused_native_layer_norm_26.run(buf65, buf68, buf69, buf70, 2048, 128, grid=grid(2048), stream=stream0)
        buf71 = empty_strided_cuda((4, 1, 1, 1, 8), (8, 32, 32, 32, 1), torch.float32)
        buf72 = empty_strided_cuda((4, 1, 1, 1, 8), (8, 32, 32, 32, 1), torch.float32)
        buf73 = empty_strided_cuda((4, 1, 1, 1, 8), (8, 32, 32, 32, 1), torch.float32)
        # Topologically Sorted Source Nodes: [output_25], Original ATen: [aten.native_layer_norm]
        triton_per_fused_native_layer_norm_27.run(buf68, buf69, buf70, buf71, buf72, buf73, 32, 64, grid=grid(32), stream=stream0)
        buf74 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf75 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf123 = empty_strided_cuda((4, 1, 1, 1), (1, 1, 1, 1), torch.float32)
        # Topologically Sorted Source Nodes: [output_25], Original ATen: [aten.native_layer_norm, aten.native_layer_norm_backward]
        triton_per_fused_native_layer_norm_native_layer_norm_backward_28.run(buf71, buf72, buf73, buf74, buf75, buf123, 4, 8, grid=grid(4), stream=stream0)
        buf77 = buf65; del buf65  # reuse
        # Topologically Sorted Source Nodes: [output_25], Original ATen: [aten.native_layer_norm]
        triton_poi_fused_native_layer_norm_29.run(buf77, buf74, buf75, 262144, grid=grid(262144), stream=stream0)
        buf78 = empty_strided_cuda((4, 256, 16, 16), (65536, 1, 4096, 256), torch.float32)
        # Topologically Sorted Source Nodes: [output_25, output_26], Original ATen: [aten.native_layer_norm, aten.relu]
        triton_poi_fused_native_layer_norm_relu_30.run(buf77, primals_24, primals_25, buf78, 1024, 256, grid=grid(1024, 256), stream=stream0)
        del primals_25
        # Topologically Sorted Source Nodes: [output_27], Original ATen: [aten.convolution]
        buf79 = extern_kernels.convolution(buf78, buf5, stride=(1, 1), padding=(1, 1), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf79, (4, 256, 16, 16), (65536, 1, 4096, 256))
        buf80 = buf70; del buf70  # reuse
        buf81 = buf69; del buf69  # reuse
        buf82 = buf68; del buf68  # reuse
        # Topologically Sorted Source Nodes: [output_28], Original ATen: [aten.native_layer_norm]
        triton_per_fused_native_layer_norm_26.run(buf79, buf80, buf81, buf82, 2048, 128, grid=grid(2048), stream=stream0)
        buf83 = buf73; del buf73  # reuse
        buf84 = buf72; del buf72  # reuse
        buf85 = buf71; del buf71  # reuse
        # Topologically Sorted Source Nodes: [output_28], Original ATen: [aten.native_layer_norm]
        triton_per_fused_native_layer_norm_27.run(buf80, buf81, buf82, buf83, buf84, buf85, 32, 64, grid=grid(32), stream=stream0)
        del buf80
        del buf81
        del buf82
        buf86 = reinterpret_tensor(buf75, (4, 1, 1, 1), (1, 1, 1, 1), 0); del buf75  # reuse
        buf87 = buf74; del buf74  # reuse
        buf89 = reinterpret_tensor(buf87, (4, 1, 1, 1), (1, 1, 1, 1), 0); del buf87  # reuse
        # Topologically Sorted Source Nodes: [output_28], Original ATen: [aten.native_layer_norm]
        triton_per_fused_native_layer_norm_31.run(buf89, buf83, buf84, buf85, buf86, 4, 8, grid=grid(4), stream=stream0)
        del buf83
        del buf84
        del buf85
        buf90 = empty_strided_cuda((4, 256, 16, 16), (65536, 1, 4096, 256), torch.float32)
        # Topologically Sorted Source Nodes: [output_28, output_29], Original ATen: [aten.native_layer_norm, aten.relu]
        triton_poi_fused_native_layer_norm_relu_32.run(buf79, buf86, buf89, primals_27, primals_28, buf90, 1024, 256, grid=grid(1024, 256), stream=stream0)
        del primals_28
        # Topologically Sorted Source Nodes: [output_30], Original ATen: [aten.convolution]
        buf91 = extern_kernels.convolution(buf90, buf6, stride=(1, 1), padding=(1, 1), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf91, (4, 512, 16, 16), (131072, 1, 8192, 512))
        buf92 = buf67; del buf67  # reuse
        # Topologically Sorted Source Nodes: [output_24, add_17, add_18, add_19, output_31, output_32], Original ATen: [aten.convolution, aten.add, aten.div]
        triton_poi_fused_add_convolution_div_33.run(buf92, primals_23, buf91, primals_30, 131072, grid=grid(131072), stream=stream0)
        del buf91
        del primals_23
        del primals_30
        buf93 = empty_strided_cuda((4, 512, 4, 4), (8192, 1, 2048, 512), torch.float32)
        # Topologically Sorted Source Nodes: [add_21, add_22, add_23, output_33], Original ATen: [aten.add, aten.div]
        triton_poi_fused_add_div_34.run(buf92, buf93, 32768, grid=grid(32768), stream=stream0)
        # Topologically Sorted Source Nodes: [output_34], Original ATen: [aten.convolution]
        buf94 = extern_kernels.convolution(buf93, primals_31, stride=(1, 1), padding=(0, 0), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf94, (4, 512, 4, 4), (8192, 1, 2048, 512))
        buf95 = empty_strided_cuda((4, 1, 1, 1, 4, 64), (256, 1024, 1024, 1024, 64, 1), torch.float32)
        buf96 = empty_strided_cuda((4, 1, 1, 1, 4, 64), (256, 1024, 1024, 1024, 64, 1), torch.float32)
        buf97 = empty_strided_cuda((4, 1, 1, 1, 4, 64), (256, 1024, 1024, 1024, 64, 1), torch.float32)
        # Topologically Sorted Source Nodes: [output_35], Original ATen: [aten.native_layer_norm]
        triton_per_fused_native_layer_norm_35.run(buf92, buf95, buf96, buf97, 1024, 128, grid=grid(1024), stream=stream0)
        buf98 = empty_strided_cuda((4, 1, 1, 1, 4), (4, 16, 16, 16, 1), torch.float32)
        buf99 = empty_strided_cuda((4, 1, 1, 1, 4), (4, 16, 16, 16, 1), torch.float32)
        buf100 = empty_strided_cuda((4, 1, 1, 1, 4), (4, 16, 16, 16, 1), torch.float32)
        # Topologically Sorted Source Nodes: [output_35], Original ATen: [aten.native_layer_norm]
        triton_per_fused_native_layer_norm_36.run(buf95, buf96, buf97, buf98, buf99, buf100, 16, 64, grid=grid(16), stream=stream0)
        buf101 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf102 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf122 = empty_strided_cuda((4, 1, 1, 1), (1, 1, 1, 1), torch.float32)
        # Topologically Sorted Source Nodes: [output_35], Original ATen: [aten.native_layer_norm, aten.native_layer_norm_backward]
        triton_per_fused_native_layer_norm_native_layer_norm_backward_37.run(buf98, buf99, buf100, buf101, buf102, buf122, 4, 4, grid=grid(4), stream=stream0)
        buf104 = buf92; del buf92  # reuse
        # Topologically Sorted Source Nodes: [output_35], Original ATen: [aten.native_layer_norm]
        triton_poi_fused_native_layer_norm_38.run(buf104, buf101, buf102, 131072, grid=grid(131072), stream=stream0)
        buf105 = empty_strided_cuda((4, 512, 8, 8), (32768, 1, 4096, 512), torch.float32)
        # Topologically Sorted Source Nodes: [output_35, output_36], Original ATen: [aten.native_layer_norm, aten.relu]
        triton_poi_fused_native_layer_norm_relu_39.run(buf104, primals_33, primals_34, buf105, 2048, 64, grid=grid(2048, 64), stream=stream0)
        del primals_34
        # Topologically Sorted Source Nodes: [output_37], Original ATen: [aten.convolution]
        buf106 = extern_kernels.convolution(buf105, buf7, stride=(1, 1), padding=(1, 1), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf106, (4, 512, 8, 8), (32768, 1, 4096, 512))
        buf107 = buf97; del buf97  # reuse
        buf108 = buf96; del buf96  # reuse
        buf109 = buf95; del buf95  # reuse
        # Topologically Sorted Source Nodes: [output_38], Original ATen: [aten.native_layer_norm]
        triton_per_fused_native_layer_norm_35.run(buf106, buf107, buf108, buf109, 1024, 128, grid=grid(1024), stream=stream0)
        buf110 = buf99; del buf99  # reuse
        buf111 = buf98; del buf98  # reuse
        buf112 = buf100; del buf100  # reuse
        # Topologically Sorted Source Nodes: [output_38], Original ATen: [aten.native_layer_norm]
        triton_per_fused_native_layer_norm_36.run(buf107, buf108, buf109, buf110, buf111, buf112, 16, 64, grid=grid(16), stream=stream0)
        del buf107
        del buf108
        del buf109
        buf113 = reinterpret_tensor(buf102, (4, 1, 1, 1), (1, 1, 1, 1), 0); del buf102  # reuse
        buf114 = buf101; del buf101  # reuse
        buf116 = reinterpret_tensor(buf114, (4, 1, 1, 1), (1, 1, 1, 1), 0); del buf114  # reuse
        # Topologically Sorted Source Nodes: [output_38], Original ATen: [aten.native_layer_norm]
        triton_per_fused_native_layer_norm_40.run(buf116, buf110, buf111, buf112, buf113, 4, 4, grid=grid(4), stream=stream0)
        del buf110
        del buf111
        del buf112
        buf117 = empty_strided_cuda((4, 512, 8, 8), (32768, 1, 4096, 512), torch.float32)
        # Topologically Sorted Source Nodes: [output_38, output_39], Original ATen: [aten.native_layer_norm, aten.relu]
        triton_poi_fused_native_layer_norm_relu_41.run(buf106, buf113, buf116, primals_36, primals_37, buf117, 2048, 64, grid=grid(2048, 64), stream=stream0)
        del primals_37
        # Topologically Sorted Source Nodes: [output_40], Original ATen: [aten.convolution]
        buf118 = extern_kernels.convolution(buf117, buf8, stride=(1, 1), padding=(1, 1), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf118, (4, 512, 8, 8), (32768, 1, 4096, 512))
        buf119 = empty_strided_cuda((4, 512, 4, 4), (8192, 16, 4, 1), torch.float32)
        # Topologically Sorted Source Nodes: [output_34, add_24, add_25, add_26, output_41, output_42], Original ATen: [aten.convolution, aten.add, aten.div]
        triton_poi_fused_add_convolution_div_42.run(buf94, primals_32, buf118, primals_39, buf119, 64, 512, grid=grid(64, 512), stream=stream0)
        del buf118
        del buf94
        del primals_32
        del primals_39
        buf120 = empty_strided_cuda((4, 128), (128, 1), torch.float32)
        # Topologically Sorted Source Nodes: [], Original ATen: []
        extern_kernels.mm(reinterpret_tensor(buf119, (4, 8192), (8192, 1), 0), reinterpret_tensor(primals_40, (8192, 128), (1, 8192), 0), out=buf120)
        buf121 = buf120; del buf120  # reuse
        # Topologically Sorted Source Nodes: [output_45], Original ATen: [aten.tanh]
        triton_poi_fused_tanh_43.run(buf121, primals_41, 512, grid=grid(512), stream=stream0)
        del primals_41
    return (buf121, buf0, primals_4, primals_6, buf1, primals_9, buf2, primals_13, primals_15, buf3, primals_18, buf4, primals_22, primals_24, buf5, primals_27, buf6, primals_31, primals_33, buf7, primals_36, buf8, buf9, buf11, buf12, buf20, buf23, buf24, buf25, buf32, buf35, buf36, buf39, buf50, buf51, buf52, buf59, buf62, buf63, buf66, buf77, buf78, buf79, buf86, buf89, buf90, buf93, buf104, buf105, buf106, buf113, buf116, buf117, reinterpret_tensor(buf119, (4, 8192), (8192, 1), 0), buf121, primals_40, buf122, buf123, buf124, )


def benchmark_compiled_module(times=10, repeat=10):
    from torch._dynamo.testing import rand_strided
    from torch._inductor.utils import print_performance
    primals_1 = rand_strided((4, 3, 64, 64), (12288, 4096, 64, 1), device='cuda:0', dtype=torch.float32)
    primals_2 = rand_strided((64, 3, 3, 3), (27, 9, 3, 1), device='cuda:0', dtype=torch.float32)
    primals_3 = rand_strided((64, ), (1, ), device='cuda:0', dtype=torch.float32)
    primals_4 = rand_strided((128, 64, 1, 1), (64, 1, 1, 1), device='cuda:0', dtype=torch.float32)
    primals_5 = rand_strided((128, ), (1, ), device='cuda:0', dtype=torch.float32)
    primals_6 = rand_strided((64, 64, 64), (4096, 64, 1), device='cuda:0', dtype=torch.float32)
    primals_7 = rand_strided((64, 64, 64), (4096, 64, 1), device='cuda:0', dtype=torch.float32)
    primals_8 = rand_strided((64, 64, 3, 3), (576, 9, 3, 1), device='cuda:0', dtype=torch.float32)
    primals_9 = rand_strided((64, 64, 64), (4096, 64, 1), device='cuda:0', dtype=torch.float32)
    primals_10 = rand_strided((64, 64, 64), (4096, 64, 1), device='cuda:0', dtype=torch.float32)
    primals_11 = rand_strided((128, 64, 3, 3), (576, 9, 3, 1), device='cuda:0', dtype=torch.float32)
    primals_12 = rand_strided((128, ), (1, ), device='cuda:0', dtype=torch.float32)
    primals_13 = rand_strided((256, 128, 1, 1), (128, 1, 1, 1), device='cuda:0', dtype=torch.float32)
    primals_14 = rand_strided((256, ), (1, ), device='cuda:0', dtype=torch.float32)
    primals_15 = rand_strided((128, 32, 32), (1024, 32, 1), device='cuda:0', dtype=torch.float32)
    primals_16 = rand_strided((128, 32, 32), (1024, 32, 1), device='cuda:0', dtype=torch.float32)
    primals_17 = rand_strided((128, 128, 3, 3), (1152, 9, 3, 1), device='cuda:0', dtype=torch.float32)
    primals_18 = rand_strided((128, 32, 32), (1024, 32, 1), device='cuda:0', dtype=torch.float32)
    primals_19 = rand_strided((128, 32, 32), (1024, 32, 1), device='cuda:0', dtype=torch.float32)
    primals_20 = rand_strided((256, 128, 3, 3), (1152, 9, 3, 1), device='cuda:0', dtype=torch.float32)
    primals_21 = rand_strided((256, ), (1, ), device='cuda:0', dtype=torch.float32)
    primals_22 = rand_strided((512, 256, 1, 1), (256, 1, 1, 1), device='cuda:0', dtype=torch.float32)
    primals_23 = rand_strided((512, ), (1, ), device='cuda:0', dtype=torch.float32)
    primals_24 = rand_strided((256, 16, 16), (256, 16, 1), device='cuda:0', dtype=torch.float32)
    primals_25 = rand_strided((256, 16, 16), (256, 16, 1), device='cuda:0', dtype=torch.float32)
    primals_26 = rand_strided((256, 256, 3, 3), (2304, 9, 3, 1), device='cuda:0', dtype=torch.float32)
    primals_27 = rand_strided((256, 16, 16), (256, 16, 1), device='cuda:0', dtype=torch.float32)
    primals_28 = rand_strided((256, 16, 16), (256, 16, 1), device='cuda:0', dtype=torch.float32)
    primals_29 = rand_strided((512, 256, 3, 3), (2304, 9, 3, 1), device='cuda:0', dtype=torch.float32)
    primals_30 = rand_strided((512, ), (1, ), device='cuda:0', dtype=torch.float32)
    primals_31 = rand_strided((512, 512, 1, 1), (512, 1, 1, 1), device='cuda:0', dtype=torch.float32)
    primals_32 = rand_strided((512, ), (1, ), device='cuda:0', dtype=torch.float32)
    primals_33 = rand_strided((512, 8, 8), (64, 8, 1), device='cuda:0', dtype=torch.float32)
    primals_34 = rand_strided((512, 8, 8), (64, 8, 1), device='cuda:0', dtype=torch.float32)
    primals_35 = rand_strided((512, 512, 3, 3), (4608, 9, 3, 1), device='cuda:0', dtype=torch.float32)
    primals_36 = rand_strided((512, 8, 8), (64, 8, 1), device='cuda:0', dtype=torch.float32)
    primals_37 = rand_strided((512, 8, 8), (64, 8, 1), device='cuda:0', dtype=torch.float32)
    primals_38 = rand_strided((512, 512, 3, 3), (4608, 9, 3, 1), device='cuda:0', dtype=torch.float32)
    primals_39 = rand_strided((512, ), (1, ), device='cuda:0', dtype=torch.float32)
    primals_40 = rand_strided((128, 8192), (8192, 1), device='cuda:0', dtype=torch.float32)
    primals_41 = rand_strided((128, ), (1, ), device='cuda:0', dtype=torch.float32)
    fn = lambda: 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, primals_32, primals_33, primals_34, primals_35, primals_36, primals_37, primals_38, primals_39, primals_40, primals_41])
    return print_performance(fn, times=times, repeat=repeat)


if __name__ == "__main__":
    from torch._inductor.wrapper_benchmark import compiled_module_main
    compiled_module_main('None', benchmark_compiled_module)

import torch
from torch import nn


class IWConv2d(nn.Module):

    def __init__(self, input_dim, output_dim, kernel_size, he_init=True,
        stride=1, bias=True):
        super(IWConv2d, self).__init__()
        self.he_init = he_init
        self.padding = int((kernel_size - 1) / 2)
        self.conv = nn.Conv2d(input_dim, output_dim, kernel_size, stride=1,
            padding=self.padding, bias=bias)

    def forward(self, input):
        output = self.conv(input)
        return output


class ConvMeanPool(nn.Module):

    def __init__(self, input_dim, output_dim, kernel_size, he_init=True):
        super(ConvMeanPool, self).__init__()
        self.he_init = he_init
        self.conv = IWConv2d(input_dim, output_dim, kernel_size, he_init=
            self.he_init)

    def forward(self, input):
        output = self.conv(input)
        output = (output[:, :, ::2, ::2] + output[:, :, 1::2, ::2] + output
            [:, :, ::2, 1::2] + output[:, :, 1::2, 1::2]) / 4
        return output


class MeanPoolConv(nn.Module):

    def __init__(self, input_dim, output_dim, kernel_size, he_init=True):
        super(MeanPoolConv, self).__init__()
        self.he_init = he_init
        self.conv = IWConv2d(input_dim, output_dim, kernel_size, he_init=
            self.he_init)

    def forward(self, input):
        output = input
        output = (output[:, :, ::2, ::2] + output[:, :, 1::2, ::2] + output
            [:, :, ::2, 1::2] + output[:, :, 1::2, 1::2]) / 4
        output = self.conv(output)
        return output


class DepthToSpace(nn.Module):

    def __init__(self, block_size):
        super(DepthToSpace, self).__init__()
        self.block_size = block_size
        self.block_size_sq = block_size * block_size

    def forward(self, input):
        output = input.permute(0, 2, 3, 1)
        batch_size, input_height, input_width, input_depth = output.size()
        output_depth = int(input_depth / self.block_size_sq)
        output_width = int(input_width * self.block_size)
        output_height = int(input_height * self.block_size)
        t_1 = output.reshape(batch_size, input_height, input_width, self.
            block_size_sq, output_depth)
        spl = t_1.split(self.block_size, 3)
        stacks = [t_t.reshape(batch_size, input_height, output_width,
            output_depth) for t_t in spl]
        output = torch.stack(stacks, 0).transpose(0, 1).permute(0, 2, 1, 3, 4
            ).reshape(batch_size, output_height, output_width, output_depth)
        output = output.permute(0, 3, 1, 2)
        return output


class UpSampleConv(nn.Module):

    def __init__(self, input_dim, output_dim, kernel_size, he_init=True,
        bias=True):
        super(UpSampleConv, self).__init__()
        self.he_init = he_init
        self.conv = IWConv2d(input_dim, output_dim, kernel_size, he_init=
            self.he_init, bias=bias)
        self.depth_to_space = DepthToSpace(2)

    def forward(self, input):
        output = input
        output = torch.cat((output, output, output, output), 1)
        output = self.depth_to_space(output)
        output = self.conv(output)
        return output


class ResidualBlock(nn.Module):

    def __init__(self, input_dim, output_dim, kernel_size, resample=None, hw=64
        ):
        super(ResidualBlock, self).__init__()
        self.input_dim = input_dim
        self.output_dim = output_dim
        self.kernel_size = kernel_size
        self.resample = resample
        self.bn1 = None
        self.bn2 = None
        self.relu1 = nn.ReLU()
        self.relu2 = nn.ReLU()
        if resample == 'down':
            self.bn1 = nn.LayerNorm([input_dim, hw, hw])
            self.bn2 = nn.LayerNorm([input_dim, hw, hw])
        elif resample == 'up':
            self.bn1 = nn.BatchNorm2d(input_dim)
            self.bn2 = nn.BatchNorm2d(output_dim)
        elif resample is None:
            self.bn1 = nn.BatchNorm2d(output_dim)
            self.bn2 = nn.LayerNorm([input_dim, hw, hw])
        else:
            raise Exception('invalid resample value')
        if resample == 'down':
            self.conv_shortcut = MeanPoolConv(input_dim, output_dim,
                kernel_size=1, he_init=False)
            self.conv_1 = IWConv2d(input_dim, input_dim, kernel_size=
                kernel_size, bias=False)
            self.conv_2 = ConvMeanPool(input_dim, output_dim, kernel_size=
                kernel_size)
        elif resample == 'up':
            self.conv_shortcut = UpSampleConv(input_dim, output_dim,
                kernel_size=1, he_init=False)
            self.conv_1 = UpSampleConv(input_dim, output_dim, kernel_size=
                kernel_size, bias=False)
            self.conv_2 = IWConv2d(output_dim, output_dim, kernel_size=
                kernel_size)
        elif resample is None:
            self.conv_shortcut = IWConv2d(input_dim, output_dim,
                kernel_size=1, he_init=False)
            self.conv_1 = IWConv2d(input_dim, input_dim, kernel_size=
                kernel_size, bias=False)
            self.conv_2 = IWConv2d(input_dim, output_dim, kernel_size=
                kernel_size)
        else:
            raise Exception('invalid resample value')

    def forward(self, input):
        if self.input_dim == self.output_dim and self.resample is None:
            shortcut = input
        else:
            shortcut = self.conv_shortcut(input)
        output = input
        output = self.bn1(output)
        output = self.relu1(output)
        output = self.conv_1(output)
        output = self.bn2(output)
        output = self.relu2(output)
        output = self.conv_2(output)
        return shortcut + output


class IWEncoder(nn.Module):

    def __init__(self, input_size=64, z_dim=128, n_image_channels=3):
        super(IWEncoder, self).__init__()
        self.size = input_size
        self.n_image_channels = n_image_channels
        self.ssize = self.size // 16
        self.conv1 = IWConv2d(n_image_channels, self.size, 3, he_init=False)
        self.rb1 = ResidualBlock(self.size, 2 * self.size, 3, resample=
            'down', hw=self.size)
        self.rb2 = ResidualBlock(2 * self.size, 4 * self.size, 3, resample=
            'down', hw=int(self.size / 2))
        self.rb3 = ResidualBlock(4 * self.size, 8 * self.size, 3, resample=
            'down', hw=int(self.size / 4))
        self.rb4 = ResidualBlock(8 * self.size, 8 * self.size, 3, resample=
            'down', hw=int(self.size / 8))
        self.ln1 = nn.Linear(self.ssize * self.ssize * 8 * self.size, z_dim)

    def forward(self, input):
        output = input.contiguous()
        output = output.view(-1, self.n_image_channels, self.size, self.size)
        output = self.conv1(output)
        output = self.rb1(output)
        output = self.rb2(output)
        output = self.rb3(output)
        output = self.rb4(output)
        output = output.view(-1, self.ssize * self.ssize * 8 * self.size)
        output = self.ln1(output)
        output = torch.tanh(output)
        return output


def get_inputs():
    return [torch.rand([4, 3, 64, 64])]


def get_init_inputs():
    return [[], {}]

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
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_poi_fused_0(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK: tl.
    constexpr, XBLOCK: tl.constexpr):
    ynumel = 192
    xnumel = 9
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = yindex < ynumel
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y3 = yindex
    y0 = yindex % 3
    y1 = yindex // 3
    tmp0 = tl.load(in_ptr0 + (x2 + 9 * y3), xmask & ymask, eviction_policy=
        'evict_last')
    tl.store(out_ptr0 + (y0 + 3 * x2 + 27 * y1), tmp0, xmask & ymask)


@triton.jit
def triton_poi_fused_1(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK: tl.
    constexpr, XBLOCK: tl.constexpr):
    xnumel = 9
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y3 = yindex
    y0 = yindex % 64
    y1 = yindex // 64
    tmp0 = tl.load(in_ptr0 + (x2 + 9 * y3), xmask, eviction_policy='evict_last'
        )
    tl.store(out_ptr0 + (y0 + 64 * x2 + 576 * y1), tmp0, xmask)


@triton.jit
def triton_poi_fused_2(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK: tl.
    constexpr, XBLOCK: tl.constexpr):
    xnumel = 9
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y3 = yindex
    y0 = yindex % 64
    y1 = yindex // 64
    tmp0 = tl.load(in_ptr0 + (x2 + 9 * y3), xmask, eviction_policy='evict_last'
        )
    tl.store(out_ptr0 + (y0 + 64 * x2 + 576 * y1), tmp0, xmask)


@triton.jit
def triton_poi_fused_3(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK: tl.
    constexpr, XBLOCK: tl.constexpr):
    xnumel = 9
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y3 = yindex
    y0 = yindex % 128
    y1 = yindex // 128
    tmp0 = tl.load(in_ptr0 + (x2 + 9 * y3), xmask, eviction_policy='evict_last'
        )
    tl.store(out_ptr0 + (y0 + 128 * x2 + 1152 * y1), tmp0, xmask)


@triton.jit
def triton_poi_fused_4(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK: tl.
    constexpr, XBLOCK: tl.constexpr):
    xnumel = 9
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y3 = yindex
    y0 = yindex % 128
    y1 = yindex // 128
    tmp0 = tl.load(in_ptr0 + (x2 + 9 * y3), xmask, eviction_policy='evict_last'
        )
    tl.store(out_ptr0 + (y0 + 128 * x2 + 1152 * y1), tmp0, xmask)


@triton.jit
def triton_poi_fused_5(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK: tl.
    constexpr, XBLOCK: tl.constexpr):
    xnumel = 9
    yoffset = (tl.program_id(1) + tl.program_id(2) * tl.num_programs(1)
        ) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y3 = yindex
    y0 = yindex % 256
    y1 = yindex // 256
    tmp0 = tl.load(in_ptr0 + (x2 + 9 * y3), xmask, eviction_policy='evict_last'
        )
    tl.store(out_ptr0 + (y0 + 256 * x2 + 2304 * y1), tmp0, xmask)


@triton.jit
def triton_poi_fused_6(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK: tl.
    constexpr, XBLOCK: tl.constexpr):
    xnumel = 9
    yoffset = (tl.program_id(1) + tl.program_id(2) * tl.num_programs(1)
        ) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y3 = yindex
    y0 = yindex % 256
    y1 = yindex // 256
    tmp0 = tl.load(in_ptr0 + (x2 + 9 * y3), xmask, eviction_policy='evict_last'
        )
    tl.store(out_ptr0 + (y0 + 256 * x2 + 2304 * y1), tmp0, xmask)


@triton.jit
def triton_poi_fused_7(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK: tl.
    constexpr, XBLOCK: tl.constexpr):
    xnumel = 9
    yoffset = (tl.program_id(1) + tl.program_id(2) * tl.num_programs(1)
        ) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y3 = yindex
    y0 = yindex % 512
    y1 = yindex // 512
    tmp0 = tl.load(in_ptr0 + (x2 + 9 * y3), xmask, eviction_policy='evict_last'
        )
    tl.store(out_ptr0 + (y0 + 512 * x2 + 4608 * y1), tmp0, xmask)


@triton.jit
def triton_poi_fused_view_8(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK: tl.
    constexpr, XBLOCK: tl.constexpr):
    ynumel = 12
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = yindex < ynumel
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    tl.full([XBLOCK, YBLOCK], True, tl.int1)
    x2 = xindex
    y3 = yindex
    y0 = yindex % 3
    y1 = yindex // 3
    tmp0 = tl.load(in_ptr0 + (x2 + 4096 * y3), ymask, eviction_policy=
        'evict_last')
    tl.store(out_ptr0 + (y0 + 3 * x2 + 12288 * y1), tmp0, ymask)


@triton.jit
def triton_poi_fused_convolution_9(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)
    x2 = xindex
    x0 = xindex % 64
    tmp0 = tl.load(in_out_ptr0 + x2, None)
    tmp1 = tl.load(in_ptr0 + x0, None, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tl.store(in_out_ptr0 + x2, tmp2, None)


@triton.jit
def triton_poi_fused_add_div_10(in_ptr0, out_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 % 64
    x1 = xindex // 64 % 32
    x2 = xindex // 2048
    x3 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + 128 * x1 + 8192 * x2), None)
    tmp1 = tl.load(in_ptr0 + (4096 + x0 + 128 * x1 + 8192 * x2), None)
    tmp3 = tl.load(in_ptr0 + (64 + x0 + 128 * x1 + 8192 * x2), None)
    tmp5 = tl.load(in_ptr0 + (4160 + x0 + 128 * x1 + 8192 * x2), None)
    tmp2 = tmp0 + tmp1
    tmp4 = tmp2 + tmp3
    tmp6 = tmp4 + tmp5
    tmp7 = 0.25
    tmp8 = tmp6 * tmp7
    tl.store(out_ptr0 + x3, tmp8, None)


@triton.jit
def triton_per_fused_native_layer_norm_11(in_ptr0, out_ptr0, out_ptr1,
    out_ptr2, xnumel, rnumel, XBLOCK: tl.constexpr):
    RBLOCK: tl.constexpr = 128
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    tl.full([XBLOCK, RBLOCK], True, tl.int1)
    rindex = tl.arange(0, RBLOCK)[None, :]
    tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r3 = rindex
    x0 = xindex % 32
    x1 = xindex // 32 % 64
    x2 = xindex // 2048
    x4 = xindex
    tmp0 = tl.load(in_ptr0 + (2 * x0 + 64 * (r3 % 64) + 4096 * ((r3 + 128 *
        x1) // 64 % 64) + 262144 * x2 + (r3 + 128 * x1) // 4096), None,
        eviction_policy='evict_last')
    tmp1 = tl.broadcast_to(tmp0, [XBLOCK, RBLOCK])
    tmp3 = tl.broadcast_to(tmp1, [XBLOCK, RBLOCK])
    tmp5 = tl.sum(tmp3, 1)[:, None]
    tmp6 = tl.full([XBLOCK, 1], 128, tl.int32)
    tmp7 = tmp6.to(tl.float32)
    tmp8 = tmp5 / tmp7
    tmp9 = tmp1 - tmp8
    tmp10 = tmp9 * tmp9
    tmp11 = tl.broadcast_to(tmp10, [XBLOCK, RBLOCK])
    tmp13 = tl.sum(tmp11, 1)[:, None]
    tl.store(out_ptr0 + x4, tmp8, None)
    tl.store(out_ptr1 + x4, tmp13, None)
    tl.store(out_ptr2 + x4, tmp7, None)


@triton.jit
def triton_per_fused_native_layer_norm_12(in_ptr0, in_ptr1, in_ptr2,
    out_ptr0, out_ptr1, out_ptr2, xnumel, rnumel, XBLOCK: tl.constexpr):
    xnumel = 128
    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)
    r2 = rindex
    x0 = xindex % 32
    x1 = xindex // 32
    x3 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + 32 * r2 + 2048 * x1), xmask, other=0.0)
    tmp1 = tl.load(in_ptr1 + (x0 + 32 * r2 + 2048 * x1), xmask, other=0.0)
    tmp2 = tl.load(in_ptr2 + (x0 + 32 * r2 + 2048 * x1), 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]
    tmp15 = tmp12[:, None]
    tl.store(out_ptr0 + x3, tmp13, xmask)
    tl.store(out_ptr1 + x3, tmp14, xmask)
    tl.store(out_ptr2 + x3, tmp15, xmask)


@triton.jit
def triton_per_fused_native_layer_norm_13(in_out_ptr0, in_ptr0, in_ptr1,
    in_ptr2, out_ptr0, xnumel, rnumel, XBLOCK: tl.constexpr):
    xnumel = 4
    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, :]
    tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r1 = rindex
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + (r1 + 32 * x0), xmask, other=0.0)
    tmp1 = tl.load(in_ptr1 + (r1 + 32 * x0), xmask, other=0.0)
    tmp2 = tl.load(in_ptr2 + (r1 + 32 * 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 = 262144.0
    tmp17 = tmp14 / tmp16
    tmp18 = 1e-05
    tmp19 = tmp17 + tmp18
    tmp20 = libdevice.rsqrt(tmp19)
    tl.debug_barrier()
    tl.store(in_out_ptr0 + x0, tmp20, xmask)
    tl.store(out_ptr0 + x0, tmp13, xmask)


@triton.jit
def triton_poi_fused_native_layer_norm_relu_14(in_ptr0, in_ptr1, in_ptr2,
    in_ptr3, in_ptr4, out_ptr0, ynumel, xnumel, YBLOCK: tl.constexpr,
    XBLOCK: tl.constexpr):
    ynumel = 256
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = yindex < ynumel
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    tl.full([XBLOCK, YBLOCK], True, tl.int1)
    x2 = xindex
    y0 = yindex % 64
    y1 = yindex // 64
    tmp0 = tl.load(in_ptr0 + (y0 + 64 * x2 + 262144 * y1), ymask,
        eviction_policy='evict_last')
    tmp1 = tl.load(in_ptr1 + y1, ymask, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr2 + y1, ymask, eviction_policy='evict_last')
    tmp5 = tl.load(in_ptr3 + (x2 + 4096 * y0), ymask, eviction_policy=
        'evict_last')
    tmp7 = tl.load(in_ptr4 + (x2 + 4096 * y0), ymask, eviction_policy=
        'evict_last')
    tmp2 = tmp0 - tmp1
    tmp4 = tmp2 * tmp3
    tmp6 = tmp4 * tmp5
    tmp8 = tmp6 + tmp7
    tmp9 = tl.full([1, 1], 0, tl.int32)
    tmp10 = triton_helpers.maximum(tmp9, tmp8)
    tl.store(out_ptr0 + (y0 + 64 * x2 + 262144 * y1), tmp10, ymask)


@triton.jit
def triton_poi_fused_add_convolution_div_15(in_out_ptr0, in_ptr0, in_ptr1,
    in_ptr2, xnumel, XBLOCK: tl.constexpr):
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    tl.full([XBLOCK], True, tl.int1)
    x3 = xindex
    x0 = xindex % 128
    x1 = xindex // 128 % 32
    x2 = xindex // 4096
    tmp0 = tl.load(in_out_ptr0 + x3, None)
    tmp1 = tl.load(in_ptr0 + x0, None, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr1 + (x0 + 256 * x1 + 16384 * x2), None)
    tmp4 = tl.load(in_ptr2 + x0, None, eviction_policy='evict_last')
    tmp6 = tl.load(in_ptr1 + (8192 + x0 + 256 * x1 + 16384 * x2), None)
    tmp9 = tl.load(in_ptr1 + (128 + x0 + 256 * x1 + 16384 * x2), None)
    tmp12 = tl.load(in_ptr1 + (8320 + x0 + 256 * x1 + 16384 * x2), None)
    tmp2 = tmp0 + tmp1
    tmp5 = tmp3 + tmp4
    tmp7 = tmp6 + tmp4
    tmp8 = tmp5 + tmp7
    tmp10 = tmp9 + tmp4
    tmp11 = tmp8 + tmp10
    tmp13 = tmp12 + tmp4
    tmp14 = tmp11 + tmp13
    tmp15 = 0.25
    tmp16 = tmp14 * tmp15
    tmp17 = tmp2 + tmp16
    tl.store(in_out_ptr0 + x3, tmp17, None)


@triton.jit
def triton_poi_fused_add_div_16(in_ptr0, out_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 % 128
    x1 = xindex // 128 % 16
    x2 = xindex // 2048
    x3 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + 256 * x1 + 8192 * x2), None)
    tmp1 = tl.load(in_ptr0 + (4096 + x0 + 256 * x1 + 8192 * x2), None)
    tmp3 = tl.load(in_ptr0 + (128 + x0 + 256 * x1 + 8192 * x2), None)
    tmp5 = tl.load(in_ptr0 + (4224 + x0 + 256 * x1 + 8192 * x2), None)
    tmp2 = tmp0 + tmp1
    tmp4 = tmp2 + tmp3
    tmp6 = tmp4 + tmp5
    tmp7 = 0.25
    tmp8 = tmp6 * tmp7
    tl.store(out_ptr0 + x3, tmp8, None)


@triton.jit
def triton_per_fused_native_layer_norm_17(in_ptr0, out_ptr0, out_ptr1,
    out_ptr2, xnumel, rnumel, XBLOCK: tl.constexpr):
    RBLOCK: tl.constexpr = 128
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    tl.full([XBLOCK, RBLOCK], True, tl.int1)
    rindex = tl.arange(0, RBLOCK)[None, :]
    tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r3 = rindex
    x0 = xindex % 16
    x1 = xindex // 16 % 64
    x2 = xindex // 1024
    x4 = xindex
    tmp0 = tl.load(in_ptr0 + (8 * x0 + 128 * (r3 % 32) + 4096 * ((r3 + 128 *
        x1) // 32 % 32) + 131072 * x2 + (r3 + 128 * x1) // 1024), None,
        eviction_policy='evict_last')
    tmp1 = tl.broadcast_to(tmp0, [XBLOCK, RBLOCK])
    tmp3 = tl.broadcast_to(tmp1, [XBLOCK, RBLOCK])
    tmp5 = tl.sum(tmp3, 1)[:, None]
    tmp6 = tl.full([XBLOCK, 1], 128, tl.int32)
    tmp7 = tmp6.to(tl.float32)
    tmp8 = tmp5 / tmp7
    tmp9 = tmp1 - tmp8
    tmp10 = tmp9 * tmp9
    tmp11 = tl.broadcast_to(tmp10, [XBLOCK, RBLOCK])
    tmp13 = tl.sum(tmp11, 1)[:, None]
    tl.store(out_ptr0 + x4, tmp8, None)
    tl.store(out_ptr1 + x4, tmp13, None)
    tl.store(out_ptr2 + x4, tmp7, None)


@triton.jit
def triton_per_fused_native_layer_norm_18(in_ptr0, in_ptr1, in_ptr2,
    out_ptr0, out_ptr1, out_ptr2, xnumel, rnumel, XBLOCK: tl.constexpr):
    xnumel = 64
    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)
    r2 = rindex
    x0 = xindex % 16
    x1 = xindex // 16
    x3 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + 16 * r2 + 1024 * x1), xmask, other=0.0)
    tmp1 = tl.load(in_ptr1 + (x0 + 16 * r2 + 1024 * x1), xmask, other=0.0)
    tmp2 = tl.load(in_ptr2 + (x0 + 16 * r2 + 1024 * x1), 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]
    tmp15 = tmp12[:, None]
    tl.store(out_ptr0 + x3, tmp13, xmask)
    tl.store(out_ptr1 + x3, tmp14, xmask)
    tl.store(out_ptr2 + x3, tmp15, xmask)


@triton.jit
def triton_per_fused_native_layer_norm_native_layer_norm_backward_19(in_ptr0,
    in_ptr1, in_ptr2, out_ptr0, out_ptr1, out_ptr2, xnumel, rnumel, XBLOCK:
    tl.constexpr):
    xnumel = 4
    RBLOCK: tl.constexpr = 16
    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 + 16 * x0), xmask, other=0.0)
    tmp1 = tl.load(in_ptr1 + (r1 + 16 * x0), xmask, other=0.0)
    tmp2 = tl.load(in_ptr2 + (r1 + 16 * 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 = 131072.0
    tmp17 = tmp14 / tmp16
    tmp18 = 1e-05
    tmp19 = tmp17 + tmp18
    tmp20 = libdevice.rsqrt(tmp19)
    tmp21 = 7.62939453125e-06
    tmp22 = tmp20 * tmp21
    tl.store(out_ptr2 + x0, tmp22, xmask)
    tl.store(out_ptr0 + x0, tmp13, xmask)
    tl.store(out_ptr1 + x0, tmp14, xmask)


@triton.jit
def triton_poi_fused_native_layer_norm_20(in_out_ptr0, in_ptr0, in_ptr1,
    xnumel, XBLOCK: tl.constexpr):
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    tl.full([XBLOCK], True, tl.int1)
    x2 = xindex
    x1 = xindex // 131072
    tmp0 = tl.load(in_out_ptr0 + x2, None)
    tmp1 = tl.load(in_ptr0 + x1, None, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr1 + x1, None, eviction_policy='evict_last')
    tmp2 = tmp0 - tmp1
    tmp4 = 131072.0
    tmp5 = tmp3 / tmp4
    tmp6 = 1e-05
    tmp7 = tmp5 + tmp6
    tmp8 = libdevice.rsqrt(tmp7)
    tmp9 = tmp2 * tmp8
    tl.store(in_out_ptr0 + x2, tmp9, None)


@triton.jit
def triton_poi_fused_native_layer_norm_relu_21(in_ptr0, in_ptr1, in_ptr2,
    out_ptr0, ynumel, xnumel, YBLOCK: tl.constexpr, XBLOCK: tl.constexpr):
    ynumel = 512
    xnumel = 1024
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = yindex < ynumel
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y0 = yindex % 128
    y1 = yindex // 128
    tmp0 = tl.load(in_ptr0 + (y0 + 128 * x2 + 131072 * y1), xmask & ymask,
        eviction_policy='evict_last')
    tmp1 = tl.load(in_ptr1 + (x2 + 1024 * y0), xmask & ymask,
        eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr2 + (x2 + 1024 * y0), xmask & ymask,
        eviction_policy='evict_last')
    tmp2 = tmp0 * tmp1
    tmp4 = tmp2 + tmp3
    tmp5 = tl.full([1, 1], 0, tl.int32)
    tmp6 = triton_helpers.maximum(tmp5, tmp4)
    tl.store(out_ptr0 + (y0 + 128 * x2 + 131072 * y1), tmp6, xmask & ymask)


@triton.jit
def triton_per_fused_native_layer_norm_22(in_out_ptr0, in_ptr0, in_ptr1,
    in_ptr2, out_ptr0, xnumel, rnumel, XBLOCK: tl.constexpr):
    xnumel = 4
    RBLOCK: tl.constexpr = 16
    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 + 16 * x0), xmask, other=0.0)
    tmp1 = tl.load(in_ptr1 + (r1 + 16 * x0), xmask, other=0.0)
    tmp2 = tl.load(in_ptr2 + (r1 + 16 * 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 = 131072.0
    tmp17 = tmp14 / tmp16
    tmp18 = 1e-05
    tmp19 = tmp17 + tmp18
    tmp20 = libdevice.rsqrt(tmp19)
    tl.debug_barrier()
    tl.store(in_out_ptr0 + x0, tmp20, xmask)
    tl.store(out_ptr0 + x0, tmp13, xmask)


@triton.jit
def triton_poi_fused_native_layer_norm_relu_23(in_ptr0, in_ptr1, in_ptr2,
    in_ptr3, in_ptr4, out_ptr0, ynumel, xnumel, YBLOCK: tl.constexpr,
    XBLOCK: tl.constexpr):
    ynumel = 512
    xnumel = 1024
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = yindex < ynumel
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y0 = yindex % 128
    y1 = yindex // 128
    tmp0 = tl.load(in_ptr0 + (y0 + 128 * x2 + 131072 * y1), xmask & ymask,
        eviction_policy='evict_last')
    tmp1 = tl.load(in_ptr1 + y1, ymask, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr2 + y1, ymask, eviction_policy='evict_last')
    tmp5 = tl.load(in_ptr3 + (x2 + 1024 * y0), xmask & ymask,
        eviction_policy='evict_last')
    tmp7 = tl.load(in_ptr4 + (x2 + 1024 * y0), xmask & ymask,
        eviction_policy='evict_last')
    tmp2 = tmp0 - tmp1
    tmp4 = tmp2 * tmp3
    tmp6 = tmp4 * tmp5
    tmp8 = tmp6 + tmp7
    tmp9 = tl.full([1, 1], 0, tl.int32)
    tmp10 = triton_helpers.maximum(tmp9, tmp8)
    tl.store(out_ptr0 + (y0 + 128 * x2 + 131072 * y1), tmp10, xmask & ymask)


@triton.jit
def triton_poi_fused_add_convolution_div_24(in_out_ptr0, in_ptr0, in_ptr1,
    in_ptr2, xnumel, XBLOCK: tl.constexpr):
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    tl.full([XBLOCK], True, tl.int1)
    x3 = xindex
    x0 = xindex % 256
    x1 = xindex // 256 % 16
    x2 = xindex // 4096
    tmp0 = tl.load(in_out_ptr0 + x3, None)
    tmp1 = tl.load(in_ptr0 + x0, None, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr1 + (x0 + 512 * x1 + 16384 * x2), None)
    tmp4 = tl.load(in_ptr2 + x0, None, eviction_policy='evict_last')
    tmp6 = tl.load(in_ptr1 + (8192 + x0 + 512 * x1 + 16384 * x2), None)
    tmp9 = tl.load(in_ptr1 + (256 + x0 + 512 * x1 + 16384 * x2), None)
    tmp12 = tl.load(in_ptr1 + (8448 + x0 + 512 * x1 + 16384 * x2), None)
    tmp2 = tmp0 + tmp1
    tmp5 = tmp3 + tmp4
    tmp7 = tmp6 + tmp4
    tmp8 = tmp5 + tmp7
    tmp10 = tmp9 + tmp4
    tmp11 = tmp8 + tmp10
    tmp13 = tmp12 + tmp4
    tmp14 = tmp11 + tmp13
    tmp15 = 0.25
    tmp16 = tmp14 * tmp15
    tmp17 = tmp2 + tmp16
    tl.store(in_out_ptr0 + x3, tmp17, None)


@triton.jit
def triton_poi_fused_add_div_25(in_ptr0, out_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 % 256
    x1 = xindex // 256 % 8
    x2 = xindex // 2048
    x3 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + 512 * x1 + 8192 * x2), None)
    tmp1 = tl.load(in_ptr0 + (4096 + x0 + 512 * x1 + 8192 * x2), None)
    tmp3 = tl.load(in_ptr0 + (256 + x0 + 512 * x1 + 8192 * x2), None)
    tmp5 = tl.load(in_ptr0 + (4352 + x0 + 512 * x1 + 8192 * x2), None)
    tmp2 = tmp0 + tmp1
    tmp4 = tmp2 + tmp3
    tmp6 = tmp4 + tmp5
    tmp7 = 0.25
    tmp8 = tmp6 * tmp7
    tl.store(out_ptr0 + x3, tmp8, None)


@triton.jit
def triton_per_fused_native_layer_norm_26(in_ptr0, out_ptr0, out_ptr1,
    out_ptr2, xnumel, rnumel, XBLOCK: tl.constexpr):
    RBLOCK: tl.constexpr = 128
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    tl.full([XBLOCK, RBLOCK], True, tl.int1)
    rindex = tl.arange(0, RBLOCK)[None, :]
    tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r3 = rindex
    x0 = xindex % 8
    x1 = xindex // 8 % 64
    x2 = xindex // 512
    x4 = xindex
    tmp0 = tl.load(in_ptr0 + (32 * x0 + 256 * (r3 % 16) + 4096 * ((r3 + 128 *
        x1) // 16 % 16) + 65536 * x2 + (r3 + 128 * x1) // 256), None,
        eviction_policy='evict_last')
    tmp1 = tl.broadcast_to(tmp0, [XBLOCK, RBLOCK])
    tmp3 = tl.broadcast_to(tmp1, [XBLOCK, RBLOCK])
    tmp5 = tl.sum(tmp3, 1)[:, None]
    tmp6 = tl.full([XBLOCK, 1], 128, tl.int32)
    tmp7 = tmp6.to(tl.float32)
    tmp8 = tmp5 / tmp7
    tmp9 = tmp1 - tmp8
    tmp10 = tmp9 * tmp9
    tmp11 = tl.broadcast_to(tmp10, [XBLOCK, RBLOCK])
    tmp13 = tl.sum(tmp11, 1)[:, None]
    tl.store(out_ptr0 + x4, tmp8, None)
    tl.store(out_ptr1 + x4, tmp13, None)
    tl.store(out_ptr2 + x4, tmp7, None)


@triton.jit
def triton_per_fused_native_layer_norm_27(in_ptr0, in_ptr1, in_ptr2,
    out_ptr0, out_ptr1, out_ptr2, xnumel, rnumel, XBLOCK: tl.constexpr):
    xnumel = 32
    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)
    r2 = rindex
    x0 = xindex % 8
    x1 = xindex // 8
    x3 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + 8 * r2 + 512 * x1), xmask, other=0.0)
    tmp1 = tl.load(in_ptr1 + (x0 + 8 * r2 + 512 * x1), xmask, other=0.0)
    tmp2 = tl.load(in_ptr2 + (x0 + 8 * r2 + 512 * x1), 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]
    tmp15 = tmp12[:, None]
    tl.store(out_ptr0 + x3, tmp13, xmask)
    tl.store(out_ptr1 + x3, tmp14, xmask)
    tl.store(out_ptr2 + x3, tmp15, xmask)


@triton.jit
def triton_per_fused_native_layer_norm_native_layer_norm_backward_28(in_ptr0,
    in_ptr1, in_ptr2, out_ptr0, out_ptr1, out_ptr2, xnumel, rnumel, XBLOCK:
    tl.constexpr):
    xnumel = 4
    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)
    tmp1 = tl.load(in_ptr1 + (r1 + 8 * x0), xmask, other=0.0)
    tmp2 = tl.load(in_ptr2 + (r1 + 8 * 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 = 65536.0
    tmp17 = tmp14 / tmp16
    tmp18 = 1e-05
    tmp19 = tmp17 + tmp18
    tmp20 = libdevice.rsqrt(tmp19)
    tmp21 = 1.52587890625e-05
    tmp22 = tmp20 * tmp21
    tl.store(out_ptr2 + x0, tmp22, xmask)
    tl.store(out_ptr0 + x0, tmp13, xmask)
    tl.store(out_ptr1 + x0, tmp14, xmask)


@triton.jit
def triton_poi_fused_native_layer_norm_29(in_out_ptr0, in_ptr0, in_ptr1,
    xnumel, XBLOCK: tl.constexpr):
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    tl.full([XBLOCK], True, tl.int1)
    x2 = xindex
    x1 = xindex // 65536
    tmp0 = tl.load(in_out_ptr0 + x2, None)
    tmp1 = tl.load(in_ptr0 + x1, None, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr1 + x1, None, eviction_policy='evict_last')
    tmp2 = tmp0 - tmp1
    tmp4 = 65536.0
    tmp5 = tmp3 / tmp4
    tmp6 = 1e-05
    tmp7 = tmp5 + tmp6
    tmp8 = libdevice.rsqrt(tmp7)
    tmp9 = tmp2 * tmp8
    tl.store(in_out_ptr0 + x2, tmp9, None)


@triton.jit
def triton_poi_fused_native_layer_norm_relu_30(in_ptr0, in_ptr1, in_ptr2,
    out_ptr0, ynumel, xnumel, YBLOCK: tl.constexpr, XBLOCK: tl.constexpr):
    xnumel = 256
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y0 = yindex % 256
    y1 = yindex // 256
    tmp0 = tl.load(in_ptr0 + (y0 + 256 * x2 + 65536 * y1), xmask,
        eviction_policy='evict_last')
    tmp1 = tl.load(in_ptr1 + (x2 + 256 * y0), xmask, eviction_policy=
        'evict_last')
    tmp3 = tl.load(in_ptr2 + (x2 + 256 * y0), xmask, eviction_policy=
        'evict_last')
    tmp2 = tmp0 * tmp1
    tmp4 = tmp2 + tmp3
    tmp5 = tl.full([1, 1], 0, tl.int32)
    tmp6 = triton_helpers.maximum(tmp5, tmp4)
    tl.store(out_ptr0 + (y0 + 256 * x2 + 65536 * y1), tmp6, xmask)


@triton.jit
def triton_per_fused_native_layer_norm_31(in_out_ptr0, in_ptr0, in_ptr1,
    in_ptr2, out_ptr0, xnumel, rnumel, XBLOCK: tl.constexpr):
    xnumel = 4
    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)
    tmp1 = tl.load(in_ptr1 + (r1 + 8 * x0), xmask, other=0.0)
    tmp2 = tl.load(in_ptr2 + (r1 + 8 * 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 = 65536.0
    tmp17 = tmp14 / tmp16
    tmp18 = 1e-05
    tmp19 = tmp17 + tmp18
    tmp20 = libdevice.rsqrt(tmp19)
    tl.debug_barrier()
    tl.store(in_out_ptr0 + x0, tmp20, xmask)
    tl.store(out_ptr0 + x0, tmp13, xmask)


@triton.jit
def triton_poi_fused_native_layer_norm_relu_32(in_ptr0, in_ptr1, in_ptr2,
    in_ptr3, in_ptr4, out_ptr0, ynumel, xnumel, YBLOCK: tl.constexpr,
    XBLOCK: tl.constexpr):
    xnumel = 256
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y0 = yindex % 256
    y1 = yindex // 256
    tmp0 = tl.load(in_ptr0 + (y0 + 256 * x2 + 65536 * y1), xmask,
        eviction_policy='evict_last')
    tmp1 = tl.load(in_ptr1 + y1, None, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr2 + y1, None, eviction_policy='evict_last')
    tmp5 = tl.load(in_ptr3 + (x2 + 256 * y0), xmask, eviction_policy=
        'evict_last')
    tmp7 = tl.load(in_ptr4 + (x2 + 256 * y0), xmask, eviction_policy=
        'evict_last')
    tmp2 = tmp0 - tmp1
    tmp4 = tmp2 * tmp3
    tmp6 = tmp4 * tmp5
    tmp8 = tmp6 + tmp7
    tmp9 = tl.full([1, 1], 0, tl.int32)
    tmp10 = triton_helpers.maximum(tmp9, tmp8)
    tl.store(out_ptr0 + (y0 + 256 * x2 + 65536 * y1), tmp10, xmask)


@triton.jit
def triton_poi_fused_add_convolution_div_33(in_out_ptr0, in_ptr0, in_ptr1,
    in_ptr2, xnumel, XBLOCK: tl.constexpr):
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    tl.full([XBLOCK], True, tl.int1)
    x3 = xindex
    x0 = xindex % 512
    x1 = xindex // 512 % 8
    x2 = xindex // 4096
    tmp0 = tl.load(in_out_ptr0 + x3, None)
    tmp1 = tl.load(in_ptr0 + x0, None, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr1 + (x0 + 1024 * x1 + 16384 * x2), None)
    tmp4 = tl.load(in_ptr2 + x0, None, eviction_policy='evict_last')
    tmp6 = tl.load(in_ptr1 + (8192 + x0 + 1024 * x1 + 16384 * x2), None)
    tmp9 = tl.load(in_ptr1 + (512 + x0 + 1024 * x1 + 16384 * x2), None)
    tmp12 = tl.load(in_ptr1 + (8704 + x0 + 1024 * x1 + 16384 * x2), None)
    tmp2 = tmp0 + tmp1
    tmp5 = tmp3 + tmp4
    tmp7 = tmp6 + tmp4
    tmp8 = tmp5 + tmp7
    tmp10 = tmp9 + tmp4
    tmp11 = tmp8 + tmp10
    tmp13 = tmp12 + tmp4
    tmp14 = tmp11 + tmp13
    tmp15 = 0.25
    tmp16 = tmp14 * tmp15
    tmp17 = tmp2 + tmp16
    tl.store(in_out_ptr0 + x3, tmp17, None)


@triton.jit
def triton_poi_fused_add_div_34(in_ptr0, out_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 % 512
    x1 = xindex // 512 % 4
    x2 = xindex // 2048
    x3 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + 1024 * x1 + 8192 * x2), None)
    tmp1 = tl.load(in_ptr0 + (4096 + x0 + 1024 * x1 + 8192 * x2), None)
    tmp3 = tl.load(in_ptr0 + (512 + x0 + 1024 * x1 + 8192 * x2), None)
    tmp5 = tl.load(in_ptr0 + (4608 + x0 + 1024 * x1 + 8192 * x2), None)
    tmp2 = tmp0 + tmp1
    tmp4 = tmp2 + tmp3
    tmp6 = tmp4 + tmp5
    tmp7 = 0.25
    tmp8 = tmp6 * tmp7
    tl.store(out_ptr0 + x3, tmp8, None)


@triton.jit
def triton_per_fused_native_layer_norm_35(in_ptr0, out_ptr0, out_ptr1,
    out_ptr2, xnumel, rnumel, XBLOCK: tl.constexpr):
    xnumel = 1024
    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, :]
    tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r2 = rindex
    x0 = xindex % 256
    x1 = xindex // 256
    x3 = xindex
    tmp0 = tl.load(in_ptr0 + (2 * x0 + 512 * (r2 % 64) + 32768 * x1 + r2 //
        64), xmask, eviction_policy='evict_last', 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], 128, 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]
    tl.store(out_ptr0 + x3, tmp10, xmask)
    tl.store(out_ptr1 + x3, tmp16, xmask)
    tl.store(out_ptr2 + x3, tmp9, xmask)


@triton.jit
def triton_per_fused_native_layer_norm_36(in_ptr0, in_ptr1, in_ptr2,
    out_ptr0, out_ptr1, out_ptr2, xnumel, rnumel, XBLOCK: tl.constexpr):
    xnumel = 16
    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
    tmp0 = tl.load(in_ptr0 + (r1 + 64 * x0), xmask, other=0.0)
    tmp1 = tl.load(in_ptr1 + (r1 + 64 * x0), xmask, other=0.0)
    tmp2 = tl.load(in_ptr2 + (r1 + 64 * 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]
    tmp15 = tmp12[:, None]
    tl.store(out_ptr0 + x0, tmp13, xmask)
    tl.store(out_ptr1 + x0, tmp14, xmask)
    tl.store(out_ptr2 + x0, tmp15, xmask)


@triton.jit
def triton_per_fused_native_layer_norm_native_layer_norm_backward_37(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)
    tmp21 = 3.0517578125e-05
    tmp22 = tmp20 * tmp21
    tl.store(out_ptr2 + x0, tmp22, xmask)
    tl.store(out_ptr0 + x0, tmp13, xmask)
    tl.store(out_ptr1 + x0, tmp14, xmask)


@triton.jit
def triton_poi_fused_native_layer_norm_38(in_out_ptr0, in_ptr0, in_ptr1,
    xnumel, XBLOCK: tl.constexpr):
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    tl.full([XBLOCK], True, tl.int1)
    x2 = xindex
    x1 = xindex // 32768
    tmp0 = tl.load(in_out_ptr0 + x2, None)
    tmp1 = tl.load(in_ptr0 + x1, None, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr1 + x1, None, eviction_policy='evict_last')
    tmp2 = tmp0 - tmp1
    tmp4 = 32768.0
    tmp5 = tmp3 / tmp4
    tmp6 = 1e-05
    tmp7 = tmp5 + tmp6
    tmp8 = libdevice.rsqrt(tmp7)
    tmp9 = tmp2 * tmp8
    tl.store(in_out_ptr0 + x2, tmp9, None)


@triton.jit
def triton_poi_fused_native_layer_norm_relu_39(in_ptr0, in_ptr1, in_ptr2,
    out_ptr0, ynumel, xnumel, YBLOCK: tl.constexpr, XBLOCK: tl.constexpr):
    xnumel = 64
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y0 = yindex % 512
    y1 = yindex // 512
    tmp0 = tl.load(in_ptr0 + (y0 + 512 * x2 + 32768 * y1), xmask,
        eviction_policy='evict_last')
    tmp1 = tl.load(in_ptr1 + (x2 + 64 * y0), xmask, eviction_policy=
        'evict_last')
    tmp3 = tl.load(in_ptr2 + (x2 + 64 * y0), xmask, eviction_policy=
        'evict_last')
    tmp2 = tmp0 * tmp1
    tmp4 = tmp2 + tmp3
    tmp5 = tl.full([1, 1], 0, tl.int32)
    tmp6 = triton_helpers.maximum(tmp5, tmp4)
    tl.store(out_ptr0 + (y0 + 512 * x2 + 32768 * y1), tmp6, xmask)


@triton.jit
def triton_per_fused_native_layer_norm_40(in_out_ptr0, in_ptr0, in_ptr1,
    in_ptr2, out_ptr0, 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.debug_barrier()
    tl.store(in_out_ptr0 + x0, tmp20, xmask)
    tl.store(out_ptr0 + x0, tmp13, xmask)


@triton.jit
def triton_poi_fused_native_layer_norm_relu_41(in_ptr0, in_ptr1, in_ptr2,
    in_ptr3, in_ptr4, out_ptr0, ynumel, xnumel, YBLOCK: tl.constexpr,
    XBLOCK: tl.constexpr):
    xnumel = 64
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y0 = yindex % 512
    y1 = yindex // 512
    tmp0 = tl.load(in_ptr0 + (y0 + 512 * x2 + 32768 * y1), xmask,
        eviction_policy='evict_last')
    tmp1 = tl.load(in_ptr1 + y1, None, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr2 + y1, None, eviction_policy='evict_last')
    tmp5 = tl.load(in_ptr3 + (x2 + 64 * y0), xmask, eviction_policy=
        'evict_last')
    tmp7 = tl.load(in_ptr4 + (x2 + 64 * y0), xmask, eviction_policy=
        'evict_last')
    tmp2 = tmp0 - tmp1
    tmp4 = tmp2 * tmp3
    tmp6 = tmp4 * tmp5
    tmp8 = tmp6 + tmp7
    tmp9 = tl.full([1, 1], 0, tl.int32)
    tmp10 = triton_helpers.maximum(tmp9, tmp8)
    tl.store(out_ptr0 + (y0 + 512 * x2 + 32768 * y1), tmp10, xmask)


@triton.jit
def triton_poi_fused_add_convolution_div_42(in_ptr0, in_ptr1, in_ptr2,
    in_ptr3, out_ptr0, ynumel, xnumel, YBLOCK: tl.constexpr, XBLOCK: tl.
    constexpr):
    ynumel = 64
    xnumel = 512
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = yindex < ynumel
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x3 = xindex
    y4 = yindex
    y0 = yindex % 4
    y5 = yindex // 4
    y2 = yindex // 16
    y6 = yindex % 16
    tmp0 = tl.load(in_ptr0 + (x3 + 512 * y4), xmask & ymask,
        eviction_policy='evict_last')
    tmp1 = tl.load(in_ptr1 + x3, xmask, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr2 + (x3 + 1024 * y0 + 8192 * y5), xmask & ymask,
        eviction_policy='evict_last')
    tmp4 = tl.load(in_ptr3 + x3, xmask, eviction_policy='evict_last')
    tmp6 = tl.load(in_ptr2 + (4096 + x3 + 1024 * y0 + 8192 * y5), xmask &
        ymask, eviction_policy='evict_last')
    tmp9 = tl.load(in_ptr2 + (512 + x3 + 1024 * y0 + 8192 * y5), xmask &
        ymask, eviction_policy='evict_last')
    tmp12 = tl.load(in_ptr2 + (4608 + x3 + 1024 * y0 + 8192 * y5), xmask &
        ymask, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tmp5 = tmp3 + tmp4
    tmp7 = tmp6 + tmp4
    tmp8 = tmp5 + tmp7
    tmp10 = tmp9 + tmp4
    tmp11 = tmp8 + tmp10
    tmp13 = tmp12 + tmp4
    tmp14 = tmp11 + tmp13
    tmp15 = 0.25
    tmp16 = tmp14 * tmp15
    tmp17 = tmp2 + tmp16
    tl.store(out_ptr0 + (y6 + 16 * x3 + 8192 * y2), tmp17, xmask & ymask)


@triton.jit
def triton_poi_fused_tanh_43(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
    x2 = xindex
    x0 = xindex % 128
    tmp0 = tl.load(in_out_ptr0 + x2, xmask)
    tmp1 = tl.load(in_ptr0 + x0, xmask, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tmp3 = libdevice.tanh(tmp2)
    tl.store(in_out_ptr0 + x2, tmp3, xmask)


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, primals_32,
        primals_33, primals_34, primals_35, primals_36, primals_37,
        primals_38, primals_39, primals_40, primals_41) = args
    args.clear()
    assert_size_stride(primals_1, (4, 3, 64, 64), (12288, 4096, 64, 1))
    assert_size_stride(primals_2, (64, 3, 3, 3), (27, 9, 3, 1))
    assert_size_stride(primals_3, (64,), (1,))
    assert_size_stride(primals_4, (128, 64, 1, 1), (64, 1, 1, 1))
    assert_size_stride(primals_5, (128,), (1,))
    assert_size_stride(primals_6, (64, 64, 64), (4096, 64, 1))
    assert_size_stride(primals_7, (64, 64, 64), (4096, 64, 1))
    assert_size_stride(primals_8, (64, 64, 3, 3), (576, 9, 3, 1))
    assert_size_stride(primals_9, (64, 64, 64), (4096, 64, 1))
    assert_size_stride(primals_10, (64, 64, 64), (4096, 64, 1))
    assert_size_stride(primals_11, (128, 64, 3, 3), (576, 9, 3, 1))
    assert_size_stride(primals_12, (128,), (1,))
    assert_size_stride(primals_13, (256, 128, 1, 1), (128, 1, 1, 1))
    assert_size_stride(primals_14, (256,), (1,))
    assert_size_stride(primals_15, (128, 32, 32), (1024, 32, 1))
    assert_size_stride(primals_16, (128, 32, 32), (1024, 32, 1))
    assert_size_stride(primals_17, (128, 128, 3, 3), (1152, 9, 3, 1))
    assert_size_stride(primals_18, (128, 32, 32), (1024, 32, 1))
    assert_size_stride(primals_19, (128, 32, 32), (1024, 32, 1))
    assert_size_stride(primals_20, (256, 128, 3, 3), (1152, 9, 3, 1))
    assert_size_stride(primals_21, (256,), (1,))
    assert_size_stride(primals_22, (512, 256, 1, 1), (256, 1, 1, 1))
    assert_size_stride(primals_23, (512,), (1,))
    assert_size_stride(primals_24, (256, 16, 16), (256, 16, 1))
    assert_size_stride(primals_25, (256, 16, 16), (256, 16, 1))
    assert_size_stride(primals_26, (256, 256, 3, 3), (2304, 9, 3, 1))
    assert_size_stride(primals_27, (256, 16, 16), (256, 16, 1))
    assert_size_stride(primals_28, (256, 16, 16), (256, 16, 1))
    assert_size_stride(primals_29, (512, 256, 3, 3), (2304, 9, 3, 1))
    assert_size_stride(primals_30, (512,), (1,))
    assert_size_stride(primals_31, (512, 512, 1, 1), (512, 1, 1, 1))
    assert_size_stride(primals_32, (512,), (1,))
    assert_size_stride(primals_33, (512, 8, 8), (64, 8, 1))
    assert_size_stride(primals_34, (512, 8, 8), (64, 8, 1))
    assert_size_stride(primals_35, (512, 512, 3, 3), (4608, 9, 3, 1))
    assert_size_stride(primals_36, (512, 8, 8), (64, 8, 1))
    assert_size_stride(primals_37, (512, 8, 8), (64, 8, 1))
    assert_size_stride(primals_38, (512, 512, 3, 3), (4608, 9, 3, 1))
    assert_size_stride(primals_39, (512,), (1,))
    assert_size_stride(primals_40, (128, 8192), (8192, 1))
    assert_size_stride(primals_41, (128,), (1,))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((64, 3, 3, 3), (27, 1, 9, 3), torch.float32)
        get_raw_stream(0)
        triton_poi_fused_0[grid(192, 9)](primals_2, buf0, 192, 9, XBLOCK=16,
            YBLOCK=64, num_warps=4, num_stages=1)
        del primals_2
        buf1 = empty_strided_cuda((64, 64, 3, 3), (576, 1, 192, 64), torch.
            float32)
        triton_poi_fused_1[grid(4096, 9)](primals_8, buf1, 4096, 9, XBLOCK=
            16, YBLOCK=64, num_warps=4, num_stages=1)
        del primals_8
        buf2 = empty_strided_cuda((128, 64, 3, 3), (576, 1, 192, 64), torch
            .float32)
        triton_poi_fused_2[grid(8192, 9)](primals_11, buf2, 8192, 9, XBLOCK
            =16, YBLOCK=64, num_warps=4, num_stages=1)
        del primals_11
        buf3 = empty_strided_cuda((128, 128, 3, 3), (1152, 1, 384, 128),
            torch.float32)
        triton_poi_fused_3[grid(16384, 9)](primals_17, buf3, 16384, 9,
            XBLOCK=16, YBLOCK=64, num_warps=4, num_stages=1)
        del primals_17
        buf4 = empty_strided_cuda((256, 128, 3, 3), (1152, 1, 384, 128),
            torch.float32)
        triton_poi_fused_4[grid(32768, 9)](primals_20, buf4, 32768, 9,
            XBLOCK=16, YBLOCK=64, num_warps=4, num_stages=1)
        del primals_20
        buf5 = empty_strided_cuda((256, 256, 3, 3), (2304, 1, 768, 256),
            torch.float32)
        triton_poi_fused_5[grid(65536, 9)](primals_26, buf5, 65536, 9,
            XBLOCK=16, YBLOCK=64, num_warps=4, num_stages=1)
        del primals_26
        buf6 = empty_strided_cuda((512, 256, 3, 3), (2304, 1, 768, 256),
            torch.float32)
        triton_poi_fused_6[grid(131072, 9)](primals_29, buf6, 131072, 9,
            XBLOCK=16, YBLOCK=64, num_warps=4, num_stages=1)
        del primals_29
        buf7 = empty_strided_cuda((512, 512, 3, 3), (4608, 1, 1536, 512),
            torch.float32)
        triton_poi_fused_7[grid(262144, 9)](primals_35, buf7, 262144, 9,
            XBLOCK=16, YBLOCK=64, num_warps=4, num_stages=1)
        del primals_35
        buf8 = empty_strided_cuda((512, 512, 3, 3), (4608, 1, 1536, 512),
            torch.float32)
        triton_poi_fused_7[grid(262144, 9)](primals_38, buf8, 262144, 9,
            XBLOCK=16, YBLOCK=64, num_warps=4, num_stages=1)
        del primals_38
        buf9 = empty_strided_cuda((4, 3, 64, 64), (12288, 1, 192, 3), torch
            .float32)
        triton_poi_fused_view_8[grid(12, 4096)](primals_1, buf9, 12, 4096,
            XBLOCK=64, YBLOCK=16, num_warps=4, num_stages=1)
        del primals_1
        buf10 = extern_kernels.convolution(buf9, buf0, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf10, (4, 64, 64, 64), (262144, 1, 4096, 64))
        buf11 = buf10
        del buf10
        triton_poi_fused_convolution_9[grid(1048576)](buf11, primals_3, 
            1048576, XBLOCK=1024, num_warps=4, num_stages=1)
        del primals_3
        buf12 = empty_strided_cuda((4, 64, 32, 32), (65536, 1, 2048, 64),
            torch.float32)
        triton_poi_fused_add_div_10[grid(262144)](buf11, buf12, 262144,
            XBLOCK=512, num_warps=8, num_stages=1)
        buf13 = extern_kernels.convolution(buf12, primals_4, stride=(1, 1),
            padding=(0, 0), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf13, (4, 128, 32, 32), (131072, 1, 4096, 128))
        buf14 = empty_strided_cuda((4, 1, 1, 1, 32, 64), (2048, 8192, 8192,
            8192, 1, 32), torch.float32)
        buf15 = empty_strided_cuda((4, 1, 1, 1, 32, 64), (2048, 8192, 8192,
            8192, 1, 32), torch.float32)
        buf16 = empty_strided_cuda((4, 1, 1, 1, 32, 64), (2048, 8192, 8192,
            8192, 1, 32), torch.float32)
        triton_per_fused_native_layer_norm_11[grid(8192)](buf11, buf14,
            buf15, buf16, 8192, 128, XBLOCK=8, num_warps=8, num_stages=1)
        buf17 = empty_strided_cuda((4, 1, 1, 1, 32), (32, 128, 128, 128, 1),
            torch.float32)
        buf18 = empty_strided_cuda((4, 1, 1, 1, 32), (32, 128, 128, 128, 1),
            torch.float32)
        buf19 = empty_strided_cuda((4, 1, 1, 1, 32), (32, 128, 128, 128, 1),
            torch.float32)
        triton_per_fused_native_layer_norm_12[grid(128)](buf14, buf15,
            buf16, buf17, buf18, buf19, 128, 64, XBLOCK=1, num_warps=2,
            num_stages=1)
        buf20 = empty_strided_cuda((4, 1, 1, 1), (1, 1, 1, 1), torch.float32)
        buf21 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf23 = reinterpret_tensor(buf21, (4, 1, 1, 1), (1, 1, 1, 1), 0)
        del buf21
        triton_per_fused_native_layer_norm_13[grid(4)](buf23, buf17, buf18,
            buf19, buf20, 4, 32, XBLOCK=1, num_warps=2, num_stages=1)
        buf24 = empty_strided_cuda((4, 64, 64, 64), (262144, 1, 4096, 64),
            torch.float32)
        triton_poi_fused_native_layer_norm_relu_14[grid(256, 4096)](buf11,
            buf20, buf23, primals_6, primals_7, buf24, 256, 4096, XBLOCK=32,
            YBLOCK=32, num_warps=4, num_stages=1)
        del primals_7
        buf25 = extern_kernels.convolution(buf24, buf1, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf25, (4, 64, 64, 64), (262144, 1, 4096, 64))
        buf26 = buf16
        del buf16
        buf27 = buf15
        del buf15
        buf28 = buf14
        del buf14
        triton_per_fused_native_layer_norm_11[grid(8192)](buf25, buf26,
            buf27, buf28, 8192, 128, XBLOCK=8, num_warps=8, num_stages=1)
        buf29 = buf19
        del buf19
        buf30 = buf18
        del buf18
        buf31 = buf17
        del buf17
        triton_per_fused_native_layer_norm_12[grid(128)](buf26, buf27,
            buf28, buf29, buf30, buf31, 128, 64, XBLOCK=1, num_warps=2,
            num_stages=1)
        del buf26
        del buf27
        del buf28
        buf32 = empty_strided_cuda((4, 1, 1, 1), (1, 1, 1, 1), torch.float32)
        buf33 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf35 = reinterpret_tensor(buf33, (4, 1, 1, 1), (1, 1, 1, 1), 0)
        del buf33
        triton_per_fused_native_layer_norm_13[grid(4)](buf35, buf29, buf30,
            buf31, buf32, 4, 32, XBLOCK=1, num_warps=2, num_stages=1)
        del buf29
        del buf30
        del buf31
        buf36 = empty_strided_cuda((4, 64, 64, 64), (262144, 1, 4096, 64),
            torch.float32)
        triton_poi_fused_native_layer_norm_relu_14[grid(256, 4096)](buf25,
            buf32, buf35, primals_9, primals_10, buf36, 256, 4096, XBLOCK=
            32, YBLOCK=32, num_warps=4, num_stages=1)
        del primals_10
        buf37 = extern_kernels.convolution(buf36, buf2, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf37, (4, 128, 64, 64), (524288, 1, 8192, 128))
        buf38 = buf13
        del buf13
        triton_poi_fused_add_convolution_div_15[grid(524288)](buf38,
            primals_5, buf37, primals_12, 524288, XBLOCK=512, num_warps=8,
            num_stages=1)
        del buf37
        del primals_12
        del primals_5
        buf39 = empty_strided_cuda((4, 128, 16, 16), (32768, 1, 2048, 128),
            torch.float32)
        triton_poi_fused_add_div_16[grid(131072)](buf38, buf39, 131072,
            XBLOCK=512, num_warps=8, num_stages=1)
        buf40 = extern_kernels.convolution(buf39, primals_13, stride=(1, 1),
            padding=(0, 0), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf40, (4, 256, 16, 16), (65536, 1, 4096, 256))
        buf41 = empty_strided_cuda((4, 1, 1, 1, 16, 64), (1024, 4096, 4096,
            4096, 1, 16), torch.float32)
        buf42 = empty_strided_cuda((4, 1, 1, 1, 16, 64), (1024, 4096, 4096,
            4096, 1, 16), torch.float32)
        buf43 = empty_strided_cuda((4, 1, 1, 1, 16, 64), (1024, 4096, 4096,
            4096, 1, 16), torch.float32)
        triton_per_fused_native_layer_norm_17[grid(4096)](buf38, buf41,
            buf42, buf43, 4096, 128, XBLOCK=8, num_warps=8, num_stages=1)
        buf44 = empty_strided_cuda((4, 1, 1, 1, 16), (16, 64, 64, 64, 1),
            torch.float32)
        buf45 = empty_strided_cuda((4, 1, 1, 1, 16), (16, 64, 64, 64, 1),
            torch.float32)
        buf46 = empty_strided_cuda((4, 1, 1, 1, 16), (16, 64, 64, 64, 1),
            torch.float32)
        triton_per_fused_native_layer_norm_18[grid(64)](buf41, buf42, buf43,
            buf44, buf45, buf46, 64, 64, XBLOCK=1, num_warps=2, num_stages=1)
        buf47 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf48 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf124 = empty_strided_cuda((4, 1, 1, 1), (1, 1, 1, 1), torch.float32)
        triton_per_fused_native_layer_norm_native_layer_norm_backward_19[grid
            (4)](buf44, buf45, buf46, buf47, buf48, buf124, 4, 16, XBLOCK=1,
            num_warps=2, num_stages=1)
        buf50 = buf38
        del buf38
        triton_poi_fused_native_layer_norm_20[grid(524288)](buf50, buf47,
            buf48, 524288, XBLOCK=1024, num_warps=4, num_stages=1)
        buf51 = empty_strided_cuda((4, 128, 32, 32), (131072, 1, 4096, 128),
            torch.float32)
        triton_poi_fused_native_layer_norm_relu_21[grid(512, 1024)](buf50,
            primals_15, primals_16, buf51, 512, 1024, XBLOCK=32, YBLOCK=32,
            num_warps=4, num_stages=1)
        del primals_16
        buf52 = extern_kernels.convolution(buf51, buf3, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf52, (4, 128, 32, 32), (131072, 1, 4096, 128))
        buf53 = buf43
        del buf43
        buf54 = buf42
        del buf42
        buf55 = buf41
        del buf41
        triton_per_fused_native_layer_norm_17[grid(4096)](buf52, buf53,
            buf54, buf55, 4096, 128, XBLOCK=8, num_warps=8, num_stages=1)
        buf56 = buf46
        del buf46
        buf57 = buf45
        del buf45
        buf58 = buf44
        del buf44
        triton_per_fused_native_layer_norm_18[grid(64)](buf53, buf54, buf55,
            buf56, buf57, buf58, 64, 64, XBLOCK=1, num_warps=2, num_stages=1)
        del buf53
        del buf54
        del buf55
        buf59 = reinterpret_tensor(buf48, (4, 1, 1, 1), (1, 1, 1, 1), 0)
        del buf48
        buf60 = buf47
        del buf47
        buf62 = reinterpret_tensor(buf60, (4, 1, 1, 1), (1, 1, 1, 1), 0)
        del buf60
        triton_per_fused_native_layer_norm_22[grid(4)](buf62, buf56, buf57,
            buf58, buf59, 4, 16, XBLOCK=1, num_warps=2, num_stages=1)
        del buf56
        del buf57
        del buf58
        buf63 = empty_strided_cuda((4, 128, 32, 32), (131072, 1, 4096, 128),
            torch.float32)
        triton_poi_fused_native_layer_norm_relu_23[grid(512, 1024)](buf52,
            buf59, buf62, primals_18, primals_19, buf63, 512, 1024, XBLOCK=
            32, YBLOCK=32, num_warps=4, num_stages=1)
        del primals_19
        buf64 = extern_kernels.convolution(buf63, buf4, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf64, (4, 256, 32, 32), (262144, 1, 8192, 256))
        buf65 = buf40
        del buf40
        triton_poi_fused_add_convolution_div_24[grid(262144)](buf65,
            primals_14, buf64, primals_21, 262144, XBLOCK=512, num_warps=8,
            num_stages=1)
        del buf64
        del primals_14
        del primals_21
        buf66 = empty_strided_cuda((4, 256, 8, 8), (16384, 1, 2048, 256),
            torch.float32)
        triton_poi_fused_add_div_25[grid(65536)](buf65, buf66, 65536,
            XBLOCK=512, num_warps=4, num_stages=1)
        buf67 = extern_kernels.convolution(buf66, primals_22, stride=(1, 1),
            padding=(0, 0), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf67, (4, 512, 8, 8), (32768, 1, 4096, 512))
        buf68 = empty_strided_cuda((4, 1, 1, 1, 8, 64), (512, 2048, 2048, 
            2048, 1, 8), torch.float32)
        buf69 = empty_strided_cuda((4, 1, 1, 1, 8, 64), (512, 2048, 2048, 
            2048, 1, 8), torch.float32)
        buf70 = empty_strided_cuda((4, 1, 1, 1, 8, 64), (512, 2048, 2048, 
            2048, 1, 8), torch.float32)
        triton_per_fused_native_layer_norm_26[grid(2048)](buf65, buf68,
            buf69, buf70, 2048, 128, XBLOCK=32, num_warps=8, num_stages=1)
        buf71 = empty_strided_cuda((4, 1, 1, 1, 8), (8, 32, 32, 32, 1),
            torch.float32)
        buf72 = empty_strided_cuda((4, 1, 1, 1, 8), (8, 32, 32, 32, 1),
            torch.float32)
        buf73 = empty_strided_cuda((4, 1, 1, 1, 8), (8, 32, 32, 32, 1),
            torch.float32)
        triton_per_fused_native_layer_norm_27[grid(32)](buf68, buf69, buf70,
            buf71, buf72, buf73, 32, 64, XBLOCK=1, num_warps=2, num_stages=1)
        buf74 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf75 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf123 = empty_strided_cuda((4, 1, 1, 1), (1, 1, 1, 1), torch.float32)
        triton_per_fused_native_layer_norm_native_layer_norm_backward_28[grid
            (4)](buf71, buf72, buf73, buf74, buf75, buf123, 4, 8, XBLOCK=1,
            num_warps=2, num_stages=1)
        buf77 = buf65
        del buf65
        triton_poi_fused_native_layer_norm_29[grid(262144)](buf77, buf74,
            buf75, 262144, XBLOCK=1024, num_warps=4, num_stages=1)
        buf78 = empty_strided_cuda((4, 256, 16, 16), (65536, 1, 4096, 256),
            torch.float32)
        triton_poi_fused_native_layer_norm_relu_30[grid(1024, 256)](buf77,
            primals_24, primals_25, buf78, 1024, 256, XBLOCK=32, YBLOCK=32,
            num_warps=4, num_stages=1)
        del primals_25
        buf79 = extern_kernels.convolution(buf78, buf5, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf79, (4, 256, 16, 16), (65536, 1, 4096, 256))
        buf80 = buf70
        del buf70
        buf81 = buf69
        del buf69
        buf82 = buf68
        del buf68
        triton_per_fused_native_layer_norm_26[grid(2048)](buf79, buf80,
            buf81, buf82, 2048, 128, XBLOCK=32, num_warps=8, num_stages=1)
        buf83 = buf73
        del buf73
        buf84 = buf72
        del buf72
        buf85 = buf71
        del buf71
        triton_per_fused_native_layer_norm_27[grid(32)](buf80, buf81, buf82,
            buf83, buf84, buf85, 32, 64, XBLOCK=1, num_warps=2, num_stages=1)
        del buf80
        del buf81
        del buf82
        buf86 = reinterpret_tensor(buf75, (4, 1, 1, 1), (1, 1, 1, 1), 0)
        del buf75
        buf87 = buf74
        del buf74
        buf89 = reinterpret_tensor(buf87, (4, 1, 1, 1), (1, 1, 1, 1), 0)
        del buf87
        triton_per_fused_native_layer_norm_31[grid(4)](buf89, buf83, buf84,
            buf85, buf86, 4, 8, XBLOCK=1, num_warps=2, num_stages=1)
        del buf83
        del buf84
        del buf85
        buf90 = empty_strided_cuda((4, 256, 16, 16), (65536, 1, 4096, 256),
            torch.float32)
        triton_poi_fused_native_layer_norm_relu_32[grid(1024, 256)](buf79,
            buf86, buf89, primals_27, primals_28, buf90, 1024, 256, XBLOCK=
            32, YBLOCK=32, num_warps=4, num_stages=1)
        del primals_28
        buf91 = extern_kernels.convolution(buf90, buf6, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf91, (4, 512, 16, 16), (131072, 1, 8192, 512))
        buf92 = buf67
        del buf67
        triton_poi_fused_add_convolution_div_33[grid(131072)](buf92,
            primals_23, buf91, primals_30, 131072, XBLOCK=512, num_warps=8,
            num_stages=1)
        del buf91
        del primals_23
        del primals_30
        buf93 = empty_strided_cuda((4, 512, 4, 4), (8192, 1, 2048, 512),
            torch.float32)
        triton_poi_fused_add_div_34[grid(32768)](buf92, buf93, 32768,
            XBLOCK=256, num_warps=4, num_stages=1)
        buf94 = extern_kernels.convolution(buf93, primals_31, stride=(1, 1),
            padding=(0, 0), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf94, (4, 512, 4, 4), (8192, 1, 2048, 512))
        buf95 = empty_strided_cuda((4, 1, 1, 1, 4, 64), (256, 1024, 1024, 
            1024, 64, 1), torch.float32)
        buf96 = empty_strided_cuda((4, 1, 1, 1, 4, 64), (256, 1024, 1024, 
            1024, 64, 1), torch.float32)
        buf97 = empty_strided_cuda((4, 1, 1, 1, 4, 64), (256, 1024, 1024, 
            1024, 64, 1), torch.float32)
        triton_per_fused_native_layer_norm_35[grid(1024)](buf92, buf95,
            buf96, buf97, 1024, 128, XBLOCK=8, num_warps=8, num_stages=1)
        buf98 = empty_strided_cuda((4, 1, 1, 1, 4), (4, 16, 16, 16, 1),
            torch.float32)
        buf99 = empty_strided_cuda((4, 1, 1, 1, 4), (4, 16, 16, 16, 1),
            torch.float32)
        buf100 = empty_strided_cuda((4, 1, 1, 1, 4), (4, 16, 16, 16, 1),
            torch.float32)
        triton_per_fused_native_layer_norm_36[grid(16)](buf95, buf96, buf97,
            buf98, buf99, buf100, 16, 64, XBLOCK=8, num_warps=4, num_stages=1)
        buf101 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf102 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf122 = empty_strided_cuda((4, 1, 1, 1), (1, 1, 1, 1), torch.float32)
        triton_per_fused_native_layer_norm_native_layer_norm_backward_37[grid
            (4)](buf98, buf99, buf100, buf101, buf102, buf122, 4, 4, XBLOCK
            =1, num_warps=2, num_stages=1)
        buf104 = buf92
        del buf92
        triton_poi_fused_native_layer_norm_38[grid(131072)](buf104, buf101,
            buf102, 131072, XBLOCK=512, num_warps=8, num_stages=1)
        buf105 = empty_strided_cuda((4, 512, 8, 8), (32768, 1, 4096, 512),
            torch.float32)
        triton_poi_fused_native_layer_norm_relu_39[grid(2048, 64)](buf104,
            primals_33, primals_34, buf105, 2048, 64, XBLOCK=32, YBLOCK=32,
            num_warps=4, num_stages=1)
        del primals_34
        buf106 = extern_kernels.convolution(buf105, buf7, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf106, (4, 512, 8, 8), (32768, 1, 4096, 512))
        buf107 = buf97
        del buf97
        buf108 = buf96
        del buf96
        buf109 = buf95
        del buf95
        triton_per_fused_native_layer_norm_35[grid(1024)](buf106, buf107,
            buf108, buf109, 1024, 128, XBLOCK=8, num_warps=8, num_stages=1)
        buf110 = buf99
        del buf99
        buf111 = buf98
        del buf98
        buf112 = buf100
        del buf100
        triton_per_fused_native_layer_norm_36[grid(16)](buf107, buf108,
            buf109, buf110, buf111, buf112, 16, 64, XBLOCK=8, num_warps=4,
            num_stages=1)
        del buf107
        del buf108
        del buf109
        buf113 = reinterpret_tensor(buf102, (4, 1, 1, 1), (1, 1, 1, 1), 0)
        del buf102
        buf114 = buf101
        del buf101
        buf116 = reinterpret_tensor(buf114, (4, 1, 1, 1), (1, 1, 1, 1), 0)
        del buf114
        triton_per_fused_native_layer_norm_40[grid(4)](buf116, buf110,
            buf111, buf112, buf113, 4, 4, XBLOCK=1, num_warps=2, num_stages=1)
        del buf110
        del buf111
        del buf112
        buf117 = empty_strided_cuda((4, 512, 8, 8), (32768, 1, 4096, 512),
            torch.float32)
        triton_poi_fused_native_layer_norm_relu_41[grid(2048, 64)](buf106,
            buf113, buf116, primals_36, primals_37, buf117, 2048, 64,
            XBLOCK=32, YBLOCK=32, num_warps=4, num_stages=1)
        del primals_37
        buf118 = extern_kernels.convolution(buf117, buf8, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf118, (4, 512, 8, 8), (32768, 1, 4096, 512))
        buf119 = empty_strided_cuda((4, 512, 4, 4), (8192, 16, 4, 1), torch
            .float32)
        triton_poi_fused_add_convolution_div_42[grid(64, 512)](buf94,
            primals_32, buf118, primals_39, buf119, 64, 512, XBLOCK=4,
            YBLOCK=64, num_warps=4, num_stages=1)
        del buf118
        del buf94
        del primals_32
        del primals_39
        buf120 = empty_strided_cuda((4, 128), (128, 1), torch.float32)
        extern_kernels.mm(reinterpret_tensor(buf119, (4, 8192), (8192, 1), 
            0), reinterpret_tensor(primals_40, (8192, 128), (1, 8192), 0),
            out=buf120)
        buf121 = buf120
        del buf120
        triton_poi_fused_tanh_43[grid(512)](buf121, primals_41, 512, XBLOCK
            =128, num_warps=4, num_stages=1)
        del primals_41
    return (buf121, buf0, primals_4, primals_6, buf1, primals_9, buf2,
        primals_13, primals_15, buf3, primals_18, buf4, primals_22,
        primals_24, buf5, primals_27, buf6, primals_31, primals_33, buf7,
        primals_36, buf8, buf9, buf11, buf12, buf20, buf23, buf24, buf25,
        buf32, buf35, buf36, buf39, buf50, buf51, buf52, buf59, buf62,
        buf63, buf66, buf77, buf78, buf79, buf86, buf89, buf90, buf93,
        buf104, buf105, buf106, buf113, buf116, buf117, reinterpret_tensor(
        buf119, (4, 8192), (8192, 1), 0), buf121, primals_40, buf122,
        buf123, buf124)


class IWConv2d(nn.Module):

    def __init__(self, input_dim, output_dim, kernel_size, he_init=True,
        stride=1, bias=True):
        super(IWConv2d, self).__init__()
        self.he_init = he_init
        self.padding = int((kernel_size - 1) / 2)
        self.conv = nn.Conv2d(input_dim, output_dim, kernel_size, stride=1,
            padding=self.padding, bias=bias)

    def forward(self, input):
        output = self.conv(input)
        return output


class ConvMeanPool(nn.Module):

    def __init__(self, input_dim, output_dim, kernel_size, he_init=True):
        super(ConvMeanPool, self).__init__()
        self.he_init = he_init
        self.conv = IWConv2d(input_dim, output_dim, kernel_size, he_init=
            self.he_init)

    def forward(self, input):
        output = self.conv(input)
        output = (output[:, :, ::2, ::2] + output[:, :, 1::2, ::2] + output
            [:, :, ::2, 1::2] + output[:, :, 1::2, 1::2]) / 4
        return output


class MeanPoolConv(nn.Module):

    def __init__(self, input_dim, output_dim, kernel_size, he_init=True):
        super(MeanPoolConv, self).__init__()
        self.he_init = he_init
        self.conv = IWConv2d(input_dim, output_dim, kernel_size, he_init=
            self.he_init)

    def forward(self, input):
        output = input
        output = (output[:, :, ::2, ::2] + output[:, :, 1::2, ::2] + output
            [:, :, ::2, 1::2] + output[:, :, 1::2, 1::2]) / 4
        output = self.conv(output)
        return output


class DepthToSpace(nn.Module):

    def __init__(self, block_size):
        super(DepthToSpace, self).__init__()
        self.block_size = block_size
        self.block_size_sq = block_size * block_size

    def forward(self, input):
        output = input.permute(0, 2, 3, 1)
        batch_size, input_height, input_width, input_depth = output.size()
        output_depth = int(input_depth / self.block_size_sq)
        output_width = int(input_width * self.block_size)
        output_height = int(input_height * self.block_size)
        t_1 = output.reshape(batch_size, input_height, input_width, self.
            block_size_sq, output_depth)
        spl = t_1.split(self.block_size, 3)
        stacks = [t_t.reshape(batch_size, input_height, output_width,
            output_depth) for t_t in spl]
        output = torch.stack(stacks, 0).transpose(0, 1).permute(0, 2, 1, 3, 4
            ).reshape(batch_size, output_height, output_width, output_depth)
        output = output.permute(0, 3, 1, 2)
        return output


class UpSampleConv(nn.Module):

    def __init__(self, input_dim, output_dim, kernel_size, he_init=True,
        bias=True):
        super(UpSampleConv, self).__init__()
        self.he_init = he_init
        self.conv = IWConv2d(input_dim, output_dim, kernel_size, he_init=
            self.he_init, bias=bias)
        self.depth_to_space = DepthToSpace(2)

    def forward(self, input):
        output = input
        output = torch.cat((output, output, output, output), 1)
        output = self.depth_to_space(output)
        output = self.conv(output)
        return output


class ResidualBlock(nn.Module):

    def __init__(self, input_dim, output_dim, kernel_size, resample=None, hw=64
        ):
        super(ResidualBlock, self).__init__()
        self.input_dim = input_dim
        self.output_dim = output_dim
        self.kernel_size = kernel_size
        self.resample = resample
        self.bn1 = None
        self.bn2 = None
        self.relu1 = nn.ReLU()
        self.relu2 = nn.ReLU()
        if resample == 'down':
            self.bn1 = nn.LayerNorm([input_dim, hw, hw])
            self.bn2 = nn.LayerNorm([input_dim, hw, hw])
        elif resample == 'up':
            self.bn1 = nn.BatchNorm2d(input_dim)
            self.bn2 = nn.BatchNorm2d(output_dim)
        elif resample is None:
            self.bn1 = nn.BatchNorm2d(output_dim)
            self.bn2 = nn.LayerNorm([input_dim, hw, hw])
        else:
            raise Exception('invalid resample value')
        if resample == 'down':
            self.conv_shortcut = MeanPoolConv(input_dim, output_dim,
                kernel_size=1, he_init=False)
            self.conv_1 = IWConv2d(input_dim, input_dim, kernel_size=
                kernel_size, bias=False)
            self.conv_2 = ConvMeanPool(input_dim, output_dim, kernel_size=
                kernel_size)
        elif resample == 'up':
            self.conv_shortcut = UpSampleConv(input_dim, output_dim,
                kernel_size=1, he_init=False)
            self.conv_1 = UpSampleConv(input_dim, output_dim, kernel_size=
                kernel_size, bias=False)
            self.conv_2 = IWConv2d(output_dim, output_dim, kernel_size=
                kernel_size)
        elif resample is None:
            self.conv_shortcut = IWConv2d(input_dim, output_dim,
                kernel_size=1, he_init=False)
            self.conv_1 = IWConv2d(input_dim, input_dim, kernel_size=
                kernel_size, bias=False)
            self.conv_2 = IWConv2d(input_dim, output_dim, kernel_size=
                kernel_size)
        else:
            raise Exception('invalid resample value')

    def forward(self, input):
        if self.input_dim == self.output_dim and self.resample is None:
            shortcut = input
        else:
            shortcut = self.conv_shortcut(input)
        output = input
        output = self.bn1(output)
        output = self.relu1(output)
        output = self.conv_1(output)
        output = self.bn2(output)
        output = self.relu2(output)
        output = self.conv_2(output)
        return shortcut + output


class IWEncoderNew(nn.Module):

    def __init__(self, input_size=64, z_dim=128, n_image_channels=3):
        super(IWEncoderNew, self).__init__()
        self.size = input_size
        self.n_image_channels = n_image_channels
        self.ssize = self.size // 16
        self.conv1 = IWConv2d(n_image_channels, self.size, 3, he_init=False)
        self.rb1 = ResidualBlock(self.size, 2 * self.size, 3, resample=
            'down', hw=self.size)
        self.rb2 = ResidualBlock(2 * self.size, 4 * self.size, 3, resample=
            'down', hw=int(self.size / 2))
        self.rb3 = ResidualBlock(4 * self.size, 8 * self.size, 3, resample=
            'down', hw=int(self.size / 4))
        self.rb4 = ResidualBlock(8 * self.size, 8 * self.size, 3, resample=
            'down', hw=int(self.size / 8))
        self.ln1 = nn.Linear(self.ssize * self.ssize * 8 * self.size, z_dim)

    def forward(self, input_0):
        primals_2 = self.conv1.conv.weight
        primals_3 = self.conv1.conv.bias
        primals_6 = self.rb1.bn1.weight
        primals_7 = self.rb1.bn1.bias
        primals_9 = self.rb1.bn2.weight
        primals_10 = self.rb1.bn2.bias
        primals_4 = self.rb1.conv_shortcut.conv.conv.weight
        primals_5 = self.rb1.conv_shortcut.conv.conv.bias
        primals_8 = self.rb1.conv_1.conv.weight
        primals_11 = self.rb1.conv_2.conv.conv.weight
        primals_12 = self.rb1.conv_2.conv.conv.bias
        primals_15 = self.rb2.bn1.weight
        primals_16 = self.rb2.bn1.bias
        primals_18 = self.rb2.bn2.weight
        primals_19 = self.rb2.bn2.bias
        primals_13 = self.rb2.conv_shortcut.conv.conv.weight
        primals_14 = self.rb2.conv_shortcut.conv.conv.bias
        primals_17 = self.rb2.conv_1.conv.weight
        primals_20 = self.rb2.conv_2.conv.conv.weight
        primals_21 = self.rb2.conv_2.conv.conv.bias
        primals_24 = self.rb3.bn1.weight
        primals_25 = self.rb3.bn1.bias
        primals_27 = self.rb3.bn2.weight
        primals_28 = self.rb3.bn2.bias
        primals_22 = self.rb3.conv_shortcut.conv.conv.weight
        primals_23 = self.rb3.conv_shortcut.conv.conv.bias
        primals_26 = self.rb3.conv_1.conv.weight
        primals_29 = self.rb3.conv_2.conv.conv.weight
        primals_30 = self.rb3.conv_2.conv.conv.bias
        primals_33 = self.rb4.bn1.weight
        primals_34 = self.rb4.bn1.bias
        primals_36 = self.rb4.bn2.weight
        primals_37 = self.rb4.bn2.bias
        primals_31 = self.rb4.conv_shortcut.conv.conv.weight
        primals_32 = self.rb4.conv_shortcut.conv.conv.bias
        primals_35 = self.rb4.conv_1.conv.weight
        primals_38 = self.rb4.conv_2.conv.conv.weight
        primals_39 = self.rb4.conv_2.conv.conv.bias
        primals_40 = self.ln1.weight
        primals_41 = self.ln1.bias
        primals_1 = 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, primals_32, primals_33, primals_34,
            primals_35, primals_36, primals_37, primals_38, primals_39,
            primals_40, primals_41])
        return output[0]
