
    VpfYJ                    2   d Z ddlmZ ddlZddlmZ ddlZddlmZ ddlm	Z
 ddlmZ ddlZd e ej         ej        d                    j                  z  Zd8dZd9dZ ej        ej        g d           ej        ej        g d          	 	 	 	 	 	 	 	 	 	 d:d;d0                        Zd<d1Zd=d2Z ej        d3          d>d4            Zd?d5Zd@d6Ze                     ee            ej        ej        ddg          	 	 dAdBd7            Z!dS )Cz<Module containing fused attention forward and backward pass.    )annotationsN)Any)lax)pallasgfffffffloat32segment_ids_refjax.Array | Noneo_refr   residual_refs	num_headsintsm_scalefloatcausalboolblock_qblock_dblock_kc                 	
 | j         d         }t          j        d          t          j        t          j                  t          d          z
  }t          j        t          j                  }t          j        	ft          j                  }t          j        z            }t          j        | |t          j        d           f          d nt          j        |f          	
fd}r"t          j
        dz   z  
z   dz
  
          }nt          j        |
          }t          j        d|||||f          \  }}}|r3|\  }}t          j        ||f|           t          j        ||f|           |                    |j                  }t          j        ||t          j        d           f|           d S )Nr   dtypeinfc                P   |\  }}}t          j        | z            }t          j        |t          d           f          } d nt          j         |f          }t          j        |j                  }!dk    r|!z  }s d }	 t          |          }	rg"z  t          j                  z   }
| z  t          j                  z   }|
d d d f         |d d d f         k    }|	|nt          j	        |	|          }	t          j
        |	|t                    }|                    d          }t          j        ||          }t          j        ||z
            }||z  }t          j        ||d d d f         z
            }|                    d          }||z   }d|z  }||d d d f         z  }||z  d d d f         |z  }t          j        #|t          j                  f          }t          j        |                    |j                  |          }||z   }|||fS )N      ?axis)pldsliceloadslicedotTsegment_maskjnparangelogical_andwhereDEFAULT_MASK_VALUEmaxmaximumexpsumastyper   )$start_kcarryo_prevm_prevl_prevcurr_k_slicekkv_segment_idsqkmaskspan_qspan_kcausal_maskm_currm_next
correctionl_prev_corrs_currl_currl_next
l_next_rcpo_prev_corrvo_curro_nextr   r   r   r   k_refqq_segment_idsr   r   start_qv_refs$                            i/var/www/html/nettyfy-visnx/env/lib/python3.11/site-packages/jax/experimental/pallas/ops/gpu/attention.pybodyz mha_forward_kernel.<locals>.bodyD   sW   "FFF9Ww.88L
eDkk233A " 	W_|o66 
 
13B2~~Hnb  3,d		$M>::	 
7"SZ%8%887"SZ%8%88QQQWoaaa8<KKS_T;-O-O 	 9T2122bVVV__F[((F&))Jv%KW
VAAAtG_ F ZZRZ  F6!FfJjD))F+QQQW5>K
bi&8&89::AVFMM!'**A..F6!F66!!       )shaper   
program_idr%   zerosr   r   r   r    r   divcdiv	fori_loopstorer.   r   )q_refrH   rL   r   r
   r   r   r   r   r   r   r   seq_lenm_il_iocurr_q_slicerN   upper_boundl_refm_refrI   rJ   rK   s    ```  `````          @@@rM   mha_forward_kernelra      s    KN'M!' 		'---e<#	'---#	i'"#+666!
 7W,g66,gelBIdOO455! 
	  d7?\O44 1" 1" 1" 1" 1" 1" 1" 1" 1" 1" 1" 1" 1" 1" 1"d  ,''Wq[1G;a?IIKK''7++KadQSMBB+!S# * LE5HU\OS)))HU\OS)))hhu{!(5<41155555rO   rJ   	jax.Arrayr6   c                   t          j        | d          } |j        dk    rt          j        |d          }nt          j        |d          }t          j        | |                              t           j                  S )Nr   r   rP   r   )r%   expand_dimsndimequalr.   bool_)rJ   r6   s     rM   r$   r$      sr    
 /-b999-A_^!<<<NN_^!<<<N	=.	1	1	8	8	C	CCrO   )
               	   
            )nondiff_argnums)
r   r   r   r   backward_pass_impl	num_warps
num_stagesgrid	interpretdebug)static_argnamesr   F   triton   segment_idsjnp.ndarray | Noners   strrt   
int | Noneru   rv   tuple[int, ...] | Nonerw   rx   c                   ~| j         \  }}}}t          ||          }t          ||          }|}|t          j        ||          ||f}|	}|
|dk    rdnd}t	          j        t          ||||||          }t          j        d |d |fd           t          j        d |d |fd           t          j        d |d |fd           g}|                    |d nt          j        d |fd                      t          j
        | j         | j        	          } t          j        |||t          j        d |d |fd
           t          t          ||
                    |||d	  	        | |||          S )N@   rh   rl   )r   r   r   r   r   r   c                    |d|dfS Nr    _jr5   s      rM   <lambda>zmha.<locals>.<lambda>       Aq!Q< rO   c                    |d|dfS r   r   r   s      rM   r   zmha.<locals>.<lambda>   r   rO   c                    |d|dfS r   r   r   s      rM   r   zmha.<locals>.<lambda>   r   rO   c                
    |dfS r   r   r   s      rM   r   zmha.<locals>.<lambda>   
    !Q rO   rQ   r   c                    |d|dfS r   r   r   s      rM   r   zmha.<locals>.<lambda>   r   rO   rt   ru   r{   mha_forwardrv   in_specs	out_specscompiler_params	out_shaperx   rw   name)rQ   minr   rU   	functoolspartialra   	BlockSpecappendjaxShapeDtypeStructr   pallas_calldict)rI   r5   rE   r}   r   r   r   r   rs   rt   ru   rv   rw   rx   
batch_sizerY   r   head_dimgrid_
num_warps_kernelr   r   s                          rM   mhar      s   D -.W**gy(!!'!!'
%
]WWg&&
I>E*"nn!J/9&.%,h$*, , ,& 	l$
)+G+G  	l$
)+G+G  	l$
)+G+G 
( //		 d<w)?)?@@  
 "@@@)
$
)+G+G  
zBBB   
 
 
 q![
 
 rO   c                   ~| j         \  }}}}t          ||          }t          ||          }|}|t          j        ||          ||f}|	}|
|dk    rdnd}t	          j        t          ||||||          }t          j        | j         | j	                  t          j        |||ft          j                  t          j        |||ft          j                  g}t          j        d |d |fd           t          j        d |d |fd           t          j        d |d |fd           g}|                    |d nt          j        d |fd	                       t          j        |||t          j        d |d |fd
           t          j        d d |fd           t          j        d d |fd           gt          t          ||
                    |||d	  	        | |||          \  }}}|| ||||||ffS )Nr   rh   rl   )r   r   r   r   r   r   r   c                    |d|dfS r   r   r   s      rM   r   z_mha_forward.<locals>.<lambda>  r   rO   c                    |d|dfS r   r   r   s      rM   r   z_mha_forward.<locals>.<lambda>  r   rO   c                    |d|dfS r   r   r   s      rM   r   z_mha_forward.<locals>.<lambda>  r   rO   c                
    |dfS r   r   r   s      rM   r   z_mha_forward.<locals>.<lambda>  r   rO   c                    |d|dfS r   r   r   s      rM   r   z_mha_forward.<locals>.<lambda>%  s    1a| rO   c                    ||dfS r   r   r   s      rM   r   z_mha_forward.<locals>.<lambda>'      q!Qi rO   c                    ||dfS r   r   r   s      rM   r   z_mha_forward.<locals>.<lambda>(  r   rO   r   r   r   r   )rQ   r   r   rU   r   r   ra   r   r   r   r%   r   r   r   r   r   )rI   r5   rE   r}   r   r   r   r   rs   rt   ru   rv   rw   rx   r   rY   r   r   r   r   r   r   r   outlms                             rM   _mha_forwardr      s     -.W**gy(!!'!!'
%
]WWg&&
I>E*"nn!J/9&.vw%,h@ @ @& 
888	*i!A!$. . .	*i!A!$. . .	) 	l$
)+G+G  	l$
)+G+G  	l$
)+G+G 
( //		 d<w)?)?@@  
bn
,WdH-/K/K  ,dG,.G.G
H
H
,dG,.G.G
H
H 
zBBB   #  $ q![% )#q!& 
q!QS!Q/	//rO   c                  t          j        d          }t          j        ||z  |          }t          j        | |t	          d           f                              t          j                  }t          j        ||t	          d           f                              t          j                  }	t          j        ||f                              t          j                  }
|	|
d d d f         z  }	t          j        ||	z  d          }t          j	        ||t	          d           f|	                    |j
                             t          j	        ||f|                    |j
                             d S )Nr   rP   r   )r   rR   dsr    r!   r.   r%   r   r-   rW   r   )out_refdout_refr_   new_dout_ref	delta_refr   pid_moff_mr\   dodenomdeltas               rM   _preprocess_backward_kernelr   5  s0    -

%
%
)
)%gguT{{+,,33CK@@!	wx%t-..55ckBB"
'%%
"
"
)
)#+
6
6%	E!!!T'N"
'!b&q
!
!
!%(<%t-99\'((* * *(9uhY_ = =>>>>>rO   preprocess_backwardc                   | j         \  }}}}	t          j        |j         |j                  t          j        |j         |j                  g}
 t	          j        t          j        t          |          t	          j	        ||          ||ft	          j
        d |d |	fd           t	          j
        d |d |	fd           t	          j
        d d |fd           gt	          j
        d |d |	fd           t	          j
        d d |fd           gt          t          dd	          
          |
||d	  	        | ||          \  }}||fS )N)r   c                    |d|dfS r   r   r   s      rM   r   z&_preprocess_backward.<locals>.<lambda>S      q!QPQl rO   c                    |d|dfS r   r   r   s      rM   r   z&_preprocess_backward.<locals>.<lambda>T  r   rO   c                    ||dfS r   r   r   s      rM   r   z&_preprocess_backward.<locals>.<lambda>U      Q1I rO   c                    |d|dfS r   r   r   s      rM   r   z&_preprocess_backward.<locals>.<lambda>X  r   rO   c                    ||dfS r   r   r   s      rM   r   z&_preprocess_backward.<locals>.<lambda>Y  r   rO   rh      r   r   mha_preprocess_backwardr   )rQ   r   r   r   r   r   r   r   r   rU   r   r   )r   r   r   r   rx   rw   r   rY   r   r   r   	do_scaledr   s                rM   _preprocess_backwardr   G  sr    .1Y**gy(	28RX..	17AG,,)&R^3WEEEGGW%%z9=
dGT846R6RSS
dGT846R6RSS
dD'*,E,EFF 	dGT846R6RSS
dD'*,E,EFF a000   $%& & &$ '*2q%2 2)U& 
E	rO   c                   
 ~~ j         d         
 fd}t          j        dt          j                  |d            d S )Nr   c                  	 t          j        
gt           j                  }t          j        
gt           j                  }t          j        t          j        | z            t          d           f          t          j        t          j        | z            t          d           f          	| z  t          j                  z   d n+t          j        t          j        | z            f          	fd}rt          j	        | z            }nd}t          j
        |t          j                  |||f          \  }}t          j        t          j        | z            t          d           f|                    j                             t          j        t          j        | z            t          d           f|                    j                             d S )Nr   c                   |\  }}t          j        t          j        | z            t          d           f          }t          j        |j                  }|                    j                  }|                    t          j	                  }dk    r|z  }d n+t          j        t          j        | z            f          }s~d }t          |          }rM| z  t          j                  z   }|d d d f         d d d f         k    }	||	nt          j        ||	          }t          j        ||t                    }t          j        t          j        | z            f          }
t          j        ||
d d d f         z
            }t          j        t          j        | z            t          d           f          }|t          j        |                    |j                  j        |          z   }t          j        t          j        | z            f          }t          j        ft          j	                  |d d d f         z
  }|t          j        |j                  z   }||z  }dk    r|z  }|t          j        |                    j                  j        |          z   }t          j        t          j        | z            t          d           fd          }|t          j        |                    j                                                |j                  z   }t          j        t          j        | z            t          d           f|d           ||fS )Nr   r   
evict_last)eviction_policy)r   r    r   r!   r"   r#   r.   r   r%   r   r$   r&   r'   r(   r)   r,   rS   rW   )rK   r0   dvdkrI   r7   rJ   r8   r9   r;   r   pr   didpr   dqr   r   r   r   do_scaled_refdq_refr5   r6   r`   rX   r   r   r:   rE   s                    rM   
inner_loopz;mha_backward_kernel.<locals>.outer_loop.<locals>.inner_loop  s_   fb"
'%"%' 17;;U4[[I
J
Ja6!QS>>b99U[!!b99S[!!b	S
h $ $ww/@')J)J(LMM  
 5?.&m^<<$ 	W$sz'':'::&qqq$w6$'?:+  k?455 
 YtR!344
'%"%' 17;;=
>
>a
'"qDz/
"
"a7=257):G#D#DeDkk"RSSbqxx))+R000b79ruWw%6@@BCCb9gw's{;;;bDkIbr13br6b	S(]ryy--/333b76BE'G"3W==!$KK):FH H Hbryy))1--44RX>>>bhvg/99d%&(,H H H HVmrO   r   )r%   rS   r   r   r    r   r!   r&   r   rT   rV   rU   rW   r.   r   )r/   r   r   r   r   lower_boundr5   r6   r:   rE   r   r   r   r   r   dk_refr   r   dv_refrH   r`   rX   r   rY   r   rL   s         @@@@rM   
outer_loopz'mha_backward_kernel.<locals>.outer_loop  s1   	GW%S[	9	9	9B	GW%S[	9	9	9B
g/995;;GHHA
g/995;;GHHAwG!4!44F " 	W_ruWw->'H'H&JKK . . . . . . . . . . . . . . . . . .^  GGg-w77kkk];(A(A:8% %FBHVbeGg-w77Dkk#$&IIfl$;$;= = =HVbeGg-w77Dkk#$&IIfl$;$;= = = = =rO   )rQ   r   rV   r   rU   )rX   rH   rL   r   r   r   r_   r`   r   r   r   r   r   r   r   r   r   r   r   rY   s   ```` ` `` ```````` @rM   mha_backward_kernelr   e  s    . uKN'E= E= E= E= E= E= E= E= E= E= E= E= E= E= E= E= E= E= E= E=L -277G,,j$?????rO   c                x   ~~~|
\  }}}}}}}|dk    rA t          j        t          j        t          | |          ||||          d         |          S |dk    r|j        \  }}}}t          ||          }t          ||          }t          |||||	|          \  }}t          j	        |j        t          j
                  }t          j        |j        |j                  t          j        |j        |j                  t          j        |j        |j                  g}t          j        d |d |fd           t          j        d |d |fd           t          j        d |d |fd           t          j        d |d |fd           t          j        d |d |fd	           t          j        d d |fd
           t          j        d d |fd           t          j        d d |fd           t          j        d |d |fd           g	}||                    dd            ddi}n0|                    dt          j        d |fd                      ddi}||f}d} t          j        t          j        t"          |||| |          |||t          j        d |d |fd           t          j        d |d |fd           t          j        d |d |fd           gd|	|t%          t%          |d                    |
  
        ||||||||||
  
        \  }}}nt'          d|           |                    |j                  ||d fS )Nxla)r   r   rP   r{   c                    | d|dfS r   r   r   r5   s     rM   r   z_mha_backward.<locals>.<lambda>      1aA, rO   c                    | d|dfS r   r   r   s     rM   r   z_mha_backward.<locals>.<lambda>  r   rO   c                    | d|dfS r   r   r   s     rM   r   z_mha_backward.<locals>.<lambda>  r   rO   c                    | d|dfS r   r   r   s     rM   r   z_mha_backward.<locals>.<lambda>  r   rO   c                    | d|dfS r   r   r   s     rM   r   z_mha_backward.<locals>.<lambda>  r   rO   c                    | |dfS r   r   r   s     rM   r   z_mha_backward.<locals>.<lambda>      !Q rO   c                    | |dfS r   r   r   s     rM   r   z_mha_backward.<locals>.<lambda>  r   rO   c                    | |dfS r   r   r   s     rM   r   z_mha_backward.<locals>.<lambda>  r   rO   c                    | d|dfS r   r   r   s     rM   r   z_mha_backward.<locals>.<lambda>  r   rO   r   rl   r   c                
    | dfS r   r   r   s     rM   r   z_mha_backward.<locals>.<lambda>  s
    QF rO   rm   )r   r   r   r   r   c                    | d|dfS r   r   r   s     rM   r   z_mha_backward.<locals>.<lambda>      q!Ql rO   c                    | d|dfS r   r   r   s     rM   r   z_mha_backward.<locals>.<lambda>  r   rO   c                    | d|dfS r   r   r   s     rM   r   z_mha_backward.<locals>.<lambda>  r   rO   mha_backwardr   r   )	rv   r   r   r   r   rx   rw   r   input_output_aliasesz&Invalid backward pass implementation: )r   vjpr   r   mha_referencerQ   r   r   r%   rS   r   r   r   r   r   insertr   r   r   
ValueErrorr.   )r   r   r   r   rs   rt   ru   rv   rw   rx   resr   rI   r5   rE   r}   r   r   r   r   rY   r   r   r   r   r   
out_shapesr   r   r   r   s                                  rM   _mha_backwardr     s    T$'!!Q;Q5  	37-(6JJJ			  	 
   X%%/0w,JH'7##G'7##G+CQ	RRIu	17CK	(	(B	28RX..	17AG,,	17AG,,J 	7D(+-F-F	
 	
 	7D(+-F-F	
 	
 	7D(+-F-F	
 	
 	7D(+-F-F	
 	
 	7D(+-F-F	
 	
 	dD'*,B,BCC
dD'*,B,BCC
dD'*,B,BCC
7D(+-F-F	
 	
'H. ooaVooatWo7J7JKKLLLV	"DI	
 	
 	
 Lwh/1J1J  Lwh/1J1J  Lwh/1J1J 

 D9$K$K$KLLL17  8 A{CAq%9= =JBBB< R>PRR
S
SS	17		RT	))rO   c                   | j         d         }|j         d         }t          j        d| |                              t          j                  }d }	|=t          j        t          ||          d          }	t          j        |	|j                   }	|rdt          j        t          j	        dd||ft                              }
t          j        |
|j                   }
|	|
nt          j        |	|
          }	|	|n"t          j        |	|t          d                    }t          j                            ||z                                | j                  }t          j        d||          S )NrP   zbqhc,bkhc->bhqkr   z-infzbhqk,bkhc->bqhc)rQ   r%   einsumr.   r   rd   r$   broadcast_totrilonesr   r'   r(   r   r   nnsoftmaxr   )rI   r5   rE   r}   r   r   	q_seq_len
kv_seq_lenlogitsr8   r;   weightss               rM   r   r   '  s3    gaj)wqz*:'A..55ckBB&	$?<[AA1EEDD&,//D O(38Q9j$ANNNOOK";==K,;;COD+,N,ND\66syvuV}}'M'M&FNN6H,--44QW=='	%w	2	22rO   )r   r	   r
   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   )rJ   rb   r6   rb   )
r   Frz   rz   r{   Nr|   NFF)r}   r~   r   r   r   r   r   r   r   r   rs   r   rt   r   ru   r   rv   r   rw   r   rx   r   )r}   r	   r   r   r   r   r   r   r   r   rs   r   rt   r   ru   r   rv   r   rw   r   rx   r   )r   r   )r   r   rx   r   rw   r   )r   r	   r   r   r   r   r   r   r   r   r   r   )r   r   r   r   r   r   r   r   rs   r   rt   r   ru   r   rv   r   rw   r   rx   r   )r   F)r}   r~   r   r   )"__doc__
__future__r   r   typingr   r   r   jax.experimentalr   r   	jax.numpynumpyr%   npr   finfor   r*   r)   ra   r$   r   
custom_vjpjitr   r   r   named_scoper   r   r   defvjpr   r   rO   rM   <module>r     s   C B " " " " " "           



       ) ) ) ) ) )          EE("(828I+>+>"?"?"CDDD e6 e6 e6 e6PD D D D N$F$F$F   G    ( & #'@ @ @ @  $@FI0 I0 I0 I0X? ? ? ?$ &''   (':`@ `@ `@ `@F[* [* [* [*x 

< ' ' ' 37Z,BCCC 3 3 3 3 DC3 3 3rO   