o
    i.	                     @   s\   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d	dfd
dZdS )    N)scale_and_clampEPS
BLOCK_SIZEHAS_IN_SCALEHAS_OUT_SCALE
HAS_OUTPUTHAS_RESIDUALreturnc           !      C   s  t jddt j}|||  }|r|||	  n|}|r"|||  nd }|r+t |nd }|r4t |
nd }t j|gt jd}td| |D ]F}|t d| }|| k }t j|| |ddt j}|rg||9 }|rt j|| |ddt j}||7 }t j	|| ||d ||| 7 }qEt 
t ||  | }|jj}td| |D ]V}|t d| }|| k }|rt j|| |dt j}nt j|| |dt j}|r||9 }t j|| |dt j}|||  } |rt| ||} t j	|| | |d qd S )Nr   )axis)dtypeg        )maskother)r   )tl
program_idtoint64loadzerosfloat32rangearangestorersqrtsumr   
element_tyr   )!nbx_ptrx_stridex_scale_ptrr_ptrr_stridew_ptro_ptro_strideo_scale_ptrr   r   r   r   r   r   ix_rowo_rowr_rowx_scaleo_scale
square_sumoffoffsetsr   xrrmsoutput_dtypewresult r5   T/home/ubuntu/vllm_env/lib/python3.10/site-packages/flashinfer/triton/kernels/norm.pyrms_norm_kernel   sD   r7   )	tritontriton.languagelanguager   flashinfer.triton.kernels.quantr   jit	constexprr7   r5   r5   r5   r6   <module>   s&    