
    sg                     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_builder)r   inline_asm_elementwiseint64r
   s    S/var/www/html/venv/lib/python3.12/site-packages/triton/language/extra/cuda/utils.pyglobaltimerr      s.    &&'BD"TXT^T^hmtu08: :    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+    &&';T2TZZaelm08: :r   c                 Z    t        j                  | j                  j                  dz        S )N    r   	constexproptions	num_warpsr   s    r   num_threadsr      s"    >>(**44r9::r   c                 T    t        j                  | j                  j                        S Nr   r   s    r   r   r      s    >>(**4455r   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\\cgno 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};zmad.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 &&sHse4CSCS]ahi08: :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_float8r5   V   s    
xx""$0xH
==  "#t||hGJ88??""$(?(?(AAAL
xx #t||%Zbc.|y[cdLr   c                 "    t        | ||d|      S )NTr*   r5   r    r2   r)   r
   s       r   convert_custom_float8_sm80r9   f   s     f.Bd]effr   c                 "    t        | ||d|      S )NFr*   r7   r8   s       r   convert_custom_float8_sm70r;   k   s     f.Be^fggr   r   )NN)triton.languager   externr   r   builtinr   r   r!   r&   r5   r9   r;    r   r   <module>r@      s      : :
 : :
 ; ; 6 6     :  :F   g g h hr   