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
import torch.onnx
from itertools import product as product
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_add_mul_rsub_sigmoid_tanh_0(in_out_ptr0, in_out_ptr1,
    in_ptr0, in_ptr1, in_ptr2, in_ptr3, in_ptr4, out_ptr0, xnumel, XBLOCK:
    tl.constexpr):
    xnumel = 256
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x2 = xindex
    x0 = xindex % 4
    tmp0 = tl.load(in_out_ptr0 + x2, xmask)
    tmp1 = tl.load(in_ptr0 + x2, xmask)
    tmp3 = tl.load(in_ptr1 + x0, xmask, eviction_policy='evict_last')
    tmp6 = tl.load(in_out_ptr1 + x2, xmask)
    tmp7 = tl.load(in_ptr2 + x2, xmask)
    tmp9 = tl.load(in_ptr3 + x0, xmask, eviction_policy='evict_last')
    tmp12 = tl.load(in_ptr4 + x2, xmask)
    tmp2 = tmp0 + tmp1
    tmp4 = tmp2 + tmp3
    tmp5 = tl.sigmoid(tmp4)
    tmp8 = tmp6 + tmp7
    tmp10 = tmp8 + tmp9
    tmp11 = libdevice.tanh(tmp10)
    tmp13 = tmp5 * tmp12
    tmp14 = 1.0
    tmp15 = tmp14 - tmp5
    tmp16 = tmp15 * tmp11
    tmp17 = tmp13 + tmp16
    tl.store(in_out_ptr0 + x2, tmp5, xmask)
    tl.store(in_out_ptr1 + x2, tmp11, xmask)
    tl.store(out_ptr0 + x2, tmp17, xmask)


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


def gen_nonlinearity(A, nonlinearity):
    """
    Returns required activation for a tensor based on the inputs

    nonlinearity is either a callable or a value in
        ['tanh', 'sigmoid', 'relu', 'quantTanh', 'quantSigm', 'quantSigm4']
    """
    if nonlinearity == 'tanh':
        return torch.tanh(A)
    elif nonlinearity == 'sigmoid':
        return torch.sigmoid(A)
    elif nonlinearity == 'relu':
        return torch.relu(A, 0.0)
    elif nonlinearity == 'quantTanh':
        return torch.max(torch.min(A, torch.ones_like(A)), -1.0 * torch.
            ones_like(A))
    elif nonlinearity == 'quantSigm':
        A = (A + 1.0) / 2.0
        return torch.max(torch.min(A, torch.ones_like(A)), torch.zeros_like(A))
    elif nonlinearity == 'quantSigm4':
        A = (A + 2.0) / 4.0
        return torch.max(torch.min(A, torch.ones_like(A)), torch.zeros_like(A))
    else:
        if not callable(nonlinearity):
            raise ValueError(
                'nonlinearity is either a callable or a value ' +
                "['tanh', 'sigmoid', 'relu', 'quantTanh', " + "'quantSigm'")
        return nonlinearity(A)


class RNNCell(nn.Module):

    def __init__(self, input_size, hidden_size, gate_nonlinearity,
        update_nonlinearity, num_W_matrices, num_U_matrices, num_biases,
        wRank=None, uRank=None, wSparsity=1.0, uSparsity=1.0):
        super(RNNCell, self).__init__()
        self._input_size = input_size
        self._hidden_size = hidden_size
        self._gate_nonlinearity = gate_nonlinearity
        self._update_nonlinearity = update_nonlinearity
        self._num_W_matrices = num_W_matrices
        self._num_U_matrices = num_U_matrices
        self._num_biases = num_biases
        self._num_weight_matrices = [self._num_W_matrices, self.
            _num_U_matrices, self._num_biases]
        self._wRank = wRank
        self._uRank = uRank
        self._wSparsity = wSparsity
        self._uSparsity = uSparsity
        self.oldmats = []

    @property
    def state_size(self):
        return self._hidden_size

    @property
    def input_size(self):
        return self._input_size

    @property
    def output_size(self):
        return self._hidden_size

    @property
    def gate_nonlinearity(self):
        return self._gate_nonlinearity

    @property
    def update_nonlinearity(self):
        return self._update_nonlinearity

    @property
    def wRank(self):
        return self._wRank

    @property
    def uRank(self):
        return self._uRank

    @property
    def num_W_matrices(self):
        return self._num_W_matrices

    @property
    def num_U_matrices(self):
        return self._num_U_matrices

    @property
    def num_weight_matrices(self):
        return self._num_weight_matrices

    @property
    def name(self):
        raise NotImplementedError()

    def forward(self, input, state):
        raise NotImplementedError()

    def getVars(self):
        raise NotImplementedError()

    def get_model_size(self):
        """
        Function to get aimed model size
        """
        mats = self.getVars()
        endW = self._num_W_matrices
        endU = endW + self._num_U_matrices
        totalnnz = 2
        for i in range(0, endW):
            mats[i].device
            totalnnz += utils.countNNZ(mats[i].cpu(), self._wSparsity)
            mats[i]
        for i in range(endW, endU):
            mats[i].device
            totalnnz += utils.countNNZ(mats[i].cpu(), self._uSparsity)
            mats[i]
        for i in range(endU, len(mats)):
            mats[i].device
            totalnnz += utils.countNNZ(mats[i].cpu(), False)
            mats[i]
        return totalnnz * 4

    def copy_previous_UW(self):
        mats = self.getVars()
        num_mats = self._num_W_matrices + self._num_U_matrices
        if len(self.oldmats) != num_mats:
            for i in range(num_mats):
                self.oldmats.append(torch.FloatTensor())
        for i in range(num_mats):
            self.oldmats[i] = torch.FloatTensor(mats[i].detach().clone())

    def sparsify(self):
        mats = self.getVars()
        endW = self._num_W_matrices
        endU = endW + self._num_U_matrices
        for i in range(0, endW):
            mats[i] = utils.hardThreshold(mats[i], self._wSparsity)
        for i in range(endW, endU):
            mats[i] = utils.hardThreshold(mats[i], self._uSparsity)
        self.copy_previous_UW()

    def sparsifyWithSupport(self):
        mats = self.getVars()
        endU = self._num_W_matrices + self._num_U_matrices
        for i in range(0, endU):
            mats[i] = utils.supportBasedThreshold(mats[i], self.oldmats[i])


class UGRNNLRCellNew(RNNCell):
    """
    UGRNN LR Cell with Both Full Rank and Low Rank Formulations
    Has multiple activation functions for the gates
    hidden_size = # hidden units

    gate_nonlinearity = nonlinearity for the gate can be chosen from
    [tanh, sigmoid, relu, quantTanh, quantSigm]
    update_nonlinearity = nonlinearity for final rnn update
    can be chosen from [tanh, sigmoid, relu, quantTanh, quantSigm]

    wRank = rank of W matrix
    (creates 3 matrices if not None else creates 2 matrices)
    uRank = rank of U matrix
    (creates 3 matrices if not None else creates 2 matrices)

    UGRNN architecture and compression techniques are found in
    UGRNN(LINK) paper

    Basic architecture is like:

    z_t = gate_nl(W1x_t + U1h_{t-1} + B_g)
    h_t^ = update_nl(W1x_t + U1h_{t-1} + B_h)
    h_t = z_t*h_{t-1} + (1-z_t)*h_t^

    Wi and Ui can further parameterised into low rank version by
    Wi = matmul(W, W_i) and Ui = matmul(U, U_i)
    """

    def __init__(self, input_size, hidden_size, gate_nonlinearity='sigmoid',
        update_nonlinearity='tanh', wRank=None, uRank=None, wSparsity=1.0,
        uSparsity=1.0, name='UGRNNLR'):
        super(UGRNNLRCellNew, self).__init__(input_size, hidden_size,
            gate_nonlinearity, update_nonlinearity, 2, 2, 2, wRank, uRank,
            wSparsity, uSparsity)
        if wRank is not None:
            self._num_W_matrices += 1
            self._num_weight_matrices[0] = self._num_W_matrices
        if uRank is not None:
            self._num_U_matrices += 1
            self._num_weight_matrices[1] = self._num_U_matrices
        self._name = name
        if wRank is None:
            self.W1 = nn.Parameter(0.1 * torch.randn([input_size, hidden_size])
                )
            self.W2 = nn.Parameter(0.1 * torch.randn([input_size, hidden_size])
                )
        else:
            self.W = nn.Parameter(0.1 * torch.randn([input_size, wRank]))
            self.W1 = nn.Parameter(0.1 * torch.randn([wRank, hidden_size]))
            self.W2 = nn.Parameter(0.1 * torch.randn([wRank, hidden_size]))
        if uRank is None:
            self.U1 = nn.Parameter(0.1 * torch.randn([hidden_size,
                hidden_size]))
            self.U2 = nn.Parameter(0.1 * torch.randn([hidden_size,
                hidden_size]))
        else:
            self.U = nn.Parameter(0.1 * torch.randn([hidden_size, uRank]))
            self.U1 = nn.Parameter(0.1 * torch.randn([uRank, hidden_size]))
            self.U2 = nn.Parameter(0.1 * torch.randn([uRank, hidden_size]))
        self.bias_gate = nn.Parameter(torch.ones([1, hidden_size]))
        self.bias_update = nn.Parameter(torch.ones([1, hidden_size]))
        self._device = self.bias_update.device

    @property
    def name(self):
        return self._name

    @property
    def cellType(self):
        return 'UGRNNLR'

    def getVars(self):
        Vars = []
        if self._num_W_matrices == 2:
            Vars.extend([self.W1, self.W2])
        else:
            Vars.extend([self.W, self.W1, self.W2])
        if self._num_U_matrices == 2:
            Vars.extend([self.U1, self.U2])
        else:
            Vars.extend([self.U, self.U1, self.U2])
        Vars.extend([self.bias_gate, self.bias_update])
        return Vars

    def forward(self, input_0, input_1):
        primals_1 = self.W1
        primals_3 = self.W2
        primals_4 = self.U1
        primals_6 = self.U2
        primals_7 = self.bias_gate
        primals_8 = self.bias_update
        primals_2 = input_0
        primals_5 = input_1
        output = call([primals_1, primals_2, primals_3, primals_4,
            primals_5, primals_6, primals_7, primals_8])
        return output[0]


Measured latency_ms: 0.037385