# AOT ID: ['0_inference']
from ctypes import c_void_p, c_long, c_int
import torch
import math
import random
import os
import tempfile
from math import inf, nan
from torch._inductor.hooks import run_intermediate_hooks
from torch._inductor.utils import maybe_profile
from torch._inductor.codegen.memory_planning import _align as align
from torch import device, empty_strided
from torch._inductor.async_compile import AsyncCompile
from torch._inductor.select_algorithm import extern_kernels
from torch._inductor.codegen.multi_kernel import MultiKernelCall
import triton
import triton.language as tl
from torch._inductor.runtime.triton_heuristics import grid, split_scan_grid, grid_combo_kernels, start_graph, end_graph
from torch._C import _cuda_getCurrentRawStream as get_raw_stream

aten = torch.ops.aten
inductor_ops = torch.ops.inductor
_quantized = torch.ops._quantized
assert_size_stride = torch._C._dynamo.guards.assert_size_stride
empty_strided_cpu = torch._C._dynamo.guards._empty_strided_cpu
empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda
empty_strided_xpu = torch._C._dynamo.guards._empty_strided_xpu
reinterpret_tensor = torch._C._dynamo.guards._reinterpret_tensor
alloc_from_pool = torch.ops.inductor._alloc_from_pool
async_compile = AsyncCompile()


# kernel path: runs/run_shard_4/inductor_cache/mr/cmr7a6ykanzkwpcsrqlcnz2nradw5cvqpademfs4557bpgql5b44.py
# Topologically Sorted Source Nodes: [attn_1, mul, result, result_1, sum_1], Original ATen: [aten.div, aten.mul, aten._softmax, aten.sum]
# Source node to ATen node mapping:
#   attn_1 => div
#   mul => mul
#   result => amax, div_1, exp, sub, sum_1
#   result_1 => mul_1
#   sum_1 => sum_2
# Graph fragment:
#   %div : [num_users=1] = call_function[target=torch.ops.aten.div.Tensor](args = (%bmm, 4), kwargs = {})
#   %mul : [num_users=2] = call_function[target=torch.ops.aten.mul.Tensor](args = (%div, %arg2_1), kwargs = {})
#   %amax : [num_users=1] = call_function[target=torch.ops.aten.amax.default](args = (%mul, [2], True), kwargs = {})
#   %sub : [num_users=1] = call_function[target=torch.ops.aten.sub.Tensor](args = (%mul, %amax), kwargs = {})
#   %exp : [num_users=2] = call_function[target=torch.ops.aten.exp.default](args = (%sub,), kwargs = {})
#   %sum_1 : [num_users=1] = call_function[target=torch.ops.aten.sum.dim_IntList](args = (%exp, [2], True), kwargs = {})
#   %div_1 : [num_users=1] = call_function[target=torch.ops.aten.div.Tensor](args = (%exp, %sum_1), kwargs = {})
#   %mul_1 : [num_users=2] = call_function[target=torch.ops.aten.mul.Tensor](args = (%div_1, %arg2_1), kwargs = {})
#   %sum_2 : [num_users=1] = call_function[target=torch.ops.aten.sum.dim_IntList](args = (%mul_1, [2], True), kwargs = {})
triton_poi_fused__softmax_div_mul_sum_0 = async_compile.triton('triton_poi_fused__softmax_div_mul_sum_0', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[16], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: '*fp32', 5: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=80, major=8, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=108), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused__softmax_div_mul_sum_0', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 8, 'num_reduction': 0, 'backend_hash': 'A9C866B4A14FD3277824029365D703C2427B2E685E54EC9B3EF4ADC8D1EEAC1D', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@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)
''', device_str='cuda')


# kernel path: runs/run_shard_4/inductor_cache/cg/ccgvbcuuaussyqitdvez74m3h3ffsx647g73pwhsr5ylvmy3s67d.py
# Topologically Sorted Source Nodes: [attn_1, mul, result, result_1, add, result_2], Original ATen: [aten.div, aten.mul, aten._softmax, aten.add]
# Source node to ATen node mapping:
#   add => add
#   attn_1 => div
#   mul => mul
#   result => amax, div_1, exp, sub
#   result_1 => mul_1
#   result_2 => div_2
# Graph fragment:
#   %div : [num_users=1] = call_function[target=torch.ops.aten.div.Tensor](args = (%bmm, 4), kwargs = {})
#   %mul : [num_users=2] = call_function[target=torch.ops.aten.mul.Tensor](args = (%div, %arg2_1), kwargs = {})
#   %amax : [num_users=1] = call_function[target=torch.ops.aten.amax.default](args = (%mul, [2], True), kwargs = {})
#   %sub : [num_users=1] = call_function[target=torch.ops.aten.sub.Tensor](args = (%mul, %amax), kwargs = {})
#   %exp : [num_users=2] = call_function[target=torch.ops.aten.exp.default](args = (%sub,), kwargs = {})
#   %div_1 : [num_users=1] = call_function[target=torch.ops.aten.div.Tensor](args = (%exp, %sum_1), kwargs = {})
#   %mul_1 : [num_users=2] = call_function[target=torch.ops.aten.mul.Tensor](args = (%div_1, %arg2_1), kwargs = {})
#   %add : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%sum_2, 1e-13), kwargs = {})
#   %div_2 : [num_users=2] = call_function[target=torch.ops.aten.div.Tensor](args = (%mul_1, %add), kwargs = {})
triton_poi_fused__softmax_add_div_mul_1 = async_compile.triton('triton_poi_fused__softmax_add_div_mul_1', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[64], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: '*fp32', 5: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=80, major=8, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=108), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused__softmax_add_div_mul_1', 'mutated_arg_names': ['in_out_ptr0'], 'no_x_dim': False, 'num_load': 5, 'num_reduction': 0, 'backend_hash': 'A9C866B4A14FD3277824029365D703C2427B2E685E54EC9B3EF4ADC8D1EEAC1D', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@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)
''', device_str='cuda')


async_compile.wait(globals())
del async_compile

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)
        # Topologically Sorted Source Nodes: [attn], Original ATen: [aten.bmm]
        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)
        # Topologically Sorted Source Nodes: [attn_1, mul, result, result_1, sum_1], Original ATen: [aten.div, aten.mul, aten._softmax, aten.sum]
        stream0 = get_raw_stream(0)
        triton_poi_fused__softmax_div_mul_sum_0.run(buf0, arg2_1, buf1, buf2, buf3, 16, grid=grid(16), stream=stream0)
        buf4 = buf0; del buf0  # reuse
        # Topologically Sorted Source Nodes: [attn_1, mul, result, result_1, add, result_2], Original ATen: [aten.div, aten.mul, aten._softmax, aten.add]
        triton_poi_fused__softmax_add_div_mul_1.run(buf4, arg2_1, buf1, buf2, buf3, 64, grid=grid(64), stream=stream0)
        del arg2_1
        del buf1
        del buf2
        del buf3
        buf5 = empty_strided_cuda((4, 4, 4), (16, 4, 1), torch.float32)
        # Topologically Sorted Source Nodes: [output], Original ATen: [aten.bmm]
        extern_kernels.bmm(buf4, arg3_1, out=buf5)
        del arg3_1
    return (buf5, buf4, )


def benchmark_compiled_module(times=10, repeat=10):
    from torch._dynamo.testing import rand_strided
    from torch._inductor.utils import print_performance
    arg0_1 = rand_strided((4, 4, 4), (16, 4, 1), device='cuda:0', dtype=torch.float32)
    arg1_1 = rand_strided((4, 4, 4), (16, 4, 1), device='cuda:0', dtype=torch.float32)
    arg2_1 = rand_strided((4, 4, 4), (16, 4, 1), device='cuda:0', dtype=torch.float32)
    arg3_1 = rand_strided((4, 4, 4), (16, 4, 1), device='cuda:0', dtype=torch.float32)
    fn = lambda: call([arg0_1, arg1_1, arg2_1, arg3_1])
    return print_performance(fn, times=times, repeat=repeat)


if __name__ == "__main__":
    from torch._inductor.wrapper_benchmark import compiled_module_main
    compiled_module_main('None', benchmark_compiled_module)

import torch


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 ScaledDotProductAttention(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, q, k, v, mask):
        attn = torch.bmm(q, k.transpose(1, 2))
        attn = attn / self.temperature
        attn = masked_softmax(attn, mask, 2)
        __attn = self.dropout(attn)
        output = torch.bmm(__attn, v)
        return output, attn


def get_inputs():
    return [torch.rand([4, 4, 4]), torch.rand([4, 4, 4]), torch.rand([4, 4,
        4]), torch.rand([4, 4, 4])]


def get_init_inputs():
    return [[], {'temperature': 4}]

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]
