+ j 0  ^RIHt^RIt^RIHt^RIHt^RIHt^RI H t H t H t ^RI t ^RI HtHtHtHtHtHtHtHt^RIHt^RIt^R IHt^R IHt^RIt^R IHt^R I H!t!H"t"H#t#H$t$]!R 4t%Rt&]PNt(RRlt)RRlt*RRlt+Rt,RRlt-]*RRl4t.!RR4t/!RR4t0!RR4t1!RR ]14t2!R!R"]04t3]3!^4t4R#t5R$t6R%t7!R&R']14t8]8t9!R(R)]84t:!R*R+]84t;!R,R-]14t<!R.R/]84t=]8!R04t>]8!R14t?]8!R24t@]8!R34tA]8!R44tB]8!R54tC]8!R64tD]8!R74tE]8!R84tF]8!R94tG]8!R:4tH]8!R;4tI]8!R<4tJ]8!R=4tK]8!R>4tL]8!R?4tM]8!R@4tN]8!RA4tO]8!RB4tP]:!]B4tQRCRDltR!RERF]04tSRRGltT!RHRI]04tU!RJRK4tV!RLRM]14tW!RNRO]04tX!RPRQ]W4tY!RRRS]X4tZ]!RRT7!RURV]144t[RWt\]'d^RXI]H^t^MRYt^]_!4t`RZta]^!R[R\7R]4tbR^tcR_R`ltdRateRbR[/RcRdlltfReRfltgRgRhlth]b!RiRj44tiR]inj]k!Rk4]in8]*RRll4tl]*RRml4tm]*RRnl4tnRo]! Rp2]nnoRqtpRrtq]*RRsl4tr]*RRtl4ts]+]*RuR/Rvl44tt]+]*RuR/RwRxll44tu]+]*RuR/Ryl44tv]*RRzl4tw]*RR{l4txRR|lty]+]*RR}R~ll44tz]+]*RuR/Rl44t{]+]*RRl44t|]+]*RR[RuRRR/Rl44t}Rt~]+]*RRl44t]+]*RRRll44t]*RRl4t]*RR[RR]OR3Rl4t]*RRl4t]*RRRll4t]*RRRll4t]+]*RRl44t]*RRRll4t])!R4]+]*RRl444t]*RRRll4tRRRllt]+]*]!RRR7RRl444t]+]*]!R4RRl444t]+]*]!R4RRl444t]+]*]!R4RRl444t]+]*]!R4RRl444t]+]*]!R4RRl444t]+]*]!R4RRl444t]+]*]!R4RRl444t]*RRl4t]*RRl4t]*RRRll4t]*RRRll4t]*RRRll4t]*](EP0R3RRll4t]*](EP0R3RRll4t]*](EP0R3RRll4tRRRllt]R4t]+]*RRl44t]*RRl4t]*RRl4tRRRllt]+]*RRl44t]+]*RRl44t]+]*RRl44t]*R^RuRRR/RRll4t]*RRl4t]*RRl4t]*RRl4t]*RRl4t]*RRl4t]*RRRRRRRR[RuR/RRll4t]*RRl4t]*RR[RuR/Rl4t]*RRl4t]*RRRll4t!RR]04t!RR]04t!RR]04tRRlt]*RRRll4tRtRt]_!4tRt]*R]RuR/Rl4t]*R]RuR/Rl4tR#)) annotationsN)warn)contextmanager)Enum)partialwrapscached_property)UnionCallableListSequenceTypeVarOptionalTuple TYPE_CHECKING) dataclass)knobs) JITCallable)ir)TRITON_MAX_TENSOR_NUMELvalidate_block_shapeget_primitive_bitwidth _tuple_createT__triton_builtin__TcNa\S\4'dV3Rl#VSnS#)z9If the result of this function is unused, throw an error.c<\VS4#N)must_use_result)fnxs&U/home/luca/screen-checking/.venv/lib/python3.14/site-packages/triton/language/core.py!must_use_result..s ?2q1) isinstancestr_must_use_result)r!ssf&r"rrs#!S12A Hr%c V^8dQhRRRR/#r rreturn)formats"r" __annotate__r0"sar%ca\S4'gQh\S4V3Rl4p\V\R4\P !S4VnV#)zMark a function as a builtin.cN<RV9g VR,f \R4hS!V/VB#) _semanticzeDid you forget to add @triton.jit ? (`_semantic` argument must be provided outside of JIT functions.) ValueErrorargskwargsr s*,r"wrapperbuiltin..wrapper&s: f ${(;(Cab b4"6""r%T)callablersetattrTRITON_BUILTINinspect signature)r r9sf r"builtinr@"sM B<<< 2Y##  G^T*))"-G Nr%c V^8dQhRRRR/#r+r.)r/s"r"r0r03s**!**r%c ,a\S4'gQh\P!S4p\VPP 4RR0, 4^8pSP 'gRSnS;P RSP RV'dRMR RSP RV'dR MR R 2 , unV3R lp\VPP44pV^,PR R 7V^&VPVR7pWSn WSnRSP R2Vn\S4'd\V\R4\\SP\!S\"4'd S4S#T4S#)aDecorator that adds this free function as a member fn on class tensor. When called as a member function on class tensor, the first argument to `fn` is `self`, i.e. the tensor object. If there are multiple decorators on a function, you probably want this one to be the highest one (i.e. furthest from the function's `def`), so it's applied last. Unfortunately you still need to add a type stub to the body of class tensor in order for pytype to know about it. r3 _generatorzb This function can also be called as a member function on :py:class:`tensor`, as :code:`x.(z...z)` instead of :code:`z(xz, ...z)`. c<S!V/VB#rr.r6s*,r"r9"_tensor_member_fn..wrapperMs4"6""r%selfname) parameterszForwards to :py:func:`z` free functionT)r;r>r?lenrKkeys__doc____name__listvaluesreplace __signature__ is_builtinr<r=tensorr&r)r orig_sighas_argsr9 new_paramsnew_sigsf r"_tensor_member_fnrZ3sX B<<<  $H8&&++-l0KKLqPH ::: JJ Qub9: KK=h7B78J # h))0023JqM))v)6JqM*5G#.r{{m?KGO"~~. FBKKz"k'B'BP IIPP Ir%c\V4p\V4^8Xd\V^,4V^,#V# \dT#i;i)z7Returns x[0] if x has one element and x[0] is iterable.)_unwrap_if_constexprrLiter TypeErrorr!s&r"_unwrap_iterabler``sPQA 1v{  1JQ4K H   H s: A A cV^8dQhRR/#r,r-boolr.)r/s"r"r0r0ws..d.r%c$\V\R4#)z-Is this a registered triton builtin function?F)getattrr=r s&r"rTrTws 2~u --r%c$VPV4#r) to_tensor)r!r3s&&r"rhrh|s   q !!r%c]tRt^tRtRtR#)constaf This class is used as a type annotation to mark pointers to constant data. The `store` function cannot be called with a pointer to const. Constness is part of the pointer type and the usual Triton type consistency rules apply. For example you cannot have a function that returns constant pointer in one return statement and non-constant pointer in another. r.N)rO __module__ __qualname____firstlineno__rN__static_attributes__r.r%r"rjrjs  r%rjc>]tRt^t$RtR]R&RRltRRltRtR #) base_valuezLBase class of values that exist in the triton IR (i.e. not constexprs). base_typetypec$V^8dQhRRRRRR/#r,builder ir.builderrJr'r-Noner.)r/s"r"r0base_value.__annotate__s!"""3"4"r%c \hrNotImplementedErrorrHrurJs&&&r" _set_namebase_value._set_name!!r%c V^8dQhRRRR/#r,handlesList[ir.value]r-rwr.)r/s"r"r0rxs"">"d"r%c \h)z^Flatten frontend value into a sequence of mlir handles, which are appended to the output list rzrHrs&&r" _flatten_irbase_value._flatten_irs "!r%r.N) rOrkrlrmrN__annotations__r}rrnr.r%r"rprps O"""r%rpcR]tRt^tRRltRRltRRltRRltR R ltR t R #) rqcV^8dQhRR/#rbr.)r/s"r"r0base_type.__annotate__sAAtAr%c \R4h)zTypes must implement __eq__rzrHothers&&r"__eq__base_type.__eq__s!"?@@r%cV^8dQhRR/#rbr.)r/s"r"r0rs##t#r%c W8X*#rr.rs&&r"__ne__base_type.__ne__ ""r%c$V^8dQhRRRRRR/#r,rrcursorintr-zTuple[base_value, int]r.)r/s"r"r0rs"""^"S"E["r%c \h)aBuild a frontend value with the current dtype, wrapping a list of existing handles. cursor is the index of the first handle relevant to this value, and the function should return the updated cursor position after any handles consumed by the created value. rzrHrrs&&&r" _unflatten_irbase_type._unflatten_irs "!r%cV^8dQhRR/#r,r-r'r.)r/s"r"r0rsSSSr%c 2\RVP 24h)zNYI: Type mangling for type )r{ __class__rHs&r"manglebase_type.mangles!$@@P"QRRr%c$V^8dQhRRRRRR/#r,rurvout List[ir.type]r-rwr.)r/s"r"r0rs!"""-"D"r%c \hrrzrHrurs&&&r"_flatten_ir_typesbase_type._flatten_ir_typesrr%r.N) rOrkrlrmrrrrrrnr.r%r"rqrqs"A#"S""r%rqcX]tRt^tRtRtRRltRtRRltRR lt R R lt R t R #)constexpr_typec WnR#rvaluerHrs&&r"__init__constexpr_type.__init__s r%c d\V\4;'dVPVP8H#r)r&rrrs&&r"rconstexpr_type.__eq__s$%0NNTZZ5;;5NNr%cV^8dQhRR/#rr.)r/s"r"r0constexpr_type.__annotate__s//#/r%c "RVP R2#)zconstexpr_type[]rrs&r"__repr__constexpr_type.__repr__s  A..r%c ,\VP4#r)hashrrs&r"__hash__constexpr_type.__hash__DJJr%cV^8dQhRR/#rr.)r/s"r"r0rsr%c \VPR4'dVPP4pM\VP4pRV 2#)rc)hasattrrrrepr)rHvals& r"rconstexpr_type.mangles? 4::x ( (**##%Ctzz"C3%yr%c$V^8dQhRRRRRR/#rr.)r/s"r"r0rs!-Dr%c R#rr.rs&&&r"r constexpr_type._flatten_ir_typesr%c$V^8dQhRRRRRR/#rr.)r/s"r"r0rs"--^-S-E[-r%c 0\VP4V3#r constexprrrs&&&r"rconstexpr_type._unflatten_irs$f,,r%rN) rOrkrlrmrrrrrrrrnr.r%r"rrs+O/ --r%rc4]tRt^tRtRtRRltRtRRltRR lt R t R t R t R t RtRtRtRtRtRtRtRtRtRtRtRtRtRtRtRtRtRtR t R!t!R"t"R#t#R$t$R%t%R&t&R't'R(t(R)t)R*t*R+t+R,t,R-t-R.t.R/t/R0t0R1t1R2#)3rzD This class is used to store a value that is known at compile-time. c z\V\4'dVPpK$Wn\V4VnR#r)r&rrrrrrs&&r"rconstexpr.__init__s) **KKE "5) r%cV^8dQhRR/#rr.)r/s"r"r0constexpr.__annotate__s**#*r%c "RVP R2#)z constexpr[rrrs&r"rconstexpr.__repr__sDJJdr%c R#rr.rs&&r"rconstexpr._flatten_irrr%c VP#rrrs&r" __index__constexpr.__index__ zzr%c L\VP\V4,4#rrrr\rs&&r"__add__constexpr.__add__&:5&AABBr%c L\\V4VP,4#rrr\rrs&&r"__radd__constexpr.__radd__-e4tzzABBr%c L\VP\V4, 4#rrrs&&r"__sub__constexpr.__sub__rr%c L\\V4VP, 4#rrrs&&r"__rsub__constexpr.__rsub__rr%c L\VP\V4,4#rrrs&&r"__mul__constexpr.__mul__rr%c L\VP\V4,4#rrrs&&r"__mod__constexpr.__mod__rr%c L\\V4VP,4#rrrs&&r"__rmul__constexpr.__rmul__rr%c L\VP\V4, 4#rrrs&&r" __truediv__constexpr.__truediv__rr%c L\\V4VP, 4#rrrs&&r" __rtruediv__constexpr.__rtruediv__rr%c L\VP\V4,4#rrrs&&r" __floordiv__constexpr.__floordiv__ ';E'BBCCr%c L\\V4VP,4#rrrs&&r" __rfloordiv__constexpr.__rfloordiv__s-e4 BCCr%c D\VP\V484#rrrs&&r"__gt__constexpr.__gt__&:5&AABBr%c D\\V4VP84#rrrs&&r"__rgt__constexpr.__rgt__-e4tzzABBr%c D\VP\V484#rrrs&&r"__ge__constexpr.__ge__';E'BBCCr%c D\\V4VP84#rrrs&&r"__rge__constexpr.__rge__-e4 BCCr%c D\VP\V484#rrrs&&r"__lt__constexpr.__lt__rr%c D\\V4VP84#rrrs&&r"__rlt__constexpr.__rlt__ rr%c D\VP\V48*4#rrrs&&r"__le__constexpr.__le__#rr%c D\\V4VP8*4#rrrs&&r"__rle__constexpr.__rle__&rr%c D\VP\V48H4#rrrs&&r"rconstexpr.__eq__)rr%c D\VP\V48g4#rrrs&&r"rconstexpr.__ne__,rr%c ,\VP4#r)rcrrs&r"__bool__constexpr.__bool__/rr%c .\VP)4#rrrs&r"__neg__constexpr.__neg__2$**%%r%c L\VP\V4,4#rrrs&&r"__and__constexpr.__and__5rr%c R\VP;'d \V44#rrrs&&r" logical_andconstexpr.logical_and8sCC(rr%c R\VP;'g \V44#rrrs&&r" logical_orconstexpr.logical_orAsBB';E'BCCr%c .\VP54#rrrs&r"__pos__constexpr.__pos__Dr1r%c .\VP(4#rrrs&r" __invert__constexpr.__invert__Gr1r%c L\VP\V4,4#rrrs&&r"__pow__constexpr.__pow__Js%9%%@@AAr%c L\\V4VP,4#rrrs&&r"__rpow__constexpr.__rpow__Ms-e4djj@AAr%c L\VP\V4, 4#rrrs&&r" __rshift__constexpr.__rshift__Prr%c L\VP\V4,4#rrrs&&r" __lshift__constexpr.__lshift__Srr%c 6\VP'*4#rrrs&r"__not__constexpr.__not__VsTZZ((r%c ,\VP4#r)r]rrs&r"__iter__constexpr.__iter__Yrr%c &VP!V/VB#rr)rHr7kwdss&*,r"__call__constexpr.__call__\szz4(4((r%c ZR\V44pVPP!V!#)c38"TFp\V4xK R#5irr\.0r!s& r" (constexpr.__getitem__..`sH1GA$Q''1G)_normalize_tupler __getitem__)rHr7s&*r"rfconstexpr.__getitem___s'H1A$1GHzz%%t,,r%)rrrN)2rOrkrlrmrNrrrr}rrrrrrrrrrrrr r rrrrrr"r%rrr,r/r3r6r9r<r?rBrErHrKrNrQrTrWr[rfrnr.r%r"rrs* *-CCCCCCCCCDDCCDDCCDDDD &CECCD&&BBDD) )-r%rc \V\4'dVUu.uFp\V4NK up#\V\P4'd%\ YUu.uFp\V4NK up4#\V\4'd0\ VUu.uFp\V4NK upVP 4#\V\4'd VP#T#uupiuupiuupir) r&rPr\builtinstuplerrrrr)or!s& r"r\r\gs!T123A$Q'33!X^^$$Q! D!Q!5a!8! DEE!Uq9q!*1-q9166BB I..1775A5 4 D9sCC$ C)c~\V4p\V\\P34'd \ V4pV#r)r\r&rPrirj)tnormalized_tuples& r"rereqs6+A."T8>>$:;; !12 r%c \V\4'dw\V\4'd_VPPP pVP V8d,\RVP RV RVP R24R#R#R#R#)zValue z exceeds the maximum bitwidth (z ) for type 'z)'. This may result in undefined behavior.N) r&rUrrrscalarprimitive_bitwidthrrdtype)r shift_valuebitwidths&& r"check_bit_widthruxs%  Z Y%G%G::$$77    ( **++J8*T`afalal`mnWX  )&H r%ca]tRtRt.R=Ot.R>Ot.R?Ot.R@OtR.t!RR] 4t !RR] 4t Rt Rt R tR tR tR tR tRtRtRtRtRtRtRtRtRtRtRtRtRtRt Rt!Rt"Rt#Rt$R t%R!t&R"t'R#t(])R$4t*])R%4t+])R&4t,])R'4t-])R(4t.R)R*lt/R+t0]1R,4t2R-R.lt3R/R0lt4R1t5R2t6]1R3R4l4t7R5t8R6R7lt9R8V3R9llt:R:R;lt;R//14 5## #"'"2"2"9"9D  $ 7 7D  U%% %"'"2"2";";D  $ 7 7D  U^^ #z!)*&%'"")*&%&"")*&%&")*&%'"#)*&%'")+&%'")*&%(")+&%(")+&%)""%EdV#LMM;$r%c RVP9#)fp8rIrs&r"is_fp8 dtype.is_fp8s !!r%c VPR8H#)rrIrs&r" is_fp8e4nvdtype.is_fp8e4nvyyI%%r%c VPR8H#)rrIrs&r" is_fp8e4b8dtype.is_fp8e4b8rr%c VPR8H#)rrIrs&r" is_fp8e4b15dtype.is_fp8e4b15yyJ&&r%c VPR8H#)rrIrs&r"is_fp8e5dtype.is_fp8e5yyG##r%c VPR8H#)rrIrs&r" is_fp8e5b16dtype.is_fp8e5b16rr%c VPR8H#)rrIrs&r"is_fp16 dtype.is_fp16yyF""r%c VPR8H#)rrIrs&r"is_bf16 dtype.is_bf16rr%c VPR8H#)rrIrs&r"is_fp32 dtype.is_fp32rr%c VPR8H#)rrIrs&r"is_fp64 dtype.is_fp64rr%c VPR8H#)int1rIrs&r"is_int1 dtype.is_int1rr%c VPR8H#)int8rIrs&r"is_int8 dtype.is_int8rr%c VPR8H#)int16rIrs&r"is_int16dtype.is_int16rr%c VPR8H#)int32rIrs&r"is_int32dtype.is_int32rr%c VPR8H#)int64rIrs&r"is_int64dtype.is_int64rr%c VPR8H#)uint8rIrs&r"is_uint8dtype.is_uint8rr%c VPR8H#)uint16rIrs&r" is_uint16dtype.is_uint16yyH$$r%c VPR8H#)uint32rIrs&r" is_uint32dtype.is_uint32rr%c VPR8H#)uint64rIrs&r" is_uint64dtype.is_uint64rr%c <VP\P9#r)rJrrrrs&r" is_floatingdtype.is_floatingsyyENN**r%c <VP\P9#r)rJrrSTANDARD_FP_TYPESrs&r"is_standard_floatingdtype.is_standard_floatingsyyE3333r%c <VP\P9#r)rJrrrrs&r" is_int_signeddtype.is_int_signedyyE,,,,r%c <VP\P9#r)rJrrrrs&r"is_int_unsigneddtype.is_int_unsignedrr%c fVP\P\P,9#r)rJrrrrrs&r"is_int dtype.is_ints"yyE,,u/?/????r%c "VP4#r)rrs&r"is_bool dtype.is_bool ||~r%c &VP4'd\PP#VP 4'd\PP #VP 4'gQh\PP#r)rrrrrrrrrrs&r"kind dtype.kind s] <<>>::%% % [[]]::&& &##%% %%::&& &r%c VP4'd"^VP^, ,^, #VP4'd^VP,^, #Qhr,rrrrs&r"get_int_max_valuedtype.get_int_max_valuesV     t((1,-1 1    ! !d'''!+ +ur%c VP4'd^VP^, ,)#VP4'd^#Qhrrrs&r"get_int_min_valuedtype.get_int_min_valuesA     ))A-.. .    ! !ur%c V\P\P,\P,\P,9#r)rrrrrr)type_strs&r"is_dtypedtype.is_dtype$s35++e.>.>>ORWRcRccccr%c \R4h)zNot implementedrr.r%r"is_void dtype.is_void(s,--r%c R#Fr.r.r%r"is_blockdtype.is_block,r%c R#r r.r.r%r"is_ptr dtype.is_ptr0r r%c R#r r.r.r%r"is_constdtype.is_const4r r%cV^8dQhRR/#rbr.)r/s"r"r0dtype.__annotate__8s''t'r%c z\V4p\V\4'gR#VPVP8H#r )r\r&rrrJrs&&r"r dtype.__eq__8s/$U+%''yyEJJ&&r%c .\VP34#r)rrJrs&r"rdtype.__hash__>sTYYM""r%c V#rr.rs&r"rp dtype.scalarA r%c$V^8dQhRRRRRR/#rr.)r/s"r"r0rEs!(((-(D(r%c FVPVPV44R#r)appendto_irrs&&&r"rdtype._flatten_ir_typesEs 4::g&'r%c V^8dQhRRRR/#)r,rurvr-zir.typer.)r/s"r"r0rHs$?$?Z$?G$?r%c VPPR4'd]\VR4'dKVPVPP9d&\ RV RVPP 24hVPR8XdVP 4#VPR8XdVP4#VPR9dVP4#VPR9dVP4#VPR9dVP4#VPR9dVP4#VPR8XdVP4#VPR8XdVP4#VPR8XdVP4#VPR 8XdVP4#VPR 8XdVP!4#VPR 8XdVP#4#VPR 8XdVP%4#VPR 8XdVP'4#VPR8XdVP)4#\ RV R24h)roptionsztype zB not supported in this architecture. The supported fp8 dtypes are rwrrrrrrrrrrzfail to convert z to ir type)rr)rr)rr)rr)rJ startswithrr$supported_fp8_dtypesr5 get_void_ty get_int1_ty get_int8_ty get_int16_ty get_int32_ty get_int64_ty get_fp8e5_tyget_fp8e5b16_tyget_fp8e4nv_tyget_fp8e4b8_tyget_fp8e4b15_ty get_half_ty get_bf16_ty get_float_ty get_double_tyrHrus&&r"r  dtype.to_irHs 99   & &w **tyy@d@d/d 5/AAHAeAe@f"hii 99 &&( ( YY& &&( ( YY+ +&&( ( YY- -'') ) YY- -'') ) YY- -'') ) YY' !'') ) YY* $**, , YY) #))+ + YY) #))+ + YY* $**, , YY& &&( ( YY& &&( ( YY& '') ) YY& ((* *+D6=>>r%c VP#rrIrs&r"__str__ dtype.__str__n yyr%c VPPR4'dRVPR,,#VPPR4'dRVPR,,#VP#)fpfloat:r,NNbfbfloat)rJr%rs&r" codegen_namedtype.codegen_nameqs\ 99   % %TYYr]* * YY ! !$ ' 'diim+ +99 r%cV^8dQhRR/#rr.)r/s"r"r0rzsr%c VP#)z"See cache_key_part() in triton.cc.rIrs&r"cache_key_partdtype.cache_key_partysyyr%c (RVP4 2#)z4Output of repr needs to be an evaluatable expressionztriton.language.)rArs&r"rdtype.__repr__~s!$"3"3"5!677r%c$V^8dQhRRRRRR/#rr.)r/s"r"r0rs"99^9S9E[9r%c 8\W,V4V^,3#)rUrs&&&r"rdtype._unflatten_irsgot,fqj88r%cV^8dQhRR/#rr.)r/s"r"r0rs     r%c T<VP4'dL\PPpVPV8XdRMRpV\ VP 4,#VP4'd \ V4#VP4'dR#\SV`)4#)iuV) rrrr|rzrr'rrrsuperr)rHrzprefixrs& r"r dtype.mangles} ;;==%%,,F //69SsFC 1 122 2     t9  <<>>w~r%cV^8dQhRR/#)r, element_tyrrr.)r/s"r"r0rs%r%c 6VP4'dQhV#r)r )rHrWs&&r"with_element_tydtype.with_element_tys==??""r%)rrrrrrJrq)rrrr)rrrrr) rrrrrrrrr)rrrr)>rOrkrlrmrrrrrrr|rrrrrrrrrrrrrrrrrrrrrrrrrrrrrr staticmethodrrr rrrrpropertyrprr r9rArErrrrYrn __classcell__rs@r"rrrrs4J@JfH8(KTt )NV"&&'$'######$$$$%%%+4--@'dd..' #($?L89   r%rrcr]tRtRtRRRlltRRltRtRtRtR t R R lt ] R 4t R Rlt RtR#) pointer_typeic$V^8dQhRRRRRR/#)r,rWrr address_spacerrjrcr.)r/s"r"r0pointer_type.__annotate__s&^^5^^^r%c \V4p\V\4'g#\R\ V4P R24hWnW nW0nV'gRV R2Vn R#RV R2Vn R#)zelement_ty has type `z`; expected `dtype`.zpointer<>zconst_pointer.s!=*Q#a&&*rdr:)rrrzip)rHkrs& r"rJtuple_type.namesx ;; !=$**!===C CSXXc$++tzz6RS6Rda!AaSz6RSTTWZZZSs.B c VP#rrIrs&r"r9tuple_type.__str__r;r%c ,\VP4#r)r]rrs&r"rWtuple_type.__iter__rr%c V^8dQhRRRR/#)r,rurvrrr.)r/s"r"r0tuple_type.__annotate__ s///-/r%c NVPFpVPW4K R#r)rr)rHrurtys&&& r"rtuple_type._flatten_ir_types s**B   .r%c V^8dQhRRRR/#)r,indexrr-rrr.)r/s"r"r0rs!!!!r%c (VPV,#r)r)rHrs&&r"rftuple_type.__getitem__szz%  r%c \V4\V4J;'d;VPVP8H;'dVPVP8H#r)rrrrrs&&r"rtuple_type.__eq__sBDzT%[(ffTZZ5;;-Fff4;;Z_ZfZfKffr%c$V^8dQhRRRRRR/#)r,rrrrr-zTuple[tuple, int]r.)r/s"r"r0rs"++^+S+EV+r%c .pVPF'pVPW4wrRVPV4K) \W04V3#r)rrrrj)rHrrrQrrs&&& r"rtuple_type._unflatten_irsE**B,,W=ME MM% V"F**r%c bRRPRVP44,R,#)rrc3@"TFqP4xK R#5ir)r)rars& r"rb$tuple_type.mangle..s?JbiikkJs)rrrs&r"rtuple_type.mangles%SXX?DJJ???#EEr%)rrr)rOrkrlrmrr rJr9rWrrfrrrrnr.r%r"rrsB[[  /!g+Fr%rc]tRtRtRtRtR#) slice_typei c RVnR#)rNrIrs&r"rslice_type.__init__"s   r%rINrOrkrlrmrrnr.r%r"rr s!r%rrwrrrrrrrrrrrrrrrrrrc$V^8dQhRRRRRR/#)r,rtrsignedrcr-rrr.)r/s"r"r0r0>s&UUCUU%Ur%cV^8Xd\#V^8XdV'd\#V^8XdV'g\#V^8XdV'd\#V^8XdV'g\#V^ 8XdV'd\ #V^ 8XdV'g\ #V^@8XdV'd\#V^@8XdV'g\#\RV RV 24h)rLzUnsupported bitwidth z and signedness ) rrrrrrrrrr5)rtrs&&r" get_int_dtyper>s1} Q6 Qv RF R RF R RF R 0 :J6(STTr%ca]tRtRtRtRV3RlltRRltRRltR R lt] RR l4t ] RR l4t ] RRl4t ] RRl4t ] RRl4t] RRl4t] RRl4t] RRl4t] RRl4t] RRl4t] RRl4t] RRl4t] RRl4t] RRl4t] RRl4t] RRl4t] RRl4t] RRl4t] RRl4t] RRl4t] RR l4t] RR!l4t] RR"l4t ] RR#l4t!] RR$l4t"] RR%l4t#] RR&l4t$] RR'l4t%] RR(l4t&] RR)l4t'] RR*l4t(] RR+l4t)] RR,l4t*] RR-l4t+] RR.l4t,] RR/l4t-] RR0l4t.] RR1l4t/] RR2l4t0] RR3l4t1]2R44t3] RR5R6ll4t4R7R8lt5R9R:lt6R;R<lt7R=R>lt8R?R@lt9RARBlt:RCRDlt;RRERFlltRRKRLllt?RRMRNllt@RRORPlltARRQRRlltBRRSRTlltCRRURVlltDRRWRXlltERRYRZlltFR[R\ltGR]R^ltHR_R`ltIRaRbltJRcRdltKReRfltLRgRhltMRRiRjlltNRRkRllltORmRnltPRoRpltQRqRrltRRsRtltSRRuRvlltTRwRxltURRyRzlltVRR{R|lltWRR}R~lltXRRRlltYRRRlltZRRRllt[RRRllt\RRRllt]RRRllt^R ]_3RRllt`RRRlltaRtbV;tc#)rUiZaRepresents an N-dimensional array of values or pointers. :code:`tensor` is the fundamental data structure in Triton programs. Most functions in :py:mod:`triton.language` operate on and return tensors. Most of the named member functions here are duplicates of the free functions in :code:`triton.language`. For example, :code:`triton.language.sqrt(x)` is equivalent to :code:`x.sqrt()`. :code:`tensor` also defines most of the magic/dunder methods, so you can write :code:`x+y`, :code:`x << 2`, etc. .. rubric:: Constructors .. For some reason Sphinx includes __init__ before printing the full table of methods. Not what I want, but I can't figure out how to fix it. Give it its own section so it looks intentional. :) cV^8dQhRR/#)r,rrrrr.)r/s"r"r0tensor.__annotate__ns ? ?U ?r%c |<\SV`4WnVP4'd VPMRVn\ \ P!VP44VnW n VPVn \VPUu.uFp\ V4NK up4VnR#uupi)Not called by user code.Nr.) rSrhandler rrmathprodrrrrprrrj)rHrrrr)rs&&& r"rtensor.__init__nsy  #'==??TZZ tyy45  [[ $**=*QIaL*=> =sB9c$V^8dQhRRRRRR/#rtr.)r/s"r"r0r{&RRR3R4Rr%c VPPVPW PP444R#rrset_loccreate_name_locget_locr|s&&&r"r}tensor._set_name{, G33D++:M:M:OPQr%c V^8dQhRRRR/#rr.)r/s"r"r0r~$$>$d$r%c <VPVP4R#rrrrs&&r"rtensor._flatten_ir~t{{#r%cV^8dQhRR/#rr.)r/s"r"r0rsSSSr%c \VP4R,RPRVP44,R,#)rrc38"TFp\V4xK R#5irr)rar)s& r"rb!tensor.__str__..s0LAQrdr)r'rrrrrs&r"r9tensor.__str__s24::$tyy0L0L'LLsRRr%Nc \WRVR7#T)sanitize_overflowr3addrHrr3s&&&r"rtensor.__add__4$)LLr%c \WRVR7#rrrs&&&r"rtensor.__radd__5$)LLr%c \WRVR7#rsubrs&&&r"rtensor.__sub__rr%c \WRVR7#rrrs&&&r"rtensor.__rsub__rr%c \WRVR7#rmulrs&&&r"rtensor.__mul__rr%c \WRVR7#rr rs&&&r"rtensor.__rmul__rr%c :\V4pVPW4#rr\truedivrs&&&r"rtensor.__truediv__s$U+  --r%c :\V4pVPW4#rrrs&&&r"rtensor.__rtruediv__s$U+  --r%c :\V4pVPW4#rr\floordivrs&&&r"rtensor.__floordiv__s$U+!!$..r%c :\V4pVPW4#rrrs&&&r"r tensor.__rfloordiv__s$U+!!%..r%c :\V4pVPW4#rr\modrs&&&r"rtensor.__mod__$U+}}T))r%c :\V4pVPW4#rrrs&&&r"__rmod__tensor.__rmod__$U+}}U))r%c $VPV4#r)minusrHr3s&&r"r/tensor.__neg__st$$r%c $VPV4#r)invertr(s&&r"rEtensor.__invert__s%%r%c :\V4pVPW4#rr\and_rs&&&r"r3tensor.__and__$U+~~d**r%c :\V4pVPW4#rr.rs&&&r"__rand__tensor.__rand__$U+~~e**r%c :\V4pVPW4#rr\or_rs&&&r"r9 tensor.__or__r!r%c :\V4pVPW4#rr7rs&&&r"__ror__tensor.__ror__r%r%c :\V4pVPW4#rr\xor_rs&&&r"r<tensor.__xor__r1r%c :\V4pVPW4#rr>rs&&&r"__rxor__tensor.__rxor__r5r%c P\W4\V4pVPW4#rrur\shlrs&&&r"rQtensor.__lshift__s#$$U+}}T))r%c P\W4\V4pVPW4#rrErs&&&r" __rlshift__tensor.__rlshift__s#$$U+}}U))r%c \W4\V4pVPP4'dVP W4#VP W4#rrur\rrrashrlshrrs&&&r"rNtensor.__rshift__sE$$U+ :: # # % %>>$. .>>$. .r%c \W4\V4pVPP4'dVP W4#VP W4#rrLrs&&&r" __rrshift__tensor.__rrshift__sE$$U+ :: # # % %>>%. .>>%. .r%c FVPV4pVPW4#rrh greater_thanrs&&&r"r  tensor.__gt__s###E*%%d22r%c FVPV4pVPW4#rrTrs&&&r"rtensor.__rgt__s###E*%%e22r%c FVPV4pVPW4#rrh greater_equalrs&&&r"r tensor.__ge__s###E*&&t33r%c FVPV4pVPW4#rrZrs&&&r"rtensor.__rge__s###E*&&u33r%c FVPV4pVPW4#rrh less_thanrs&&&r"r tensor.__lt__###E*""4//r%c FVPV4pVPW4#rr`rs&&&r"rtensor.__rlt__###E*""5//r%c FVPV4pVPW4#rrh less_equalrs&&&r"r" tensor.__le__$###E*##D00r%c FVPV4pVPW4#rrhrs&&&r"r%tensor.__rle__)s###E*##E00r%c FVPV4pVPW4#rrhequalrs&&&r"r tensor.__eq__/s!##E*t++r%c FVPV4pVPW4#rrors&&&r"__req__tensor.__req__4s!##E*u++r%c FVPV4pVPW4#rrh not_equalrs&&&r"r tensor.__ne__9rcr%c FVPV4pVPW4#rrvrs&&&r"__rne__tensor.__rne__>rfr%c FVPV4pVPW4#r)rhr6rs&&&r"r6tensor.logical_andCs###E*$$T11r%c FVPV4pVPW4#r)rhr?rs&&&r"r?tensor.logical_orHrkr%c $VPV4#r)not_r(s&&r"rTtensor.__not__Os~~d##r%c z\V\P\\34'gVfV.p\V\4'd VP pTp\ V4FwrE\V4fVPW44pK%\V\P\34'd\;QJd=RVPVPVP34F 'dK RM2 RM.!RVPVPVP344'dK\RV 24h V#)Nc3<"TFp\V4RJxK R#5irr_)raargs& r"rb%tensor.__getitem__..]s!A_A]#(-5A]sFTzunsupported tensor index: )r&rislicerrjrQ enumerater\ expand_dimsallstartstopstepr5)rHslicesr3retdimsls&&& r"rftensor.__getitem__Ss fx~~ui@ A AV^XF fe $ $]]F (GC#B'/++C5B 788SSA_BD((BGGUWU\U\A]A_SSSA_BD((BGGUWU\U\A]A_>_>_ #=bT!BCC) r%c QR4h)zTransposes a 2D tensor.z0Transposition must be created by the AST Visitorr.rs&r"rtensor.Tds IHHur%c$V^8dQhRRRRRR/#r,rrfp_downcast_roundingz Optional[str]bitcastrcr.)r/s"r"r0rjs'UUU]UTXUr%c \WW#VR7#)z# Alias for :py:func:`tensor.cast`. r3)cast)rHrrrrr3s&&&&&r"to tensor.tois D!5)TTr%cV^8dQhRR/#r,r-rUr.)r/s"r"r0rus  f r%c R#rr.rHrs&*r" broadcast_totensor.broadcast_tou r%cV^8dQhRR/#rr.)r/s"r"r0rx  f r%c R#rr.rHdimss&*r"trans tensor.transxrr%cV^8dQhRR/#rr.)r/s"r"r0r{   r%c R#rr.rs&*r"permutetensor.permute{rr%cV^8dQhRR/#r,r-ztuple[tensor, tensor]r.)r/s"r"r0r~s  , r%c R#rr.rs&r"split tensor.split~rr%cV^8dQhRR/#rr.)r/s"r"r0rrr%c R#rr.rs&*r"view tensor.viewrr%cV^8dQhRR/#rr.)r/s"r"r0rs   r%c R#rr.rs&*r"reshapetensor.reshaperr%cV^8dQhRR/#rr.)r/s"r"r0rs  6 r%c R#rr.)rHaxiss&&r"rtensor.expand_dimsrr%cV^8dQhRR/#rr.)r/s"r"r0rs  v r%c R#rr.)rHrrrrs&&&&r"r tensor.castrr%cV^8dQhRR/#rr.)r/s"r"r0rs  ci r%c R#rr.)rHrmaskboundary_checkcache_modifiereviction_policys&&&&&&r"store tensor.storerr%cV^8dQhRR/#rr.)r/s"r"r0rs  & r%c R#rr.)rHoffsetss&&r"advancetensor.advancerr%cV^8dQhRR/#rr.)r/s"r"r0rs  F r%c R#rr.)rHcmprsemscopes&&&&&r" atomic_castensor.atomic_casrr%cV^8dQhRR/#rr.)r/s"r"r0r  6 r%c R#rr.rHrrrrs&&&&&r" atomic_xchgtensor.atomic_xchgrr%cV^8dQhRR/#rr.)r/s"r"r0r  & r%c R#rr.rs&&&&&r" atomic_addtensor.atomic_addrr%cV^8dQhRR/#rr.)r/s"r"r0rrr%c R#rr.rs&&&&&r" atomic_maxtensor.atomic_maxrr%cV^8dQhRR/#rr.)r/s"r"r0rrr%c R#rr.rs&&&&&r" atomic_mintensor.atomic_minrr%cV^8dQhRR/#rr.)r/s"r"r0rrr%c R#rr.rs&&&&&r" atomic_andtensor.atomic_andrr%cV^8dQhRR/#rr.)r/s"r"r0rs   r%c R#rr.rs&&&&&r" atomic_ortensor.atomic_orrr%cV^8dQhRR/#rr.)r/s"r"r0rrr%c R#rr.rs&&&&&r" atomic_xortensor.atomic_xorrr%cV^8dQhRR/#rr.)r/s"r"r0r  V r%c R#rr.rs&r"exp tensor.exprr%cV^8dQhRR/#rr.)r/s"r"r0rrr%c R#rr.rs&r"log tensor.logrr%cV^8dQhRR/#rr.)r/s"r"r0rrr%c R#rr.rs&r"cos tensor.cosrr%cV^8dQhRR/#rr.)r/s"r"r0rrr%c R#rr.rs&r"sin tensor.sinrr%cV^8dQhRR/#rr.)r/s"r"r0rs  f r%c R#rr.rs&r"sqrt tensor.sqrtrr%cV^8dQhRR/#rr.)r/s"r"r0r  v r%c R#rr.rs&r"rsqrt tensor.rsqrtrr%cV^8dQhRR/#rr.)r/s"r"r0rrr%c R#rr.rs&r"abs tensor.absrr%cV^8dQhRR/#rr.)r/s"r"r0rs  6 r%c R#rr.)rHr combine_fn keep_dimss&&&&r"reduce tensor.reducerr%cV^8dQhRR/#rr.)r/s"r"r0rrr%c R#rr.)rHrrreverses&&&&r"associative_scantensor.associative_scanrr%cV^8dQhRR/#rr.)r/s"r"r0rs  v r%c R#rr.)rHindicesrs&&&r"gather tensor.gatherrr%cV^8dQhRR/#rr.)r/s"r"r0rs  V r%c R#rr.)rHnum_binss&&r" histogramtensor.histogramrr%cV^8dQhRR/#rr.)r/s"r"r0rs  6 r%c R#rr.)rHdivs&&r"cdiv tensor.cdivrr%cV^8dQhRR/#rr.)r/s"r"r0rs   r%c R#rr.rs&r"sigmoidtensor.sigmoidrr%cV^8dQhRR/#rr.)r/s"r"r0rs   r%c R#rr.)rHrr ieee_roundings&&&&r"softmaxtensor.softmaxrr%cV^8dQhRR/#rr.)r/s"r"r0rrr%c R#rr.rs&r"ravel tensor.ravelrr%cV^8dQhRR/#rr.)r/s"r"r0r  kq r%c R#rr.rHrreturn_indicesreturn_indices_tie_break_leftrs&&&&&r"max tensor.maxrr%cV^8dQhRR/#rr.)r/s"r"r0r  F r%c R#rr.rHrtie_break_leftrs&&&&r"argmax tensor.argmaxrr%cV^8dQhRR/#rr.)r/s"r"r0rr>r%c R#rr.r@s&&&&&r"min tensor.minrr%cV^8dQhRR/#rr.)r/s"r"r0rrFr%c R#rr.rHs&&&&r"argmin tensor.argminrr%cV^8dQhRR/#rr.)r/s"r"r0rs  V r%c R#rr.)rHrrrrs&&&&r"sum tensor.sumrr%cV^8dQhRR/#rr.)r/s"r"r0rs  V r%c R#rr.rHrrs&&&r"xor_sumtensor.xor_sumrr%cV^8dQhRR/#rr.)r/s"r"r0rs  v r%c R#rr.rZs&&&r" reduce_ortensor.reduce_orrr%cV^8dQhRR/#rr.)r/s"r"r0rs  v r%c R#rr.rHrrs&&&r"cumsum tensor.cumsumrr%cV^8dQhRR/#rr.)r/s"r"r0rs   r%c R#rr.rcs&&&r"cumprodtensor.cumprodrr%c$V^8dQhRRRRRR/#)r,rr descendingr-rUr.)r/s"r"r0rs"   i RX r%c R#rr.)rHrrks&&&r"sort tensor.sortrr%cV^8dQhRR/#rr.)r/s"r"r0rrr%c R#rr.)rHrs&&r"flip tensor.fliprr%)rrrrrrrrNFN)NF)Nr.rDrDNNNNNr )NFF)NFTF)TF)rF)drOrkrlrmrNrr}rr9r@rrrrrrrrrr rr#r/rEr3r3r9r;r<rBrQrIrNrQr rrrrrr"r%rrsrrzr6r?rTrfr\rrrrrrrrrrrrrrrrrrrrrrrrr rrrrr$r)r.r2r7r;rCrJrNrRrVr[r_rdrh CONSTEXPR_0rmrqrnr]r^s@r"rUrUZsc& ? ?R$S M M M M M M M M M M M M . . . . / / / / * * * *  % % & &  + + + + * * * * + + + + * *  * *  / / / / 3 3 3 3  4 4 4 4  0 0 0 0  1 1 1 1  , , , , 0 0 0 0 2 2 1 1  $ $   II U U                                          %)+   r%rUc \VUu.uF;p\V\\\34'd \ V4M VP NK= upV4#uupir)rr&rr>rrrrr)rQrr!s&& r"_type_for_tuple_valuesrxsJ djkdj_`Jq3u:M,N,N~a(TUTZTZZdjkms ttksAAc]tRtRtRRRlltRRltRtRtR tR t R t R t R t Rt RtRRltRRltRtRtR#)rjiNc V^8dQhRRRR/#)r,r7r rrzOptional[tuple_type]r.)r/s"r"r0tuple.__annotate__s<>#:;; ;;SYYsxx%@AB Br%c VPPpVeW9d\RV 24hVPVP V4,#)Nz 'tuple' object has no attribute )rrrAttributeErrorrQr)rHrJrs&& r" __getattr__tuple.__getattr__sF!! >T/ #CD6!JK K{{6<<-..r%c \V4p\V\4'gQhW PV&\ VPVP P 4VnR#r)r\r&rrQrxrrr)rHrrs&&&r"_setitemtuple._setitemsG"3'#s#### C*4;; 8H8HI r%c d\V4p\VPVP,4#r)rerjrQrs&&r"r tuple.__add__s$ 'T[[5<</00r%c ~\V\4'gQh\VPVP,4#r)r&rrjrQrrs&&r"r tuple.__mul__$s-%++++T[[5;;.//r%c \\V4p\VPVP8H4#r)rerrQrs&&r"r tuple.__eq__(s$ ' 455r%c T\\P!VP44#r)rrirjrQrs&r"rtuple.__hash__,sHNN4;;/00r%c h\VPUu.uFp\V4NK up4#uupir)r'rQ)rHr!s& r"r9 tuple.__str__/s'DKK0KqCFK0110s/c ,\VP4#r)r]rQrs&r"rWtuple.__iter__2sDKK  r%c ,\VP4#r)rLrQrs&r"__len__ tuple.__len__5s4;;r%c$V^8dQhRRRRRR/#rtr.)r/s"r"r0r{8s!4443444r%c VPPpVe7\W0P4FwrEVP W RV 24K R#\ VP4FwreVP W RV 24K R#)N.)rrrrrQr}r)rHrurJrfieldrrPs&&& r"r}tuple._set_name8sp!!   4 GvQug%675"$++. GvQqc]3/r%cV^8dQhRR/#)r,rrr.)r/s"r"r0r{As##>#r%c NVPFpVPV4K R#r)rQr)rHrrs&& r"rtuple._flatten_irAsA MM' "r%c NRRPRVP44 R2#)rErc38"TFp\V4xK R#5ir)rr`s& r"rb!tuple.__repr__..Fs:kT!WWkrd))rrQrs&r"rtuple.__repr__Es%499:dkk::;1==r%)rrrQr)rOrkrlrmrrfrrrrrrr9rWrr}rrrnr.r%r"rjrjsJ<C/J 1 0612! 4#>r%rjc]tRtRtRtRtR#)riIc HWnW nW0n\4VnR#r)rrrrrr)rHrrrs&&&&r"rslice.__init__Ks   L r%)rrrrrNrr.r%r"rrIs!r%rcj]tRtRtRRltRRltRRltRR ltR R ltR R lt RRlt Rt R#)tensor_descriptor_base_typeiRcV^8dQhRR/#r,rr.)r/s"r"r0(tensor_descriptor_base_type.__annotate__Ts%%:%r%c WnR#rr)rHrs&&r"r$tensor_descriptor_base_type.__init__Ts$r%c$V^8dQhRRRRRR/#r,rrrrr-z"Tuple[tensor_descriptor_base, int]r.)r/s"r"r0rWs"!!^!S!Eg!r%c N\W,VP4pW2^,3#rK)tensor_descriptor_baser)rHrrrs&&& r"r)tensor_descriptor_base_type._unflatten_irWs!&wHqj  r%c$V^8dQhRRRRRR/#rr.)r/s"r"r0r[s&eee-eDer%c VPPP4pVPVP VPP V4V44R#r)rrWrrcreate_tensor_descriptor_typer )rHrur is_signeds&&& r"r-tensor_descriptor_base_type._flatten_ir_types[sCOO..<<>  7889N9Nw9WYbcdr%cV^8dQhRR/#rr.)r/s"r"r0r_s777r%c "RVP R2#)ztensor_descriptorEvEbhEr%c &VPWV4#r)descriptor_atomic_addrs&&&&r"r!tensor_descriptor_base.atomic_add..tGDDr%c$V^8dQhRRRRRR/#rr.)r/s"r"r0rrr%c &VPWV4#r)descriptor_atomic_minrs&&&&r"r!tensor_descriptor_base.atomic_minrr%c$V^8dQhRRRRRR/#rr.)r/s"r"r0rrr%c &VPWV4#r)descriptor_atomic_maxrs&&&&r"r!tensor_descriptor_base.atomic_maxrr%c$V^8dQhRRRRRR/#rr.)r/s"r"r0rrr%c &VPWV4#r)descriptor_atomic_andrs&&&&r"r!tensor_descriptor_base.atomic_andrr%c$V^8dQhRRRRRR/#rr.)r/s"r"r0rs(DD!=DfDagDr%c &VPWV4#r)descriptor_atomic_orrs&&&&r"r tensor_descriptor_base.atomic_ors--d7CCr%c$V^8dQhRRRRRR/#rr.)r/s"r"r0rrr%c &VPWV4#r)descriptor_atomic_xorrs&&&&r"r!tensor_descriptor_base.atomic_xorrr%r3cV^8dQhRR/#rr.)r/s"r"r0rsNNvNr%c \V4^8XgQR\V4 24hV^,pV^,pVPWVRR4#)z)Gather multiple descriptors worth of dataz5descriptor gather only supports 2D indexing, but got rD)rLdescriptor_gather)rHr3r7 x_offsetsy_offsets&$* r"r$tensor_descriptor_base.gathersR4yA~b!VWZ[_W`Vabb~G 7**4Hb"MMr%cV^8dQhRR/#rr.)r/s"r"r0rsNNvNr%c \V4^8XgQR\V4 24hV^,pV^,pVPWWE4#)z*Scatter multiple descriptors worth of dataz6descriptor scatter only supports 2D indexing, but got )rLdescriptor_scatter)rHrr3r7rrs&&$* r"scattertensor_descriptor_base.scattersN4yA~c!WX[\`XaWbcc~G 7++DMMr%)rrrr)rOrkrlrmrNrr}rr\rrrrr9r@rrrrrrrrr$r rnr]r^s@r"rrosJ<<R$$$**// @ @ @ @ E E E E E E E E D D E E NdN N NdN Nr%rcRa]tRtRtRRltRRltRV3RlltV3RltR tV;t #) tensor_descriptor_typeic$V^8dQhRRRRRR/#)r,r shape_typer strides_typer.)r/s"r"r0#tensor_descriptor_type.__annotate__s")):):)U_)r%c *WnW nW0nR#rrrr)rHrrrs&&&&r"rtensor_descriptor_type.__init__s$$(r%c$V^8dQhRRRRRR/#rr.)r/s"r"r0rs"^SEgr%c W,pV^, pVPPW4wrBVPPW4wrRVPpVPp\ W4WPP 4pWb3#rK)rrrrQtensor_descriptorr)rHrrrrstridesrs&&& r"r$tensor_descriptor_type._unflatten_irsj! 55gF ++99'J ..!&//J}r%c$V^8dQhRRRRRR/#rr.)r/s"r"r0rs!:::-:D:r%c <\SV`W4VPPW4VPPW4R#r)rSrrr)rHrurrs&&&r"r(tensor_descriptor_type._flatten_ir_typess6 !'/ ))'7 ++G9r%c <\SV`V4;'d;VPVP8H;'dVPVP8H#r)rSrrr)rHrrs&&r"rtensor_descriptor_type.__eq__sWw~e$kk$//U=M=M*MkkTXTeTeW\WiWiUj kr%r) rOrkrlrmrrrrrnr]r^s@r"rrs#) :: kkr%rcVa]tRtRtRtRV3RlltRV3RlltRV3RlltR tV;t #) riz9A descriptor representing a tensor in global memory. c$V^8dQhRRRRRR/#)r,r List[tensor]rrr.)r/s"r"r0tensor_descriptor.__annotate__s"   l  \  Wa  r%c <\SV`W4\V4Vn\V4Vn\ VVPP VPP R7VnR#)r)rrN)rSrrjrrrrr)rHrrrrrs&&&&&r"rtensor_descriptor.__init__sO ,5\ W~ * zz**  r%c$V^8dQhRRRRRR/#rtr.)r/s"r"r0r%s!:::3:4:r%c <\SV`W4VPPWR,4VPPWR,4R#)z.shapez.strideN)rSr}rr)rHrurJrs&&&r"r}tensor_descriptor._set_names= '( WXo6 wy(89r%c V^8dQhRRRR/#rr.)r/s"r"r0r%s**>*d*r%c <\SV`V4VPPV4VPPV4R#r)rSrrr)rHrrs&&r"rtensor_descriptor._flatten_irs4 G$ w'   )r%)rrrr) rOrkrlrmrNrr}rrnr]r^s@r"rrs&   :: **r%r)frozencT]tRtRt$RtR]R&R]R&RRltR R ltR R ltR t R#)_aggregate_typei zA generic base type for all Triton aggregate types. This class contains a reference to the original user-defined Python class and a list of class fields with their Triton types. rrbase_clszList[Tuple[str, base_type]]rc$V^8dQhRRRRRR/#)r,rrrrr-zTuple[ir.value, int]r.)r/s"r"r0_aggregate_type.__annotate__s"  ^ S EY r%c VPP4pVPF$wrEVPW4wrb\ W4V4K& W23#r)r1 _get_instancerrr<)rHrrinstancerJrrs&&& r"r_aggregate_type._unflatten_irsK==..0 HD,,W=ME HE *$r%c$V^8dQhRRRRRR/#rr.)r/s"r"r0r3s!///-/D/r%c RVPFwr4VPW4K R#r)rr)rHrurrJrs&&& r"r!_aggregate_type._flatten_ir_typess  HD   .$r%cV^8dQhRR/#rr.)r/s"r"r0r3s...r%c VPP RVPP 2pVPUUu.uFwrVP 4NK pppX RRP V4 R2#uuppi)rrrre)r1rkrlrrr)rHrJrrs& r"r_aggregate_type.manglesj--**+1T]]-G-G,HI04 < *4"))+ <q6*+1--=sA4r.N) rOrkrlrmrNrrrrrnr.r%r"r0r0 s( N '' /..r%r0c\V\4'd^RIHpV!V\4#\V\ P4'd[\ P!RV44p\ VRR4p\VUu.uFqDPNK upV4p\W%4#\V\4'dV#\V4#uupi)r)_apply_to_tuple_valuesc38"TFp\V4xK R#5ir)_wrap_init_args)rarPs& r"rb"_wrap_init_args..*s ?Q!3!3Qrd_fieldsN) r&rjtriton.compiler.code_generatorr?rArirerrrrpr)r!r?wrappedrrrs& r"rArA%s!UI%a99!X^^$$.. ?Q ??It, 1A16 :W!!!Z   Q< 2s=C)dataclass_transformc R#)cV#rr.)objs&r"r#%dataclass_transform..8s3r%r.r8s,r"rFrF7sr%cv/p/p\VPR,4F{pV\Jg V\JdK\ VRR4'g\ RV 24hVP \P!V44VP \ VR/44K} \P!V4p\P!V4pVFfpWV,W&VPPV\4pV\JdK9\V4'd\V\4'gKbWrV&Kh W3#)rLNN__triton_aggregate__FzOAggregates can only inherit from other aggregates, but got non-aggregate base: __aggregate_defaults__)reversed__mro__rpobjectrer^updater>get_annotationstypingget_type_hints__dict__get_AGGREGATE_MISSINGr;r&)clsall_annotations all_defaultsbase own_nameshintsrJrs& r"_resolve_aggregate_fieldsr`>s OLR) :  t3U;;mnrmstu uw66t<=GD*BBGH *'',I  ! !# &E % llt%78 $ $ }} 3 ; ;!$   ((r%F) eq_defaultcaaaaaa\S4wooSPPRR4oSf4\P!SP 44oVVV3RloRSn!VVVV3RlR\4oS.p\P!S4Fwr#\P!V4'g5\P!V4'g\V\4'gKUVR8XdK^VR8XdKgVSP9dKz\SW#4VR8wgKVPV4K VSnSP"SnSP$SnSP&SnSP(Sn\P!SP 44Sn\-S4SnS#)rNc v<\V4\S 48d?\SP R\S 4^, R\V4^, R24h\S 4Fwr4V\V48d+WB9d\SP RV R24hW,pMDWB9dVP V4pM,VS9d SV,pM\SP RV R24h\ V4p\ WV4K V'd0\\V44p\SP RV R24hR#)z.__init__() takes z positional arguments but z were givenz..__init__() got multiple values for argument ''z(.__init__() missing required argument: 'z0.__init__() got an unexpected keyword argument 'N) rLr^rOrpoprAr<nextr]) rHr7r8rrJr unexpectedr\rZ field_namess &*, r"init_aggregate..initas54y3{++3<<.0B3{CSVWCWBXY''*4y1}o[!BCC )5 3t9$~'3<<.8fgkfllm(noo KE^"JJt,E\)(.E#s||n4\]a\bbc$dee'.E* 6!$v,/ 3<<.0`ak`llm noor%Tc<a]tRtRtRtRtZt]V3Rl4tRRRR/VVV3Rllt VVV3Rlt R V3R llt R V3R llt ] VV3R l4tRtV;t#)#_aggregate..aggregate_valuei|Tc "<\SV`V4#r)rS__new__)this_clsrs&r"r51_aggregate..aggregate_value._get_instances7?8, ,r%r3NrCc <VP4p\PVRR4/p\S \4'dMRR\ P !S 4P9dWR&R\ P !S 4P9dW&R&S !V.VO5/VBVBSP4F0p\WW4'dK\RS P RV R24h \PVRR4V#)_aggregate_init_completeFr3rCzconstructor for z did not initialize attribute 'rdT) r5rR __setattr__r&rr>r?rKrMrrrO) ror3rCr7r8r6 extra_kwargsrJr[rZris &$$*, r"rn+_aggregate..aggregate_value.__new__s--/H   x)CU KL$ ,,'"3"3D"9"D"DD09-7#4#4T#:#E#EE1;.  ;D ;L ;F ;(,,.x..(+;CLL>Ihimhnno)pqq/   x)CT JOr%c  B<VS9d\SP RV R24h\VSV,4'g%\RSV, RV R\ V4 24h\ VRR4'd\RV RSP R 24h\ SV`W4R #) z has no attribute 'rdz Expected z for attribute 'z', got rrFzcannot assign to field 'z' on immutable aggregate z6; use aggregate_replace() to construct a modified copyN)rrOr&r^rrrerSrs)rHrJrrr[rZs&&&r"rs/_aggregate..aggregate_value.__setattr__s?*$ ~5Ha%PQQe_T%:;;)OD,A+BBRSWRXX_`dej`k_l mnnt7??$'?vE^_b_k_k^lm\&]^^ G  ,r%c$V^8dQhRRRRRR/#rtr.)r/s"r"r00_aggregate..aggregate_value.__annotate__s& Q QZ Qs Qt Qr%c t<SP4F"p\W4PW RV 24K$ R#)rN)rMrer})rHrurJkey_namer[s&&& r"r}-_aggregate..aggregate_value._set_names4+002'11'V1XJ;OP3r%c V^8dQhRRRR/#rr.)r/s"r"r0rys 9 9~ 9$ 9r%c j<SP4Fp\W4PV4K R#r)rMrer)rHrrJr[s&& r"r/_aggregate..aggregate_value._flatten_irs)',,.#//8/r%c  <\SSP4Uu.uFq\W4P3NK up4#uupir)r0rMrerr)rHrJaggregate_valuer[s& r"rr(_aggregate..aggregate_value.typesB"?Q`QeQeQg#hQg74+>+C+C$DQg#hj j#hsA r.)rOrkrlrmrrNr classmethodr5rnrsr}rr\rrrnr])rrr[rZris@r"rrl|so!#)  -  - t   8 - Q Q 9 9  j  jr%rr0__annotate_func__)r`rWrXrirjrMrrpr> getmembers isfunctionismethodr&rr<r hash_attrsrOrkrlrN__aggregate_fields__dictrO) rZrrJmemberrr[r\rhris f @@@@@r" _aggregaterXso$=c$B!O\ <<  J -D |nn_%9%9%;<  p0#':j:j*:jxJ!,,S1   f % %)9)9&)A)AZPVXcEdEdz!~%/// OT 2**!!&)!2$",O"||O!$O##3O !kkO+3>>/:N:N:P+QO(-1,-?O* r%c z\\V4RR4'g\R\V4 24h\V4PpVF-pW29gK \\V4P RV R24h VUu/uFq3\W4bK ppVP V4\V4!R/VB#uupi)uCreate a copy of an aggregate instance with specified fields replaced. Similar to dataclasses.replace() — returns a new instance of the same aggregate type with the given fields updated and all other fields copied from the original instance. :param instance: The aggregate instance to copy :param changes: Keyword arguments for fields to replace :return: A new aggregate instance with the specified changes rNFz7aggregate_replace() expects an aggregate instance, got z has no field 'rdr.)rerrr^rrOrS)r6changesrhrJr8s&, r"aggregate_replacers 4>#95 A AQRVW_R`Qabccx.55K  "tH~667tfANO O9D D GH++ F D MM' > #F ##EsB8cV^8dQhRR/#rbr.)r/s"r"r0r0s[[D[r%cL\V\4;'d\VRR4#)__triton_block_ptr__F)r&rprers&r" _is_block_ptrrs" eZ ( Z ZWU**BCC Nr%c]tRtRt$R]R&R]R&R]R&R]R&R]R&R]R &R tRR ltR tRRltRRlt RRlt RRlt Rt R #) _block_ptri$rUr]rjrrrrorderTNc  VPP4'd!VPP4'd \R4h\ VPP \ 4'd \R4hWn\VRV4Vn \VRV4Vn \VRV4Vn \VRRR7Vn \VR4Vn\VP4pV^8Xd \R 4hRVP3RVP3RVP3RVP33F2wr\V 4V8wgK\R V R \V 4 R V 24h VPU u.uFp \!V 4NK p p \#V 4\%\&P(!V448wd\R V^, 24hR#uup i)z+Expected `base` to be a scalar pointer typez1Expected `base` to point to a scalar element typerrrrT)rrzMsGc VVPUu.uFp\V4NK up#uupir)rr\)rHextents& r" _tile_shape_block_ptr._tile_shapeIs(;?;K;KL;K$V,;KLLLs&c  VP4p\V\V44pVPpRp\ V4EF/wrx\ VP V,\^WR7VR7p \P!V4Fp \V ^VR7p K \P!V^,\V44Fp \V RVR7p K \WVR7p \ V\WPV,VR7VR7pWt9gKVPVPV ^4VP!WP"V,44p VfT MVPWk4pEK2 WV3#)Nr)rrrLr]rrrarangerirrrr rr/r[rar) rHrr3 tile_shape checked_dimsptrsrrrcoordrvalids &&& r" _materialize_block_ptr._materializeLs/%%' =ncR\o^ yy$Z0KC S)6!V+Q]fgE^^C(#E1 B)^^C!GS_=#E2C> iHEtS S(9YO[deD"!y'>'>ua'H)J]J]^ceoeopsetJuv $ u)..2M1zr%c  .p\VRV4p\V4\VP48wd.\R\VP4 R\V4 24h\ VPV4F!wrEVP \ WEVR74K# \VPVPVP\V4VPVPVR7#)rz"Expected `offsets` to have length z but received r)rrLrr5rrrrr]rrrjrr)rHrr3 new_offsets old_offsetdeltas&&& r"r_block_ptr.advance^s 7IV w<3t||, ,A#dllBSATTbcfgncobpqr r!$T\\7!; J   s: J K"<$))TZZu[?QSWScSceieoeo$-/ /r%c  (VfVe \R4h\V4p\V4p\V4p\V4pVfRpVPW8R7wrVR8XdRp MVR8Xd:VPPP P 4'dRM^p MaVR8XdLVPPP P4'd \R4h\R4p M\RV R 24h\WWWgVR 7#) NzK`mask` and `other` arguments cannot be specified for loading block pointersrDrzeronanz@Padding option `nan` is not supported for integer block pointerszPadding option z not supported)rrrrvolatiler3) r5r\rr]rrrWrrr>r) rHrrrpadding_optionrrrr3rgenerated_others &&&&&&&&& r"r_block_ptr.loadhs  u0jk k-n=-n=.?'1  !N&&~&K  R "O v %&*YY__%?%?%K%K%M%McSTO u $yy))0022 !cdd#ElO~.>nMN ND?$3R[] ]r%c  Ve \R4h\V4p\V4pVPW6R7wrr\WqW$VVR7#)Nz>`mask` argument cannot be specified for storing block pointersr)rrrr3)r5r\rr)rHrrrrrr3rs&&&&&&& r"r_block_ptr.storesT  ]^ ^-n=.?&&~&K Tt\k(* *r%)r]rrrrrr)r.NNNr.rDrDrDFNNr.rDrDN) rOrkrlrmrrrrrrrrrnr.r%r"rr$sM L L N N LW4M$/]:**r%rc.VPP#r)r]rrrs&r"r#r#s r%c:\V4pVPV4#)z Returns the id of the current program instance along the given :code:`axis`. :param axis: The axis of the 3D launch grid. Must be 0, 1 or 2. :type axis: int )r\ program_idrr3s&&r"rrs  %D    %%r%c:\V4pVPV4#)z Returns the number of program instances launched along the given :code:`axis`. :param axis: The axis of the 3D launch grid. Must be 0, 1 or 2. :type axis: int )r\ num_programsrs&&r"rrs  %D  ! !$ ''r%cP\V4p\V4pVPW4#r)r\r)rendr3s&&&r"rrs'  'E s #C   E ''r%z Returns contiguous values within the half-open interval :code:`[start, end)`. :code:`end - start` must be less than or equal to :code:`TRITON_MAX_TENSOR_NUMEL = z` :param start: Start of the interval. Must be a power of two. :type start: int32 :param end: End of the interval. Must be a power of two greater than :code:`start`. :type end: int32 cX\V4pVUu.uFp\V4NK up#uupirr_)rr)s& r"rrs*  'E-2 3U  #U 33 3s'c2\V4p\V4V#r)rrrs&r"_shape_check_implrs % E Lr%ch\V4p\V4p\V4pVPWV4#)aX Returns a tensor filled with the scalar value for the given :code:`shape` and :code:`dtype`. :param shape: Shape of the new array, e.g., (8, 16) or (8, ) :type shape: tuple of ints :param value: A scalar value to fill the array with :type value: scalar :param dtype: Data type of the new array, e.g., :code:`tl.float16` :type dtype: tl.dtype )rr\full)rrrrr3s&&&&r"rrs3 e $E  'E  'E >>% ..r%c$VPW4#)z Tries to broadcast the two given blocks to a common compatible shape. :param input: The first input tensor. :type input: Block :param other: The second input tensor. :type other: Block )broadcast_impl_value)inputrr3s&&&r" broadcastrs  ) )% 77r%r3cL\\V44pVPW4#)a= Tries to broadcast the given tensor to a new :code:`shape`. :param input: The input tensor. :type input: Block :param shape: The desired shape. :type shape: :code:`shape` can be passed as a tuple or as individual parameters: :: # These are equivalent broadcast_to(x, (32, 32)) broadcast_to(x, 32, 32) )rr`broadcast_impl_shaperr3rs&$*r"rrs%" .u5 6E  ) )% 77r%cV^8dQhRR/#)r,rrUr.)r/s"r"r0r0s***r%c\V4pV'ge\VP4pV^8d \R4h\ \ P !V^, 44V^, V^, .,pVPW4#)ar Permutes the dimensions of a tensor. If the parameter :code:`dims` is not specified, the function defaults to swapping the last two axes, thereby performing an (optionally batched) 2D transpose. :param input: The input tensor. :param dims: The desired ordering of dimensions. For example, :code:`(2, 1, 0)` reverses the order dims in a 3D tensor. :code:`dims` can be passed as a tuple or as individual parameters: :: # These are equivalent trans(x, (2, 1, 0)) trans(x, 2, 1, 0) :py:func:`permute` is equivalent to this function, except it doesn't have the special case when no permutation is specified. z2tl.trans invoked with a 0- or 1-dimensional tensor)r`rLrr5rPrirr)rr3rns&$* r"rrsk. D !D    q5QR RHNN1q5)*a!eQU^;   U ))r%c:\V4pVPW4#)a Permutes the dimensions of a tensor. :param input: The input tensor. :type input: Block :param dims: The desired ordering of dimensions. For example, :code:`(2, 1, 0)` reverses the order dims in a 3D tensor. :code:`dims` can be passed as a tuple or as individual parameters: :: # These are equivalent permute(x, (2, 1, 0)) permute(x, 2, 1, 0) :py:func:`trans` is equivalent to this function, except when :code:`dims` is empty, it tries to swap the last two axes. )r`r)rr3rs&$*r"rr4s( D !D   U ))r%c^aaaV'dVPSSV4#\SP4pV\SP48Xg!QRV R\SP4 24h\\ S4V4o\ ;QJd5VVV3Rl\ P!V44F 'dK RM* RM&!VVV3Rl\ P!V444'g%QRS RSP RSP 24h\SSVR7p\\ P!V44pVPSV4\WgVR7p\SP4pSPS,SPS,,VS&\WhVR7#) a Concatenate the given blocks :param input: The first input tensor. :type input: Tensor :param other: The second input tensor. :type other: Tensor :param can_reorder: Compiler hint. If true, the compiler is allowed to reorder elements while concatenating inputs. Only use if the order does not matter (e.g., result is only used in reduction ops). :type can_reorder: bool :param dim: The dimension to concatenate along (used when can_reorder is False). :type dim: int z%tensors must have the same rank, got z and c3<"TF4qS8wgK SPV,SPV,8HxK6 R#5irr)rarPrrrs& r"rbcat..bs69MAR/u{{1~Q/9Ms?/?FTz6tensor dims must match except in the concat dimension , got r) catrLr _wrap_axisr\rrirrrPinsertrr) rr can_reorderrr3rrr new_shapes ff&f& r"rrLsd }}UE;77 u{{ D 3u{{# #j'LTFRWX[\a\g\gXhWi%jj # )#. 5C 39M3339M  yMcURXY^YdYdXeejkpkvkvjwxy  UEY/A % &E LLdI.AU[[!I[[% C(88IcN 19 55r%c$VPW4#)a& Join the given tensors in a new, minor dimension. For example, given two tensors of shape (4,8), produces a new tensor of shape (4,8,2). Given two scalars, returns a tensor of shape (2). The two inputs are broadcasted to be the same shape. If you want to join more than two elements, you can use multiple calls to this function. This reflects the constraint in Triton that tensors must have power-of-two sizes. join is the inverse of split. :param a: The first input tensor. :type a: Tensor :param b: The second input tensor. :type b: Tensor )r)abr3s&&&r"rros* >>! r%c\VP4^8XdV#^pVPF pW4,pK V^8XgQR4hVPV4#)z. Convert a single-element tensor to a scalar. z'can only unsplat single-element tensors)rLrunsplat)r!r3rCrds&&& r"_unsplatrsV 177|q E WW  A:@@@:   Q r%cV^8dQhRR/#rr.)r/s"r"r0r0s1Fr%c\VP4^8HpV'dVPV^4pVPV4wrEV'd\ WAVR7p\ WQVR7pWE3#)a Split a tensor in two along its last dim, which must have size 2. For example, given a tensor of shape (4,8,2), produces two tensors of shape (4,8). Given a tensor of shape (2), returns two scalars. If you want to split into more than two pieces, you can use multiple calls to this function (probably plus calling reshape). This reflects the constraint in Triton that tensors must have power-of-two sizes. split is the inverse of join. :param a: The tensor to split. :type a: Tensor r3rC)rLrrrr)rr3rC was_rank_1out_lhsout_rhss&&& r"rrs`*QWW"J  ! !!Q ' q)G7JO7JO  r%cf\R4\\V44pVPWRR7#)a^ Returns a tensor with the same elements as `input` but a different shape. The order of the elements may not be preserved. :param input: The input tensor. :type input: Block :param shape: The desired shape. :code:`shape` can be passed as a tuple or as individual parameters: :: # These are equivalent view(x, (32, 32)) view(x, 32, 32) zCview is deprecated, please use reshape with can_reorder being true.Tr)rrr`rrs&$*r"rrs3" NO .u5 6E   Ut  < Returns a tensor with the same number of elements as input but with the provided shape. :param input: The input tensor. :type input: Block :param shape: The new shape. :code:`shape` can be passed as a tuple or as individual parameters: :: # These are equivalent reshape(x, (32, 32)) reshape(x, 32, 32) r)rr`rLrr)rrr3rCrs&$$$*r"rrs?" .u5 6E 5zQzJJ   U; 77r%crV)Tu;8:dV8gM\RV RV) RV 24hV^8dV#W,#)z invalid axis z . Expected z <= axis < r4)rndims&&r"rrsD ET D =k4% D6RSS194-$+-r%cVPV4p\V4p\V\\34'd \ V4MV.p\ VP4\ V4,pVUu.uFp\\V4V4NK pp\ \V44\ V48wd\RV 24hTp\V4FpVPWg4pK V#uupi)a2 Expand the shape of a tensor, by inserting new length-1 dimensions. Axis indices are with respect to the resulting tensor, so ``result.shape[axis]`` will be 1 for each axis. :param input: The input tensor. :type input: tl.tensor :param axis: The indices to add new axes :type axis: int | Sequence[int] z7expand_dims received duplicate axes, normalized axes = ) rhr\r&r rjrPrLrrrr5rr)rrr3axesnew_ndimrrrs&&& r"rrs    &E  %D#D8U*;<<4:4&D5;;#d)+HCG H4aJ+A. 94D H 3t9~T"RSWRXYZZ C D\##C+ J Is0C'c$V^8dQhRRRRRR/#rr.)r/s"r"r0r0 s">>u>M>SW>r%cVPV4p\V4p\V4p\V4pV'dVPW4#VPWV4#)au Casts a tensor to the given :code:`dtype`. :param dtype: The target data type. :type dtype: tl.dtype :param fp_downcast_rounding: The rounding mode for downcasting floating-point values. This parameter is only used when self is a floating-point tensor and dtype is a floating-point type with a smaller bitwidth. Supported values are :code:`"rtne"` (round to nearest, ties to even) and :code:`"rtz"` (round towards zero). :type fp_downcast_rounding: str, optional :param bitcast: If true, the tensor is bitcasted to the given :code:`dtype`, instead of being numerically casted. :type bitcast: bool, optional )rhr\rr)rrrrrr3s&&&&&r"rr sZ$    &E  'E/0DE"7+G  .. >>%(< ==r%cVe VeQR4hVf^RVPPP9p\PP ;'gV'dV'gVfRMRp\ V4p\ V4p\ V4p\ V4p\VP4p \VP4p \V 4\V 4u;8Xd^8gQR4hQR4hV RR V RR 8XgQR4hV R ,V R ,8XgQR4hV RR V R ,.,p Ve"\VP4V 8XgQR4h\V 4p V ^8d^p \P!V ^, 4FpWV,,p K VPW .V R R,R R 7pVPW.V R R,R R 7pVeVPW-.V R R,R R 7pVPWW#WV4pV ^8dVPWR R 7p\VP4V 8XgQR 4hV#)a Returns the matrix product of two blocks. The two blocks must both be two-dimensional or three-dimensional and have compatible inner dimensions. For three-dimensional blocks, `tl.dot` performs the batched matrix product, where the first dimension of each block represents the batch dimension. .. warning:: When using TF32 precision, the float32 inputs may be truncated to TF32 format (19-bit floating point) without rounding which may bias the result. For best results, you must round to TF32 explicitly, or load the data using `TensorDescriptor` with `round_f32_to_tf32=True`. :param input: The first tensor to be multiplied. :type input: 2D or 3D tensor of scalar-type in {:code:`int8`, :code:`float8_e5m2`, :code:`float16`, :code:`bfloat16`, :code:`float32`} :param other: The second tensor to be multiplied. :type other: 2D or 3D tensor of scalar-type in {:code:`int8`, :code:`float8_e5m2`, :code:`float16`, :code:`bfloat16`, :code:`float32`} :param acc: The accumulator tensor. If not None, the result is added to this tensor. :type acc: 2D or 3D tensor of scalar-type in {:code:`float16`, :code:`float32`, :code:`int32`} :param input_precision: How to exercise the Tensor Cores for f32 x f32. If the device does not have Tensor Cores or the inputs are not of dtype f32, this option is ignored. For devices that do have tensor cores, the default precision is tf32. :type input_precision: string. Available options for nvidia: :code:`"tf32"`, :code:`"tf32x3"`, :code:`"ieee"`. Default: :code:`"tf32"`. Available options for amd: :code:`"ieee"`, (CDNA3 only) :code:`"tf32"`. :param allow_tf32: *Deprecated.* If true, input_precision is set to "tf32". Only one of :code:`input_precision` and :code:`allow_tf32` can be specified (i.e. at least one must be :code:`None`). Nz;Only one of input_precision and allow_tf32 can be specifiedtf32ieeez*input and other must have equal ranks >= 2z,input and other must have equal batch shapesz4input and other must have equal reduction dimensionsz!accumulator shape is incompatibleFrzoutput shape is unexpectedr)rur$allowed_dot_input_precisionsrlanguage fp32_defaultr\rPrrLrirrdot)rraccinput_precision allow_tf32max_num_imprecise_acc out_dtyper3 supports_tf32a_shapeb_shapec_shaper batch_sizerPress&&&&&&&& r"rr2 sQ<  "j&8w:ww 8)"3"3";";"X"XX ..55vv]FPT^Tf;Ant +?;O$Y/I01FG s #C5;;G5;;G w<3w< ,1 ,Z.ZZ ,Z.ZZ , 3B<73B< 'W)WW ' 2;'"+ %]']] %crlgbk]*G CII')N+NN) w`_ for more details. :param eviction_policy: changes eviction policy in NVIDIA PTX :type eviction_policy: str, optional :param volatile: changes volatile option in NVIDIA PTX :type volatile: bool, optional )rrrrrrrr3)rrr\rh) pointerrrrrrrrr3s &&&&&&&&&r"rr s\W||>+9em&/1 1  %D  'E ""4( ##E*).9N).9N*?;O#H-H >>'`o" $$r%c$V^8dQhRRRRRR/#)r,descrrrr-rUr.)r/s"r"r0r0 s$33!73B^3.43r%c&VPWR7#)z.Load a block of data from a tensor descriptor.r)r)r+rr3s&&&r"load_tensor_descriptorr- s 99W9 22r%c(V^8dQhRRRRRRRR/#)r,r+rrrrrUr-r.)r/s"r"r0r0 s,;;"8;C_;hn;/5;r%c(VPWVR7#)z-Store a block of data to a tensor descriptor.r)r)r+rrr3s&&&&r"store_tensor_descriptorr0 s ::g : ::r%c \V4'dVPWW4WVR7#VPV4p\V4pVeVPV4p\V4p\V4pVPWW#WE4#)a Store a tensor of data into memory locations defined by `pointer`. (1) If `pointer` is a single element pointer, a scalar is stored. In this case: - `mask` must also be scalar, and - `boundary_check` and `padding_option` must be empty. (2) If `pointer` is an N-dimensional tensor of pointers, an N-dimensional block is stored. In this case: - `mask` is implicitly broadcast to `pointer.shape`, and - `boundary_check` must be empty. (3) If `pointer` is a block pointer defined by `make_block_ptr`, a block of data is stored. In this case: - `mask` must be None, and - `boundary_check` can be specified to control the behavior of out-of-bound access. `value` is implicitly broadcast to `pointer.shape` and typecast to `pointer.dtype.element_ty`. :param pointer: The memory location where the elements of `value` are stored :type pointer: `triton.PointerType`, or block of `dtype=triton.PointerType` :param value: The tensor of elements to be stored :type value: Block :param mask: If `mask[idx]` is false, do not store `value[idx]` at `pointer[idx]` :type mask: Block of triton.int1, optional :param boundary_check: tuple of integers, indicating the dimensions which should do the boundary check :type boundary_check: tuple of ints, optional :param cache_modifier: changes cache option in NVIDIA PTX :type cache_modifier: str, optional, should be one of {"", ".wb", ".cg", ".cs", ".wt"}, where ".wb" stands for cache write-back all coherent levels, ".cg" stands for cache global, ".cs" stands for cache streaming, ".wt" stands for cache write-through, see `cache operator `_ for more details. :param eviction_policy: changes eviction policy in NVIDIA PTX :type eviction_policy: str, optional, should be one of {"", "evict_first", "evict_last"} )rrrrr3)rrrhr\)r)rrrrrr3s&&&&&&&r"rr sRW}}Un-<S S    &E  %D ""4().9N*?;O ??74 aar%cV^8dQhRR/#)r,r]rUr.)r/s"r"r0r06 s ^ ^ ^r%c 6\R4\WW#WEVR7#)aK Returns a pointer to a block in a parent tensor :param base: The base pointer to the parent tensor :param shape: The shape of the parent tensor :param strides: The strides of the parent tensor :param offsets: The offsets to the block :param block_shape: The shape of the block :param order: The order of the original data format z[tl.make_block_ptr is deprecated. Use TensorDescriptor or tl.make_tensor_descriptor instead.r)rr)r]rrrrrr3s&&&&&&&r"make_block_ptrr45 s fg d7[S\ ]]r%zNote that tl.advance does not have any side effects. To move the block pointer, you need to assign the result of tl.advance to a variable.c^\V4'dVPWR7#\R4h)z Advance a block pointer :param base: the block pointer to advance :param offsets: the offsets to advance, a tuple by dimension rzH`tl.advance` only supports block pointers created by `tl.make_block_ptr`)rrr5)r]rr3s&&&r"rrE s,T||G|99 _ ``r%c ,V^8dQhRRRRRRRRRR /#) r,r]rUrr$rrzList[constexpr]r-rr.)r/s"r"r0r0W sA9_9_ 9_ 9_9_! 9_9_r%c>\V4pVPWW#V4#)aMake a tensor descriptor object :param base: the base pointer of the tensor, must be 16-byte aligned :param shape: A list of non-negative integers representing the tensor shape :param strides: A list of tensor strides. Leading dimensions must be multiples of 16-byte strides and the last dimension must be contiguous. :param block_shape: The shape of block to be loaded/stored from global memory Notes ***** On NVIDIA GPUs with TMA support, this will result in a TMA descriptor object and loads and stores from the descriptor will be backed by the TMA hardware. Currently only 2-5 dimensional tensors are supported. Example ******* .. code-block:: python @triton.jit def inplace_abs(in_out_ptr, M, N, M_BLOCK: tl.constexpr, N_BLOCK: tl.constexpr): desc = tl.make_tensor_descriptor( in_out_ptr, shape=[M, N], strides=[N, 1], block_shape=[M_BLOCK, N_BLOCK], ) moffset = tl.program_id(0) * M_BLOCK noffset = tl.program_id(1) * N_BLOCK value = desc.load([moffset, noffset]) desc.store([moffset, noffset], tl.abs(value)) # TMA descriptors require a global memory allocation def alloc_fn(size: int, alignment: int, stream: Optional[int]): return torch.empty(size, device="cuda", dtype=torch.int8) triton.set_allocator(alloc_fn) M, N = 256, 256 x = torch.randn(M, N, device="cuda") M_BLOCK, N_BLOCK = 32, 32 grid = (M / M_BLOCK, N / N_BLOCK) inplace_abs[grid](x, M, N, M_BLOCK, N_BLOCK) )r\make_tensor_descriptor)r]rrrrr3s&&&&&&r"r8r8V s$r*.9N  + +D~ ^^r%c$V^8dQhRRRRRR/#)r,rJr'has_cmprcr-Callable[[T], T]r.)r/s"r"r0r0 s"S4.__annotate__ sqr%cT<RS R2pS'd VR, pVR, pWnV#)z Performs an atomic z at the memory location specified by :code:`pointer`. Return the data stored at :code:`pointer` before the atomic operation. :param pointer: The memory locations to operate on :type pointer: Block of dtype=triton.PointerDTypez| :param cmp: The values expected to be found in the atomic object :type cmp: Block of dtype=pointer.dtype.element_tya :param val: The values with which to perform the atomic operation :type val: Block of dtype=pointer.dtype.element_ty :param sem: Specifies the memory semantics for the operation. Acceptable values are "acquire", "release", "acq_rel" (stands for "ACQUIRE_RELEASE"), and "relaxed". If not provided, the function defaults to using "acq_rel" semantics. :type sem: str, optional :param scope: Defines the scope of threads that observe the synchronizing effect of the atomic operation. Acceptable values are "gpu" (default), "cta" (cooperative thread array, thread block), or "sys" (stands for "SYSTEM"). The default value is "gpu". :type scope: str, optional )rN)r?docstrr:rJs& r" _decorator&_add_atomic_docstr.._decorator sJv69  : :F     r%r.)rJr:rCsff r"_add_atomic_docstrrE s4 r%zcompare-and-swap)r:cVPV4pVPV4p\V4p\V4pVPWW#V4#r)rhr\r)r)rrrrr3s&&&&&&r"rr sM   c "C   c "C s #C  'E   c >>r%exchangecVPV4p\V4p\V4p\V4pVPWW#V4#r)rhr\rr)rrrrr3s&&&&&&r"rr sG   c "C s #C  'E  %D  t% @@r%rcVPV4p\V4p\V4p\V4pVPWW#V4#r)rhr\rrIs&&&&&&r"rr G   c "C s #C  'E  %D   d ??r%rCcVPV4p\V4p\V4p\V4pVPWW#V4#r)rhr\rrIs&&&&&&r"rr rKr%rNcVPV4p\V4p\V4p\V4pVPWW#V4#r)rhr\rrIs&&&&&&r"rr rKr%z logical andcVPV4p\V4p\V4p\V4pVPWW#V4#r)rhr\rrIs&&&&&&r"rr rKr%z logical orcVPV4p\V4p\V4p\V4pVPWW#V4#r)rhr\rrIs&&&&&&r"rr sG   c "C s #C  'E  %D   wT >>r%z logical xorcVPV4p\V4p\V4p\V4pVPWW#V4#r)rhr\rrIs&&&&&&r"rr rKr%ctVPV4p\V4p\V4pVPWV4#)a Returns a tensor of elements from either :code:`x` or :code:`y`, depending on :code:`condition`. Note that :code:`x` and :code:`y` are always evaluated regardless of the value of :code:`condition`. If you want to avoid unintended memory operations, use the :code:`mask` arguments in `triton.load` and `triton.store` instead. The shape of :code:`x` and :code:`y` are both broadcast to the shape of :code:`condition`. :code:`x` and :code:`y` must have the same data type. :param condition: When True (nonzero), yield x, otherwise yield y. :type condition: Block of triton.bool :param x: values selected at indices where condition is True. :param y: values selected at indices where condition is False. )rhr\where) conditionr!yr3s&&&&r"rRrR s9"##I.IQAQA ??9 ++r%c\V4pVPV4p\VPPRR4pRV9dVP V^V4#VPPP 'dfVPV4pVP^4pVPVPWE4VPV44pVPVRR4V#)a Mark values that are expected to have underflowed to zero. In regular compilation this preserves :code:`x`. Debug builds assert that :code:`x` is zero wherever :code:`mask` is true. Under FPSAN this becomes :code:`where(mask, 0, x)` so sanitized execution observes the intended floating-point underflow. :param x: values to preserve outside FPSAN mode. :param mask: positions where :code:`x` is expected to be zero. instrumentation_moderDfpsanz.expect_zero expected x == 0 where mask is trueN) r\rhrerur$rRdebugr8rpr device_assert)r!rr3rVx_tensorrconds&&& r" expect_zeror\+ s QA   t $D"9#4#4#<#<>TVXY&&tQ**  &&&&&q)""1%}}Y__XRS&VX\] Hr%cV^8dQhRR/#r,rrr.)r/s"r"r0r0K 222r%cR\V4p\V4pVPWV4#r)r\rr!rTrr3s&&&&r"rrJ (QAQA ==0 11r%cV^8dQhRR/#r^r.)r/s"r"r0r0R r_r%cR\V4p\V4pVPWV4#r)r\rras&&&&r"rrQ rbr%cV^8dQhRR/#r^r.)r/s"r"r0r0Y r_r%cR\V4p\V4pVPWV4#r)r\r ras&&&&r"r r X rbr%cV^8dQhRR/#r, propagate_nanrr.)r/s"r"r0r0` 222r%cVPV4pVPV4p\WR7p\WR7p\V4pVPWV4#)a$ Computes the element-wise minimum of :code:`x` and :code:`y`. :param x: the first input tensor :type x: Block :param y: the second input tensor :type y: Block :param propagate_nan: whether to propagate NaN values. :type propagate_nan: tl.PropagateNan .. seealso:: :class:`tl.PropagateNan` r)rh_promote_bfloat16_to_float32r\minimumr!rTrir3s&&&&r"rmrm_ U AAAA$Q :code:`max` is undefined. :param x: the input tensor :type x: Block :param min: the lower bound for clamping :type min: Block :param max: the upper bound for clamping :type max: Block :param propagate_nan: whether to propagate NaN values. Applies only to the :code:`x` tensor. If either :code:`min` or :code:`max` is NaN, the result is undefined. :type propagate_nan: tl.PropagateNan .. seealso:: :class:`tl.PropagateNan` r)rhrlr\clamp)r!rNrCrir3s&&&&&r"ruru sm$ AA   c "C   c "C$Qr.)r/s"r"r0+_add_reduction_docstr..__annotate__ sqr%c<RpSeVRS RS RS R2, pSeVRS RS RS R2, pSeVRS RS RS R 2, pVPSR 7VnV#) a Returns the {name} of all elements in the :code:`input` tensor along the provided :code:`axis` The reduction operation should be associative and commutative. :param input: the input values :type input: Tensor :param axis: the dimension along which the reduction should be done. If None, reduce all dimensions :type axis: int :param keep_dims: if true, keep the reduced dimensions with length 1 :type keep_dims: bool :param z-: if true, return index corresponding to the z value :type z: boolzC: if true, in case of a tie (i.e., multiple elements have the same zI value), return the left-most index for values that aren't NaN :type c: the desired data type of the returned tensor. If specified, the input tensor is casted to :code:`z` before the operation is performed. This is useful for preventing data overflows. If not specified, integer and bool dtypes are upcasted to :code:`tl.int32` and float dtypes are upcasted to at least :code:`tl.float32`. :type : tl.dtyperIr/rN)r?rBryrJrwrxs& r"rC)_add_reduction_docstr.._decorator s   )   LTFS  f( (F  $  ?]^b]cd /# #F   ;y{DzEE +Z# #F}}$}/  r%r.)rJrwrxryrCsffff r"_add_reduction_docstrr s8 r%c#X"VP4pRxVPV4R#5ir)get_insertion_pointrestore_insertion_point)ruips& r"_insertion_guardr s$  $ $ &B  ##B's(*c $aaaaaa\S\4'd\S3SSVSSR7^,#VVVV3RlpV3Rlo\S4o\V4pSe'\ S\ S^,P 44oSPSSV4pV'dqSe9\;QJd.VV3RlV4FNK 5pV#!VV3RlV44pV#\;QJd.VV3RlV4FNK 5M!VV3RlV44pV#)aApplies the combine_fn to all elements in :code:`input` tensors along the provided :code:`axis` :param input: the input tensor, or tuple of tensors :type input: Tensor :param axis: the dimension along which the reduction should be done. If None, reduce all dimensions :type axis: int | None :param combine_fn: a function to combine two groups of scalar tensors (must be marked with @triton.jit) :type combine_fn: Callable :param keep_dims: if true, keep the reduced dimensions with length 1 :type keep_dims: bool rr3rCc <a SUu.uFqPPNK up^,pVP^4pSPo \ S 4;_uu_4V 3RlpS P V\ \WB444p\V4UUu.uF wrg\VPV4V4NK" pppS PSV/R7p \V \4'dV P.p MV U u.uFqPNK p p S P!V !RRR4R#uupiuuppiuup i +'giR#;i)r,c&<VPS4#rr rrus&r"r#5reduce..make_combine_region.. aggg.r%rKN)rrrp get_regionrurcreate_block_with_parentrPrrrUrcall_JitFunctionr&rcreate_reduce_ret) reduce_oprm param_typesregionr blockrPrr7resultsrrrurCr3rrs& @r"make_combine_region#reduce..make_combine_region s .34evv}}e4q8 %%a(## g & &.E44VT#eBY=Z[E:CK:PQ:PF599Q<,:PDQ 11*d21NG'6**">>*-45W88W5  % %w /' &5 R 6' & &/D-9D=&D2 >?D==D8D=2 D== E c^<\P!V4Fp\V^SR7pK V#)rr)rirr)rmndimsrr3s&& r" expand_ndimsreduce..expand_ndims s)&AAqI6A'r%c3@<"TFp\VSSR7xK R#5i)rN)r)rarmr3rs& r"rbreduce.. sO3a AtyAA3sc3j<"TF(pS!V\S^,P44xK* R#5irN)rLr)rarmrrs& r"rbr s(Jc QE!HNN(;< ?C  %O3O%C JO3OOC J%JcJ%%JcJJC Jr%ctVPPpV\JdVP\VR7#V#)r)rrrpbfloat16rr)rmr3rs&& r"rlrl s1 IHttGyt11 Hr%c \V4pVPV,p\^WdR7p\VP4^8dh\P !\VP44Uu.uFp\ V4NK p pW\WyVR7p\WpPVR7p\W3WW4VR7wrW3#uupi)rrr) r\rrrLrirrrrr) rrrrr3rCrrraxes_to_expandrvaluerindicess &&&&&& r"_reduce_with_indicesr s  %D DA 1a -E 5;;!08s5;;?O0PQ0P1)A,0PQ  EYGUKK9Eund))35F  Rs1Cc$V^8dQhRRRRRR/#)r,rJr'ryr-r;r.)r/s"r"r0r0+ s"33:Jr%caaRVV3RllpV#)c V^8dQhRRRR/#r>r.)r/s"r"r0&_add_scan_docstr..__annotate__- sqr%ch<RpSeVRS RS RS R2, pVPSR7VnV#)aZ Returns the {name} of all elements in the :code:`input` tensor along the provided :code:`axis` :param input: the input values :type input: Tensor :param axis: the dimension along which the scan should be done :type axis: int :param reverse: if true, the scan is performed in the reverse direction :type reverse: boolr~rz` before the operation is performed. If not specified, small integer types (< 32 bits) are upcasted to prevent overflow. Note that :code:`tl.bfloat16` inputs are automatically promoted to :code:`tl.float32`. :type rrIr)r?rBryrJs& r"rC$_add_scan_docstr.._decorator- sb   ;y{DzEE +Z# #F}}$}/  r%r.)rJryrCsff r"_add_scan_docstrr+ s& r%c  aaaa\S\4'd\S3VSVSSR7^,#VVVV3Rlp\V4p\V4pVe'\ V\ S^,P 44pSPSWV4#)aApplies the combine_fn to each elements with a carry in :code:`input` tensors along the provided :code:`axis` and update the carry :param input: the input tensor, or tuple of tensors :type input: Tensor :param axis: the dimension along which the reduction should be done :type axis: int :param combine_fn: a function to combine two groups of scalar tensors (must be marked with @triton.jit) :type combine_fn: Callable :param reverse: whether to apply the associative scan in the reverse direction along axis :type reverse: bool rc <a SUu.uFqPPNK up^,pVP^4pSPo \ S 4;_uu_4V 3RlpS P V\ \WB444p\V4UUu.uF wrg\VPV4V4NK" pppS PSV/R7p \V \4'dV P.p MV U u.uFqPNK p p S P!V !RRR4R#uupiuuppiuup i +'giR#;i)r,c&<VPS4#rrrs&r"r#?associative_scan..make_combine_region..Z rr%rKN)rrrprrurrrPrrrUrrr&rcreate_scan_ret)scan_oprmrrr rrPrr7rrrrurCr3rrs& @r"r-associative_scan..make_combine_regionU s .34evv}}e4q8 ##A&## g & &.E44VT#eBY=Z[E:CK:PQ:PF599Q<,:PDQ 11*d21NG'6**">>*-45W88W5  # #W -' &5 R 6' & &r)r&rUrr\rrLr)rrrrr3rCrsf&f&ff r"rrC s%   4WPYfpqrstt . .  %D"7+G $E!HNN 34  % %eT PPr%c|\V4p\V4pVeVPV4pVPWV4#)aUcomputes an histogram based on input tensor with num_bins bins, the bins have a width of 1 and start at 0. :param input: the input tensor :type input: Tensor :param num_bins: number of histogram bins :type num_bins: int :param mask: if `mask[idx]` is false, exclude `input[idx]` from histogram :type mask: Block of `triton.int1`, optional )r\rhr))rr(rr3rCs&&&&&r"r)r)k sA$H-H  %D ""4(   u 55r%ch\V4p\V4p\V4pVPWV4#)zGather from a tensor along a given dimension. :param src: the source tensor :type src: Tensor :param index: the index tensor :type index: Tensor :param axis: the dimension to gather along :type axis: int )r\r$)srcrrr3s&&&&r"r$r$ s5 s #C  'E  %D   C --r%packc V^8dQhRRRR/#)r, scalar_fnz!Callable[..., Tuple[tensor, ...]]r7rUr.)r/s"r"r0r0 s H.H.0H. H.r%c aV^8g QRV 24hVUu.uFqUPPNK ppVPpVP4o.pVP 4p \ V4Fzwr\ P!V4F[p SPV PV4V 4VP\SPW,V ,4V 44K] K| \V4;_uu_4VPS4VPW/R7p \!V \4pV'dV 3p V Uu.uFqP"NK ppVP%V 4VP'V4RRR4X Uu.uFpVPNK ppTpV^8dCVRRV1,p\ P!^V4FpVVVRV1,8XdKQR4h V3RlpVP%V 4VP)VVVV4pX'd V^,#T#uupiuupi +'giL;iuupi)ad Map a scalar function over a tensor. The input tensors :code:`args` are implicitly broadcasted to the same shape. This may be useful in allowing control flow over single elements in a tensor, for example a multi-branch function where one branch is more expensive. With :code:`tl.where` you are forced to calculate both sides of the branch, but with an if we only execute one side. .. highlight:: python .. code-block:: python @triton.jit def selu_scalar(x, alpha): if x > 0: return a else: return alpha * (tl.exp(x) - 1) @triton.jit def selu(x, alpha): return tl.map_elementwise(selu_scalar, x, alpha) :param scalar_fn: the function to map over. :param pack: the number of elements to be processed by one function call. :return: one tensor or a tuple of tensors, depending on the mapped function. zpack must be >= 1, got rKNz!type mismatch in unpacked resultscL<VP^4pVPS4R#r)r push_back)elementwise_oprrs& r"make_elementwise_region0map_elementwise..make_elementwise_region s!**1-r%)rrrpru new_blockrrriradd_argument_atr rrUrrset_insertion_point_to_startrr&rrcreate_map_elementwise_retmap_elementwise)rrr3rCr7rm in_scalar_tysru scalar_args original_locrPrjscalar_results is_singlerrr!fn_result_typesscalar_result_typesoffsetrresultrs&$$$* @r"rr sJ 196/v669,01DqVV]]DM1G    EK??$L=)%A  ! !"((7"3\ B   veii1 &=rB C&* ' " ",,U3#44YTV4W~v6 +_N%34^88^4 %**73 #(66~!qvv~O6) ax-ff5nnQ-F&/&,$,*GG lIl lG.  OOL!  & &t-@$H_ `F!6!9-v-G2$5 # "7s*H)AH3H.$H3 I.H33 I c"VP4#)z9 Insert a barrier to synchronize all threads in a block. ) debug_barrierrs&r"rr s  " " $$r%c \V\4'dV.p\V4Fowr4\V\4'g\RV R24h\VP\ 4'dKL\RV R\ VP4 R24h VUu.uFqUPNK ppVPW4#uupi)a Let the compiler know that ``values[d]`` is the largest power of two that divides the first element of every contiguous group along dimension ``d`` of :code:`input` (see :func:`max_contiguous` for the definition of contiguous group). ``values`` must have one entry per dimension of :code:`input`. For a 1D input with contiguity 1, this is equivalent to saying that every element of :code:`input` is a multiple of ``values[0]``. For example, if :code:`values` is ``[16]`` and :code:`input` is ``[64, 80, 96, 112]``, the hint is valid because 16 divides 64. This hint enables alignment-dependent optimizations such as vectorized memory accesses. values element  must have type `constexpr`1 must have type `constexpr[int]`, got `constexpr[r)r&rrr^rrrr multiple_ofrrQr3rPrr!s&&& r"rr s &)$$&!!Y''oaS0KLM M!''3''oaS0abfghgngnboappqrs s " & &v!ggvF &   //'Cc \V\4'dV.p\V4Fowr4\V\4'g\RV R24h\VP\ 4'dKL\RV R\ VP4 R24h VUu.uFqUPNK ppVPW4#uupi)a Let the compiler know that the elements of :code:`input` along dimension ``d`` form contiguous groups of length ``values[d]``. ``values`` must have one entry per dimension of :code:`input` and each entry must be a power of two. A 1D array of ``N`` elements with contiguity ``C`` is viewed as ``N/C`` runs of ``C`` integers each, where the integers in a run are sequentially contiguous. For example, if :code:`values` is ``[4]``, the array ``[0, 1, 2, 3, 8, 9, 10, 11]`` satisfies the hint because it consists of two runs of 4 contiguous values. Together with :func:`multiple_of`, this hint enables vectorized loads and stores of contiguous, aligned regions. rrrr)r&rrr^rrrrmax_contiguousrs&&& r"rr s"&)$$&!!Y''oaS0KLM M!''3''oaS0abfghgngnboappqrs s " & &v!ggvF &  # #E 22'rc \V\4'dV.p\V4Fowr4\V\4'g\RV R24h\VP\ 4'dKL\RV R\ VP4 R24h VUu.uFqUPNK ppVPW4#uupi)a Let the compiler know that the elements of :code:`input` along dimension ``d`` form constant groups of length ``values[d]``. ``values`` must have one entry per dimension of :code:`input` and each entry must be a power of two. A 1D array of ``N`` elements with constancy ``C`` is viewed as ``N/C`` runs of ``C`` identical values. For example, if :code:`values` is ``[4]``, the array ``[0, 0, 0, 0, 1, 1, 1, 1]`` satisfies the hint. rrrr)r&rrr^rrrr max_constancyrs&&& r"rr" s&)$$&!!Y''oaS0KLM M!''3''oaS0abfghgngnboappqrs s " & &v!ggvF &  " "5 11'rcBVPVPV44#)z4 Allow compiler to assume the :code:`cond` is True. )assumerh)r[r3s&&r"rr9 s   I//5 66r%sep r fileflushc V^8dQhRRRR/#)r,rr'rr.)r/s"r"r0r0G s   s  s  r%cR#)ai Print the values at compile time. The parameters are the same as the builtin :code:`print`. NOTE: Calling the Python builtin :code:`print` is not the same as calling this, it instead maps to :code:`device_print`, which has special requirements for the arguments. .. highlight:: python .. code-block:: python tl.static_print(f"BLOCK_SIZE={BLOCK_SIZE}") Nr.)rrrrr3rQs$$$$$*r" static_printrF s r%cR#)z Assert the condition at compile time. Does not require that the :code:`TRITON_DEBUG` environment variable is set. .. highlight:: python .. code-block:: python tl.static_assert(BLOCK_SIZE == 1024) Nr.)r[msgr3s&&&r" static_assertrV s r%hexc<^RIp\V4p\V\4'g QV R24hRpVFpWdP9gKRpM V'g QV R24h.pVF#pVP VP V44K% VPWV4#)a[ Print the values at runtime from the device. String formatting does not work for runtime values, so you should provide the values you want to print as arguments. The first value must be a string, all following values must be scalars or tensors. Calling the Python builtin :code:`print` is the same as calling this function, and the requirements for the arguments will match this function (not the normal requirements for :code:`print`). .. highlight:: python .. code-block:: python tl.device_print("pid", pid) print("pid", pid) On CUDA, printfs are streamed through a buffer of limited size (on one host, we measured the default as 6912 KiB, but this may not be consistent across GPUs and CUDA versions). If you notice some printfs are being dropped, you can increase the buffer size by calling .. highlight:: python .. code-block:: python triton.runtime.driver.active.utils.set_printf_fifo_size(size_bytes) CUDA may raise an error if you try to change this value after running a kernel that uses printfs. The value set here may only affect the current device (so if you have multiple GPUs, you'd need to call it multiple times). :param prefix: a prefix to print before the values. This is required to be a string literal. :param args: the values to print. They can be any tensor or scalar. :param hex: print all values as hex instead of decimal Nz is not stringTFz is not an ascii string)stringr\r&r' printablerrh device_print) rTrr3r7rb_asciichnew_argsrs &$$* r"rrd sD !& )F fc " "=vhn$== "G %% %G  6vh5667H ++C01  ! !&C 88r%c\V4p\V4pVeVPV4pVPVPV4W4#)a Assert the condition at runtime from the device. Requires that the environment variable :code:`TRITON_DEBUG` is set to a value besides :code:`0` in order for this to have any effect. Using the Python :code:`assert` statement is the same as calling this function, except that the second argument must be provided and must be a string, e.g. :code:`assert pid == 0, "pid != 0"`. The environment variable must be set for this :code:`assert` statement to have any effect. .. highlight:: python .. code-block:: python tl.device_assert(pid == 0) assert pid == 0, f"pid != 0" :param cond: the condition to assert. This is required to be a boolean tensor. :param msg: the message to print if the assertion fails. This is required to be a string literal. )r\rhrY)r[rrr3s&&&&r"rYrY sL& s #C  %D ""4(  " "9#6#6t#`_ assembly: .. highlight:: python .. code-block:: python @triton.jit def kernel(A, B, C, D, BLOCK: tl.constexpr): a = tl.load(A + tl.arange(0, BLOCK)) # uint8 tensor b = tl.load(B + tl.arange(0, BLOCK)) # float32 tensor # For each (a,b) in zip(a,b), perform the following: # - Let ai be `a` converted to int32. # - Let af be `a` converted to float. # - Let m be the max of ai and b. # - Return ai and mi. # Do the above 4 elements at a time. (c, d) = tl.inline_asm_elementwise( asm=""" { // Unpack `a` into `ai`. .reg .b8 tmp<4>; mov.b32 {tmp0, tmp1, tmp2, tmp3}, $8; cvt.u32.u8 $0, tmp0; cvt.u32.u8 $1, tmp1; cvt.u32.u8 $2, tmp2; cvt.u32.u8 $3, tmp3; } // Convert `ai` to float. cvt.rn.f32.s32 $4, $0; cvt.rn.f32.s32 $5, $1; cvt.rn.f32.s32 $6, $2; cvt.rn.f32.s32 $7, $3; // Take max of `ai` and `b`. max.f32 $4, $4, $9; max.f32 $5, $5, $10; max.f32 $6, $6, $11; max.f32 $7, $7, $12; """, constraints=( # 8 output registers, namely # $0=ai0, $1=ai1, $2=ai2, $3=ai3, # $4=m0, $5=m1, $6=m2, $7=m3. "=r,=r,=r,=r,=r,=r,=r,=r," # 5 input registers, namely # $8=ai, # $9=b0, $10=b1, $11=b2, $12=b3. # The four elements from `a` are all packed into one register. "r,r,r,r,r"), args=[a, b], dtype=(tl.int32, tl.float32), is_pure=True, pack=4, ) tl.store(C + tl.arange(0, BLOCK), c) tl.store(D + tl.arange(0, BLOCK), d) :param asm: assembly to run. Must match target's assembly format. :param constraints: asm constraints in `LLVM format `_ :param args: the input tensors, whose values are passed to the asm block :param dtype: the element type(s) of the returned tensor(s) :param is_pure: if true, the compiler assumes the asm block has no side-effects :param pack: the number of elements to be processed by one instance of inline assembly :return: one tensor or a tuple of tensors of the given dtypes TF)arithmetic_check allow_lhs_ptr allow_rhs_ptrc3^<"TF"wr\SPV4V4xK$ R#5ir)rU get_result)rarPrcalls& r"rb)inline_asm_elementwise...s(P=OEA*B//=Os*-)r\r]r^rUrr _DtypeClassrhrbinary_op_type_checking_implrrrrrYrrucreate_inline_asmr rUrrj)rrr7rrrrr3has_multiple_outputsres_tysr dispatch_argsbin_op_type_checking broadcast_argrrrPdtrmrrurrs&&&&&&& @r"inline_asm_elementwiser sr s #C&{3K  %D"7+G U # KK-u 5EG=ABTc,,S1TBB}B&  2 2"   &a( !D3DH A}"    $]3&:4&O#  !4HMN"}))99"=GN!./AqxxG/G  $ $Sw]d@e]dWY'AR]d@egn uD dooa('!*55 5PYw=OP5P5PYw=OP PP; $ CO/@es) G&(G=8$H"HH &G:9G:c0]tRtRtRtRRltRtRtRtR#) static_rangei6a Iterator that counts upward forever. .. highlight:: python .. code-block:: python @triton.jit def kernel(...): for i in tl.static_range(10): ... :note: This is a special iterator used to implement similar semantics to Python's :code:`range` in the context of :code:`triton.jit` functions. In addition, it also guides the compiler to unroll the loop aggressively. :param arg1: the start value. :param arg2: the end value. :param step: the step value. Nc L\V\4'g QV R24hVf\^4VnM&\V\4'g QV R24hW0nVf\^4VnWnR#\V\4'g QV R24hWnW nR#)z7 used as tl.static_range start value is not a constexprNz6 used as tl.static_range step value is not a constexprz5 used as tl.static_range end value is not a constexpr)r&rrrr)rHarg1arg2rs&&&&r"rstatic_range.__init__Hs$ **ltf4k,ll* <!! DIdI.. o4&8n0o o.I <"1DJHdI.. n4&8m0n n.JHr%c \R4hz8static_range can only be used in @triton.jit'd functionsrrs&r"rWstatic_range.__iter__WUVVr%c \R4hrrrs&r"__next__static_range.__next__Zrr%)rrrrt rOrkrlrmrNrrWrrnr.r%r"rr6s" WWr%rc0]tRtRtRtRRltRtRtRtR#) ri^a Iterator that counts upward forever. .. highlight:: python .. code-block:: python @triton.jit def kernel(...): for i in tl.range(10, num_stages=3): ... :note: This is a special iterator used to implement similar semantics to Python's :code:`range` in the context of :code:`triton.jit` functions. In addition, it allows user to pass extra attributes to the compiler. :param arg1: the start value. :param arg2: the end value. :param step: the step value. :param num_stages: pipeline the loop into this many stages (so there are :code:`num_stages` iterations of the loop in flight at once). Note this is subtly different than passing :code:`num_stages` as a kernel argument. The kernel argument only pipelines loads that feed into :code:`dot` operations, while this attribute tries to pipeline most (though not all) loads in this loop. :param loop_unroll_factor: Tells the Triton IR level loop unroller how many times to unroll a for loop that this range is used with. Less than 2 for this value implies no unrolling. :param disallow_acc_multi_buffer: If true, prevent the accumulator of the dot operation in the loop to be multi-buffered, if applicable. :param flatten: automatically flatten the loop nest starting at this loop to create a single flattened loop. The compiler will try to pipeline the flattened loop which can avoid stage stalling. :param warp_specialize: Enable automatic warp specialization on the loop. The compiler will attempt to partition memory, MMA, and vector operations in the loop into separate async partitions. This will increase the total number of warps required by the kernel. :param disable_licm: Tells the compiler it shouldn't hoist loop invariant code outside the loop. This is often useful to avoid creating long liveranges within a loop. Note that warp specialization is only supported on Blackwell GPUs and only works on simple matmul loops. Support for arbitrary loops will be expanded over time. Nc  Vf\^4VnMW0nVf\^4VnWnM WnW nW@nWPnW`nWpnWnWn R#r) rrrr num_stagesloop_unroll_factordisallow_acc_multi_bufferflattenwarp_specialize disable_licm) rHrrrr r!r"r#r$r%s &&&&&&&&&&r"rrange.__init__sX <!! DII <"1DJHJH$"4)B& .(r%c \R4hz4tl.range can only be used in @triton.jit'd functionsrrs&r"rWrange.__iter__QRRr%c \R4hr(rrs&r"rrange.__next__r*r%) r%r"rr#r!r rrr$)NNNNFFFFrr.r%r"rr^s)V)&SSr%rc$]tRtRtRtRRltRtR#)rSia  While loop condition wrapper. .. highlight:: python .. code-block:: python @triton.jit def kernel(...): while tl.condition(c, disable_licm) ... :note: This is a special wrapper used to annotate while loops in the context of :code:`triton.jit` functions. It allows user to pass extra attributes to the compiler. :param disable_licm: Tells the compiler it shouldn't hoist loop invariant code outside the loop. This is often useful to avoid creating long liveranges within a loop. c WnW nR#rrSr%)rHrr%s&&&r"rcondition.__init__s (r%r/Nr )rOrkrlrmrNrrnr.r%r"rSrSs"))r%rSc 0V^8dQhRRRRRRRRRR R R /#) r,lib_namer'lib_pathr7rParg_type_symbol_dictrret_typerrrrcr.)r/s"r"r0r0sA%n%nS%nC%nt%nSW%nch%nsw%nr%c \V4^8Xd \R4h\\VP44^,4p\V4V8wd\R\V4 RV 24h.p .p VF|p \ V \ 4'd9V P V P4V P V P4KQV P \V 44V P V 4K~ \V 4p W9d \RVP4 RV 24hWI,^,p VPp \ V!WWVPV 4V4V4#)aT Dispatch a function to a library :param func: the function to dispatch :param lib_name: the name of the library :param lib_path: the path of the library :param args: the arguments of the function :param arg_type_symbol_dict: the type of the arguments :param ret_type: the type of the return value :return: the return value of the function zarg_type_symbol_dict is emptyz+length of input args does not match.Expect rz,input arg type does not match.Expect one of ) rLr5rPrMr&rUrrrrrrrjrur )r?r2r3r7r4r5rr3num_args arg_typesarg_listrsymbolrus&&&&&&&& r"dispatchr;sR  A%8994,1134Q78H 4yH##&t9+VH:?@ @IH c6 " "   SYY ' OOCJJ '   T#Y ' OOC i I,**>*C*C*E)FfYKYZ Z&03##d8vPWAXZabdlmmr%c ,V^8dQhRRRRRRRRRR /#) r,r2r'r3r7rPr4rrrcr.)r/s"r"r0r0s8'q'q'q'q4'qW['qfj'qr%c  VP4pRp.p\P!\V44Fgp VP Wi,4Wi&VP Wi,P 4Wi,PP4'gKeRpKi \V4pW8,^,p \V4^8dRp W9dRp V^,p VFp VPWV R7wrK \P!\V44F"p VPWi,V V R7wWi&pK$ V'gV PPV 4p VPPp\WWW:WE4#)a: Dispatch an elementwise function to a library :param lib_name: the name of the library :param lib_path: the path of the library :param args: the arguments of the function :param arg_type_symbol_dict: the type of the arguments :param is_pure: whether the function is pure :return: the return value of the function TF)r)copyrirrLrhrrrrrr rjrrYrucreate_extern_elementwiser;)r2r3r7r4rr3r  all_scalarr8rPr5rr rrr?s&&&&&& r"extern_elementwiserAsjIIKMJI ^^C . /$..}/?@ )//0   ) ) + +J 0 i I#.q1H 9~  ,$ %a( !D(EEdWg F i A}"M 23A"+"H"HIY[hZj#I#l M a4$))99(CH    6 6D DH=Q]d ppr%c$VPW4#)z Convert both operands to a single common type :param lhs: the left operand :param rhs: the right operand :param builder: the builder )r)rr"semantics&&&r"binary_op_type_legalizationrDs  0 0 ::r%c\V4#)z#A decorator for external functions.)r@rfs&r"externrFs 2;r%cPVR8H;'d\P!RV4^8#)rg?)rcopysignr_s&r"is_negative_zerorI%s$ 8 1 1 c1-11r%ric&\V4p\;QJdRV4F 'dK RM RM !RV44pV'dV\JgQR4h\;QJdRV4F 'gK RM RM !RV44'dQR4h\;QJdRV4F 'gK RM RM !RV44'dQR4h\ \ P !\V444#V\Jd\PpM \R4\V4^8gQR 4hV^,pVR ,Fp\WEWR 7pK VPP4'd \R 4V#) c3L"TFp\V\4'*xK R#5irr&rpr`s& r"rbbuiltin_max..,Cd:a444d"$FTz-propagate_nan is not supported on builtin maxc3N"TFp\P!V4xK R#5irrisnanr`s& r"rbrM/3dtzz!}}d#%z)constexpr max does not support NaN valuesc38"TFp\V4xK R#5irrIr`s& r"rbrM09Dq'**Drdz,constexpr max does not support negative zerozJpassing propagate_nan to builtin max is deprecated, use tl.minimum insteadmin requires at least 2 valuesrMrir3zMbuiltin max on non-scalar tensor values is deprecated, use tl.maximum instead)r\r_NOTHINGanyrrirC PropagateNanNONErrLrrrrr )rir3r7 is_constexprmax_valrs$$* r" builtin_maxr`)/  %D3CdC333CdCCL(Y*YY(33d33333d333`5``339D93339D999i;ii9&:4&@ABB $))  YZ t9>;;;>1gGBxx'mY|| \] Nr%c&\V4p\;QJdRV4F 'dK RM RM !RV44pV'dV\JgQR4h\;QJdRV4F 'gK RM RM !RV44'dQR4h\;QJdRV4F 'gK RM RM !RV44'dQR4h\ \ P !\V444#V\Jd\PpM \R4\V4^8gQR 4hV^,pVR ,Fp\WEWR 7pK VPP4'd \R 4V#) c3L"TFp\V\4'*xK R#5irrLr`s& r"rbbuiltin_min..DrNrOFTz-propagate_nan is not supported on builtin minc3N"TFp\P!V4xK R#5irrQr`s& r"rbrdGrSrTz)constexpr min does not support NaN valuesc38"TFp\V4xK R#5irrVr`s& r"rbrdHrWrdz,constexpr min does not support negative zerozJpassing propagate_nan to builtin min is deprecated, use tl.minimum insteadrXrMrYzMbuiltin min on non-scalar tensor values is deprecated, use tl.minimum instead)r\rrZr[rrirNr\r]rrLrmrrr )rir3r7r^min_valrs$$* r" builtin_minrhArar%rqr)FrNrtrs)NNNNNNrr)rNr ru)NNNN)TN)FNN)rDN)rDNN) __future__rrwarningsr contextlibrenumr functoolsrrr rUr r r r rrrr dataclassesrrirDr runtime.jitrr> _C.libtritonr_utilsrrrrrr= PROPAGATE_NANr\rr@rZr`rTrhrjrprqrrrvr\rerurrrr`rrrrwrrrrrrrrrfloat8e5 float8e5b16 float8e4nv float8e4b8 float8e4b15float16rrfloat64pi32_trrUrxrjrrrrrr0rAtyping_extensionsrFrRrYr`rrrrrrrrrr\rrrrNrrrrrrrrrrrrrrrrrrrrr-r0rr4rr8rErrrrrrrrrRr\rrr r]rmrrrurrrrlrrrr)r$rrrrrrrrrrYrrrrSr;rArDrFrZrIr`rhr.r%r"r|s " %55 [[[!%ii CL%  "*Z ..  " "   " """,-Y-8O- O-dl 6LILd %.5%.P3!3!l$F$FN!! V} V} V} g g g g x x x >J 9  9  J - = - - e U8] Z] @ uH>JH>V!!/)/:[NZ[N|k8k6*.*B $.i..6 5X)4&{'{|$2[ PU    h*h* h*V#' 89  & &$ ( (  ( ( &'>%>? 4   / /,  8  88$8 8&*$* *<*D* *, 6 6D    .   B=$= =(G G8u888 8*. 8> >< D DN KO[`mq'T+7 +7f >$ >$B 3 3  ; ; 2b 2bj  ^  ^Q a  a 9_ 9_B>&5?6 ?JA  AE@ @E@ @E@ @M"@# @L!?" ?M"@# @ , ,,    < 2 2  2 2  2 2  -9->->$2 2* -9->->$2 2* 2>2C2Ct7 7DD(( . .b      ,0#Q #QL6 6$. ." H. H. H.  H. H.` % % 0 04 3 36 2 2, 7 7   S  T    E  ]a            -9E-9T-9 -9` I I2 ~Q ~QL%W:%WPCSJCSL) )8%nP 'q 'qT; 82 X . X r%