o
    پi                     @   s^   d dl Z d dlm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dS )
    Nc                 C   s   | d|> d  |? S )N    )nlog2_kr   r   ]/home/ubuntu/.local/lib/python3.10/site-packages/triton_kernels/routing_details/_expt_data.py
_cdiv_pow2   s   r   SIZESBLOCKc                 C   s"  t d}||krw||| 7 }t j|g|jjd}	|t d| }
t |dkd|| d }td|d |D ]>}t d|| }||k }t j| | |dd}t	||}t 
|d|	 }|	t |d|jj7 }	t |
||  |
|7 }
q6d S ||d 8 }|||  t d| }t |d d S )Nr   )dtyper   )maskotherl    )tl
program_idzerosr
   
element_tyarangewhererangeloadr   cumsumsumtostore)Histn_expts_totMDStartstile_starts_stridem
MDTileInfofirst_tile_dim_log2r   r	   pidx_tile	Tile_ptrstile_dim_log2ioffs_nmask_n0hist_tok	hist_tiletile_startsTileInfoOutr   r   r   _expt_data_memset
   s&   


r*   c                 C   s   t d}|| }	|| }
||
| 7 }||
| 7 }t | |	 }||
 }t||}t ||	 }||7 }td||D ]}|t d| }|d> |	 }t j|| |||k d q:d S )Nr      )r   )r   r   r   r   r   r   r   )r   MDTileStartsr   r   tile_info_stridemr   r   r	   r   expt_idbuff_idn_tokensr"   n_blockstile_off	block_off
block_offsdatar   r   r   _expt_data_compute*   s   

r6   )	tritontriton.languagelanguager   jitr   	constexprr*   r6   r   r   r   r   <module>   s     
