
    "h                        d dl Z d dlZd dlZd dlZd dlZd dlmZ d dl	Z	d dl
mZ d dlmZmZ d dlmZmZ g dZg dZeez   Zg dZeez   Zedgz   Zd	d
gZdgez   dgz   ez   dgz   Zd Zd Zd Zd Zd Zddee   fdZddej@                  deeejB                  f   fdZ"de#fdZ$d Z%d Z&e	jN                  jQ                   e&        d      Z)y)    N)RandomState)OptionalUnion)TensorWrapperreinterpret)int8int16int32int64)uint8uint16uint32uint64)float16float32float64bfloat16float8_e4m3fnfloat8_e5m2boolr   c                  H    t         j                  j                  dd      dk(  S )NTRITON_INTERPRET01)osenvironget     U/var/www/html/sandstorm/venv/lib/python3.12/site-packages/triton/_internal_testing.pyis_interpreterr!      s    ::>>,c2c99r   c                  |    t               ry t        j                  j                  j                  j                         S N)r!   tritonruntimedriveractiveget_current_targetr   r   r    r(   r(      s*    >>  ''::<<r   c                  <    t               } | dS | j                  dk(  S )NFcudar(   backendtargets    r    is_cudar/   !   s"    !FN5@&(@@r   c                  <    t               } | dS | j                  dk(  S )NFhipr+   r-   s    r    is_hipr2   &   s"    !FN5?%(??r   c                  H    t               } | dS t        | j                        S )N )r(   strarchr-   s    r    get_archr7   +   s"    !F25S%55r   rsc                 F   t        | t              r| f} |t        d      }|t        t        z   v rt        j                  t        t
        |            }||j                  nt        ||j                        }||j                  nt        ||j                        }t        t
        |      }|j                  ||| |      }d||dk(  <   |S |r)d|v r%|j                  dd| t
        j                        }|S |t        v r"|j                  dd|       j                  |      S |d	k(  rV|j                  dd|       j                  d
      j                  d      t        j                   d      z  j                  d
      S |dv r|j                  dd|       dkD  S t#        d|       )zp
    Override `rs` if you're calling this function twice and don't want the same
    result for both calls.
       )seed)dtype   r   float8   (   r   r   r   l      )r   int1bool_g        zUnknown dtype )
isinstanceintr   
int_dtypesuint_dtypesnpiinfogetattrminmaxrandintr   float_dtypesnormalastypeviewr   RuntimeError)shape	dtype_strr8   lowhighrH   r<   xs           r    numpy_randomrW   0   s|   
 %		zb!J,,Y/0;eiiCUYY,? Luyyc$		.BI&JJsD%uJ5!q&		x9,JJr2uBGGJ4	l	"yyAu%,,Y77	j	 		!Q&--i8==hG"))T^J__eefopp	/	/yyAu%++^I;788r   rV   returnc                    | j                   j                  }|t        v r_|j                  d      }| j	                  t        t        |            }t        t        j                  ||      t        t        |            S |r3d|v r/t        t        j                  | |      t        t        |            S |dk(  r*|dk(  r%t        j                  | |      j                         S t        j                  | |      S )z
    Note: We need dst_type because the type of x can be different from dst_type.
          For example: x is of type `float32`, dst_type is `bfloat16`.
          If dst_type is None, we infer dst_type from x.
    u)devicer>   r   r   )r<   namerF   lstriprO   rI   rG   r   torchtensortlr   )rV   r[   dst_typetsigned_type_namex_signeds         r    	to_tritonre   N   s     	
AK88C=88GB(89:5<<@'"a.QQH,u||Af=wr8?TUU	>h*4<<&1::<<||Af--r   c                 $   t        | t        j                  j                        r| j                  S t        | t
        j                        r0t        j                  dt        |             }|j                  d      S t        dt        |              )Nz^torch\.(\w+)$r=   znot a triton or torch dtype: )rC   r$   languager<   r\   r^   rematchr5   group	TypeErrortype)r<   ms     r    torch_dtype_namern   a   sh    %../zz	E5;;	'HH&E
3wwqz7U}EFFr   c                    t        | t              rX| j                  j                         j	                         j                  t        t        t        | j                                    S t        | t        j                        rf| j                  t        j                  u r,| j                         j                         j	                         S | j                         j	                         S t        d|        )Nz Not a triton-compatible tensor: )rC   r   basecpunumpyrO   rI   rG   rn   r<   r^   Tensorr   float
ValueError)rV   s    r    to_numpyrv   l   s    !]#vvzz|!!#**727G7P+QRR	Au||	$77enn$557==?((**uuw}};A3?@@r   c                  b    t               xr$ t        j                  j                         d   dk\  S )Nr   	   )r/   r^   r*   get_device_capabilityr   r   r    supports_tmarz   w   s&    9C99;A>!CCr   z.Requires TMA support (NVIDIA Hopper or higher))reason)NNNr#   )*r   rh   rr   rG   r^   r$   triton.languagerg   r`   pytestnumpy.randomr   typingr   r   triton.runtime.jitr   r   rE   rF   integral_dtypesrM   dtypesdtypes_with_bfloat16torch_float8_dtypestorch_dtypesr!   r(   r/   r2   r7   rW   ndarrayrs   re   r5   rn   rv   rz   markskipifrequires_tmar   r   r    <module>r      s	   	 	      $ " 90
5{*0	<	', &6 x*$y0<?:,N:=A
@
6
9x'< 9<. .u]ELL=X7Y .&Gs GAD {{!!ln"4=m!nr   