import torch
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 numpy as np
from torch import nn
from torch.nn.parameter import Parameter
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_arange_0(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 = 4 * x0
    tl.store(out_ptr0 + x0, tmp0, xmask)


@triton.jit
def triton_poi_fused__to_copy_div_remainder_1(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 = 0.25
    tmp2 = tmp0 * tmp1
    tmp3 = tmp2.to(tl.int64)
    tmp4 = 4.0
    tmp5 = tmp0 % tmp4
    tmp6 = tl.full([1], 0, tl.int32)
    tmp7 = tmp5 != tmp6
    tmp8 = libdevice.signbit(tmp5) if tmp5.dtype is tl.float32 else tmp5 < 0
    tmp9 = libdevice.signbit(tmp4) if tmp4.dtype is tl.float32 else tmp4 < 0
    tmp10 = tmp8 != tmp9
    tmp11 = tmp7 & tmp10
    tmp12 = tmp5 + tmp4
    tmp13 = tl.where(tmp11, tmp12, tmp5)
    tmp14 = tmp13.to(tl.int64)
    tl.store(out_ptr0 + x0, tmp3, xmask)
    tl.store(out_ptr1 + x0, tmp14, xmask)


@triton.jit
def triton_poi_fused_mul_2(in_ptr0, in_ptr1, 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 + x0, xmask)
    tmp1 = tl.load(in_ptr1 + x0, xmask)
    tmp2 = tmp0 * tmp1
    tl.store(out_ptr0 + x0, tmp2, xmask)


def call(args):
    primals_1, primals_2, primals_3 = args
    args.clear()
    assert_size_stride(primals_1, (4, 4), (4, 1))
    assert_size_stride(primals_2, (1, 4), (4, 1))
    assert_size_stride(primals_3, (4, 4), (4, 1))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((4,), (1,), torch.int64)
        get_raw_stream(0)
        triton_poi_fused_arange_0[grid(4)](buf0, 4, XBLOCK=4, num_warps=1,
            num_stages=1)
        buf1 = empty_strided_cuda((4, 4), (4, 1), torch.int64)
        buf7 = empty_strided_cuda((4, 4), (4, 1), torch.int64)
        triton_poi_fused__to_copy_div_remainder_1[grid(16)](primals_1, buf1,
            buf7, 16, XBLOCK=16, num_warps=1, num_stages=1)
        del primals_1
        buf2 = torch.ops.aten._embedding_bag.default(primals_2,
            reinterpret_tensor(buf1, (16,), (1,), 0), buf0, False, 1)
        del primals_2
        buf3 = buf2[0]
        buf4 = buf2[1]
        buf5 = buf2[2]
        buf6 = buf2[3]
        del buf2
        buf8 = torch.ops.aten._embedding_bag.default(primals_3,
            reinterpret_tensor(buf7, (16,), (1,), 0), buf0, False, 1)
        del primals_3
        buf9 = buf8[0]
        buf10 = buf8[1]
        buf11 = buf8[2]
        buf12 = buf8[3]
        del buf8
        buf13 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        triton_poi_fused_mul_2[grid(16)](buf3, buf9, buf13, 16, XBLOCK=16,
            num_warps=1, num_stages=1)
    return buf13, buf0, reinterpret_tensor(buf1, (16,), (1,), 0
        ), buf3, buf4, buf5, buf6, reinterpret_tensor(buf7, (16,), (1,), 0
        ), buf9, buf10, buf11, buf12


class QREmbeddingBagNew(nn.Module):
    """Computes sums or means over two 'bags' of embeddings, one
    using the quotient of the indices and the other using the remainder
    of the indices, without instantiating the intermediate embeddings,
    then performsan operation to combine these.

    For bags of constant length and no :attr:`per_sample_weights`, this class

        * with ``mode="sum"`` is equivalent to :class:`~torch.nn.Embedding` followed by ``sum(dim=1)``,
        * with ``mode="mean"`` is equivalent to :class:`~torch.nn.Embedding` followed by ``torch.mean(dim=1)``,
        * with ``mode="max"`` is equivalent to :class:`~torch.nn.Embedding` followed by ``torch.max(dim=1)``.

    However, :class:`~torch.nn.EmbeddingBag` is much more time and memory
    efficient than using a chain of these operations.

    QREmbeddingBag also supports per-sample weights as an argument
    to the forward pass. This scales the output of the Embedding
    before performing a weighted reduction as specified by ``mode``.

    If :attr:`per_sample_weights`` is passed, the only supported ``mode`` is
    ``"sum"``, which computes a weighted sum according to :attr:
    `per_sample_weights`.

    Known Issues:
    Autograd breaks with multiple GPUs. It breaks only with
    multiple embeddings.

    Args:
        num_categories (int): total number of unique categories.
            The input indices must be in 0, 1, ..., num_categories - 1.
        embedding_dim (list): list of sizes for each embedding vector in each table.
            If ``"add"`` or ``"mult"`` operation are used, these embedding
            dimensions must be the same.
            If a single embedding_dim is used, then it will use this
            embedding_dim for both embedding tables.
        num_collisions (int): number of collisions to enforce.
        operation (string, optional):
            ``"concat"``, ``"add"``, or ``"mult". Specifies the operation
            to compose embeddings. ``"concat"`` concatenates the embeddings,
            ``"add"`` sums the embeddings, and ``"mult"`` multiplies
            (component-wise) the embeddings.
            Default: ``"mult"``
        max_norm (float, optional):
            If given, each embedding vector with norm larger than
            :attr:`max_norm` is renormalized to have norm
            :attr:`max_norm`.
        norm_type (float, optional):
            The p of the p-norm to compute for the :attr:`max_norm` option.
            Default ``2``.
        scale_grad_by_freq (boolean, optional):
            if given, this will scale gradients by the inverse
            of frequency of the words in the mini-batch.
            Default ``False``.

            .. note::
                This option is not supported when ``mode="max"``.

        mode (string, optional):
            ``"sum"``, ``"mean"`` or ``"max"``. Specifies the way to reduce the
            bag.

            * ``"sum"`` computes the weighted sum, taking `per_sample_weights` into consideration.
            * ``"mean"`` computes the average of the values in the bag,
            * ``"max"`` computes the max value over each bag.

            Default: ``"mean"``

        sparse (bool, optional):
            if ``True``, gradient w.r.t. :attr:`weight` matrix
            will be a sparse tensor.
            See Notes for more details regarding sparse gradients.

            .. note::
                This option is not supported when ``mode="max"``.

    Attributes:
        weight (Tensor):
            the learnable weights of each embedding table
            is the module of shape `(num_embeddings, embedding_dim)`
            initialized using a uniform distribution
            with sqrt(1 / num_categories).

    Inputs:
        :attr:`input` (LongTensor), :attr:`offsets` (LongTensor, optional), and
            :attr:`per_index_weights` (Tensor, optional)
            If :attr:`input` is 2D of shape `(B, N)`,
            it will be treated as ``B`` bags (sequences) each of
            fixed length ``N``, and this will return ``B`` values
            aggregated in a way depending on the :attr:`mode`.
            :attr:`offsets` is ignored and required to be ``None``
            in this case.
            If :attr:`input` is 1D of shape `(N)`,
            it will be treated as a concatenation of multiple bags (sequences).
            :attr:`offsets` is required to be a 1D tensor containing the
            starting index positions of each bag in :attr:`input`. Therefore,
            for :attr:`offsets` of shape `(B)`, :attr:`input` will be viewed as
            having ``B`` bags. Empty bags (i.e., having 0-length) will have
            returned vectors filled by zeros.
        per_sample_weights (Tensor, optional):
            a tensor of float / double weights, or None
            to indicate all weights should be taken to be ``1``.
            If specified, :attr:`per_sample_weights` must have exactly the
            same shape as input and is treated as having the same
            :attr:`offsets`, if those are not ``None``.
            Only supported for ``mode='sum'``.

    Returns:
        The output tensor of shape `(B, embedding_dim)`

    """
    __constants__ = ['num_embeddings', 'embedding_dim', 'num_collisions',
        'operation', 'max_norm', 'norm_type', 'scale_grad_by_freq', 'mode',
        'sparse']

    def __init__(self, num_embeddings, embedding_dim, num_collisions,
        operation='mult', max_norm=None, norm_type=2.0, scale_grad_by_freq=
        False, mode='mean', sparse=False, _weight=None):
        super(QREmbeddingBagNew, self).__init__()
        assert operation in ['concat', 'mult', 'add'], 'Not valid operation!'
        self.num_categories = num_embeddings
        if isinstance(embedding_dim, int) or len(embedding_dim) == 1:
            self.embedding_dim = [embedding_dim, embedding_dim]
        else:
            self.embedding_dim = embedding_dim
        self.num_collisions = num_collisions
        self.operation = operation
        self.max_norm = max_norm
        self.norm_type = norm_type
        self.scale_grad_by_freq = scale_grad_by_freq
        if self.operation == 'add' or self.operation == 'mult':
            assert self.embedding_dim[0] == self.embedding_dim[1
                ], 'Embedding dimensions do not match!'
        self.num_embeddings = [int(np.ceil(num_embeddings / num_collisions)
            ), num_collisions]
        if _weight is None:
            self.weight_q = Parameter(torch.Tensor(self.num_embeddings[0],
                self.embedding_dim[0]))
            self.weight_r = Parameter(torch.Tensor(self.num_embeddings[1],
                self.embedding_dim[1]))
            self.reset_parameters()
        else:
            assert list(_weight[0].shape) == [self.num_embeddings[0], self.
                embedding_dim[0]
                ], 'Shape of weight for quotient table does not' + 'match num_embeddings and embedding_dim'
            assert list(_weight[1].shape) == [self.num_embeddings[1], self.
                embedding_dim[1]
                ], 'Shape of weight for remainder table does not' + 'match num_embeddings and embedding_dim'
            self.weight_q = Parameter(_weight[0])
            self.weight_r = Parameter(_weight[1])
        self.mode = mode
        self.sparse = sparse

    def reset_parameters(self):
        nn.init.uniform_(self.weight_q, np.sqrt(1 / self.num_categories))
        nn.init.uniform_(self.weight_r, np.sqrt(1 / self.num_categories))

    def extra_repr(self):
        s = '{num_embeddings}, {embedding_dim}'
        if self.max_norm is not None:
            s += ', max_norm={max_norm}'
        if self.norm_type != 2:
            s += ', norm_type={norm_type}'
        if self.scale_grad_by_freq is not False:
            s += ', scale_grad_by_freq={scale_grad_by_freq}'
        s += ', mode={mode}'
        return s.format(**self.__dict__)

    def forward(self, input_0):
        primals_2 = self.weight_q
        primals_1 = self.weight_r
        primals_3 = input_0
        output = call([primals_1, primals_2, primals_3])
        return output[0]


Measured latency_ms: 0.061695