
    )`i.	                         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
            ZdS )    N)scale_and_clampEPS
BLOCK_SIZEHAS_IN_SCALEHAS_OUT_SCALE
HAS_OUTPUTHAS_RESIDUALreturnc                    t          j        d                              t           j                  }|||z  z   }|r|||	z  z   n|}|r|||z  z   nd }|rt          j        |          nd }|rt          j        |
          nd }t          j        |gt           j                  }t          d| |          D ]}|t          j        d|          z   }|| k     }t          j        ||z   |d                              t           j                  }|r||z  }|rVt          j        ||z   |d                              t           j                  }||z  }t          j	        ||z   ||           |||z  z  }t          j
        t          j        |          | z  |z             }|j        j        }t          d| |          D ]}|t          j        d|          z   }|| k     }|r7t          j        ||z   |                              t           j                  }n=t          j        ||z   |                              t           j                  }|r||z  }t          j        ||z   |                              t           j                  }|||z  z  } |rt          | ||          } t          j	        ||z   | |           d 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results!                                    r/home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/flashinfer/triton/kernels/norm.pyrms_norm_kernelr8      s   ( 	1  **A AL E$.9EAL  EE$0:EAL  dE&2<bgk"""G&3=bgk"""G :,bj999JQ:&&  	!Z000{GEGO$c:::==bjII 	LA 	4d#>>>AA"*MMAFAHUW_ad3333a!e

 (26*%%)C/
0
0C ;)LQ:&& 5 5	!Z000{ 	d33366rzBBAAd33366rzBBA WGEGO$///222:>> a#g 	D$VWlCCF
&t44444%5 5    )	tritontriton.languagelanguager   flashinfer.triton.kernels.quantr   jit	constexprr8    r9   r7   <module>rA      s           ; ; ; ; ; ; E5 
E5 E5 ,E5 <E5  !E5" ,#E5$ 
%E5 E5 E5 E5 E5 E5r9   