o
    پif                     @   s  d Z ddlmZmZmZmZ ddlZddlZddlm	Z
 ddlmZ ejde
jde
jde
jde
jd	e
jd
e
jfddZdejdejdejdejfddZejde
jde
jde
jde
jd	e
jd
e
jfddZdejdejdejdejfddZejde
jde
jde
jde
jd	e
jf
ddZdejdejdejdejfddZdedeeee ee f fddZdee deeeef  fd d!ZdS )"zCommon utilities.    )AnyListOptionalTupleN)envsbuffer_stridenope_striderope_stridenope_dimrope_dimBLOCKc
                 C   s8  t d}
t d}||	 }|t d|	 }|| }||k }t ||
 t j}| ||  | }||	 |krEt j||
|  | |d}nM||kr[|| }t j||
|  | |d}n7||k }||k||| k @ }t j||
|  | ||@ dd}t j||
|  ||  ||@ dd}t |||}t j|||d d S )Nr      maskr   other)tl
program_idarangeloadtoint64wherestore)kv_buffer_ptrcache_k_nope_ptrcache_k_rope_ptrloc_ptrr   r   r	   r
   r   r   pid_locpid_blkbaseoffs	total_dimr   locdst_ptrsrc	offs_ropeis_nopeis_ropesrc_nopesrc_rope r+   N/home/ubuntu/.local/lib/python3.10/site-packages/sglang/srt/mem_cache/utils.pyset_mla_kv_buffer_kernel   sB   

r-   	kv_bufferr#   cache_k_nopecache_k_ropec           
      C   n   |j d }|j d }|| }d}| }|t||f}	t|	 | |||| d|d|d|||d
 d S N   r   )r   )shapenumeltritoncdivr-   stride
r.   r#   r/   r0   r
   r   r"   r   n_locgridr+   r+   r,   set_mla_kv_buffer_tritonU   $   


r=   c
                 C   s   t d}
t d}||	 }|t d|	 }|| }||k }t ||
 }| ||  | }||k }t j||
|  | ||@ dd}t j||
|  ||  || @ dd}|| }t j|||d d S )Nr   r   g        r   r   )r   r   r   r   r   )r   r   r   r   r   r   r	   r
   r   r   r   r   r    r!   r"   r   r#   r$   r'   r)   r*   r%   r+   r+   r,   set_mla_kv_scale_buffer_kernelp   s&   

r?   c           
      C   r1   r2   )r5   r6   r7   r8   r?   r9   r:   r+   r+   r,   set_mla_kv_scale_buffer_triton   r>   r@   c	                 C   s   t d}	t ||	 t j}
| |
|  }t d|}|| }t |}t ||	|  | | t d|}|| | }t |}t ||	|  | | d S )Nr   )r   r   r   r   r   r   r   )r   r   r   r   r   r   r	   r
   r   r   r#   loc_src_ptr	nope_offsnope_src_ptrnope_src	rope_offsrope_src_ptrrope_srcr+   r+   r,   get_mla_kv_buffer_kernel   s"   


rH   c              
   C   sT   |j d }|j d }| }|f}t| | |||| d|d|d||	 d S )Nr3   r   )r5   r6   rH   r9   )r.   r#   r/   r0   r
   r   r;   r<   r+   r+   r,   get_mla_kv_buffer_triton   s   

rI   devicereturnc                 C   s2   t j dur	dnd}|rddlm} || S dS )a@  
    Initialize custom memory pool based on environment variable.

    This function can be modified to support more features that require a custom memory pool.

    Args:
        device: The device to allocate memory on

    Returns:
        Tuple of (enable_custom_mem_pool, custom_mem_pool, custom_mem_pool_type)
    NTFr   )init_mooncake_custom_mem_pool)FNN)r   SGLANG_MOONCAKE_CUSTOM_MEM_POOLget(sglang.srt.disaggregation.mooncake.utilsrL   )rJ   enable_custom_mem_poolrL   r+   r+   r,   maybe_init_custom_mem_pool   s   rQ   tokensc                    sH   t  rt d tr S t  dk rg S  fddtt  d D S )Nr      c                    s    g | ]} |  |d   fqS )r   r+   ).0irR   r+   r,   
<listcomp>  s     z)convert_to_bigram_key.<locals>.<listcomp>r   )len
isinstancetuplerangerV   r+   rV   r,   convert_to_bigram_key  s
   r\   )__doc__typingr   r   r   r   torchr7   triton.languagelanguager   sglang.srt.environr   jit	constexprr-   Tensorr=   r?   r@   rH   rI   strboolrQ   intr\   r+   r+   r+   r,   <module>   s   	
;
	
'
	!

&