o
    ¸iÂ  ã                   @   sD  d dl Z d dlmZ d dlmZ d dlZd dlmZ e  	de j
 ¡Ze  	d¡Zde  	de j
 ¡ ZG dd	„ d	eeƒZd
ee fdd„Zd
ee fdd„Zejdd„ ƒZejdd„ ƒZejdd„ ƒZejdd„ ƒZejdd„ ƒZejdd„ ƒZejdd„ ƒZejdd„ ƒZejdd „ ƒZejd!d"„ ƒZejd#d$„ ƒZejd%d&„ ƒZdS )'é    N)ÚEnum)ÚOptionalç       @ç      à?ç      ð?é   c                   @   s    e Zd ZdZdZdZdZdZdS )Ú
ActivationÚsquared_reluÚgeluÚgelu_approxÚ
leaky_reluÚreluN)Ú__name__Ú
__module__Ú__qualname__ÚSquaredReLUÚGeLUÚ
GeLUApproxÚ	LeakyReLUÚReLU© r   r   úY/home/ubuntu/vllm_env/lib/python3.10/site-packages/flash_attn/ops/triton/k_activations.pyr      s    r   Ú
activationc              
   C   ó.   | rt jtt jtt jtt jtt j	t
i|  S d S ©N)r   r   r   r   r   r   r
   r   r   r   r	   ©r   r   r   r   Úget_triton_activation_kernel   ó   	ùûúÿ	÷r   c              
   C   r   r   )r   r   Ú	relu_gradr   Úleaky_relu_gradr   Ú	gelu_gradr   Úgelu_approx_gradr   Úsquared_relu_gradr   r   r   r   Ú get_triton_activation_bwd_kernel)   r   r#   c                 C   s   dt  d|  ¡ d S )Nr   é   )ÚtlÚsigmoid©Úxr   r   r   Útanh7   s   r)   c                 C   s   t  | ¡}|d|  d S )Nr   r   )r%   Úexp)r(   Úexp_xr   r   r   Úcosh=   s   
r,   c                 C   s   d}t  | dk| | | j¡¡S )zo
    ReLU_ activation function

    .. _ReLU: https://pytorch.org/docs/stable/generated/torch.nn.ReLU.html
    ç        r   ©r%   ÚwhereÚtoÚdtype)r(   Úzeror   r   r   r   G   s   r   c                 C   s*   d}d}t  | dk| | j¡| | j¡¡S )Nr-   r   r   r.   )r(   r2   Úoner   r   r   r   R   s   "r   c                 C   s   t | ƒ}||  | j¡S )zv
    Squared ReLU activation, as proposed in the Primer_ paper.

    .. _Primer: https://arxiv.org/abs/2109.08668
    )r   r0   r1   )r(   Úx_r   r   r   r	   \   s   r	   c                 C   s   t  | dkd|  d¡S )Nr   r   r-   )r%   r/   r'   r   r   r   r"   g   s   r"   c                 C   s&   d}|  | j¡}t | dk| ||  ¡S )zu
    LeakyReLU_ activation

    .. _LeakyReLU: https://pytorch.org/docs/stable/generated/torch.nn.LeakyReLU.html
    ç{®Gáz„?r   ©r0   r1   r%   r/   )r(   Úscaler   r   r   r   m   s   r   c                 C   s2   d}d}|  | j¡}|  | j¡}t | dk||¡S )Nr5   r$   r   r6   )r(   Úmin_gradÚmax_gradr   r   r   r   y   s
   r   c                 C   s   | d dt j | t ¡  S )z!Gaussian Error Linear Unit (GELU)r   r   )r%   Ú	libdeviceÚerfÚ_sqrt1_2r'   r   r   r   r
   „   s   r
   c                 C   s:   ddt j | t ¡  }t  d|  |  ¡t }|| |  S )Nr   r   g      à¿)r%   r:   r;   r<   r*   Ú_gaussian_pdf_normalization)r(   ÚcdfÚpdfr   r   r   r    Š   s   r    c                 C   s(   d|  dt t|  dd|  |    ƒ  S )z„
    GeLU_ activation - Gaussian error linear unit, with tanh approximation

    .. _GeLU: https://arxiv.org/pdf/1606.08415.pdf
    r   r   ç÷Hmâä¦?)r)   Ú_sqrt2pir'   r   r   r   r   ‘   s   (r   c                 C   sL   t d|  dd|  |    ƒ}d|  d||  dd|  |     dd|   S )Ng Þe3Eˆé?r$   r@   r   g6”ü¾vf»?)r)   )r(   Útanh_outr   r   r   r!   ›   s   $ÿr!   )ÚmathÚenumr   Útypingr   ÚtritonÚtriton.languageÚlanguager%   ÚsqrtÚpirA   r<   r=   Ústrr   r   r#   Újitr)   r,   r   r   r	   r"   r   r   r
   r    r   r!   r   r   r   r   Ú<module>   sF   


	


	








	