
    wi                        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j!        f         fdZ"de#fdZ$d Z%d Z&e	j'        (                     e&             d          Z)dS )    N)RandomState)OptionalUnion)TensorWrapperreinterpret)int8int16int32int64)uint8uint16uint32uint64)float16float32float64bfloat16float8_e4m3fnfloat8_e5m2boolr   c                  J    t           j                            dd          dk    S )NTRITON_INTERPRET01)osenvironget     h/root/.openclaw/workspace/chatterbox_venv_py311/lib/python3.11/site-packages/triton/_internal_testing.pyis_interpreterr!      s    :>>,c22c99r   c                  r    t                      rd S t          j        j        j                                        S N)r!   tritonruntimedriveractiveget_current_targetr   r   r    r(   r(      s/     t> '::<<<r   c                  <    t                      } | dn
| j        dk    S )NFcudar(   backendtargets    r    is_cudar/   !   s"    !!FN55&(@@r   c                  <    t                      } | dn
| j        dk    S )NFhipr+   r-   s    r    is_hipr2   &   s"    !!FN55%(??r   c                  N    t                      } | dnt          | j                  S )N )r(   strarchr-   s    r    get_archr7   +   s%    !!F22S%5%55r   rsc                    t          | t                    r| f} |t          d          }|t          t          z   v rt          j        t          t
          |                    }||j        nt          ||j                  }||j	        nt          ||j	                  }t          t
          |          }|
                    ||| |          }d||dk    <   |S |r)d|v r%|
                    dd	| t
          j                  }|S |t          v r*|                    dd|                               |          S |d
k    re|                    dd|                               d                              d          t          j        d          z                      d          S |dv r|                    dd|           dk    S t#          d|           )zp
    Override `rs` if you're calling this function twice and don't want the same
    result for both calls.
    N   )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0;eiiCUY,?,? Luyyc$	.B.BI&&JJsD%uJ55!q&		 
9x9,,JJr2uBGJ44	l	"	"yyAu%%,,Y777	j	 	 		!Q&&--i88==hGG")T^J_J__eefoppp	/	/	/yyAu%%++7I77888r   rV   returnc                 "   | j         j        }|t          v rt|                    d          }|                     t          t          |                    }t          t          j	        ||          t          t          |                    S |r;d|v r7t          t          j	        | |          t          t          |                    S |dk    r.|dk    r(t          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9::5<@@@'"a..QQQ 	VH,,u|Af===wr8?T?TUUU	>>h*44<&111::<<<|Af----r   c                 .   t          | t          j        j                  r| j        S t          | t
          j                  r7t          j        dt          |                     }|	                    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   s|    %.// Gz	E5;	'	' GH&E

33wwqzzEUEEFFFr   c                 ,   t          | t                    rc| j                                                                                            t          t          t          | j	                                      S t          | t          j                  rq| j	        t          j        u r8|                                                                                                 S |                                                                 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    !]## Avzz||!!##**727G7P7P+Q+QRRR	Au|	$	$ A7en$$5577==??((***uuww}}?A??@@@r   c                  n    t                      o't          j                                        d         dk    S )Nr   	   )r/   r^   r*   get_device_capabilityr   r   r    supports_tmarz   w   s)    99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 9 9 9 9 9 9 9000
555{*000	<	', &6 x*$y0<?:,N: : := = =A A A
@ @ @
6 6 6
9 9x'< 9 9 9 9<. . .u]EL=X7Y . . . .&Gs G G G GA A AD D D {!!llnn"4=m!nnr   