OtjldddlmZddlZddlmZddlmZddlmZddl m Z m Z m Z ddl Z ddl mZmZmZmZmZmZmZddlmZddlZd d lmZd d lmZddlZd d lmZd d lm Z m!Z!m"Z"m#Z#edZ$dZ%ej&Z'ddZ(ddZ)ddZ*dZ+ddZ,e)ddZ-GddZ.GddZ/GddZ0Gd d!e0Z1Gd"d#e/Z2e2dZ3d$Z4d%Z5d&Z6Gd'd(e0Z7e7Z8Gd)d*e7Z9Gd+d,e7Z:Gd-d.e0Z;Gd/d0e7Z<e7d1Z=e7d2Z>e7d3Z?e7d4Z@e7d5ZAe7d6ZBe7d7ZCe7d8ZDe7d9ZEe7d:ZFe7d;ZGe7d<ZHe7d=ZIe7d>ZJe7d?ZKe7d@ZLe7dAZMe7dBZNe7dCZOe9eAZPddGZQGdHdIe/ZRddJZSGdKdLe/ZTGdMdNZUGdOdPe0ZVGdQdRe/ZWGdSdTeVZXGdUdVeWZYedWGdXdYe0ZZdZZ[e)dd[Z\e)dd\Z]e)dd]Z^d^e d_e^__d`Z`daZae)ddbZbe)ddcZce*e)ddddeZde*e)dddddgZee*e)ddddhZfe)ddjZge)ddkZhddlZie*e)dddnZje*e)ddddoZke*e)ddpZle*e)didddqdrZmdsZne*e)ddtZoe*e)dddxZpe)ddddeNdfdyZqe)ddiddeNdfdzZre) dd}Zse) dddZte) dddZue*e)ddZve)dddZwe(de*e)ddZxe) dddZydddZze*e)ezddddZ{e*e)ezdddZ|e*e)ezdddZ}e*e)ezdddZ~e*e)ezdddZe*e)ezdddZe*e)ezdddZe*e)ezdddZe)ddZe)dddZe)dddZe)dddZe)e'jdfddZe)e'jdfddZe)e'jdfddZ dddZedZe*e)ddZe)ddZe)ddZdddZe*e)ddZe*e)ddZe*e)ddZe)ddddddZe)ddZe)dd„Ze)ddÄZe)ddĄZe)ddńZe)ddddiddȜdd˄Ze)d d̄Ze)didd͜d΄Ze)d dτZe) dd dքZGdׄde/ZGdلde/ZGdۄde/Zd dZe) dd dZdZdZeZdZe)edddZe)edddZdS() annotationsN)warn)contextmanager)Enum)partialwrapscached_property)UnionCallableListSequenceTypeVarOptionalTuple) dataclass)knobs) JITCallable)ir)TRITON_MAX_TENSOR_NUMELvalidate_block_shapeget_primitive_bitwidth _tuple_createT__triton_builtin__TcJttrfdS|_S)z9If the result of this function is unused, throw an error.c$t|SN)must_use_result)fnxs R/var/www/html/CA-Chatbot/venv/lib/python3.11/site-packages/triton/language/core.pyz!must_use_result..s?2q11) isinstancestr_must_use_result)r!ss` r"rrs1!S311112A Hr$r returnctsJtfd}t|tdt j|_|S)zMark a function as a builtin.cJd|vs|dtd|i|S)N _semanticzeDid you forget to add @triton.jit ? (`_semantic` argument must be provided outside of JIT functions.) ValueErrorargskwargsr s r"wrapperzbuiltin..wrapper&sH f $ ${(;(Cabb br4"6"""r$T)callablersetattrTRITON_BUILTINinspect signature)r r2s` r"builtinr8"si B<<< 2YY####Y#  G^T***)"--G Nr$c tsJtj}t|jddhz dk}jsd_xjdjd|rdnddjd |rd ndd z c_fd }t|j }|d  d|d <| |}||_ ||_djd|_trt|tdttjt!t"rn|S)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. r, _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|i|Srr/s r"r2z"_tensor_member_fn..wrapperMsr4"6"""r$rselfname) parameterszForwards to :py:func:`z` free functionT)r3r6r7lenrCkeys__doc____name__listvaluesreplace __signature__ is_builtinr4r5tensorr%r)r orig_sighas_argsr2 new_paramsnew_sigs` r"_tensor_member_fnrR3s B<<< $$H8&++--l0KKLLqPH : JJ(08uub K'/677BJJ ##### h)002233JqM))v)66JqM*55G#GGKr{KKKGO"~~/... FBKz"k'B'B!OPPP Ir$ct|dkr. t|d|dS#t$rYnwxYw|S)z7Returns x[0] if x has one element and x[0] is iterable.r;r)rDiter TypeErrorr!s r"_unwrap_iterablerW`sU 1vv{{  1JJJQ4K    D  Hs 2 ??boolc.t|tdS)z-Is this a registered triton builtin function?F)getattrr5r s r"rLrLvs 2~u - --r$c,||Sr to_tensor)r!r,s r"r^r^{s   q ! !!r$ceZdZdZdS)consta~ 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. N)rG __module__ __qualname__rFr?r$r"r`r`s Dr$r`c&eZdZUdZded<d dZd S) base_valuezLBase class of values that exist in the triton IR (i.e. not constexprs). base_typetypehandlesList[ir.value]r)Nonect)znFlatten frontend value into a sequence of mlir handles, which are appended to the output list NotImplementedErrorr@rgs r" _flatten_irzbase_value._flatten_irs "!r$Nrgrhr)ri)rGrarbrF__annotations__rnr?r$r"rdrds=OOO""""""r$rdc6eZdZddZddZdd Zdd ZddZdS)rer)rXc td)NzTypes must implement __eq__rkr@others r"__eq__zbase_type.__eq__s!"?@@@r$c||k Srr?rss r"__ne__zbase_type.__ne__EM""r$rgrhcursorintTuple[base_value, int]ct)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. rkr@rgrys r" _unflatten_irzbase_type._unflatten_irs "!r$r&c0td|j)NzNYI: Type mangling for type )rl __class__r@s r"manglezbase_type.mangles!"Q"Q"QRRRr$builder ir.builderout List[ir.type]rictrrkr@rrs r"_flatten_ir_typeszbase_type._flatten_ir_typess!!r$Nr)rXrgrhryrzr)r{r)r&rrrrr)ri)rGrarbrurwr~rrr?r$r"reresAAAA####""""SSSS""""""r$rec@eZdZdZdZddZdZddZdd ZddZ dS)constexpr_typec||_dSrvaluer@rs r"__init__zconstexpr_type.__init__s  r$cLt|to|j|jkSr)r%rrrss r"ruzconstexpr_type.__eq__s %00NTZ5;5NNr$r)r&cd|jdS)Nzconstexpr_type[]rrs r"__repr__zconstexpr_type.__repr__s.....r$c*t|jSr)hashrrs r"__hash__zconstexpr_type.__hash__DJr$ct|jdr|j}nt|j}d|S)Nrc)hasattrrrrepr)r@vals r"rzconstexpr_type.manglesF 4:x ( ( #*##%%CCtz""C3yyr$rrrrricdSrr?rs r"rz constexpr_type._flatten_ir_typesr$rgrhryrzr{c.t|j|fSr constexprrr}s r"r~zconstexpr_type._unflatten_irs$$f,,r$Nrrr) rGrarbrrurrrrr~r?r$r"rrsOOO////   ------r$rceZdZdZdZd3dZdZd4d Zd Zd Z d Z dZ dZ dZ dZdZdZdZdZdZdZdZdZdZdZdZdZdZdZd Zd!Zd"Zd#Z d$Z!d%Z"d&Z#d'Z$d(Z%d)Z&d*Z'd+Z(d,Z)d-Z*d.Z+d/Z,d0Z-d1Z.d2S)5rzL This class is used to store a value that is known at compile-time. ct|tr|j}t|t||_t||_dSr)r%rrrrfrs r"rzconstexpr.__init__sI ** KE **  "5)) r$r)r&cd|jdS)Nz constexpr[rrrs r"rzconstexpr.__repr__s)DJ))))r$c8t|j|jfSr)rrrfrs r"rzconstexpr.__hash__sTZ+,,,r$rgrhricdSrr?rms r"rnzconstexpr._flatten_irrr$c|jSrrrs r" __index__zconstexpr.__index__ zr$cJt|jt|zSrrr_unwrap_if_constexprrss r"__add__zconstexpr.__add__ &:5&A&AABBBr$cJtt||jzSrrrrrss r"__radd__zconstexpr.__radd__ -e44tzABBBr$cJt|jt|z Srrrss r"__sub__zconstexpr.__sub__rr$cJtt||jz Srrrss r"__rsub__zconstexpr.__rsub__rr$cJt|jt|zSrrrss r"__mul__zconstexpr.__mul__rr$cJt|jt|zSrrrss r"__mod__zconstexpr.__mod__rr$cJtt||jzSrrrss r"__rmul__zconstexpr.__rmul__rr$cJt|jt|z Srrrss r" __truediv__zconstexpr.__truediv__rr$cJtt||jz Srrrss r" __rtruediv__zconstexpr.__rtruediv__rr$cJt|jt|zSrrrss r" __floordiv__zconstexpr.__floordiv__ ';E'B'BBCCCr$cJtt||jzSrrrss r" __rfloordiv__zconstexpr.__rfloordiv__s -e44 BCCCr$cLt|jt|kSrrrss r"__gt__zconstexpr.__gt__  &:5&A&AABBBr$cLtt||jkSrrrss r"__rgt__zconstexpr.__rgt__  -e44tzABBBr$cLt|jt|kSrrrss r"__ge__zconstexpr.__ge__ ';E'B'BBCCCr$cLtt||jkSrrrss r"__rge__zconstexpr.__rge__ -e44 BCCCr$cLt|jt|kSrrrss r"__lt__zconstexpr.__lt__rr$cLtt||jkSrrrss r"__rlt__zconstexpr.__rlt__rr$cLt|jt|kSrrrss r"__le__zconstexpr.__le__rr$cLtt||jkSrrrss r"__rle__zconstexpr.__rle__rr$cLt|jt|kSrrrss r"ruzconstexpr.__eq__"rr$cLt|jt|kSrrrss r"rwzconstexpr.__ne__%rr$c*t|jSr)rXrrs r"__bool__zconstexpr.__bool__(rr$c,t|j Srrrs r"__neg__zconstexpr.__neg__+$*%%%r$cJt|jt|zSrrrss r"__and__zconstexpr.__and__.rr$cHt|jot|Srrrss r" logical_andzconstexpr.logical_and1s C(z(constexpr.__getitem__..Ys+HHA$Q''HHHHHHr$)_normalize_tupler __getitem__)r@r0s r"r zconstexpr.__getitem__Xs0HH1A$1G1GHHH%tz%t,,r$Nrro)/rGrarbrFrrrrnrrrrrrrrrrrrrrrrrrrrrurwrrrrrrrrrrrrrrrrr r?r$r"rrs*** ****---CCCCCCCCCCCCCCCCCCCCCCCCCCCDDDDDDCCCCCCDDDDDDCCCCCCDDDDDDDDDDDD   &&&CCCEEECCCCCCDDD&&&&&&BBBBBBDDDDDD)))   )))-----r$rcPt|tr d|DSt|tjrt |d|DSt|trtd|D|jSt|t r|jn|S)Nc,g|]}t|Sr?rrs r" z(_unwrap_if_constexpr..bs!333A$Q''333r$c,g|]}t|Sr?rrs r"r z(_unwrap_if_constexpr..ds! D D DQ!5a!8!8 D D Dr$c,g|]}t|Sr?rrs r"r z(_unwrap_if_constexpr..fs!999!*1--999r$)r%rHbuiltinstuplerrfrr)os r"rr`s!T4333333!X^$$FQ D D! D D DEEE!UC99q99916BBB I.. 5177A5r$ct|}t|ttjfrt |}|Sr)rr%rHrr)tnormalized_tuples r"r r js?+A.."T8>$:;;3 !122 r$c t|trVt|trC|jjj}|j|kr)td|jd|d|jddSdSdSdS)NzValue z exceeds the maximum bitwidth (z ) for type 'z)'. This may result in undefined behavior.) r%rMrrfscalarprimitive_bitwidthrrdtype)r shift_valuebitwidths r"check_bit_widthrqs%  Z Y%G%G:$7   ( ( X*XX8XXafalXXX      ( (r$ceZdZgdZgdZgdZgdZdgZGddeZ Gdd eZ d Z d Z d Z d ZdZdZdZdZdZdZdZdZdZdZdZdZdZdZdZdZdZdZ d Z!d!Z"d"Z#d#Z$d$Z%d%Z&d&Z'e(d'Z)e(d(Z*e(d)Z+e(d*Z,e(d+Z-dGd.Z.d/Z/e0d0Z1dHd6Z2dId8Z3d9Z4d:Z5e0dJd<Z6d=Z7dKdCZ8dJfdD Z9dLdFZ:xZ;S)Mr)int8int16int32int64)int1uint8uint16uint32uint64) fp8e4b15fp8e4nvfp8e4b8fp8e5fp8e5b16fp16bf16fp32fp64)r,r-r.r/voidceZdZdZdZdS)dtype.SIGNEDNESSrr;N)rGrarbSIGNEDUNSIGNEDr?r$r" SIGNEDNESSr2sr$r5ceZdZdZdZdZdS) dtype.KINDrr;rN)rGrarbBOOLEANINTEGRALFLOATINGr?r$r"KINDr7sr$r;ct|}||_|tjtjztjztjzvs J|t||_|jdz|_ |tjvr$tj j |_ |j|_ dS|tjvr$tj j|_ |j|_ dS|tjvr|dkrd|_d|_dS|dkrd|_d|_dS|dkrd|_d|_dS|dkrd |_d|_dS|d krd |_d |_dS|d krd |_d|_dS|dkrd|_d|_dS|dkrd|_d|_dS|dkrd|_d|_dSt#d|dS)Nr'r(r)r*rr+r, r-r.r/4iz Unsupported floating-point type )rrBr SINT_TYPES UINT_TYPESFP_TYPES OTHER_TYPESrritemsizer5r3int_signedness int_bitwidthr4fp_mantissa_width exponent_bias RuntimeError)r@rBs r"rzdtype.__init__s#D)) u'%*::U^KeN_____ae___"8">">/14 5# # #"'"2"9D  $ 7D    U% % %"'"2";D  $ 7D    U^ # #z!!)*&%'""""")*&%&""""")*&%&""")*&%'"""##)*&%'""")+&%'""")*&%(""")+&%(""")+&%)""""#Ld#L#LMMM;$ #r$cd|jvS)Nfp8rArs r"is_fp8z dtype.is_fp8s !!r$c|jdkS)Nr(rArs r" is_fp8e4nvzdtype.is_fp8e4nvyI%%r$c|jdkS)Nr)rArs r" is_fp8e4b8zdtype.is_fp8e4b8rUr$c|jdkS)Nr'rArs r" is_fp8e4b15zdtype.is_fp8e4b15yJ&&r$c|jdkS)Nr*rArs r"is_fp8e5zdtype.is_fp8e5yG##r$c|jdkS)Nr+rArs r" is_fp8e5b16zdtype.is_fp8e5b16rZr$c|jdkS)Nr,rArs r"is_fp16z dtype.is_fp16yF""r$c|jdkS)Nr-rArs r"is_bf16z dtype.is_bf16rbr$c|jdkS)Nr.rArs r"is_fp32z dtype.is_fp32rbr$c|jdkS)Nr/rArs r"is_fp64z dtype.is_fp64rbr$c|jdkS)Nr"rArs r"is_int1z dtype.is_int1rbr$c|jdkS)NrrArs r"is_int8z dtype.is_int8rbr$c|jdkS)NrrArs r"is_int16zdtype.is_int16r]r$c|jdkS)Nr rArs r"is_int32zdtype.is_int32r]r$c|jdkS)Nr!rArs r"is_int64zdtype.is_int64r]r$c|jdkS)Nr#rArs r"is_uint8zdtype.is_uint8r]r$c|jdkS)Nr$rArs r" is_uint16zdtype.is_uint16yH$$r$c|jdkS)Nr%rArs r" is_uint32zdtype.is_uint32rwr$c|jdkS)Nr&rArs r" is_uint64zdtype.is_uint64rwr$c(|jtjvSr)rBrrHrs r" is_floatingzdtype.is_floatingsyEN**r$c(|jtjvSr)rBrSTANDARD_FP_TYPESrs r"is_standard_floatingzdtype.is_standard_floatingsyE333r$c(|jtjvSr)rBrrFrs r" is_int_signedzdtype.is_int_signedyE,,,r$c(|jtjvSr)rBrrGrs r"is_int_unsignedzdtype.is_int_unsignedrr$cB|jtjtjzvSr)rBrrFrGrs r"is_intz dtype.is_intsyE,u/????r$c*|Sr)rjrs r"is_boolz dtype.is_bool||~~r$c|rtjjS|rtjjS|sJtjjSr)rrr;r8rr9r}r:rs r"kindz dtype.kinds[ <<>> ':% % [[]] ':& &##%% % %%:& &r$c|rd|jdz zdz S|r d|jzdz SJ)Nrr;rrLrrs r"get_int_max_valuezdtype.get_int_max_valuesZ      2t(1,-1 1    ! ! ,d''!+ +ur$cv|rd|jdz z S|rdSJ)Nrr;rrrs r"get_int_min_valuezdtype.get_int_min_valuesJ      /)A-.. .    ! ! 1ur$cl|tjtjztjztjzvSr)rrFrGrHrI)type_strs r"is_dtypezdtype.is_dtypes(5+e.>>ORWRccccr$c td)NzNot implementedrOr?r$r"is_voidz dtype.is_void!s,---r$cdSNFr?r?r$r"is_blockzdtype.is_block%ur$cdSrr?r?r$r"is_ptrz dtype.is_ptr)rr$cdSrr?r?r$r"is_constzdtype.is_const-rr$r)rXcnt|}t|tsdS|j|jkSr)rr%rrBrss r"ruz dtype.__eq__1s5$U++%'' 5yEJ&&r$c,t|jfSr)rrBrs r"rzdtype.__hash__7sTYM"""r$c|Srr?rs r"rz dtype.scalar: r$rrrrricV|||dSr)appendto_irrs r"rzdtype._flatten_ir_types>s& 4::g&&'''''r$ir.typecr|jdrBt|dr2|j|jjvrt d|d|jj|jdkr|S|jdkr|S|jdvr|S|jdvr| S|jd vr| S|jd vr| S|jd kr| S|jd kr| S|jd kr|S|jdkr|S|jdkr|S|jdkr|S|jdkr|S|jdkr|S|jdkr|St d|d)NrQoptionsztype zB not supported in this architecture. The supported fp8 dtypes are r0r")rr#)rr$)r r%)r!r&r*r+r(r)r'r,r-r.r/zfail to convert z to ir type)rB startswithrrsupported_fp8_dtypesr. 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_tyr@rs r"rz dtype.to_irAsm 9   & & iw ** ity@d/d/d "h"h"hAHAe"h"hiii 9  &&(( ( Y& &&(( ( Y+ + +&&(( ( Y- - -'')) ) Y- - -'')) ) Y- - -'')) ) Y' ! !'')) ) Y* $ $**,, , Y) # #))++ + Y) # #))++ + Y* $ $**,, , Y& &&(( ( Y& &&(( ( Y& '')) ) Y& ((** *=D===>>>r$c|jSrrArs r"__str__z dtype.__str__g yr$c|jdrd|jddzS|jdrd|jddzS|jS)Nfpfloatrbfbfloat)rBrrs r" codegen_namezdtype.codegen_namejsd 9   % % TYqrr]* * Y ! !$ ' ' dim+ +9 r$r&c|jS)z"See cache_key_part() in triton.cc.rArs r"cache_key_partzdtype.cache_key_partrs yr$c0d|S)z4Output of repr needs to be an evaluatable expressionztriton.language.)rrs r"rzdtype.__repr__ws7$"3"3"5"5777r$rgrhryrzr{c8t||||dzfSNr;)rMr}s r"r~zdtype._unflatten_ir{sgfot,,fqj88r$cL|r7tjj}|j|krdnd}|t |jzS|rt |S|rdSt S)NiuV) rrr5r3rKr&rLr}rsuperr)r@r3prefixrs r"rz dtype.mangle~s ;;== 3%,F /699SSsFC 1222 2      t99  <<>> 3ww~~r$ element_tyc2|rJ|Sr)r)r@rs r"with_element_tyzdtype.with_element_tys==??"""r$rr)rrr)rrr)rr)zconst_pointer<) rr%rrUrfrGrrr`rB)r@rrr`s r"rzpointer_type.__init__s)*55 *e,, ecD4D4D4Mcccdd d$* 49],z,,,,?]PZ?]?]?] r$rrr)ir.pointer_typech||j||jSr) get_ptr_tyrrrrs r"rzpointer_type.to_irs+!!$/"7"7"@"@$BTUUUr$c|jSrrArs r"rzpointer_type.__str__rr$c*|Srrrs r"rzpointer_type.__repr__rr$cdSNTr?rs r"rzpointer_type.is_ptrtr$c|jSr)r`rs r"rzpointer_type.is_constrr$ct|}t|tsdS|j|jko|j|jko|j|jkSr)rr%rrrr`rss r"ruzpointer_type.__eq__sX$U++%.. 5%"22~t7IUM`7`~eieosxs~e~~r$c|Srr?rs r"rzpointer_type.scalarrr$r&c:d|jS)NP)rrrs r"rzpointer_type.mangles-4?))++---r$N)r;F)rrrrzr`rX)rrr)rrr) rGrarbrrrrrrrurrrr?r$r"rrs^^^^^VVVV X......r$rc|eZdZddZdd Zd Zd Zd ZddZddZ ddZ e dZ e dZ ddZdS) block_typerrshaper c ||_t|ttfsJtt ||_|jst dt|j|_d|jd|jd|_ dS)Nz0d block_type is forbidden<, r) rr%rHr _unwrap_shaperrUrnumelrB)r@rrs r"rzblock_type.__init__s$54-00110=//00 z :899 9)$*55 8 88do888 r$rrr) ir.block_typech||j||jSr) get_block_tyrrrrs r"rzblock_type.to_irs*##DO$9$9'$B$BDJOOOr$c|jSrrArs r"rzblock_type.__str__rr$c*|Srrrs r"rzblock_type.__repr__rr$cdSrr?rs r"rzblock_type.is_blockrr$ Tuple[int]c|jSrrrs r"get_block_shapeszblock_type.get_block_shapesrr$ scalar_tyc,t||jSr)rr)r@rs r"rzblock_type.with_element_tys)TZ000r$rXcpt|tsdS|j|jko|j|jkSr)r%rrrrss r"ruzblock_type.__eq__s7%,, 5%"22PtzU[7PPr$c|jSr)rrs r"rzblock_type.scalars r$c0|j|jjdzzS)Nr=)rrrrs r"nbyteszblock_type.nbytesszT_?1DEEr$r&c|j}dtt|j}|d|dS)N_S)rrjoinmapr&r)r@eltrs r"rzblock_type.manglesIk  ""S$*--..      r$N)rrrr )rrr)r)r)r)rrr)rrr)rGrarbrrrrrrrrurrrrr?r$r"rrs 9 9 9 9PPPP1111QQQQ XFFXF!!!!!!r$rc\eZdZddZedZdZdZdd ZddZ dZ ddZ dZ dS) tuple_typeNc"||_||_dSr)typesfields)r@r r s r"rztuple_type.__init__s  r$c|j*ddd|jDzdzSdddt|j|jDzdzS)N[,c34K|]}t|VdSrr&)rvs r"rz"tuple_type.name..s(!=!=Q#a&&!=!=!=!=!=!=r$rc"g|] \}}|d| S):r?)rkrs r"r z#tuple_type.name..s&SSSda!zzazzSSSr$)r rr ziprs r"rBztuple_type.namesq ; !=!=$*!=!=!====C CSXXSSc$+tz6R6RSSSTTTWZZZr$c|jSrrArs r"rztuple_type.__str__rr$c*t|jSr)rTr rs r"rztuple_type.__iter__rr$rrrrcF|jD]}|||dSr)r r)r@rrtys r"rztuple_type._flatten_ir_typess6* / /B  # . . . . / /r$indexrzr)rc|j|Sr)r )r@rs r"r ztuple_type.__getitem__sz%  r$c~t|t|uo|j|jko|j|jkSr)rfr r rss r"ruztuple_type.__eq__ s7DzzT%[[(fTZ5;-Ff4;Z_ZfKffr$rgrhryTuple[tuple, int]cg}|jD]0}|||\}}||1t|||fSr)r r~rr)r@rgryrIrrs r"r~ztuple_type._unflatten_irs[* ! !B,,Wf==ME6 MM% VT""F**r$cVddd|jDzdzS)Nrrc3>K|]}|VdSrr)rrs r"rz$tuple_type.mangle..s*??bbiikk??????r$)rr rs r"rztuple_type.mangles/SXX??DJ??????#EEr$r)rrrr)rrzr)r)rgrhryrzr)r) rGrarbrr rBrrrr rur~rr?r$r"r r s[[_[    ////!!!!ggg++++FFFFFr$r ceZdZdZdS) slice_typecd|_dS)Nr$rArs r"rzslice_type.__init__s   r$NrGrarbrr?r$r"r$r$s#!!!!!r$r$r0r"rrr r!r#r$r%r&r*r+r(r)r'r,r-r.r/rrzsignedc6|dkrtS|dkr |rtS|dkr |stS|dkr |rtS|dkr |stS|dkr |rt S|dkr |st S|dkr |rtS|dkr |stStd|d|)Nr;r=rA @zUnsupported bitwidth z and signedness ) r"rr#rr$r r%r!r&r.)rr's r" get_int_dtyper+7s1}} Q6 Qv RF R RF R RF R SSS6SSTTTr$ceZdZdZdpfd Zdqd Zdrd Zedsd ZedsdZ edsdZ edsdZ edsdZ edsdZ edsdZedsdZedsdZedsdZedsdZedsdZedsdZedsdZedsdZedsdZedsdZedsdZedsdZedsd Zedsd!Zedsd"Zedsd#Zedsd$Zedsd%Z edsd&Z!edsd'Z"edsd(Z#edsd)Z$edsd*Z%edsd+Z&edsd,Z'edsd-Z(edsd.Z)edsd/Z*edsd0Z+edsd1Z,edsd2Z-edsd3Z.edsd4Z/e0d5Z1edtdud;Z2dvd<Z3dvd=Z4dvd>Z5dwd@Z6dvdAZ7dvdBZ8dvdCZ9dxdvdDZ:dydvdGZ;dvdHZd{dvdKZ?d{dvdLZ@d{dvdMZAd{dvdNZBd{dvdOZCd{dvdPZDdvdQZEdvdRZFdvdSZGdvdTZHdvdUZIdvdVZJdvdWZKd|dvdXZLd|dvdYZMdvdZZNdvd[ZOdvd\ZPdvd]ZQd}dvd^ZRdvd_ZSd~dvdaZTddvdbZUd~dvdcZVddvddZWdtdvdeZXdxdvdfZYdxdvdgZZddvdiZ[ddvdjZ\d e]fddnZ^dsdvdoZ_xZ`S)rMaRepresents 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. :) rfrc\t||_|r|jnd|_t t j|j|_||_ |j |_ td|jD|_dS)Not called by user code.r?c,g|]}t|Sr?rrr(s r"r z#tensor.__init__..rs===QIaLL===r$N) rrhandlerrrmathprodrrfrrr)r@r2rfrs r"rztensor.__init__gs  #'==??:TZZ ty4455  [ ==$*===>> r$rgrhr)ric:||jdSrrr2rms r"rnztensor._flatten_irtt{#####r$r&ct|jdzdd|jDzdzS)Nrrc34K|]}t|VdSrrr1s r"rz!tensor.__str__..ys(0L0LAQ0L0L0L0L0L0Lr$r)r&rrrrs r"rztensor.__str__ws<4:$tyy0L0L0L0L0L'L'LLsRRr$Nc(t||d|SNT)sanitize_overflowr,addr@rtr,s r"rztensor.__add__{4$)LLLLr$c(t||d|Sr;r=r?s r"rztensor.__radd__5$$)LLLLr$c(t||d|Sr;subr?s r"rztensor.__sub__r@r$c(t||d|Sr;rDr?s r"rztensor.__rsub__rBr$c(t||d|Sr;mulr?s r"rztensor.__mul__r@r$c(t||d|Sr;rHr?s r"rztensor.__rmul__rBr$cLt|}|||Srrtruedivr?s r"rztensor.__truediv__s%$U++  u---r$cLt|}|||SrrLr?s r"rztensor.__rtruediv__s%$U++  ---r$cLt|}|||Srrfloordivr?s r"rztensor.__floordiv__s%$U++!!$...r$cLt|}|||SrrPr?s r"rztensor.__rfloordiv__s%$U++!!%...r$cLt|}|||Srrmodr?s r"rztensor.__mod__#$U++}}T5)))r$cLt|}|||SrrTr?s r"__rmod__ztensor.__rmod__#$U++}}UD)))r$c,||Sr)minusr@r,s r"rztensor.__neg__st$$$r$c,||Sr)invertr\s r"rztensor.__invert__s%%%r$cLt|}|||Srrand_r?s r"rztensor.__and__#$U++~~dE***r$cLt|}|||Srr`r?s r"__rand__ztensor.__rand__#$U++~~eT***r$cLt|}|||Srror_r?s r"rz tensor.__or__rVr$cLt|}|||Srrgr?s r"__ror__ztensor.__ror__rYr$cLt|}|||Srrxor_r?s r"rztensor.__xor__rbr$cLt|}|||Srrlr?s r"__rxor__ztensor.__rxor__rer$clt||t|}|||Srrrshlr?s r"rztensor.__lshift__s3e$$$$U++}}T5)))r$clt||t|}|||Srrqr?s r" __rlshift__ztensor.__rlshift__s3t$$$$U++}}UD)))r$ct||t|}|jr|||S|||Srrrrrashrlshrr?s r"rztensor.__rshift__s]e$$$$U++ : # # % % />>$.. .>>$.. .r$ct||t|}|jr|||S|||Srrvr?s r" __rrshift__ztensor.__rrshift__s]t$$$$U++ : # # % % />>%.. .>>%.. .r$cX||}|||Srr^ greater_thanr?s r"rz tensor.__gt__s+##E**%%dE222r$cX||}|||Srr|r?s r"rztensor.__rgt__s+##E**%%eT222r$cX||}|||Srr^ greater_equalr?s r"rz tensor.__ge__s+##E**&&tU333r$cX||}|||Srrr?s r"rztensor.__rge__ s+##E**&&ud333r$cX||}|||Srr^ less_thanr?s r"rz tensor.__lt__+##E**""4///r$cX||}|||Srrr?s r"rztensor.__rlt__+##E**""5$///r$cX||}|||Srr^ less_equalr?s r"rz tensor.__le__+##E**##D%000r$cX||}|||Srrr?s r"rztensor.__rle__s+##E**##E4000r$cX||}|||Srr^equalr?s r"ruz tensor.__eq__%s)##E**tU+++r$cX||}|||Srrr?s r"__req__ztensor.__req__*s)##E**ud+++r$cX||}|||Srr^ not_equalr?s r"rwz tensor.__ne__/rr$cX||}|||Srrr?s r"__rne__ztensor.__rne__4rr$cX||}|||Sr)r^rr?s r"rztensor.logical_and9s+##E**$$T5111r$cX||}|||Sr)r^rr?s r"rztensor.logical_or>rr$c,||Sr)not_r\s r"rztensor.__not__Es~~d###r$ct|tjttfs||g}t|tr|j}|}t |D]\}}t||||}+t|tjtfr,td|j |j |j fDrxtd||S)Nc38K|]}t|duVdSrr)rargs r"rz%tensor.__getitem__..SsQA_A_:=(--5A_A_A_A_A_A_r$zunsupported tensor index: )r%rslicerrrI enumerater expand_dimsallstartstopstepr.)r@slicesr,retdimsls r"r ztensor.__getitem__Is  fx~ui@ A A V^XF fe $ $ #]F (( D DGC#B''/++C55B 788 DSA_A_BD(BGUWU\A]A_A_A_>_>_ D !Bb!B!BCCC r$cJd)zTransposes a 2D tensor.Fz0Transposition must be created by the AST Visitorr?rs r"rztensor.TZs IHHHur$Ffp_downcast_rounding Optional[str]bitcastrXc*t|||||S)z3 Alias for :py:func:`tensor.cast`. r,)cast)r@rrrr,s r"toz tensor.to_s D%!5w)TTTTr$cdSrr?r@rs r" broadcast_toztensor.broadcast_tok r$cdSrr?r@dimss r"transz tensor.transnrr$cdSrr?rs r"permuteztensor.permuteqrr$tuple[tensor, tensor]cdSrr?rs r"splitz tensor.splittrr$cdSrr?rs r"viewz tensor.viewwrr$cdSrr?rs r"reshapeztensor.reshapezrr$cdSrr?)r@axiss r"rztensor.expand_dims}rr$cdSrr?)r@rrrs r"rz tensor.castrr$r?r<cdSrr?)r@rmaskboundary_checkcache_modifiereviction_policys r"storez tensor.storerr$cdSrr?)r@offsetss r"advanceztensor.advancerr$cdSrr?)r@cmprsemscopes r" atomic_casztensor.atomic_casrr$cdSrr?r@rrrrs r" atomic_xchgztensor.atomic_xchgrr$cdSrr?rs r" atomic_addztensor.atomic_addrr$cdSrr?rs r" atomic_maxztensor.atomic_maxrr$cdSrr?rs r" atomic_minztensor.atomic_minrr$cdSrr?rs r" atomic_andztensor.atomic_andrr$cdSrr?rs r" atomic_orztensor.atomic_orrr$cdSrr?rs r" atomic_xorztensor.atomic_xorrr$cdSrr?rs r"expz tensor.exprr$cdSrr?rs r"logz tensor.logrr$cdSrr?rs r"cosz tensor.cosrr$cdSrr?rs r"sinz tensor.sinrr$cdSrr?rs r"sqrtz tensor.sqrtrr$cdSrr?rs r"rsqrtz tensor.rsqrtrr$cdSrr?rs r"absz tensor.absrr$cdSrr?)r@r combine_fn keep_dimss r"reducez tensor.reducerr$cdSrr?)r@rrreverses r"associative_scanztensor.associative_scanrr$cdSrr?)r@indicesrs r"gatherz tensor.gatherrr$cdSrr?)r@num_binss r" histogramztensor.histogramrr$cdSrr?)r@divs r"cdivz tensor.cdivrr$cdSrr?rs r"sigmoidztensor.sigmoidrr$cdSrr?)r@rr ieee_roundings r"softmaxztensor.softmaxrr$cdSrr?rs r"ravelz tensor.ravelrr$TcdSrr?r@rreturn_indicesreturn_indices_tie_break_leftrs r"maxz tensor.maxrr$cdSrr?r@rtie_break_leftrs r"argmaxz tensor.argmaxrr$cdSrr?r s r"minz tensor.minrr$cdSrr?rs r"argminz tensor.argminrr$cdSrr?)r@rrrs r"sumz tensor.sumrr$cdSrr?r@rrs r"xor_sumztensor.xor_sumrr$cdSrr?rs r" reduce_orztensor.reduce_orrr$rcdSrr?r@rrs r"cumsumz tensor.cumsumrr$cdSrr?rs r"cumprodztensor.cumprodrr$rr descendingcdSrr?)r@rr!s r"sortz tensor.sortrr$cdSrr?)r@rs r"flipz tensor.fliprr$)rfrrorrNFNrrrrrrXr)rMr)rr)Nr?r<r<NNNNNF)NFF)NFTF)TF)rF)rrr!rr)rM)arGrarbrFrrnrr8rrrrrrrrrrrrXrrrrdrrjrrorrtrrzrrrrrrrrrurrwrrrrr rrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrr rrrrrrrr  CONSTEXPR_0r#r%rrs@r"rMrMSs & ? ? ? ? ? ?$$$$SSSS MMM WM MMM WM MMM WM MMM WM MMM WM MMM WM ... W. ... W. /// W/ /// W/ *** W* *** W*  %%% W% &&& W&  +++ W+ +++ W+ *** W* *** W* +++ W+ +++ W+ *** W*  *** W*  /// W/ /// W/ 333 W3 333 W3  444 W4 444 W4  000 W0 000 W0  111 W1 111 W1  ,,, W, ,,, W, 000 W0 000 W0 222 W2 111 W1  $$$ W$  W IIXI UUUU WU                                                                                                                                                                                              %)+              r$rMc6td|D|S)Nc~g|]:}t|tttfrt |n|j;Sr?)r%rzrrrrfrs r"r z*_type_for_tuple_values..s@kkk_`Jq3u:M,N,NZ~a(((TUTZkkkr$)r )rIr s r"_type_for_tuple_valuesr0s# kkdjkkkms t ttr$cdeZdZdddZdd Zd Zd Zd Zd ZdZ dZ dZ dZ dZ ddZdZdS)rNr0r rfOptional[tuple_type]cd|D|_t|tr ||_dS|t||_dSt |j|_dS)Ncg|]}|Sr?r?)rrs r"r z"tuple.__init__..s'''Qq'''r$)rIr%r rfr0)r@r0rfs r"rztuple.__init__sb''$''' dJ ' ' <DIII  "4((DIII.t{;;DIIIr$idxrc,t|trt|}t|tr |j|St|tt jfsJt |j|j|j|j Sr) r%rzrrIrrrrrr)r@r5s r"r ztuple.__getitem__s~ c3   !C..C c9 % % C;s# #cE8>#:;; ; ;;SYsx%@ABB Br$c|jj}|||vrtd||j||S)Nz 'tuple' object has no attribute )rfr AttributeErrorrIrr@rBr s r" __getattr__ztuple.__getattr__sK! >T// !JD!J!JKK K{6<<--..r$ct|}t|tsJ||j|<t |j|jj|_dSr)rr%rzrIr0rfr )r@r5rs r"_setitemztuple._setitemsM"3''#s#####  C*4; 8HII r$cXt|}t|j|jzSr)r rrIrss r"rz tuple.__add__s& ''T[5</000r$cht|tsJt|j|jzSr)r%rrrIrrss r"rz tuple.__mul__s/%+++++T[5;.///r$cZt|}t|j|jkSr)r rrIrss r"ruz tuple.__eq__s& '' 4555r$cNttj|jSr)rrrrIrs r"rztuple.__hash__"sHN4;//000r$c>td|jDS)Nc,g|]}t|Sr?rrs r"r z!tuple.__str__..&s000qCFF000r$)r&rIrs r"rz tuple.__str__%s!00DK000111r$c*t|jSr)rTrIrs r"rztuple.__iter__(sDK   r$c*t|jSr)rDrIrs r"__len__z tuple.__len__+s4;r$rgrhcD|jD]}||dSr)rIrn)r@rgrs r"rnztuple._flatten_ir.s2 # #A MM' " " " " # #r$cRddd|jDdS)Nr=rc34K|]}t|VdSr)rrs r"rz!tuple.__repr__..3s(::T!WW::::::r$))rrIrs r"rztuple.__repr__2s0=499::dk:::::====r$r)r0r rfr2)r5r)rgrh)rGrarbrr r:r<rrrurrrrErnrr?r$r"rrs<<<<<CCCC///JJJ 111 000666111222!!!   ####>>>>>r$rceZdZdZdS)rcV||_||_||_t|_dSr)rrrr$rf)r@rrrs r"rzslice.__init__8s&   LL r$Nr&r?r$r"rr6s#!!!!!r$rcFeZdZddZdd ZddZddZddZddZddZ dS)tensor_descriptor_base_typerc||_dSrr)r@rs r"rz$tensor_descriptor_base_type.__init__As $r$rgrhryrzr)"Tuple[tensor_descriptor_base, int]cFt|||j}||dzfSr)tensor_descriptor_baser)r@rgryrs r"r~z)tensor_descriptor_base_type._unflatten_irDs&&wvHHfqj  r$rrrrric|jj}|||j||dSr)rrrrcreate_tensor_descriptor_typer)r@rr is_signeds r"rz-tensor_descriptor_base_type._flatten_ir_typesHsQO.<<>>  7889N9Nw9W9WYbccdddddr$r&cd|jdS)Nztensor_descriptor.s$<<<*4"))++<<.aggregate_valueTcFt|Sr)r__new__)this_clsrs r"rz1_aggregate..aggregate_value._get_instances77??8,, ,r$Nr,r:c|}i}tjtrnJdt jjjvr||d<dt jjjvr||d<j|g|Ri||jD]-}t||stdj d|d.|S)Nr,r:zconstructor for z did not initialize attribute '') rr%rrr6r7rCrprErr8rG) rr,r:r0r1r extra_kwargsrBclss r"rz+_aggregate..aggregate_value.__new__s--//HL#, 44 <'"3CL"A"A"LLL09L-7#4S\#B#B#MMM1;L. CL CD C C CL CF C C C+0022 r rx..r()pCL)p)pim)p)p)pqqqrOr$c (|jvrtjd|dt|j|s0t dj|d|dt |t ||dS)Nz has no attribute 'rz Expected z for attribute 'z', got )rpr8rGr%rUrfr __setattr__)r@rBrrrs r"rz/_aggregate..aggregate_value.__setattr__+s3...$ %P%P%P%P%PQQQeS%8%>?? s qC,?,E q qW[ q qdhindodo q qrrr GG  e , , , , ,r$rgrhr)ricjD]%}t|||&dSr)rprErZrn)r@rgrBrs r"rnz/_aggregate..aggregate_value._flatten_ir2sL+0022 9 9d##//8888 9 9r$cltfdjDS)Nc>g|]}|t|jfSr?)rZrf)rrBr@s r"r z<_aggregate..aggregate_value.type..9s+#l#l#lT74+>+>+C$D#l#l#lr$)rrprE)r@aggregate_valuers`r"rfz(_aggregate..aggregate_value.type6sB"?#l#l#l#lQTQdQiQiQkQk#l#l#lnn nr$ro) rGrarbr__triton_aggregate__ classmethodrrrrnrrfr)rrrs@r"rr s!#  - - - -  -04       , - - - - - - 9 9 9 9 9 9  n n n n n  n n n n nr$rr)rdrr6 getmembers isfunctionismethodr%rr4r hash_attrsrGrarbrF)rrrBmemberrs` @r" _aggregater s,n,n,n,n,n,n,n,n*,n,n,n\,J!,S11**v  f % % *)9&)A)A *ZPVXcEdEd *z!!v666!!&)))!+O"|O!$O#O !kO r$cJt|}||S)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_idrr,s r"rrQs%  % %D    % %%r$cJt|}||S)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"rrds%  % %D  ! !$ ' ''r$cjt|}t|}|||Sr)rarange)rendr,s r"rrus2  ' 'E s # #C   E3 ' ''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 c8t|}d|DS)Nc,g|]}t|Sr?rr1s r"r z!_unwrap_shape..s! 3 3 3  # # 3 3 3r$rrs r"rrs#  ' 'E 3 3U 3 3 33r$cBt|}t||Sr)rrrs r"_shape_check_implrs" % E Lr$ct|}t|}t|}||||S)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 )rrfull)rrrr,s r"rrsA e $ $E  ' 'E  ' 'E >>% . ..r$c.|||S)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)inputrtr,s r" broadcastrs  ) )% 7 77r$rcftt|}|||S)ae 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) )rrWbroadcast_impl_shaperr,rs r"rrs0" .u55 6 6E  ) )% 7 77r$rct|}|sXt|j}|dkrtdt t j|dz |dz |dz gz}|||S)a 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. rz2tl.trans invoked with a 0- or 1-dimensional tensorr;)rWrDrr.rHrranger)rr,rns r"rrs. D ! !D <    q55QRR RHN1q5))**a!eQU^;   UD ) ))r$cLt|}|||S)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. )rWr)rr,rs r"rrs'( D ! !D   UD ) ))r$Fc|r||Stj}|tjks"Jd|dtjtt |t fdt j|DsJddjdjt|}tt j|}| |t|||}tj}jjz|<t|||S)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\K|]&}|kj|j|kV'dSrr)rrrrrts r"rzcat..sZAQRRRu{1~Q/RRRRr$z6tensor dims must match except in the concat dimension , got r) catrDr _wrap_axisrrrrrrHinsertrr) rrt can_reorderrr,rankrorder new_shapes `` ` r"rrs 8}}UE;777 u{  D 3u{## # # #%jT%j%jX[\a\gXhXh%j%j # # # )#.. 5 5C 9M9M  yyxcxxY^Ydxxkpkvxxyy  UEY///A %% & &E LLd5I...AU[!!I[% C(88IcN 1i9 5 5 55r$c.|||S)aZ 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)abr,s r"rr(s* >>!Q  r$ct|jdkr|Sd}|jD]}||z}|dks Jd||S)z6 Convert a single-element tensor to a scalar. rr;z'can only unsplat single-element tensors)rDrunsplat)r!r,r:rds r"_unsplatr@se 17||q E W   A:::@:::   Q  r$rct|jdk}|r||d}||\}}|r$t |||}t |||}||fS)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 r;rr)rDrrrr)rr,r: was_rank_1out_lhsout_rhss r"rrMs*QW"J(  ! !!Q ' ' q))GWP7iJOOO7iJOOO G r$ctdtt|}|||dS)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)rrrWrrs r"rrpsD" NOOO .u55 6 6E   UEt  < <.s, H H HaJ+A.. 9 9 H H Hr$z7expand_dims received duplicate axes, normalized axes = ) r^rr%r rrHrDrsetr.sortedr)rrr,axesrrrs @r"rrs    & &E  % %D#D8U*;<< H4:::4&D5;#d))+H H H H H4 H H HD 3t99~~T""YSWYYZZZ C D\\,,##C++ Jr$rrrc||}t|}t|}t|}|r|||S||||S)a 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 )r^rrr)rrrrr,s r"rrsu$    & &E  ' 'E/0DEE"7++G/  ... >>%(< = ==r$c| | Jd|.d|jjjv}tjjp |r|s|dnd}t |}t |}t |}t |}t|j} t|j} t| t| cxkrdks nJd| dd| ddks Jd| d | dks Jd | dd | d gz} |"t|j| ks Jd t| } | d krd } tj | dz D] }| | |z} | || g| ddzd}| || g| ddzd}|$| || g| ddzd}| ||||||}| d kr| || d}t|j| ks Jd|S)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 specifiedtf32ieeerz*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 incompatibler;Frzoutput shape is unexpected)rrallowed_dot_input_precisionsrlanguage fp32_defaultrrHrrDrrrdot)rrtaccinput_precision allow_tf32max_num_imprecise_acc out_dtyper, supports_tf32a_shapeb_shapec_shaper batch_sizerress r"rrs<  "j&8&8:w&8&8 8)"3";"XX .5v];uFP;uT^Tf;A&nt +?;;O$Y//I01FGG s # #C5;G5;G w<<3w<< , , , ,1 , , , , ,.Z , , , 3B3<73B3< ' ' ')W ' ' ' 2;'"+ % % %'] % % %crclgbk]*G CI')))+N))) 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 )rr^rj) pointerrrtrpadding_optionrrvolatiler,s r"rjrjbs^  % %D  ' 'E ""4(( ##E**).99N).99N*?;;O#H--H >>'4P^`o" $ $$r$descrrgc0|||S)z.Load a block of data from a tensor descriptor.r)rj)rrr,s r"load_tensor_descriptorrs 99W 9 2 22r$rc2||||S)z-Store a block of data to a tensor descriptor.r)r)rrrr,s r"store_tensor_descriptorrs ::gu : : ::r$c||}t|}|||}t|}t|}|||||||S)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"} )r^rr)rrrrrrr,s r"rrsuT    & &E  % %D ""4(().99N*?;;O ??7E4Q` a aar$basecTtd|||||||S)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.)rmake_block_ptr)rrrrrdrr,s r"rrs2 fggg  # #D%';PU V VVr$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.|||S)z Advance a block pointer :param base: the block pointer to advance :param offsets: the offsets to advance, a tuple by dimension )r)rrr,s r"rrs   T7 + ++r$zerorrrrdList[constexpr]cRt|}||||||S)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) )rmake_tensor_descriptor)rrrrdrr,s r"r#r#s.r*.99N  + +D%+~ ^ ^^r$rBr&has_cmpCallable[[T], T]cdfd }|S)Nfuncrr)c:dd}r|dz }|dz }||_|S)Nz 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 )rF)r'docstrr$rBs r" _decoratorz&_add_atomic_docstr.._decoratorC sT9999  : : :F     r$r'rr)rr?)rBr$r*s`` r"_add_atomic_docstrr,A s/4 r$zcompare-and-swap)r$c||}||}t|}t|}||||||Sr)r^rr)rrrrrr,s r"rr` s`   c " "C   c " "C s # #C  ' 'E   c3 > >>r$exchangec||}t|}t|}t|}||||||Sr)r^rrrrrrrr,s r"rrk sZ   c " "C s # #C  ' 'E  % %D  #tS% @ @@r$r>c||}t|}t|}t|}||||||Sr)r^rrr0s r"rrv Z   c " "C s # #C  ' 'E  % %D   dC ? ??r$r c||}t|}t|}t|}||||||Sr)r^rrr0s r"rr r2r$rc||}t|}t|}t|}||||||Sr)r^rrr0s r"rr r2r$z logical andc||}t|}t|}t|}||||||Sr)r^rrr0s r"rr r2r$z logical orc||}t|}t|}t|}||||||Sr)r^rrr0s r"rr sZ   c " "C s # #C  ' 'E  % %D   wT3 > >>r$z logical xorc||}t|}t|}t|}||||||Sr)r^rrr0s r"rr r2r$c||}t|}t|}||||S)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. )r^rwhere) conditionr!yr,s r"r9r9 sG"##I..IQAQA ??9a + ++r$r<clt|}t|}||||Sr)rr>r!r;r<r,s r"r>r> 3QAQA ==A0 1 11r$clt|}t|}||||Sr)rrEr=s r"rErE r>r$clt|}t|}||||Sr)rrIr=s r"rIrI r>r$ propagate_nanc||}||}t||}t||}t|}||||S)aH 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)r^_promote_bfloat16_to_float32rminimumr!r;rAr,s r"rDrD s 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)r^rCrclamp)r!rr rAr,s r"rJrJ s$ AA   c " "C   c " "C$Q)<<._decorator@ s  ) ( ((MQ(( ((( (F  $ # ##^b## ### #F  # ##{D## ### #F}}$}//  r$r+r?)rBrKrLrMr*s```` r"_add_reduction_docstrrU= s;4 r$c#dK|}dV||dSr)get_insertion_pointrestore_insertion_point)rips r"_insertion_guardrZ] s:  $ $ & &B EEE ##B'''''r$cttrtf|dSfd}fdtt|}(t t dj|}|r;tfd|D}ntfd|D}|S)a'Applies 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 rr,r:rc d Ddz}|d} jt5fd}|t t ||fdt |D}  |i}t|tr |j g}n d|D}j |ddddS#1swxYwYdS)Nc&g|]}|jjSr?rfrrrs r"r z7reduce..make_combine_region..w 444qv}444r$rrc.|Srrrrs r"r#z5reduce..make_combine_region..{ aggg..r$c\g|](\}}t||)Sr?rMrrrrblocks r"r z7reduce..make_combine_region..} 1QQQBF599Q<<,,QQQr$r1cg|] }|j Sr?r2rrs r"r z7reduce..make_combine_region.. 55518555r$) get_regionrrZcreate_block_with_parentrHrrcall_JitFunctionr%rMr2create_reduce_ret) reduce_op param_typesregionrr0resultsrgrirr:r,rrs @@r"make_combine_regionz#reduce..make_combine_regionv s[44e444q8 %%a((# g & & 0 0....E44VT#e[BYBY=Z=Z[[EQQQQ)K:P:PQQQD 11*d21NNG'6** 6">*55W555 %G %w / / 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0BC))C-0C-cZtj|D]}t|d}|S)Nrr)rrr)rndimsrr,s r" expand_ndimszreduce..expand_ndims s7&& 7 7AAqI666AAr$Nc3<K|]}t|VdS)rN)r)rrr,rs r"rzreduce.. s2OOa AtyAAAOOOOOOr$c3`K|](}|tdjV)dS)rN)rDr)rrr}rs r"rzreduce.. s=JJ QE!HN(;(;<<JJJJJJr$) r%rMrrrrDr reductionr) rrrrr,r:ryrr}s ``` `` @r"rrd s:%  wuizYR[hrssstuvv 0 0 0 0 0 0 0 0  % %D$Y//I $E!HN 3 344   eT+> ? ?CK  OOOOO3OOOOOCCJJJJJcJJJJJC Jr$ch|jj}|tur|t|S|S)Nr)rfrbfloat16rr)rr,rs r"rCrC s3 IHttGyt111 Hr$ct|}|j|}td||}t|jdkr\dt jt|jD}||=t |||}t||j|}t||f|||||\} } | | fS)Nrrr;c,g|]}t|Sr?r0)rrs r"r z(_reduce_with_indices.. sQQQ1)A,,QQQr$r\) rrrrDrrrrr) rrrrr,r:rraxes_to_expandrvaluerindicess r"_reduce_with_indicesr s  % %D DA 1a9 - - -E 5;!QQs5;?O?O0P0PQQQ 4 E>YGGGUEK9EEEuendJ)_h)3555FH 8 r$cdfd }|S)Nr'rr)cdd}|ddddz }||_|S)NaZ 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: boolrPrQz` 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 rRrArS)r'r)rMrBs r"r*z$_add_scan_docstr.._decorator sm  # ##{D## ### #F}}$}//  r$r+r?)rBrMr*s`` r"_add_scan_docstrr s/& r$cttrtf||dSfd}t|}|(t |t dj}|||S)a6Applies 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 rrc d Ddz}|d} jt5fd}|t t ||fdt |D}  |i}t|tr |j g}n d|D}j |ddddS#1swxYwYdS)Nc&g|]}|jjSr?r_r`s r"r zAassociative_scan..make_combine_region.. rar$rrc.|Srrcrds r"r#z?associative_scan..make_combine_region.. rer$c\g|](\}}t||)Sr?rgrhs r"r zAassociative_scan..make_combine_region.. rjr$rkcg|] }|j Sr?rmrns r"r zAassociative_scan..make_combine_region.. rpr$) rqrrZrrrHrrrsr%rMr2create_scan_ret) scan_oprvrwrr0rxrgrirr:r,rrs @@r"ryz-associative_scan..make_combine_region s[44e444q8 ##A&&# g & & . .....E44VT#e[BYBY=Z=Z[[EQQQQ)K:P:PQQQD 11*d21NNG'6** 6">*55W555 #G #W - - . . . . . . . . . . . . . . . . . .rz)r%rMrrrrDr)rrrrr,r:rys` ` `` r"rr s%  u 4WPYfpqqqrstt . . . . . . . .  % %D $E!HN 3 344  % %eT3F P PPr$ct|}t|}|||}||||S)aqcomputes 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 )rr^r)rrrr,r:s r"rr sO$H--H  % %D ""4((   uh 5 55r$ct|}t|}t|}||||S)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 )rr)srcrrr,s r"rr sC s # #C  ' 'E  % %D   C - --r$r;)packr,r: scalar_fn!Callable[..., Tuple[tensor, ...]]r0c |dksJd|D}|j}|g}|}t|D]\} } t j|D]g} | |||t | |z| z| ht|5| | ||i} t| t} | r| f} d| D}||||dddn #1swxYwYd| D}|}|dkr<|dd|}t jd|D]}|||d|ks Jdfd}|||||||}| r|d n|S) a 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. r;c&g|]}|jjSr?r_r`s r"r z#map_elementwise..D s111qQV]111r$rkcg|] }|j Sr?rmrns r"r z#map_elementwise..V s44418444r$Ncg|] }|j Sr?)rfrs r"r z#map_elementwise..Z s666!qv666r$z!type mismatch in unpacked resultsc\|d}|dS)Nr)rq push_back)elementwise_oprwris r"make_elementwise_regionz0map_elementwise..make_elementwise_regiona s0**1--r$r)r new_blockget_locrrradd_argument_atrrrMrrZset_insertion_point_to_startrsr%set_loccreate_map_elementwise_retmap_elementwise)rrr,r:r0 in_scalar_tysr scalar_args original_locrrjscalar_results is_singlergfn_result_typesscalar_result_typesoffsetrresultris @r"rr sJ 1999911D111MG     EK??$$L=))DD2%% D DA  ! !"((7"3"3\ B B B   veiiD1 &=&=rBB C C C C D ' " " 4 4,,U333#44Y TV4WW~v66  -+_N44^444 %%%**7333 4 4 4 4 4 4 4 4 4 4 4 4 4 4 476~666O) axx-fff5nQ-- m mF&/&,$,*GGGGIlGGGG      OOL!!!  & &t-@$H_ ` `F! -6!99v-s'A>E11E58E5c*|S)zA Insert a barrier to synchronize all threads in a block. ) debug_barrierrs r"rro s  " " $ $$r$c tt|tr|g}t|D]o\}}t|tstd|dt|jt s(td|dt |jdpd|D}|||S)zd Let the compiler know that the values in :code:`input` are all multiples of :code:`value`. values element  must have type `constexpr`1 must have type `constexpr[int]`, got `constexpr[rcg|] }|j Sr?rrs r"r zmultiple_of..  & & &!ag & & &r$)r%rrrUrrzrf multiple_ofrrIr,rrs r"rrw s &)$$&!!tt1!Y'' NLaLLLMM M!'3'' trarrbfghgnboborrrss s t & &v & & &F   / //r$c tt|tr|g}t|D]o\}}t|tstd|dt|jt s(td|dt |jdpd|D}|||S)z^ Let the compiler know that the `value` first values in :code:`input` are contiguous. rrrrcg|] }|j Sr?rrs r"r z"max_contiguous.. rr$)r%rrrUrrzrfmax_contiguousrs r"rr s &)$$&!!tt1!Y'' NLaLLLMM M!'3'' trarrbfghgnboborrrss s t & &v & & &F  # #E6 2 22r$c tt|tr|g}t|D]o\}}t|tstd|dt|jt s(td|dt |jdpd|D}|||S)z Let the compiler know that the `value` first values in :code:`input` are constant. e.g. if :code:`values` is [4], then each group of 4 values in :code:`input` should all be equal, for example [0, 0, 0, 0, 1, 1, 1, 1]. rrrrcg|] }|j Sr?rrs r"r z!max_constancy.. rr$)r%rrrUrrzrf max_constancyrs r"rr s&)$$&!!tt1!Y'' NLaLLLMM M!'3'' trarrbfghgnboborrrss s t & &v & & &F  " "5& 1 11r$cR|||S)z< Allow compiler to assume the :code:`cond` is True. )assumer^)condr,s r"rr s&   I//55 6 66r$  )seprfileflushr,rrcdS)a 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?)rrrrr,rIs r" static_printr s  Dr$cdS)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?)rmsgr,s r" static_assertr s  Dr$)hexr,c>ddl}t|}t|ts J|dd}|D]}||jvrd}n|s J|dg}|D]*}|||+||||S)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 rNz is not stringTFz is not an ascii string)stringrr%r& printablerr^ device_print) rrr,r0rb_asciichnew_argsrs r"rr sDMMM !& ) )F fc " "==v$=$=$=== "G V% % %G E & 66v666667H22 ++C001111  ! !&(C 8 88r$ct|}t|}|||}|||||S)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. )rr^ device_assert)rrrr,s r"rr s]& s # #C  % %D ""4((  " "9#6#6t#<# Execute inline assembly over a tensor. Essentially, this is :code:`map` where the function is inline assembly. The input tensors :code:`args` are implicitly broadcasted to the same shape. :code:`dtype` can be a tuple of types, in which case the output is a tuple of tensors. Each invocation of the inline asm processes :code:`pack` elements at a time. Exactly which set of inputs a block receives is unspecified. Input elements of size less than 4 bytes are packed into 4-byte registers. This op does not support empty :code:`dtype` -- the inline asm must return at least one tensor, even if you don't need it. You can work around this by returning a dummy tensor of arbitrary type; it shouldn't cost you anything if you don't use it. Example using `PTX `_ 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 TFc:g|]}|Sr?r])rrr,s r"r z*inline_asm_elementwise.. s'BBBc,,S11BBBr$)arithmetic_check allow_lhs_ptr allow_rhs_ptrrcDg|]}j|Sr?)rfr)rdt broadcast_args r"r z*inline_asm_elementwise.. s*NNN"})99"==NNNr$cg|] }|j Sr?rmr`s r"r z*inline_asm_elementwise.. s///Aqx///r$c:g|]}|Sr?rc)rrrs r"r z*inline_asm_elementwise.. s&@e@e@eWY'ARAR@e@e@er$c3dK|]*\}}t||V+dSr)rM get_result)rrrcalls r"rz)inline_asm_elementwise.. s=PPEAr**B//PPPPPPr$)rrTrUtypingrr _DtypeClassrbinary_op_type_checking_implrrrcreate_inline_asmrMrr)rrr0rrrr,has_multiple_outputsres_tys dispatch_argsbin_op_type_checkingrrrrgrrrs ` @@@r"inline_asm_elementwiser sr s # #C&{33K  % %D"7++G U # $  K-u 5 5EGBBBBTBBBB}O&  2"     &a( ! I ID33D-HH A}}   O$]33 P P4&:&:4&O&O# a !!NNNNNNNG/////GG  $ $S+w@e@e@e@e]d@e@e@egnpt u uD 6dooa(('!*555 PPPPYw=O=OPPP P PPsAA&%A&c&eZdZdZddZdZdZdS) static_rangea 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. Ncht|ts J|d|td|_n)t|ts J|d||_|td|_||_dSt|ts J|d||_||_dS)Nz7 used as tl.static_range start value is not a constexprr;z6 used as tl.static_range step value is not a constexprrz5 used as tl.static_range end value is not a constexpr)r%rrrr)r@arg1arg2rs r"rzstatic_range.__init__ s$ **llt,l,l,lll* <!! DIIdI.. o o40o0o0o o o.DI <"1DJDHHHdI.. n n40n0n0n n n.DJDHHHr$c tdNz8static_range can only be used in @triton.jit'd functionsrrs r"rzstatic_range.__iter__ UVVVr$c tdrrrs r"__next__zstatic_range.__next__ rr$r*rGrarbrFrrrr?r$r"rr sX"    WWWWWWWWr$rc*eZdZdZ ddZdZdZdS)ra 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. NFc |td|_n||_|td|_||_n||_||_||_||_||_||_||_| |_ dS)Nr;r) rrrr num_stagesloop_unroll_factordisallow_acc_multi_bufferflattenwarp_specialize disable_licm) r@rrrrrrrrrs r"rzrange.__init__ sz <!! DIIDI <"1DJDHHDJDH$"4)B& .(r$c tdNz4tl.range can only be used in @triton.jit'd functionsrrs r"rzrange.__iter__ QRRRr$c tdrrrs r"rzrange.__next__ r r$)NNNNFFFFrr?r$r"rr sc))VX\ej))))&SSSSSSSSr$rceZdZdZddZdS)r:a= 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. Fc"||_||_dSr)r:r)r@rrs r"rzcondition.__init__' s(r$Nr,)rGrarbrFrr?r$r"r:r: s2"))))))r$r:lib_namelib_pathrHarg_type_symbol_dictdictret_typec ,t|dkrtdtt|d}t||kr"tdt|d|g} g} |D]} t | t r5| | j| | jL| t| | | t| } | |vr'td|d| || d} |j } t |||| | | | ||S)a 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 rzarg_type_symbol_dict is emptyz+length of input args does not match.Expect rz,input arg type does not match.Expect one of ) rDr.rHrEr%rMrrr2rfrrr)r'r rr0rrrr,num_args arg_typesarg_listrsymbolrs r"dispatchr1 s   A%%89994,113344Q788H 4yyH?#&t99??4<??@@ @IH!! c6 " " !   SY ' ' ' OOCJ ' ' ' '   T#YY ' ' ' OOC i  I,,,Y*>*C*C*E*EYYMVYYZZ Z&i03#dd8XvxPWAXAXZabbdlmmmr$c |}d}g}tjt|D]a} ||| || <||| j|| jrd}bt|}||d} t|dkrd} ||vrd} |d} |D]} | | | | \}} tjt|D]&} | || | | \|| <}'|s| j | } |j j }t|||||| ||S)av 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 TFr;r)r)copyrrrDr^rrrfrrrrrcreate_extern_elementwiser)r rr0rrr,r all_scalarrrrrrrrr's r"extern_elementwiserY sIIKKMJI ^C .. / /$..}Q/?@@ aq)/000   ) ) + + Ji  I#I.q1H 9~~ , , ,$ %a( ! i iD(EEdMWg F i i A}}M 2 233 l lA"+"H"HWXIY[hZj#I#l#l M! aa D$)99(CCH   6D D(Hm=QS[]dfo p ppr$c.|||S)z Convert both operands to a single common type :param lhs: the left operand :param rhs: the right operand :param builder: the builder )r)r r semantics r"binary_op_type_legalizationr s  0 0c : ::r$c t|S)z#A decorator for external functions.)r8r[s r"externr! s 2;;r$c@|dkotjd|dkS)Ngg?r)r3copysignrVs r"is_negative_zeror$ s" 8 1 c1--11r$rAr,clt|}td|D}|rw|tus Jdtd|DrJtd|DrJt t jt|S|tur tj}ntdt|dks Jd|d}|d dD]}t|||| }|j rtd |S) Nc3BK|]}t|t VdSrr%rdrs r"rzbuiltin_max.. /CC:a444CCCCCCr$z-propagate_nan is not supported on builtin maxc3>K|]}tj|VdSrr3isnanrs r"rzbuiltin_max.. *33tz!}}333333r$c34K|]}t|VdSrr$rs r"rzbuiltin_max.. +99q'**999999r$zJpassing propagate_nan to builtin max is deprecated, use tl.minimum insteadrmin requires at least 2 valuesrr;r%zMbuiltin max on non-scalar tensor values is deprecated, use tl.maximum instead)rr_NOTHINGanyrrr  PropagateNanNONErrDrHrfr)rAr,r0 is_constexprmax_valrs r" builtin_maxr8 X  % %DCCdCCCCCLC(((*Y(((33d3333333399D99999999&:4&@&@AABBB  $)  YZZZ t99>>>;>>>1gGABBxZZ'3myYYY|^ \]]] Nr$clt|}td|D}|rw|tus Jdtd|DrJtd|DrJt t jt|S|tur tj}ntdt|dks Jd|d}|d dD]}t|||| }|j rtd |S) Nc3BK|]}t|t VdSrr(rs r"rzbuiltin_min.. r)r$z-propagate_nan is not supported on builtin minc3>K|]}tj|VdSrr+rs r"rzbuiltin_min.. r-r$c34K|]}t|VdSrr/rs r"rzbuiltin_min.. r0r$zJpassing propagate_nan to builtin min is deprecated, use tl.minimum insteadrr1rr;r%zMbuiltin min on non-scalar tensor values is deprecated, use tl.minimum instead)rrr2r3rrrr4r5rrDrDrfr)rAr,r0r6min_valrs r" builtin_minr? r9r$)T)r rr)rrr)rrzr'rXr)r)rrM)FrNr*r)r&r')NNr?r<r<r<FN)rrRrrgr)rM)rrRrrgrrMr)rM)Nr?r<r<N)rrM)r N) rrMrrrrrdr!r)rr,)rBr&r$rXr)r%r+)NNNN)TN)r<r)rAr) rBr&rKr&rLr&rMr&r)r%)FNN)rBr&rMr&r)r%)rrr0rM)rr&rr&)r<N)r<NN) rr&rr&r0r rrrrXrrz) r r&rr&r0rHrrrrrrX) r r&rr&r0rHrrrrX) __future__rr3warningsr contextlibrenumr functoolsrrr rr r r r rrr dataclassesrrr<r runtime.jitrr6 _C.libtritonr_utilsrrrrrr5 PROPAGATE_NANr4rr8rRrWrLr^r`rdrerrr-rr rrrrrr r$r0r"rrr r!r#r$r%r&float8e5 float8e5b16 float8e4nv float8e4b8 float8e4b15float16rrfloat64pi32_tr+rMr0rrrMrRrrrrrrrrFrrrrrrrrrrrrrrrrrrrrjrrrrrr#r,rrrrrrrrr9r>rErIr5rDrHrJrUrZrrCrrrrrrrrrrrrrrrrrrr:rrrr!objectr2r$r8r?r?r$r"rSs)"""""" %%%%%%5555555555 LLLLLLLLLLLLLLLLLL!!!!!!%%%%%%iiiiiiiiiiii GCLL%     "****Z   ,....  """ "         " " " " " " " """"""""",-----Y---8L-L-L-L-L- L-L-L-^ill 666LLLLLILLLd %.%.%.%.%.5%.%.%.P3!3!3!3!3!3!3!3!l$F$F$F$F$F$F$F$FN!!!!!!!! uV}} uV}} uV}} g g g g x x x 5>>eJ U9   U9   eJ %-- 5== %-- %-- e  UUUU8Z Z Z Z Z ZZ Z Z z uuuu?>?>?>?>?>J?>?>?>D!!!!!!!!/////)///:XNXNXNXNXNZXNXNXNvkkkkk8kkk6*****.***8 $.....i...6???N &&& &$ ((( (  ((( ( '>   444   /// /,  8 8 8  8*.8888 8&*.***** *<$(**** *, 666 6D      .      B"&==== =(GGG G',8888 8*... 8>>>> >< ]amtDDD DN KO[`mq'T&7&7&7 &7\ rt#'9$9$9$ 9$x %)3333 3  &*;;;; ; .b.b.b .bb  W W W W  WQ,,,  ,  9_9_9_9_ 9_B>&555???65 ?JAAA  AE@@@ @E@@@ @E@@@ @M""@@@#" @L!!???"! ?M""@@@#" @ ,,, ,6 2222 2  2222 2  2222 2  -9->$2222 2* -9->$2222 2* 2>2Ct7777 7D[_+/@((( ... .b        ,0"Q"Q"Q "QJ666 6$... ."   H.H.H.H.H. H.` %%% %  0 0 0  0  3 3 3  3 222 2$ 777 7 %(TE]a                  $)T-9-9-9-9 -9` III I2 ?C~Q~Q~Q~Q ~QL%W%W%W%W%W:%W%W%WPCSCSCSCSCSJCSCSCSL))))) )))8%n%n%n%nP !%'q'q'q'q 'qT;;; 688222 %- . %- r$