
    )`i                     r   d dl Z d dlmZ e j        d             Ze j        d             Ze j        d             Ze j        dej        dej        fd            Z	e j        dej        dej        fd            Z
e j        dej        dej        fd	            Ze j        dej        dej        fd
            ZdS )    Nc                 
   t          j        ||          }|t          j        ||z
            z  |t          j        ||z
            z  z   }| t          j        ||z
            z  |t          j        ||z
            z  z   } | ||fS N)tlmaximumexp2)omdother_oother_mother_dm_maxs          u/home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/flashinfer/triton/kernels/cascade.pystate_merger      s    Jq'""E	BGAI277U?+C+C!CCA	BGAI277U?+C+C!CCAeQ;    c                     | |z  } | ||fS r    r   r	   r
   s      r   state_normalizer      s    	AAa7Nr   c                 0    |t          j        |          z   S r   )r   log2r   s      r   state_get_lser      s    rwqzz>r   bdxbdyc
           
         t          j        d          }
t          j        |          D ]}t          j        |	          D ]}t          j        ||
|z  z   |z             }t          j        ||
|z  z   |z             }|
|z  |z   |z  |z   }t          j        | |z             }t          j        ||z             }t	          ||d||d          \  }}}t          |||          \  }}}|
|z  |z   |z  |z   }t          j        ||z   |           |r3t          j        ||
|z  z   |z   t          j        |          |z              d S )Nr   axis   )r   r	   r
   r   r   r   )r   
program_idrangeloadr   r   storer   )v_a_ptrs_a_ptrv_b_ptrs_b_ptrv_merged_ptrs_merged_ptr	num_headshead_dimr   r   postxhead_idxs_a_vals_b_valoffsetsv_av_bv_mergeds_maxr
   v_merged_offsets                         r   merge_state_kernelr6      s|    -Q


Chsmm   	 	Hggi7(BCCGggi7(BCCGY1X=BG''G+,,C''G+,,C!,AsGQ" " "HeQ "15!!D!DHeQ"Y9XEJOH\O3X>>>  3?2X=GAJJ&  !	 r   c	                    t          j        d          }	|rt          j        ||	z             dk    rd S t          j        |          D ]>}
t          j        ||	|z  z   |
z             }t          j        ||	|z  z   |
z             }t          j        ||          }t          j        ||z
            }t          j        ||z
            }|||z   z  }|||z   z  }t          j        |          D ]a}|	|z  |
z   |z  |z   }t          j        | |z             }t          j        ||z             }||z  ||z  z   }t          j        | |z   |           b|r6t          j        ||	|z  z   |
z   t          j        ||z             |z              @d S )Nr   r   )r   r   r!   r    r   r   r"   r   )v_ptrs_ptrv_other_ptrs_other_ptrr)   r*   mask_ptrr   r   r+   r-   s_vals_other_valr4   scaleother_scaler,   offsetv_vecv_other_vecs                       r   merge_state_in_place_kernelrD   =   s    -Q


C 78c>""a''FHSMM  i/(:;;gkC)O;hFGG
5+..&&gkE122,-!U[%89(3-- 	, 	,BIo0H<rAFGEFN++E'+"677KEMK+$==EHUV^U++++ 	Hi'(2+,,u4   r   c	                 b   t          j        d          }	t          j        |          D ]}
t          j        |          D ]}d\  }}}t          j        |          D ]f}t          j        ||	|z  |z   |z  z   |z             }t          j        | |	|z  |z   |z  |z   |z  z   |
z             }t	          |||||d          \  }}}gt          |||          \  }}}t          j        ||	|z  |z   |z  z   |
z   |           |r-t          j        ||	|z  z   |z   t          |||                     d S Nr   r   )g        g     jg      ?r   )r   r   r    r!   r   r   r"   r   )r8   r9   r'   r(   num_index_setsr)   r*   r   r   r+   r,   r-   r   r	   r
   itersvs                     r   merge_states_kernelrK   c   s    -Q


Chsmm   	 	H$GAq!00 	8 	8GS>1D8IEEP  Gn,t3y@8KxWX 
 &aAq!Q771aa%aA..GAq!H\S9_x%?8$KKbPRSTTT  3?2X=}QPQST?U?U  	 r   c	                    t          j        d          }	t          j        |          D ]F}
t          j        |          D ]-}d\  }}}t          j        t          j        ||	z             t          j        ||	z   dz                       D ]y}|                    t           j                  }t          j        |||z  z   |z             }t          j        | ||z  |z   |z  z   |
z             }t          |||||d          \  }}}zt          |||          \  }}}t          j        ||	|z  |z   |z  z   |
z   |           |r-t          j        ||	|z  z   |z   t          |||                     /Hd S rF   )
r   r   r    r!   toint64r   r   r"   r   )r8   r9   indptrr'   r(   r)   r*   r   r   r+   r,   r-   r   r	   r
   rH   iter_i64rI   rJ   s                      r   #variable_length_merge_states_kernelrQ      s    -Q


Chsmm   	 	H$GAq!#!6!6q@P8Q8QRR 8 87728,,GEHy$888CDDGEX	%9H%D$PPSUUVV%aAq!Q771aa%aA..GAq!H\S9_x%?8$KKbPRSTTT  3?2X=}QPQST?U?U  	 r   )tritontriton.languagelanguager   jitr   r   r   	constexprr6   rD   rK   rQ   r   r   r   <module>rW      sl               
    ! 
! 
! ! ! !H " 
" 
" " " "J  
 
   D  
 
     r   