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 torch.nn as nn
from torch.nn import Parameter
import torch.optim
import torch.utils.data
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_mul_softplus_0(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 = 20.0
    tmp3 = tmp1 > tmp2
    tmp4 = tl_math.exp(tmp1)
    tmp5 = libdevice.log1p(tmp4)
    tmp6 = tl.where(tmp3, tmp1, tmp5)
    tmp7 = tmp0 * tmp6
    tl.store(out_ptr0 + x0, tmp7, xmask)


@triton.jit
def triton_poi_fused_new_full_scatter_1(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 = 0.0
    tl.store(out_ptr0 + x0, tmp0, xmask)


@triton.jit
def triton_poi_fused__softmax_new_full_scatter_2(in_ptr0, in_ptr1, out_ptr0,
    out_ptr1, xnumel, XBLOCK: tl.constexpr):
    xnumel = 8
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x2 = xindex
    x1 = xindex // 2
    tmp0 = tl.load(in_ptr0 + x2, xmask)
    tmp1 = tl.load(in_ptr0 + 2 * x1, xmask, eviction_policy='evict_last')
    tmp2 = tl.load(in_ptr0 + (1 + 2 * x1), xmask, eviction_policy='evict_last')
    tmp12 = tl.load(in_ptr1 + x2, xmask)
    tmp3 = triton_helpers.maximum(tmp1, tmp2)
    tmp4 = tmp0 - tmp3
    tmp5 = tl_math.exp(tmp4)
    tmp6 = tmp1 - tmp3
    tmp7 = tl_math.exp(tmp6)
    tmp8 = tmp2 - tmp3
    tmp9 = tl_math.exp(tmp8)
    tmp10 = tmp7 + tmp9
    tmp11 = tmp5 / tmp10
    tl.device_assert((0 <= tmp12) & (tmp12 < 4) | ~xmask,
        'index out of bounds: 0 <= tmp12 < 4')
    tl.store(out_ptr0 + x2, tmp11, xmask)
    tl.store(out_ptr1 + (tmp12 + 4 * x1), tmp11, 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, (4, 4), (4, 1))
    assert_size_stride(primals_3, (4, 4), (4, 1))
    buf0 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf1 = torch.ops.aten.normal_functional.default(buf0)
        buf2 = buf1
        del buf1
        buf3 = buf0
        del buf0
        extern_kernels.mm(primals_1, reinterpret_tensor(primals_3, (4, 4),
            (1, 4), 0), out=buf3)
        del primals_3
        buf4 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        get_raw_stream(0)
        triton_poi_fused_mul_softplus_0[grid(16)](buf2, buf3, buf4, 16,
            XBLOCK=16, num_warps=1, num_stages=1)
        buf5 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        extern_kernels.addmm(buf4, primals_1, reinterpret_tensor(primals_2,
            (4, 4), (1, 4), 0), alpha=1, beta=1, out=buf5)
        del buf4
        del primals_2
        buf6 = torch.ops.aten.topk.default(buf5, 2)
        buf7 = buf6[0]
        buf8 = buf6[1]
        del buf6
        buf10 = buf5
        del buf5
        triton_poi_fused_new_full_scatter_1[grid(16)](buf10, 16, XBLOCK=16,
            num_warps=1, num_stages=1)
        buf9 = empty_strided_cuda((4, 2), (2, 1), torch.float32)
        triton_poi_fused__softmax_new_full_scatter_2[grid(8)](buf7, buf8,
            buf9, buf10, 8, XBLOCK=8, num_warps=1, num_stages=1)
        del buf7
    return buf10, primals_1, buf2, buf3, buf8, buf9


class SparseGateNew(nn.Module):

    def __init__(self, in_features, n_experts, k=2):
        """
        Returns a sparsely gated noisy softmax.
        See OUTRAGEOUSLY LARGE NEURAL NETWORKS:
            THE SPARSELY-GATED MIXTURE-OF-EXPERTS LAYER
            Shazeer et. al
            Link: https://arxiv.org/pdf/1701.06538.pdf
        """
        assert k > 1, 'Need k >= 1. If k == 1, then derivatives are zero everywhere.'
        super(SparseGateNew, self).__init__()
        self.gate_weights = Parameter(torch.Tensor(n_experts, in_features))
        self.noise_weights = Parameter(torch.Tensor(n_experts, in_features))
        self.n_experts = n_experts
        self.n_selected = k
        self.reset_parameters()

    def reset_parameters(self):
        nn.init.constant_(self.gate_weights, 0.0)
        nn.init.constant_(self.noise_weights, 0.0)

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


Measured latency_ms: 0.042535