o
    پiN                     @   s   d dl Z d dlZd dlm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dejfd
dZ
ejdejfddZdd ZdS )    Nrreturnc                 C   s$   |  tj} | |> | d| ? B d@ S )zY
    rotate left 32-bit integer x by r bits
    e.g. x = 01110001, r = 2 -> 11000101
            )totluint64)xr    r
   P/home/ubuntu/.local/lib/python3.10/site-packages/sglang/srt/layers/utils/hash.pyrotl32   s   r   hc                 C   s@   | | d? N } | d d@ } | | d? N } | d d@ } | | d? N } | S )z7
    final mix of 32-bit hash value for MurmurHash
       l   kJ r      l   5.e r
   )r   r
   r
   r   fmix32   s   r   kc                 C   s`   d}d}d}d}d}d}|| d@ }t ||}|| d@ }| |N } t | |} | | | d@ } | S )z1
    Mixes a 32-bit key into the hash state.
    l   Q-< i5   r      l   dkL r   )r   )r   r   c1c2r1r2mmnnr
   r
   r   murmur3_mix   s   

r   
BLOCK_SIZEc                 C   s   t d}t d}|}	|| t d| }
|
|k }t | |	 t j}t ||	 t j}t j||
 |ddt j}d}|d@ t j}t||}|d? d@ t j}t||}t||}t||}|dN }t|}t j	||	|  |
 ||d dS )	ab  
    MurmurHash 32-bit implementation for Triton.
    Reference:
    - https://medium.com/@thealonemusk/murmurhash-the-scrappy-algorithm-that-secretly-powers-half-the-internet-2d3f79b4509b
    - https://en.wikipedia.org/wiki/MurmurHash

    We treat 64-bit seed, 32-bit position, and 32-bit col_index as 4 4-byte blocks, and bit-blend them together.
    r      )maskotherr   r   r   )r   N)
r   
program_idarangeloadr   r   uint32r   r   store)seed_ptrpositions_ptrcol_indices_ptr
output_ptrnum_rowsnum_colsr   pid_rowpid_colrow_idxcol_offsetsr   seedposcolr   r   r
   r
   r   murmur_hash32_kernel2   s$   





 r1   c           	   	   C   s   | j |j ks
J dt| j dkrt|j dks$J d| j d|j | j d }|j d }| j}tj||ftj|d}d}|t||f}t| | ||||||d |S )	Nz0Seed and positions must have the same shape (n,)r   z%Inputs must be 1D tensors seed.shape=z col_indices.shape=r   )dtypedevicei   )r   )	shapelenr3   torchemptyr"   tritoncdivr1   )	r.   	positionscol_indicesnmr3   hashedr   gridr
   r
   r   murmur_hash32h   s    

r@   )r6   r8   triton.languagelanguager   jit	constexprr"   r   r   r   r1   r@   r
   r
   r
   r   <module>   s    	5