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 torch.nn as nn
from torch.autograd import Variable
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


@triton.jit
def triton_poi_fused__to_copy_add_elu_ge_mul_rsub_0(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
    tmp0 = tl.load(in_ptr0 + x0, xmask)
    tmp1 = 0.0
    tmp2 = tmp0 >= tmp1
    tmp3 = tmp2.to(tl.float32)
    tmp4 = tmp3 * tmp0
    tmp5 = 1.0
    tmp6 = tmp5 - tmp3
    tmp7 = tmp0 > tmp1
    tmp8 = tmp0 * tmp5
    tmp9 = libdevice.expm1(tmp8)
    tmp10 = tmp9 * tmp5
    tmp11 = tl.where(tmp7, tmp8, tmp10)
    tmp12 = 1.6732632423543772
    tmp13 = tmp11 * tmp12
    tmp14 = tmp6 * tmp13
    tmp15 = tmp4 + tmp14
    tmp16 = 1.0507009873554805
    tmp17 = tmp15 * tmp16
    tl.store(out_ptr0 + x0, tmp17, xmask)


def call(args):
    arg0_1, = args
    args.clear()
    assert_size_stride(arg0_1, (4, 4, 4, 4), (64, 16, 4, 1))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.float32)
        get_raw_stream(0)
        triton_poi_fused__to_copy_add_elu_ge_mul_rsub_0[grid(256)](arg0_1,
            buf0, 256, XBLOCK=128, num_warps=4, num_stages=1)
        del arg0_1
    return buf0,


def where(condition, if_true, if_false):
    """
    Torch equivalent of numpy.where.

    Parameters
    ----------
    condition : torch.ByteTensor or torch.cuda.ByteTensor or torch.autograd.Variable
        Condition to check.
    if_true : torch.Tensor or torch.cuda.Tensor or torch.autograd.Variable
        Output value if condition is true.
    if_false: torch.Tensor or torch.cuda.Tensor or torch.autograd.Variable
        Output value if condition is false

    Returns
    -------
    torch.Tensor

    Raises
    ------
    AssertionError
        if if_true and if_false are not both variables or both tensors.
    AssertionError
        if if_true and if_false don't have the same datatype.
    """
    if isinstance(if_true, Variable) or isinstance(if_false, Variable):
        assert isinstance(condition, Variable
            ), 'Condition must be a variable if either if_true or if_false is a variable.'
        assert isinstance(if_false, Variable) and isinstance(if_false, Variable
            ), 'Both if_true and if_false must be variables if either is one.'
        assert if_true.data.type() == if_false.data.type(
            ), 'Type mismatch: {} and {}'.format(if_true.data.type(),
            if_false.data.type())
    else:
        assert not isinstance(condition, Variable
            ), 'Condition must not be a variable because neither if_true nor if_false is one.'
        assert if_true.type() == if_false.type(
            ), 'Type mismatch: {} and {}'.format(if_true.data.type(),
            if_false.data.type())
    casted_condition = condition.type_as(if_true)
    output = casted_condition * if_true + (1 - casted_condition) * if_false
    return output


class SELUNew(nn.Module):

    @staticmethod
    def selu(x):
        alpha = 1.6732632423543772
        scale = 1.0507009873554805
        return scale * where(x >= 0, x, alpha * F.elu(x))

    def forward(self, input_0):
        arg0_1 = input_0
        output = call([arg0_1])
        return output[0]


Measured latency_ms: 0.03234