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 libdevice, math as tl_math
import numpy as np
import torch.utils.data
import torch
import torch.nn as nn
import torch.nn.functional as F
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_native_layer_norm_0(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 + 4 * x0, xmask, eviction_policy='evict_last')
    tmp1 = tl.load(in_ptr0 + (1 + 4 * x0), xmask, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr0 + (2 + 4 * x0), xmask, eviction_policy='evict_last')
    tmp5 = tl.load(in_ptr0 + (3 + 4 * x0), xmask, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tmp4 = tmp2 + tmp3
    tmp6 = tmp4 + tmp5
    tmp7 = 4.0
    tmp8 = tmp6 / tmp7
    tmp9 = tmp0 - tmp8
    tmp10 = tmp9 * tmp9
    tmp11 = tmp1 - tmp8
    tmp12 = tmp11 * tmp11
    tmp13 = tmp10 + tmp12
    tmp14 = tmp3 - tmp8
    tmp15 = tmp14 * tmp14
    tmp16 = tmp13 + tmp15
    tmp17 = tmp5 - tmp8
    tmp18 = tmp17 * tmp17
    tmp19 = tmp16 + tmp18
    tmp20 = tmp19 / tmp7
    tmp21 = 1e-06
    tmp22 = tmp20 + tmp21
    tmp23 = libdevice.rsqrt(tmp22)
    tl.store(out_ptr0 + x0, tmp8, xmask)
    tl.store(out_ptr1 + x0, tmp23, xmask)


@triton.jit
def triton_poi_fused_native_layer_norm_1(in_ptr0, in_ptr1, in_ptr2, in_ptr3,
    in_ptr4, out_ptr0, 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
    x0 = xindex % 4
    tmp0 = tl.load(in_ptr0 + x2, xmask)
    tmp1 = tl.load(in_ptr1 + x1, xmask, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr2 + x1, xmask, eviction_policy='evict_last')
    tmp5 = tl.load(in_ptr3 + x0, xmask, eviction_policy='evict_last')
    tmp7 = tl.load(in_ptr4 + x0, xmask, eviction_policy='evict_last')
    tmp2 = tmp0 - tmp1
    tmp4 = tmp2 * tmp3
    tmp6 = tmp4 * tmp5
    tmp8 = tmp6 + tmp7
    tl.store(out_ptr0 + x2, tmp8, xmask)


@triton.jit
def triton_poi_fused_clone_div_2(in_ptr0, out_ptr0, xnumel, XBLOCK: tl.
    constexpr):
    xnumel = 256
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex % 4
    x1 = xindex // 4 % 4
    x2 = xindex // 16 % 4
    x3 = xindex // 64
    x4 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + 4 * x2 + 16 * x1 + 64 * x3), xmask)
    tmp1 = 0.5
    tmp2 = tmp0 * tmp1
    tl.store(out_ptr0 + x4, tmp2, xmask)


@triton.jit
def triton_poi_fused_clone_3(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK: tl.
    constexpr, XBLOCK: tl.constexpr):
    ynumel = 64
    xnumel = 4
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = yindex < ynumel
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y0 = yindex % 16
    y1 = yindex // 16
    y3 = yindex
    tmp0 = tl.load(in_ptr0 + (y0 + 16 * x2 + 64 * y1), xmask & ymask,
        eviction_policy='evict_last')
    tl.store(out_ptr0 + (x2 + 4 * y3), tmp0, xmask & ymask)


@triton.jit
def triton_poi_fused__softmax_4(in_ptr0, 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
    x1 = xindex // 4
    tmp0 = tl.load(in_ptr0 + x2, xmask)
    tmp1 = tl.load(in_ptr0 + 4 * x1, xmask, eviction_policy='evict_last')
    tmp2 = tl.load(in_ptr0 + (1 + 4 * x1), xmask, eviction_policy='evict_last')
    tmp4 = tl.load(in_ptr0 + (2 + 4 * x1), xmask, eviction_policy='evict_last')
    tmp6 = tl.load(in_ptr0 + (3 + 4 * x1), xmask, eviction_policy='evict_last')
    tmp3 = triton_helpers.maximum(tmp1, tmp2)
    tmp5 = triton_helpers.maximum(tmp3, tmp4)
    tmp7 = triton_helpers.maximum(tmp5, tmp6)
    tmp8 = tmp0 - tmp7
    tmp9 = tl_math.exp(tmp8)
    tl.store(out_ptr0 + x2, tmp9, xmask)


@triton.jit
def triton_poi_fused__softmax_5(in_ptr0, 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
    x1 = xindex // 4
    tmp0 = tl.load(in_ptr0 + x2, xmask)
    tmp1 = tl.load(in_ptr0 + 4 * x1, xmask, eviction_policy='evict_last')
    tmp2 = tl.load(in_ptr0 + (1 + 4 * x1), xmask, eviction_policy='evict_last')
    tmp4 = tl.load(in_ptr0 + (2 + 4 * x1), xmask, eviction_policy='evict_last')
    tmp6 = tl.load(in_ptr0 + (3 + 4 * x1), xmask, eviction_policy='evict_last')
    tmp3 = tmp1 + tmp2
    tmp5 = tmp3 + tmp4
    tmp7 = tmp5 + tmp6
    tmp8 = tmp0 / tmp7
    tl.store(out_ptr0 + x2, tmp8, xmask)


@triton.jit
def triton_poi_fused_clone_6(in_ptr0, out_ptr0, xnumel, XBLOCK: tl.constexpr):
    xnumel = 256
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex % 4
    x1 = xindex // 4 % 4
    x2 = xindex // 16 % 4
    x3 = xindex // 64
    x4 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + 4 * x2 + 16 * x1 + 64 * x3), xmask)
    tl.store(out_ptr0 + x4, tmp0, xmask)


@triton.jit
def triton_poi_fused_add_7(in_out_ptr0, in_ptr0, xnumel, XBLOCK: tl.constexpr):
    xnumel = 64
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex
    tmp0 = tl.load(in_out_ptr0 + x0, xmask)
    tmp1 = tl.load(in_ptr0 + x0, xmask)
    tmp2 = tmp0 + tmp1
    tl.store(in_out_ptr0 + x0, tmp2, xmask)


def call(args):
    (primals_1, primals_2, primals_3, primals_4, primals_5, primals_6,
        primals_7, primals_8, primals_9) = args
    args.clear()
    assert_size_stride(primals_1, (4, 4, 4), (16, 4, 1))
    assert_size_stride(primals_2, (4, 4, 4), (16, 4, 1))
    assert_size_stride(primals_3, (4, 4, 4), (16, 4, 1))
    assert_size_stride(primals_4, (4,), (1,))
    assert_size_stride(primals_5, (4,), (1,))
    assert_size_stride(primals_6, (16, 4), (4, 1))
    assert_size_stride(primals_7, (16, 4), (4, 1))
    assert_size_stride(primals_8, (16, 4), (4, 1))
    assert_size_stride(primals_9, (4, 16), (16, 1))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((4, 4, 1), (4, 1, 16), torch.float32)
        buf1 = empty_strided_cuda((4, 4, 1), (4, 1, 16), torch.float32)
        get_raw_stream(0)
        triton_poi_fused_native_layer_norm_0[grid(16)](primals_1, buf0,
            buf1, 16, XBLOCK=16, num_warps=1, num_stages=1)
        buf2 = empty_strided_cuda((4, 4, 4), (16, 4, 1), torch.float32)
        triton_poi_fused_native_layer_norm_1[grid(64)](primals_1, buf0,
            buf1, primals_4, primals_5, buf2, 64, XBLOCK=64, num_warps=1,
            num_stages=1)
        del buf0
        del buf1
        del primals_4
        del primals_5
        buf3 = empty_strided_cuda((16, 16), (16, 1), torch.float32)
        extern_kernels.mm(reinterpret_tensor(buf2, (16, 4), (4, 1), 0),
            reinterpret_tensor(primals_6, (4, 16), (1, 4), 0), out=buf3)
        buf4 = empty_strided_cuda((16, 16), (16, 1), torch.float32)
        extern_kernels.mm(reinterpret_tensor(primals_2, (16, 4), (4, 1), 0),
            reinterpret_tensor(primals_7, (4, 16), (1, 4), 0), out=buf4)
        del primals_7
        buf5 = empty_strided_cuda((16, 16), (16, 1), torch.float32)
        extern_kernels.mm(reinterpret_tensor(primals_3, (16, 4), (4, 1), 0),
            reinterpret_tensor(primals_8, (4, 16), (1, 4), 0), out=buf5)
        del primals_8
        buf6 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.float32)
        triton_poi_fused_clone_div_2[grid(256)](buf3, buf6, 256, XBLOCK=128,
            num_warps=4, num_stages=1)
        buf7 = reinterpret_tensor(buf3, (4, 4, 4, 4), (64, 16, 4, 1), 0)
        del buf3
        triton_poi_fused_clone_3[grid(64, 4)](buf4, buf7, 64, 4, XBLOCK=4,
            YBLOCK=32, num_warps=4, num_stages=1)
        buf8 = reinterpret_tensor(buf4, (16, 4, 4), (16, 4, 1), 0)
        del buf4
        extern_kernels.bmm(reinterpret_tensor(buf6, (16, 4, 4), (16, 4, 1),
            0), reinterpret_tensor(buf7, (16, 4, 4), (16, 4, 1), 0), out=buf8)
        buf9 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.float32)
        triton_poi_fused__softmax_4[grid(256)](buf8, buf9, 256, XBLOCK=128,
            num_warps=4, num_stages=1)
        buf10 = reinterpret_tensor(buf8, (4, 4, 4, 4), (64, 16, 4, 1), 0)
        del buf8
        triton_poi_fused__softmax_5[grid(256)](buf9, buf10, 256, XBLOCK=256,
            num_warps=4, num_stages=1)
        buf11 = buf9
        del buf9
        triton_poi_fused_clone_6[grid(256)](buf5, buf11, 256, XBLOCK=128,
            num_warps=4, num_stages=1)
        buf12 = reinterpret_tensor(buf5, (16, 4, 4), (16, 4, 1), 0)
        del buf5
        extern_kernels.bmm(reinterpret_tensor(buf10, (16, 4, 4), (16, 4, 1),
            0), reinterpret_tensor(buf11, (16, 4, 4), (16, 4, 1), 0), out=buf12
            )
        buf13 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.float32)
        triton_poi_fused_clone_6[grid(256)](buf12, buf13, 256, XBLOCK=128,
            num_warps=4, num_stages=1)
        del buf12
        buf14 = empty_strided_cuda((16, 4), (4, 1), torch.float32)
        extern_kernels.mm(reinterpret_tensor(buf13, (16, 16), (16, 1), 0),
            reinterpret_tensor(primals_9, (16, 4), (1, 16), 0), out=buf14)
        buf15 = reinterpret_tensor(buf14, (4, 4, 4), (16, 4, 1), 0)
        del buf14
        triton_poi_fused_add_7[grid(64)](buf15, primals_1, 64, XBLOCK=64,
            num_warps=1, num_stages=1)
    return buf15, buf10, primals_1, reinterpret_tensor(buf2, (16, 4), (4, 1), 0
        ), reinterpret_tensor(primals_2, (16, 4), (4, 1), 0
        ), reinterpret_tensor(primals_3, (16, 4), (4, 1), 0
        ), buf10, reinterpret_tensor(buf13, (16, 16), (16, 1), 0
        ), primals_9, reinterpret_tensor(buf11, (16, 4, 4), (16, 1, 4), 0
        ), reinterpret_tensor(buf6, (16, 4, 4), (16, 1, 4), 0
        ), reinterpret_tensor(buf7, (16, 4, 4), (16, 1, 4), 0), primals_6


class PositionalEncoding(nn.Module):

    def __init__(self, max_pos, d_k):
        super().__init__()
        self.w_rpr = nn.Linear(d_k, max_pos + 1, bias=False)

    def __call__(self, q, dist_matrices):
        return self.forward(q, dist_matrices)

    def forward(self, q, dist_matrices):
        """
        :param q:  [batch, heads, seq, d_k]
        :param dist_matrices:  list of dist_matrix , each of size n_edges X nedges
        :return: resampled_q_dot_rpr: [batch, heads, seq, seq]
        """
        q_dot_rpr = self.w_rpr(q)
        attn_rpr = self.resample_rpr_product(q_dot_rpr, dist_matrices)
        return attn_rpr

    @staticmethod
    def resample_rpr_product(q_dot_rpr, dist_matrices):
        """
        :param q_dot_rpr:  [batch, heads, seq, max_pos+1]
        :param dist_matrices: list of dist_matrix , each of size n_edges X nedges
        :return: resampled_q_dot_rpr: [batch, heads, seq, seq]
        """
        bs, _n_heads, max_seq, _ = q_dot_rpr.shape
        max_pos = q_dot_rpr.shape[-1] - 1
        seq_lens = np.array([d.shape[0] for d in dist_matrices])
        if (seq_lens == max_seq).all():
            pos_inds = np.stack(dist_matrices)
        else:
            pos_inds = np.ones((bs, max_seq, max_seq), dtype=np.int32
                ) * np.iinfo(np.int32).max
            for i_b in range(bs):
                dist_matrix = dist_matrices[i_b]
                n_edges = dist_matrix.shape[0]
                pos_inds[i_b, :n_edges, :n_edges] = dist_matrix
        pos_inds[pos_inds > max_pos] = max_pos
        batch_inds = np.arange(bs)[:, None, None]
        edge_inds = np.arange(max_seq)[None, :, None]
        resampled_q_dot_rpr = q_dot_rpr[batch_inds, :, edge_inds, pos_inds
            ].permute(0, 3, 1, 2)
        return resampled_q_dot_rpr


class PositionalScaledDotProductAttention(nn.Module):
    """ Scaled Dot-Product Attention with optional positional encodings """

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

    def forward(self, q, k, v, mask=None, dist_matrices=None):
        """
        q: [batch, heads, seq, d_k]  queries
        k: [batch, heads, seq, d_k]  keys
        v: [batch, heads, seq, d_v]  values
        mask: [batch, 1, seq, seq]   for each edge, which other edges should be accounted for. "None" means all of them.
                                mask is important when using local attention, or when the meshes are of different sizes.
        rpr: [batch, seq, seq, d_k]  positional representations
        """
        attn_k = torch.matmul(q / self.temperature, k.transpose(2, 3))
        if self.positional_encoding is None:
            attn = attn_k
        else:
            attn_rpr = self.positional_encoding(q / self.temperature,
                dist_matrices)
            attn = attn_k + attn_rpr
        if mask is not None:
            attn = attn.masked_fill(mask == 0, -1000000000.0)
        attn = self.dropout(F.softmax(attn, dim=-1))
        output = torch.matmul(attn, v)
        return output, attn


class MultiHeadAttentionNew(nn.Module):
    """
    Multi-Head Attention module
    from https://github.com/jadore801120/attention-is-all-you-need-pytorch
    by Yu-Hsiang Huang.
    use_values_as_is is our addition.
    """

    def __init__(self, n_head, d_model, d_k, d_v, dropout=0.1,
        use_values_as_is=False, use_positional_encoding=False,
        max_relative_position=8):
        super().__init__()
        self.attention_type = type
        self.n_head = n_head
        self.d_k = d_k
        self.w_qs = nn.Linear(d_model, n_head * d_k, bias=False)
        self.w_ks = nn.Linear(d_model, n_head * d_k, bias=False)
        if not use_values_as_is:
            self.d_v = d_v
            self.w_vs = nn.Linear(d_model, n_head * d_v, bias=False)
            self.fc = nn.Linear(n_head * d_v, d_model, bias=False)
        else:
            self.d_v = d_model
            self.w_vs = lambda x: self.__repeat_single_axis(x, -1, n_head)
            self.fc = lambda x: self.__average_head_results(x, n_head)
        self.dropout = nn.Dropout(dropout)
        self.layer_norm = nn.LayerNorm(d_model, eps=1e-06)
        positional_encoding = None
        if use_positional_encoding:
            positional_encoding = PositionalEncoding(max_relative_position, d_k
                )
        self.attention = PositionalScaledDotProductAttention(temperature=
            d_k ** 0.5, positional_encoding=positional_encoding)

    @staticmethod
    def __repeat_single_axis(x, axis, n_rep):
        rep_sizes = [1] * x.ndim
        rep_sizes[axis] = n_rep
        x_rep = x.repeat(rep_sizes)
        return x_rep

    @staticmethod
    def __average_head_results(x, n_head):
        shape = list(x.shape)[:-1] + [n_head, -1]
        avg_x = x.view(shape).mean(-2)
        return avg_x

    def forward(self, input_0, input_1, input_2):
        primals_6 = self.w_qs.weight
        primals_7 = self.w_ks.weight
        primals_8 = self.w_vs.weight
        primals_9 = self.fc.weight
        primals_4 = self.layer_norm.weight
        primals_5 = self.layer_norm.bias
        primals_1 = input_0
        primals_2 = input_1
        primals_3 = input_2
        output = call([primals_1, primals_2, primals_3, primals_4,
            primals_5, primals_6, primals_7, primals_8, primals_9])
        return output[0], output[1]


Measured latency_ms: 0.051644