
    wi7                      d dl mZ d dlZd dlmZmZmZmZmZ d dl	Z	ddl
mZ ddlmZ ddlmZ  ed	          Z G d
 de          ZddZddZddZddZdddZdd$Z	 	 ddd*Zdd-Zdd1Zdd2Zdd3Zdd4Zdd5Zdd7Z dd8Z!dd=Z"dd>Z#ddAZ$ddBZ%ddCZ&ddDZ'ddEZ(ddFZ)ddGZ*ddHZ+ddIZ,ddJZ-ddKZ.ddLZ/ddMZ0ddNZ1ddQZ2ddRZ3ddSZ4ddTZ5ddUZ6ddVZ7ddWZ8ddZZ9dd^Z:dd`Z;ddcZ<dddZ=ddeZ>ddhZ?ddiZ@ddlZAddmZBddnZCddqZDddsZE	 ddduZFdv ZGdw ZHdx ZIdy ZJdz ZKd{ ZLd| ZMd} ZNd~ ZOddZPddZQddZRddZSddZTd ZUd ZVddZWddZXddZYd dZZd dZ[d dZ\d dZ]d dZ^d dZ_d dZ`d ZaddZbddZcddZdddZed ZfddZgddZhddZiddÄZjddĄZkddńZld	dƄZmd
dʄZndd̈́Zod	d΄Zpdτ ZqddЄZrdd҄ZsddӄZtdS (      )annotationsN)ListOptionalSequenceTupleTypeVar   )ir   )core)mathTc                       e Zd Z fdZ xZS )IncompatibleTypeErrorImplc                    || _         || _        d| j                                         z   dz   | j                                        z   | _        t	          t
          |                               | j                   d S )Nzinvalid operands of type  and )type_atype_b__repr__messagesuperr   __init__)selfr   r   	__class__s      h/root/.openclaw/workspace/chatterbox_venv_py311/lib/python3.11/site-packages/triton/language/semantic.pyr   z"IncompatibleTypeErrorImpl.__init__   sl    2T[5I5I5K5KKgUX\XcXlXlXnXnn'..77EEEEE    )__name__
__module____qualname__r   __classcell__)r   s   @r   r   r      sA        F F F F F F F F Fr   r   axisintbuilder
ir.builderreturn	tl.tensorc                    | dvrt          d|            t          j        |                    |           t          j                  S )Nr   r   r	   z+program_id axis must be 0, 1, or 2 but got )
ValueErrortltensorcreate_get_program_idint32r!   r#   s     r   
program_idr/      sF    9MtMMNNN9W22488"(CCCr   c                    | dvrt          d|            t          j        |                    |           t          j                  S )Nr(   z-num_programs axis must be 0, 1, or 2 but got )r)   r*   r+   create_get_num_programsr-   r.   s     r   num_programsr2   "   sF    9OOOPPP9W44T::BHEEEr   a_tytl.dtypeb_tyc                   | j         }|j         }| j        }|j        }||k    r
||k    r| n|S |t          j        j        j        k    r
||k    r| n|S |t          j        j        j        k    r
||k    r|n| S t          d| d|           )Nzunexpected signedness r   )int_bitwidthint_signednessr*   dtype
SIGNEDNESSUNSIGNED	TypeError)r3   r5   a_rankb_ranka_snb_sns         r   integer_promote_implrA   -   s    FFDD t||ttD0	$-	-	-''ttT1	$-	-	-''ttT1
>T>>>>
?
??r   a_is_scalarboolb_is_scalar
div_or_modc                   ||k    rk|r| |fn|| f\  }}|                                 j        |                                 j        k    r*|r&|t          j        t          j        fv rt          j        S |S |                                 s|                                rt          j        S |                                 s|                                rt          j        S | 	                                s|	                                r|rt          j        S t          j        S | 
                                s|
                                rN|rt          j        S | 
                                r |
                                rt          j        S t          j        S |                                 r(|                                r| |k    r| nt          j        S |                                 r|                                st          d|  d|           |rO| j        |j        k    r?t          d|                                 z   dz   |                                z   dz             t!          | |          S )Nunexpected type r   zCannot use /, #, or % with x because they have different signedness;this is unlikely to result in a useful answer. Cast them to the same signedness.)kindvaluer*   float16bfloat16float32is_fp64float64is_fp32is_fp16is_bf16is_fp8is_intr<   r8   r   rA   )r3   rB   r5   rD   rE   	scalar_ty	tensor_tys          r   computation_type_implrW   =   sG   
 k!!/:Ld||t	9>>!Y^^%5%5%;;; "yRZ,EEEz! ||~~  z ||~~  z ||~~   	::||~~   	:<<>> 	dllnn 	;z{{}} 4 4t||tt3;;== > ><4<<d<<===  ld)T-@@@5G'QTXTaTaTcTcckk l l 	l  d+++r   T
check_typec                   t          | t                    r2t          j        |                    |           t          j                  S t          | t                    rd| cxk    rdk     rn nt          j        }njd| cxk    rdk     rn nt          j        }nMd| cxk    rdk     rn nt          j	        }n0d| cxk    rdk     rn nt          j
        }nt          d|  d          t          d	| ||
          S t          | t                    rzd}dddz  z  }t          d         |           }|t          d          k    s|dk    s| | k    s||cxk    r|k    rn nt          j        }nt          j        }t          d	| ||
          S t          | t          j                  rt%          | j        |          S t          | t          j                  r| S |r#t)          d|  dt+          |            d          | S )N           l                             l            zNonrepresentable integer . r9   r#   g      8g   ?r	      absinf        zcannot convert z	 of type z
 to tensor)
isinstancerC   r*   r+   get_int1int1r"   r-   uint32int64uint64r)   fullfloat__builtins__rM   rO   	constexpr	to_tensorrJ   r<   type)xr#   rX   r9   min_float32max_float32abs_xs          r   ro   ro   o   s=   !T y))!,,bg666	As		 QHEEa%IEEq    5     HEEa%IEE====>>>B8888	Au		 !QV+U#A&&E%LL  C<<66%....;.....JEEJEB8888	Ar|	$	$ '***	Ary	!	!  KI!IId1ggIIIJJJHr   r   r   allow_ptr_aNonec                    |                                  r`|st          | |          |                                 r| |k    rt          | |          |                                rt          | |          d S d S N)is_ptrr   is_floating)r   r   ru   s      r   check_ptr_type_implr{      s    }} < 	<+FF;;;==?? 	<& 0 0+FF;;; 	<+FF;;;< <	< 	<r   Flhstl.tensor | numbers.NumberrhsTuple[tl.tensor, tl.tensor]c                   t          | t          j                  }t          |t          j                  }|r| }	t          | |          } |r|}
t          ||          }| j        j        }|j        j        }t          |||           t          |||           |r|                                s|                                st          |||||          }|r|	dk     r|	                                s|r)|
dk     r#|	                                rt          d          |rt          d|	||          nt          | ||          } |rt          d|
||          nt          |||          }t          | ||          \  } }| |fS )Nr   z{Cannot perform a binary operation between an unsigned tensor and a negative scalar. Perform a explicit cast on one of them.r_   r`   )re   numbersNumberro   rp   scalarr{   ry   rW   is_int_unsignedr)   rk   castbroadcast_impl_value)r|   r~   r#   allow_lhs_ptrallow_rhs_ptrarithmetic_checkrE   lhs_is_scalarrhs_is_scalar
lhs_scalar
rhs_scalar
lhs_sca_ty
rhs_sca_ty
ret_sca_tys                 r   binary_op_type_checking_implr      s    sGN33MsGN33M &
W%% &
W%% JJ
J>>>
J>>> 	t
 1 1 3 3 	tJ<M<M<O<O 	t*:}jR_akll
 	Hj1nn1K1K1M1Mn  -%/!^^
8R8R8T8T^ G H H H CPtd
*g? ? ? ?UYZ]_ikrUsUs 	 CPtd
*g? ? ? ?UYZ]_ikrUsUs 	 $Cg66HC8Or   	binary_opcallablec                   | j         j        j        dk    s|j        j        sd S | j         j        }|j         j        }||k    sJ |                                sJ t          | t          j        |          } t          |t          j        |          } || |d|          }|	                                }t          j
        |                    |          t          j                  }|                                }t          j
        |                    |          t          j                  }t          t          |||          t          |||          |          }	d|j         d|j         }
t#          |	|
|           d S )N@   Fr"   z! overflow detected for operation )rp   r   r7   optionssanitize_overflowrT   r   r*   ri   get_int_max_valuer+   	get_int64get_int_min_valueand_
less_equalgreater_equalr   device_assert)r|   r~   r#   r   r   r   ret	max_value	min_valuecondmsgs              r    binary_op_sanitize_overflow_implr      s[   
x#r))1R)JJ####
sBHg
&
&C
sBHg
&
&C
)CeW
-
-C,,..I	'++I66AAI,,..I	'++I66AAI
3	733]3	SZ5[5[]deeD
^
'
^
^)J\
^
^C$W%%%%%r   inputotherr   c                r   t          | ||dd          \  } }| j        j        }|j        j        }|                                r#|                                rt	          d          |                                r0|                                s|| }} | j        j        }|j        j        }|                                r8t          j        |                    | j        |j                  | j                  S |	                                r8t          j        |
                    | j        |j                  | j                  S |                                rQ|rt          | ||t                     t          j        |                    | j        |j                  | j                  S t	          d|           )NTzcannot add pointers togetherrG   )r   rp   r   ry   r<   r*   r+   create_addptrhandlerz   create_faddrT   r   add
create_add)r   r   r   r#   input_scalar_tyother_scalar_tys         r   r   r      s   /ugtTRRLE5j'Oj'O 8O$:$:$<$< 86777  ,(>(>(@(@ ,eu*+*+ 	Uy..u|U\JJEJWWW		$	$	&	& Uy,,U\5<HH%*UUU				!	! U 	I,UE7CHHHy++EL%,GGTTT
888
9
99r   c           	        t          | ||dd          \  } }| j        j        }|                                rFt	          j        |                    | j        t          ||          j                  | j                  S |	                                r8t	          j        |
                    | j        |j                  | j                  S |                                rQ|rt          | ||t                     t	          j        |                    | j        |j                  | j                  S t          d|           )NTFrG   )r   rp   r   ry   r*   r+   r   r   minusrz   create_fsubrT   r   sub
create_subr<   r   r   r   r#   rU   s        r   r   r      s   /ugtUSSLE5
!I hy..u|U5'=R=R=YZZ\a\fggg Uy,,U\5<HH%*UUU					 U 	I,UE7CHHHy++EL%,GGTTT
2y22
3
33r   c                   t          | ||          \  } }| j        j        }|                                r8t	          j        |                    | j        |j                  | j                  S |                                rQ|rt          | ||t                     t	          j        |                    | j        |j                  | j                  S t          d|           NrG   )r   rp   r   rz   r*   r+   create_fmulr   rT   r   mul
create_mulr<   r   s        r   r   r     s    /ugFFLE5
!I Uy,,U\5<HH%*UUU					 U 	I,UE7CHHHy++EL%,GGTTT
2y22
3
33r   c           	     V   t          | ||dddd          \  } }| j        j        }|j        j        }|                                r'|                                rt          |||          }n|                                r&|                                rt          | ||          } n|                                rK|                                r7t          | t          j        |          } t          |t          j        |          }nn|                                rH|                                r4|j        |j        k    rt          |||          }n$t          | ||          } nt          d|           t          j
        |                    | j        |j                  | j                  S NFTrG   )r   rp   r   rz   rT   r   r*   rM   fp_mantissa_widthr<   r+   create_fdivr   )r   r   r#   r   r   s        r   truedivr     s   /ugueUY[_``LE5j'Oj'O""$$ >)?)?)A)A >UOW55				!	! >o&A&A&C&C >UOW55				!	! >o&<&<&>&> >UBJ00UBJ00		$	$	&	& >?+F+F+H+H >,/PPP99EE99EE <?<<===9W((u|DDejQQQr   c           	     B   t          | ||dddd          \  } }| j        j        }|j        j        }|                                r|                                rt	          ||          }t          | ||          } t          |||          }|                                r8t          j        |	                    | j
        |j
                  | j                  S t          j        |                    | j
        |j
                  | j                  S t          d|           r   )r   rp   r   rT   rA   r   is_int_signedr*   r+   create_sdivr   create_udivr<   )r   r   r#   r   r   ret_tys         r   floordivr   0  s   /ugueUY[_``LE5j'Oj'O ZO$:$:$<$< Z%oGGUFG,,UFG,,!! 	Z9W00u|LLejYYY9W00u|LLejYYY
888
9
99r   ieee_roundingc           	     D   | j         j        }|j         j        }|                                r|                                st          d          t	          | ||dddd          \  } }|                    | j        |j                  }t          j        || j                   S )Nz4both operands of fdiv must have floating scalar typeFT)	rp   r   rz   r<   r   r   r   r*   r+   )r   r   r   r#   r   r   r   s          r   fdivr   ?  s    j'Oj'O&&(( P0K0K0M0M PNOOO/ugueUZ\`aaLE5


elEL
9
9C9S%*%%%r   c           	        t          | ||dddd          \  } }| j        j        }|j        j        }|                                rJt	          j        t          | |d|          |          }t          | t          ||d|          d|          }|S |	                                r|j
        |j
        k    r?t          d|                                z   dz   |                                z   dz             |                                r8t          j        |                    | j        |j                  | j                  S t          j        |                    | j        |j                  | j                  S t          d|           )NFT_builderzCannot mod z by rH   rG   )r   rp   r   rz   r   floorr   r   r   rT   r8   r<   r   r   r*   r+   create_sremr   create_urem)r   r   r#   rU   r   r   r   s          r   modr   J  s   /ugueUY[_``LE5
!Ij'O Z
4ueW==PPP%UE4994II
					 Z#'EEEMI,>,>,@,@@6IOLdLdLfLff jo o p p p ""$$ 	Z9W00u|LLejYYY9W00u|LLejYYY
2y22
3
33r   rq   ypropagate_nantl.PropagateNanc                   t          | ||          \  } }| j        }|                                r|t          j        j        k    r8t          j        |                    | j        |j                  | j	                  S |t          j        j
        k    r8t          j        |                    | j        |j                  | j	                  S t          d|           |                                r8t          j        |                    | j        |j                  | j	                  S |                                r8t          j        |                    | j        |j                  | j	                  S t#          d|           NzUnexpected propagate_nan Unexpected dtype )r   r9   rz   r*   PropagateNanALLr+   create_minimumfr   rp   NONEcreate_minnumfr)   r   create_minsir   create_minuir<   rq   r   r   r#   r9   s        r   minimumr   f  H   '1g66DAqGE 5BO///9W44QXqxHH!&QQQbo2229W33AHahGGPPPHHHIII					 5y--ahAA16JJJ				 	  5y--ahAA16JJJ3E33444r   c                   t          | ||          \  } }| j        }|                                r|t          j        j        k    r8t          j        |                    | j        |j                  | j	                  S |t          j        j
        k    r8t          j        |                    | j        |j                  | j	                  S t          d|           |                                r8t          j        |                    | j        |j                  | j	                  S |                                r8t          j        |                    | j        |j                  | j	                  S t#          d|           r   )r   r9   rz   r*   r   r   r+   create_maximumfr   rp   r   create_maxnumfr)   r   create_maxsir   create_maxuir<   r   s        r   maximumr   x  r   r   minmaxc                T   t          |||          \  }}t          | ||          \  } }t          | ||          \  } }| j        }|                                r?t          j        |                    | j        |j        |j        |          | j                  S t          d| d          )Nr   z(. Only floating point clamp is supported)	r   r9   rz   r*   r+   create_clampfr   rp   r<   )rq   r   r   r   r#   r9   s         r   clampr     s    +Cg>>HC)!S'::FAs)!S'::FAsGE ]y..qxSZQ^__abaghhh[E[[[\\\r   c                N   t          | ||          \  } }| j        j        }|j        j        }|                                r|                                st	          ||          t          ||          }||k    rt          | ||          } ||k    rt          |||          }| |fS rx   )r   rp   r   rT   r   rA   r   )r   r   r#   input_sca_tyother_sca_tyr   s         r   bitwise_op_type_checking_implr     s    /ugFFLE5:$L:$L   D(;(;(=(= D'lCCC%lLAAJ\!!UJ00\!!UJ00%<r   c                    t          | ||          \  } }t          j        |                    | j        |j                  | j                  S rx   )r   r*   r+   
create_andr   rp   r   r   r#   s      r   r   r     ?    0wGGLE59W''elCCUZPPPr   c                    t          | ||          \  } }t          j        |                    | j        |j                  | j                  S rx   )r   r*   r+   	create_orr   rp   r   s      r   or_r     s?    0wGGLE59W&&u|U\BBEJOOOr   c                    t          | ||          \  } }t          j        |                    | j        |j                  | j                  S rx   )r   r*   r+   
create_xorr   rp   r   s      r   xor_r     r   r   c                   | j                                         s#t          | t          j        d          |          } |j                                         s#t          |t          j        d          |          }t          | ||          S Nrg   )rp   is_int1bitcastr*   r9   r   r   s      r   logical_andr     sv    : :rx//99: :rx//99ug&&&r   c                   | j                                         s#t          | t          j        d          |          } |j                                         s#t          |t          j        d          |          }t          | ||          S r   )rp   r   r   r*   r9   r   r   s      r   
logical_orr     sv    : :rx//99: :rx//99ueW%%%r   c                    | j                                         s#t          | t          j        d          |          } t          | |          S r   )rp   r   r   r*   r9   invert)r   r#   s     r   not_r    sC    : :rx//99%!!!r   c                    t          | ||          \  } }t          j        |                    | j        |j                  | j                  S rx   )r   r*   r+   create_lshrr   rp   r   s      r   lshrr    ?    0wGGLE59W((u|DDejQQQr   c                    t          | ||          \  } }t          j        |                    | j        |j                  | j                  S rx   )r   r*   r+   create_ashrr   rp   r   s      r   ashrr    r  r   c                    t          | ||          \  } }t          j        |                    | j        |j                  | j                  S rx   )r   r*   r+   
create_shlr   rp   r   s      r   shlr    r   r   c                    | S rx   r_   )r   s    r   plusr    s    Lr   c                *   | j         j        }|                                r't          d|                                z   dz             t          j        |                    |                    |                    |          }t          || d|          S )Nz$wrong type argument to unary minus ()T)
rp   r   ry   r)   r   r*   r+   get_null_valueto_irr   )r   r#   r   _0s       r   r   r     s    :$L a?,BWBWBYBYY\__```	7)),*<*<W*E*EFF	U	UBr5$(((r   c                P   | j         j        }|                                s|                                r't	          d|                                z   dz             t          j        |                    |	                    |                    |          }t          | ||          S )Nz%wrong type argument to unary invert (r  )rp   r   ry   rz   r)   r   r*   r+   get_all_ones_valuer  r   )r   r#   r   _1s       r   r   r     s    :$L b 8 8 : : b@<CXCXCZCZZ]``aaa	7--l.@.@.I.IJJL	Y	YBr7###r   vtl.block_typec                    | j                                         st          j        S | j         j        }t          j        t          j        |          S rx   )rp   is_blockr*   rg   shape
block_type)r  r  s     r   
_bool_liker    s:    6?? wFLE=%(((r   c                ^   t          | ||          \  } }| j        j        }|                                r@t	          j        |                    | j        |j                  t          |                     S |	                                r|
                                r@t	          j        |                    | j        |j                  t          |                     S t	          j        |                    | j        |j                  t          |                     S t          d|           r   )r   rp   r   rz   r*   r+   create_fcmpOGTr   r  rT   r   create_icmpSGTcreate_icmpUGTr<   r   r   r#   rU   s       r   greater_thanr"       /ugFFLE5
!I dy//elKKZX]M^M^___					 d""$$ 	d9W33EL%,OOQ[\aQbQbccc9W33EL%,OOQ[\aQbQbccc
2y22
3
33r   c                ^   t          | ||          \  } }| j        j        }|                                r@t	          j        |                    | j        |j                  t          |                     S |	                                r|
                                r@t	          j        |                    | j        |j                  t          |                     S t	          j        |                    | j        |j                  t          |                     S t          d|           r   )r   rp   r   rz   r*   r+   create_fcmpOGEr   r  rT   r   create_icmpSGEcreate_icmpUGEr<   r!  s       r   r   r     r#  r   c                ^   t          | ||          \  } }| j        j        }|                                r@t	          j        |                    | j        |j                  t          |                     S |	                                r|
                                r@t	          j        |                    | j        |j                  t          |                     S t	          j        |                    | j        |j                  t          |                     S t          d|           r   )r   rp   r   rz   r*   r+   create_fcmpOLTr   r  rT   r   create_icmpSLTcreate_icmpULTr<   r!  s       r   	less_thanr,    r#  r   c                ^   t          | ||          \  } }| j        j        }|                                r@t	          j        |                    | j        |j                  t          |                     S |	                                r|
                                r@t	          j        |                    | j        |j                  t          |                     S t	          j        |                    | j        |j                  t          |                     S t          d|           r   )r   rp   r   rz   r*   r+   create_fcmpOLEr   r  rT   r   create_icmpSLEcreate_icmpULEr<   r!  s       r   r   r   .  r#  r   c                   t          | ||          \  } }| j        j        }|                                r@t	          j        |                    | j        |j                  t          |                     S |	                                r@t	          j        |
                    | j        |j                  t          |                     S t          d|           r   )r   rp   r   rz   r*   r+   create_fcmpOEQr   r  rT   create_icmpEQr<   r!  s       r   equalr4  =      /ugFFLE5
!I _y//elKKZX]M^M^___					 _y..u|U\JJJW\L]L]^^^
2y22
3
33r   c                   t          | ||          \  } }| j        j        }|                                r@t	          j        |                    | j        |j                  t          |                     S |	                                r@t	          j        |
                    | j        |j                  t          |                     S t          d|           r   )r   rp   r   rz   r*   r+   create_fcmpUNEr   r  rT   create_icmpNEr<   r!  s       r   	not_equalr9  I  r5  r   startendc                   t          | t                    rt          |t                    st          d          t          | dz	            }t          |dz	            }|s|rt          d          || k    rt          d          || z
  }||dz
  z  dk    rt          d          |g}t	          j        t          j        |          }t	          j        |                    | |          |          S )Nz/arange's arguments must be of type tl.constexpr    zarange must fit in int32z=arange's end argument must be greater than the start argumentr   r   z#arange's range must be a power of 2)	re   r"   r)   rC   r*   r  r-   r+   create_make_range)r:  r;  r#   is_start_int64is_end_int64ranger  r   s           r   arangerB  Z  s    eS!! LC)=)= LJKKK%2+&&Nr	??L 5 53444
e||XYYY%KE!!>???GE]28U++F9W..uc::FCCCr   r  	List[int]r9   c                   t          |t          j                  r,|j        j        dk    s
J d            t          |||          }nx|t          d          |dk    r)|                    |                    |                    }n#t          |d|j
                   } ||          }t          j        ||          }t          || |          S )Nr   zonly accepts size-1 tensorz2dtype must be specified when value is not a tensorr   get_)re   r*   r+   numelrJ   r   r)   r  r  getattrnamesplat)r  rJ   r9   r#   get_value_fns        r   rk   rk   k  s    %## ({ A%%%'C%%%UE7++ =QRRRA::**5;;w+?+?@@EE"7,?5:,?,?@@L L''E	%''w'''r   rJ   c                   | j                                         r
J d            t          |          dk    r| S t          j        | j        |          }t          j        |                    | j        |          |          S )NzCannot splat a block tensorr   )	rp   r  lenr*   r  r9   r+   create_splatr   )rJ   r  r#   r   s       r   rI  rI    sp    z""$$CC&CCC$
5zzQ]5;..F9W))%,>>GGGr   	dst_shapecan_reorderc                    d}|D ]}||z  }| j         j        |k    rt          d          t          j        | j         j        |          }t          j        |                    | j        ||          |          S )Nr   z:reshape() cannot change total number of elements in tensor)	rp   rF  r)   r*   r  r   r+   create_reshaper   )r   rN  rO  r#   rF  sr   s          r   reshaperS    s}    E  
z5  UVVV]5:,i88F9W++EL)[QQSYZZZr   c                @   d | j         D             }|                    |d           | j                                        st	          | ||          S t          j        | j        j        |          }t          j        |	                    | j
        |          |          S )Nc                6    g | ]}t          j        |          S r_   r*   _constexpr_to_value.0rq   s     r   
<listcomp>zexpand_dims.<locals>.<listcomp>  s#    @@@q'**@@@r   r   )r  r#   )r  insertrp   r  rI  r*   r  r   r+   create_expand_dimsr   )r   r!   r#   rN  r   s        r   expand_dimsr]    s    @@EK@@@IT1:   >U)W====]5:,i88F9W//dCCVLLLr   c                &   |s
J d            t          | j                  dk    sJ t          j        | j        j        | j        d         |j        d         z   g          }t          j        |                    | j        |j                  |          S )Nz;current implementation of `cat` always may reorder elementsr   r   )	rL  r  r*   r  rp   r   r+   
create_catr   )r|   r~   rO  r#   ret_types        r   catra    s~    UUUUU;sy>>Q}SX_sy|cil/J.KLLH9W''
CJ??JJJr   abc                   t          | ||          \  } }| j        g k    }|r"t          | d|          } t          |d|          }t          | j        d         t          j                  rt	          j        d          }nd}| j        |gz   }t	          j        | j        j        |          }t	          j	        |
                    | j        |j                  |          }|rt          |dgd|          }|S )Nr   r	   FrO  r#   )r   r  r]  re   r*   rn   r  rp   r   r+   create_joinr   rS  )rb  rc  r#   
was_rank_1two	new_shaper`  r   s           r   joinrk    s    1g..DAq BJ '1g&&1g&&!'"+r|,, l1oo3%I}QV]I66H
)G''!(;;X
F
FC DcA3E7CCCJr   c                j   t          | j                  dk    sJ t          j        | j        d                   dk    sJ | j        d d         }t          j        | j        j        |          }|                    | j                  \  }}t          j	        ||          t          j	        ||          fS )Nr   re  r	   )
rL  r  r*   rW  r  rp   r   create_splitr   r+   )rb  r#   rj  r`  outLHSoutRHSs         r   splitrp    s    LL1"172;//14444I}QV]I66H))!(33NFF
	&(##
	&(## r   dims
Tuple[int]c                    t           j                  t          |          k    rt          d          t          d |D                       t	          t          t          |                              k    rt          d|           t          j         j        j	         fd|D                       }t          j
        |                     j        |          |          S )Nz5permute dims must have the same length as input shapec              3  >   K   | ]}t          j        |          V  d S rx   rV  )rY  ds     r   	<genexpr>zpermute.<locals>.<genexpr>  s-      66Ab$Q''666666r   z?permute dims must be a permutation of 0, 1, ..., n-1, but were c                *    g | ]}j         |         S r_   )r  )rY  ru  r   s     r   rZ  zpermute.<locals>.<listcomp>  s    0N0N0NAQ0N0N0Nr   )rL  r  r)   sortedlistrA  r*   r  rp   r   r+   create_transr   )r   rq  r#   r`  s   `   r   permuter{    s    
5;3t99$$PQQQ6666666$uSYY?O?O:P:PPPa[_aabbb}UZ.0N0N0N0N0N0N0NOOH9W))%,==xHHHr   c                   | j                                         sHt          j        | j         |          }t          j        |                    | j        |          |          S | j                                         }t          |          t          |          k    rt          d| d|           ||k    r| S t          |          D ];\  }}||         |k    r*|dk    r$t          d||          d| d| d| d| 
          <t          j        | j         j        |          }t          j        |                    | j        |          |          S )Nz!Cannot broadcast, rank mismatch: z, r   z3Cannot broadcast, the expanded size of the tensor (z ) must match the existing size (z) at non-singleton dimension : )rp   r  r*   r  r+   rM  r   get_block_shapesrL  r)   	enumerater   create_broadcast)r   r  r#   r   	src_shapeiitems          r   broadcast_impl_shaper    sv   :   Luz511y--elEBBFKKK
++--I
9~~U##QYQQ%QQRRR	Y'' < <48t		 ;SXYZS[ ; ;?C; ;!"; ;&/; ;38; ; < < < ]5:,e44F9W--elEBBFKKKr   c           	     v   | j         }|j         }|                                ru|                                sat          j        |j        |j                  }t          j        |                    |j        |	                                          |          }n|                                su|                                rat          j        |j        |j                  }t          j        |                    | j        |	                                          |          } n|                                r|                                r|	                                }|	                                }t          |          t          |          k     rt          t          |          t          |                    D ]g}t          j        |                    | j        d          t          j        |j        dg|z                       } | j         }|	                                }hnt          |          t          |          k     rt          t          |          t          |                    D ]g}t          j        |                    |j        d          t          j        |j        dg|z                       }|j         }|	                                }ht          |          t          |          k    sJ g }t          |          D ]\  }	}
||	         }|
dk    r|                    |           )|dk    s||
k    r|                    |
           Kt          dt!          |	          z   dz   t!          |
          z   dz   t!          |          z             ||k    rHt          j        |j        |          }t          j        |                    | j        |          |          } ||k    rHt          j        |j        |          }t          j        |                    |j        |          |          }| |fS )Nr   r   z?Cannot make_shape_compatible: incompatible dimensions at index r}  r   )rp   r  r*   r  r   r  r+   rM  r   r~  rL  rA  r\  r  appendr)   strr  )r|   r~   r#   lhs_tyrhs_ty	lhs_shape	rhs_shape_	ret_shaper  leftrightr   s                r   r   r     s   XFXF  +U!2!2 +Uv}fl;;i,,SZ9P9P9R9RSSU[\\__ 'U6??#4#4 'Uv}fl;;i,,SZ9P9P9R9RSSU[\\			 #Uv00 #U++--	++--	y>>C	NN**3y>>3y>>:: 6 6i : :3:q I I "fmaS9_ M MO O"3355			6
 ^^c)nn,,3y>>3y>>:: 6 6i : :3:q I I "fmaS9_ M MO O"3355		9~~Y////	 ++ 	a 	aGAtaLEqyy  ''''1**%4--  &&&&  "-/21vv"68<"=?B4yy"IKR"SUXY^U_U_"` a a a	!!]6=)<<F)G44SZKKVTTC	!!]6=)<<F)G44SZKKVTTC8Or   rounding_modeOptional[str]c                    | d S | dk    rt           j        j        S | dk    rt           j        j        S t	          d|  d          )NrtnertzzInvalid rounding mode: z0. Supported rounding modes are 'rtne' and 'rtz'.)r
   ROUNDING_MODERTNERTZr)   )r  s    r   _str_to_rounding_moder  "  sU    t$$##
n}nnn
o
oor   dst_tyc                F   | j         }|                                r1t          j        |j        | j                                                   }||k    r| S |j        }|j        }|                                s|                                rt          | ||          S |j        }|j        }||k    r2t          dt          |          z   dz   t          |          z             t          j        |                    | j        |                    |                    |          S )Nz!Cannot bitcast data-type of size z to data-type of size )rp   r  r*   r  r   r~  ry   r   primitive_bitwidthr)   r  r+   create_bitcastr   r  )r   r  r#   src_ty
src_sca_ty
dst_sca_tysrc_bitsdst_bitss           r   r   r   ,  s   ZF Mv}ej.I.I.K.KLLJJ ,j//11 ,E67+++,H,H8<s8}}L P. .03H> ? ? 	?9W++EL&,,w:O:OPPRXYYYr   fp_downcast_roundingc                   | j         }t          |t          j                  r|j        }t          |t          j                  r|j        }|                                r1t          j        |j        | j                                                   }||k    r| S |j        }|j        }t          |          }d}|
                                rP|
                                r<|j        |j        k     r,|t          j        j        }nL|t          j        j        k    rd}n4|2t          dt!          |          z   dz   t!          |          z             |                                s|                                r?|j                            d          	 
J d             |j        d         | |||          S |                                r|
                                s*|
                                r|                                s|rBt          j        |                    | j        |                    |          |          |          S |                                r|                                r(|                                r>|                                s*t9          t9          | t          j        |          ||          S |
                                o#|
                                o|j        |j        k    }|rAt          j        |                    | j        |                    |                    |          S |
                                o#|
                                o|j        |j        k     }	|	rAt          j        |                    | j        |                    |                    |          S |                                 r|                                 r|j!        |j!        k    s|j"        |j"        k    r|#                                o|$                                 }
|$                                rX| j%                            |          }t          j        |&                    |          | j%                  }tO          | ||          S t          j        |(                    | j        |                    |          |
          |          S |)                                r|                                 r|$                                rX| j%                            |          }t          j        |&                    |          | j%                  }tO          | ||          S |#                                rAt          j        |*                    | j        |                    |                    |          S t          j        |+                    | j        |                    |                    |          S |                                 r|)                                r|$                                s|#                                sAt          j        |,                    | j        |                    |                    |          S t          j        |-                    | j        |                    |                    |          S |.                                r|                                 r|j!        }|dk    rAt          j        |/                    | j        |                    |                    |          S |d	k    rZtO          t9          | t          j0        |          t          j        |1                    d
          t          j0                  |          S |                                 rU|.                                rAt          j        |2                    | j        |                    |                    |          S |.                                rU|.                                rAt          j        |3                    | j        |                    |                    |          S J d|  d|             )NFTz]fp_downcast_rounding should be set only for truncating fp conversions. Source scalar type is z and destination type is convert_custom_typesz0target doesn't provide conversion for this type.r   r   r   r   zcannot cast z to )4rp   re   r*   rn   rJ   r  r  r   r~  r  rz   r  r
   r  r  r)   r  is_fp8e4b15codegen_fnsgetrS   r+   create_fp_to_fpr   r  rQ   rP   rR   r   rM   create_fp_trunccreate_fp_extrT   r7   r8   r   is_boolr9   r  r9  create_int_castis_standard_floatingcreate_fp_to_sicreate_fp_to_uicreate_ui_to_fpcreate_si_to_fpry   create_ptr_to_intri   r   create_int_to_ptrr  )r   r  r#   r  r  r  r  use_custom_roundingtruncate_fpext_fpsign_extendtyr  bitwidths                 r   r   r   ?  s/   ZF&",'' &55 :39 Mv}ej.I.I.K.KLLJJ 11EFF yJ$:$: % % y

'**G
G
G'@P@U)=)=!R%5%:::RV<O+ 68;JHJefhklvhwhwx y y y 	   rJ$:$:$<$< r"&&"$ $+/0 01c0 0 0:w"#9:5&J^ipqqqq 	 u
 6 6 8 8 u  u%/%6%6%8%8uu y00v||G?T?TVjkkmsttt 	 KZ%7%7%9%9 KK%/%7%7%9%9KD
G44j'JJJ
 ((** F  F%
(EE   _y00v||G?T?TUUW]^^^ ##%% F  F%
(EE   ]y..u|V\\'=R=RSSU[\\\  pz0022 p:#:::j>W[e[t>t>t ..00M9K9K9M9M5M 	p""7++B711"55u{CCBUB0009W44U\6<<PWCXCXZeffhnooo &&(( cZ->->-@-@ c 	c""7++B711"55u{CCBUB000%%'' 	c9W44U\6<<PWCXCXYY[abbb9W44U\6<<PWCXCXYY[abbb  cz>>@@ c 	cz'?'?'A'A 	c9W44U\6<<PWCXCXYY[abbb9W44U\6<<PWCXCXYY[abbb  qz0022 q*r>>9W66u|V\\RYEZEZ[[]cdddq==T%7;;RYwGXGXYZG[G[]_]e=f=fhoppp  az0022 ay225<gAVAVWWY_```  ^z0022 ^y//fll7>S>STTV\]]]4444F44445r   c                    t           j        j        }| r[| dk    rt           j        j        }nC| dk    rt           j        j        }n+| dk    rt           j        j        }nt          d|  d          |S )Nz.ca.cgz.cvCache modifier  not supported)r
   CACHE_MODIFIERr   CACGCVr)   cache_modifiercaches     r   _str_to_load_cache_modifierr    s}    "E OU""%(EEu$$%(EEu$$%(EEM~MMMNNNLr   c                   t           j        j        }| rs| dk    rt           j        j        }n[| dk    rt           j        j        }nC| dk    rt           j        j        }n+| dk    rt           j        j        }nt          d|  d          |S )Nz.wbr  z.csz.wtr  r  )r
   r  r   WBr  CSWTr)   r  s     r   _str_to_store_cache_modifierr    s    "E 
OU""%(EEu$$%(EEu$$%(EEu$$%(EEM~MMMNNNLr   c                    t           j        j        }| rC| dk    rt           j        j        }n+| dk    rt           j        j        }nt          d|  d          |S )N
evict_lastevict_firstzEviction policy r  )r
   EVICTION_POLICYNORMAL
EVICT_LASTEVICT_FIRSTr)   )eviction_policyevictions     r   _str_to_eviction_policyr    se    !(H Ql**)4HH--)5HHOOOOPPPOr   c                    d }| rC| dk    rt           j        j        }n+| dk    rt           j        j        }nt	          d|  d          |S )NzeronanzPadding option r  )r
   PADDING_OPTIONPAD_ZEROPAD_NANr)   )padding_optionpaddings     r   _str_to_padding_optionr    s_    G OV##'0GGu$$'/GGM~MMMNNNNr   c                   t           j        j        }| rs| dk    rt           j        j        }n[| dk    rt           j        j        }nC| dk    rt           j        j        }n+| dk    rt           j        j        }nt          d|  d          |S )Nacquirereleaseacq_relrelaxedMemory semantic r  )r
   MEM_SEMANTICACQUIRE_RELEASEACQUIRERELEASERELAXEDr)   )
sem_optionsems     r   _str_to_semr    s    
/
)C 
L""/)CC9$$/)CC9$$/1CC9$$/)CCJ
JJJKKKJr   c                    t           j        j        }| r[| dk    rt           j        j        }nC| dk    rt           j        j        }n+| dk    rt           j        j        }nt          d|  d          |S )Ngpuctasysr  r  )r
   MEM_SYNC_SCOPEGPUCTASYSTEMr)   )scope_optionscopes     r   _str_to_scoper    s}    !E N5  %)EEU""%)EEU""%,EELLLLMMMLr   c                n   | rt          | d          s| g} d | D             } | D ]5}t          |t                    rd|cxk    rt          |          k     sn J 6t          |           dk    sJ t          |           t          t	          |                     k    s
J d            t          |           S dS )N__iter__c                T    g | ]%}t          |t          j                  r|j        n|&S r_   re   r*   rn   rJ   rY  elems     r   rZ  z0_canonicalize_boundary_check.<locals>.<listcomp>  s0    lllUY
4(F(FP$**Dlllr   r   z'Duplicate dimension in `boundary_check`r_   )hasattrre   r"   rL  setrx  )boundary_checkblock_shapedims      r   _canonicalize_boundary_checkr    s     &~z22 	.,-Nll]klll! 	H 	HCc3''GA,G,G,G,Gs;7G7G,G,G,G,G,GG,G>""Q&&&&>""c#n*=*=&>&>>>>@i>>>n%%%2r   c	           
        ||t          d          | j        j        j        }	|	t          j        k    s
J d            |	                                r$|t          j        j        k    rt          d          | j        j        }
t          ||

                                          }t          j        |                    | j        |||||          |
          S )NK`mask` and `other` arguments cannot be specified for loading block pointers3`tl.int1` should be rewrited in `tl.make_block_ptr`z@Padding option `nan` is not supported for integer block pointers)r)   rp   
element_tyr*   rg   rT   r
   r  r  r  r~  r+   create_tensor_pointer_loadr   )ptrmaskr   r  r  r  r  is_volatiler#   elt_tyr  s              r   _load_block_pointerr    s     5,fgggX +FRWS}} ]7b&7&???[\\\ X F 2.&BYBYB[B[\\N 9**3:~wPUW_almmouw w wr   c	           
     (   | j         j                                        s*t          d| j                                          d          ||t          d          |s|rt          d          | j                                         sT|r(|j                                         rt          d          |r(|j                                         rt          d          | j                                         rT|(t          || j                                         |          }|(t          || j                                         |          }| j         j        }	|	j        }
|
t          j
        k    }|r7t          j        }
t          j        |
|	j                  }	t          | |	|          } |t          ||
|          }| j                                         r/| j                                         }t          j        |
|          }n|
}|1t          j        |                    | j        |||          |          }n@t          j        |                    | j        |j        |r|j        nd |||          |          }|rt          |t          j
        |          }|S )NUnsupported ptr type z in `tl.load`z)`other` cannot be provided without `mask`z`padding_option` or `boundary_check` argument is not supported for loading a tensor ofpointers or loading a scalar. Because the compiler does not know the boundary; please use block pointers (defined by `make_block_ptr`) insteadEMask argument cannot be block type if pointer argument is not a blockzFOther argument cannot be block type if pointer argument is not a block)rp   r   ry   r)   r   r  r  r~  r  r*   rg   int8pointer_typeaddress_spacer   r  r+   create_loadr   create_masked_load)r  r	  r   r  r  r  r  r
  r#   ptr_tyr  r  r  r  r   s                  r   _load_legacyr  *  s   8?!!## US1B1B1D1DSSSTTT |)DEEE U. U T U U 	U
 8 g 	fDI&&(( 	fdeee 	gUZ((** 	gefff x V'ch.G.G.I.I7SSD(0I0I0K0KWUUE X_FF G ))=>>3(( UFG,, x ))++vu--  |i++CJxUUW]^^i&&sz4;PU@_[_afhp'24 45;= =  *3))Jr   r  r	  Optional[tl.tensor]r  r   r  r  r  r  r
  c	                &   t          |          }	t          |          }
t          |          }| j                                        r5| j        j                                        rt          | |||||	|
||	  	        S t          | |||||	|
||	  	        S rx   )	r  r  r  rp   ry   r  r  r  r  )r  r	  r   r  r  r  r  r
  r#   r  r  r  s               r   loadr  h  s     (77E&77H$^44G
x nSX099;; n"3e^WeU]_jlsttt CunguhXcelmmmr   desc_ptrc           	         t          ||d          }|                    | j        ||                    |          t	          |          t          |                    }t          j        ||          S NFrequire_i64)_convert_to_ir_valuescreate_descriptor_loadr   r  r  r  r*   r+   )r  offsetsr  r  rp   r#   rq   s          r   descriptor_loadr"  x  sk    #GW%HHHG&&xGATAT'B>'R'R'>'O'O	Q 	QA 9Qr   c                    t          ||d          }t          j        |                    | j        |j        |          t          j                  S r  )r  r*   r+   create_descriptor_storer   void)r  rJ   r!  r#   s       r   descriptor_storer&    sC    #GW%HHHG9W44X_elT[\\^`^efffr   global_addressbox_dimList[tl.tensor]
global_dimglobal_strideelement_stride	elem_typeinterleave_layoutswizzle_mode	fill_modec                   |r|d         j         t          j        k    sJ t          j        |
                    | j        |j        d |D             d |D             d |D             d |D             ||||	
  
        t          j                  S )Nr   c                    g | ]	}|j         
S r_   r   rX  s     r   rZ  z$tensormap_create.<locals>.<listcomp>  s    '''!QX'''r   c                    g | ]	}|j         
S r_   r3  rX  s     r   rZ  z$tensormap_create.<locals>.<listcomp>  s    ***!QX***r   c                    g | ]	}|j         
S r_   r3  rX  s     r   rZ  z$tensormap_create.<locals>.<listcomp>  s    ---!QX---r   c                    g | ]	}|j         
S r_   r3  rX  s     r   rZ  z$tensormap_create.<locals>.<listcomp>  s    ...!QX...r   )r9   r*   ri   r+   create_tensormap_creater   r%  )r  r'  r(  r*  r+  r,  r-  r.  r/  r0  r#   s              r   tensormap_creater8    s     Ba 0 6"( B B BB9''O!''w'''**z***--}---..~...	
 	
 	  r   c                p    t          j        |                    | j                  t           j                  S rx   )r*   r+   #create_tensormap_fenceproxy_acquirer   r%  )r  r#   s     r   tensormap_fenceproxy_acquirer;    s(    9W@@QQSUSZ[[[r   c           	     .   |t          d          | j        j                                        }|j                                        st          |||          }|j                                        s
J d            ||j                                        k    s(J d| d|j                                         d            | j        j        j        |j        j        k    s*J d| j        j        j         d|j        j         d            | j        j        j        }|t          j        k    s
J d            t          ||          }t          |||          }t          j
        |                    | j        |j        |||          t          j                  S )	Nr  z-Value argument must be block type or a scalarzBlock shape(z) and value shape(z
) mismatchzBlock element type(z) and value element type(r  )r)   rp   r  r~  r  r  r*   rg   r  r   r+   create_tensor_pointer_storer   r%  )	r  valr	  r  r  r  r#   r   r  s	            r   _store_block_pointerr?    s    fggg (%6688K8 >"3W==8OO OOO#(33     \k\\SX5N5N5P5P\\\  8)SX-@@@@  CqX[X`XkXv  Cq  Cq  RU  RZ  Re  Cq  Cq  Cq@@@X +FRWS 2.+NNN sFG
$
$C 9W88SZQ_afhpqqW  r   c           	        | j         j                                        s*t          d| j                                          d          |rt          d          | j                                         sR|j                                         rt          d          |r(|j                                         rt          d          | j                                         rRt          || j                                         |          }|(t          || j                                         |          }| j         j        }|j        }|t          j
        k    r7t          j        }t          j        ||j                  }t          | ||          } t          |||          }|s?t          j        |                    | j        |j        ||          t          j                  S |j         j                                        st          d          t          j        |                    | j        |j        |j        ||          t          j                  S )Nr  z in `tl.store`z`boundary_check` argument is not supported for storing a tensor of pointers or storing a scalar. Because the compiler does not know the boundary; please use block pointers (defined by `make_block_ptr`) insteadzFValue argument cannot be block type if pointer argument is not a blockr  z"Mask must have boolean scalar type)rp   r   ry   r)   r   r  r  r~  r  r*   rg   r  r  r  r   r+   create_storer   r%  r  create_masked_store)	r  r>  r	  r  r  r  r#   r  r  s	            r   _store_legacyrC    s'   8?!!## VT1B1B1D1DTTTUUU  B A B B 	B
 8 f8 	gefff 	fDI&&(( 	fdeee x T"3(A(A(C(CWMM'ch.G.G.I.I7SSDX_FF )=>>3(( sFG
$
$C  ay--cj#*eXVVXZX_```9##%% ?=>>>9W00SZV[]effhjhopppr   r>  c           	        t          |          }t          |          }| j                                        s| j        j                                        rt          d          | j                                        r3| j        j                                        rt          | ||||||          S t          | ||||||          S )N"Cannot store to a constant pointer)r  r  rp   is_constr   r)   ry   r  r  r?  rC  )	r  r>  r	  r  r  r  r#   r  r  s	            r   storerG    s     )88E&77H
x ?cho6688 ?=>>>
x WSX099;; W#CdNE8U\]]] S#t^UHgVVVr   cmpr  r  c           	        t          |          }t          |          }| j        j        j        }|j        dvrt          d          t          j        |	                    | j
        |j
        |j
        ||          |j                  S )N)   r=  r   z9atomic_cas only supports elements with width {16, 32, 64})r  r  rp   r   r  r  r)   r*   r+   create_atomic_casr   )r  rH  r>  r  r  r#   r  s          r   
atomic_casrL    sy    
c

C%  E+J$L88TUUU9W..sz3:szSVX]^^`c`hiiir   op&Tuple[tl.tensor, tl.tensor, tl.tensor]c                   | j         j                                        s)t          d| j                                         z             | j                                         s| j         j                                        rt          d          | j         j        j        }|t          j        u r|dk    rt          d|z   dz             |t          j	        t          j
        t          j        t          j        fv r%t          d|z   dz   t          |          z             | j                                         rT|(t          || j                                         |          }|(t          || j                                         |          }t#          || j         j        j        |          }|s|                    d          }t          j	        }| j                                         rc|                    || j                                                   }t          j        t          j	        | j                                                   }t          j        ||          }| ||fS )Nz)Pointer argument of store instruction is rE  r   atomic_z does not support fp16z does not support T)rp   r   ry   r)   r   rF  r  r*   rK   rg   r  int16rL   r  r  r  r~  r   rf   rM  r  r+   )r  r>  r	  rM  r#   r  mask_irmask_tys           r   atom_red_typechecking_implrT    s   8?!!## \DsxGXGXGZGZZ[[[
x ?ch1::<< ?=>>>+JRZB%KKR*BBCCCbgrw"+>>>R*>>ZPQQQ
x R'ch.G.G.I.I7SSD?&sCH,E,E,G,GQQC
sCHO.
8
8C +""4(('8 	J**7CH4M4M4O4OPPGmBGSX-F-F-H-HIIGy'**T>r   c                
   t          | ||d|          \  } }}t          |          }t          |          }|j        j        }|                                r|                                rPt          j        |	                    t          j        j        | j        |j        |j        ||          |j                  S t          j        |	                    t          j        j        | j        |j        |j        ||          |j                  S |t          j        t          j        hvrt#          d|           t%          g d||          }|t          j        k    rt          j        nt          j        }t+          |||          }	t+          | t          j        |d          |          }
|t          j        k    rt          j        nt          j        }t+          |||          }t+          | t          j        |d          |          }t3          |||          }t5          |||          }t          j        |	                    t          j        j        |
j        |	j        t7          |||          j        ||          |	j                  }t          j        |	                    t          j        j        |j        |j        t7          |||          j        ||          |j                  }t;          ||||          }t+          |||          S )Nr   z#atomic_max not supported for dtype rd   r   )rT  r  r  rp   r   rT   r   r*   r+   create_atomic_rmwr
   	ATOMIC_OPMAXr   UMAXrM   rO   r<   rk   r-   ri   r   r  rh   rj   r   r,  r   UMINwherer  r>  r	  r  r  r#   sca_tyr  i_typei_vali_ptrui_typeui_valui_ptrposnegpos_retneg_retr   s                      r   
atomic_maxrh  ,     /S$wOONCd
c

C%  EX_F}} y!! 	y9))",*:CJ
TXT_adfkllnqnvx x x 9))",*;SZUYU`beglmmorowy y y
 bj"*---FfFFGGGC))D2:--RXX28FC))EC33W==E!RZ//biiRYGS'7++FS"/'155w??F
T7
+
+C
Cw
'
'Ci!!","2EL%,"&tS'":":"A3	O 	OPUPZ\ \G i!!","3V]FM"&tS'":":"A3	O 	OPVP[] ]G Wgw
/
/C3(((r   c                
   t          | ||d|          \  } }}t          |          }t          |          }|j        j        }|                                r|                                rPt          j        |	                    t          j        j        | j        |j        |j        ||          |j                  S t          j        |	                    t          j        j        | j        |j        |j        ||          |j                  S |t          j        t          j        hvrt#          d|           t%          g d||          }|t          j        k    rt          j        nt          j        }t+          |||          }	t+          | t          j        |d          |          }
|t          j        k    rt          j        nt          j        }t+          |||          }t+          | t          j        |d          |          }t3          |||          }t5          |||          }t          j        |	                    t          j        j        |
j        |	j        t7          |||          j        ||          |	j                  }t          j        |	                    t          j        j        |j        |j        t7          |||          j        ||          |j                  }t;          ||||          }t+          |||          S )Nr   z#atomic_min not supported for dtype rd   r   )rT  r  r  rp   r   rT   r   r*   r+   rV  r
   rW  MINr   rZ  rM   rO   r<   rk   r-   ri   r   r  rh   rj   r   r,  r   rY  r[  r\  s                      r   
atomic_minrl  S  ri  r   c           
     r   t          | ||d|          \  } }}t          |          }t          |          }|j        j        }|                                rt          j        j        nt          j        j	        }t          j        |                    || j        |j        |j        ||          |j                  S )Nr   )rT  r  r  rp   r   rz   r
   rW  FADDADDr*   r+   rV  r   )r  r>  r	  r  r  r#   r]  rM  s           r   
atomic_addrp  z  s    /S$wOONCd
c

C%  EX_F$0022	H		8HB9W..r3:sz4;X[]bccehemnnnr   c           
        t          | ||d|          \  } }}t          |          }t          |          }t          j        |                    t          j        j        | j	        |j	        |j	        ||          |j
                  S )Nand)rT  r  r  r*   r+   rV  r
   rW  ANDr   rp   r  r>  r	  r  r  r#   s         r   
atomic_andru    y    /S$wOONCd
c

C%  E9W..r|/?SZY]YdfikpqqX  r   c           
        t          | ||d|          \  } }}t          |          }t          |          }t          j        |                    t          j        j        | j	        |j	        |j	        ||          |j
                  S )Nor)rT  r  r  r*   r+   rV  r
   rW  ORr   rp   rt  s         r   	atomic_orrz    sx    /S$gNNNCd
c

C%  E9W..r|
CJX\XcehjoppX  r   c           
        t          | ||d|          \  } }}t          |          }t          |          }t          j        |                    t          j        j        | j	        |j	        |j	        ||          |j
                  S )Nxor)rT  r  r  r*   r+   rV  r
   rW  XORr   rp   rt  s         r   
atomic_xorr~    rv  r   c           
        t          | ||d|          \  } }}t          |          }t          |          }t          j        |                    t          j        j        | j	        |j	        |j	        ||          |j
                  S )Nxchg)rT  r  r  r*   r+   rV  r
   rW  XCHGr   rp   rt  s         r   atomic_xchgr    sy    /S$PPNCd
c

C%  E9W..r|/@#*cjZ^ZegjlqrrX  r   c                    |                                  |j        j        v sJ d|j        j         d|              |                                 } | dk    rd} t	          t
          j        |           S )Nzinput_precision must be one of z. Got TF32X3TF32x3)lowerr   allowed_dot_input_precisionsupperrG  r
   INPUT_PRECISION)input_precisionr#   s     r   _str_to_dot_input_precisionr    sz      ""go&RRRRo'/*Voo^moo SRR%++--O("""2%777r   accr  max_num_imprecise_acc	out_dtypec           
        | j                                         r|j                                         sJ | j                                        r|j                                        rn| j        t          j        t          j        t          j        t          j        t          j	        fv sJ d| j                     |j        t          j        t          j        t          j        t          j        t          j	        fv sJ d|j                     | j        |j        k    sJ d| j         d|j                     | j        
                                s|j        
                                r6t          | t          j        |          } t          |t          j        |          }||j        j        }t          ||          }t          | j                  }t          |j                  }||cxk    rdk    s,n ||cxk    rdk    sn J d| j         d|j         d	            | j        d
         j        |j        d         j        k    sAJ d| j         d|j         d| j        d
         j         d|j        d         j         d		            |j                            d          
J d             |j        d         | j         |j                   }	| j        d         j        |	d         k    r8| j        d
         j        |	d         k    r|j        d
         j        |	d         k    s%J d|	d          d|	d          d|	d                      | j         j                                        rF| j         j        t          j        k    s
J d            |                    d          }
t          j        }n|                                rt3          d          | j         j                                        s| j         j                                        r"|                    d          }
t          j	        }n@|                                r|                    d          n|                    d          }
|}| j         j        d         }|j         j        d
         }| j         j        d
         }|dk    r| j         j        d         nd }t	          j        ||r|||gn||g          }| |                    |
|r|||gn||g          }n|j         }|j         |k    sJ |B| j                                        r&|j                                        r|j        j!        }nQd}nN| j                                        r5|j                                        r||k    rt3          d| d| d	          t	          j"        |#                    | j         |j         |||          |          S )NzUnsupported lhs dtype zUnsupported rhs dtype z&Both operands must be same dtype. Got r   r	      +Both inputs must be either 2D or 3D; (lhs: 	 vs rhs: r  re  zFirst input shape (z) and second input shape z= are not compatible for matmul (second index of first shape (z0) must be equal to first index of second shape (min_dot_sizez2target doesn't provide lower shape bounds for dot.r   r   zInput shapes should have M >= z, N >= z
 and K >= zonly int8 supported!zhout_dtype=bfloat16 is unsupported. Please use out_dtype=float32/float16 and cast with `.to(tl.bfloat16)`zmax_num_imprecise_acc (z) must be <= K ()$rp   r  r9   rS   r*   r  uint8rK   rL   rM   r  r   r   default_dot_input_precisionr  rL  r  rJ   r  r  r   rT   	get_int32r-   rR   r)   rP   get_fp32rQ   get_fp16r  rM  r   max_num_imprecise_acc_defaultr+   
create_dot)r|   r~   r  r  r  r  r#   lhs_rankrhs_rankr  r  ret_scalar_tyMNKBr   
acc_handles                     r   dotr    sY   8638#4#4#6#6666
y lci..00 lyRWbh
BKZ) ) ) )*N39*N*N) ) )yRWbh
BKZ) ) ) )*N39*N*N) ) )yCI%%%'kPSPY'k'k`c`i'k'k%%%
y -#)"7"7"9"9 -3
G,,3
G,,!/E1/7KKO39~~H39~~Hx$$$$1$$$$H(A(A(A(A(A(A(A(A(A  DRqtqz  DR  DR  FI  FO  DR  DR  DR(A(AA9R=#)
#   q  q  qSY  q  q  VY  V_  `b  Vc  Vi  q  q  [^  [d  eg  [h  [n  q  q  q  "">22>>@t>>>67&~6sxJJL9R=,q/11cim6I\Z[_6\6\IbM<?222q\!_qq\RS_qq`lmn`oqq 323 x "x"')))+A)))q!!					 "vx x 	x		 	 	"	" "cho&=&=&?&? "a  
$-$5$5$7$7PWa   W=M=Ma=P=P!rArArA%]]qA]=q*D1a))q!fEEF
{))"1.Hq!Qii1a&II

Z
x6!!!! $9 	&#)"2"2"4"4 	&$+O$Q!!$%!!9 	d#)"2"2"4"4 	d9NQR9R9Rb7Lbb^_bbbccc9W''
CJ
O]rss  r   float_formatc                   | dk    rt           j        j        S | dk    rt           j        j        S | dk    rt           j        j        S | dk    rt           j        j        S | dk    rt           j        j        S t          d|  d          )Ne4m3e5m2e2m3e3m2e2m1zInvalid float format: r^   )r
   F8F6F4TYE4M3E5M2E2M3E3M2E2M1r)   )r  s    r   _str_to_fp_typer    s    v{v{v{v{v{
=l===
>
>>r   	lhs_scale	rhs_scaletl.tensor | Nonec	                   | j                                         r|j                                         sJ t          | j                  }	t          |j                  }
|	|
cxk    rdk    s,n |	|
cxk    rdk    sn J d| j         d|j         d            t	          |          }t	          |          }|dv sJ d|             |dv sJ d	|             t          |t          j                  o|j        d u }|s
J d
            | j         j        d         }|j         j        dd          \  }}|dk    rdnd}||| j         j        d         z  k    sJ d| j         d|j         d            |dk    sJ d|            |	dk    r| j         j        d         nd }t          j	        ||r|||gn||g          }|
                    d          }| |                    ||r|||gn||g          }n|j        }|j         |k    sJ |rd n|j        }t          j        |                    | j        |j        ||j        |||          |          S )Nr	   r  r  r  r  )r  r  r  zNYI: lhs_format )r  r  zNYI: rhs_format zNYI: rhs_scale not supportedr  r  r   re  zCReduction dimension should pack the same number of elements; (lhs: r   z!scaled_dot NYI for K < 64. Got K=r   )rp   r  rL  r  r  re   r*   rn   rJ   r  r  rM  r   r+   create_dot_scaled)r|   r  
lhs_formatr~   r  
rhs_formatr  r  r#   r  r  lhs_format_enumrhs_format_enumrhs_scale_is_noner  r  r  PACKEDr  r   r  r  rhs_scale_handles                          r   
dot_scaledr    s   8638#4#4#6#666639~~H39~~Hx$$$$1$$$$H(A(A(A(A(A(A(A(A(A  DRqtqz  DR  DR  FI  FO  DR  DR  DR(A(AA%j11O%j11O11113Rj3R3R111))))+Jj+J+J)))"9bl;;W	SW@W<<<<<rA8>"##DAq&&QQAF
    sSVS\ssgjgpsss  777:a::777%]]qA]91&@q!Qii1a&AAF			!		B
{))"1.Hq!Qii1a&II

Z
x6!!!!0Ftti6F9!!#*i.>QTQ[]m"1:	? 	?@FH H Hr   	conditionc                   | j         t          j        k    rt          j        d| j                     t          | t          j        |          } t          |||dd          \  }}| j                                        r)t          | ||          \  } }t          |||          \  }}nt          | ||          \  } }|j        }t          j
        |                    | j        |j        |j                  |          S )Nzgtl.where with a non-boolean condition is deprecated and will error out in a future triton release. Got T)r9   r*   rg   warningswarnr   r   rp   r  r   r+   create_selectr   )r  rq   r   r#   r  r   s         r   r[  r[  0  s    "'!! Hv  wF  H  H	
 	
 	
 Y11I'1gtTBBDAq~   C+Iq'BB	1#Aq'2211+Iq'BB	1VF9W**9+;QXqxPPRXYYYr   c                `    |rt          j        ||          }n|}t          j        | |          S rx   )r*   r  r+   )rq   rU   r  res_tys       r   wrap_tensorr  F  s6     y)44 9Qr   inputsSequence[tl.tensor]Tuple[tl.tensor, ...]c                "    t          fd D                        d d         j        j        t                    }|k     sJ d| d            fdt	                    D             t          fd D                       s
J d                                d  D                        |                                            t           fd	t          t                               D                       S )
Nc              3  T   K   | ]"}t          ||j        j        gd           V  #dS )Trf  N)rS  rF  rJ   )rY  tr#   s     r   rv  zreduction.<locals>.<genexpr>Q  s;      ffZ[wq17=/tWUUUffffffr   r   z&reduction axis must be < inputs rank (r  c                &    g | ]\  }}|k    |S r_   r_   )rY  r  rR  r!   s      r   rZ  zreduction.<locals>.<listcomp>W  s"    ===tq!199999r   c              3  8   K   | ]}|j         j        k    V  d S rx   )rp   r  )rY  r  r  s     r   rv  zreduction.<locals>.<genexpr>X  s,      55qv|u$555555r   z-all reduction inputs must have the same shapec                    g | ]	}|j         
S r_   r3  rY  r  s     r   rZ  zreduction.<locals>.<listcomp>Z  s    &@&@&@Aqx&@&@&@r   c              3     K   | ]8}t                              |          |         j        j                  V  9d S rx   r  
get_resultrp   r   )rY  r  r  	reduce_opr  s     r   rv  zreduction.<locals>.<genexpr>^  sG      tt\]Y11!44fQin6KYWWttttttr   )	tuplerp   r  rL  r  allcreate_reduceverifyrA  )r  r!   region_builder_fnr#   rankr  r  r  s   `` ` @@@r   	reductionr  O  sC   |ffff_efffff1IN Eu::D$;;;HHHH;;;====y//===I5555f55555ff7fff5%%&@&@&@&@&@$GGIi   ttttttafgjkqgrgrasasttttttr   reversec                     d         j         j        t                    }| |cxk    r|k     sn J d| d| d            |dk     r||z  } D ]}|j         j        k    s
J d            |                    d  D             ||           |                                            t           fdt          t                               D                       S )Nr   z
scan axis z must be < inputs rank (r  z(all scan inputs must have the same shapec                    g | ]	}|j         
S r_   r3  r  s     r   rZ  z$associative_scan.<locals>.<listcomp>s  s    "<"<"<18"<"<"<r   c              3     K   | ]8}t                              |          |         j        j                  V  9d S rx   r  )rY  r  r  scan_opr  s     r   rv  z#associative_scan.<locals>.<genexpr>w  sG      nnVWW//22F1IN4I5QQnnnnnnr   )rp   r  rL  create_scanr  r  rA  )	r  r!   r  r  r#   r  r  r  r  s	   `      @@r   associative_scanr  f  s(   1IN Eu::D5D4!Sd!S!SD!S!S!Saxx Q Qv|u$$$&P$$$$!!"<"<V"<"<"<dGLLGgNNnnnnnn[`adekalal[m[mnnnnnnr   num_binsc                $   t          | j                  dk    s
J d            | j                                        s
J d            t	          j        |                    | j        |          t	          j        t          j	        |f                    S )Nr   z histogram only supports 1D inputz%histogram only supports integer input)
rL  r  r9   rT   r*   r+   create_histogramr   r  r-   )r   r  r#   s      r   	histogramr    s    u{q   "D   ;HH!HHH9W--elHEEr}UWU]`h_kGlGlmmmr   valuesc                   t          dt          | j                            t          |          k    rt          d          | j                            dt          j        || j                                                             | S )Nr   zAShape of input to multiple_of does not match the length of valuesztt.divisibility)	r   rL  r  r)   r   set_attrr
   	make_attrget_contextrq   r  s     r   multiple_ofr    sl    
1c!'lls6{{**\]]]H'fah>R>R>T>T)U)UVVVHr   c                    t          | j                  t          |          k    rt          d          | j                            dt          j        || j                                                             | S )NzDShape of input to max_contiguous does not match the length of valuesztt.contiguityrL  r  r)   r   r  r
   r  r  r  s     r   max_contiguousr    sa    
17||s6{{""_```Hor|FAH<P<P<R<R'S'STTTHr   c                    t          | j                  t          |          k    rt          d          | j                            dt          j        || j                                                             | S )NzCShape of input to max_constancy does not match the length of valuesztt.constancyr  r  s     r   max_constancyr    sa    
17||s6{{""^___Hnbl618;O;O;Q;Q&R&RSSSHr   c                d    t          j        |                                 t           j                  S rx   )r*   r+   create_barrierr%  )r#   s    r   debug_barrierr    s"    9W++--rw777r   prefixargshexc                v   |                      d          s|r| dz  } |                      d          s|r| d d         dz   } t          |           dk    r|                     d          sd| z   } d |D             }d |D             }t          j        |                    | |||          t          j                  S )N r}  re  r	   c                    g | ]	}|j         
S r_   r3  rY  args     r   rZ  z device_print.<locals>.<listcomp>  s    +++s
+++r   c                    g | ]B}|j         t          j        t          j        t          j        t          j        t          j        fv CS r_   )r9   r*   rg   r  rQ  r-   ri   r  s     r   rZ  z device_print.<locals>.<listcomp>  s5    ___SVrw28RXNN___r   )endswithrL  
startswithr*   r+   create_printr%  )r  r  r  r#   new_args	is_signeds         r   device_printr    s     ??3 D #??4   $T $t#
6{{Qv0055v++d+++H__Z^___I9W))&#xKKRWUUUr   r   r   c                    |j         j        sd S t          j        |                    | j        |          t          j                  S rx   )r   debugr*   r+   create_assertr   r%  )r   r   r#   s      r   r   r     s:    ?  9W**4;<<bgFFFr   c                p    t          j        |                    | j                  t           j                  S rx   )r*   r+   create_assumer   r%  )r   r#   s     r   assumer    s&    9W**4;77AAAr   c                z   t          |t                    rt          j        |          }t          |t          j                  r|rAd|j        cxk    rdk     sn J d|j         d            |                     |j                  S d|j        cxk    rdk     sn J d|j         d            |                     |j                  S t          |t          j                  r|j        j        dk    s
J d	            |j	        
                                s
J d
            |j	        t          j        k    rG|rE|                     |j        |                                 |j	                                                  S |j	        t          j        k    r|s
J d            |j        S J dt#          |                       )Nr\   r]   z@Block pointers only support 64 bit `shape/strides`, got a value z which is out of the rangerZ   r[   zFBlock pointers only support 32 bit `offsets/block_shape`, got a value r   z*Expected a scalar in shape/strides/offsetsz8Expected an integer scalar type in shape/strides/offsetsFzzBlock pointers only support 32 bit `offsets/block_shape`, add a `.to(tl.int32)` or use regular indexing for 64 bit supportz3Unsupported element type in shape/strides/offsets: )re   r"   r*   rn   rJ   r   r  r+   rF  r9   rT   ri   r  r   get_int64_tyr   r-   rp   )r#   r  r  s      r   _convert_elem_to_ir_valuer    s   $ "|D!!$%%  	1TZ////%///// 2F#z2F 2F 2F///$$TZ000TZ////%///// 2F#z2F 2F 2F///$$TZ000	D")	$	$ z1$$$&R$$$z  ""^^$^^^":!!k!**4;8L8L8N8NPTPZPhPhPjPjkkkZ28##K#S S S S5{TTT

TTTT5r   c                h     t          |d          r fd|D             S t           |          gS )Nr  c                2    g | ]}t          |          S r_   )r  )rY  r  r#   r  s     r   rZ  z)_convert_to_ir_values.<locals>.<listcomp>  s&    \\\$)'4EE\\\r   )r  r  )r#   	list_liker  s   ` `r   r  r    sJ    y*%% ]\\\\\R[\\\\%gy+FFGGr   basec           	        t          ||          }t          ||          }t          ||d          }| j                                        r| j        j                                        rt          d          | j        j        t          j        k    r8t          | t          j	        t          j
        | j        j                  |          } t          d          sgd D             t          d D                       s
J d            t          |d          s|g}d |D             }t          |          t          t!          t#          |                              k    s
J d	            t          fd
||||fD                       s
J d            |                    | j        ||||          }t          j        |t          j	        t          j        | j        j                                      S )NFr  zMExpected `base` to be a pointer type (but not a block pointer type or others)r  c                T    g | ]%}t          |t          j                  r|j        n|&S r_   r  r  s     r   rZ  z"make_block_ptr.<locals>.<listcomp>  s/    bbbdD",!?!?I4::Tbbbr   c              3  `   K   | ])}t          |t                    od |cxk    odk     nc V  *dS )rZ   r[   N)re   r"   r  s     r   rv  z!make_block_ptr.<locals>.<genexpr>  sM      XXDz$$$?4)?)?)?)?%)?)?)?)?XXXXXXr   zGExpected a list of constant integers (`int32_t` range) in `block_shape`c                T    g | ]%}t          |t          j                  r|j        n|&S r_   r  r  s     r   rZ  z"make_block_ptr.<locals>.<listcomp>  s/    VVV:dBL99CTZZtVVVr   z<Expected a permutation of (0, 1, ..., len(order)-1) in orderc              3  X   K   | ]$}t                    t          |          k    V  %d S rx   )rL  )rY  r  r   s     r   rv  z!make_block_ptr.<locals>.<genexpr>  s6      ddis;3y>>1ddddddr   zBExpected shape/strides/offsets/block_shape to have the same length)r  rp   ry   r  r  r)   r*   rg   r   r  r  r  r  r  rx  ry  rA  rL  create_make_block_ptrr   r+   r  )r  r  stridesr!  r   orderr#   r   s       `   r   make_block_ptrr    s    "'511E#GW55G#GW%HHHG 9 j!5!>!>!@!@ jhiii yrw&&D"/"'493JKKWUU ;
++ $"mbbVabbbKXXKXXXXX R RQR RX 5*%% VVPUVVVE%==Ds5zz!2!2333335s333 ddddE7T[]bCcddddd M MLM Md **4;wQ\^cddF9VR_R]49;OQ\-]-]^^___r   c                    t          ||d          }t          j        |                    | j        |          | j                  S r  )r  r*   r+   create_advancer   rp   )r  r!  r#   s      r   advancer    s>    #GW%HHHG 9W++DKAA49MMMr   )r!   r"   r#   r$   r%   r&   )r3   r4   r5   r4   r%   r4   )r3   r4   rB   rC   r5   r4   rD   rC   rE   rC   r%   r4   )T)rX   rC   )r   r4   r   r4   ru   rC   r%   rv   )FFTF)r|   r}   r~   r}   r#   r$   r%   r   )r|   r&   r~   r&   r#   r$   r   r   )
r   r}   r   r}   r   rC   r#   r$   r%   r&   )r   r}   r   r}   r#   r$   r%   r&   )
r   r}   r   r}   r   rC   r#   r$   r%   r&   )rq   r&   r   r&   r   r   r#   r$   )
rq   r&   r   r&   r   r&   r   r   r#   r$   )r   r&   r   r&   r#   r$   r%   r   )r   r&   r   r&   r#   r$   r%   r&   )r   r&   r#   r$   )r   r&   r%   r&   )r   r&   r#   r$   r%   r&   )r   r&   r#   r&   r%   r&   )r  r&   r%   r  )r:  r"   r;  r"   r#   r$   r%   r&   )r  rC  r9   r4   r#   r$   r%   r&   )rJ   r&   r  rC  r#   r$   r%   r&   )
r   r&   rN  rC  rO  rC   r#   r$   r%   r&   )r   r&   r!   r"   r#   r$   r%   r&   )
r|   r&   r~   r&   rO  rC   r#   r$   r%   r&   )rb  r&   rc  r&   r#   r$   r%   r&   )rb  r&   r#   r$   r%   r   )r   r&   rq  rr  r#   r$   r%   r&   )r   r&   r  rC  r#   r$   r%   r&   )r|   r&   r~   r&   r#   r$   r%   r&   )r  r  )r   r&   r  r4   r#   r$   r%   r&   rx   )
r   r&   r  r4   r#   r$   r  r  r%   r&   )r  r&   r	  r  r   r  r  r   r  r  r  r  r  r  r
  rC   r#   r$   r%   r&   )
r  r&   r  r  r  r  r#   r$   r%   r&   )r  r&   rJ   r&   r#   r$   r%   r&   )r  r&   r'  r&   r(  r)  r*  r)  r+  r)  r,  r)  r-  r"   r.  r"   r/  r"   r0  r"   r#   r$   r%   r&   )r  r&   r#   r$   r%   r&   )r  r&   r>  r&   r	  r  r  r  r  r  r#   r$   r%   r&   )r  r&   rH  r&   r>  r&   r  r  r  r  r#   r$   r%   r&   )r  r&   r>  r&   r	  r&   rM  r  r#   r$   r%   rN  )r  r&   r>  r&   r	  r&   r  r  r  r  r#   r$   r%   r&   )r|   r&   r~   r&   r  r&   r  r  r  r"   r  r4   r#   r$   r%   r&   )r  r  )r|   r&   r  r&   r~   r&   r  r  r  r  r  r4   r#   r$   r%   r&   )
r  r&   rq   r&   r   r&   r#   r$   r%   r&   )r  r  r!   r"   r#   r$   r%   r  )
r  r  r!   r"   r  rC   r#   r$   r%   r  )r   r&   r  r"   r#   r$   r%   r&   )rq   r&   r  rC  r%   r&   )r#   r$   r%   r&   )
r  r  r  r)  r  rC   r#   r$   r%   r&   )r   r&   r   r  r#   r$   r%   r&   )r  r&   r#   r$   r%   r&   )u
__future__r   r  typingr   r   r   r   r   r   _C.libtritonr
    r   r*   r   r   	Exceptionr   r/   r2   rA   rW   ro   r{   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r  r  r  r  r  r   r   r  r"  r   r,  r   r4  r9  rB  rk   rI  rS  r]  ra  rk  rp  r{  r  r   r  r   r   r  r  r  r  r  r  r  r  r  r  r"  r&  r8  r;  r?  rC  rG  rL  rT  rh  rl  rp  ru  rz  r~  r  r  r  r  r  r[  r  r  r  r  r  r  r  r  r  r   r  r  r  r  r  r_   r   r   <module>r$     s	   " " " " " "  ; ; ; ; ; ; ; ; ; ; ; ; ; ;                   GCLLF F F F F	 F F FD D D DF F F F@ @ @ @ /, /, /, /,d# # # # #V	< 	< 	< 	< ]a,1    B& & & &&: : : :64 4 4 4$4 4 4 4R R R R4: : : :& & & &4 4 4 485 5 5 5$5 5 5 5$	] 	] 	] 	]"   Q Q Q Q
P P P P
Q Q Q Q
' ' ' '& & & &" " " "R R R R
R R R R
Q Q Q Q   ) ) ) )$ $ $ $) ) ) )4 4 4 44 4 4 44 4 4 44 4 4 4	4 	4 	4 	4	4 	4 	4 	4"D D D D"( ( ( (.H H H H[ [ [ [M M M MK K K K   2
 
 
 
I I I IL L L L$2 2 2 2tp p p pZ Z Z Z( 04l5 l5 l5 l5 l5h     	 	 		 	 	     
 
 
w w w,; ; ;|n n n n    g g g g
   >\ \ \ \  :)q )q )qXW W W W,j j j j   6$) $) $) $)N$) $) $) $)No o o o            8 8 8E E E EP? ? ? ? H  H  H  HPZ Z Z Z,     u u u u.o o o o2n n n n         8 8 8 8V V V VG G G GB B B BU U U0H H H H$` $` $` $`NN N N N N Nr   