
    3j#                     H   S SK r S SKrS SKrS SKrS SKrS SKJr  S SKJ	r	  S SK
JrJrJr  S SKrS SKJr  S SKJrJrJr  / SQr/ SQr\\-   r/ SQr\S	/-   r\\-   r\S	/-   rS
S/rS/\-   S/-   \-   S	/-   r\" \" \5      1 Sk-
  5      r S r!S r"S r#S r$S r%S r&S r'S r(S r)S r*S r+S r,S r-S r.S r/S r0S r1S  r2S! r3S" r4S# r5S7S$\\   4S% jjr6S8S&\Rn                  S'\\\Rp                  4   4S( jjr9S&\:S'\Rv                  4S) jr<S'\:4S* jr=S+ r>S9S, jr?S- r@S9S. jrA\R                  R                  \?" 5       (       + \A" 5       S/9rDS0\ES1\E4S2 jrFS3\\Rp                  \R                  R                  R$                  4   S'\Rp                  4S4 jrIS8S5\\\:      4S6 jjrJg):    Nknobs)OptionalSetUnion)RandomState)TensorWrapperreinterprettype_canonicalisation_dict)int8int16int32int64)uint8uint16uint32uint64)float16float32float64bfloat16float8_e4m3fnfloat8_e5m2boolr   >   r   r   r   c                  H    [         R                  R                  SS5      S:H  $ )NTRITON_INTERPRET01)osenvironget     R/home/wildlama/miniconda3/lib/python3.13/site-packages/triton/_internal_testing.pyis_interpreterr%      s    ::>>,c2c99r#   c                      [        5       (       a  g [        R                  R                  R                  R                  5       $ N)r%   tritonruntimedriveractiveget_current_targetr"   r#   r$   r,   r,      s-    >>  ''::<<r#   c                  >    [        5       n U c  S$ U R                  S:H  $ )NFcudar,   backendtargets    r$   is_cudar3   $   s"    !FN5@&(@@r#   c                  l    [        5       =(       a$    [        R                  R                  5       S   S:  $ )Nr      r3   torchr.   get_device_capabilityr"   r#   r$   is_ampere_or_newerr9   )   &    9C99;A>!CCr#   c                  l    [        5       =(       a$    [        R                  R                  5       S   S;   $ )Nr   )
      r6   r"   r#   r$   is_blackwellr>   -   s&    9J99;A>(JJr#   c                  l    [        5       =(       a$    [        R                  R                  5       SS S:H  $ )Nr      )r<      r6   r"   r#   r$   is_blackwell_ultrarB   1   s(    9K99;Aa@GKKr#   c                  l    [        5       =(       a$    [        R                  R                  5       S   S:  $ Nr   	   r6   r"   r#   r$   is_hopper_or_newerrF   5   r:   r#   c                  l    [        5       =(       a$    [        R                  R                  5       S   S:H  $ rD   r6   r"   r#   r$   	is_hopperrH   9   r:   r#   c                  l    [        5       =(       a$    [        R                  R                  5       S   S:H  $ )Nr      r6   r"   r#   r$   is_sm12xrK   =   s&    9D99;A>"DDr#   c                  >    [        5       n U c  S$ U R                  S:H  $ )NFhipr/   r1   s    r$   is_hiprN   A   "    !FN5?%(??r#   c                  v    [        5       n U S L=(       a%    U R                  S:H  =(       a    U R                  S:H  $ )NrM   gfx90ar,   r0   archr1   s    r$   is_hip_cdna2rT   F   1    !FU&..E"9UfkkX>UUr#   c                  v    [        5       n U S L=(       a%    U R                  S:H  =(       a    U R                  S:H  $ )NrM   gfx942rR   r1   s    r$   is_hip_cdna3rX   K   rU   r#   c                  v    [        5       n U S L=(       a%    U R                  S:H  =(       a    U R                  S:H  $ )NrM   gfx950rR   r1   s    r$   is_hip_cdna4r[   P   rU   r#   c                  v    [        5       n U S L=(       a%    U R                  S:H  =(       a    SU R                  ;   $ )NrM   gfx11rR   r1   s    r$   is_hip_rdna3r^   U   1    !FT&..E"9Tg>TTr#   c                  v    [        5       n U S L=(       a%    U R                  S:H  =(       a    SU R                  ;   $ )NrM   gfx12rR   r1   s    r$   is_hip_rdna4rb   Z   r_   r#   c                  v    [        5       n U S L=(       a%    U R                  S:H  =(       a    SU R                  ;   $ )NrM   gfx1250rR   r1   s    r$   is_hip_gfx1250re   _   s1    !FV&..E"9Vi6;;>VVr#   c                  Z    [        5       =(       d    [        5       =(       d
    [        5       $ r'   )rT   rX   r[   r"   r#   r$   is_hip_cdnarg   d   s    >=\^=|~=r#   c                  8    [        5       =(       d
    [        5       $ r'   )r^   rb   r"   r#   r$   is_hip_rdnari   h   s    >+\^+r#   c                  (    [        5       (       a  S$ S$ )Ni  i   )r[   r"   r#   r$   get_hip_lds_sizerk   l   s    !^^6..r#   c                  >    [        5       n U c  S$ U R                  S:H  $ )NFxpur/   r1   s    r$   is_xpurn   p   rO   r#   c                  J    [        5       n U c  S$ [        U R                  5      $ )N )r,   strrS   r1   s    r$   get_archrr   u   s"    !F25S%55r#   rsc                 Z   [        U [        5      (       a  U 4n Uc	  [        SS9nU[        [        -   ;   a  [
        R                  " [        [
        U5      5      nUc  UR                  O[        X5R                  5      nUc  UR                  O[        XER                  5      n[        [
        U5      nUR                  X4XS9nSXwS:H  '   U$ U(       a(  SU;   a"  UR                  SSU [
        R                  S9nU$ U[        ;   a"  UR                  SSU 5      R                  U5      $ US	:X  aW  UR                  SSU 5      R                  S
5      R                  S5      [
        R                   " S5      -  R                  S
5      $ US;   a  UR                  SSU 5      S:  $ [#        SU 35      e)zd
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_strrs   lowhighr   rw   xs           r$   numpy_randomr   z   su   
 %		zb!J,,Y/0;eiiCYY,? Luyyc$		.BI&JJs%J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#   r   returnc                    U R                   R                  nU[        ;   a\  UR                  S5      nU R	                  [        [        U5      5      n[        [        R                  " XQS9[        [        U5      5      $ U(       a2  SU;   a,  [        [        R                  " XS9[        [        U5      5      $ US:X  a(  US:X  a"  [        R                  " XS9R                  5       $ [        R                  " XS9$ )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)devicery   r   r   )rw   namer   lstripr   r   r   r
   r7   tensortlr   )r   r   dst_typetsigned_type_namex_signeds         r$   	to_tritonr      s     	
AK88C=88GB(89:5<<@'"a.QQH,u||A=wr8?TUU	>h*4<<1::<<||A--r#   c                 >    [         R                  " [        U    S 5      $ r'   )r   	str_to_tyr   r   s    r$   str_to_triton_dtyper      s    <<215t<<r#   c                 :   [        U [        R                  R                  5      (       a  U R                  $ [        U [
        R                  5      (       a1  [        R                  " S[        U 5      5      nUR                  S5      $ [        S[        U 5       35      e)Nz^torch\.(\w+)$rx   znot a triton or torch dtype: )r~   r(   languagerw   r   r7   rematchrq   group	TypeErrortype)rw   ms     r$   torch_dtype_namer      sn    %..//zz	E5;;	'	'HH&E
3wwqz7U}EFFr#   c                    [        U [        5      (       aX  U R                  R                  5       R	                  5       R                  [        [        [        U R                  5      5      5      $ [        U [        R                  5      (       ag  U R                  [        R                  L a,  U R                  5       R                  5       R	                  5       $ U R                  5       R	                  5       $ [        SU  35      e)Nz Not a triton-compatible tensor: )r~   r	   basecpunumpyr   r   r   r   rw   r7   Tensorr   float
ValueErrorr   s    r$   to_numpyr      s    !]##vvzz|!!#**727G7P+QRR	Au||	$	$77enn$557==?((**uuw}};A3?@@r#   c                 z   [        5       (       a  g[        5       (       d  g[        R                  R                  R
                  nU (       a  SOSn[        [        [        UR                  S5      5      5      n[        U5      S:X  d   U5       e[        R                  R                  5       S   S:  =(       a    X2:  $ )	NTF)rJ   r   )rJ   rA   .r@   r   rE   )r%   r3   r   nvidiaptxasversiontuplemapr   splitlenr7   r.   r8   )
byval_onlycuda_versionmin_cuda_versioncuda_version_tuples       r$   supports_tmar      s    99<<%%--L",w's3(:(:3(?@A!"a';);;'::++-a0A5`:L:``r#   c                      [        5       (       a  g[        5       (       d  g[        R                  R	                  5       S   S:  $ )NTFr   rE   )r%   r3   r7   r.   r8   r"   r#   r$   supports_wsr      s5    99::++-a0A55r#   c                     U (       a  gg)NzURequires __grid_constant__ TMA support (NVIDIA Hopper or higher, CUDA 12.0 or higher)zLRequires advanced TMA support (NVIDIA Hopper or higher, CUDA 12.3 or higher)r"   )r   s    r$   tma_skip_msgr      s    f]r#   )reasonsizealignc                 J    [         R                  " U [         R                  SS9$ )Nr.   )rw   r   )r7   emptyr   )r   r   _s      r$   default_alloc_fnr      s    ;;t5::f==r#   r   c                     [        U [        R                  R                  R                  5      (       a  U R
                  $ U $ r'   )r~   r(   r)   jitr	   r   )r   s    r$   unwrap_tensorr      s-    !V^^''5566vvHr#   skipped_attrc                 z  ^^^^^	 SSK Jm  U c
  [        5       n [        R                  " 5       mTR
                  R                  5        VVs0 s H=  u  p[        UTR                  5      (       d  M"  UTR                  :w  d  M4  X;  d  M;  X_M?     snnm/ mTR                  m	UUUU4S jnUUUUU	4S jnX44$ s  snnf )Nr   r   c                    > TR                  5        H  u  p[        TXR                  5       R                  5       5        UR                  R                  5        HX  nUR                  [        R                  ;   a  TR                  UR                  SS9  M=  TR                  UR                  5        MZ     M     STl        T$ )NF)raisingT)itemssetattrcopyresetknob_descriptorsvalueskeyr   r    delenvappendpropagate_env)r   knobsetknobenv_to_unsetr   	knobs_mapmonkeypatchs      r$   fresh_function)_fresh_knobs_impl.<locals>.fresh_function  s    &__.MDE4!5!5!780077988rzz)&&txx&? ''1	 : / #r#   c                     > TR                  5        H  u  p[        TX5        M     TR                  5         T H*  nU[        R                  ;   d  M  [        R                  U	 M,     TTl        g r'   )r   r   undor   r    r   )r   r   kr   r   r   r   prev_propagate_envs      r$   reset_function)_fresh_knobs_impl.<locals>.reset_function  sX    &__.MDE4) / 	ABJJJJqM  1r#   )
r(   r   setpytestMonkeyPatch__dict__r   r~   
base_knobsr   )
r   r   r   r   r   r   r   r   r   r   s
        @@@@@r$   _fresh_knobs_implr      s    u$$&K #^^1133MDgu//0 	5<@P@P5P 	UYUm 	3I L,,
 
	1 	1 ))Gs   !B70B7B7	B7)NNNr'   )F)Kr   r   r   r   r7   r(   triton.languager   r   r   typingr   r   r   r   numpy.randomr   triton.runtime.jitr	   r
   r   r   r   integral_dtypesr   float_dtypes_with_bfloat16dtypesdtypes_with_bfloat16torch_float8_dtypestorch_dtypessortedr   
tma_dtypesr%   r,   r3   r9   r>   rB   rF   rH   rK   rN   rT   rX   r[   r^   rb   re   rg   ri   rk   rn   rr   r   ndarrayr   r   rq   rw   r   r   r   r   r   r   markskipifrequires_tmar   r   r)   r   r   r   r"   r#   r$   <module>r     s   	 	      ' '  $ U U0
5{*0)ZL8 	<	', &6 x*$y0<?:,NC,-0NNO
:=A
DKLDDE@
V
V
V
U
U
W
>,/@
6
9x'< 9<. .u]ELL=X7Y .&=3 =288 =Gs GA	a6^ {{!!ln"4\^!L>3 >s >U5<<););)I)IIJ u|| +*HSX$6 +*r#   