Ë
    ÛéñiJ  ã                   óf  — d dl mZ ej                  dd„«       Zej                  dd„«       Zej
                  dd„«       Zej
                  dd„«       Zej
                  dd„«       Zej
                  dd„«       Z	ej
                  dd	„«       Z
ej
                  dd
„«       Zej
                  dd„«       Zy)é    )ÚcoreNc           	      óV   — t        j                  ddg t         j                  dd| ¬«      S )Nzmov.u64 $0, %globaltimer;z=lFé   ©ÚdtypeÚis_pureÚpackÚ	_semantic)r   Úinline_asm_elementwiseÚint64©r
   s    úb/var/www/pod-logistic/pod-ai/venv/lib/python3.12/site-packages/triton/language/extra/cuda/utils.pyÚglobaltimerr      s.   € ä×&Ñ&Ð'BÀDÈ"ÔTX×T^ÑT^ÐhmÐtuØ1:ô<ð <ó    c           	      óV   — t        j                  ddg t         j                  dd| ¬«      S )Nzmov.u32 $0, %smid;z=rTr   r   )r   r   Úint32r   s    r   Úsmidr   
   s+   € ä×&Ñ&Ð';¸TÀ2ÌTÏZÉZÐaeÐlmØ1:ô<ð <r   c                 ón   — t        j                  | j                  j                  j                  dz  «      S )Né    ©r   Ú	constexprÚbuilderÚoptionsÚ	num_warpsr   s    r   Únum_threadsr      s(   € ä�>‰>˜)×+Ñ+×3Ñ3×=Ñ=ÀÑBÓCÐCr   c                 óh   — t        j                  | j                  j                  j                  «      S ©Nr   r   s    r   r   r      s#   € ä�>‰>˜)×+Ñ+×3Ñ3×=Ñ=Ó>Ð>r   c           	      óX   — t        j                  dd| gt         j                  dd|¬«      S )Na¸  {                                      
.reg .b32 a<2>, b<2>;                  
prmt.b32 a0, 0, $2, 0x5746;            
and.b32 b0, a0, 0x7f007f00;            
and.b32 b1, a0, 0x00ff00ff;            
and.b32 a1, a0, 0x00800080;            
shr.b32  b0, b0, 1;                    
add.u32 b1, b1, a1;                    
lop3.b32 $0, b0, 0x80008000, a0, 0xf8; 
shl.b32 $1, b1, 7;                     
}                                      
z=r,=r,rTé   r   )r   r   Úfloat16)Úargr
   s     r   Úconvert_fp8e4b15_to_float16r"   !   s2   € ä×&Ñ&ð
	4ð 6?ÀÀÌTÏ\É\ÐcgÐnoØôð r   c           	      ó€   — d}|r|dz  }n|dz  }|dz  }t        j                  |d| gt         j                  dd|¬«      S )	NaN  {
            .reg .pred p<4>;
            .reg .b32 a<2>, b<2>;
            .reg .b16 c<4>;
            .reg .b16 max_val_f16;
            .reg .b32 max_val_f16x2;
            mov.b16 max_val_f16,   0x3F00;
            mov.b32 max_val_f16x2, 0x3F003F00;
            and.b32 a0, $1, 0x7fff7fff;
            and.b32 a1, $2, 0x7fff7fff;zSmin.f16x2 a0, a0, max_val_f16x2;
                  min.f16x2 a1, a1, max_val_f16x2;a×  setp.lt.f16x2  p0|p1, a0, max_val_f16x2;
                  setp.lt.f16x2  p2|p3, a1, max_val_f16x2;
                  mov.b32 {c0, c1}, a0;
                  mov.b32 {c2, c3}, a1;
                  selp.b16  c0, c0, max_val_f16, p0;
                  selp.b16  c1, c1, max_val_f16, p1;
                  selp.b16  c2, c2, max_val_f16, p2;
                  selp.b16  c3, c3, max_val_f16, p3;
                  mov.b32 a0, {c0, c1};
                  mov.b32 a1, {c2, c3};zömad.lo.u32 a0, a0, 2, 0x00800080;
              mad.lo.u32 a1, a1, 2, 0x00800080;
              lop3.b32 b0, $1, 0x80008000, a0, 0xea;
              lop3.b32 b1, $2, 0x80008000, a1, 0xea;
              prmt.b32 $0, b0, b1, 0x7531;
              }z=r,r,rTr   r   )r   r   Úfloat8e4b15)r!   Ú	has_minx2r
   Úasms       r   Úconvert_float16_to_fp8e4b15r'   2   sm   € ð	+€Cñ Øð 6ñ 	6‰ð 	ð 	+ñ 		+ˆð ð ñ €Cô ×&Ñ& s¨H°s°eÄ4×CSÑCSÐ]aÐhiØ1:ô<ð <r   c                 ó"  — | j                   j                  j                  «       rJt        | |¬«      }|j                  j	                  «       r!|j                  t        j                  |¬«      }|S | j                   j                  j                  «       s&| j                   j                  j	                  «       sJ ‚| }| j                   j                  j	                  «       r"|j                  t        j                  d|¬«      }t        |||¬«      }|S )Nr   Úrtz)Úfp_downcast_roundingr
   ©r%   r
   )ÚtypeÚscalarÚis_fp8e4b15r"   Úis_fp32Útor   Úfloat32Úis_fp16r    r'   )r!   Údst_tyr*   r%   r
   Ú
upcast_valÚdowncast_vals          r   Úconvert_custom_float8r6   V   sÂ   € à
‡x�x‡�×"Ñ"Ô$Ü0°À	ÔJˆ
Ø�=‰=× Ñ Ô"Ø#Ÿ™¤t§|¡|¸y˜ÓIˆJØÐà�8‰8�?‰?×"Ñ"Ô$¨¯©¯©×(?Ñ(?Ô(AÐAÐAØ€LØ
‡x�x‡�×ÑÔ Ø#—‘¤t§|¡|È%Ð[d�ÓeˆÜ.¨|ÀyÐ\eÔf€LØÐr   c                 ó"   — t        | ||d|¬«      S )NTr+   ©r6   ©r!   r3   r*   r
   s       r   Úconvert_custom_float8_sm80r:   f   s   € ä   fÐ.BÈdÐ^gÔhÐhr   c                 ó"   — t        | ||d|¬«      S )NFr+   r8   r9   s       r   Úconvert_custom_float8_sm70r<   k   s   € ä   fÐ.BÈeÐ_hÔiÐir   r   )NN)Útriton.languager   Úexternr   r   Úbuiltinr   r   r"   r'   r6   r:   r<   © r   r   ú<module>rA      sç   ðÝ  ð ‡�ò<ó ð<ð
 ‡�ò<ó ð<ð
 ‡�òDó ðDð ‡�ò?ó ð?ð ‡�òó ðð  ‡�ò <ó ð <ðF ‡�òó ðð ‡�òió ðið ‡�òjó ñjr   