a
    zd                     @   s   d dl Z d dlZd dlmZ dd Zej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G dd dZdS )    Nc                 C   s4   | dkrdS | dkrdS | dkr$dS | dkr0dS d	S )
N            i      i          )nr	   r	   g/var/www/html/stable-diffusion-webui/venv/lib/python3.9/site-packages/triton/ops/blocksparse/softmax.py	num_warps   s    r   )ROW_SIZE
BLOCK_SIZEIS_DENSEc                  C   s  t d}t d}t d}|t d | }t d|
| }t d|
| }||| d  }t |d }t |d }|| }||| | | 7 }||| | 7 }|rt d|
}nH|dt d t d |  }t j|| | ||k dd}|| | }||k }t j|| | |td d}|t j}|}||9 }|d ur||| 7 }||| 7 }|| d | }|dk||k @ }t j|||  | |dd}||7 }|t j}t ||k|	@ td |}t 	|}t j
| | | ||d d S )Nr   r   r   maskotherinf        r   )tl
program_idnum_programsarangeloadfloattofloat32wheresoftmaxstore) OutAZ	stride_xzLUTRextent	stride_zr	stride_hrscale	is_causalr   r   r   hmzhmlane_nblock_nheadersizeoffsetZoff_ansoff_lutstart_nr   aoutoff_lomask_lo
rel_logitsr	   r	   r   _blocksparse_softmax_fwd   sB    	


 

r;   c           '      C   s"  t d}t d}t d}|t d | }t d|| }t d|| }||| d  }t |d }t |d }|| | | }||| | 7 }||k }|||  | }|||  | }|rt d|}nD|dt d t d |  }t j|| | |dd} | | | }t j|| |dd}!|!t j}!t j|| |dd}"|"t j}"t ||k|@ |!|!k@ d|!}!|!|"t |!|" d  }#|d ur|||
 7 }||| 7 }|	| d | }$|$dk|$|	k @ |@ }%t j	|||	  |$ |#|%d |#| }#| ||  | }&t j	|&| |#|d d S )Nr   r   r   r   r   r   )
r   r   r   r   r   r   r   r   sumr    )'ZDAZ
stride_zdxZDOutZstride_zdoutr!   Zstride_zoutr(   r#   ZDRr%   r&   r'   Z	stride_err)   r   r   r   r*   r+   r,   r-   r.   r/   r0   r1   r2   Zoff_mnr   ZAsZDOutsr3   r4   r5   r6   doutdar8   r9   ZDAsr	   r	   r   _blocksparse_softmax_bwdJ   sD    


 
r?   c                   @   s0   e Zd Zedd Zedd Zedd ZdS )_softmaxc              	   C   s   t jg t j| jd}| }t| jd D ]*}t || |d d d d f df}q*|| }t 	|}t j
|d d dd|dd < | jddd d df }t j||fddd}	t |	|ft j|}
|
t| fS )	Ndtypedevicer   )Zdimr   F)as_tupler   )torchZtensorint64rC   clonerangeshapecatr<   Z
zeros_likeZcumsumZnonzerostackviewtypeint32r   intmax)layoutblockrC   _emptysizesr*   Ztotal_sizesoffsetscolumnsr0   lutr	   r	   r   make_lut   s    (
z_softmax.make_lutc
                 C   s  |d ur,t |tjr,|jjdks$J | }|jd }
|d |d | |
g}|d u rXdn|j}|d u rjdn| }t|}t	| |||d|||d |d |d |||t
||	t|d | || || _|| _|| _|| _|| _|| _|j| _|	| _|| _|S )Ncpur   r   )r   r   r   r   rD   r   r   r   r   )
isinstancerF   ZTensorrC   rN   itemrJ   stride
empty_liker;   tritonnext_power_of_2r   Zsave_for_backwardspdimsrS   maxlutr(   	rel_shaperel_stridesrB   	rel_dtypeis_denser)   )ctxr6   r(   r:   r)   rb   rS   rX   rc   rg   Mgridrd   re   r7   r	   r	   r   forward   s:    

z_softmax.forwardc                 C   s   | j \}}d }| jd r.tj| j| j|jd}|jd }| jd | jd | j	 |f}t
|}t| ||d||d||d| j||| jd | jd | jd | jd | j| j	t| j| jt| jd |d d |d d d d d d d d d d d d d d fS )N   rA   r   r   rD   r   r[   )Zsaved_tensorsZneeds_input_gradrF   zerosrd   rf   rC   rJ   rb   rS   r_   r?   r^   r(   re   r)   r`   ra   rc   rg   r   )rh   r=   r7   rX   Zdrri   rj   r>   r	   r	   r   backward   s6    






"


z_softmax.backwardN)__name__
__module____qualname__staticmethodrY   rk   rn   r	   r	   r	   r   r@      s   

%r@   c                   @   s(   e Zd Zd	ddZddddddZdS )
r   Fc                 C   s8   |j | _|| _|| _t| j| j|\| _| _|| _d S )N)	rJ   rb   rR   rS   r@   rY   rX   rc   rg   )selfrR   rS   rC   rg   r	   r	   r   __init__   s
    zsoftmax.__init__g      ?N)r(   r:   r)   c                C   sL   |d ur$|j |j kr$td|j  t||||| j| j| j| j| j	}|S )Nz$relative position embedding must be )	rB   
ValueErrorr@   applyrb   rS   rX   rc   rg   )rs   r6   r(   r:   r)   r	   r	   r   __call__   s    zsoftmax.__call__)F)ro   rp   rq   rt   rw   r	   r	   r	   r   r      s   
r   )rF   r`   Ztriton.languagelanguager   r   ZjitZ	constexprr;   r?   ZautogradZFunctionr@   r   r	   r	   r	   r   <module>   s   6	;Z