# 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_4/inductor_cache/2y/c2yegdgextnbn5jgy4cz2lj6jlwuiko5e2ct3dyp33g6oilmvbqk.py
# Topologically Sorted Source Nodes: [cat], Original ATen: [aten.cat]
# Source node to ATen node mapping:
#   cat => cat
# Graph fragment:
#   %cat : [num_users=2] = call_function[target=torch.ops.aten.cat.default](args = ([%full_default, %primals_1], 1), kwargs = {})
triton_poi_fused_cat_0 = async_compile.triton('triton_poi_fused_cat_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=[32], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=80, major=8, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=108), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_cat_0', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 1, 'num_reduction': 0, 'backend_hash': 'A9C866B4A14FD3277824029365D703C2427B2E685E54EC9B3EF4ADC8D1EEAC1D', '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_cat_0(in_ptr0, out_ptr0, xnumel, XBLOCK : tl.constexpr):
    xnumel = 32
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex % 8
    x1 = (xindex // 8)
    x2 = xindex
    tmp0 = x0
    tmp1 = tl.full([1], 0, tl.int64)
    tmp2 = tmp0 >= tmp1
    tmp3 = tl.full([1], 4, tl.int64)
    tmp4 = tmp0 < tmp3
    tmp5 = 0.0
    tmp6 = tl.full(tmp5.shape, 0.0, tmp5.dtype)
    tmp7 = tl.where(tmp4, tmp5, tmp6)
    tmp8 = tmp0 >= tmp3
    tmp9 = tl.full([1], 8, tl.int64)
    tmp10 = tmp0 < tmp9
    tmp11 = tl.load(in_ptr0 + ((4*x1) + ((-4) + x0)), tmp8 & xmask, eviction_policy='evict_last', other=0.0)
    tmp12 = tl.where(tmp4, tmp7, tmp11)
    tl.store(out_ptr0 + (x2), tmp12, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_4/inductor_cache/2c/c2comiq2kwhy77qu2gl2txvotf5jwlmcczdxgj76reytmfyt6ea4.py
# Topologically Sorted Source Nodes: [hx, hardtanh, mul, hardsigmoid, mul_1], Original ATen: [aten.zeros, aten.hardtanh, aten.mul, aten.hardsigmoid, aten.hardsigmoid_backward, aten.hardtanh_backward]
# Source node to ATen node mapping:
#   hardsigmoid => add, clamp_max_1, clamp_min_1, div
#   hardtanh => clamp_max, clamp_min
#   hx => full_default
#   mul => mul
#   mul_1 => mul_1
# Graph fragment:
#   %full_default : [num_users=2] = call_function[target=torch.ops.aten.full.default](args = ([4, 4], 0), kwargs = {dtype: torch.float32, layout: torch.strided, device: cuda:0, pin_memory: False})
#   %add_tensor_1 : [num_users=3] = call_function[target=torch.ops.aten.add.Tensor](args = (%mm_default_1, %primals_3), kwargs = {})
#   %clamp_min : [num_users=1] = call_function[target=torch.ops.aten.clamp_min.default](args = (%add_tensor_1, -1.0), kwargs = {})
#   %clamp_max : [num_users=2] = call_function[target=torch.ops.aten.clamp_max.default](args = (%clamp_min, 1.0), kwargs = {})
#   %mul : [num_users=3] = call_function[target=torch.ops.aten.mul.Tensor](args = (%clamp_max, 3), kwargs = {})
#   %add : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%mul, 3), kwargs = {})
#   %clamp_min_1 : [num_users=1] = call_function[target=torch.ops.aten.clamp_min.default](args = (%add, 0), kwargs = {})
#   %clamp_max_1 : [num_users=1] = call_function[target=torch.ops.aten.clamp_max.default](args = (%clamp_min_1, 6), kwargs = {})
#   %div : [num_users=3] = call_function[target=torch.ops.aten.div.Tensor](args = (%clamp_max_1, 6), kwargs = {})
#   %mul_1 : [num_users=3] = call_function[target=torch.ops.aten.mul.Tensor](args = (%div, %full_default), kwargs = {})
#   %gt : [num_users=1] = call_function[target=torch.ops.aten.gt.Scalar](args = (%mul, -3.0), kwargs = {})
#   %lt : [num_users=1] = call_function[target=torch.ops.aten.lt.Scalar](args = (%mul, 3.0), kwargs = {})
#   %bitwise_and : [num_users=1] = call_function[target=torch.ops.aten.bitwise_and.Tensor](args = (%gt, %lt), kwargs = {})
#   %le_2 : [num_users=1] = call_function[target=torch.ops.aten.le.Scalar](args = (%add_tensor_1, -1.0), kwargs = {})
#   %ge_2 : [num_users=1] = call_function[target=torch.ops.aten.ge.Scalar](args = (%add_tensor_1, 1.0), kwargs = {})
#   %bitwise_or_2 : [num_users=1] = call_function[target=torch.ops.aten.bitwise_or.Tensor](args = (%le_2, %ge_2), kwargs = {})
triton_poi_fused_hardsigmoid_hardsigmoid_backward_hardtanh_hardtanh_backward_mul_zeros_1 = async_compile.triton('triton_poi_fused_hardsigmoid_hardsigmoid_backward_hardtanh_hardtanh_backward_mul_zeros_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=[16], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*i1', 4: '*fp32', 5: '*fp32', 6: '*i1', 7: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=80, major=8, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=108), '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_hardsigmoid_hardsigmoid_backward_hardtanh_hardtanh_backward_mul_zeros_1', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 2, 'num_reduction': 0, 'backend_hash': 'A9C866B4A14FD3277824029365D703C2427B2E685E54EC9B3EF4ADC8D1EEAC1D', '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_hardsigmoid_hardsigmoid_backward_hardtanh_hardtanh_backward_mul_zeros_1(in_ptr0, in_ptr1, out_ptr0, out_ptr1, out_ptr2, out_ptr3, out_ptr4, xnumel, XBLOCK : tl.constexpr):
    xnumel = 16
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x2 = xindex
    x0 = xindex % 4
    tmp0 = tl.load(in_ptr0 + (x2), xmask)
    tmp1 = tl.load(in_ptr1 + (x0), xmask, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tmp3 = -1.0
    tmp4 = triton_helpers.maximum(tmp2, tmp3)
    tmp5 = 1.0
    tmp6 = triton_helpers.minimum(tmp4, tmp5)
    tmp7 = tmp2 <= tmp3
    tmp8 = tmp2 >= tmp5
    tmp9 = tmp7 | tmp8
    tmp10 = 3.0
    tmp11 = tmp6 * tmp10
    tmp12 = tmp11 + tmp10
    tmp13 = 0.0
    tmp14 = triton_helpers.maximum(tmp12, tmp13)
    tmp15 = 6.0
    tmp16 = triton_helpers.minimum(tmp14, tmp15)
    tmp17 = 0.16666666666666666
    tmp18 = tmp16 * tmp17
    tmp19 = tmp18 * tmp13
    tmp20 = -3.0
    tmp21 = tmp11 > tmp20
    tmp22 = tmp11 < tmp10
    tmp23 = tmp21 & tmp22
    tl.store(out_ptr0 + (x2), tmp6, xmask)
    tl.store(out_ptr1 + (x2), tmp9, xmask)
    tl.store(out_ptr2 + (x2), tmp18, xmask)
    tl.store(out_ptr3 + (x2), tmp19, xmask)
    tl.store(out_ptr4 + (x2), tmp23, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_4/inductor_cache/zd/czdeq2ohbgubcyeps2ukquvfhigxtyega57i24ketclusfgmyedi.py
# Topologically Sorted Source Nodes: [cat_1], Original ATen: [aten.cat]
# Source node to ATen node mapping:
#   cat_1 => cat_1
# Graph fragment:
#   %cat_1 : [num_users=2] = call_function[target=torch.ops.aten.cat.default](args = ([%mul_1, %primals_1], 1), kwargs = {})
triton_poi_fused_cat_2 = async_compile.triton('triton_poi_fused_cat_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=[32], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=80, major=8, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=108), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_cat_2', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 2, 'num_reduction': 0, 'backend_hash': 'A9C866B4A14FD3277824029365D703C2427B2E685E54EC9B3EF4ADC8D1EEAC1D', '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_cat_2(in_ptr0, in_ptr1, out_ptr0, xnumel, XBLOCK : tl.constexpr):
    xnumel = 32
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex % 8
    x1 = (xindex // 8)
    x2 = xindex
    tmp0 = x0
    tmp1 = tl.full([1], 0, tl.int64)
    tmp2 = tmp0 >= tmp1
    tmp3 = tl.full([1], 4, tl.int64)
    tmp4 = tmp0 < tmp3
    tmp5 = tl.load(in_ptr0 + ((4*x1) + x0), tmp4 & xmask, eviction_policy='evict_last', other=0.0)
    tmp6 = tmp0 >= tmp3
    tmp7 = tl.full([1], 8, tl.int64)
    tmp8 = tmp0 < tmp7
    tmp9 = tl.load(in_ptr1 + ((4*x1) + ((-4) + x0)), tmp6 & xmask, eviction_policy='evict_last', other=0.0)
    tmp10 = tl.where(tmp4, tmp5, tmp9)
    tl.store(out_ptr0 + (x2), tmp10, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_4/inductor_cache/rf/crfihv5y2ucyci2za4g5si66yngly3ssf24mqweq4gl5qni2ztax.py
# Topologically Sorted Source Nodes: [hardtanh_1, mul_2, sub, add, add_1, hy], Original ATen: [aten.hardtanh, aten.mul, aten.rsub, aten.add, aten.hardtanh_backward]
# Source node to ATen node mapping:
#   add => add_1
#   add_1 => add_2
#   hardtanh_1 => clamp_max_2, clamp_min_2
#   hy => clamp_max_3, clamp_min_3
#   mul_2 => mul_2
#   sub => sub
# Graph fragment:
#   %add_tensor : [num_users=3] = call_function[target=torch.ops.aten.add.Tensor](args = (%mm_default, %primals_5), kwargs = {})
#   %clamp_min_2 : [num_users=1] = call_function[target=torch.ops.aten.clamp_min.default](args = (%add_tensor, -1.0), kwargs = {})
#   %clamp_max_2 : [num_users=3] = call_function[target=torch.ops.aten.clamp_max.default](args = (%clamp_min_2, 1.0), kwargs = {})
#   %mul_2 : [num_users=2] = call_function[target=torch.ops.aten.mul.Tensor](args = (%div, %clamp_max_2), kwargs = {})
#   %sub : [num_users=2] = call_function[target=torch.ops.aten.sub.Tensor](args = (0, %mul_2), kwargs = {})
#   %add_1 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%clamp_max_2, %sub), kwargs = {})
#   %add_2 : [num_users=3] = call_function[target=torch.ops.aten.add.Tensor](args = (%add_1, %mul_1), kwargs = {})
#   %clamp_min_3 : [num_users=1] = call_function[target=torch.ops.aten.clamp_min.default](args = (%add_2, -1.0), kwargs = {})
#   %clamp_max_3 : [num_users=1] = call_function[target=torch.ops.aten.clamp_max.default](args = (%clamp_min_3, 1.0), kwargs = {})
#   %le : [num_users=1] = call_function[target=torch.ops.aten.le.Scalar](args = (%add_2, -1.0), kwargs = {})
#   %ge : [num_users=1] = call_function[target=torch.ops.aten.ge.Scalar](args = (%add_2, 1.0), kwargs = {})
#   %bitwise_or : [num_users=1] = call_function[target=torch.ops.aten.bitwise_or.Tensor](args = (%le, %ge), kwargs = {})
#   %le_1 : [num_users=1] = call_function[target=torch.ops.aten.le.Scalar](args = (%add_tensor, -1.0), kwargs = {})
#   %ge_1 : [num_users=1] = call_function[target=torch.ops.aten.ge.Scalar](args = (%add_tensor, 1.0), kwargs = {})
#   %bitwise_or_1 : [num_users=1] = call_function[target=torch.ops.aten.bitwise_or.Tensor](args = (%le_1, %ge_1), kwargs = {})
triton_poi_fused_add_hardtanh_hardtanh_backward_mul_rsub_3 = async_compile.triton('triton_poi_fused_add_hardtanh_hardtanh_backward_mul_rsub_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=[16], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: '*fp32', 5: '*i1', 6: '*fp32', 7: '*fp32', 8: '*fp32', 9: '*i1', 10: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=80, major=8, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=108), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_add_hardtanh_hardtanh_backward_mul_rsub_3', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 4, 'num_reduction': 0, 'backend_hash': 'A9C866B4A14FD3277824029365D703C2427B2E685E54EC9B3EF4ADC8D1EEAC1D', '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_hardtanh_hardtanh_backward_mul_rsub_3(in_ptr0, in_ptr1, in_ptr2, in_ptr3, out_ptr0, out_ptr1, out_ptr2, out_ptr3, out_ptr4, out_ptr5, xnumel, XBLOCK : tl.constexpr):
    xnumel = 16
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x2 = xindex
    x0 = xindex % 4
    tmp0 = tl.load(in_ptr0 + (x2), xmask)
    tmp1 = tl.load(in_ptr1 + (x0), xmask, eviction_policy='evict_last')
    tmp10 = tl.load(in_ptr2 + (x2), xmask)
    tmp15 = tl.load(in_ptr3 + (x2), xmask)
    tmp2 = tmp0 + tmp1
    tmp3 = -1.0
    tmp4 = triton_helpers.maximum(tmp2, tmp3)
    tmp5 = 1.0
    tmp6 = triton_helpers.minimum(tmp4, tmp5)
    tmp7 = tmp2 <= tmp3
    tmp8 = tmp2 >= tmp5
    tmp9 = tmp7 | tmp8
    tmp11 = tmp10 * tmp6
    tmp12 = 0.0
    tmp13 = tmp12 - tmp11
    tmp14 = tmp6 + tmp13
    tmp16 = tmp14 + tmp15
    tmp17 = triton_helpers.maximum(tmp16, tmp3)
    tmp18 = triton_helpers.minimum(tmp17, tmp5)
    tmp19 = tmp16 <= tmp3
    tmp20 = tmp16 >= tmp5
    tmp21 = tmp19 | tmp20
    tl.store(out_ptr0 + (x2), tmp6, xmask)
    tl.store(out_ptr1 + (x2), tmp9, xmask)
    tl.store(out_ptr2 + (x2), tmp11, xmask)
    tl.store(out_ptr3 + (x2), tmp13, xmask)
    tl.store(out_ptr4 + (x2), tmp18, xmask)
    tl.store(out_ptr5 + (x2), tmp21, xmask)
''', device_str='cuda')


async_compile.wait(globals())
del async_compile

def call(args):
    primals_1, primals_2, primals_3, primals_4, primals_5 = args
    args.clear()
    assert_size_stride(primals_1, (4, 4), (4, 1))
    assert_size_stride(primals_2, (4, 8), (8, 1))
    assert_size_stride(primals_3, (4, ), (1, ))
    assert_size_stride(primals_4, (4, 8), (8, 1))
    assert_size_stride(primals_5, (4, ), (1, ))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((4, 8), (8, 1), torch.float32)
        # Topologically Sorted Source Nodes: [cat], Original ATen: [aten.cat]
        stream0 = get_raw_stream(0)
        triton_poi_fused_cat_0.run(primals_1, buf0, 32, grid=grid(32), stream=stream0)
        buf1 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        # Topologically Sorted Source Nodes: [], Original ATen: []
        extern_kernels.mm(buf0, reinterpret_tensor(primals_2, (8, 4), (1, 8), 0), out=buf1)
        del primals_2
        buf2 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        buf14 = empty_strided_cuda((4, 4), (4, 1), torch.bool)
        buf3 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        buf4 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        buf13 = empty_strided_cuda((4, 4), (4, 1), torch.bool)
        # Topologically Sorted Source Nodes: [hx, hardtanh, mul, hardsigmoid, mul_1], Original ATen: [aten.zeros, aten.hardtanh, aten.mul, aten.hardsigmoid, aten.hardsigmoid_backward, aten.hardtanh_backward]
        triton_poi_fused_hardsigmoid_hardsigmoid_backward_hardtanh_hardtanh_backward_mul_zeros_1.run(buf1, primals_3, buf2, buf14, buf3, buf4, buf13, 16, grid=grid(16), stream=stream0)
        del primals_3
        buf5 = empty_strided_cuda((4, 8), (8, 1), torch.float32)
        # Topologically Sorted Source Nodes: [cat_1], Original ATen: [aten.cat]
        triton_poi_fused_cat_2.run(buf4, primals_1, buf5, 32, grid=grid(32), stream=stream0)
        del primals_1
        buf6 = buf1; del buf1  # reuse
        # Topologically Sorted Source Nodes: [], Original ATen: []
        extern_kernels.mm(buf5, reinterpret_tensor(primals_4, (8, 4), (1, 8), 0), out=buf6)
        buf7 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        buf12 = empty_strided_cuda((4, 4), (4, 1), torch.bool)
        buf8 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        buf9 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        buf10 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        buf11 = empty_strided_cuda((4, 4), (4, 1), torch.bool)
        # Topologically Sorted Source Nodes: [hardtanh_1, mul_2, sub, add, add_1, hy], Original ATen: [aten.hardtanh, aten.mul, aten.rsub, aten.add, aten.hardtanh_backward]
        triton_poi_fused_add_hardtanh_hardtanh_backward_mul_rsub_3.run(buf6, primals_5, buf3, buf4, buf7, buf12, buf8, buf9, buf10, buf11, 16, grid=grid(16), stream=stream0)
        del buf6
        del primals_5
    return (buf10, buf9, buf8, buf7, buf5, buf4, buf3, buf2, buf0, buf0, buf3, buf5, buf7, buf11, buf12, primals_4, buf13, buf14, )


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, 4), (4, 1), device='cuda:0', dtype=torch.float32)
    primals_2 = rand_strided((4, 8), (8, 1), device='cuda:0', dtype=torch.float32)
    primals_3 = rand_strided((4, ), (1, ), device='cuda:0', dtype=torch.float32)
    primals_4 = rand_strided((4, 8), (8, 1), device='cuda:0', dtype=torch.float32)
    primals_5 = rand_strided((4, ), (1, ), device='cuda:0', dtype=torch.float32)
    fn = lambda: call([primals_1, primals_2, primals_3, primals_4, primals_5])
    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 math
import torch
from torch import Tensor
import torch.nn as nn
import torch.nn.functional as F
from typing import Optional
import torch.nn.parallel
import torch.optim
import torch.utils.data
import torch.utils.data.distributed


def truncated_normal(t, mean=0.0, std=0.01):
    torch.nn.init.normal_(t, mean=mean, std=std)
    while True:
        cond = torch.logical_or(t < mean - 2 * std, t > mean + 2 * std)
        if torch.sum(cond):
            t = torch.where(cond, torch.nn.init.normal_(torch.ones_like(t),
                mean=mean, std=std), t)
        else:
            break
    return t


class HUBHardsigmoid(torch.nn.Module):
    """
    This is a hub scaled addition (x+1)/2.
    """

    def __init__(self, scale=3):
        super(HUBHardsigmoid, self).__init__()
        self.scale = scale

    def forward(self, x) ->str:
        return torch.nn.Hardsigmoid()(x * self.scale)


class HUBHardtanh(torch.nn.Hardtanh):
    """
    Inputs within range [-1, +1] directly pass through, while inputs outsides will be clipped to -1 and +1.
    This module is used for training and inference in binary domain.
    """

    def __init__(self):
        super(HUBHardtanh, self).__init__()


class HardMGUCell(torch.nn.Module):
    """
    This is a minimal gated unit by replacing sigmoid and tanh with hubhardsigmoid and hubhardtanh if hard is set to True.
    Refer to "Simplified Minimal Gated Unit Variations for Recurrent Neural Networks" and "Gate-Variants of Gated Recurrent Unit (GRU) Neural Networks" for more details.
    This module is fully unary computing aware, i.e., all intermediate data are bounded to the legal unary range.
    This module follows the uBrain implementation style to maximize hardware reuse.
    This modeule assigns batch to dim[0].
    This module applies floating-point data.
    """

    def __init__(self, input_size: 'int', hidden_size: 'int', bias: 'bool'=
        True, hard: 'bool'=True) ->None:
        super(HardMGUCell, self).__init__()
        self.input_size = input_size
        self.hidden_size = hidden_size
        self.bias = bias
        self.hard = hard
        if hard is True:
            self.fg_sigmoid = HUBHardsigmoid()
            self.ng_tanh = HUBHardtanh()
        else:
            self.fg_sigmoid = nn.Sigmoid()
            self.ng_tanh = nn.Tanh()
        self.weight_f = nn.Parameter(torch.empty((hidden_size, hidden_size +
            input_size)))
        self.weight_n = nn.Parameter(torch.empty((hidden_size, hidden_size +
            input_size)))
        if bias:
            self.bias_f = nn.Parameter(torch.empty(hidden_size))
            self.bias_n = nn.Parameter(torch.empty(hidden_size))
        else:
            self.register_parameter('bias_f', None)
            self.register_parameter('bias_n', None)
        self.reset_parameters()

    def reset_parameters(self):
        stdv = 1.0 / math.sqrt(self.hidden_size)
        for weight in self.parameters():
            weight.data = truncated_normal(weight, mean=0, std=stdv)

    def forward(self, input: 'Tensor', hx: 'Optional[Tensor]'=None) ->Tensor:
        if hx is None:
            hx = torch.zeros(input.size()[0], self.hidden_size, dtype=input
                .dtype, device=input.device)
        self.fg_ug_in = torch.cat((hx, input), 1)
        self.fg_in = HUBHardtanh()(F.linear(self.fg_ug_in, self.weight_f,
            self.bias_f))
        self.fg = self.fg_sigmoid(self.fg_in)
        self.fg_hx = self.fg * hx
        self.ng_ug_in = torch.cat((self.fg_hx, input), 1)
        self.ng = self.ng_tanh(F.linear(self.ng_ug_in, self.weight_n, self.
            bias_n))
        self.fg_ng = self.fg * self.ng
        self.fg_ng_inv = 0 - self.fg_ng
        hy = HUBHardtanh()(self.ng + self.fg_ng_inv + self.fg_hx)
        return hy


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


def get_init_inputs():
    return [[], {'input_size': 4, 'hidden_size': 4}]

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
import math
import torch.nn as nn
import torch.nn.parallel
import torch.optim
import torch.utils.data
import torch.utils.data.distributed
assert_size_stride = torch._C._dynamo.guards.assert_size_stride
empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda
reinterpret_tensor = torch._C._dynamo.guards._reinterpret_tensor


@triton.jit
def triton_poi_fused_cat_0(in_ptr0, out_ptr0, xnumel, XBLOCK: tl.constexpr):
    xnumel = 32
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex % 8
    x1 = xindex // 8
    x2 = xindex
    tmp0 = x0
    tl.full([1], 0, tl.int64)
    tmp3 = tl.full([1], 4, tl.int64)
    tmp4 = tmp0 < tmp3
    tmp5 = 0.0
    tmp6 = tl.full(tmp5.shape, 0.0, tmp5.dtype)
    tmp7 = tl.where(tmp4, tmp5, tmp6)
    tmp8 = tmp0 >= tmp3
    tl.full([1], 8, tl.int64)
    tmp11 = tl.load(in_ptr0 + (4 * x1 + (-4 + x0)), tmp8 & xmask,
        eviction_policy='evict_last', other=0.0)
    tmp12 = tl.where(tmp4, tmp7, tmp11)
    tl.store(out_ptr0 + x2, tmp12, xmask)


@triton.jit
def triton_poi_fused_hardsigmoid_hardsigmoid_backward_hardtanh_hardtanh_backward_mul_zeros_1(
    in_ptr0, in_ptr1, out_ptr0, out_ptr1, out_ptr2, out_ptr3, out_ptr4,
    xnumel, XBLOCK: tl.constexpr):
    xnumel = 16
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x2 = xindex
    x0 = xindex % 4
    tmp0 = tl.load(in_ptr0 + x2, xmask)
    tmp1 = tl.load(in_ptr1 + x0, xmask, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tmp3 = -1.0
    tmp4 = triton_helpers.maximum(tmp2, tmp3)
    tmp5 = 1.0
    tmp6 = triton_helpers.minimum(tmp4, tmp5)
    tmp7 = tmp2 <= tmp3
    tmp8 = tmp2 >= tmp5
    tmp9 = tmp7 | tmp8
    tmp10 = 3.0
    tmp11 = tmp6 * tmp10
    tmp12 = tmp11 + tmp10
    tmp13 = 0.0
    tmp14 = triton_helpers.maximum(tmp12, tmp13)
    tmp15 = 6.0
    tmp16 = triton_helpers.minimum(tmp14, tmp15)
    tmp17 = 0.16666666666666666
    tmp18 = tmp16 * tmp17
    tmp19 = tmp18 * tmp13
    tmp20 = -3.0
    tmp21 = tmp11 > tmp20
    tmp22 = tmp11 < tmp10
    tmp23 = tmp21 & tmp22
    tl.store(out_ptr0 + x2, tmp6, xmask)
    tl.store(out_ptr1 + x2, tmp9, xmask)
    tl.store(out_ptr2 + x2, tmp18, xmask)
    tl.store(out_ptr3 + x2, tmp19, xmask)
    tl.store(out_ptr4 + x2, tmp23, xmask)


@triton.jit
def triton_poi_fused_cat_2(in_ptr0, in_ptr1, out_ptr0, xnumel, XBLOCK: tl.
    constexpr):
    xnumel = 32
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex % 8
    x1 = xindex // 8
    x2 = xindex
    tmp0 = x0
    tl.full([1], 0, tl.int64)
    tmp3 = tl.full([1], 4, tl.int64)
    tmp4 = tmp0 < tmp3
    tmp5 = tl.load(in_ptr0 + (4 * x1 + x0), tmp4 & xmask, eviction_policy=
        'evict_last', other=0.0)
    tmp6 = tmp0 >= tmp3
    tl.full([1], 8, tl.int64)
    tmp9 = tl.load(in_ptr1 + (4 * x1 + (-4 + x0)), tmp6 & xmask,
        eviction_policy='evict_last', other=0.0)
    tmp10 = tl.where(tmp4, tmp5, tmp9)
    tl.store(out_ptr0 + x2, tmp10, xmask)


@triton.jit
def triton_poi_fused_add_hardtanh_hardtanh_backward_mul_rsub_3(in_ptr0,
    in_ptr1, in_ptr2, in_ptr3, out_ptr0, out_ptr1, out_ptr2, out_ptr3,
    out_ptr4, out_ptr5, xnumel, XBLOCK: tl.constexpr):
    xnumel = 16
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x2 = xindex
    x0 = xindex % 4
    tmp0 = tl.load(in_ptr0 + x2, xmask)
    tmp1 = tl.load(in_ptr1 + x0, xmask, eviction_policy='evict_last')
    tmp10 = tl.load(in_ptr2 + x2, xmask)
    tmp15 = tl.load(in_ptr3 + x2, xmask)
    tmp2 = tmp0 + tmp1
    tmp3 = -1.0
    tmp4 = triton_helpers.maximum(tmp2, tmp3)
    tmp5 = 1.0
    tmp6 = triton_helpers.minimum(tmp4, tmp5)
    tmp7 = tmp2 <= tmp3
    tmp8 = tmp2 >= tmp5
    tmp9 = tmp7 | tmp8
    tmp11 = tmp10 * tmp6
    tmp12 = 0.0
    tmp13 = tmp12 - tmp11
    tmp14 = tmp6 + tmp13
    tmp16 = tmp14 + tmp15
    tmp17 = triton_helpers.maximum(tmp16, tmp3)
    tmp18 = triton_helpers.minimum(tmp17, tmp5)
    tmp19 = tmp16 <= tmp3
    tmp20 = tmp16 >= tmp5
    tmp21 = tmp19 | tmp20
    tl.store(out_ptr0 + x2, tmp6, xmask)
    tl.store(out_ptr1 + x2, tmp9, xmask)
    tl.store(out_ptr2 + x2, tmp11, xmask)
    tl.store(out_ptr3 + x2, tmp13, xmask)
    tl.store(out_ptr4 + x2, tmp18, xmask)
    tl.store(out_ptr5 + x2, tmp21, xmask)


def call(args):
    primals_1, primals_2, primals_3, primals_4, primals_5 = args
    args.clear()
    assert_size_stride(primals_1, (4, 4), (4, 1))
    assert_size_stride(primals_2, (4, 8), (8, 1))
    assert_size_stride(primals_3, (4,), (1,))
    assert_size_stride(primals_4, (4, 8), (8, 1))
    assert_size_stride(primals_5, (4,), (1,))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((4, 8), (8, 1), torch.float32)
        get_raw_stream(0)
        triton_poi_fused_cat_0[grid(32)](primals_1, buf0, 32, XBLOCK=32,
            num_warps=1, num_stages=1)
        buf1 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        extern_kernels.mm(buf0, reinterpret_tensor(primals_2, (8, 4), (1, 8
            ), 0), out=buf1)
        del primals_2
        buf2 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        buf14 = empty_strided_cuda((4, 4), (4, 1), torch.bool)
        buf3 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        buf4 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        buf13 = empty_strided_cuda((4, 4), (4, 1), torch.bool)
        triton_poi_fused_hardsigmoid_hardsigmoid_backward_hardtanh_hardtanh_backward_mul_zeros_1[
            grid(16)](buf1, primals_3, buf2, buf14, buf3, buf4, buf13, 16,
            XBLOCK=16, num_warps=1, num_stages=1)
        del primals_3
        buf5 = empty_strided_cuda((4, 8), (8, 1), torch.float32)
        triton_poi_fused_cat_2[grid(32)](buf4, primals_1, buf5, 32, XBLOCK=
            32, num_warps=1, num_stages=1)
        del primals_1
        buf6 = buf1
        del buf1
        extern_kernels.mm(buf5, reinterpret_tensor(primals_4, (8, 4), (1, 8
            ), 0), out=buf6)
        buf7 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        buf12 = empty_strided_cuda((4, 4), (4, 1), torch.bool)
        buf8 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        buf9 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        buf10 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        buf11 = empty_strided_cuda((4, 4), (4, 1), torch.bool)
        triton_poi_fused_add_hardtanh_hardtanh_backward_mul_rsub_3[grid(16)](
            buf6, primals_5, buf3, buf4, buf7, buf12, buf8, buf9, buf10,
            buf11, 16, XBLOCK=16, num_warps=1, num_stages=1)
        del buf6
        del primals_5
    return (buf10, buf9, buf8, buf7, buf5, buf4, buf3, buf2, buf0, buf0,
        buf3, buf5, buf7, buf11, buf12, primals_4, buf13, buf14)


def truncated_normal(t, mean=0.0, std=0.01):
    torch.nn.init.normal_(t, mean=mean, std=std)
    while True:
        cond = torch.logical_or(t < mean - 2 * std, t > mean + 2 * std)
        if torch.sum(cond):
            t = torch.where(cond, torch.nn.init.normal_(torch.ones_like(t),
                mean=mean, std=std), t)
        else:
            break
    return t


class HUBHardsigmoid(torch.nn.Module):
    """
    This is a hub scaled addition (x+1)/2.
    """

    def __init__(self, scale=3):
        super(HUBHardsigmoid, self).__init__()
        self.scale = scale

    def forward(self, x) ->str:
        return torch.nn.Hardsigmoid()(x * self.scale)


class HUBHardtanh(torch.nn.Hardtanh):
    """
    Inputs within range [-1, +1] directly pass through, while inputs outsides will be clipped to -1 and +1.
    This module is used for training and inference in binary domain.
    """

    def __init__(self):
        super(HUBHardtanh, self).__init__()


class HardMGUCellNew(torch.nn.Module):
    """
    This is a minimal gated unit by replacing sigmoid and tanh with hubhardsigmoid and hubhardtanh if hard is set to True.
    Refer to "Simplified Minimal Gated Unit Variations for Recurrent Neural Networks" and "Gate-Variants of Gated Recurrent Unit (GRU) Neural Networks" for more details.
    This module is fully unary computing aware, i.e., all intermediate data are bounded to the legal unary range.
    This module follows the uBrain implementation style to maximize hardware reuse.
    This modeule assigns batch to dim[0].
    This module applies floating-point data.
    """

    def __init__(self, input_size: 'int', hidden_size: 'int', bias: 'bool'=
        True, hard: 'bool'=True) ->None:
        super(HardMGUCellNew, self).__init__()
        self.input_size = input_size
        self.hidden_size = hidden_size
        self.bias = bias
        self.hard = hard
        if hard is True:
            self.fg_sigmoid = HUBHardsigmoid()
            self.ng_tanh = HUBHardtanh()
        else:
            self.fg_sigmoid = nn.Sigmoid()
            self.ng_tanh = nn.Tanh()
        self.weight_f = nn.Parameter(torch.empty((hidden_size, hidden_size +
            input_size)))
        self.weight_n = nn.Parameter(torch.empty((hidden_size, hidden_size +
            input_size)))
        if bias:
            self.bias_f = nn.Parameter(torch.empty(hidden_size))
            self.bias_n = nn.Parameter(torch.empty(hidden_size))
        else:
            self.register_parameter('bias_f', None)
            self.register_parameter('bias_n', None)
        self.reset_parameters()

    def reset_parameters(self):
        stdv = 1.0 / math.sqrt(self.hidden_size)
        for weight in self.parameters():
            weight.data = truncated_normal(weight, mean=0, std=stdv)

    def forward(self, input_0):
        primals_2 = self.weight_f
        primals_4 = self.weight_n
        primals_3 = self.bias_f
        primals_5 = self.bias_n
        primals_1 = input_0
        output = call([primals_1, primals_2, primals_3, primals_4, primals_5])
        return output[0]
