import torch
from torch._inductor.select_algorithm import extern_kernels
import triton
import triton.language as tl
from torch._inductor.runtime.triton_heuristics import grid
from torch._C import _cuda_getCurrentRawStream as get_raw_stream
from torch._inductor.runtime.triton_helpers import libdevice
import torch.nn as nn
from torch.optim import *
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_mul_sigmoid_tanh_0(in_ptr0, in_ptr1, out_ptr0,
    out_ptr1, out_ptr2, xnumel, XBLOCK: tl.constexpr):
    xnumel = 64
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex % 4
    x1 = xindex // 4
    x2 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + 12 * x1), xmask)
    tmp1 = tl.load(in_ptr1 + x0, xmask, eviction_policy='evict_last')
    tmp4 = tl.load(in_ptr0 + (4 + x0 + 12 * x1), xmask)
    tmp5 = tl.load(in_ptr1 + (4 + x0), xmask, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tmp3 = libdevice.tanh(tmp2)
    tmp6 = tmp4 + tmp5
    tmp7 = tl.sigmoid(tmp6)
    tmp8 = tmp7 * tmp3
    tl.store(out_ptr0 + x2, tmp3, xmask)
    tl.store(out_ptr1 + x2, tmp7, xmask)
    tl.store(out_ptr2 + x2, tmp8, xmask)


@triton.jit
def triton_poi_fused_add_mul_rsub_1(in_ptr0, in_ptr1, out_ptr0, out_ptr1,
    out_ptr2, out_ptr3, 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 + (16 + x0), xmask)
    tmp3 = tl.load(in_ptr1 + (16 + x0), xmask)
    tmp4 = tl.load(in_ptr1 + x0, xmask)
    tmp7 = tl.load(in_ptr0 + (32 + x0), xmask)
    tmp9 = tl.load(in_ptr1 + (32 + x0), xmask)
    tmp1 = 1.0
    tmp2 = tmp1 - tmp0
    tmp5 = tmp2 * tmp4
    tmp6 = tmp3 + tmp5
    tmp8 = tmp1 - tmp7
    tmp10 = tmp8 * tmp6
    tmp11 = tmp9 + tmp10
    tl.store(out_ptr0 + x0, tmp2, xmask)
    tl.store(out_ptr1 + x0, tmp6, xmask)
    tl.store(out_ptr2 + x0, tmp8, xmask)
    tl.store(out_ptr3 + x0, tmp11, xmask)


@triton.jit
def triton_poi_fused_rsub_2(in_ptr0, out_ptr0, 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 + (48 + x0), xmask)
    tmp1 = 1.0
    tmp2 = tmp1 - tmp0
    tl.store(out_ptr0 + x0, tmp2, xmask)


@triton.jit
def triton_poi_fused_mul_sigmoid_stack_3(in_ptr0, in_ptr1, in_ptr2, in_ptr3,
    in_ptr4, in_ptr5, out_ptr0, out_ptr1, out_ptr2, xnumel, XBLOCK: tl.
    constexpr):
    xnumel = 64
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x1 = xindex // 4
    x0 = xindex % 4
    x2 = xindex
    tmp29 = tl.load(in_ptr4 + (8 + x0 + 12 * x1), xmask)
    tmp30 = tl.load(in_ptr5 + (8 + x0), xmask, eviction_policy='evict_last')
    tmp0 = x1
    tl.full([1], 0, tl.int64)
    tmp3 = tl.full([1], 4, tl.int64)
    tmp4 = tmp0 < tmp3
    tmp5 = tl.load(in_ptr0 + (x0 + 4 * x1), tmp4 & xmask, other=0.0)
    tmp6 = tmp0 >= tmp3
    tmp7 = tl.full([1], 8, tl.int64)
    tmp8 = tmp0 < tmp7
    tmp9 = tmp6 & tmp8
    tmp10 = tl.load(in_ptr1 + (x0 + 4 * (-4 + x1)), tmp9 & xmask, other=0.0)
    tmp11 = tmp0 >= tmp7
    tmp12 = tl.full([1], 12, tl.int64)
    tmp13 = tmp0 < tmp12
    tmp14 = tmp11 & tmp13
    tmp15 = tl.load(in_ptr2 + (x0 + 4 * (-8 + x1)), tmp14 & xmask, other=0.0)
    tmp16 = tmp0 >= tmp12
    tl.full([1], 16, tl.int64)
    tmp19 = tl.load(in_ptr0 + (48 + x0 + 4 * (-12 + x1)), tmp16 & xmask,
        other=0.0)
    tmp20 = tl.load(in_ptr3 + (x0 + 4 * (-12 + x1)), tmp16 & xmask, other=0.0)
    tmp21 = tl.load(in_ptr2 + (x0 + 4 * (-12 + x1)), tmp16 & xmask, other=0.0)
    tmp22 = tmp20 * tmp21
    tmp23 = tmp19 + tmp22
    tmp24 = tl.full(tmp23.shape, 0.0, tmp23.dtype)
    tmp25 = tl.where(tmp16, tmp23, tmp24)
    tmp26 = tl.where(tmp14, tmp15, tmp25)
    tmp27 = tl.where(tmp9, tmp10, tmp26)
    tmp28 = tl.where(tmp4, tmp5, tmp27)
    tmp31 = tmp29 + tmp30
    tmp32 = tl.sigmoid(tmp31)
    tmp33 = tmp32 * tmp28
    tl.store(out_ptr0 + x2, tmp28, xmask)
    tl.store(out_ptr1 + x2, tmp32, xmask)
    tl.store(out_ptr2 + x2, tmp33, xmask)


def call(args):
    primals_1, primals_2, primals_3 = args
    args.clear()
    assert_size_stride(primals_1, (4, 4, 4), (16, 4, 1))
    assert_size_stride(primals_2, (12, 4), (4, 1))
    assert_size_stride(primals_3, (12,), (1,))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((16, 12), (12, 1), torch.float32)
        extern_kernels.mm(reinterpret_tensor(primals_1, (16, 4), (4, 1), 0),
            reinterpret_tensor(primals_2, (4, 12), (1, 4), 0), out=buf0)
        del primals_2
        buf1 = empty_strided_cuda((4, 4, 4), (16, 4, 1), torch.float32)
        buf2 = empty_strided_cuda((4, 4, 4), (16, 4, 1), torch.float32)
        buf3 = empty_strided_cuda((4, 4, 4), (16, 4, 1), torch.float32)
        get_raw_stream(0)
        triton_poi_fused_mul_sigmoid_tanh_0[grid(64)](buf0, primals_3, buf1,
            buf2, buf3, 64, XBLOCK=64, num_warps=1, num_stages=1)
        buf4 = empty_strided_cuda((1, 4, 4), (16, 4, 1), torch.float32)
        buf5 = empty_strided_cuda((1, 4, 4), (16, 4, 1), torch.float32)
        buf6 = empty_strided_cuda((1, 4, 4), (16, 4, 1), torch.float32)
        buf7 = empty_strided_cuda((1, 4, 4), (16, 4, 1), torch.float32)
        triton_poi_fused_add_mul_rsub_1[grid(16)](buf2, buf3, buf4, buf5,
            buf6, buf7, 16, XBLOCK=16, num_warps=1, num_stages=1)
        buf8 = empty_strided_cuda((1, 4, 4), (16, 4, 1), torch.float32)
        triton_poi_fused_rsub_2[grid(16)](buf2, buf8, 16, XBLOCK=16,
            num_warps=1, num_stages=1)
        buf9 = empty_strided_cuda((16, 4), (4, 1), torch.float32)
        buf10 = empty_strided_cuda((4, 4, 4), (16, 4, 1), torch.float32)
        buf11 = empty_strided_cuda((4, 4, 4), (16, 4, 1), torch.float32)
        triton_poi_fused_mul_sigmoid_stack_3[grid(64)](buf3, buf5, buf7,
            buf8, buf0, primals_3, buf9, buf10, buf11, 64, XBLOCK=64,
            num_warps=1, num_stages=1)
        del buf0
        del primals_3
    return buf11, reinterpret_tensor(buf9, (1, 4, 4), (16, 4, 1), 48
        ), reinterpret_tensor(primals_1, (16, 4), (4, 1), 0
        ), buf1, buf2, reinterpret_tensor(buf3, (4, 4), (4, 1), 0
        ), buf4, reinterpret_tensor(buf5, (4, 4), (4, 1), 0
        ), buf6, reinterpret_tensor(buf7, (4, 4), (4, 1), 0), buf8, buf9, buf10


class ForgetMult(torch.nn.Module):
    """ForgetMult computes a simple recurrent equation:
    h_t = f_t * x_t + (1 - f_t) * h_{t-1}

    This equation is equivalent to dynamic weighted averaging.

    Inputs: X, hidden
        - X (seq_len, batch, input_size): tensor containing the features of the input sequence.
        - F (seq_len, batch, input_size): tensor containing the forget gate values, assumed in range [0, 1].
        - hidden_init (batch, input_size): tensor containing the initial hidden state for the recurrence (h_{t-1}).
    """

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

    def forward(self, f, x, hidden_init=None):
        result = []
        forgets = f.split(1, dim=0)
        prev_h = hidden_init
        for i, h in enumerate((f * x).split(1, dim=0)):
            if prev_h is not None:
                h = h + (1 - forgets[i]) * prev_h
            h = h.view(h.size()[1:])
            result.append(h)
            prev_h = h
        return torch.stack(result)


class QRNNLayerNew(nn.Module):
    """Applies a single layer Quasi-Recurrent Neural Network (QRNN) to an input sequence.

    Args:
        input_size: The number of expected features in the input x.
        hidden_size: The number of features in the hidden state h. If not specified, the input size is used.
        save_prev_x: Whether to store previous inputs for use in future convolutional windows (i.e. for a continuing sequence such as in language modeling). If true, you must call reset to remove cached previous values of x. Default: False.
        window: Defines the size of the convolutional window (how many previous tokens to look when computing the QRNN values). Supports 1 and 2. Default: 1.
        zoneout: Whether to apply zoneout (i.e. failing to update elements in the hidden state) to the hidden state updates. Default: 0.
        output_gate: If True, performs QRNN-fo (applying an output gate to the output). If False, performs QRNN-f. Default: True.

    Inputs: X, hidden
        - X (seq_len, batch, input_size): tensor containing the features of the input sequence.
        - hidden (batch, hidden_size): tensor containing the initial hidden state for the QRNN.

    Outputs: output, h_n
        - output (seq_len, batch, hidden_size): tensor containing the output of the QRNN for each timestep.
        - h_n (1, batch, hidden_size): tensor containing the hidden state for t=seq_len
    """

    def __init__(self, input_size, hidden_size=None, save_prev_x=False,
        zoneout=0, window=1, output_gate=True):
        super(QRNNLayerNew, self).__init__()
        assert window in [1, 2
            ], 'This QRNN implementation currently only handles convolutional window of size 1 or size 2'
        self.window = window
        self.input_size = input_size
        self.hidden_size = hidden_size if hidden_size else input_size
        self.zoneout = zoneout
        self.save_prev_x = save_prev_x
        self.prevX = None
        self.output_gate = output_gate
        self.linear = nn.Linear(self.window * self.input_size, 3 * self.
            hidden_size if self.output_gate else 2 * self.hidden_size)

    def reset(self):
        self.prevX = None

    def forward(self, input_0):
        primals_2 = self.linear.weight
        primals_3 = self.linear.bias
        primals_1 = input_0
        output = call([primals_1, primals_2, primals_3])
        return output[0], output[1]


Measured latency_ms: 0.019826