o
    (ºiÍ  ã                   @   s²   d dl Z d dlmZ e jdd„ ƒZe jdd„ ƒZe jdd„ ƒZe jdejd	ejfd
d„ƒZ	e jdejd	ejfdd„ƒZ
e jdejd	ejfdd„ƒZe jdejd	ejfdd„ƒZdS )é    Nc                 C   s^   t  ||¡}|t  || ¡ |t  || ¡  }| t  || ¡ |t  || ¡  } | ||fS ©N)ÚtlÚmaximumÚexp2)ÚoÚmÚdÚother_oÚother_mÚother_dÚm_max© r   ú^/home/ubuntu/veenaModal/venv/lib/python3.10/site-packages/flashinfer/triton/kernels/cascade.pyÚstate_merge   s   $$
r   c                 C   s   | | } | ||fS r   r   ©r   r   r   r   r   r   Ústate_normalize   s   
r   c                 C   s   |t  |¡ S r   )r   Úlog2r   r   r   r   Ústate_get_lse   s   r   ÚbdxÚbdyc
              
   C   s  t jdd}
t  |¡D ]s}t  |	¡D ]k}t  ||
|  | ¡}t  ||
|  | ¡}|
| | | | }t  | | ¡}t  || ¡}t||d||dd\}}}t|||ƒ\}}}|
| | | | }t  || |¡ |r}t  ||
|  | t  |¡| ¡ qq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_offsetr   r   r   Úmerge_state_kernel   s,   ÿþ€ðÿr0   c	                 C   s0  t jdd}	|rt  ||	 ¡dkrd S t  |¡D ]}}
t  ||	|  |
 ¡}t  ||	|  |
 ¡}t  ||¡}t  || ¡}t  || ¡}|||  }|||  }t  |¡D ]*}|	| |
 | | }t  | | ¡}t  || ¡}|| ||  }t  | | |¡ qU|r•t  ||	|  |
 t  || ¡| ¡ q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_valr.   ÚscaleÚother_scaler&   ÚoffsetÚv_vecÚv_other_vecr   r   r   Úmerge_state_in_place_kernel=   s2   þ€ñr=   c	              
   C   sü   t jdd}	t  |¡D ]p}
t  |¡D ]h}d\}}}t  |¡D ]0}t  ||	| | |  | ¡}t  | |	| | | | |  |
 ¡}t|||||dƒ\}}}qt|||ƒ\}}}t  ||	| | |  |
 |¡ |rzt  ||	|  | t|||ƒ¡ qqd S ©Nr   r   )g        g     jèÀg      ð?r   )r   r   r   r   r   r   r   r   )r1   r2   r!   r"   Únum_index_setsr#   r$   r   r   r%   r&   r'   r   r   r   ÚiterÚsÚvr   r   r   Úmerge_states_kernelc   s2   
ÿÿþÿ ÿ€ñÿrC   c	              
   C   s  t jdd}	t  |¡D ]{}
t  |¡D ]s}d\}}}t  t  ||	 ¡t  ||	 d ¡¡D ].}| t j¡}t  |||  | ¡}t  | || | |  |
 ¡}t|||||dƒ\}}}q+t|||ƒ\}}}t  ||	| | |  |
 |¡ |r…t  ||	|  | t	|||ƒ¡ qqd S r>   )
r   r   r   r   ÚtoÚint64r   r   r   r   )r1   r2   Úindptrr!   r"   r#   r$   r   r   r%   r&   r'   r   r   r   r@   Úiter_i64rA   rB   r   r   r   Ú#variable_length_merge_states_kernel†   s$   
( ÿ€öÿrH   )ÚtritonÚtriton.languageÚlanguager   Újitr   r   r   Ú	constexprr0   r=   rC   rH   r   r   r   r   Ú<module>   s@    


	÷
ö$ø	÷%ø	÷"ø	÷