+
    ړLj                       ^ RI Ht ^ RIt^ RIt^ RIHtHtHtHtH	t	H
t
Ht ^ RIt^ RIHt ^RIHt ^RIHt ]	! R4      t]	! R4      t ! R	 R
]4      t ! R R]
],          4      tR# )    )annotationsN)ListOptionalSequenceTupleTypeVarGenericType)driver)ir)coreTTensorTyc                  (   a  ] tR t^tV 3R ltRtV ;t# )IncompatibleTypeErrorImplc                	   < Wn         W n        R V P                   P                  4       ,           R,           V P                  P                  4       ,           V n        \        \
        V `  V P                  4       R# )zinvalid operands of type  and N)type_atype_b__repr__messagesuperr   __init__)selfr   r   	__class__s   &&&D/app/.local/lib/python3.14/site-packages/triton/language/semantic.pyr   "IncompatibleTypeErrorImpl.__init__   sT    2T[[5I5I5KKgUX\XcXcXlXlXnn'7E    )r   r   r   )__name__
__module____qualname____firstlineno__r   __static_attributes____classcell__)r   s   @r   r   r      s    F Fr   r   c                  t   ] tR t^t$ ]P
                  tR]R&   ]tR]R&   R tR R lt	R R	 lt
R
 R ltR R ltRR ltR tR R ltRR R lltR R ltR R ltR R ltR R ltR R ltR R ltR  R! ltR" R# ltR$ R% ltR& R' ltR( R) ltR* R+ ltR, R- ltR. R/ ltR0 R1 ltR2 R3 lt R4 R5 lt!R6 R7 lt"R8 R9 lt#R: R; lt$R< R= lt%R> R? lt&R@ RA lt'RB RC lt(RD RE lt)RF RG lt*RH RI lt+RJ RK lt,RL RM lt-RN RO lt.RP RQ lt/RRRS/RT RU llt0RV RW lt1RX RY lt2RZ R[ lt3R\ R] lt4R^ R_ lt5R` Ra lt6Rb Rc lt7Rd Re lt8Rf Rg lt9Rh Ri lt:Rj Rk lt;Rl Rm lt<Rn Ro lt=Rp Rq lt>Rr Rs lt?RRt Ru llt@Rv tARw tBRx tCRy tDRz tER{ tFR| tGR} tHR~ tIR R ltJR R ltKR R ltLR R ltMR R ltNR tOR tPR R ltQR R ltRR R ltSR R ltTR R ltUR R ltVR R ltWR tXR tYR R ltZR R lt[R R lt\R R lt]R R lt^R R lt_R R lt`R R ltaR R ltbR R ltcR R ltdR teR R ltfR R ltgR R lthR tiR R ltjR R ltkR tlR R ltmR R ltnR R ltoR tpR R ltqR R ltrR R ltsR R lttR R ltuR R ltvR R ltwR R ltxR R ltyR tzRR lt{R R lt|R R lt}RR R llt~RtRS# )TritonSemanticzType[TensorTy]tensorz
ir.builderbuilderc                	    Wn         R # Nr(   )r   r(   s   &&r   r   TritonSemantic.__init__    s    r   c                    V ^8  d   QhRRRR/#    axisintreturnr    )formats   "r   __annotate__TritonSemantic.__annotate__'   s     O Os Ox Or   c                	    VR9  d   \        RV 24      hV P                  V P                  P                  V4      \        P
                  4      # )r   z+program_id axis must be 0, 1, or 2 but got r      r/   )
ValueErrorr'   r(   create_get_program_idtlint32r   r0   s   &&r   
program_idTritonSemantic.program_id'   sA    y J4&QRR{{4<<==dCRXXNNr   c                    V ^8  d   QhRRRR/# r.   r3   )r4   s   "r   r5   r6   ,   s     Q Q Q Qr   c                	    VR9  d   \        RV 24      hV P                  V P                  P                  V4      \        P
                  4      # )r   z-num_programs axis must be 0, 1, or 2 but got r8   )r:   r'   r(   create_get_num_programsr<   r=   r>   s   &&r   num_programsTritonSemantic.num_programs,   sA    y LTFSTT{{4<<??ErxxPPr   c               $    V ^8  d   QhRRRRRR/# )r/   a_tytl.dtypeb_tyr2   r3   )r4   s   "r   r5   r6   5   s&     D D D Dh Dr   c                	p   VP                   pVP                   pVP                  pVP                  pWV8X  d   W48  d   V# T# V\        P                  P                  P
                  8X  d   W48  d   V# T# V\        P                  P                  P
                  8X  d   WC8  d   V# T# \        R V RV 24      h)zunexpected signedness r   )int_bitwidthint_signednessr<   dtype
SIGNEDNESSUNSIGNED	TypeError)r   rG   rI   a_rankb_ranka_snb_sns   &&&    r   integer_promote_impl#TritonSemantic.integer_promote_impl5   s    """""""" <!?444RXX((111!+455RXX((111!+4550eD6BCCr   c               0    V ^8  d   QhRRRRRRRRRRRR/# )	r/   rG   rH   a_is_scalarboolrI   b_is_scalar
div_or_modr2   r3   )r4   s   "r   r5   r6   D   s9     05 05( 05 05X 05dh 05*.053;05r   c                	   W$8w  d   V'       d   W3MW13w  rgVP                  4       P                  VP                  4       P                  8:  d@   V'       d6   V\        P                  \        P                  39   d   \        P
                  # V# VP                  4       '       g   VP                  4       '       d   \        P                  # VP                  4       '       g   VP                  4       '       d   \        P
                  # VP                  4       '       g   VP                  4       '       d)   V'       d   \        P
                  # \        P                  # VP                  4       '       d?   VP                  4       '       d)   V'       d   \        P
                  # \        P                  # VP                  4       '       g   VP                  4       '       d   \        P
                  # VP                  4       '       d/   VP                  4       '       d   W8X  d   V# \        P                  # VP                  4       '       d   VP                  4       '       g   \        R V RV 24      hV'       d_   VP                  VP                  8w  dD   \        RVP                  4       ,           R,           VP                  4       ,           R,           4      hV P!                  W4      # )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_intrP   rL   r   rU   )r   rG   rX   rI   rZ   r[   	scalar_ty	tensor_tys   &&&&&&  r   computation_type_impl$TritonSemantic.computation_type_implD   s   
 %3>D<TL I~~%%)9)?)??9R[[0I#I::%   <<>>T\\^^:: <<>>T\\^^:: <<>>T\\^^zz!zz!<<>>dllnnzz!{{"<<>>T\\^^::;;==T[[]]<47RZZ7{{}}DKKMM.tfE$@AA $--1D1DD9DMMOKgUX\XeXeXggoo p p ((44r   c                	j   \        WP                  4      '       d   V# \        V\        P                  4      '       d   VP                  MTp\        V\
        \        \        34      '       d$   V P                  V4      pV P                  WR 7      # V'       d   \        RV R\        V4       R24      hV# )rM   cannot convert 	 of type 
 to tensor)
isinstancer'   r<   	constexprr`   r1   floatrY   to_tensor_typescalar_constantrP   type)r   x
check_typerM   s   &&& r   	to_tensorTritonSemantic.to_tensorv   s    a%%H!!R\\22AGGa#ud+,,''*E'''77oaS	$q'*MNNr   c                	   \        V\        P                  4      '       d   V# \        V\        P                  4      '       d   VP                  p\        V\
        4      '       d   \        P                  # \        V\        4      '       d   R	Tu;8:  d   R8  d    \        P                  #  RTu;8:  d   R
8  d    \        P                  #  RTu;8:  d   R8  d    \        P                  #  RTu;8:  d   R8  d    \        P                  #  \        RV R24      h\        V\        4      '       d   RpR^^,          ,          p\        P                  ! V4      pV\        R4      8X  g>   VR8X  g7   W8w  g1   Y$u;8:  d   V8:  d3    \        P                   #  \        P"                  # \        P                   # \        P"                  # \%        RV R\'        V4       R24      h)r/   zNonrepresentable integer .infg        rq   rr   rs              l                             l            g      8g   ?)rt   r<   rM   constexpr_typer`   rY   int1r1   r=   uint32int64uint64r:   rv   builtinsabsrc   re   rP   ry   )r   rz   min_float32max_float32abs_xs   &&   r   rw   TritonSemantic.to_tensor_type   sy   a""H2,,--Aa77N3"U"xx #!#e#yy  $1$u$xx %!#e#yy  $81=>>5!!!K%C/KLLOEe$|v2{2zz! 3 zz! zz!zz!/!Id1gYjIJJr   c               (    V ^8  d   QhRRRRRRRR/# )r/   r   rH   r   allow_ptr_arY   r2   Noner3   )r4   s   "r   r5   r6      s0     	@ 	@( 	@H 	@SW 	@\` 	@r   c                	    VP                  4       '       d^   V'       g   \        W4      hVP                  4       '       d   W8w  d   \        W4      hVP                  4       '       d   \        W4      hR # R # r*   )is_ptrr   is_floating)r   r   r   r   s   &&&&r   check_ptr_type_impl"TritonSemantic.check_ptr_type_impl   sZ    ==??/??}}F$4/??!!##/?? $ r   c               $    V ^8  d   QhRRRRRR/# )r/   lhsTensorTy | numbers.Numberrhsr2   Tuple[TensorTy, TensorTy]r3   )r4   s   "r   r5   r6      s$     # #0I #Pi #:S#r   c                	   \        V\        P                  4      p\        V\        P                  4      pV'       d   Tp	V P                  V4      pV'       d   Tp
V P                  V4      pVP                  P
                  pVP                  P
                  pV P                  WV4       V P                  WV4       V'       Ed   VP                  4       '       Egz   VP                  4       '       Egc   V P                  WWV4      pV'       d   X	^ 8  d   VP                  4       '       g&   V'       d)   X
^ 8  d"   VP                  4       '       d   \        R4      hVP                  4       '       d   V'       d?   VP                  4       X	u;8:  d   VP                  4       8:  g   M \        RV	 RV 24      hV'       d?   VP                  4       X
u;8:  d   VP                  4       8:  g   M \        RV
 RV 24      hV'       d   V P                  X	VR7      MV P                  W4      pV'       d   V P                  X
VR7      MV P                  W-4      pV P!                  W4      w  rW3# )r   z{Cannot perform a binary operation between an unsigned tensor and a negative scalar. Perform a explicit cast on one of them.zScalar z is out of range for type rp   )rt   numbersNumberr|   ry   scalarr   r   rm   is_int_unsignedr:   rj   get_int_min_valueget_int_max_valuerx   castbroadcast_impl_value)r   r   r   allow_lhs_ptrallow_rhs_ptrarithmetic_checkr[   lhs_is_scalarrhs_is_scalar
lhs_scalar
rhs_scalar
lhs_sca_ty
rhs_sca_ty
ret_sca_tys   &&&&&&&       r   binary_op_type_checking_impl+TritonSemantic.binary_op_type_checking_impl   s    #37"37J..%CJ..%C XX__
XX__
  G  GJ$5$5$7$7
@Q@Q@S@S33JzjtuJ*q.Z5O5O5Q5Q$aJ<V<V<X<X  "K L L  "" **F*F*HJ +I*4*F*F*H+I$wzl:TU_T`%abb **F*F*HJ +I*4*F*F*H+I$wzl:TU_T`%abbHU$&&z&D[_[d[deh[uCHU$&&z&D[_[d[deh[uC ,,S6xr   c               $    V ^8  d   QhRRRRRR/# )r/   r   r   r   	binary_opcallabler3   )r4   s   "r   r5   r6      s"     , ,H ,8 ,X` ,r   c                	F   VP                   P                  P                  ^@8  g'   V P                  P                  P
                  '       g   R# VP                   P                  pVP                   P                  pWE8X  g   Q hVP                  4       '       g   Q hV P                  V\        P                  4      pV P                  V\        P                  4      pV! WR4      pVP                  4       pV P                  V\        P                  4      pVP                  4       pV P                  V\        P                  4      pV P                  V P                  Wg4      V P                  Wh4      4      p	RVP                   RVP                    2p
V P#                  WR4       R# )@   NFr1   z! overflow detected for operation )ry   r   rK   r(   optionssanitize_overflowrj   r   r<   r   r   rx   r   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_impl/TritonSemantic.binary_op_sanitize_overflow_impl   s<   88??''2-T\\5I5I5[5[5[XX__
XX__
'''  """"iiRXX&iiRXX&%(002	((BHH=	002	((BHH=	yy8$:L:LS:\]J++,,MiN`N`Mab4d+r   c               (    V ^8  d   QhRRRRRRRR/# r/   inputr   otherr   rY   r2   r   r3   )r4   s   "r   r5   r6      s,     > >2 >;T >#>(0>r   c                	   V P                  WR R 4      w  rVP                  P                  pVP                  P                  pVP                  4       '       d"   VP                  4       '       d   \	        R4      hVP                  4       '       dE   VP                  4       '       g/   Y!r!VP                  P                  pVP                  P                  pVP                  4       '       d   VP
                  pVP                  P                  4       '       d   VP                  P                  ^@8  dj   VP                  P                  \        P                  4      P                  V P                  4      pV P                  P                  VP
                  VR4      pV P                  V P                  P!                  VP
                  V4      VP                  4      # VP#                  4       '       dK   V P                  V P                  P%                  VP
                  VP
                  4      VP                  4      # VP'                  4       '       do   V'       d   V P)                  WV P*                  4       V P                  V P                  P-                  VP
                  VP
                  4      VP                  4      # \	        RV 24      h)Tzcannot add pointers togetherFr]   )r   ry   r   r   rP   handlerM   r   rK   with_element_tyr<   r   to_irr(   create_int_castr'   create_addptrr   create_faddrj   r   add
create_add)r   r   r   r   input_scalar_tyother_scalar_tyother_handlei64_tys   &&&&    r   r   TritonSemantic.add   s   88tTR**++**++!!##(>(>(@(@:;; !!##O,B,B,D,D 5#jj//O#jj//O!!## <<L{{**,,1I1IB1N33BHH=CCDLLQ#||;;ELL&RWX;;t||99%,,UW\WaWabb((**;;t||77ellSUZU_U_``##%% 55eDHHM;;t||66u||U\\RTYT^T^__*?*;<==r   c               (    V ^8  d   QhRRRRRRRR/# r   r3   )r4   s   "r   r5   r6     s,     8 82 8;T 8#8(08r   c                	   V P                  WR R4      w  rVP                  P                  pVP                  4       '       d#   V P	                  WP                  V4      RR7      # VP                  4       '       dK   V P                  V P                  P                  VP                  VP                  4      VP                  4      # VP                  4       '       do   V'       d   V P                  WV P                  4       V P                  V P                  P                  VP                  VP                  4      VP                  4      # \        RV 24      h)TF)r   r]   )r   ry   r   r   r   minusr   r'   r(   create_fsubr   rj   r   sub
create_subrP   r   r   r   r   rk   s   &&&& r   r   TritonSemantic.sub  s    88tUSJJ%%	88E::e#48NN  "";;t||77ellSUZU_U_`` 55eDHHM;;t||66u||U\\RTYT^T^__*9+677r   c               (    V ^8  d   QhRRRRRRRR/# r   r3   )r4   s   "r   r5   r6     s,     8 82 8;T 8#8(08r   c                	8   V P                  W4      w  rVP                  P                  pVP                  4       '       dK   V P	                  V P
                  P                  VP                  VP                  4      VP                  4      # VP                  4       '       do   V'       d   V P                  WV P                  4       V P	                  V P
                  P                  VP                  VP                  4      VP                  4      # \        R V 24      hr]   )r   ry   r   r   r'   r(   create_fmulr   rj   r   mul
create_mulrP   r   s   &&&& r   r   TritonSemantic.mul  s    88FJJ%%	  "";;t||77ellSUZU_U_`` 55eDHHM;;t||66u||U\\RTYT^T^__*9+677r   c               $    V ^8  d   QhRRRRRR/# r/   r   r   r   r2   r   r3   )r4   s   "r   r5   r6   $  s)     ] ]6 ]?X ]]e ]r   c                	   V P                  WR R RR4      w  rVP                  P                  pVP                  P                  pVP                  4       '       d*   VP	                  4       '       d   V P                  W#4      pEM$VP	                  4       '       d)   VP                  4       '       d   V P                  W4      pMVP	                  4       '       dX   VP	                  4       '       dB   V P                  V\        P                  4      pV P                  V\        P                  4      pMyVP                  4       '       dV   VP                  4       '       d@   VP                  VP                  8  d   V P                  W#4      pM V P                  W4      pM\        RV 24      hV P                  V P                  P                  VP                  VP                  4      VP                  4      # FTr]   )r   ry   r   r   rj   r   r<   rc   fp_mantissa_widthrP   r'   r(   create_fdivr   )r   r   r   r   r   s   &&&  r   truedivTritonSemantic.truediv$  sk   88ueUY[_`**++**++&&((_-C-C-E-EIIe5E##%%/*E*E*G*GIIe5E##%%/*@*@*B*BIIeRZZ0EIIeRZZ0E((**/J/J/L/L00?3T3TT		%9		%9 ..?@AA{{4<<33ELL%,,OQVQ[Q[\\r   c               $    V ^8  d   QhRRRRRR/# r   r3   )r4   s   "r   r5   r6   =  s$     > >7 >@Y >^f >r   c                	   V P                  WR R RR4      w  rVP                  P                  pVP                  P                  pVP                  4       '       d   VP                  4       '       d   V P	                  W44      pV P                  W4      pV P                  W%4      pVP                  4       '       dK   V P                  V P                  P                  VP                  VP                  4      VP                  4      # V P                  V P                  P                  VP                  VP                  4      VP                  4      # \        RV 24      hr   )r   ry   r   rj   rU   r   is_int_signedr'   r(   create_sdivr   create_udivrP   )r   r   r   r   r   ret_tys   &&&   r   floordivTritonSemantic.floordiv=  s
   88ueUY[_`**++**++!!##(>(>(@(@..PFIIe,EIIe,E##%%{{4<<#;#;ELL%,,#WY^YcYcdd{{4<<#;#;ELL%,,#WY^YcYcdd*?*;<==r   c               (    V ^8  d   QhRRRRRRRR/# )r/   r   r   r   ieee_roundingrY   r2   r   r3   )r4   s   "r   r5   r6   K  s,     , ,3 ,<U ,fj ,ow ,r   c                	   VP                   P                  pVP                   P                  pVP                  4       '       d   VP                  4       '       g   \        R 4      hV P	                  WRRRR4      w  rV P
                  P                  VP                  VP                  4      pV P                  WaP                   4      # )z4both operands of fdiv must have floating scalar typeFT)	ry   r   r   rP   r   r(   r   r   r'   )r   r   r   r   r   r   r   s   &&&&   r   fdivTritonSemantic.fdivK  s    **++**++**,,O4O4O4Q4QRSS88ueUZ\`all&&u||U\\B{{3

++r   c               $    V ^8  d   QhRRRRRR/# r   r3   )r4   s   "r   r5   r6   T  s$     8 82 8;T 8Ya 8r   c                	   V P                  WR R RR4      w  rVP                  P                  pVP                  P                  pVP                  4       '       dK   V P	                  V P
                  P                  VP                  VP                  4      VP                  4      # VP                  4       '       Ed	   VP                  VP                  8w  dD   \        RVP                  4       ,           R,           VP                  4       ,           R,           4      hVP                  4       '       dK   V P	                  V P
                  P                  VP                  VP                  4      VP                  4      # V P	                  V P
                  P                  VP                  VP                  4      VP                  4      # \        RV 24      h)FTzCannot mod z by r^   r]   )r   ry   r   r   r'   r(   create_fremr   rj   rL   rP   r   r   create_sremcreate_urem)r   r   r   rk   r   s   &&&  r   modTritonSemantic.modT  se   88ueUY[_`JJ%%	**++  "";;t||77ellSUZU_U_``''?+I+II	0B0B0D Dv MP_PhPhPj j ns !s t t &&(({{4<<#;#;ELL%,,#WY^YcYcdd{{4<<#;#;ELL%,,#WY^YcYcdd*9+677r   c               $    V ^8  d   QhRRRRRR/# r/   rz   r   ypropagate_nantl.PropagateNanr3   )r4   s   "r   r5   r6   k  !     9 9 9h 9 9r   c                	   V P                  W4      w  rVP                  pVP                  4       '       d   V\        P                  P
                  8X  dK   V P                  V P                  P                  VP                  VP                  4      VP                  4      # V\        P                  P                  8X  dK   V P                  V P                  P                  VP                  VP                  4      VP                  4      # \        R V 24      hVP                  4       '       dK   V P                  V P                  P                  VP                  VP                  4      VP                  4      # VP!                  4       '       dK   V P                  V P                  P#                  VP                  VP                  4      VP                  4      # \%        RV 24      hzUnexpected propagate_nan Unexpected dtype )r   rM   r   r<   PropagateNanALLr'   r(   create_minimumfr   ry   NONEcreate_minnumfr:   r   create_minsir   create_minuirP   r   rz   r  r  rM   s   &&&& r   minimumTritonSemantic.minimumk  R   006 3 33{{4<<#?#?!((#SUVU[U[\\"//"6"66{{4<<#>#>qxx#RTUTZTZ[[ #<]O!LMM  "";;t||88188LaffUU""$$;;t||88188LaffUU/w788r   c               $    V ^8  d   QhRRRRRR/# r  r3   )r4   s   "r   r5   r6   |  r  r   c                	   V P                  W4      w  rVP                  pVP                  4       '       d   V\        P                  P
                  8X  dK   V P                  V P                  P                  VP                  VP                  4      VP                  4      # V\        P                  P                  8X  dK   V P                  V P                  P                  VP                  VP                  4      VP                  4      # \        R V 24      hVP                  4       '       dK   V P                  V P                  P                  VP                  VP                  4      VP                  4      # VP!                  4       '       dK   V P                  V P                  P#                  VP                  VP                  4      VP                  4      # \%        RV 24      hr  )r   rM   r   r<   r
  r  r'   r(   create_maximumfr   ry   r  create_maxnumfr:   r   create_maxsir   create_maxuirP   r  s   &&&& r   maximumTritonSemantic.maximum|  r  r   c               (    V ^8  d   QhRRRRRRRR/# )r/   rz   r   minmaxr  r  r3   )r4   s   "r   r5   r6     s.     	a 	ax 	ah 	aX 	ao 	ar   c                	   V P                  W#4      w  r#V P                  W4      w  rV P                  W4      w  rVP                  pVP                  4       '       dW   V P                  V P                  P                  VP                  VP                  VP                  V4      VP                  4      # \        R V R24      h)r	  z(. Only floating point clamp is supported)	r   rM   r   r'   r(   create_clampfr   ry   rP   )r   rz   r  r  r  rM   s   &&&&& r   clampTritonSemantic.clamp  s    44S>221:221:;;t||99!((CJJPSPZPZ\ijlmlrlrss/w6^_``r   c               $    V ^8  d   QhRRRRRR/# )r/   r   r   r   r2   r   r3   )r4   s   "r   r5   r6     s"      8 H Qj r   c                	r   V P                  W4      w  rVP                  P                  pVP                  P                  pVP                  4       '       d   VP                  4       '       g   \	        W44      hV P                  W44      pWS8w  d   V P                  W4      pWT8w  d   V P                  W%4      pW3# r*   )r   ry   r   rj   r   rU   r   )r   r   r   input_sca_tyother_sca_tyr   s   &&&   r   bitwise_op_type_checking_impl,TritonSemantic.bitwise_op_type_checking_impl  s    88Fzz((zz((""$$L,?,?,A,A+LGG..|J
%IIe0E%IIe0E|r   c               $    V ^8  d   QhRRRRRR/# r/   r   r   r   r2   r3   )r4   s   "r   r5   r6     &     \ \( \8 \ \r   c                	    V P                  W4      w  rV P                  V P                  P                  VP                  VP                  4      VP
                  4      # r*   )r(  r'   r(   
create_andr   ry   r   r   r   s   &&&r   r   TritonSemantic.and_  E    99%G{{4<<225<<NPUPZPZ[[r   c               $    V ^8  d   QhRRRRRR/# r+  r3   )r4   s   "r   r5   r6     s&     [ [ [( [x [r   c                	    V P                  W4      w  rV P                  V P                  P                  VP                  VP                  4      VP
                  4      # r*   )r(  r'   r(   	create_orr   ry   r/  s   &&&r   or_TritonSemantic.or_  sB    99%G{{4<<11%,,MuzzZZr   c               $    V ^8  d   QhRRRRRR/# r+  r3   )r4   s   "r   r5   r6     r,  r   c                	    V P                  W4      w  rV P                  V P                  P                  VP                  VP                  4      VP
                  4      # r*   )r(  r'   r(   
create_xorr   ry   r/  s   &&&r   xor_TritonSemantic.xor_  r1  r   c               $    V ^8  d   QhRRRRRR/# r+  r3   )r4   s   "r   r5   r6     s!     ' ' '( 'x 'r   c                	$   VP                   P                  4       '       g!   V P                  V\        P                  4      pVP                   P                  4       '       g!   V P                  V\        P                  4      pV P                  W4      # r*   )ry   is_int1bitcastr<   r   r   r/  s   &&&r   logical_andTritonSemantic.logical_and  s_    zz!!##LL0Ezz!!##LL0Eyy&&r   c               $    V ^8  d   QhRRRRRR/# r+  r3   )r4   s   "r   r5   r6     s!     & & & &h &r   c                	$   VP                   P                  4       '       g!   V P                  V\        P                  4      pVP                   P                  4       '       g!   V P                  V\        P                  4      pV P                  W4      # r*   )ry   r>  r?  r<   r   r5  r/  s   &&&r   
logical_orTritonSemantic.logical_or  s_    zz!!##LL0Ezz!!##LL0Exx%%r   c                   V ^8  d   QhRR/# )r/   r   r   r3   )r4   s   "r   r5   r6     s     " "( "r   c                	    VP                   P                  4       '       g!   V P                  V\        P                  4      pV P                  V4      # r*   )ry   r>  r?  r<   r   invertr   r   s   &&r   not_TritonSemantic.not_  s8    zz!!##LL0E{{5!!r   c               $    V ^8  d   QhRRRRRR/# r+  r3   )r4   s   "r   r5   r6     &     ] ]( ]8 ] ]r   c                	    V P                  W4      w  rV P                  V P                  P                  VP                  VP                  4      VP
                  4      # r*   )r(  r'   r(   create_lshrr   ry   r/  s   &&&r   lshrTritonSemantic.lshr  E    99%G{{4<<33ELL%,,OQVQ[Q[\\r   c               $    V ^8  d   QhRRRRRR/# r+  r3   )r4   s   "r   r5   r6     rM  r   c                	    V P                  W4      w  rV P                  V P                  P                  VP                  VP                  4      VP
                  4      # r*   )r(  r'   r(   create_ashrr   ry   r/  s   &&&r   ashrTritonSemantic.ashr  rR  r   c               $    V ^8  d   QhRRRRRR/# r+  r3   )r4   s   "r   r5   r6     s&     \ \ \( \x \r   c                	    V P                  W4      w  rV P                  V P                  P                  VP                  VP                  4      VP
                  4      # r*   )r(  r'   r(   
create_shlr   ry   r/  s   &&&r   shlTritonSemantic.shl  r1  r   c                    V ^8  d   QhRRRR/# r/   r   r   r2   r3   )r4   s   "r   r5   r6     s      ( x r   c                	    V# r*   r3   rI  s   &&r   plusTritonSemantic.plus  s    r   c                    V ^8  d   QhRRRR/# r^  r3   )r4   s   "r   r5   r6     s     ) )8 ) )r   c                	T   VP                   P                  pVP                  4       '       d(   \        R VP	                  4       ,           R,           4      hV P                  V P                  P                  VP                  V P                  4      4      V4      pV P                  W1R4      # )z$wrong type argument to unary minus ()T)
ry   r   r   r:   r   r'   r(   get_null_valuer   r   )r   r   r&  _0s   &&  r   r   TritonSemantic.minus  s    zz((  ClF[F[F]]`ccdd[[44\5G5G5UVXdexx4((r   c                    V ^8  d   QhRRRR/# r^  r3   )r4   s   "r   r5   r6     s     $ $H $ $r   c                	~   VP                   P                  pVP                  4       '       g   VP                  4       '       d(   \	        R VP                  4       ,           R,           4      hV P                  V P                  P                  VP                  V P                  4      4      V4      pV P                  W4      # )z%wrong type argument to unary invert (rd  )ry   r   r   r   r:   r   r'   r(   get_all_ones_valuer   r:  )r   r   r&  _1s   &&  r   rH  TritonSemantic.invert  s    zz((  L$<$<$>$>D|G\G\G^^addee[[889K9KDLL9YZ\hiyy##r   c                    V ^8  d   QhRRRR/# )r/   vr   r2   tl.block_typer3   )r4   s   "r   r5   r6     s     / /H / /r   c                	T    VP                   P                  \        P                  4      # r*   )ry   r   r<   r   )r   rn  s   &&r   
_bool_likeTritonSemantic._bool_like  s    vv%%bgg..r   c               $    V ^8  d   QhRRRRRR/# r+  r3   )r4   s   "r   r5   r6     s!     8 8( 88 8 8r   c                	   V P                  W4      w  rVP                  P                  pVP                  4       '       dP   V P	                  V P
                  P                  VP                  VP                  4      V P                  V4      4      # VP                  4       '       d   VP                  4       '       dP   V P	                  V P
                  P                  VP                  VP                  4      V P                  V4      4      # V P	                  V P
                  P                  VP                  VP                  4      V P                  V4      4      # \        R V 24      hr   )r   ry   r   r   r'   r(   create_fcmpOGTr   rq  rj   r   create_icmpSGTcreate_icmpUGTrP   r   r   r   rk   s   &&& r   greater_thanTritonSemantic.greater_than     88FJJ%%	  "";;t||::5<<VX\XgXghmXnoo&&(({{4<<#>#>u||U\\#Z\`\k\klq\rss{{4<<#>#>u||U\\#Z\`\k\klq\rss*9+677r   c               $    V ^8  d   QhRRRRRR/# r+  r3   )r4   s   "r   r5   r6     s!     8 88 8H 8 8r   c                	   V P                  W4      w  rVP                  P                  pVP                  4       '       dP   V P	                  V P
                  P                  VP                  VP                  4      V P                  V4      4      # VP                  4       '       d   VP                  4       '       dP   V P	                  V P
                  P                  VP                  VP                  4      V P                  V4      4      # V P	                  V P
                  P                  VP                  VP                  4      V P                  V4      4      # \        R V 24      hr   )r   ry   r   r   r'   r(   create_fcmpOGEr   rq  rj   r   create_icmpSGEcreate_icmpUGErP   rx  s   &&& r   r   TritonSemantic.greater_equal  r{  r   c               $    V ^8  d   QhRRRRRR/# r+  r3   )r4   s   "r   r5   r6     s!     8 8x 8 8X 8r   c                	   V P                  W4      w  rVP                  P                  pVP                  4       '       dP   V P	                  V P
                  P                  VP                  VP                  4      V P                  V4      4      # VP                  4       '       d   VP                  4       '       dP   V P	                  V P
                  P                  VP                  VP                  4      V P                  V4      4      # V P	                  V P
                  P                  VP                  VP                  4      V P                  V4      4      # \        R V 24      hr   )r   ry   r   r   r'   r(   create_fcmpOLTr   rq  rj   r   create_icmpSLTcreate_icmpULTrP   rx  s   &&& r   	less_thanTritonSemantic.less_than  r{  r   c               $    V ^8  d   QhRRRRRR/# r+  r3   )r4   s   "r   r5   r6     s!     8 8 8 8h 8r   c                	   V P                  W4      w  rVP                  P                  pVP                  4       '       dP   V P	                  V P
                  P                  VP                  VP                  4      V P                  V4      4      # VP                  4       '       d   VP                  4       '       dP   V P	                  V P
                  P                  VP                  VP                  4      V P                  V4      4      # V P	                  V P
                  P                  VP                  VP                  4      V P                  V4      4      # \        R V 24      hr   )r   ry   r   r   r'   r(   create_fcmpOLEr   rq  rj   r   create_icmpSLEcreate_icmpULErP   rx  s   &&& r   r   TritonSemantic.less_equal  r{  r   c               $    V ^8  d   QhRRRRRR/# r+  r3   )r4   s   "r   r5   r6   (  s!     	8 	88 	8H 	8 	8r   c                	   V P                  W4      w  rVP                  P                  pVP                  4       '       dP   V P	                  V P
                  P                  VP                  VP                  4      V P                  V4      4      # VP                  4       '       dP   V P	                  V P
                  P                  VP                  VP                  4      V P                  V4      4      # \        R V 24      hr   )r   ry   r   r   r'   r(   create_fcmpOEQr   rq  rj   create_icmpEQrP   rx  s   &&& r   equalTritonSemantic.equal(      88FJJ%%	  "";;t||::5<<VX\XgXghmXnoo;;t||99%,,UW[WfWfglWmnn*9+677r   c               $    V ^8  d   QhRRRRRR/# r+  r3   )r4   s   "r   r5   r6   3  s!     	8 	8x 	8 	8X 	8r   c                	   V P                  W4      w  rVP                  P                  pVP                  4       '       dP   V P	                  V P
                  P                  VP                  VP                  4      V P                  V4      4      # VP                  4       '       dP   V P	                  V P
                  P                  VP                  VP                  4      V P                  V4      4      # \        R V 24      hr   )r   ry   r   r   r'   r(   create_fcmpUNEr   rq  rj   create_icmpNErP   rx  s   &&& r   	not_equalTritonSemantic.not_equal3  r  r   r   Nc               (    V ^8  d   QhRRRRRRRR/# )r/   startr1   endr   ro  r2   r   r3   )r4   s   "r   r5   r6   B  s.     Z ZC Zc Zm Zx Zr   c               	D   \        V\        4      '       d   \        V\        4      '       g   \        R 4      h\        V^ ,	          4      p\        V^ ,	          4      pV'       g	   V'       d   \        R4      hW!8:  d   \        R4      hW!,
          pWf^,
          ,          ^ 8w  d   \        R4      hV.pVf&   \        P
                  ! \        P                  V4      pVP                  V P                  4      pV P                  V P                  P                  WV4      V4      # )z/arange's arguments must be of type tl.constexprzarange must fit in int32z=arange's end argument must be greater than the start argumentz#arange's range must be a power of 2)rt   r1   r:   rY   r<   
block_typer=   r   r(   r'   create_make_range)	r   r  r  r   is_start_int64is_end_int64rangeshape	ret_ty_irs	   &&&$     r   arangeTritonSemantic.arangeB  s    %%%ZS-A-ANOOerk*C2I\788<\]]QYA%BCC>]]288U3FLL.	{{4<<99)CPRXYYr   c                    V ^8  d   QhRRRR/# r/   rM   rH   r2   r   r3   )r4   s   "r   r5   r6   T  s     ) )H ) )r   c                	   Vf   \        R4      hV^ 8X  d6   V P                  P                  VP                  V P                  4      4      pMVP	                  4       '       dQ   V P                  P                  V4      pV P                  P                  WP                  V P                  4      4      pM+\        V P                  RVP                   24      pV! V4      pV P                  W4      # )Nz2dtype must be specified when value is not a tensorget_)
r:   r(   re  r   ri   get_fp32create_fp_truncgetattrnamer'   )r   r`   rM   get_value_fns   &&& r   rx   TritonSemantic.scalar_constantT  s    =QRRA:LL//DLL0IJE\\^^LL))%0ELL00DLL8QRE"4<<4

|1DEL 'E{{5((r   c                    V ^8  d   QhRRRR/# r  r3   )r4   s   "r   r5   r6   b  s     2 2 2X 2r   c                	    \        V\        P                  4      '       d4   VP                  P                  ^8X  g   Q R4       hV P                  W4      # V P                  W4      # )r9   zonly accepts size-1 tensor)rt   r<   r'   numelr`   r   rx   )r   r`   rM   s   &&&r   make_scalarTritonSemantic.make_scalarb  sQ    eRYY'';;$$)G+GG)99U**##E11r   c               $    V ^8  d   QhRRRRRR/# )r/   r  	List[int]rM   rH   r2   r   r3   )r4   s   "r   r5   r6   i  s&     A A) A8 A Ar   c                	D    V P                  V P                  W#4      V4      # r*   )splatr  )r   r  r`   rM   s   &&&&r   fullTritonSemantic.fulli  s    zz$**58%@@r   c               $    V ^8  d   QhRRRRRR/# )r/   r`   r   r  r  r2   r3   )r4   s   "r   r5   r6   p  s&     h h8 hI h( hr   c                	T   VP                   P                  4       '       d   Q R 4       h\        V4      ^ 8X  d   V# \        P                  ! VP
                  V4      pV P                  V P                  P                  VP                  V P                  4      VP                  4      V4      # )zCannot splat a block tensor)ry   is_blocklenr<   r  rM   r'   r(   create_splatr   r   )r   r`   r  r   s   &&& r   r  TritonSemantic.splatp  s|    ::&&((G*GG(u:?Lu{{E2{{4<<44V\\$,,5OQVQ]Q]^`fggr   c                    V ^8  d   QhRRRR/# r/   r`   r   r2   r3   )r4   s   "r   r5   r6   w  s     S SX S( Sr   c                	    V P                  V P                  P                  VP                  4      VP                  4      # r*   )r'   r(   create_unsplatr   rM   )r   r`   s   &&r   unsplatTritonSemantic.unsplatw  s*    {{4<<66u||DekkRRr   c               (    V ^8  d   QhRRRRRRRR/# )r/   r   r   	dst_shaper  can_reorderrY   r2   r3   )r4   s   "r   r5   r6   z  s/     f fX f) f$ fS[ fr   c                	6   ^pV F  pWE,          pK  	  VP                   P                  V8w  d   \        R4      h\        P                  ! VP                   P
                  V4      pV P                  V P                  P                  VP                  W#4      V4      # )r9   z:reshape() cannot change total number of elements in tensor)
ry   r  r:   r<   r  r   r'   r(   create_reshaper   )r   r   r  r  r  sr   s   &&&&   r   reshapeTritonSemantic.reshapez  sw    AJE ::u$YZZuzz00)<{{4<<66u||Y\^deer   c               $    V ^8  d   QhRRRRRR/# )r/   r   r   r0   r1   r2   r3   )r4   s   "r   r5   r6     s&     X X X X Xr   c                	   VP                    Uu. uF  p\        P                  ! V4      NK  	  ppVP                  V^4       VP                  P                  4       '       g   V P                  WR7      # \        P                  ! VP                  P                  V4      pV P                  V P                  P                  VP                  V4      V4      # u upi )r9   )r  )r  r<   _unwrap_if_constexprinsertry   r  r  r  r   r'   r(   create_expand_dimsr   )r   r   r0   rz   r  r   s   &&&   r   expand_dimsTritonSemantic.expand_dims  s    9>EAR,,Q/	Eq!zz""$$::e:55uzz00)<{{4<<::5<<NPVWW Fs   Cc               (    V ^8  d   QhRRRRRRRR/# )r/   r   r   r   r  rY   r2   r3   )r4   s   "r   r5   r6     s.     V Vx Vh VT Vh Vr   c                	   V'       g   Q R 4       h\        VP                  4      ^8X  g   Q h\        P                  ! VP                  P
                  VP                  ^ ,          VP                  ^ ,          ,           .4      pV P                  V P                  P                  VP                  VP                  4      V4      # )z;current implementation of `cat` always may reorder elements)
r  r  r<   r  ry   r   r'   r(   
create_catr   )r   r   r   r  ret_types   &&&& r   catTritonSemantic.cat  s    YYY{399~"""==399Q<#))A,3N2OP{{4<<223::szzJHUUr   c               $    V ^8  d   QhRRRRRR/# )r/   ar   br2   r3   )r4   s   "r   r5   r6     s!      h 8  r   c                	n   V P                  W4      w  rVP                  . 8H  pV'       d%   V P                  V^ 4      pV P                  V^ 4      p\        VP                  R,          \        P
                  4      '       d   \        P
                  ! ^4      pM^pVP                  V.,           p\        P                  ! VP                  P                  V4      pV P                  V P                  P                  VP                  VP                  4      V4      pV'       d   V P                  V^.RR7      pV# )r   Fr  )r   r  r  rt   r<   ru   r  ry   r   r'   r(   create_joinr   r  )r   r  r  
was_rank_1two	new_shaper  r   s   &&&     r   joinTritonSemantic.join  s    ((. WW]
  A&A  A&Aaggbk2<<00,,q/CCGGseO	==	:kk$,,22188QXXFQ,,sQCU,;C
r   c                    V ^8  d   QhRRRR/# )r/   r  r   r2   r   r3   )r4   s   "r   r5   r6     s     

 

x 

$= 

r   c                	   \        VP                  4      ^ 8  g   Q h\        P                  ! VP                  R,          4      ^8X  g   Q hVP                  RR p\        P                  ! VP
                  P                  V4      pV P                  P                  VP                  4      w  rEV P                  WC4      V P                  WS4      3# )r   Nr  )r  r  r<   r  r  ry   r   r(   create_splitr   r'   )r   r  r  r  outLHSoutRHSs   &&    r   splitTritonSemantic.split  s    AGGq ! ''49:9GGCRL	==	:22188<KK)KK)
 	
r   c               $    V ^8  d   QhRRRRRR/# )r/   r   r   dims
Tuple[int]r2   r3   )r4   s   "r   r5   r6     s&     T TX TZ TH Tr   c                	.   \        VP                  4      \        V4      8w  d.   \        R \        VP                  4       R\        V4       24      h\        R V 4       4      \	        \        \        V4      4      4      8w  d   \        RV 24      h\        P                  ! VP                  P                  V Uu. uF  q1P                  V,          NK  	  up4      pV P                  V P                  P                  VP                  V4      V4      # u upi )z;permute dims must have the same length as input shape, got r   c              3  N   "   T F  p\         P                  ! V4      x  K  	  R # 5ir*   )r<   r  ).0ds   & r   	<genexpr>)TritonSemantic.permute.<locals>.<genexpr>  s     ;d"))!,,ds   #%z?permute dims must be a permutation of 0, 1, ..., n-1, but were )r  r  r:   sortedlistr  r<   r  ry   r   r'   r(   create_transr   )r   r   r  r  r  s   &&&  r   permuteTritonSemantic.permute  s    u{{s4y(McRWR]R]N^M__dehimendopr r;d;;tE#d)DT?UU^_c^deff==!2!2T4RT[[^^T4RS{{4<<44U\\4H(SS 5Ss   :D
c               $    V ^8  d   QhRRRRRR/# )r/   r   r   r  r  r2   r3   )r4   s   "r   r5   r6     s&     W W( W: W( Wr   c                	R   VP                   P                  4       '       g   V P                  W4      # VP                   P                  4       p\	        V4      \	        V4      8w  d   \        R V RV 24      hW#8X  d   V# \        V4       F;  w  rEW$,          V8w  g   K  V^8w  g   K  \        RW$,           RV RV RV RV 2
4      h	  \        P                  ! VP                   P                  V4      pV P                  V P                  P                  VP                  V4      V4      # )z!Cannot broadcast, rank mismatch: , z3Cannot broadcast, the expanded size of the tensor (z ) must match the existing size (z) at non-singleton dimension : )ry   r  r  get_block_shapesr  r:   	enumerater<   r  r   r'   r(   create_broadcastr   )r   r   r  	src_shapeiitemr   s   &&&    r   broadcast_impl_shape#TritonSemantic.broadcast_impl_shape  s   zz""$$::e++JJ//1	y>SZ'@2eWUVVL +GAx4DAI #VW\W_V` aCCG& I%%&Cr)Bug"? @ @ ,
 uzz00%8{{4<<88uMvVVr   c               $    V ^8  d   QhRRRRRR/# )r/   r   r   r   r2   r3   )r4   s   "r   r5   r6     s!     2 2 2x 2H 2r   c           	     	&	   VP                   pVP                   pVP                  4       '       d   VP                  4       '       gn   VP                  VP                  4      pV P	                  V P
                  P                  VP                  V P
                  4      VP                  4      V4      pW3# VP                  4       '       g   VP                  4       '       dn   VP                  VP                  4      pV P	                  V P
                  P                  VP                  V P
                  4      VP                  4      V4      pW3# VP                  4       '       Ed/   VP                  4       '       Ed   VP                  4       pVP                  4       p\        V4      \        V4      8  d   \        \        V4      \        V4      4       F  pV P	                  V P
                  P                  VP                  ^ 4      \        P                  ! VP                  ^.VP                  ,           4      4      pVP                   pVP                  4       pK  	  M\        V4      \        V4      8  d   \        \        V4      \        V4      4       F  pV P	                  V P
                  P                  VP                  ^ 4      \        P                  ! VP                  ^.VP                  ,           4      4      pVP                   pVP                  4       pK  	  \        V4      \        V4      8X  g   Q h. p\!        V4       F  w  rWi,          pV
^8X  d   VP#                  V4       K'  V^8X  g   W8X  d   VP#                  V
4       KG  \%        R\'        V	4      ,           R,           \'        V
4      ,           R,           \'        V4      ,           4      h	  WX8w  dX   \        P                  ! VP                  V4      pV P	                  V P
                  P)                  VP                  V4      V4      pWh8w  dX   \        P                  ! VP                  V4      pV P	                  V P
                  P)                  VP                  V4      V4      pW3# )r   z?Cannot make_shape_compatible: incompatible dimensions at index r  r   )ry   r  r   r   r'   r(   r  r   r   r  r  r  r  r<   r  valuesr  appendr:   strr  )r   r   r   lhs_tyrhs_ty	lhs_shape	rhs_shape_	ret_shaper  leftrightr   s   &&&          r   r   #TritonSemantic.broadcast_impl_value  ss    ??V__%6%6++FMM:F++dll77T\\8RTWT^T^_aghCV xS ""v'8'8++FMM:F++dll77T\\8RTWT^T^_aghCN xK __6??#4#4//1I//1I9~I.s9~s9~>A++dll&E&EcjjRS&T&(mmFMMA3IYIYCY&Z\C XXF & 7 7 9I	 ?
 Y#i.0s9~s9~>A++dll&E&EcjjRS&T&(mmFMMA3IYIYCY&Z\C XXF & 7 7 9I	 ?
 y>S^333I$Y/!19$$U+qjem$$T*$ &136q6&:<@&ACFt9&MOV&WY\]bYc&d e e 0 %v}}i@kk$,,"?"?

I"VX^_%v}}i@kk$,,"?"?

I"VX^_xr   c                   V ^8  d   QhRR/# )r/   rounding_modeOptional[str]r3   )r4   s   "r   r5   r6   	  s     t t= tr   c                	    Vf   R # VR8X  d   \         P                  P                  # VR8X  d   \         P                  P                  # \	        RV R24      h)NrtnertzzInvalid rounding mode: z0. Supported rounding modes are 'rtne' and 'rtz'.)r   ROUNDING_MODERTNERTZr:   )r   r  s   &&r   _str_to_rounding_mode$TritonSemantic._str_to_rounding_mode	  sU     F"##(((E!##'''2=/Aqrssr   c               $    V ^8  d   QhRRRRRR/# )r/   r   r   dst_tyrH   r2   r3   )r4   s   "r   r5   r6     s&     j jX jx jH jr   c                	t   VP                   pVP                  4       '       d   VP                  VP                  4      pW28X  d   V# VP                  pVP                  pVP	                  4       '       g   VP	                  4       '       d   V P                  W4      # VP                  pVP                  pWg8w  d3   \        R \        V4      ,           R,           \        V4      ,           4      hV P                  V P                  P                  VP                  VP                  V P                  4      4      V4      # )z!Cannot bitcast data-type of size z to data-type of size )ry   r  r   r   r   r   primitive_bitwidthr:   r  r'   r(   create_bitcastr   r   )r   r   r&  src_ty
src_sca_ty
dst_sca_tysrc_bitsdst_bitss   &&&     r   r?  TritonSemantic.bitcast  s    ??++FMM:FL]]
]]
*"3"3"5"599U++0000@3x=P T2 247MB C C{{4<<66u||V\\RVR^R^E_`bhiir   c               (    V ^8  d   QhRRRRRRRR/# )r/   r   r   r&  rH   fp_downcast_roundingr  r2   r3   )r4   s   "r   r5   r6   $  s/     i9 i9( i9H i9M i9em i9r   c                	   VP                   pVP                  pVP                  pWV8X  d   V# VP                  4       '       d   VP                  V4      pV P	                  V4      pR pVP                  4       '       ds   VP                  4       '       d]   VP                  VP                  8  dB   Vf   \        P                  P                  pMXV\        P                  P                  8w  d   RpM6Ve3   \        R\        V4      ,           R,           \        V4      ,           4      hVP                  4       '       g   VP                  4       '       dU   V P                  P                  P                  R4       f   Q R4       hV P                  P                  R,          ! WW0R7      # VP!                  4       '       d   VP                  4       '       g5   VP                  4       '       d   VP!                  4       '       g	   V'       dQ   V P#                  V P                  P%                  VP&                  VP)                  V P                  4      V4      V4      # VP+                  4       '       d   VP-                  4       '       d-   VP/                  4       '       dG   VP-                  4       '       g1   V P1                  V P1                  V\2        P4                  4      V4      # VP                  4       ;'       d2    VP                  4       ;'       d    VP                  VP                  8  pV'       dP   V P#                  V P                  P7                  VP&                  VP)                  V P                  4      4      V4      # VP                  4       ;'       d2    VP                  4       ;'       d    VP                  VP                  8  p	V	'       dP   V P#                  V P                  P9                  VP&                  VP)                  V P                  4      4      V4      # VP;                  4       '       EdL   VP;                  4       '       Ed5   VP<                  VP<                  8w  g   VP>                  VP>                  8w  d   VPA                  4       ;'       d    VPC                  4       '       * p
VPC                  4       '       dl   VPD                  P)                  V P                  4      pV P#                  V P                  PG                  V4      VPD                  4      pV PI                  W4      # V P#                  V P                  PK                  VP&                  VP)                  V P                  4      V
4      V4      # VPM                  4       '       EdM   VP;                  4       '       Ed6   VPC                  4       '       dl   VPD                  P)                  V P                  4      pV P#                  V P                  PG                  V4      VPD                  4      pV PI                  W4      # VPA                  4       '       dP   V P#                  V P                  PO                  VP&                  VP)                  V P                  4      4      V4      # V P#                  V P                  PQ                  VP&                  VP)                  V P                  4      4      V4      # VP;                  4       '       d   VPM                  4       '       d   VPC                  4       '       g   VPA                  4       '       gP   V P#                  V P                  PS                  VP&                  VP)                  V P                  4      4      V4      # V P#                  V P                  PU                  VP&                  VP)                  V P                  4      4      V4      # VPW                  4       '       d   VP;                  4       '       d   VP<                  pV^@8X  dP   V P#                  V P                  PY                  VP&                  VP)                  V P                  4      4      V4      # V^8X  dh   V PI                  V P1                  V\2        PZ                  4      V P#                  V P                  P]                  ^ 4      \2        PZ                  4      4      # VP;                  4       '       df   VPW                  4       '       dP   V P#                  V P                  P_                  VP&                  VP)                  V P                  4      4      V4      # VPW                  4       '       df   VPW                  4       '       dP   V P#                  V P                  Pa                  VP&                  VP)                  V P                  4      4      V4      # Q RT RT 24       h)	FTz]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.)	_semanticzcannot cast z to )1ry   r   r  r   r#  r   r(  r   r   r!  r:   r  is_fp8e4b15r(   codegen_fnsgetri   r'   create_fp_to_fpr   r   rg   rf   rh   r   r<   rc   r  create_fp_extrj   rK   rL   r   is_boolrM   re  r  r   is_standard_floatingcreate_fp_to_sicreate_fp_to_uicreate_ui_to_fpcreate_si_to_fpr   create_ptr_to_intr   	get_int64create_int_to_ptrr)  )r   r   r&  r1  r*  r+  r,  use_custom_roundingtruncate_fpext_fpsign_extendtyrf  bitwidths   &&&&          r   r   TritonSemantic.cast$  sq   ]]
]]
#L??++J7F  $99:NO#!!##
(>(> )
 )
++j.K.KK#+BDTDTDYDY-A%)9)9)>)>>VZ@S#/  ":<?
O"LNi"j!$Z"1 2 2 ""$$
(>(>(@(@<<++//&(/34 h5gh 4<<++,BCESgxx J$:$:$<$<""$$):):)<)<;;,,U\\6<<;UWklntv v   ););)=)=  ););)=)=99TYYubjj9:FF
 !,,. J J""$J J))J,I,II 	 ;;t||;;ELL&,,W[WcWcJdegmnn '') J J""$J J))J,I,II 	 ;;t||99%,,UYUaUaHbcekll :#4#4#6#6##z'>'>>*B[B[_i_x_xBx$224QQZ=O=O=Q9QK!!##[[&&t||4[[!<!<R!@%++N~~e00{{4<<#?#?fll[_[g[gNhju#v#)+ + **,,1B1B1D1D!!##[[&&t||4[[!<!<R!@%++N~~e00))++{{4<<#?#?fll[_[g[gNh#ikqrr{{4<<#?#?fll[_[g[gNh#ikqrr :#B#B#D#D!!##:+C+C+E+E{{4<<#?#?fll[_[g[gNh#ikqrr{{4<<#?#?fll[_[g[gNh#ikqrr :#4#4#6#6!..H2~{{4<<#A#A%,,PVP\P\]a]i]iPj#kmstt1}~~diirxx&@$++dllNdNdefNgikiqiqBrss :#4#4#6#6;;t||==ellFLLY]YeYeLfgiopp :#4#4#6#6;;t||::5<<VZVbVbIcdflmm8UG4x88ur   c                	:   \         P                  P                  pV'       dy   VR 8X  d   \         P                  P                  pV# VR8X  d   \         P                  P                  pV# VR8X  d   \         P                  P
                  pV# \        RV R24      hV# )z.ca.cgz.cvCache modifier  not supported)r   CACHE_MODIFIERr  CACGCVr:   r   cache_modifiercaches   && r   _str_to_load_cache_modifier*TritonSemantic._str_to_load_cache_modifier  s    !!&&&)),,   5()),,
 	  5()),,  !?>2B.!QRRr   c                	   \         P                  P                  pV'       d   VR 8X  d   \         P                  P                  pV# VR8X  d   \         P                  P                  pV# VR8X  d   \         P                  P
                  pV# VR8X  d   \         P                  P                  pV# \        RV R24      hV# )z.wbrK  z.csz.wtrL  rM  )r   rN  r  WBrP  CSWTr:   rR  s   && r   _str_to_store_cache_modifier+TritonSemantic._str_to_store_cache_modifier  s    !!&&&)),,   5()),,   5()),,
 	  5()),,  !?>2B.!QRRr   c                	    \         P                  P                  pV'       dV   VR 8X  d   \         P                  P                  pV# VR8X  d   \         P                  P                  pV# \        RV R24      hV# )
evict_lastevict_firstzEviction policy rM  )r   EVICTION_POLICYNORMAL
EVICT_LASTEVICT_FIRSTr:   )r   eviction_policyevictions   && r   _str_to_eviction_policy&TritonSemantic._str_to_eviction_policy  su    %%,,,.--88
 	 !M1--99  !#3O3DN!STTr   c                	    R pV'       dV   VR8X  d   \         P                  P                  pV# VR8X  d   \         P                  P                  pV# \	        RV R24      hV# )NzeronanzPadding option rM  )r   PADDING_OPTIONPAD_ZEROPAD_NANr:   )r   padding_optionpaddings   && r   _str_to_padding_option%TritonSemantic._str_to_padding_option  sh    '++44
 	  5(++33  !?>2B.!QRRr   c                	   \         P                  P                  pV'       d   VR 8X  d   \         P                  P                  pV# VR8X  d   \         P                  P                  pV# VR8X  d   \         P                  P                  pV# VR8X  d   \         P                  P
                  pV# \        RV R24      hV# )acquirereleaseacq_relrelaxedMemory semantic rM  )r   MEM_SEMANTICACQUIRE_RELEASEACQUIRERELEASERELAXEDr:   )r   
sem_optionsems   && r   _str_to_semTritonSemantic._str_to_sem  s    oo--Y&oo-- 
 y(oo-- 
 y(oo55
 
	 y(oo-- 
 !#3J<~!NOO
r   c                	:   \         P                  P                  pV'       dy   VR 8X  d   \         P                  P                  pV# VR8X  d   \         P                  P                  pV# VR8X  d   \         P                  P                  pV# \        RV R24      hV# )gpuctasysrw  rM  )r   MEM_SYNC_SCOPEGPUCTASYSTEMr:   )r   scope_optionscopes   && r   _str_to_scopeTritonSemantic._str_to_scope  s    !!%%u$))--  &))--
 	 &))00  !#3L>!PQQr   c                	   V'       d   \        VR 4      '       g   V.pV Uu. uF1  p\        V\        P                  4      '       d   VP                  MTNK3  	  ppV F6  p\        V\
        4      '       d   ^ Tu;8:  d   \        V4      8  d   K3  Q h Q h	  \        V4      ^ 8  g   Q h\        V4      \        \        V4      4      8X  g   Q R4       h\        V4      # R# u upi )__iter__z'Duplicate dimension in `boundary_check`r3   )	hasattrrt   r<   ru   r`   r1   r  setr  )r   boundary_checkblock_shapeelemdims   &&&  r   _canonicalize_boundary_check+TritonSemantic._canonicalize_boundary_check  s    >:66"0!1aopaoY]JtR\\,J,JdjjPTTaoNp%!#s++S0K3{;K0KKK0KKK &~&***~&#c..A*BBmDmmB.))	 qs   7Cc	           
     	   Vf   Ve   \        R4      hVP                  P                  P                  p	V	\        P                  8w  g   Q R4       hV	P                  4       '       d+   V\        P                  P                  8X  d   \        R4      hVP                  P                  p
V P                  WJP                  4       4      pV P                  V P                  P                  VP                  WEWgV4      V
4      # )NK`mask` and `other` arguments cannot be specified for loading block pointers4`tl.int1` should be rewritten in `tl.make_block_ptr`z@Padding option `nan` is not supported for integer block pointers)r:   ry   
element_tyr<   r   rj   r   rk  rm  r  r  r'   r(   create_tensor_pointer_loadr   )r   ptrmaskr   r  ro  rT  re  is_volatileelt_tyr&  s   &&&&&&&&&  r   _load_block_pointer"TritonSemantic._load_block_pointer  s     u0jkk$$// X"XX ==??w"*;*;*C*CC_`` $$ ::>KbKbKde {{LL33CJJY^juv 	r   c	           
     	   VP                   P                  P                  4       '       g(   \        R VP                   P	                  4        R24      hVf   Ve   \        R4      hV'       g	   V'       d   \        R4      hVP                   P                  4       '       gg   V'       d,   VP                   P                  4       '       d   \        R4      hV'       d,   VP                   P                  4       '       d   \        R4      hVP                   P                  4       '       d/   Ve   V P                  W4      w  rVe   V P                  W4      w  rVP                   P                  p	V	P                  p
V
\        P                  8H  pV'       dB   \        P                  p
\        P                  ! WP                  4      p	V P                  W4      pVe   V P                  W:4      pVP                   P                  4       '       d   VP                   P                  V
4      pMT
pVf9   V P                  V P                   P#                  VP$                  WgV4      V4      pMWT P                  V P                   P'                  VP$                  VP$                  V'       d   VP$                  MRVWx4      V4      pV'       d!   V P                  V\        P                  4      pV# )Unsupported ptr type z in `tl.load`Nz)`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)ry   r   r   r:   r   r  r   r  r<   r   int8pointer_typeaddress_spacer   r   r'   r(   create_loadr   create_masked_load)r   r  r  r   r  ro  rT  re  r  ptr_tyr  r:  r&  r   s   &&&&&&&&&     r   _load_legacyTritonSemantic._load_legacy  s9   xx%%''4SXX5F5F5H4IWXX <E-HIIn X Y Y
 xx  ""		**,, !hii,,.. !ijj 88 55c@	 !66sB
 "" BGG#WWF__V-A-ABF))C(C IIe,E 88XX--f5F F <++dll66szz5T_`bhiC++//

DKKY^dhjo08GHNPC ))C)C
r   c               <    V ^8  d   QhRRRRRRRRRR	R
R	RR	RRRR/	# )r/   r  r   r  Optional[TensorTy]r   r  r   rn  r  rS  rd  r  rY   r2   r3   )r4   s   "r   r5   r6   ?  s^     n n n(: nCU ngl n n25nHKnZ^ncknr   c	           
     	L   V P                  V4      p	V P                  V4      p
V P                  V4      pVP                  P	                  4       '       d?   VP                  P
                  P                  4       '       d   V P                  WW4WW4      # V P                  WW4WW4      # r*   )	rU  rf  rp  ry   r   r  r  r  r  )r   r  r  r   r  rn  rS  rd  r  rT  re  ro  s   &&&&&&&&&   r   loadTritonSemantic.load?  s     00@//@--n=88??!4!4!=!=!?!?++Cug^ftt $$SwW_mmr   c               (    V ^8  d   QhRRRRRRRR/# )r/   desctl.tensor_descriptor_baserS  r  rd  r2   r   r3   )r4   s   "r   r5   r6   M  s,     	/ 	/$= 	/X[ 	/),	/19	/r   c                	   \        V\        P                  4      '       g   Q h\        VP                  4      p\        V4      V8X  g   Q R V R\        V4       24       hV P                  VRR7      pV P                  P                  VP                  W P                  V4      V P                  V4      4      pV P                  WaP                  4      # )	expected  offsets, but got Frequire_i64)rt   r<   tensor_descriptor_baser  r  _convert_to_ir_valuesr(   create_descriptor_loadr   rU  rf  r'   r  )r   r  offsetsrS  rd  ndimrz   s   &&&&&  r   descriptor_loadTritonSemantic.descriptor_loadM  s    $ 9 9::::4##$7|t#Wy6HW%WW#,,W%,HLL//WFfFfguFv040L0L_0]_{{1oo..r   c               $    V ^8  d   QhRRRRRR/# )r/   r  r  r`   r   r2   r   r3   )r4   s   "r   r5   r6   X  s#     / /(A /( /`d /r   c                	    \        V\        P                  4      '       g   Q h\        VP                  4      p\        V4      V8X  g   Q R V R\        V4       24       hVP
                  VP                  8X  g   Q hR# )r  r  N)rt   r<   r  r  r  r  )r   r  r`   r  r  s   &&&& r   validate_store_like"TritonSemantic.validate_store_likeX  sl    $ 9 9::::4##$7|t#Wy6HW%WW#{{d.....r   c               $    V ^8  d   QhRRRRRR/# r/   r  r  r`   r   r2   r3   )r4   s   "r   r5   r6   ^  s(     n n%> nx n]e nr   c                	    V P                  WV4       V P                  W!P                  4      pV P                  VR R7      pV P	                  V P
                  P                  VP                  VP                  V4      \        P                  4      # Fr  )
r  r   rM   r  r'   r(   create_descriptor_storer   r<   void)r   r  r`   r  s   &&&&r   descriptor_storeTritonSemantic.descriptor_store^  sm      g6		%,,,W%,H{{4<<??U\\[bcegelelmmr   c               $    V ^8  d   QhRRRRRR/# r  r3   )r4   s   "r   r5   r6   e  (     u u*C uH ubj ur   c                	   V P                  WV4       VP                  \        P                  \        P                  \        P
                  \        P                  \        P                  \        P                  09   g   Q R 4       hV P                  VRR7      p\        P                  P                  pV P                  V P                  P                  WAP                   VP                   V4      \        P"                  4      # Unsupported dtypeFr  )r  rM   r<   r   r=   r   rc   ra   rb   r  r   DESCRIPTOR_REDUCE_KINDADDr'   r(   create_descriptor_reducer   r  r   r  r`   r  r_   s   &&&& r   descriptor_atomic_add$TritonSemantic.descriptor_atomic_adde  s      g6zzbii299bjj"**VXVaVabbwdwwb,,W%,H((,,{{4<<@@{{TYT`T`bijlnlslsttr   c                	    \         P                  P                  4       pVP                  R 8H  ;'       d    VP                  ^Z8  # )cuda)r   activeget_current_targetbackendarch)r   targets   & r   _has_native_tmaTritonSemantic._has_native_tmal  s4    113&(>>V[[B->?r   c                	^   V\         P                  \         P                  \         P                  \         P                  \         P
                  \         P                  09   g   Q R 4       hV\         P
                  \         P                  09   d    V P                  4       '       g   Q R4       hR# R# )r  z-16-bit float types require native tma supportN)r<   r   r=   r   r   ra   rb   r  )r   rM   s   &&r   $_descriptor_atomic_min_max_supported3TritonSemantic._descriptor_atomic_min_max_supportedp  sp    BHHbii2::r{{[[p]pp[RZZ--''))Z+ZZ) .r   c               $    V ^8  d   QhRRRRRR/# r  r3   )r4   s   "r   r5   r6   u  r  r   c                	T   V P                  WV4       V P                  VP                  4       V P                  VR R7      p\        P
                  P                  pV P                  V P                  P                  WAP                  VP                  V4      \        P                  4      # r  )r  r  rM   r  r   r  MINr'   r(   r  r   r<   r  r  s   &&&& r   descriptor_atomic_min$TritonSemantic.descriptor_atomic_minu        g611$**=,,W%,H((,,{{4<<@@{{TYT`T`bijlnlslsttr   c               $    V ^8  d   QhRRRRRR/# r  r3   )r4   s   "r   r5   r6   |  r  r   c                	T   V P                  WV4       V P                  VP                  4       V P                  VR R7      p\        P
                  P                  pV P                  V P                  P                  WAP                  VP                  V4      \        P                  4      # r  )r  r  rM   r  r   r  MAXr'   r(   r  r   r<   r  r  s   &&&& r   descriptor_atomic_max$TritonSemantic.descriptor_atomic_max|  r  r   c               $    V ^8  d   QhRRRRRR/# r  r3   )r4   s   "r   r5   r6     r  r   c                	   V P                  WV4       VP                  \        P                  \        P                  \        P
                  \        P                  09   g   Q R 4       hV P                  VRR7      p\        P                  P                  pV P                  V P                  P                  WAP                  VP                  V4      \        P                  4      # r  )r  rM   r<   r   r=   r   r   r  r   r  ANDr'   r(   r  r   r  r  s   &&&& r   descriptor_atomic_and$TritonSemantic.descriptor_atomic_and        g6zzbii299bhhGG\I\\G,,W%,H((,,{{4<<@@{{TYT`T`bijlnlslsttr   c               $    V ^8  d   QhRRRRRR/# r  r3   )r4   s   "r   r5   r6     s(     u u)B u8 uai ur   c                	   V P                  WV4       VP                  \        P                  \        P                  \        P
                  \        P                  09   g   Q R 4       hV P                  VRR7      p\        P                  P                  pV P                  V P                  P                  WAP                  VP                  V4      \        P                  4      # r  )r  rM   r<   r   r=   r   r   r  r   r  ORr'   r(   r  r   r  r  s   &&&& r   descriptor_atomic_or#TritonSemantic.descriptor_atomic_or  s      g6zzbii299bhhGG\I\\G,,W%,H((++{{4<<@@{{TYT`T`bijlnlslsttr   c               $    V ^8  d   QhRRRRRR/# r  r3   )r4   s   "r   r5   r6     r  r   c                	   V P                  WV4       VP                  \        P                  \        P                  \        P
                  \        P                  09   g   Q R 4       hV P                  VRR7      p\        P                  P                  pV P                  V P                  P                  WAP                  VP                  V4      \        P                  4      # r  )r  rM   r<   r   r=   r   r   r  r   r  XORr'   r(   r  r   r  r  s   &&&& r   descriptor_atomic_xor$TritonSemantic.descriptor_atomic_xor  r  r   c               $    V ^8  d   QhRRRRRR/# )r/   rS  r  rd  r2   r   r3   )r4   s   "r   r5   r6     s#     $ $3 $ad $iq $r   c                	   \        V\        P                  4      '       g   Q hVR 8X  g   Q R4       hVR 8X  g   Q R4       h\        VP                  4      ^8X  g   Q RVP                   24       hVP                  ^ ,          ^8X  g   Q RVP                   24       h\        VP
                  4      ^8X  g   Q RVP
                   24       hVP
                  ^ ,          ^8  g   Q RVP
                   24       hVP                  p^ VP                  ,          ^,          pVP                  ^,          V8  g"   Q RV RV R	VP                  ^,           24       h\        P                  ! VP                  VP
                  ^ ,          VP                  ^,          .4      pV P                  V3R
R7      ^ ,          pV P                  P                  VP                  VP                  W8P                  V P                  4      4      p	V P                  W4      # ) z#cache modifier is not supported yetz$eviction policy is not supported yetdescriptor must be 2D, but got *descriptor block must have 1 row, but got x offsets must be 1D, but got z5descriptor gather must have at least 8 rows, but got zdescriptor gather of  must have at least  columns, but got Fr  )rt   r<   r  r  r  r  rM   r(  r  r  r(   create_descriptor_gatherr   r   r'   )
r   r  	x_offsetsy_offsetrS  rd  rM   min_colsry   rz   s
   &&&&&&    r   descriptor_gather TritonSemantic.descriptor_gather  s   $ 9 9::::#J%JJ#"$L&LL$ 4##$)_-LTM]M]L^+__)"a'h+UVZVfVfUg)hh' 9??#q(\,J9??J[*\\( q!Q&q*_`i`o`o_p(qq&

111A5 	A3E7:NxjXjkok{k{|}k~j  A	A  }}TZZ)//!*<d>N>Nq>Q)RS--xl-NqQLL11$++y?O?OQY[e[efjfrfr[st{{1##r   c                    V ^8  d   QhRRRR/# r  r3   )r4   s   "r   r5   r6     s     * *h *PX *r   c                	Z   \        V\        P                  4      '       g   Q h\        VP                  4      ^8X  g   Q RVP                   24       hVP                  ^ ,          ^8X  g   Q RVP                   24       h\        VP
                  4      ^8X  g   Q RVP                   24       hVP
                  ^ ,          ^8  g   Q RVP
                   24       hVP                  p^ VP                  ,          ^,          pVP                  ^,          V8  g"   Q RV RV RVP                  ^,           24       hV P                  V3RR	7      ^ ,          pV P                  P                  VP                  VP                  VP                  V4       V P                  R
\        P                  4      # )r/   r  r  r  z6descriptor scatter must have at least 8 rows, but got zdescriptor scatter of r  r  Fr  N)rt   r<   r  r  r  r  shapaerM   r(  r  r(   create_descriptor_scatterr   r'   r  )r   r  r`   r  r	  rM   r
  s   &&&&&  r   descriptor_scatter!TritonSemantic.descriptor_scatter  s   $ 9 9:::: 4##$)_-LTM]M]L^+__)"a'h+UVZVfVfUg)hh' 9??#q(],J9K[K[J\*]]( q!Q&r*`ajapap`q(rr&

111A5 	B4UG;OPXzYklpl|l|}~l  lA  B	B  --xl-NqQ..t{{ELL)JZJZ\de{{4))r   c           	     	   Ve   \        R4      hVP                  P                  P                  4       pVP                  P	                  4       '       g   V P                  W'4      pVP                  P	                  4       '       g   Q R4       hWrP                  P                  4       8X  g'   Q RV RVP                  P                  4        R24       hVP                  P                  P                  VP                  P                  8X  gA   Q RVP                  P                  P                   RVP                  P                   R24       hVP                  P                  P                  pV\        P                  8w  g   Q R4       hV P                  WG4      pV P                  W(4      pV P                  V P                  P                  VP                  VP                  WEV4      \        P                  4      # )	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:   ry   r  r  r  r
  r<   r   r  r   r'   r(   create_tensor_pointer_storer   r  )	r   r  valr  r  rT  re  r  r  s	   &&&&&&&  r   _store_block_pointer#TritonSemantic._store_block_pointer  s    jkk hh))::<xx  ""++C=Cxx  ""S$SS"hh77 
 
 	a+&89R9R9T8UU_`	a 
xx""--1D1DD  	uH[\_\d\d\o\o\z\z[{  |U  VY  V^  V^  Vi  Vi  Uj  jt  Gu  	uD$$// X"XX  ::>W ii$ {{LL44SZZ^dlmoqovovx 	xr   c           	     	   VP                   P                  P                  4       '       g(   \        R VP                   P	                  4        R24      hV'       d   \        R4      hVP                   P                  4       '       g_   VP                   P                  4       '       d   \        R4      hV'       d,   VP                   P                  4       '       d   \        R4      hVP                   P                  4       '       de   VP                  pVf   V P                  W4      w  rMV P                  WV4      w  rpWqP                  8w  d   \        RVP                   RV 24      hVP                   P                  pVP                  p	V	\        P                  8X  dB   \        P                  p	\        P                  ! WP                  4      pV P                  W4      pV P                  W)4      pVfP   V P                  V P                   P#                  VP$                  VP$                  WV4      \        P&                  4      # VP                   P                  P)                  4       '       g   \        R4      hV P                  V P                   P+                  VP$                  VP$                  VP$                  WV4      \        P&                  4      # )r  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(Expected pointer argument to have shape z	 but got "Mask must have boolean scalar type)ry   r   r   r:   r   r  r  broadcast_tensorsr  r<   r   r  r  r  r   r'   r(   create_storer   r  r:  create_masked_store)
r   r  r  r  r  rT  re  	ptr_shaper  r  s
   &&&&&&&   r   _store_legacyTritonSemantic._store_legacy  s"   xx%%''4SXX5F5F5H4IXYY  E F F
 xx  ""xx  "" !ijj		**,, !hii 88		I|11#;S!%!7!7$!G$II% #KCII;V_`i_j!kll"" RWWWWF__V-A-ABF))C(C ii$ <;;t||88SZZQVacecjcjkkyy''))ABB{{4<<;;CJJ

TXT_T_afq77$ 	$r   c               0    V ^8  d   QhRRRRRRRRRRR	R/# )
r/   r  r   r  r  r  rS  r  rd  r2   r3   )r4   s   "r   r5   r6     sB     W W W W8J Wlo W"W'/Wr   c                	   V P                  V4      pV P                  V4      pVP                  P                  4       '       g+   VP                  P                  P                  4       '       d   \        R 4      hVP                  P                  4       '       d>   VP                  P                  P                  4       '       d   V P                  WW4Wx4      # V P                  WW4Wx4      # )"Cannot store to a constant pointer)r[  rf  ry   is_constr   r:   r   r  r  r  r  )	r   r  r  r  r  rS  rd  rT  re  s	   &&&&&&&  r   storeTritonSemantic.store  s     11.A//@88#((//":":"<"<ABB88??!4!4!=!=!?!?,,StU]] %%ceVVr   c               0    V ^8  d   QhRRRRRRRRRRRR/# )	r/   r  r   cmpr  r~  r  r  r2   r3   )r4   s   "r   r5   r6   %  s@     u uh uX uH u3 uWZ u_g ur   c           	     	j   V P                  V4      pV P                  V4      pVP                  P                  P                  pVP
                  R9  d   \        R4      hV P                  V P                  P                  VP                  VP                  VP                  WE4      VP                  4      # )   z9atomic_cas only supports elements with width {16, 32, 64})r)      r   )r  r  ry   r   r  r(  r:   r'   r(   create_atomic_casr   )r   r  r'  r  r~  r  r  s   &&&&&& r   
atomic_casTritonSemantic.atomic_cas%  s    s#""5)XX__//
((<XYY{{4<<99#**cjjRUR\R\^aiknksksttr   c          
     ,    V ^8  d   QhRRRRRRRRRR/# )	r/   r  r   r  r  opr  r2   z#Tuple[TensorTy, TensorTy, TensorTy]r3   )r4   s   "r   r5   r6   -  s1      h X X '*/Rr   c                	   VP                   P                  P                  4       '       g+   \        R VP                   P	                  4       ,           4      hVP                   P                  4       '       g+   VP                   P                  P                  4       '       d   \        R4      hVP                   P                  P                  pV\        P                  J d!   VR8w  d   \        RV,           R,           4      hV\        P                  J d!   VR8w  d   \        RV,           R,           4      hV\        P                  \        P                  39   g   VP                  ^8  d*   \        RV,           R,           \        V4      ,           4      hVP                   P                  4       '       d[   Ve*   V P                  W1P                   P!                  4       4      pVe*   V P                  W!P                   P!                  4       4      pV P#                  W!P                   P                  P                  4      pVf   V P$                  P'                  R4      p\        P(                  pVP                   P                  4       '       d_   VP                   P+                  \        P(                  4      pV P$                  P-                  VP/                  V P$                  4      V4      pV P1                  Wg4      pWV3# )z)Pointer argument of store instruction is r"  r   atomic_z does not support fp16z does not support bf16z does not support T)ry   r   r   r:   r   r#  r  r<   ra   rb   int16uint16r(  r  r  r
  r  r   r(   get_int1r   r   r  r   r'   )r   r  r  r  r/  r  mask_irmask_tys   &&&&&   r   atom_red_typechecking_impl)TritonSemantic.atom_red_typechecking_impl-  s   xx%%''H388K\K\K^^__88#(("5"5">">"@"@ABBXX__//
#eY^.FFGG$uY^.FFGG"((BII..*2O2ORT2TY^.BBS_TUU8800xx7P7P7RS//XX5N5N5PQiiXX__778<ll++D1GggGxx  ""((22277;,,33GMM$,,4OQXY;;w0D~r   c                    V ^8  d   QhRRRR/# )r/   rz   r   r2   r3   )r4   s   "r   r5   r6   I  s     + +( +x +r   c                	    VP                   P                  p\        P                  ! VR R7      pV P	                  W4      pV P                  WB^,
          4      pV P                  V\        P                  4      # )F)rH  signed)rM   r(  r<   get_int_dtyper?  rP  r   r   )r   rz   rH  idtypeixsignbits   &&    r   _signbitTritonSemantic._signbitI  sX    77--!!8EB\\!$))B1-yy"''**r   c               0    V ^8  d   QhRRRRRRRRRRRR/# 	r/   r  r   r  r  r~  r  r  r2   r3   )r4   s   "r   r5   r6   P  8     $) $)h $)X $)X $)C $)X[ $)`h $)r   c                	P   V P                  WVR 4      w  rpV P                  V4      pV P                  V4      pVP                  P                  pVP                  4       '       d   VP                  4       '       dp   V P                  V P                  P                  \        P                  P                  VP                  VP                  VP                  WE4      VP                  4      # V P                  V P                  P                  \        P                  P                  VP                  VP                  VP                  WE4      VP                  4      # V\        P                   \        P"                  09  d   \%        RV 24      hV\        P                   8X  d   \        P&                  M\        P(                  pV P+                  W'4      pV P+                  V\        P,                  ! V^4      4      p	V\        P                   8X  d   \        P.                  M\        P0                  p
V P+                  W*4      pV P+                  V\        P,                  ! V
^4      4      pV P3                  V4      pV P5                  V4      pV P                  V P                  P                  \        P                  P                  V	P                  VP                  V P7                  W>4      P                  WE4      VP                  4      pV P                  V P                  P                  \        P                  P8                  VP                  VP                  V P7                  W=4      P                  WE4      VP                  4      pV P;                  WV4      pV P+                  VV4      # )r  z#atomic_max not supported for dtype )r7  r  r  ry   r   rj   r   r'   r(   create_atomic_rmwr   	ATOMIC_OPr  r   UMAXr<   rc   re   rP   r=   r   r?  r  r   r   r@  rJ  r   UMINwherer   r  r  r  r~  r  sca_tyi_typei_vali_ptrui_typeui_valui_ptrnegpospos_retneg_retr   s   &&&&&&            r   
atomic_maxTritonSemantic.atomic_maxP     884O$s#""5)==??##%%{{LL222<<3C3CSZZQTQ[Q[]a]h]hjmuHH  {{LL222<<3D3DcjjRUR\R\^b^i^iknvHH  "**bjj11A&JKK#rzz1rxxS)S"//&!"<=%3"))c+c2??7A#>?mmC iin++LL**2<<+;+;U\\5<<+/99T+?+F+FTUZU_U_a ++LL**2<<+<+<fmmV]]+/99T+?+F+FTU[U`U`b jjw/||C((r   c               0    V ^8  d   QhRRRRRRRRRRRR/# rC  r3   )r4   s   "r   r5   r6   v  rD  r   c                	P   V P                  WVR 4      w  rpV P                  V4      pV P                  V4      pVP                  P                  pVP                  4       '       d   VP                  4       '       dp   V P                  V P                  P                  \        P                  P                  VP                  VP                  VP                  WE4      VP                  4      # V P                  V P                  P                  \        P                  P                  VP                  VP                  VP                  WE4      VP                  4      # V\        P                   \        P"                  09  d   \%        RV 24      hV\        P                   8X  d   \        P&                  M\        P(                  pV P+                  W'4      pV P+                  V\        P,                  ! V^4      4      p	V\        P                   8X  d   \        P.                  M\        P0                  p
V P+                  W*4      pV P+                  V\        P,                  ! V
^4      4      pV P3                  V4      pV P5                  V4      pV P                  V P                  P                  \        P                  P                  V	P                  VP                  V P7                  W>4      P                  WE4      VP                  4      pV P                  V P                  P                  \        P                  P8                  VP                  VP                  V P7                  W=4      P                  WE4      VP                  4      pV P;                  WV4      pV P+                  VV4      # )r  z#atomic_min not supported for dtype )r7  r  r  ry   r   rj   r   r'   r(   rF  r   rG  r  r   rI  r<   rc   re   rP   r=   r   r?  r  r   r   r@  rJ  r   rH  rJ  rK  s   &&&&&&            r   
atomic_minTritonSemantic.atomic_minv  rY  r   c               0    V ^8  d   QhRRRRRRRRRRRR/# rC  r3   )r4   s   "r   r5   r6     s8     % %h %X %X %C %X[ %`h %r   c           
     	   V P                  WVR 4      w  rpV P                  V4      pV P                  V4      pVP                  P                  pVP                  4       '       d   \        P                  P                  M\        P                  P                  pV P                  V P                  P                  WqP                  VP                  VP                  WE4      VP                  4      # )r   )r7  r  r  ry   r   r   r   rG  FADDr  r'   r(   rF  r   )r   r  r  r  r~  r  rL  r/  s   &&&&&&  r   
atomic_addTritonSemantic.atomic_add  s    884O$s#""5)"("4"4"6"6R\\BLL<L<L{{4<<99"jj#**VZVaVacfn88% 	%r   c               0    V ^8  d   QhRRRRRRRRRRRR/# rC  r3   )r4   s   "r   r5   r6     @     y yh yX yX yC yX[ y`h yr   c           
     	P   V P                  WVR 4      w  rpV P                  V4      pV P                  V4      pV P                  V P                  P                  \        P                  P                  VP                  VP                  VP                  WE4      VP                  4      # )and)r7  r  r  r'   r(   rF  r   rG  r  r   ry   r   r  r  r  r~  r  s   &&&&&&r   
atomic_andTritonSemantic.atomic_and      884O$s#""5){{LL**2<<+;+;SZZUYU`U`bemorowowy 	yr   c               0    V ^8  d   QhRRRRRRRRRRRR/# rC  r3   )r4   s   "r   r5   r6     s@     x xX xH xH x3 xWZ x_g xr   c           
     	P   V P                  WVR 4      w  rpV P                  V4      pV P                  V4      pV P                  V P                  P                  \        P                  P                  VP                  VP                  VP                  WE4      VP                  4      # )or)r7  r  r  r'   r(   rF  r   rG  r  r   ry   rg  s   &&&&&&r   	atomic_orTritonSemantic.atomic_or  s    884N$s#""5){{LL**2<<??CJJ

TXT_T_adlnqnvnvx 	xr   c               0    V ^8  d   QhRRRRRRRRRRRR/# rC  r3   )r4   s   "r   r5   r6     rd  r   c           
     	P   V P                  WVR 4      w  rpV P                  V4      pV P                  V4      pV P                  V P                  P                  \        P                  P                  VP                  VP                  VP                  WE4      VP                  4      # )xor)r7  r  r  r'   r(   rF  r   rG  r  r   ry   rg  s   &&&&&&r   
atomic_xorTritonSemantic.atomic_xor  rj  r   c               0    V ^8  d   QhRRRRRRRRRRRR/# rC  r3   )r4   s   "r   r5   r6     s8      x h h S Y\ ai r   c           
     	P   V P                  WVR 4      w  rpV P                  V4      pV P                  V4      pV P                  V P                  P                  \        P                  P                  VP                  VP                  VP                  WE4      VP                  4      # )xchg)r7  r  r  r'   r(   rF  r   rG  XCHGr   ry   rg  s   &&&&&&r   atomic_xchgTritonSemantic.atomic_xchg  s    884P$s#""5){{LL**2<<+<+<cjj#**VZVaVacfnHH 	r   c                	H   VP                  4       V P                  P                  P                  9   g,   Q R V P                  P                  P                   RV 24       hVP	                  4       pVR8X  d   RpVR8X  d   RpVR8X  d   Rp\        \        P                  V4      # )zinput_precision must be one of . Got TF32X3TF32x3BF16X3BF16x3BF16X6BF16x6)lowerr(   r   allowed_dot_input_precisionsupperr  r   INPUT_PRECISION)r   input_precisions   &&r   _str_to_dot_input_precision*TritonSemantic._str_to_dot_input_precision  s    $$&$,,*>*>*[*[[ 	y-dll.B.B._._-``fgvfwx	y[)//1h&&Oh&&Oh&&Or))?;;r   c               4    V ^8  d   QhRRRRRRRRRRR	R
RR/# )r/   r   r   r   accr  r  max_num_imprecise_accr1   	out_dtyperH   r2   r3   )r4   s   "r   r5   r6     sS     [y [yx [yh [yX [yP] [y#&[y3;[y@H[yr   c           
     	   VP                   P                  4       '       d!   VP                   P                  4       '       g   Q hVP                  P                  4       '       d#   VP                  P                  4       '       d   EM:VP                  \        P
                  \        P                  \        P                  \        P                  \        P                  \        P                  39   g   Q R VP                   24       hVP                  \        P
                  \        P                  \        P                  \        P                  \        P                  \        P                  39   g   Q RVP                   24       hVP                  VP                  8X  g"   Q RVP                   RVP                   24       hVP                  P                  4       '       g!   VP                  P                  4       '       d|   RV P                  P                  P                  9   d   \        P                   ! R4       V P#                  V\        P                  4      pV P#                  V\        P                  4      pVP                  P%                  4       ;'       g    VP                  P%                  4       pVP                  P'                  4       ;'       g    VP                  P'                  4       pV'       g	   V'       d   V'       d   RMRp	WP                  P                  P                  9   d   V P                  P                  P(                  p
\        P                   ! V	 RV
 R	V
 R
24       V P#                  V\        P                  4      pV P#                  V\        P                  4      pVf!   V P                  P                  P*                  pV P-                  V4      p\/        VP0                  4      p\/        VP0                  4      pYu;8X  d   ^8X  g5   M Yu;8X  d   ^8X  g%   M Q RVP0                   RVP0                   R24       hVP0                  R,          P2                  VP0                  R,          P2                  8X  g_   Q RVP0                   RVP0                   RVP0                  R,          P2                   RVP0                  R,          P2                   R2	4       hV P                  P4                  P7                  R4       f   Q R4       hV P                  P4                  R,          ! VP                   VP                   4      pVP0                  R,          P2                  V^ ,          8  dS   VP0                  R,          P2                  V^,          8  d*   VP0                  R,          P2                  V^,          8  g&   Q RV^ ,           RV^,           RV^,           24       hVP                   P8                  P;                  4       '       d^   VP                   P8                  \        P
                  8X  g   Q R4       hV P                  P=                  ^ 4      p\        P>                  pEMEVPA                  4       '       d   \C        R4      hVP                   P8                  PE                  4       '       g+   VP                   P8                  PA                  4       '       d-   V P                  PG                  ^ 4      p\        P                  pMVP                   P8                  PI                  4       '       d-   V P                  PK                  ^ 4      p\        P                  pMNVPM                  4       '       d   V P                  PO                  ^ 4      MV P                  PG                  ^ 4      pTpVP                   P0                  R,          pVP                   P0                  R,          pVP                   P0                  R,          pV^8X  d   VP                   P0                  ^ ,          MRp\        PP                  ! TV'       d   VVV.MVV.4      pVf7   V P                  PS                  VPU                  V P                  4      V4      pMNVPV                  pVP                   P0                  VP0                  8X  d   VP                   PX                  V8X  g   Q hVfe   VP                  P                  4       '       dB   VP                  P                  4       '       d"   V P                  P                  PZ                  pM\^ pMYVP                  P                  4       '       d:   VP                  P                  4       '       d   VV8  d   \C        RV RV R24      hV P]                  V P                  P_                  VPV                  VPV                  VWE4      V4      # )zUnsupported lhs dtype zUnsupported rhs dtype z&Both operands must be same dtype. Got r   fp8e4b15zthe use of fp8e4b15 is deprecated on Hopper and later architectures and can cause significant slow down. It will be removed in a future triton releasefp8e4b8fp8e5b16z- is AMD gfx942 specific and not supported on z^ so it's upcasted to fp16 and can cause significant slow down. Please use OCP fp8 variants on z for performanceN+Both inputs must be either 2D or 3D; (lhs: 	 vs rhs: rd  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.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 (r  )0ry   r  rM   ri   r<   r  uint8ra   rb   rc   re   r5  r(   r   !deprecated_fp8_dot_operand_dtypeswarningswarnr   
is_fp8e4b8is_fp8e5b16r  default_dot_input_precisionr  r  r  r`   r6  r7  r   rj   	get_int32r=   rh   r:   rf   r  rd   get_fp64rg   get_fp16r  r  r   r   r  max_num_imprecise_acc_defaultr'   
create_dot)r   r   r   r  r  r  r  uses_fp8e4b8uses_fp8e5b16	type_namer  lhs_rankrhs_rankr  rf  ret_scalar_tyMNKBr   
acc_handles   &&&&&&&               r   dotTritonSemantic.dot  s   xx  ""sxx'8'8':':::99#))"2"2"4"499"((BJJRZZ!#!- - S0Fsyyk.RS -99"((BJJRZZ!#!- - S0Fsyyk.RS -99		)o-STWT]T]S^^cdgdmdmcn+oo)99  ""cii&;&;&=&=T\\11SSS m ))C,C))C,Cyy++-GG1E1E1G		--/JJ3993H3H3J=%1	zILL00RRR||++00 k!Ntf U66:V;KMN iiRZZ0iiRZZ0""ll22NNO::?Ksyy>syy>(q(H,EA,E  	VItuxu~u~t  @I  JM  JS  JS  IT  TU  HV  	VEyy}""cii' 	u,SYYK7PQTQZQZP[  \Y  Z]  Zc  Zc  df  Zg  Zm  Zm  Yn  n^  _b  _h  _h  ik  _l  _r  _r  ^s  st  u	u ||''++#'( 	^)]	^ (||//?#((Syy}""l1o5#))B-:M:MQ]^_Q`:`		"##|A6	v0a0AVWHYYcdpqrdsctu	v 7 88??!!##88??bgg-E/EE-''*BHHM  z  XX__$$&&#((//*A*A*C*C&&q)BJJMXX__$$&&&&q)BJJM-6->->-@-@&&q)dllF[F[\]F^B%MHHNN2HHNN2HHNN2!)QCHHNN1D}1q!Qi1a&I;226<<3MrRJJ88>>V\\1chh6I6IY6VVV !(yy!!cii&6&6&8&8(,(<(<(Z(Z%()%yy!!cii&6&6&8&8=RUV=V #:;P:QQabcadde!fgg{{LL##CJJ

Joqwy 	yr   c                   V ^8  d   QhRR/# )r/   float_formatr  r3   )r4   s   "r   r5   r6   /  s      C r   c                	~    \        \        P                  VP                  4       R 4      pVf   \	        RV R24      hV# )NzInvalid float format: r   )r  r   ScaleDotElemTypeTYr  r:   )r   r  ty_enums   && r   _str_to_fp_typeTritonSemantic._str_to_fp_type/  s>    "//1C1C1EtL?5l^1EFFr   c                    V ^8  d   QhRRRR/# )r/   r  r   r  r  r3   )r4   s   "r   r5   r6   5  s     0 0x 0s 0r   c                b   R\         P                  R\         P                  R\         P                  R\         P                  /P                  V4      pVfG   VR8X  g   Q RV 24       hVP                  \         P                  8X  g   Q RVP                   24       hV# VP                  V8X  d   V# R\         P                  R\         P                  R\         P                  R\         P                  /V,          pVP                  V8X  g   Q RV R	VP                   24       hV P                  W4      # )
z
If float_format is subbyte, make sure it's packed as uint8 and return it.
Otherwise, return a tensor (perhaps bitcasting) of the specified float format.
e5m2e4m3bf16fp16e2m1z)Internal Error: Unexpected float format: z)e2m1 format must be packed as uint8. Got zUnexpected dtype for r|  )
r<   float8e5
float8e4nvrb   ra   r7  rM   r  r3  r?  )r   r  r  	triton_tyunsigned_tys   &&&  r   _bitcast_to_fp_type"TritonSemantic._bitcast_to_fp_type5  s   
 R[[&"--V\ZZ!!$\!2 	6)e-VWcVd+ee)99(a,UVYV_V_U`*aa(J99	!J!288VRXXvryyRXZ\ZcZcdeqrK99+d/D\NRXY\YbYbXc-dd+<<//r   c                	   Vec   VP                   P                  4       '       d   ^M^ pVP                  P                  pVRR  WV,          .8X  g   Q RV RW6,           RV 24       hVee   VP                   P                  4       '       d   ^M^ pVP                  P                  pVRR  W#V,          .8X  g   Q RV RW6,           RV 24       hR # R # )Nz*lhs_scale must be a tensor of shape [..., r  z]. Got z*rhs_scale must be a tensor of shape [..., r  )rM   
is_fp8e4nvry   r  )	r   r  r  r  	lhs_scale	rhs_scalescale_factorlhs_scale_shaperhs_scale_shapes	   &&&&&&   r   verify_scaled_shape"TritonSemantic.verify_scaled_shapeG  s     !*!;!;!=!=22L'nn22O"23'$,  m;A3bAR@SSZ[jZklm   !*!;!;!=!=22L'nn22O"23'$,  m;A3bAR@SSZ[jZklm  !r   c               H    V ^8  d   Qh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  
lhs_formatr  r   r  r  
rhs_formatr  zTensorTy | None	fast_mathrY   
lhs_k_pack
rhs_k_packr  rH   r2   r3   )r4   s   "r   r5   r6   U  sw     3t 3th 3t8 3t 3tS[ 3t03t>A3tHW3tdh3t#3t153tBJ3tOW3tr   c                	R   VP                   P                  4       '       d!   VP                   P                  4       '       g   Q h\        VP                  4      p\        VP                  4      pYu;8X  d   ^8X  g5   M Yu;8X  d   ^8X  g%   M Q RVP                   RVP                   R24       hVP                  pVP                  pV P                  V4      pV P                  V4      p0 R
mpVV9   g   Q RV 24       hVV9   g   Q RV 24       hVRJ ;'       g1    \        V\        P                  4      ;'       d    VP                  RJ pVRJ ;'       g1    \        V\        P                  4      ;'       d    VP                  RJ pV P                  W4      pV P                  WF4      pV	'       g   VR8X  g   Q R4       hV
'       g   VR8X  g   Q R4       hVP                   P                  RR w  ppVP                   P                  RR w  ppVR8X  d   ^M^pVR8X  d   ^M^pV	'       d
   VV,          MTpV
'       d
   VV,          MTpVV8X  g#   Q R	VP                   RVP                   R24       hV^8X  d   VP                   P                  ^ ,          MRpTpV	'       g   VV,          pM	VV,          pV
'       g
   VV,          p\        P                  ! TV'       d   VVV.MVV.4      pV P                  P                  ^ 4      pVf7   V P                  P                  VP                  V P                  4      V4      pMNVP                  pVP                   P                  VP                  8X  d   VP                   P                   V8X  g   Q hV'       d   RMVP                  p V'       d   RMVP                  p!T P#                  TTTV'       d   RMTV'       d   RMT4       V P%                  V P                  P'                  VP                  V!WP                  V WWV4
      V4      # )r/   r  r  rd  r  zNYI: lhs_format zNYI: rhs_format NzBonly mxfp4 inputs can be packed along a dimension different than KzCReduction dimension should pack the same number of elements; (lhs: >   r  r  r  r  r  r  )ry   r  r  r  r`   r  rt   r<   ru   r  r  r(   r  r  r   r   r  r  r'   create_dot_scaled)"r   r   r  r  r   r  r  r  r  r  r  r  r  r  lhs_format_enumrhs_format_enumallowed_formatsrhs_scale_is_nonelhs_scale_is_noner  K_LHSK_RHSr  PACKED_APACKED_BPACKED_A_DIMPACKED_B_DIMr  r  r   rf  r  rhs_scale_handlelhs_scale_handles"   &&&&&&&&&&&&                      r   
dot_scaledTritonSemantic.dot_scaledU  s    xx  ""sxx'8'8':':::syy>syy>(q(H,EA,E  	VItuxu~u~t  @I  JM  JS  JS  IT  TU  HV  	VE$**
$**
..z:..z:B_,M0@.MM,_,M0@.MM,%-rr*Y2U2q2qZcZiZimqZq%-rr*Y2U2q2qZcZiZimqZq&&s7&&s7Z61w3ww1Z61w3ww188>>"#&588>>"#&q"f,1!"f,1!+5x%'5+5x%'5|+  	T/rsvs|s|r}  ~G  HK  HQ  HQ  GR  RS  .T  	T+!)QCHHNN1DHAHAHAyq1a)q!fE\\""1%;226<<3MrRJJ88>>V\\1chh6I6IY6VVV#44):J:J#44):J:J  Aq2C$):		K{{LL**3::7GZdZdfv+:z_iklrt 	tr   c               (    V ^8  d   QhRRRRRRRR/# )r/   	conditionr   rz   r  r2   r3   )r4   s   "r   r5   r6     s.     e ex eH e eh er   c                	R   VP                   \        P                  8w  d$   \        P                  ! R VP                    24       V P                  V\        P                  4      pV P                  W#RR4      w  r#VP                  P                  4       '       d(   V P                  W4      w  rV P                  W#4      w  r#MV P                  W4      w  rVP                  pV P                  V P                  P                  VP                  VP                  VP                  4      V4      # )zgtl.where with a non-boolean condition is deprecated and will error out in a future triton release. Got T)rM   r<   r   r  r  r   r   ry   r  r   r'   r(   create_selectr   )r   r  rz   r  r  r   s   &&&&  r   rJ  TritonSemantic.where  s    ??bgg%MMy  {D  {J  {J  zK  L IIi1	00tTB>>""$$44YBLI,,Q2DAq44YBLI{{4<<55i6F6FRSRZRZ[]cddr   c                	f    V'       d   \         P                  ! W#4      pMTpV P                  W4      # r*   )r<   r  r'   )r   rz   rk   r  res_tys   &&&& r   wrap_tensorTritonSemantic.wrap_tensor  s)    ]]98F F{{1%%r   c               $    V ^8  d   QhRRRRRR/# )r/   inputsSequence[TensorTy]r0   r1   r2   Tuple[TensorTy, ...]r3   )r4   s   "r   r5   r6     s(     u u 2 u# uUi ur   c                	B  a aaa	a
 Vf5   \         ;QJ d    . V 3R lS 4       F  NK  	  5M! V 3R lS 4       4      o^ pS^ ,          P                  P                  o
\        S
4      pW$8  g   Q RV R24       h\	        S
4       UUu. uF  w  rVWR8w  g   K  VNK  	  uppo	\
        ;QJ d    V
3R lS 4       F  '       d   K   RM	  RM! V
3R lS 4       4      '       g   Q R4       hS P                  P                  S Uu. uF  qwP                  NK  	  upV4      oV! S4       SP                  4       '       g   Q h\         ;QJ d,    . VVV	V 3R l\        \        S4      4       4       F  NK  	  5# ! VVV	V 3R l\        \        S4      4       4       4      # u uppi u upi )	Nc              3  t   <"   T F-  pSP                  WP                  P                  .R R7      x  K/  	  R# 5i)Tr  N)r  r  r`   )r  tr   s   & r   r  +TritonSemantic.reduction.<locals>.<genexpr>  s.     ^W]RS4<<GGMM?<MMW]s   58z&reduction axis must be < inputs rank (rd  c              3  T   <"   T F  qP                   P                  S8H  x  K  	  R # 5ir*   )ry   r  )r  r  r  s   & r   r  r    s     9&Q66<<5(&s   %(FTz-all reduction inputs must have the same shapec              3     <"   T FB  pSP                  SP                  V4      SV,          P                  P                  S4      x  KD  	  R # 5ir*   r  
get_resultry   r   )r  r  r  	reduce_opr  r   s   & r   r  r    sB      uas\]DY11!4fQinn6K6KYWWas   A
A)tuplery   r  r  r  allr(   create_reducer   verifyr  )r   r  r0   region_builder_fnrankr  r  r  r  r  r  s   ff&&    @@@r   	reductionTritonSemantic.reduction  sS   <U^W]^UU^W]^^FDq	$$5z{LDTF!LL{#,U#3A#341qyQQ#3A	s9&9sss9&999j;jj9LL..&/I&Q&/I4P	)$!!!!u uafgjkqgrasuu 	uu uafgjkqgrasu u 	u B 0Js   
FF>Fc               (    V ^8  d   QhRRRRRRRR/# )	r/   r  r  r0   r1   reverserY   r2   r  r3   )r4   s   "r   r5   r6     s1     x x'9 x x"&x+?xr   c                	z  a aaa S^ ,          P                   P                  o\        S4      pV) Tu;8:  d   V8  g   M Q RV RV R24       hV^ 8  d	   W%,          pS F%  pVP                   P                  S8X  d   K   Q R4       h	  S P                  P	                  S Uu. uF  qfP
                  NK  	  upW$4      oV! S4       SP                  4       '       g   Q h\        ;QJ d,    . VVV V3R l\        \        S4      4       4       F  NK  	  5# ! VVV V3R l\        \        S4      4       4       4      # u upi )r   z
scan axis z must be < inputs rank (rd  z(all scan inputs must have the same shapec              3     <"   T FB  pSP                  SP                  V4      SV,          P                  P                  S4      x  KD  	  R # 5ir*   r  )r  r  r  scan_opr   r  s   & r   r  2TritonSemantic.associative_scan.<locals>.<genexpr>  s@     wdv_`T%%g&8&8&;VAY^^=R=RTYZZdvr  )	ry   r  r  r(   create_scanr   r  r  r  )	r   r  r0   r  r  r  r  r  r  s	   ff&&&  @@r   associative_scanTritonSemantic.associative_scan  s   q	$$5zu#t#Wz$7OPTvUV%WW#!8LDA66<<5(T*TT(  ,,**f+EfHHf+EtU'"~~uwdijmntjudvwuwuwdijmntjudvwww	 ,Fs    D8c               (    V ^8  d   QhRRRRRRRR/# )r/   srcr   indexr0   r1   r2   r3   )r4   s   "r   r5   r6     s.     K K( K8 K3 K8 Kr   c                	   VP                   P                  4       '       g   Q R 4       h\        VP                  P                  4      p\        VP                  P                  4      V8X  g   Q R4       hV) Tu;8:  d   V8  g   M Q RV RV R24       hV^ 8  d	   W4,          p\        V4       FS  pWS8X  d   K  VP                  P                  V,          VP                  P                  V,          8X  d   KJ  Q RV R24       h	  V P                  P                  VP                  VP                  V4      pV P                  WaP                  P                  VP                  P                  4      # )zindex must be an integer tensorz0source and index tensors must have the same rankzgather axis z must be < source rank (rd  z
index dim z( must match the corresponding source dim)rM   rj   r  ry   r  r  r(   create_gatherr   r  r   )r   r  r  r0   r  r  gathers   &&&&   r   r  TritonSemantic.gather  s'   {{!!##F%FF#388>>"5::##$,`.``,u#t#Y|D69QRVQWWX%YY#!8LDtAy::##A&#((..*;;xz$Ow=xx; 
 ++CJJdK9I9IJJr   c                	   V'       g   R# Vvr#\        \        V4      4       F   pV P                  W#V,          4      w  q#V&   K"  	  \        \        V4      ^,
          4       F   pV P                  W#V,          4      w  q#V&   K"  	  V.VO5# )r9   r3   )r  r  r   )r   r  headtailr  s   &*   r   r   TritonSemantic.broadcast_tensors  s{    Is4y!A 55dGDMDq' "s4y1}%A 55dGDMDq' &}t}r   c               (    V ^8  d   QhRRRRRRRR/# )	r/   r  zSequence[tl.tensor]result_typeszSequence[tl.dtype]packr1   r2   zTuple[tl.tensor, ...]r3   )r4   s   "r   r5   r6     s2     i i&9 iI[ icf i.Cir   c           	     	n  a a S P                   ! V!  p\        V4      ^ 8  g   Q R4       hV Uu. uF.  qQ^ ,          P                  P                  VP                  4      NK0  	  ppS P
                  P                  V Uu. uF  qfP                  NK  	  upV Uu. uF  qUP                  S P
                  4      NK  	  upV4      oV! S4       SP                  4       '       g   Q h\        ;QJ d!    . VV 3R l\        V4       4       F  NK  	  5# ! VV 3R l\        V4       4       4      # u upi u upi u upi )r   z1map_elementwise must have at least 1 input tensorc              3  j   <"   T F(  w  rSP                  SP                  V4      V4      x  K*  	  R # 5ir*   )r'   r  )r  r  rG  elementwise_opr   s   &  r   r  1TritonSemantic.map_elementwise.<locals>.<genexpr>  s.     hPguqT[[!:!:1!=rBBPgs   03)r  r  ry   r   r   r(   create_map_elementwiser   r   r  r  r  )r   r  r  r  r  rG  r  r  s   f&&&&  @r   map_elementwiseTritonSemantic.map_elementwise  s    ''06{QS SSLXYLbq	66ryyALY<<%&v!XXv&.:;lXXdll#l;

 	.)$$&&&&uhPYZfPghuhuhPYZfPghhh Z&;s   4D(=D-#D2
c               (    V ^8  d   QhRRRRRRRR/# )r/   r   r   num_binsr1   r  r  r2   r3   )r4   s   "r   r5   r6     s0     	@ 	@x 	@3 	@>P 	@U] 	@r   c                	   \        VP                  4      ^8X  g   Q R4       hVP                  P                  4       '       g   Q R4       hVe]   V P	                  W1P                  4      pVP
                  P                  P                  4       '       g   \        R4      hVP                  pV P                  V P                  P                  VP                  W#4      \        P                  ! \        P                  V.4      4      # )r9   z histogram only supports 1D inputz%histogram only supports integer inputr  )r  r  rM   rj   r
  ry   r   r:  r:   r   r'   r(   create_histogramr<   r  r=   )r   r   r!  r  s   &&&&r   	histogramTritonSemantic.histogram  s    5;;1$H&HH${{!!##L%LL#,,T;;?D99##++-- !EFF;;D{{4<<88xV==H:>@ 	@r   c               $    V ^8  d   QhRRRRRR/# r/   rz   r   r  r  r2   r3   )r4   s   "r   r5   r6     s!      X y X r   c                	   \        ^\        VP                  4      4      \        V4      8w  d   \        R4      hVP                  P                  R\        P                  ! W!P                  P                  4       4      4       V# )r9   zAShape of input to multiple_of does not match the length of valuesztt.divisibility)	r  r  r  r:   r   set_attrr   	make_attrget_contextr   rz   r  s   &&&r   multiple_ofTritonSemantic.multiple_of  sY    q#agg,3v;.`aa	+R\\&((BVBVBX-YZr   c               $    V ^8  d   QhRRRRRR/# r'  r3   )r4   s   "r   r5   r6     s!       )  r   c                	    \        VP                  4      \        V4      8w  d   \        R 4      hVP                  P	                  R\
        P                  ! W!P                  P                  4       4      4       V# )zDShape of input to max_contiguous does not match the length of valuesztt.contiguityr  r  r:   r   r)  r   r*  r+  r,  s   &&&r   max_contiguousTritonSemantic.max_contiguous  sQ    qww<3v;&cdd	/2<<@T@T@V+WXr   c               $    V ^8  d   QhRRRRRR/# r'  r3   )r4   s   "r   r5   r6   #  s!      x  x r   c                	    \        VP                  4      \        V4      8w  d   \        R 4      hVP                  P	                  R\
        P                  ! W!P                  P                  4       4      4       V# )zCShape of input to max_constancy does not match the length of valuesztt.constancyr1  r,  s   &&&r   max_constancyTritonSemantic.max_constancy#  sQ    qww<3v;&bcc	.",,vxx?S?S?U*VWr   c                   V ^8  d   QhRR/# r/   r2   r   r3   )r4   s   "r   r5   r6   )  s     C Cx Cr   c                	r    V P                  V P                  P                  4       \        P                  4      # r*   )r'   r(   create_barrierr<   r  )r   s   &r   debug_barrierTritonSemantic.debug_barrier)  s$    {{4<<668"''BBr   c               (    V ^8  d   QhRRRRRRRR/# )	r/   prefixr  argsList[TensorTy]hexrY   r2   r   r3   )r4   s   "r   r5   r6   ,  s.     a a3 an a4 aH ar   c                	   VP                  R 4      '       g   V'       d
   VR ,          pVP                  R4      '       g   V'       d   VRR R,           p\        V4      ^8  d!   VP                  R 4      '       g
   R V,           pV Uu. uF  qDP                  NK  	  ppV Uu. uF  qDP                  P                  4       NK  	  ppV P                  V P                  P                  WWV4      \        P                  4      # u upi u upi ) r  Nr  )endswithr  
startswithr   rM   r   r'   r(   create_printr<   r  )r   r?  r@  rB  argnew_args	is_signeds   &&&&   r   device_printTritonSemantic.device_print,  s     s##cMFt$$CR[4'Fv;?6#4#4S#9#96\F*./$3JJ$/:>?$3YY,,.$	?{{4<<44V(VXZX_X_`` 0?s   D ""Dc               (    V ^8  d   QhRRRRRRRR/# )r/   r   r   r   r  r  r  r2   r3   )r4   s   "r   r5   r6   :  s0     R R( R R<N RS[ Rr   c                	"   V P                   P                  P                  '       g   R # Ve!   V P                  WP	                  V4      4      pV P                  V P                   P                  VP                  V4      \        P                  4      # r*   )
r(   r   debugr5  rJ  r'   create_assertr   r<   r  )r   r   r   r  s   &&&&r   r   TritonSemantic.device_assert:  sa    ||##)))88D))D/2D{{4<<55dkk3GQQr   c                   V ^8  d   QhRR/# r9  r3   )r4   s   "r   r5   r6   A  s     M Mh Mr   c                	    V P                  V P                  P                  VP                  4      \        P
                  4      # r*   )r'   r(   create_assumer   r<   r  )r   r   s   &&r   assumeTritonSemantic.assumeA  s*    {{4<<55dkkBBGGLLr   c                	   \        V\        4      '       d   \        P                  ! V4      p\        V\        P                  4      '       d   \        VP                  \
        4      '       d&   V P                  P                  VP                  4      # V'       dV   R	VP                  u;8:  d   R8  g   M Q RVP                   R24       hV P                  P                  VP                  4      # RVP                  u;8:  d   R
8  g   M Q RVP                   R24       hV P                  P                  VP                  4      # \        V\        P                  4      '       Ed   VP                  P                  ^8X  g   Q R4       hVP                  P                  4       '       g   Q R4       hVP                  \        P                  8w  d`   V'       dX   V P                  P                  VP                   V P                  P#                  4       VP                  P%                  4       4      # VP                  \        P                  8X  d   V'       g   Q R4       hVP                   # Q R\'        T4       24       h)r/   z@Block pointers only support 64 bit `shape/strides`, got a value z which is out of the rangezFBlock pointers only support 32 bit `offsets/block_shape`, got a value z*Expected a scalar in shape/strides/offsetsz8Expected an integer scalar type in shape/strides/offsetszzBlock 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: r   r   r   r   )rt   r1   r<   ru   r`   rY   r(   r4  rA  r  r'   r  rM   rj   r   r   r   get_int64_tyr   ry   )r   r  r  s   &&&r   _convert_elem_to_ir_value(TritonSemantic._convert_elem_to_ir_valueD  s   dC  <<%DdBLL))$**d++||,,TZZ883e3 J 8##'::,.H6J J3||--djj993e3 J 8##'::,.H6J J3||--djj99bii((::##q(V*VV(::$$&&b(bb&zzRXX%+||33DKKAZAZA\48JJ4L4L4NP Prxx'W W Wu;;XKDQUJ<XXur   c                	    \        VR 4      '       d!   V Uu. uF  q0P                  W24      NK  	  up# V P                  W4      .# u upi )r  )r  rY  )r   	list_liker  r  s   &&& r   r  $TritonSemantic._convert_to_ir_values^  sI    9j))R[\R[$224ER[\\..yFGG ]s   Ac                    V ^8  d   QhRRRR/# r/   baser   r2   r3   )r4   s   "r   r5   r6   c  s     $f $f8 $f]e $fr   c           	     	  a V P                  V4      pV P                  V4      pV P                  VR R7      pVP                  P                  4       '       d+   VP                  P                  P	                  4       '       d   \        R4      hVP                  P                  \        P                  8X  dJ   V P                  V\        P                  ! \        P                  VP                  P                  4      4      p\        SR4      '       g   S.oS Uu. uF1  p\        V\        P                  4      '       d   VP                  MTNK3  	  upo\         ;QJ d    R S 4       F  '       d   K   R M	  RM! R S 4       4      '       g   Q R4       h\        VR4      '       g   V.pV Uu. uF1  p\        V\        P                  4      '       d   VP                  MTNK3  	  pp\#        V4      \%        \'        \)        V4      4      4      8X  g   Q R4       h\         ;QJ d!    V3R lW#WF3 4       F  '       d   K   R M	  RM! V3R lW#WF3 4       4      '       g   Q R	4       hV P*                  P-                  VP.                  W#VSV4      pV P1                  V\        P                  ! \        P2                  ! VP                  P                  S4      4      4      # u upi u upi )
Fr  zMExpected `base` to be a pointer type (but not a block pointer type or others)r  c              3  ~   "   T F3  p\        V\        4      ;'       d    RTu;8*  ;'       d    R8  Mu x  K5  	  R# 5i)r/   Nr   r   )rt   r1   )r  r  s   & r   r  0TritonSemantic.make_block_ptr.<locals>.<genexpr>v  s1     \P[:dC(CCVt-C-Ce-CCP[s   ==TzGExpected a list of constant integers (`int32_t` range) in `block_shape`z<Expected a permutation of (0, 1, ..., len(order)-1) in orderc              3  R   <"   T F  p\        S4      \        V4      8H  x  K  	  R # 5ir*   )r  )r  r\  r  s   & r   r  rc    s!     hGg)3{#s9~5Ggs   $'zBExpected shape/strides/offsets/block_shape to have the same length)r  ry   r   r  r  r:   r<   r   r   r  r  r  r  rt   ru   r`   r  r  r  r  r  r(   create_make_block_ptrr   r'   r  )	r   r`  r  stridesr  r  orderr  r   s	   &&&&&f&  r   make_block_ptrTritonSemantic.make_block_ptrc  s<    **51,,W5,,W%,H yy!!TYY%9%9%B%B%D%Dlmm 99277*99T2??277DII<S<S#TUD {J//&-KZefZeRVZbll%C%CtzzMZefs\P[\sss\P[\\\ 	VU	V\ uj))GETYZTYDz$==4GTYZe}U3u:%6 77w9ww7 shX_GghssshX_Gghhh 	QP	Qh 33DKKQXZeglm{{62??2==AUAUWb3c#dee% g [s   7K+7Kc                    V ^8  d   QhRRRR/# r_  r3   )r4   s   "r   r5   r6     s     Y YH Y( Yr   c                	    V P                  VR R7      pV P                  V P                  P                  VP                  V4      VP
                  4      # r  )r  r'   r(   create_advancer   ry   )r   r`  r  s   &&&r   advanceTritonSemantic.advance  sC    ,,W%,H {{4<<66t{{GLdiiXXr   c               0    V ^8  d   QhRRRRRRRRRR	R
R/# )r/   r`  r   r  rA  rf  r  zList[tl.constexpr]rn  r  r2   ztl.tensor_descriptorr3   )r4   s   "r   r5   r6     sB     (B (B8 (BN (BUc (B,>(BPS(Bau(Br   c           	     	p   \        V4      p^Tu;8:  d   ^8:  g   M \        RV R24      h\        V4      V8w  d   \        RV R\        V4       24      h\        V4      V8w  d   \        RV R\        V4       24      h\        VP                  \        P
                  4      '       g   Q hVP                  P                  P                  ^,          p\        P                  ! VR,          4      pW,          ^8  d   \        RV RV R	W,           R
24      h\        P                  ! VR,          4      p	V	^8w  d   \        RV	 24      hV U
u. uF"  qP                  V
\        P                  4      NK$  	  pp
V U
u. uF6  qP                  \        P                  ! V
4      \        P                  4      NK8  	  pp
\        P                  ! V4      p\        VP                  \        P
                  4      '       g   Q h\        P                  ! VP                  P                  V4      pVP                  pVP                  P                  P!                  4       pV P#                  V4      pVP                  P                  P%                  4       '       d+   V\&        P(                  P*                  8X  d   \        R4      hV P,                  P/                  Y Uu. uF  qP                  NK  	  upV Uu. uF  qP                  NK  	  upWMV4      p\        P0                  ! VW#V4      # u up
i u up
i u upi u upi )r9   z Expected 1 <= ndim <= 5 but got z dimensionsz	Expected z strides but got zExpected block_shape to have z dimensions but got zRDescriptor block shape must have at least 16 bytes in the last dimension, but got z * z = z bytesz-Tensor descriptor last dim must be 1 but got z8Padding option `nan` is not supported for integer blocksr  )r  r:   rt   rM   r<   r  r  r(  r  r  r=   r   _unwrap_shapery   r  r   r   rp  rj   r   rk  rm  r(   create_make_tensor_descriptortensor_descriptor)r   r`  r  rf  r  rn  r  	elem_sizecontig_dim_sizelast_striderz   ry   base_handleis_signed_intro  r  r   s   &&&&&&           r   make_tensor_descriptor%TritonSemantic.make_tensor_descriptor  s   5zTQ?v[QRRw<4y.?G~NOO{t#<TFBVWZ[bWcVdeff$**boo6666JJ))<<A	11+b/B&+detduux  zC  yD  DG  HW  Hc  Gd  dj  k  --gbk:!L[MZ[[8=>1!!!RXX.>SZ[SZa##B$;$;A$>ISZ[ &&{3$))R__5555}}TYY11;?kk		,,::<--n=99&&((W8I8I8Q8Q-QWXX;;K\aIb\aWX((\aIbOV<Ww!XXw<WYd<CE ##FEDAA' ?[ Jc<Ws   (L$<<L)L.
/L3r+   )T)FFTFr*   )ri  )r   r    r!   r"   r<   r'   __annotations__langr   r?   rD   rU   rm   r|   rw   r   r   r   r   r   r   r   r   r   r   r  r  r"  r(  r   r5  r:  r@  rD  rJ  rP  rV  r[  r`  r   rH  rq  ry  r   r  r   r  r  r  rx   r  r  r  r  r  r  r  r  r  r  r
  r   r#  r?  r   rU  r[  rf  rp  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r$  r,  r7  r@  rW  r\  ra  rh  rn  rs  ry  r  r  r  r  r  r  rJ  r  r   r	  r  r  r  r$  r-  r2  r6  r<  rK  r   rU  rY  r  rh  rm  ry  r#   r3   r   r   r&   r&      s   YYFN&DO
QD05d	KF	@#J,$>>8"8]2>,8.9"9"	a\[\'&"
]]\)$/8888	8	8Zd Z$)2AhSfXV0

TW 2ptj$i9^		
,:xn	//nu@[
uuuuu$0**x8.$`W(u8+$)L$)L%yxy
<[yz0$m3tre(&u,x.K,i*	@CaRMY4H
$fLY(B (Br   r&   )
__future__r   r   r  typingr   r   r   r   r   r	   r
   r   triton.runtimer   _C.libtritonr   r  r   r<   r   r   	Exceptionr   r&   r3   r   r   <module>r     sY    "   J J J  !  CL:F	 F^BWX& ^Br   