
    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
r
S SKJr  S SKJr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+S\\   4S jjr,S,S\RZ                  S\\\R\                  4   4S jjr/S\0S\Rb                  4S  jr2S\04S! jr3S" r4S-S# jr5S-S$ jr6\
Rn                  Rq                  \5" 5       (       + \6" 5       S%9r9S&\:S'\:4S( jr;S)\\R\                  \Rx                  Rz                  R"                  4   S\R\                  4S* jr>g).    N)knobs)RandomState)OptionalUnion)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/var/www/fran/franai/venv/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_cudar1   $   s"    !FN5@&(@@r!   c                  l    [        5       =(       a$    [        R                  R                  5       S   S:  $ )Nr   	   )r1   torchr,   get_device_capabilityr    r!   r"   	is_hopperr6   )   s&    9C99;A>!CCr!   c                  >    [        5       n U c  S$ U R                  S:H  $ )NFhipr-   r/   s    r"   is_hipr9   -   "    !FN5?%(??r!   c                  v    [        5       n U S L=(       a%    U R                  S:H  =(       a    U R                  S:H  $ )Nr8   gfx90ar*   r.   archr/   s    r"   is_hip_cdna2r?   2   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  $ )Nr8   gfx942r=   r/   s    r"   is_hip_cdna3rC   7   r@   r!   c                  v    [        5       n U S L=(       a%    U R                  S:H  =(       a    U R                  S:H  $ )Nr8   gfx950r=   r/   s    r"   is_hip_cdna4rF   <   r@   r!   c                      [        5       n [        U R                  5        U S L=(       a%    U R                  S:H  =(       a    SU R                  ;   $ )Nr8   gfx12)r*   printr>   r.   r/   s    r"   is_hip_gfx12rJ   A   s=    !F	&++T&..E"9Tg>TTr!   c                  Z    [        5       =(       d    [        5       =(       d
    [        5       $ r%   )r?   rC   rF   r    r!   r"   is_hip_cdnarL   G   s    >=\^=|~=r!   c                  >    [        5       n U c  S$ U R                  S:H  $ )NFxpur-   r/   s    r"   is_xpurO   K   r:   r!   c                  J    [        5       n U c  S$ [        U R                  5      $ )N )r*   strr>   r/   s    r"   get_archrS   P   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_strrT   lowhighrd   rX   xs           r"   numpy_randomrs   U   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!   rr   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)devicerZ   r   r   )rX   namerb   lstriprk   re   rc   r   r4   tensortlr   )rr   rw   dst_typetsigned_type_namex_signeds         r"   	to_tritonr   s   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    5      $ r%   )r{   	str_to_tyr	   rr   s    r"   str_to_triton_dtyper      s    <<21566r!   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+)$rY   znot a triton or torch dtype: )r_   r&   languagerX   rx   r4   rematchrR   group	TypeErrortype)rX   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numpyrk   re   rc   r   rX   r4   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)   r   )r      .   r   r3   )r#   r1   r   nvidiaptxasversiontuplemapr`   splitlenr4   r,   r5   )
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                     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,   )rX   rw   )r4   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!   )NNNr%   )F)?r   r   r   rc   r4   r&   triton.languager   r{   r   pytestnumpy.randomr   typingr   r   triton.runtime.jitr   r   r	   ra   rb   integral_dtypesri   float_dtypes_with_bfloat16dtypesdtypes_with_bfloat16torch_float8_dtypestorch_dtypessortedset
tma_dtypesr#   r*   r1   r6   r9   r?   rC   rF   rJ   rL   rO   rS   rs   ndarrayr   r   rR   rX   r   r   r   r   r   markskipifrequires_tmar`   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@
V
V
V
U>@
6
9x'< 9<. .u]ELL=X7Y .&73 7288 7Gs GA	a^ {{!!ln"4\^!L>3 >s >U5<<););)I)IIJ u|| r!   