çg\ ddlmZddlmZddlmZddlmZddlm Z m Z ddl Z ddl m Z m Z mZmZmZmZddlZdd lmZddlZddlZdd lmZd d lmZed ZdZdZejZddZ ddZ!dZ"ddZ#e ddZ$dZ%GddZ&e&Z'Gdde&Z(Gdde(Z)Gdd e&Z*Gd!d"e&Z+e&d#Z,e&d$Z-e&d%Z.e&d&Z/e&d'Z0e&d(Z1e&d)Z2e&d*Z3e&d+Z4e&d,Z5e&d-Z6e&d.Z7e&d/Z8e&d0Z9e&d1Z:e&d2Z;e&d3Z<e&d4Z=e&d5Z>e(e0Z?dd9Z@Gd:d;ZAGd<d=ZBeBdZCd>ZDGd?d@ZEdAZFdBZGe ddCZHe ddDZIe ddEZJdFZKe ddGZLe ddHZMe!e ddIdJZNe!e ddIddLZOe!e ddIdMZPe ddOZQe ddPZRedQZSe!e dddSZTe!e ddIdTZUe!e dNddUdVZVdWZWe!e ddXZXe!e ddd\ZYe dddde=dfd]ZZe dd`Z[e ddaZ\e ddbZ]e!e ddcZ^e dddeZ_e!e ddfZ`dddkZae!e eadldmnddoZbe!e eadpddqZce!e eadrddsZde!e eadtdduZee!e eadvddwZfe!e eadxddyZge!e eadzdd{Zhe!e ead|dd}Zie dd~Zje ejkdfddZle ejkdfddZme ejkdfddZndddZoedZpe!e ddZqe ddZre ddZsddZte!e ddZue!e ddZve ddZwe ddZxe ddZye ddZze ddddNddddZ{e ddZ|e dNdddZ}e ddZ~e dddZGddZGddZ dddZe dddZdZdZdS)) annotations)warn)contextmanager)Enum)partialwrapsN)UnionCallableListSequenceTypeVarOptional)jit)ir)semanticTi__triton_builtin__fnreturnctsJtfd}t|td|S)zMark a function as a builtin.cJd|vs|dtd|i|S)N_builderzdDid you forget to add @triton.jit ? (`_builder` argument must be provided outside of JIT functions.) ValueErrorargskwargsrs P/var/www/html/ai-engine/env/lib/python3.11/site-packages/triton/language/core.pywrapperzbuiltin..wrappersH V # #vj'9'A`aa ar4"6"""T)callablersetattrTRITON_BUILTIN)rr!s` r builtinr&sW B<< 2YY####Y#  G^T*** Nr"c tsJtj}t|jddhz dk}jsd_xjdjd|rdnddjd |rd ndd z c_fd }t|j }|d  d|d <| |}||_ djd|_trt|tdttj|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 _generatorrzb 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|SNrs r r!z"_tensor_member_fn..wrapperDsr4"6"""r"rselfname) parameterszForwards to :py:func:`z` free functionT)r#inspect signaturelenr1keys__doc____name__listvaluesreplace __signature__ is_builtinr$r%tensor)rorig_sighas_argsr! new_paramsnew_sigs` r _tensor_member_fnrB*s B<< $$H8&++--\0JJKKaOH : JJ(08uub K'/677BJJ ##### h)002233JqM))v)66JqM*55G#GKr{KKKGO"~~/... FBK))) Ir"ct|dkr. t|d|dS#t$rYnwxYw|S)z7Returns x[0] if x has one element and x[0] is iterable.rr)r4iter TypeError)xs r _unwrap_iterablerGVsU 1vv{{  1JJJQ4K    D  Hs 2 ??boolc.t|tdS)z-Is this a registered triton builtin function?F)getattrr%rs r r<r<ls 2~u - --r"c"t||Sr, _to_tensor)rFrs r to_tensorrOqs a " ""r"ct|tr(t||tSt|t rd|cxkrdkr+nn(t||tSd|cxkrdkr+nn(t||tSd|cxkrdkr+nn(t| |tSd|cxkrdkr+nn(t| |tStd|dt|trd }d d d zz}t!d |}|tdks|dks||ks||cxkr|kr+nn(t||t$St||t(St|t*rt-|j|St|tr|SJd|dt1|d)NilllllzNonrepresentable integer .g8g?rabsinfgFzcannot convert z of type z to tensor) isinstancerHr=get_int1int1int get_int32int32 get_uint32uint32 get_int64int64 get_uint64uint64 RuntimeErrorfloat __builtins__get_fp32float32get_fp64float64 constexprrNvaluetype)rFbuilder min_float32 max_float32abs_xs r rNrNvsy!Tg&&q))4000 As   Q         '++A..66 6 a    %     ',,Q//88 8 q 5 '++A..66 6 a    %     ',,Q//88 8?1???@@ @ Au   !QV+ U#A&& E%LL C<< 66 % . . . .; . . . . .'**1--w77 7'**1--w77 7 Ay ! !!'7+++ Av  CCACCQCCCCCCr"ceZdZgdZgdZgdZgdZdgZGddeZ dZ d Z d Z d Z d Zd ZdZdZdZdZdZdZdZdZdZdZdZdZdZdZdZdZdZ dZ!d Z"d!Z#e$d"Z%e$d#Z&e$d$Z'e$d%Z(e$d&Z)d7d(Z*d7d)Z+d*Z,e-d+Z.d8d0Z/d1Z0d2Z1e-d9d4Z2d5Z3d6S):dtype)int8int16rZr^)rWuint8uint16r\r`) fp8e4b15fp8e4nvfp8e4b8fp8e5fp8e5b16fp16bf16fp32fp64)rzr{r|r}voidceZdZdZdZdS)dtype.SIGNEDNESSrrN)r7 __module__ __qualname__SIGNEDUNSIGNEDr-r"r SIGNEDNESSrsr"rct|dr|j}||_|tjtjztjztjzvs J||tjvrQtjj |_ t| dd|_ |j |_dS|tjvrQtjj|_ t| dd|_ |j |_dS|tjvr|dkrd|_d|_d|_dS|dkrd|_d|_d |_dS|d krd|_d|_d|_dS|d krd |_d|_d|_dS|d krd |_d|_d|_dS|dkrd|_d|_d|_dS|dkrd |_d|_d|_dS|dkrd|_d|_d|_dS|dkrd|_d|_d|_dSt%d||dkr d|_dSdS)NrirXrurvrwrxrryrz r{rRr| r}5@iz Unsupported floating-point type r~r)hasattrrir0rp SINT_TYPES UINT_TYPESFP_TYPES OTHER_TYPESrrint_signednessrXsplit int_bitwidthprimitive_bitwidthrfp_mantissa_width exponent_biasra)r.r0s r __init__zdtype.__init__s 4 ! ! :D u'%*::U^KeN_____ae___ 5# # #"'"2"9D  #DJJu$5$5b$9 : :D &*&7D # # # U% % %"'"2";D  #DJJu$5$5b$9 : :D &*&7D # # # U^ # #z!!)*&*+'%'""""")*&*+'%&""""")*&*+'%&""")*&*+'%'"""##)*&*+'%'""")+&*,'%'""")*&*,'%(""")+&*,'%(""")+&*,'%)""""#Ld#L#LMMM V^^&'D # # #^r"cd|jvS)Nfp8r/r.s r is_fp8z dtype.is_fp8s !!r"c|jdkS)Nrvr/rs r is_fp8e4nvzdtype.is_fp8e4nvyI%%r"c|jdkS)Nrwr/rs r is_fp8e4b8zdtype.is_fp8e4b8rr"c|jdkS)Nrur/rs r is_fp8e4b15zdtype.is_fp8e4b15yJ&&r"c|jdkS)Nrxr/rs r is_fp8e5zdtype.is_fp8e5yG##r"c|jdkS)Nryr/rs r is_fp8e5b16zdtype.is_fp8e5b16rr"c|jdkS)Nrzr/rs r is_fp16z dtype.is_fp16yF""r"c|jdkS)Nr{r/rs r is_bf16z dtype.is_bf16rr"c|jdkS)Nr|r/rs r is_fp32z dtype.is_fp32rr"c|jdkS)Nr}r/rs r is_fp64z dtype.is_fp64rr"c|jdkS)NrWr/rs r is_int1z dtype.is_int1rr"c|jdkS)Nrqr/rs r is_int8z dtype.is_int8rr"c|jdkS)Nrrr/rs r is_int16zdtype.is_int16rr"c|jdkS)NrZr/rs r is_int32zdtype.is_int32rr"c|jdkS)Nr^r/rs r is_int64zdtype.is_int64rr"c|jdkS)Nrsr/rs r is_uint8zdtype.is_uint8rr"c|jdkS)Nrtr/rs r is_uint16zdtype.is_uint16 yH$$r"c|jdkS)Nr\r/rs r is_uint32zdtype.is_uint32 rr"c|jdkS)Nr`r/rs r is_uint64zdtype.is_uint64rr"c(|jtjvSr,)r0rprrs r is_floatingzdtype.is_floatingsyEN**r"c(|jtjvSr,)r0rpSTANDARD_FP_TYPESrs r is_standard_floatingzdtype.is_standard_floatingsyE333r"c(|jtjvSr,)r0rprrs r is_int_signedzdtype.is_int_signedyE,,,r"c(|jtjvSr,)r0rprrs r is_int_unsignedzdtype.is_int_unsignedrr"cB|jtjtjzvSr,)r0rprrrs r is_intz dtype.is_intsyE,u/????r"c*|Sr,)rrs r is_boolz dtype.is_bool"||~~r"cl|tjtjztjztjzvSr,)rprrrr)type_strs r is_dtypezdtype.is_dtype%s(5+e.>>ORWRccccr"c td)NzNot implementedrar-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_ptr1rr"cdSrr-r-r"r is_constzdtype.is_const5rr"othercPt|tsdS|j|jkSr)rUrpr0r.rs r __eq__z dtype.__eq__9s(%'' 5yEJ&&r"c.|| Sr,rrs r __ne__z dtype.__ne__>;;u%%%%r"c,t|jfSr,)hashr0rs r __hash__zdtype.__hash__AsTYM"""r"c|Sr,r-rs r scalarz dtype.scalarD r"rk ir.builderrir.typec|jdkr|S|jdkr|S|jdvr|S|jdvr|S|jdvr|S|jdvr|S|jdkr|S|jdkr|S|jd kr| S|jd kr| S|jd kr| S|jd kr| S|jd kr| S|jdkr|S|jdkr|St!d|d)Nr~rW)rqrs)rrrt)rZr\)r^r`rxryrvrwrurzr{r|r}zfail to convert z to ir type)r0 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_tyrr.rks r to_irz dtype.to_irHs 9  &&(( ( Y& &&(( ( Y+ + +&&(( ( Y- - -'')) ) Y- - -'')) ) Y- - -'')) ) Y' ! !'')) ) Y* $ $**,, , Y) # #))++ + Y) # #))++ + Y* $ $**,, , Y& &&(( ( Y& &&(( ( Y& '')) ) Y& ((** *=D===>>>r"c|jSr,r/rs r __str__z dtype.__str__i yr"c|jdrd|jddzS|jdrd|jddzS|jS)Nfprbrbfbfloat)r0 startswithrs r codegen_namezdtype.codegen_namelsd 9   % % TYqrr]* * Y ! !$ ' ' dim+ +9 r"strc|jS)z"See cache_key_part() in triton.cc.r/rs r cache_key_partzdtype.cache_key_partts yr"c0d|S)z4Output of repr needs to be an evaluatable expressionztriton.language.)rrs r __repr__zdtype.__repr__ys7$"3"3"5"5777r"N)rrp)rkrrrrr)4r7rrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrr staticmethodrrrrrrrrpropertyrr r rrrr-r"r rprps:444J@@@JfffH888(KT5(5(5(n"""&&&&&&'''$$$'''##################$$$$$$$$$$$$%%%%%%%%%+++444------@@@dd\d..\.\\\'''' &&&&###X????BX88888r"rpcXeZdZdddZdd Zd Zd ZdZddZddZ e dZ dS) pointer_typer element_tyrp address_spacerXct|ts%tdt|jd||_||_d|d|_dS)Nzelement_ty is a rQzpointer<>)rUrprErjr7rrr0)r.rrs r rzpointer_type.__init__s_*e,, MKtJ/?/?/HKKKLL L$*,z,,, r"rkrrir.pointer_typec^||j|dS)Nr) get_ptr_tyrr rs r r zpointer_type.to_irs(!!$/"7"7"@"@!DDDr"c|jSr,r/rs r r zpointer_type.__str__r r"c*|Sr,r rs r rzpointer_type.__repr__rr"cdSNTr-rs r rzpointer_type.is_ptrtr"rrHcpt|tsdS|j|jko|j|jkSr)rUrrrrs r rzpointer_type.__eq__s9%.. 5%"22`t7IUM`7``r"c.|| Sr,rrs r rzpointer_type.__ne__rr"c|Sr,r-rs r rzpointer_type.scalarrr"NrrrprrX)rkrrr!)rrrrH) r7rrrr r rrrrrrr-r"r rrs-----EEEEaaaa &&&&Xr"rc6eZdZd d fd ZdZdZdd ZxZS)const_pointer_typerrrprrXcLt||dSr,)superr)r.rr __class__s r rzconst_pointer_type.__init__s# ]33333r"cd|jdS)Nzconst_pointerz'block_type.__init__..s,,,QW,,,r"rznumel (z') exceeds triton maximum tensor numel ()<, r ) rrErUrhr<numelTRITON_MAX_TENSOR_NUMELrr0)r.rr<rBs r rzblock_type.__init__s$  :899 9 eAh * * -,,e,,,E    A JJ!OJJJ :/ / /ttzttZqtttuu u8 88do888 r"rkrr ir.block_typech||j||jSr,) get_block_tyrr r<rs r r zblock_type.to_irs*##DO$9$9'$B$BDJOOOr"c|jSr,r/rs r r zblock_type.__str__r r"c*|Sr,r&rs r rzblock_type.__repr__rr"cdSr(r-rs r rzblock_type.is_blockr)r" List[int]c|jSr,)r<rs r get_block_shapeszblock_type.get_block_shapes zr"rrHcpt|tsdS|j|jko|j|jkSr)rUr;rr<rs r rzblock_type.__eq__s7%,, 5%"22PtzU[7PPr"c.|| Sr,rrs r rzblock_type.__ne__rr"c|jSr,r5rs r rzblock_type.scalars r"N)rrpr<r )rkrrrI)rrO)rr;rrH) r7rrrr r rrrQrrrrr-r"r r;r;s9999*PPPPQQQQ &&&&Xr"r;c$eZdZd dZdZd d Zd S) function_type ret_types List[dtype] param_typesrNonec"||_||_dSr,)rXrZ)r.rXrZs r rzfunction_type.__init__s"&r"c&d|jd|jS)Nzfn (z) -> )rZrXrs r r zfunction_type.__str__s=d&==T^===r"rkrc|fd|jD}fd|jD}||S)Nc:g|]}|Sr-r )rAtyrks r rCz'function_type.to_ir..s%GGG"((7++GGGr"c:g|]}|Sr-r`)rAret_typerks r rCz'function_type.to_ir..s%LLLX^^G,,LLLr")rZrXget_function_ty)r.rkir_param_typesrXs ` r r zfunction_type.to_irsPGGGGd6FGGGLLLLT^LLL &&~yAAAr"N)rXrYrZrYrr[)rkr)r7rrrr r r-r"r rWrWsQ''''>>>BBBBBBr"rWr~rWrqrrrZr^rsrtr\r`rxryrvrwrurzr{r|r}bitwidthrXsignedc6|dkrtS|dkr |rtS|dkr |stS|dkr |rtS|dkr |stS|dkr |rt S|dkr |st S|dkr |rtS|dkr |stStd|d|)NrrrrrzUnsupported bitwidth z and signedness ) rWrqrsrrrtrZr\r^r`r)rfrgs r get_int_dtyperis1}} Q6 Qv RF R RF R RF R SSS6SSTTTr"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)r7rrr6r-r"r rkrk.s Dr"rkceZdZdZdZd-dZdZdZdZd Z d Z d Z d Z d Z dZdZdZdZdZdZdZdZdZdZdZdZdZdZdZdZdZdZd Z d!Z!d"Z"d#Z#d$Z$d%Z%d&Z&d'Z'd(Z(d)Z)d*Z*d+Z+d,S).rhzL This class is used to store a value that is known at compile-time. cZt|tr|j|_dS||_dSr,rUrhri)r.ris r rzconstexpr.__init__>s, eY ' ' DJJJDJJJr"rrcd|jdS)Nz constexpr[]r?rs r rzconstexpr.__repr__Ds)DJ))))r"c|jSr,r?rs r __index__zconstexpr.__index__GrRr"cJt|jt|zSr,rhri_constexpr_to_valuers r __add__zconstexpr.__add__N &9%&@&@@AAAr"cJtt||jzSr,rhrurirs r __radd__zconstexpr.__radd__Q ,U33dj@AAAr"cJt|jt|z Sr,rtrs r __sub__zconstexpr.__sub__Trwr"cJtt||jz Sr,ryrs r __rsub__zconstexpr.__rsub__Wr{r"cJt|jt|zSr,rtrs r __mul__zconstexpr.__mul__Zrwr"cJt|jt|zSr,rtrs r __mod__zconstexpr.__mod__]rwr"cJtt||jzSr,ryrs r __rmul__zconstexpr.__rmul__`r{r"cJt|jt|z Sr,rtrs r __truediv__zconstexpr.__truediv__crwr"cJtt||jz Sr,ryrs r __rtruediv__zconstexpr.__rtruediv__fr{r"cJt|jt|zSr,rtrs r __floordiv__zconstexpr.__floordiv__i ':5'A'AABBBr"cJtt||jzSr,ryrs r __rfloordiv__zconstexpr.__rfloordiv__ls ,U33tzABBBr"cLt|jt|kSr,rtrs r __gt__zconstexpr.__gt__o &9%&@&@@AAAr"cLtt||jkSr,ryrs r __rgt__zconstexpr.__rgt__r ,U33dj@AAAr"cLt|jt|kSr,rtrs r __ge__zconstexpr.__ge__u ':5'A'AABBBr"cLtt||jkSr,ryrs r __rge__zconstexpr.__rge__x ,U33tzABBBr"cLt|jt|kSr,rtrs r __lt__zconstexpr.__lt__{rr"cLtt||jkSr,ryrs r __rlt__zconstexpr.__rlt__~rr"cLt|jt|kSr,rtrs r __le__zconstexpr.__le__rr"cLtt||jkSr,ryrs r __rle__zconstexpr.__rle__rr"cLt|jt|kSr,rtrs r rzconstexpr.__eq__rr"cLt|jt|kSr,rtrs r rzconstexpr.__ne__rr"c*t|jSr,)rHrirs r __bool__zconstexpr.__bool__DJr"c,t|j Sr,rhrirs r __neg__zconstexpr.__neg__$*%%%r"cJt|jt|zSr,rtrs r __and__zconstexpr.__and__rwr"cHt|jot|Sr,rtrs r logical_andzconstexpr.logical_ands B(;E(B(BCCCr"cJt|jt|zSr,rtrs r __or__zconstexpr.__or__rwr"cJt|jt|z Sr,rtrs r __xor__zconstexpr.__xor__rwr"cHt|jpt|Sr,rtrs r logical_orzconstexpr.logical_ors A':5'A'ABBBr"c,t|j Sr,rrs r __pos__zconstexpr.__pos__rr"c,t|jSr,rrs r __invert__zconstexpr.__invert__rr"cJt|jt|zSr,rtrs r __pow__zconstexpr.__pow__s %8%?%??@@@r"cJtt||jzSr,ryrs r __rpow__zconstexpr.__rpow__s ,U33TZ?@@@r"cJt|jt|z Sr,rtrs r __rshift__zconstexpr.__rshift__rr"cJt|jt|zSr,rtrs r __lshift__zconstexpr.__lshift__rr"c,t|j Sr,rrs r __not__zconstexpr.__not__sTZ(((r"c*t|jSr,)rDrirs r __iter__zconstexpr.__iter__rr"c|j|i|Sr,r?)r.rkwdss r __call__zconstexpr.__call__stz4(4(((r"Nr),r7rrr6rrrrrvrzr}rrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrr-r"r rhrh9s ****BBBBBBBBBBBBBBBBBBBBBBBBBBBCCCCCCBBBBBBCCCCCCBBBBBBCCCCCCCCCCCC   &&&BBBDDDBBBBBBCCC&&&&&&AAAAAACCCCCC)))   )))))r"rhc 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.) rUr=rhrjrrrirrp)ri shift_valuerfs r check_bit_widthrs%  Z Y%G%G:$7   ( ( X*XX8XXafalXXX      ( (r"ceZdZdZdjdZdkdZedld Zedld Zedld Z edld Z edld Z edldZ edldZ edldZedldZedldZedldZedldZedldZedldZedldZedldZedldZedldZedldZedldZedldZedldZedldZedld Zedld!Zedld"Z edld#Z!edld$Z"edld%Z#edld&Z$edld'Z%edld(Z&edld)Z'edld*Z(edld+Z)edld,Z*edld-Z+edld.Z,edld/Z-edld0Z.e/d1Z0edmdnd7Z1dod8Z2dod9Z3dod:Z4dpd<Z5dod=Z6dod>Z7dod?Z8dqdod@Z9drdodCZ:dodDZ;dsdodEZdtdodHZ?dtdodIZ@dtdodJZAdtdodKZBdtdodLZCdodMZDdodNZEdodOZFdodPZGdodQZHdodRZIdodSZJdudodTZKdudodUZLdodVZMdodWZNdodXZOdudodYZPdodZZQdvdod\ZRdwdod]ZSdvdod^ZTdwdod_ZUdqdod`ZVdqdodaZWdxdodcZXdxdoddZYdeZfdydhZ[dldodiZ\dS)zr=aRepresents 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. :) rjrpc||_|r|jnd|_d|_|jD]}|xj|zc_t |j|_||_|j|_d|jD|_dS)zNot called by user code.r-rc,g|]}t|Sr-rhr@s r rCz#tensor.__init__..s777qill777r"N)handlerr<rGrhrjrrp)r.rrjrBs r rztensor.__init__s #'==??:TZZ    A JJ!OJJJtz**  [ 77DJ777 r"rrct|jdzdd|jDzdzS)N[rFc34K|]}t|VdSr,)rr@s r z!tensor.__str__..s(0L0LAQ0L0L0L0L0L0Lr"rp)rrpjoinr<rs r r ztensor.__str__s<4:$tyy0L0L0L0L0L'L'LLsRRr"NcNt||}tj|||Sr,)rNraddr.rrs r rvztensor.__add__%5(++|D%222r"c0|||SNr)rvrs r rzztensor.__radd__||EH|555r"cNt||}tj|||Sr,rNrsubrs r r}ztensor.__sub__rr"cNt||}tj|||Sr,rrs r rztensor.__rsub__%5(++|E4222r"cNt||}tj|||Sr,)rNrmulrs r rztensor.__mul__rr"c0|||Sr)rrs r rztensor.__rmul__ rr"cNt||}tj|||Sr,rNrtruedivrs r rztensor.__truediv__ s&5(++eX666r"cNt||}tj|||Sr,rrs r rztensor.__rtruediv__s&5(++tX666r"cNt||}tj|||Sr,rNrfloordivrs r rztensor.__floordiv__s&5(++ uh777r"cNt||}tj|||Sr,rrs r rztensor.__rfloordiv__s&5(++ h777r"cNt||}tj|||Sr,rNrmodrs r rztensor.__mod__!rr"cNt||}tj|||Sr,rrs r __rmod__ztensor.__rmod__&rr"c,tj||Sr,)rminusr.rs r rztensor.__neg__,s~dH---r"c,tj||Sr,)rinvertrs r rztensor.__invert__0stX...r"cNt||}tj|||Sr,rNrand_rs r rztensor.__and__6%5(++}T5(333r"cNt||}tj|||Sr,rrs r __rand__ztensor.__rand__;%5(++}UD(333r"cNt||}tj|||Sr,rNror_rs r rz tensor.__or__@rr"cNt||}tj|||Sr,rrs r __ror__ztensor.__ror__Err"cNt||}tj|||Sr,rNrxor_rs r rztensor.__xor__Jrr"cNt||}tj|||Sr,rrs r __rxor__ztensor.__rxor__Orr"cnt||t||}tj|||Sr,rrNrshlrs r rztensor.__lshift__Ts5e$$$5(++|D%222r"cnt||t||}tj|||Sr,r rs r __rlshift__ztensor.__rlshift__Zs5t$$$5(++|E4222r"ct||t||}|jrt j|||St j|||Sr,rrNrprrashrlshrrs r rztensor.__rshift__`s_e$$$5(++ : # # % % 8=uh77 7=uh77 7r"ct||t||}|jrt j|||St j|||Sr,rrs r __rrshift__ztensor.__rrshift__is_t$$$5(++ : # # % % 8=h77 7=h77 7r"cNt||}tj|||Sr,rNr greater_thanrs r rz tensor.__gt__ss&5(++$T5(;;;r"cNt||}tj|||Sr,rrs r rztensor.__rgt__xs&5(++$UD(;;;r"cNt||}tj|||Sr,rNr greater_equalrs r rz tensor.__ge__~s&5(++%dE8<<>B&& D28+;TVT[Tc !Bb!B!BCCC r"cJd)zTransposes a 2D tensor.Fz0Transposition must be created by the AST Visitorr-rs r rztensor.Ts IHHHHr"Ffp_downcast_rounding Optional[str]bitcastrHct|tr|j}|rtj|||Stj||||S)z3 Alias for :py:func:`tensor.cast`. )rUrhrirrDcast)r.rprBrDrs r toz tensor.tosQ gy ) ) $mG  ;#D%:: :}T5(4HIIIr"cdSr,r-r.r<s r broadcast_toztensor.broadcast_to r"cdSr,r-r.dimss r transz tensor.transrKr"cdSr,r-rMs r permuteztensor.permuterKr"tuple[tensor, tensor]cdSr,r-rs r rz tensor.splitrKr"cdSr,r-rIs r viewz tensor.viewrKr"cdSr,r-rIs r reshapeztensor.reshaperKr"cdSr,r-)r.axiss r r8ztensor.expand_dimsrKr"cdSr,r-)r.rprBrDs r rFz tensor.castrKr"r-r)cdSr,r-)r.rimaskboundary_checkcache_modifiereviction_policys r storez tensor.storerKr"cdSr,r-)r.offsetss r advanceztensor.advancerKr"cdSr,r-)r.cmpvalsemscopes r atomic_casztensor.atomic_casrKr"cdSr,r-r.rfr\rgrhs r atomic_xchgztensor.atomic_xchg rKr"cdSr,r-rks r atomic_addztensor.atomic_add rKr"cdSr,r-rks r atomic_maxztensor.atomic_maxrKr"cdSr,r-rks r atomic_minztensor.atomic_minrKr"cdSr,r-rks r atomic_andztensor.atomic_andrKr"cdSr,r-rks r atomic_orztensor.atomic_orrKr"cdSr,r-rks r atomic_xorztensor.atomic_xorrKr"cdSr,r-rs r expz tensor.exprKr"cdSr,r-rs r logz tensor.log!rKr"cdSr,r-rs r cosz tensor.cos$rKr"cdSr,r-rs r sinz tensor.sin'rKr"cdSr,r-rs r sqrtz tensor.sqrt*rKr"cdSr,r-rs r rsqrtz tensor.rsqrt-rKr"cdSr,r-rs r rSz tensor.abs0rKr"cdSr,r-)r.rY combine_fn keep_dimss r reducez tensor.reduce3rKr"cdSr,r-)r.rYrreverses r associative_scanztensor.associative_scan6rKr"cdSr,r-)r.num_binss r histogramztensor.histogram9rKr"cdSr,r-)r.divs r cdivz tensor.cdiv<rKr"cdSr,r-rs r sigmoidztensor.sigmoid?rKr"cdSr,r-)r. ieee_roundings r softmaxztensor.softmaxBrKr"cdSr,r-rs r ravelz tensor.ravelErKr"TcdSr,r-r.rYreturn_indicesreturn_indices_tie_break_leftrs r maxz tensor.maxHrKr"cdSr,r-r.rYtie_break_leftrs r argmaxz tensor.argmaxKrKr"cdSr,r-rs r minz tensor.minNrKr"cdSr,r-rs r argminz tensor.argminQrKr"cdSr,r-r.rYrs r sumz tensor.sumTrKr"cdSr,r-rs r xor_sumztensor.xor_sumWrKr"rcdSr,r-r.rYrs r cumsumz tensor.cumsumZrKr"cdSr,r-rs r cumprodztensor.cumprod]rKr"r>rh descendingcdSr,r-)r.r>rs r sortz tensor.sort`rKr"cdSr,r-)r.r>s r flipz tensor.flipcrKr")rjrprr,NFNrprprBrCrDrH)rr=rrRr)Nr-r)r)NNNNNF)NFTF)TF)rF)r>rhrrhrr=)]r7rrr6rr r&rvrzr}rrrrrrrrrrrrrrrrrrr rrrrrrrrrrrr+rr0rrrr@rrrGrJrOrQrrUrWr8rFr`rcrirlrnrprrrtrvrxrzr|r~rrrrSrrrrrrrrrrrrrrr CONSTEXPR_0rrr-r"r r=r=sT & 8 8 8 8SSSS 333 W3 666 W6 333 W3 333 W3 333 W3 666 W6 777 W7 777 W7 888 W8 888 W8 333 W3 333 W3  ... W. /// W/  444 W4 444 W4 333 W3 333 W3 444 W4 444 W4 333 W3  333 W3  888 W8 888 W8 <<< W< <<< W<  === W= === W=  999 W9 999 W9  ::: W: ::: W:  555 W5 555 W5 999 W9 999 W9 ;;; W; ::: W:  --- W-     W IIXI  J J J J W J"                                                                                                                                                                                     %)+            r"r=cPtj|d}|dkp |dkp|dkS)N01trueon)osgetenv)var_namevs r get_bool_env_varrgs/ (C  A 8 /qF{ /a4i/r"c>t|tr|jS|Sr,rn)rs r ruruos !Yw Hr"cJt|}tj||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 )rur program_idrYrs r rrus$ t $ $D  tX . ..r"cJt|}tj||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 )rur num_programsrs r rrs$ t $ $D  x 0 00r"cjt|}t|}tj|||S)a 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 = 131072` :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 )rurarange)r9endrs r rrs2  & &E c " "C ?5#x 0 00r"c t|}t|D]\}}t|trt |}t|tst d|dt|jts(t d|dt|jd|j|jdz zdkrtd|dd|DS) NzShape element  must have type `constexpr`1 must have type `constexpr[int]`, got `constexpr[rprrz must be a power of 2c,g|]}t|Sr-)rurArFs r rCz%_shape_check_impl..s! 2 2 2q  " " 2 2 2r") rur7rUrXrhrErirjr)r<ids r _shape_check_implrs  & &E%  HH1 a   ! A!Y'' MKQKKKLL L!'3'' sqQqqaefgfmananqqqrr r 7agk "a ' 'FaFFFGG G ( 2 2E 2 2 22r"ct|}t|}t|}tj||||S)a] 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, ) :value value: A scalar value to fill the array with :type shape: tuple of ints :param dtype: Data-type of the new array, e.g., :code:`tl.float16` :type dtype: DType )rrurfull)r<rirprs r rrsA e $ $E  & &E  & &E =uh 7 77r"c.tj|||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 )rbroadcast_impl_value)inputrrs r broadcastrs  (x @ @@r"rcftt|}tj|||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) )rrGrbroadcast_impl_shaperrr<s r rJrJs/" .u55 6 6E  (x @ @@r"rc6|sd}tj|||S)ap Permutes the dimensions of a tensor. If no permutation is specified, tries to do a (1,0) permutation, i.e. tries to transpose a 2D tensor. :param input: The input tensor. :param dims: The desired ordering of dimensions. For example, :code:`(2, 1, 0)` reverses the order dims in a 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. )rr)rrQrrrNs r rOrOs&,   E4 2 22r"cLt|}tj|||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 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 do a (1,0) permutation. )rGrrQrs r rQrQ s&( D ! !D  E4 2 22r"Fc0tj||||S)av Concatenate the given blocks :param input: The first input tensor. :type input: :param other: The second input tensor. :type other: :param 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) )rcat)rr can_reorderrs r rr$s <uk8 < <.s, G G GQJ*1--x 8 8 G G Gr"z7expand_dims received duplicate axes, normalized axes = ) rNrurUr r8r4r<setrsortedrr8)rrYraxesr=rrs @r r8r8s uh ' 'E t $ $D#D(33 ?4:::$D5;#d))+H G G G G$ G G GD 3t99~~T""YSWYYZZZ C D\\55"3844 Jr"rBrCrDct||}t|tr|j}|rt j|||St j||||S)a@ Casts a tensor to the given :code:`dtype`. :param dtype: The target data type. :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). :param bitcast: If true, the tensor is bitcasted to the given :code:`dtype`, instead of being numerically casted. )rNrUrhrirrDrF)rrprBrDrs r rFrFs` uh ' 'E'9%% -8uh777 =x1E F FFr"c | | Jd|/|o d|jjv}|r|s|dnd} tjd| }t |}t |}t |}t j|||||||S)a Returns the matrix product of two blocks. The two blocks must be two-dimensional and have compatible inner dimensions. :param input: The first tensor to be multiplied. :type input: 2D 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 tensor of scalar-type in {:code:`int8`, :code: `float8_e5m2`, :code:`float16`, :code:`bfloat16`, :code:`float32`} :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"`. Avaliable options for amd: :code:`"ieee"`. :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 specifiedtf32ieeeTRITON_F32_DEFAULT)optionsallowed_dot_input_precisionsrrrurdot) rraccinput_precision allow_tf32max_num_imprecise_acc out_dtyper supports_tf32default_precisions r rrs*  "j&8&8:w&8&8&8 \Vx/?/\%\ '4h*h HZFFbh)$8:KLL)/::O#I..I/0EFF <uc?'5$P_ai j jjr"basec 6tj|||||||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 )rmake_block_ptr)rr<stridesrb block_shapeorderrs r rrs"  "4+uV^ _ __r"c.tj|||S)z Advance a block pointer :param base: the block pointer to advance :param offsets: the offsets to advance, a tuple by dimension )rrc)rrbrs r rcrcs  D'8 4 44r"r0rhas_cmpCallable[[T], T]cdfd }|S)Nfuncrrc: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: Memory semantics to use ("ACQUIRE_RELEASE" (default), "ACQUIRE", "RELEASE", or "RELAXED") :type sem: str :param scope: Scope of threads that observe synchronizing effect of the atomic operation ("GPU" (default), "CTA", or "SYSTEM") :type scope: str )r6)r#docstrr r0s r _decoratorz&_add_atomic_docstr.._decoratorsT9999  : : :F     r"r#rrrr-)r0r r&s`` r _add_atomic_docstrr(s/2 r"zcompare-and-swapT)r ct||}t||}t|}t|}tj||||||Sr,)rNrurri)rrerfrgrhrs r ririsW S( # #C S( # #C c " "C  & &E  wS#uh G GGr"exchangect||}t|}t|}t|}tj||||||Sr,)rNrurrlrrfr\rgrhrs r rlrlsU S( # #C c " "C  & &E t $ $D  dC I IIr"rct||}t|}t|}t|}tj||||||Sr,)rNrurrnr,s r rnrnU S( # #C c " "C  & &E t $ $D  wT3x H HHr"rct||}t|}t|}t|}tj||||||Sr,)rNrurrpr,s r rprpr.r"rct||}t|}t|}t|}tj||||||Sr,)rNrurrrr,s r rrrrr.r"z logical andct||}t|}t|}t|}tj||||||Sr,)rNrurrtr,s r rtrtr.r"z logical orct||}t|}t|}t|}tj||||||Sr,)rNrurrvr,s r rvrv sU S( # #C c " "C  & &E t $ $D  gsD#uh G GGr"z logical xorct||}t|}t|}t|}tj||||||Sr,)rNrurrxr,s r rxrxr.r"ct||}t||}t||}tj||||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. )rNrwhere) conditionrFyrs r r5r5$sG"9h//I1hA1hA >)Q8 4 44r" propagate_nanct||}t||}t||}t||}t|}tj||||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)rN_promote_bfloat16_to_float32rurminimumrFr7r8rs r r;r;@j 1hA1hA$Q:::A$Q:::A' 66M  Aq- : ::r"ct||}t||}t||}t||}t|}tj||||S)aH Computes the element-wise maximum 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)rNr:rurmaximumr<s r r?r?Vr=r"ct||}t||}t||}t||}t||}t||}t|}tj|||||S)a< Clamps the input tensor :code:`x` within the range [min, max]. Behavior when :code:`min` > :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)rNr:rurclamp)rFrrr8rs r rArAls$ 1hA S( # #C S( # #C$Q:::A &sX > > >C &sX > > >C' 66M >!S#}h ? ??r"return_indices_arg tie_break_argcdfd }|S)Nr#rrctd} |dddz } |ddz }||_|S)Na Returns the {name} of all elements in the :code:`input` tensor along the provided :code:`axis` :param input: the input values :param axis: the dimension along which the reduction should be done :param keep_dims: if true, keep the reduced dimensions with length 1z :param z-: if true, return index corresponding to the z valuezR: if true, return the left-most indices in case of ties for values that aren't NaNr/formatr6)r#r%r0rBrCs r r&z)_add_reduction_docstr.._decoratorsL  ) \ \\MQ\\\ \F  $ p ppp pF}}$}//  r"r'r-)r0rBrCr&s``` r _add_reduction_docstrrHs5" r"c#dK|}dV||dSr,)get_insertion_pointrestore_insertion_point)rkips r _insertion_guardrMs:  $ $ & &B EEE ##B'''''r"cttrtf|dSfd}fdtt|}(t t djtj|}|r;tfd|D}ntfd|D}|S)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 :param axis: the dimension along which the reduction should be done. If None, reduce all dimensions :param combine_fn: a function to combine two groups of scalar tensors (must be marked with @triton.jit) :param keep_dims: if true, keep the reduced dimensions with length 1 rrr(rcd D}t||dz}|d}t 5 fd|jD} ||fdt |jD}  |i}t|tr |j g}n d|D} j |ddddS#1swxYwYdS)Nc&g|]}|jjSr-rjrrAts r rCz7reduce..make_combine_region..6661666r"rrc:g|]}|Sr-r`rArars r rCz7reduce..make_combine_region..%NNN"288H--NNNr"c\g|](\}}t||)Sr-r=argrArrablocks r rCz7reduce..make_combine_region..1[[[BF599Q<<,,[[[r"rcg|] }|j Sr-rrArs r rCz7reduce..make_combine_region..55518555r") rW get_regionrMrZcreate_block_with_parentr7call_JitFunctionrUr=rcreate_reduce_ret) reduce_op in_scalar_tys prototyperegionrZrresultshandlesr]rr(rrs @r make_combine_regionz#reduce..make_combine_regions`66666 !-1BCC %%a(( h ' ' 1 1NNNN 8MNNNK55fkJJE[[[[)IDY:Z:Z[[[D 11*d21NNG'6** 6">*55W555 &H & 0 0 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1BC))C-0C-cZtj|D]}t|d}|S)Nrr)builtinsranger8)rTndims_rs r expand_ndimszreduce..expand_ndimss7&& 5 5AAq8444AAr"Nc3<K|]}t|VdS)rN)r8)rArTrrYs r rzreduce..s2MMA Ath???MMMMMMr"c3`K|](}|tdjV)dS)rN)r4r<)rArTrvrs r rzreduce..s=JJ QE!HN(;(;<<JJJJJJr") rUr=rrurr4r<r reductiontuple) rrYrrrr(ror=rvs ``` `` @r rrs9%  uuizYQYfpqqqrstt11111111  t $ $D#I..I $E!HN 3 344  UD*=x H HCK  MMMMMMMMMMCCJJJJJcJJJJJC Jr"ch|jj}|tur|t|S|Sr)rjrbfloat16rGre)rTr scalar_tys r r:r:s3 IHttGht/// Hr"ct|}|j|}td||}t|jdkr\dt jt|jD}||=t |||}t||j|}t||f|||||\} } | | fS)Nrrrc,g|]}t|Sr-r)rArs r rCz(_reduce_with_indices..sQQQ1)A,,QQQr"rO) rur<rr4rrrsr8rJr) rrYrrrr(nindexaxes_to_expandrvaluerindicess r _reduce_with_indicesrs t $ $D DA 1a( + + +E 5;!QQs5;?O?O0P0PQQQ 4 E>HEEEUEK(CCCuendJ)^f)3555FH 8 r"cdfd }|S)Nr#rrcBd}||_|S)Nz Returns the {name} of all elements in the :code:`input` tensor along the provided :code:`axis` :param input: the input values :param axis: the dimension along which the scan should be doner/rF)r#r%r0s r r&z$_add_scan_docstr.._decorators'F }}$}//  r"r'r-)r0r&s` r _add_scan_docstrrs) r"cttrtf||dSfd}t|}|(t |t dj}tj|||S)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 :param axis: the dimension along which the reduction should be done :param combine_fn: a function to combine two groups of scalar tensors (must be marked with @triton.jit) :param reverse: apply the associative scan in the reverse direction along axis. rrcd D}t||dz}|d}t 5 fd|jD} ||fdt |jD}  |i}t|tr |j g}n d|D} j |ddddS#1swxYwYdS)Nc&g|]}|jjSr-rRrSs r rCzAassociative_scan..make_combine_region..rUr"rrc:g|]}|Sr-r`rWs r rCzAassociative_scan..make_combine_region..rXr"c\g|](\}}t||)Sr-rZr\s r rCzAassociative_scan..make_combine_region..!r^r"r_cg|] }|j Sr-rarbs r rCzAassociative_scan..make_combine_region..&rdr") rWrerMrZrfr7rgrUr=rcreate_scan_ret) scan_oprjrkrlrZrrmrnr]rr(rrs @r roz-associative_scan..make_combine_regions`66666 !-1BCC ##A&& h ' ' / /NNNN 8MNNNK55fkJJE[[[[)IDY:Z:Z[[[D 11*d21NNG'6** 6">*55W555 $H $g . . / / / / / / / / / / / / / / / / / /rp)rUr=rrurr4r<r)rrYrrrr(ros` ` `` r rr s%  s 4Wxdnooopqrr//////// t $ $D $E!HN 3 344  $UD2EwPX Y YYr"cLt|}tj|||S)zcomputes 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 :param num_bins: number of histogram bins )rurr)rrrr(s r rr/s&#8,,H  eXx 8 88r"c*tj|S)zA Insert a barrier to synchronize all threads in a block. )r debug_barrierrs r rrAs  !( + ++r"c rt|tr|g}t|D]o\}}t|tstd|dt|jt s(td|dt |jdpd|D}tj||S)zd Let the compiler know that the values in :code:`input` are all multiples of :code:`value`. values element rrrpcg|] }|j Sr-r?rs r rCzmultiple_of..U & & &!ag & & &r") rUrhr7rErirXrjr multiple_ofrr9rrrs r rrIs &)$$&!!tt1!Y'' NLaLLLMM M!'3'' trarrbfghgnboborrrss s t & &v & & &F  v . ..r"c rt|tr|g}t|D]o\}}t|tstd|dt|jt s(td|dt |jdpd|D}tj||S)z^ Let the compiler know that the `value` first values in :code:`input` are contiguous. rrrrpcg|] }|j Sr-r?rs r rCz"max_contiguous..err") rUrhr7rErirXrjrmax_contiguousrs r rrYs &)$$&!!tt1!Y'' NLaLLLMM M!'3'' trarrbfghgnboborrrss s t & &v & & &F  "5& 1 11r"c rt|tr|g}t|D]o\}}t|tstd|dt|jt s(td|dt |jdpd|D}tj||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]. rrrrpcg|] }|j Sr-r?rs r rCz!max_constancy..xrr") rUrhr7rErirXrjr max_constancyrs r rris&)$$&!!tt1!Y'' NLaLLLMM M!'3'' trarrbfghgnboborrrss s t & &v & & &F  !% 0 00r"  )seprfileflushrrrcdS)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=}") Nr-)rrrrrr9s r static_printrs  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-)condmsgrs r static_assertrs  Dr")hexrc4ddl}t|}t|ts J|dd}|D]}||jvrd}n|s J|dg}|D]%}|t ||&tj||||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) stringrurUr printableappendrNr device_print) prefixrrrrb_asciichnew_argsr[s r rrsDMMM  ( (F fc " "==v$=$=$====G V% % %G E & 66v666666H33 3112222  3 A AAr"ct|}ddl}|j}|j|}t |dr'|j}|j|}t |d'd}d}d}|0|j)|jj}|jjj}|jj }tj 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. rNr7unknown) rur2 currentframe getmodulerf_backf_codeco_name co_filenamef_linenor device_assertrN) rrrr2framemodulelineno func_name file_names r rrs& c " "CNNN G " "E W u % %F &* % %* ""5)) &* % %*FII U\5L( L'3 &  !*T8"<"; 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](https://llvm.org/docs/LangRef.html#inline-asm-constraint-string) :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 :param _builder: the builder :return: one tensor or a tuple of tensors of the given dtypes TFc0g|]}t|Sr-rM)rAr[rs r rCz*inline_asm_elementwise..c s#CCCsC22CCCr")rkarithmetic_check allow_lhs_ptr allow_rhs_ptrrc:g|]}t|jSr-)r;r<)rAdt broadcast_args r rCz*inline_asm_elementwise..s s&KKKrz"m&9::KKKr"cg|] }|j Sr-rarSs r rCz*inline_asm_elementwise..t s///Aqx///r"c:g|]}|Sr-r`rWs r rCz*inline_asm_elementwise..u s&AgAgAgY["((8BTBTAgAgAgr"c3dK|]*\}}t||V+dSr,)r= get_result)rArracalls r rz)inline_asm_elementwise..y s=PPEAr**B//PPPPPPr")rurDrErrFr _DtypeClassrrbinary_op_type_checking_implr<r7create_inline_asmr=rrz)rrrrprrrhas_multiple_outputsres_tys dispatch_argsbin_op_type_checkingitemrurrnrrs ` @@r inline_asm_elementwiserst c " "C%k22K t $ $D!'**G U # $  K-u 5 5EGCCCCdCCCC}L&  1"    &a( ! I ID33D-HH A}}   L$]33 P P4&:&:4&O&O# a !!KKKKUKKKG/////G  % %c;AgAgAgAg_fAgAgAgiprv w wD 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. Nc&t|tsJ|td|_nt|tsJ||_|td|_||_dSt|tsJ||_||_dSNrr)rUrhr;r9r)r.arg1arg2r;s r rzstatic_range.__init__ s$ ***** <!! DIIdI.. . . .DI <"1DJDHHHdI.. . . .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"rr7rrr6rrrr-r"r rr sX"    WWWWWWWWr"rc&eZdZdZddZdZdZdS)rsa 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. Nc|td|_n||_|td|_||_n||_||_||_dSr)rhr;r9r num_stages)r.rrr;rs r rzrange.__init__ sQ <!! DIIDI <"1DJDHHDJDH$r"c tdNz4tl.range can only be used in @triton.jit'd functionsrrs r rzrange.__iter__ QRRRr"c tdrrrs r rzrange.__next__ rr"rrr-r"r rsrs sX0 % % % %SSSSSSSSr"rslib_namelib_pathr8arg_type_symbol_dictdict ret_shaperzc ^t|dkrtdtt|d}t||kr"tdt|d|g} g} |D]} t | t r5| | j| | jL| t| | | t| } | |vr'td|d| || d} || d} |rt| |} 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_shape: the shape of the return value :param _builder: the builder :return: the return value of the function rzarg_type_symbol_dict is emptyz+length of input args does not match.Expect z, got z,input arg type does not match.Expect one of r) r4rr8r5rUr=rrprrjrzr;r )r#rrrrrrrnum_args arg_typesarg_listr[symbolrcs r dispatchr s   A%%89994,113344Q788H 4yyH?#&t99??4<??@@ @IH!! c6 " " !   SY ' ' ' OOCJ ' ' ' '   T#YY ' ' ' OOC i  I,,,Y*>*C*C*E*EYYMVYYZZ Z&i03' 215  7!(I66Hdd8XvxPXAYAY[bccemnnnr"c |}d}d}g} tjt|D]\} t || ||| <| || j|| jrd}]t| dkrt| } d} | |vrd} |d} |D]} tj | | || \}} tjt|D]&} tj || | || \|| <}'|s| j }|j }t||||||||S)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 :param _builder: the builder :return: the return value of the function TNFr)r)copyrrrsr4rNrrprjrrzrrr<create_extern_elementwiser)rrrrrrr all_scalarrrrrrrrur#s r extern_elementwiser sIIKKMJII ^C .. / /%mA&6AA aq)/000   ) ) + + J 9~~)$$  , , ,$ %a( ! h hD'DT=ZbVf h h h A}}M 2 233 k kA"*"G VWHXZgiqYi#k#k#k M! aa ,%+I  -D D(Hm=QS\^ego p ppr"c.tj|||S)z Convert both operands to a single common type :param lhs: the left operand :param rhs: the right operand :param builder: the builder )rr)lhsrhsrks r binary_op_type_legalizationr1 s  0c7 C CCr"c t|S)z#A decorator for external functions.)r&rKs r externr; s 2;;r")rrrrr8r,)rfrXrgrHrrp)rr=)FNrrrr)NNr-r)r)r)FN)Nr-r)r)N)rr=r)r0rr rHrr!r)NNNN)r8rh)r0rrBrrCrrr!)FNN)r0rrr!)rrrr)r)N) rrrrrr rprrrHrrX) rrrrrr8rrrrzrrH) rrrrrr8rrrrH) __future__rwarningsr contextlibrenumr functoolsrrrr r r r r rrr runtime.jitrr2r _C.libtritonrr)rrrHr% PROPAGATE_NAN PropagateNanr&rBrGr<rOrNrprrr0r;rWr~rWrqrrrZr^rsrtr\r`float8e5 float8e5b16 float8e4nv float8e4b8 float8e4b15float16r|rergpi32_trirkrhrrr=rrurrrrrrrJrOrQrrrrrUrWrr8rFrr rrr`rrcr(rirlrnrprrrtrvrxr5NONEr;r?rArHrMrr:rrrrrrrrrrrrrrrsrrrrr-r"r rs""""""%%%%%%$$$$$$$$ EEEEEEEEEEEEEEEE  GCLL!%      ))))X   ,....  ### #DDDDc8c8c8c8c8c8c8c8R      5   Faaaaaaaa"00000000f B B B B BE B B B  uV}} uV}} uV}} g g g g x x x 5>>eJ U9   U9   eJ %-- 5== %-- %-- e  UUUU8        B)B)B)B)B)B)B)B)Jill Z Z Z Z Z Z Z Z z 000    /// /$ 111 1  111 1" 3 3 3  8 8 8  8*  A A A  A)-AAAA A&)-33333 32#'3333 3,  = = =  = ))) ).    B!%NNNN N(',tAAAA A&... 8GGGG G4 ]amthhh hJ rt"&7-7-7- 7-t SSS S MMM M,k,k,k ,k^  ` ` ` `  `555 5<&555HHH65 HJJJJ  JEIII IEIII IEIII IM""III#" IL!!HHH"! HM""III#" I 555 56 -9->;;;; ;* -9->;;;; ;* 2>2Cd@@@@ @D,((( +++ +\        ,    ZZZ ZD999 9  ,,, ,  / / /  /  2 2 2  2 111 1. %(TE\`                  $)D-B-B-B-B -B` %k%k%k %kP >BQQQQ QN%W%W%W%W%W%W%W%WP*S*S*S*S*S*S*S*Sf&*(o(o(o(o(oV $'q'q'q'q 'qTDDDr"