a
    zþdó&  ã                   @   s”   d Z ddlZddlZddlmZ ejejejejdœdd„ƒZejejejdœdd„ƒZ	ejejejejdœd	d
„ƒZ
G dd„ dejjƒZejZdS )zÔ
Fused Attention
===============
This is a Triton implementation of the Flash Attention algorithm
(see: Dao et al., https://arxiv.org/pdf/2205.14135v2.pdf; Rabe and Staats https://arxiv.org/pdf/2112.05682v2.pdf)
é    N)ÚBLOCK_MÚBLOCK_DMODELÚBLOCK_Nc           8   	   C   s   t  d¡}t  d¡}|| t  d|¡ }t  d|¡} t  d|¡}!|| |d d …d f |	  |!d d d …f |
  }"|| | d d d …f |  |!d d …d f |  }#|| | d d …d f |	  |!d d d …f |
  }$| |" }%||# }&||$ }'t j|gt jdtdƒ }(t j|gt jd})t j||gt jd}*t  |%¡}+td|d | |ƒD ],},t  |&¡}-t j||gt jd}.|.t  |+|-¡7 }.|.|9 }.t  	|d d …d f |,| d d d …f  k|.tdƒ¡}.t  
t  |.d¡|(¡}/|)t  |(|/ ¡9 })t  |.|/d d …d f  ¡}0t  |0d¡|) }1d|1 }2|0|29 }0|*|)|2 d d …d f 9 }*|0 t j¡}0t  |'¡}3|*t  |0|3¡7 }*|1})|/}(|&|| 7 }&|'|| 7 }'qFt  d¡}|| t  d|¡ }|||  | }4|||  | }5t  |4|)¡ t  |5|(¡ t  d|¡} || |d d …d f |  | d d d …f |  }6||6 }7t  |7|*¡ d S )Nr   é   ©ÚdtypeÚinfú-infg      ð?)ÚtlÚ
program_idÚarangeÚzerosÚfloat32ÚfloatÚloadÚrangeÚdotÚwhereÚmaximumÚmaxÚexpÚsumÚtoÚfloat16Ústore)8ÚQÚKÚVÚsm_scaleÚLÚMÚOutÚ	stride_qzÚ	stride_qhÚ	stride_qmÚ	stride_qkÚ	stride_kzÚ	stride_khÚ	stride_knÚ	stride_kkÚ	stride_vzÚ	stride_vhÚ	stride_vkÚ	stride_vnZ	stride_ozZ	stride_ohZ	stride_omZ	stride_onÚZÚHÚN_CTXr   r   r   Ústart_mÚoff_hzÚoffs_mÚoffs_nZoffs_dZoff_qZoff_kZoff_vÚq_ptrsÚk_ptrsÚv_ptrsZm_prevZl_prevÚaccÚqÚstart_nÚkÚqkZm_currÚpZl_currZl_rcpÚvZl_ptrsÚm_ptrsZoff_oZout_ptrs© r@   úc/var/www/html/stable-diffusion-webui/venv/lib/python3.9/site-packages/triton/ops/flash_attention.pyÚ_fwd_kernel   sZ    

000

2

0rB   ©r   ÚD_HEADc                 C   s  t  d¡| t  d|¡ }t  d|¡}t  | |d d …d f |  |d d d …f  ¡ t j¡}	t  ||d d …d f |  |d d d …f  ¡ t j¡}
t  || ¡ t j¡}|
|d d …d f  }
t j|	|
 dd}t  ||d d …d f |  |d d d …f  |
¡ t  || |¡ d S )Nr   r   )Zaxis)r
   r   r   r   r   r   r   r   )r!   ÚDOr   ZNewDOÚDeltar   rD   Zoff_mZoff_nÚoÚdoZdenomÚdeltar@   r@   rA   Ú_bwd_preprocessZ   s    660rJ   c           @   
   C   s$  t  d¡}|| } || }!| | | |!|  7 } || | |!|  7 }|| | |!|  7 }|| | |!|  7 }|| | |!|  7 }|| | |!|  7 }|| | |!|  7 }td|ƒD ]l}"|"| }#|#t  d|¡ }$|"| t  d|¡ }%t  d|¡}&t  d|¡}'| |$d d …d f | |'d d d …f |   }(||%d d …d f | |'d d d …f |   })||%d d …d f | |'d d d …f |   }*||$d d …d f | |'d d d …f |   }+||$d d …d f | |'d d d …f |   },|||  }-|
||  }.t j||gt jd}/t j||gt jd}0t  |)¡}1t  |*¡}2t|#|| |ƒD ]p}3|3|& }4t  |(¡}5t  |5t  |1¡¡}6t  	|4d d …d f |%d d d …f k|6t
dƒ¡}6t  |.|4 ¡}7t  |6| |7d d …d f  ¡}8t  |+¡}9|/t  t  |8 t j¡¡|9¡7 }/t  |-|4 ¡}:t j||gt jd|:d d …d f  };|;t  |9t  |2¡¡7 };|8|; | }<|0t  t  |< t j¡¡|5¡7 }0t  |,¡}=|=t  |< t j¡|1¡7 }=t  |,|=¡ |,|| 7 },|(|| 7 }(|+|| 7 }+q:||%d d …d f | |'d d d …f |   }>||%d d …d f | |'d d d …f |   }?t  |>|/¡ t  |?|0¡ q°d S )Nr   r   r	   )r
   r   r   r   r   r   r   r   Ztransr   r   r   r   r   r   )@r   r   r   r   r!   rE   ZDQZDKZDVr   r    ÚDr"   r#   r$   r%   r&   r'   r(   r)   r*   r+   r,   r-   r.   r/   r0   Z	num_blockr   r   r   r2   Zoff_zZoff_hr:   ÚloZoffs_qmr4   r3   Zoffs_kr5   r6   r7   Zdo_ptrsZdq_ptrsZD_ptrsr?   ÚdvÚdkr;   r>   r1   Zoffs_m_currr9   r<   Úmr=   rH   ZDiZdpZdsÚdqZdv_ptrsZdk_ptrsr@   r@   rA   Ú_bwd_kerneln   sf    
,,,,,


.
$
,,rQ   c                   @   s$   e Zd Zedd„ ƒZedd„ ƒZdS )Ú
_attentionc              !   C   sè  t j ¡ }|d dk rtdƒ‚d}|jd |jd |jd   }}}	||krT||	ksXJ ‚|dv sdJ ‚t  |¡}
t |jd |¡|jd |jd  df}t j|jd |jd  |jd f|j	t j
d	}t j|jd |jd  |jd f|j	t j
d	}|d
krþdnd}t| |||||||
| d¡| d¡| d¡| d¡| d¡| d¡| d¡| d¡| d¡| d¡| d¡| d¡|
 d¡|
 d¡|
 d¡|
 d¡|jd |jd |jd ||||dd |  ||||
||¡ || _|| _|| _|
S )Nr   é   zDFlash attention currently only supported for compute capability < 80é€   éÿÿÿÿ>   é@   é   r   )Údevicer   rV   é   é   ©r   r   r   Ú	num_warpsZ
num_stages)ÚtorchZcudaZget_device_capabilityÚRuntimeErrorÚshapeÚ
empty_likeÚtritonZcdivÚemptyrX   r   rB   ÚstrideZsave_for_backwardÚgridr   r   )Úctxr9   r;   r>   r   Z
capabilityÚBLOCKZLqZLkZLvrG   rd   r   rO   r\   r@   r@   rA   ÚforwardÄ   s<    
"
(..    õz_attention.forwardc              #   C   sJ  d}| j \}}}}}}| ¡ }tj|tjd}	t |¡}
t |¡}t |¡}t |¡}t| jd | jd  f ||||||| jd t	| jd f |||| j
|||	|
||||| d¡| d¡| d¡| d¡| d¡| d¡| d¡| d¡| d¡| d¡| d¡| d¡|jd |jd |jd | jd ||| jddd	! |	|
|d fS )
NrT   r   r   r   rC   rW   rZ   rS   r[   )Zsaved_tensorsÚ
contiguousr]   Z
zeros_liker   r`   rJ   rd   r   rQ   r   rc   r_   )re   rH   rf   r9   r;   r>   rG   ÚlrO   rP   rN   rM   Z	do_scaledrI   r@   r@   rA   Úbackwardê   s:    



ý
   óz_attention.backwardN)Ú__name__Ú
__module__Ú__qualname__Ústaticmethodrg   rj   r@   r@   r@   rA   rR   Â   s   
%rR   )Ú__doc__r]   ra   Ztriton.languageÚlanguager
   ZjitZ	constexprrB   rJ   rQ   ZautogradZFunctionrR   ÚapplyZ	attentionr@   r@   r@   rA   Ú<module>   s    
öKýõSI