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
from torch._inductor.runtime.triton_helpers import math as tl_math
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__softmax_div_mul_sum_0(in_ptr0, in_ptr1, out_ptr0,
    out_ptr1, out_ptr2, 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 + 4 * x0, xmask, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr1 + 4 * x0, xmask, eviction_policy='evict_last')
    tmp5 = tl.load(in_ptr0 + (1 + 4 * x0), xmask, eviction_policy='evict_last')
    tmp7 = tl.load(in_ptr1 + (1 + 4 * x0), xmask, eviction_policy='evict_last')
    tmp10 = tl.load(in_ptr0 + (2 + 4 * x0), xmask, eviction_policy='evict_last'
        )
    tmp12 = tl.load(in_ptr1 + (2 + 4 * x0), xmask, eviction_policy='evict_last'
        )
    tmp15 = tl.load(in_ptr0 + (3 + 4 * x0), xmask, eviction_policy='evict_last'
        )
    tmp17 = tl.load(in_ptr1 + (3 + 4 * x0), xmask, eviction_policy='evict_last'
        )
    tmp1 = 0.25
    tmp2 = tmp0 * tmp1
    tmp4 = tmp2 * tmp3
    tmp6 = tmp5 * tmp1
    tmp8 = tmp6 * tmp7
    tmp9 = triton_helpers.maximum(tmp4, tmp8)
    tmp11 = tmp10 * tmp1
    tmp13 = tmp11 * tmp12
    tmp14 = triton_helpers.maximum(tmp9, tmp13)
    tmp16 = tmp15 * tmp1
    tmp18 = tmp16 * tmp17
    tmp19 = triton_helpers.maximum(tmp14, tmp18)
    tmp20 = tmp4 - tmp19
    tmp21 = tl_math.exp(tmp20)
    tmp22 = tmp8 - tmp19
    tmp23 = tl_math.exp(tmp22)
    tmp24 = tmp21 + tmp23
    tmp25 = tmp13 - tmp19
    tmp26 = tl_math.exp(tmp25)
    tmp27 = tmp24 + tmp26
    tmp28 = tmp18 - tmp19
    tmp29 = tl_math.exp(tmp28)
    tmp30 = tmp27 + tmp29
    tmp31 = tmp21 / tmp30
    tmp32 = tmp31 * tmp3
    tmp33 = tmp23 / tmp30
    tmp34 = tmp33 * tmp7
    tmp35 = tmp32 + tmp34
    tmp36 = tmp26 / tmp30
    tmp37 = tmp36 * tmp12
    tmp38 = tmp35 + tmp37
    tmp39 = tmp29 / tmp30
    tmp40 = tmp39 * tmp17
    tmp41 = tmp38 + tmp40
    tl.store(out_ptr0 + x0, tmp19, xmask)
    tl.store(out_ptr1 + x0, tmp30, xmask)
    tl.store(out_ptr2 + x0, tmp41, xmask)


@triton.jit
def triton_poi_fused__softmax_add_div_mul_1(in_out_ptr0, in_ptr0, in_ptr1,
    in_ptr2, in_ptr3, xnumel, XBLOCK: tl.constexpr):
    xnumel = 64
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x2 = xindex
    x1 = xindex // 4
    tmp0 = tl.load(in_out_ptr0 + x2, xmask)
    tmp3 = tl.load(in_ptr0 + x2, xmask)
    tmp5 = tl.load(in_ptr1 + x1, xmask, eviction_policy='evict_last')
    tmp8 = tl.load(in_ptr2 + x1, xmask, eviction_policy='evict_last')
    tmp11 = tl.load(in_ptr3 + x1, xmask, eviction_policy='evict_last')
    tmp1 = 0.25
    tmp2 = tmp0 * tmp1
    tmp4 = tmp2 * tmp3
    tmp6 = tmp4 - tmp5
    tmp7 = tl_math.exp(tmp6)
    tmp9 = tmp7 / tmp8
    tmp10 = tmp9 * tmp3
    tmp12 = 1e-13
    tmp13 = tmp11 + tmp12
    tmp14 = tmp10 / tmp13
    tl.store(in_out_ptr0 + x2, tmp14, xmask)


def call(args):
    arg0_1, arg1_1, arg2_1, arg3_1 = args
    args.clear()
    assert_size_stride(arg0_1, (4, 4, 4), (16, 4, 1))
    assert_size_stride(arg1_1, (4, 4, 4), (16, 4, 1))
    assert_size_stride(arg2_1, (4, 4, 4), (16, 4, 1))
    assert_size_stride(arg3_1, (4, 4, 4), (16, 4, 1))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((4, 4, 4), (16, 4, 1), torch.float32)
        extern_kernels.bmm(arg1_1, reinterpret_tensor(arg0_1, (4, 4, 4), (
            16, 1, 4), 0), out=buf0)
        del arg0_1
        del arg1_1
        buf1 = empty_strided_cuda((4, 4, 1), (4, 1, 16), torch.float32)
        buf2 = empty_strided_cuda((4, 4, 1), (4, 1, 16), torch.float32)
        buf3 = empty_strided_cuda((4, 4, 1), (4, 1, 16), torch.float32)
        get_raw_stream(0)
        triton_poi_fused__softmax_div_mul_sum_0[grid(16)](buf0, arg2_1,
            buf1, buf2, buf3, 16, XBLOCK=16, num_warps=1, num_stages=1)
        buf4 = buf0
        del buf0
        triton_poi_fused__softmax_add_div_mul_1[grid(64)](buf4, arg2_1,
            buf1, buf2, buf3, 64, XBLOCK=64, num_warps=1, num_stages=1)
        del arg2_1
        del buf1
        del buf2
        del buf3
        buf5 = empty_strided_cuda((4, 4, 4), (16, 4, 1), torch.float32)
        extern_kernels.bmm(buf4, arg3_1, out=buf5)
        del arg3_1
    return buf5, buf4


def masked_softmax(vector: 'torch.Tensor', mask: 'torch.Tensor', dim: 'int'
    =-1, memory_efficient: 'bool'=False, mask_fill_value: 'float'=-1e+32
    ) ->torch.Tensor:
    """
    https://github.com/allenai/allennlp/blob/b6cc9d39651273e8ec2a7e334908ffa9de5c2026/allennlp/nn/util.py#L231
    ``torch.nn.functional.softmax(vector)`` does not work if some elements of ``vector`` should be
    masked.  This performs a softmax on just the non-masked portions of ``vector``.  Passing
    ``None`` in for the mask is also acceptable; you'll just get a regular softmax.
    ``vector`` can have an arbitrary number of dimensions; the only requirement is that ``mask`` is
    broadcastable to ``vector's`` shape.  If ``mask`` has fewer dimensions than ``vector``, we will
    unsqueeze on dimension 1 until they match.  If you need a different unsqueezing of your mask,
    do it yourself before passing the mask into this function.
    If ``memory_efficient`` is set to true, we will simply use a very large negative number for those
    masked positions so that the probabilities of those positions would be approximately 0.
    This is not accurate in math, but works for most cases and consumes less memory.
    In the case that the input vector is completely masked and ``memory_efficient`` is false, this function
    returns an array of ``0.0``. This behavior may cause ``NaN`` if this is used as the last layer of
    a model that uses categorical cross-entropy loss. Instead, if ``memory_efficient`` is true, this function
    will treat every element as equal, and do softmax over equal numbers.
    """
    if mask is None:
        result = torch.nn.functional.softmax(vector, dim=dim)
    else:
        mask = mask.float()
        while mask.dim() < vector.dim():
            mask = mask.unsqueeze(1)
        if not memory_efficient:
            result = torch.nn.functional.softmax(vector * mask, dim=dim)
            result = result * mask
            result = result / (result.sum(dim=dim, keepdim=True) + 1e-13)
        else:
            masked_vector = vector.masked_fill((1 - mask).byte(),
                mask_fill_value)
            result = torch.nn.functional.softmax(masked_vector, dim=dim)
    return result


class ScaledDotProductAttentionNew(torch.nn.Module):
    """ Scaled Dot-Product Attention """

    def __init__(self, temperature, attn_dropout=0.1):
        super().__init__()
        self.temperature = temperature
        self.dropout = torch.nn.Dropout(attn_dropout)

    def forward(self, input_0, input_1, input_2, input_3):
        arg0_1 = input_0
        arg1_1 = input_1
        arg2_1 = input_2
        arg3_1 = input_3
        output = call([arg0_1, arg1_1, arg2_1, arg3_1])
        return output[0], output[1]


Measured latency_ms: 0.04518