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
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_relu_threshold_backward_0(in_ptr0, out_ptr0, out_ptr1,
    xnumel, XBLOCK: tl.constexpr):
    xnumel = 16
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + x0, xmask)
    tmp1 = tl.full([1], 0, tl.int32)
    tmp2 = triton_helpers.maximum(tmp1, tmp0)
    tmp3 = 0.0
    tmp4 = tmp2 <= tmp3
    tl.store(out_ptr0 + x0, tmp2, xmask)
    tl.store(out_ptr1 + x0, tmp4, xmask)


@triton.jit
def triton_poi_fused_add_mv_new_zeros_1(in_ptr0, in_ptr1, out_ptr0, xnumel,
    XBLOCK: tl.constexpr):
    xnumel = 4
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + x0, xmask)
    tmp1 = tl.load(in_ptr1 + 4 * x0, xmask, eviction_policy='evict_last')
    tmp4 = tl.load(in_ptr1 + (1 + 4 * x0), xmask, eviction_policy='evict_last')
    tmp7 = tl.load(in_ptr1 + (2 + 4 * x0), xmask, eviction_policy='evict_last')
    tmp10 = tl.load(in_ptr1 + (3 + 4 * x0), xmask, eviction_policy='evict_last'
        )
    tmp2 = 0.0
    tmp3 = tmp1 * tmp2
    tmp5 = tmp4 * tmp2
    tmp6 = tmp3 + tmp5
    tmp8 = tmp7 * tmp2
    tmp9 = tmp6 + tmp8
    tmp11 = tmp10 * tmp2
    tmp12 = tmp9 + tmp11
    tmp13 = tmp0 + tmp12
    tl.store(out_ptr0 + x0, tmp13, xmask)


@triton.jit
def triton_poi_fused_add_mul_mv_new_zeros_relu_rsub_sigmoid_threshold_backward_2(
    in_out_ptr0, in_ptr0, in_ptr1, in_ptr2, out_ptr0, out_ptr1, xnumel,
    XBLOCK: tl.constexpr):
    xnumel = 16
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x1 = xindex // 4
    x2 = xindex
    tmp0 = tl.load(in_ptr0 + x1, xmask, eviction_policy='evict_last')
    tmp1 = tl.load(in_out_ptr0 + x2, xmask)
    tmp3 = tl.load(in_ptr1 + x1, xmask, eviction_policy='evict_last')
    tmp5 = tl.load(in_ptr2 + 0)
    tmp6 = tl.broadcast_to(tmp5, [XBLOCK])
    tmp2 = tmp0 + tmp1
    tmp4 = tmp2 + tmp3
    tmp7 = tl.sigmoid(tmp6)
    tmp8 = 0.0
    tmp9 = tmp7 * tmp8
    tmp10 = 1.0
    tmp11 = tmp10 - tmp7
    tmp12 = tmp11 * tmp4
    tmp13 = tmp9 + tmp12
    tmp14 = tl.full([1], 0, tl.int32)
    tmp15 = triton_helpers.maximum(tmp14, tmp13)
    tmp16 = tmp15 <= tmp8
    tl.store(in_out_ptr0 + x2, tmp4, xmask)
    tl.store(out_ptr0 + x2, tmp15, xmask)
    tl.store(out_ptr1 + x2, tmp16, xmask)


def call(args):
    (primals_1, primals_2, primals_3, primals_4, primals_5, primals_6,
        primals_7) = args
    args.clear()
    assert_size_stride(primals_1, (4, 4), (4, 1))
    assert_size_stride(primals_2, (4, 4), (4, 1))
    assert_size_stride(primals_3, (4, 4), (4, 1))
    assert_size_stride(primals_4, (4,), (1,))
    assert_size_stride(primals_5, (4, 4), (4, 1))
    assert_size_stride(primals_6, (4,), (1,))
    assert_size_stride(primals_7, (), ())
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        buf7 = empty_strided_cuda((4, 4), (4, 1), torch.bool)
        get_raw_stream(0)
        triton_poi_fused_relu_threshold_backward_0[grid(16)](primals_2,
            buf0, buf7, 16, XBLOCK=16, num_warps=1, num_stages=1)
        del primals_2
        buf1 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        extern_kernels.mm(buf0, primals_3, out=buf1)
        buf2 = buf0
        del buf0
        extern_kernels.mm(primals_5, primals_1, out=buf2)
        del primals_5
        buf3 = empty_strided_cuda((4,), (1,), torch.float32)
        triton_poi_fused_add_mv_new_zeros_1[grid(4)](primals_6, buf1, buf3,
            4, XBLOCK=4, num_warps=1, num_stages=1)
        del primals_6
        buf4 = reinterpret_tensor(buf2, (4, 4), (1, 4), 0)
        del buf2
        buf5 = reinterpret_tensor(buf1, (4, 4), (1, 4), 0)
        del buf1
        buf6 = empty_strided_cuda((4, 4), (1, 4), torch.bool)
        triton_poi_fused_add_mul_mv_new_zeros_relu_rsub_sigmoid_threshold_backward_2[
            grid(16)](buf4, primals_4, buf3, primals_7, buf5, buf6, 16,
            XBLOCK=16, num_warps=1, num_stages=1)
        del buf3
        del primals_4
    return buf5, primals_7, buf4, buf6, reinterpret_tensor(primals_1, (4, 4
        ), (1, 4), 0), reinterpret_tensor(primals_3, (4, 4), (1, 4), 0), buf7


class DecayModuleNew(nn.Module):

    def __init__(self, input_size, hidden_size, bias=True, num_chunks=1,
        activation='relu', nodiag=False):
        super(DecayModuleNew, self).__init__()
        self.sigmoid = nn.Sigmoid()
        self.tanh = nn.Tanh()
        self.relu = nn.ReLU()
        self.input_size = input_size
        self.hidden_size = hidden_size
        self.nodiag = nodiag
        self.bias = bias
        self.num_chunks = num_chunks
        self.rgate = nn.Parameter(torch.tensor(0.8), requires_grad=True)
        self.weight_ih = nn.Parameter(torch.Tensor(num_chunks * hidden_size,
            input_size))
        self.weight_hh = nn.Parameter(torch.Tensor(num_chunks * hidden_size,
            hidden_size))
        self.d_rec = nn.Parameter(torch.zeros(num_chunks * hidden_size,
            hidden_size), requires_grad=False)
        self.activation = activation
        if bias:
            self.bias_ih = nn.Parameter(torch.Tensor(num_chunks * hidden_size))
            self.bias_hh = nn.Parameter(torch.Tensor(num_chunks * hidden_size))
        else:
            self.register_parameter('bias_ih', None)
            self.register_parameter('bias_hh', None)
        self.reset_parameters()
        if self.nodiag:
            for i in range(hidden_size):
                self.weight_hh.data[i, i] = 0

    def reset_parameters(self):
        stdv = 1.0 / math.sqrt(self.hidden_size)
        for weight in self.parameters():
            nn.init.uniform_(weight, -stdv, stdv)
        for name, param in self.named_parameters():
            if name == 'rgate':
                param.data = torch.tensor(1.4)
        for i in range(self.num_chunks):
            x = i * self.hidden_size
            for j in range(self.hidden_size):
                if j < 0.8 * self.hidden_size:
                    self.d_rec[x + j][j] = 1.0
                else:
                    self.d_rec[x + j][j] = -1.0

    def forward(self, input_0):
        primals_7 = self.rgate
        primals_1 = self.weight_ih
        primals_2 = self.weight_hh
        primals_3 = self.d_rec
        primals_4 = self.bias_ih
        primals_6 = self.bias_hh
        primals_5 = input_0
        output = call([primals_1, primals_2, primals_3, primals_4,
            primals_5, primals_6, primals_7])
        return output[0]


Measured latency_ms: 0.12541