nVh2 ddlmZddlmZddlmZddlmZddlm Z m Z ddl Z ddl m Z m Z mZmZmZmZmZddlZdd lmZddlZddlZdd lmZd d lmZd d lmZmZedZdZ ejBZ"ddZ#ddZ$dZ%ddZ&e#ddZ'GddZ(GddZ)e)dZ*dZ+dZ,GddZ-GddZ.Gdd e.Z/e/Z0Gd!d"e/Z1Gd#d$e1Z2Gd%d&e/Z3Gd'd(e.Z4Gd)d*e/Z5e/d+Z6e/d,Z7e/d-Z8e/d.Z9e/d/Z:e/d0Z;e/d1Z<e/d2Z=e/d3Z>e/d4Z?e/d5Z@e/d6ZAe/d7ZBe/d8ZCe/d9ZDe/d:ZEe/d;ZFe/d<ZGe/d=ZHe1e:ZIdd>ZJGd?d@e-ZKGdAdBe-ZLGdCdDZMGdEdFe.ZNGdGdHe-ZOGdIdJeNZPGdKdLeOZQdMZRdNZSe#ddOZTe#ddPZUe#ddQZVdRedSeV_WdTZXdUZYe#ddVZZe#ddWZ[e$e#ddXdYZ\e$e#ddXddZZ]e$e#ddXd[Z^e#dd]Z_e#dd^Z`ed_Zae$e#ddd`Zbe$e#ddXdaZce$e#d\ddbdcZdddZee$e#ddeZfe$e#dddfZge#ddddeGdfdgZhe#dd\eGdfdhZie# ddiZje# d ddjZke#ddkZle#ddlZme$e#ddmZne#dddnZoe$e#ddoZpe# d ddpZqdddqZre$e#erdrdstdduZse$e#erdvddwZte$e#erdxddyZue$e#erdzdd{Zve$e#erd|dd}Zwe$e#erd~ddZxe$e#erdddZye$e#erdddZze#ddZ{e#dddZ|e#dddZ}e#dddZ~e#e"jdfddZe#e"jdfddZe#e"jdfddZ d ddZedZe$e#ddZe#ddZe#ddZddZe$e#ddZe$e#ddZe$e#ddZe#ddZe#ddZe#ddZe#ddZe#ddZe#dddd\ddddZe#ddZe#d\dddZe#ddZe# d ddZGddZGddZ d ddZe# dddZdZdZy)) annotations)warn)contextmanager)Enum)partialwrapsN)UnionCallableListSequenceTypeVarOptionalTuple)jit)ir)semantic)TRITON_MAX_TENSOR_NUMELvalidate_block_shapeT__triton_builtin__cjtsJtfd}t|td|S)zMark a function as a builtin.c<d|vs|d td|i|S)N_builderzdDid you forget to add @triton.jit ? (`_builder` argument must be provided outside of JIT functions.) ValueErrorargskwargsfns D/home/dcms/DCMS/lib/python3.12/site-packages/triton/language/core.pywrapperzbuiltin..wrappers: V #vj'9'A`a a4"6""T)callablersetattrTRITON_BUILTIN)r!r#s` r"builtinr(s8 B<< 2Y##  G^T* Nr$c tsJtj}t|jj ddhz dkD}j sd_xj djd|rdnddjd |rd ndd z c_fd }t|jj}|d jd|d <|j|}||_ 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..wrapperCs4"6""r$rselfname) parameterszForwards to :py:func:`z` free functionT)r%inspect signaturelenr3keys__doc____name__listvaluesreplace __signature__ is_builtinr&r'tensor)r!orig_sighas_argsr# new_paramsnew_sigs` r"_tensor_member_fnrD)s3 B<<  $H8&&++-\0JJKaOH :: JJ Qub9: KK=h7B78J # h))0023JqM))v)6JqM*5G#G.r{{m?KGO"~. FBKK) Ir$cjt|dk(r t|d|dS|S#t$rY|SwxYw)z7Returns x[0] if x has one element and x[0] is iterable.rr)r6iter TypeErrorxs r"_unwrap_iterablerJUsF 1v{  1JQ4K H   H s % 22c$t|tdS)z-Is this a registered triton builtin function?F)getattrr'r!s r"r>r>ks 2~u --r$c.tj||Sr.)r to_tensor)rIrs r"rOrOps   a **r$ceZdZdZy)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)r9 __module__ __qualname__r8r/r$r"rQrQzs  r$rQceZdZdZdZd+dZdZdZdZdZ dZ d Z d Z d Z d Zd ZdZdZdZdZdZdZdZdZdZdZdZdZdZdZdZdZdZ dZ!d Z"d!Z#d"Z$d#Z%d$Z&d%Z'd&Z(d'Z)d(Z*d)Z+y*), constexprzL This class is used to store a value that is known at compile-time. ct|tr|j|_t|_y||_t|_yr.) isinstancerUvaluetype)r0rXs r"__init__zconstexpr.__init__s0 eY 'DJ DJ r$c"d|jdS)Nz constexpr[]rXr0s r"__repr__zconstexpr.__repr__sDJJt|tr |jS|Sr.rWrUrX)os r"_unwrap_if_constexprrs I.1775A5r$c t|trkt|trZ|jjj }|j |k\r*td|j d|d|jdyyyy)NzValue z exceeds the maximum bitwidth (z ) for type 'z)'. This may result in undefined behavior.) rWr?rUrYscalarprimitive_bitwidthrXrdtype)rX shift_valuebitwidths r"check_bit_widthrs% Z Y%G::$$77    ( **++J8*T`afalal`mnWX  )&H r$c$eZdZUdZded<ddZy) base_valuezLBase class of values that exist in the triton IR (i.e. not constexprs). base_typerYct)znFlatten frontend value into a sequence of mlir handles, which are appended to the output list NotImplementedErrorr0handless r" _flatten_irzbase_value._flatten_ir s "!r$NrList[ir.value]rNone)r9rRrSr8__annotations__rr/r$r"rrs O"r$rc eZdZdZdZddZy)rctd)NzTypes must implement __eq__rrfs r"rzbase_type.__eq__)s!"?@@r$c||k( Sr.r/rfs r"rzbase_type.__ne__,EM""r$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. rr0rcursors r" _unflatten_irzbase_type._unflatten_ir/s "!r$NrrrintrzTuple[base_value, int])r9rRrSrrrr/r$r"rr'sA#"r$rceZdZgdZgdZgdZgdZdgZGddeZ 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-e(d,Z.d7d-Z/d.Z0e1d/Z2d8d0Z3d1Z4d2Z5e1d9d3Z6d4Z7d:d5Z8y6);r)int8int16int32int64)int1uint8uint16uint32uint64) fp8e4b15fp8e4nvfp8e4b8fp8e5fp8e5b16fp16bf16fp32fp64)rrrrvoidceZdZdZdZy)dtype.SIGNEDNESSrrN)r9rRrSSIGNEDUNSIGNEDr/r$r" SIGNEDNESSrCs r$rceZdZdZdZdZy) dtype.KINDrrrN)r9rRrSBOOLEANINTEGRALFLOATINGr/r$r"KINDrGsr$rct|}||_|tjtjztj ztj zvsJ||tjvrStjj|_ t|jdd|_ |j|_ y|tjvrStjj|_ t|jdd|_ |j|_ y|tj vr|dk(rd|_d|_ d|_y|dk(rd|_d|_ d|_y|d k(rd|_d|_ d|_y|d k(rd |_d|_ d|_y|d k(rd |_d|_ d |_y|dk(rd|_d |_ d|_y|dk(rd|_d |_ d|_y|dk(rd|_d|_ d|_y|dk(rd|_d|_ d|_yt#d||dk(rd|_ yy)Nrrrrrrrr rr r4@iz Unsupported floating-point type rr)rr2r SINT_TYPES UINT_TYPESFP_TYPES OTHER_TYPESrrint_signednessrsplit int_bitwidthrrfp_mantissa_width exponent_bias RuntimeErrorr0r2s r"rZzdtype.__init__Ls0#D) u''%*:*::U^^KeN_N___eaee_ 5## #"'"2"2"9"9D  #DJJu$5b$9 :D &*&7&7D # U%% %"'"2"2";";D  #DJJu$5b$9 :D &*&7&7D # U^^ #z!)*&*+'%'"")*&*+'%&"")*&*+'%&")*&*+'%'"#)*&*+'%'")+&*,'%'")*&*,'%(")+&*,'%(")+&*,'%)""%EdV#LMM V^&'D #r$cd|jvS)Nfp8r1r^s r"is_fp8z dtype.is_fp8s !!r$c |jdk(S)Nrr1r^s r" is_fp8e4nvzdtype.is_fp8e4nvyyI%%r$c |jdk(S)Nrr1r^s r" is_fp8e4b8zdtype.is_fp8e4b8rr$c |jdk(S)Nrr1r^s r" is_fp8e4b15zdtype.is_fp8e4b15yyJ&&r$c |jdk(S)Nrr1r^s r"is_fp8e5zdtype.is_fp8e5yyG##r$c |jdk(S)Nrr1r^s r" is_fp8e5b16zdtype.is_fp8e5b16r!r$c |jdk(S)Nrr1r^s r"is_fp16z dtype.is_fp16yyF""r$c |jdk(S)Nrr1r^s r"is_bf16z dtype.is_bf16r)r$c |jdk(S)Nrr1r^s r"is_fp32z dtype.is_fp32r)r$c |jdk(S)Nrr1r^s r"is_fp64z dtype.is_fp64r)r$c |jdk(S)Nrr1r^s r"is_int1z dtype.is_int1r)r$c |jdk(S)Nrr1r^s r"is_int8z dtype.is_int8r)r$c |jdk(S)Nrr1r^s r"is_int16zdtype.is_int16r$r$c |jdk(S)Nrr1r^s r"is_int32zdtype.is_int32r$r$c |jdk(S)Nrr1r^s r"is_int64zdtype.is_int64r$r$c |jdk(S)Nrr1r^s r"is_uint8zdtype.is_uint8r$r$c |jdk(S)Nrr1r^s r" is_uint16zdtype.is_uint16yyH$$r$c |jdk(S)Nrr1r^s r" is_uint32zdtype.is_uint32r>r$c |jdk(S)Nrr1r^s r" is_uint64zdtype.is_uint64r>r$c:|jtjvSr.)r2rrr^s r" is_floatingzdtype.is_floatingsyyENN**r$c:|jtjvSr.)r2rSTANDARD_FP_TYPESr^s r"is_standard_floatingzdtype.is_standard_floatingsyyE3333r$c:|jtjvSr.)r2rr r^s r" is_int_signedzdtype.is_int_signedyyE,,,,r$c:|jtjvSr.)r2rr r^s r"is_int_unsignedzdtype.is_int_unsignedrJr$c\|jtjtjzvSr.)r2rr r r^s r"is_intz dtype.is_ints"yyE,,u/?/????r$c"|jSr.)r1r^s r"is_boolz dtype.is_bool||~r$c|jrtjjS|j rtjj S|j sJtjjSr.)rPrrrrNrrDrr^s r"kindz dtype.kindsV <<>::%% % [[]::&& &##% %%::&& &r$c|jrd|jdz zdz S|jrd|jzdz SJ)NrrrIrrLr^s r"get_int_max_valuezdtype.get_int_max_valuesP    t((1,-1 1    !d'''!+ +ur$cn|jrd|jdz z S|jryJ)NrrrrUr^s r"get_int_min_valuezdtype.get_int_min_values;    ))A-.. .    !ur$c|tjtjztjztjzvSr.)rr r rr)type_strs r"is_dtypezdtype.is_dtypes35++e.>.>>ORWRcRccccr$ctd)NzNot implementedrr/r$r"is_voidz dtype.is_voids,--r$cyNFr/r/r$r"is_blockzdtype.is_blockr$cyr`r/r/r$r"is_ptrz dtype.is_ptrrbr$cyr`r/r/r$r"is_constzdtype.is_constrbr$cyr`r/r/r$r"is_tuplezdtype.is_tuplerbr$cVt|tsy|j|jk(Sr`)rWrr2rfs r"rz dtype.__eq__s"%'yyEJJ&&r$c.t|jfSr.)hashr2r^s r"__hash__zdtype.__hash__sTYYM""r$c|Sr.r/r^s r"rz dtype.scalar r$c|jjdr|j|jjvr%t d|d|jj|j|jj vrt |jd|jdk(r|jS|jdk(r|jS|jdvr|jS|jdvr|jS|jd vr|jS|jd vr|jS|jd k(r|jS|jd k(r|jS|jd k(r|jS|jdk(r|j!S|jdk(r|j#S|jdk(r|j%S|jdk(r|j'S|jdk(r|j)S|jdk(r|j+St d|d)Nrztype zB not supported in this architecture. The supported fp8 dtypes are zR is deprecated in this architecture and will be removed in a future triton releaserr)rr)rr)rr)rrrrrrrrrrrzfail to convert z to ir type)r2 startswithoptionssupported_fp8_dtypesrdeprecated_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_tyr0builders r"to_irz dtype.to_ir s3 99   &yy D DD 5/AAHAeAe@f"hiiyyGOOAAA {"tuv 99 &&( ( YY& &&( ( YY+ +&&( ( YY- -'') ) YY- -'') ) YY- -'') ) YY' !'') ) YY* $**, , YY) #))+ + YY) #))+ + YY* $**, , YY& &&( ( YY& &&( ( YY& '') ) YY& ((* *+D6=>>r$c|jSr.r1r^s r"__str__z dtype.__str__1 yyr$c|jjdrd|jddzS|jjdrd|jddzS|jS)Nfpfloatrbfbfloat)r2rpr^s r" codegen_namezdtype.codegen_name4sZ 99   %TYYqr]* * YY ! !$ 'diim+ +99 r$c|jS)z"See cache_key_part() in triton.cc.r1r^s r"cache_key_partzdtype.cache_key_part<syyr$c(d|jS)z4Output of repr needs to be an evaluatable expressionztriton.language.)rr^s r"r_zdtype.__repr__As!$"3"3"5!677r$c*t||||dzfSNr)r?rs r"rzdtype._unflatten_irEsgfot,fqj88r$N)rgr)r ir.builderrzir.typerr)9r9rRrSr r rrFrrrrrZrrrr r#r&r(r+r-r/r1r3r5r7r9r;r=r@rBrDrGrIrLrNrPrSrVrX staticmethodr[r^rardrfrhrrlpropertyrrrrrr_rr/r$r"rr<s4J@JfH8(KTt 4(l"&&'$'######$$$$%%%+4--@'dd..' #&?P89r$rcNeZdZd d dZd dZdZdZdZdZd dZ e dZ y ) pointer_typect|}t|ts"tdt |j d||_||_||_|s d|d|_ yd|d|_ y)Nzelement_ty has type `z`; expected `dtype`.zpointer<>zconst_pointer<) rrWrrGrYr9 element_ty address_spacerQr2)r0rrrQs r"rZzpointer_type.__init__Qsq)*5 *e,3D4D4M4M3NNbcd d$* 49hzl!, PZ|[\?] r$cl|j|jj||jSr.) get_ptr_tyrrrrs r"rzpointer_type.to_irZs*!!$//"7"7"@$BTBTUUr$c|jSr.r1r^s r"rzpointer_type.__str__]rr$c"|jSr.rr^s r"r_zpointer_type.__repr__`rQr$cyNTr/r^s r"rdzpointer_type.is_ptrcr$c|jSr.)rQr^s r"rfzpointer_type.is_constfrbr$ct|tsy|j|jk(xr4|j|jk(xr|j|jk(Sr`)rWrrrrQrfs r"rzpointer_type.__eq__isR%.%"2"22~t7I7IUM`M`7`~eieoeosxs~s~e~~r$c|Sr.r/r^s r"rzpointer_type.scalarnrnr$N)rF)rrrrrQr)rrrzir.pointer_type)rgrrr) r9rRrSrZrrr_rdrfrrrr/r$r"rrOs<^V r$rc eZdZdfd ZxZS)nv_tma_desc_typec@t|t||d|_y)N)rQrr)superrZrr2)r0rQr __class__s r"rZznv_tma_desc_type.__init__us e=I& r$)Tr)r9rRrSrZ __classcell__rs@r"rrss ''r$rcNeZdZd dZd dZdZdZdZd dZd dZ e dZ y ) block_typec||_t|ttfsJtt ||_|j s t dt|j |_d|j d|jd|_ y)Nz0d block_type is forbidden<, r) rrWr:tuple _unwrap_shapeshaperGrnumelr2)r0rrs r"rZzblock_type.__init__|st$54-010=/0 zz89 9)$**5  |2doo%6a8 r$cl|j|jj||jSr.) get_block_tyrrrrs r"rzblock_type.to_irs(##DOO$9$9'$BDJJOOr$c|jSr.r1r^s r"rzblock_type.__str__rr$c"|jSr.rr^s r"r_zblock_type.__repr__rQr$cyrr/r^s r"razblock_type.is_blockrr$c|jSr.rr^s r"get_block_shapeszblock_type.get_block_shapesrbr$ct|tsy|j|jk(xr|j|jk(Sr`)rWrrrrfs r"rzblock_type.__eq__s6%,%"2"22PtzzU[[7PPr$c|jSr.)rr^s r"rzblock_type.scalars r$N)rrrr )rrrz ir.block_type)rz List[int]rr) r9rRrSrZrrr_rarrrrr/r$r"rrzs; 9PQ r$rcDeZdZd dZdZdZd dZd dZdZdZ d d Z y) tuple_typeNc ||_|xsdgt|z|_ddjt |j|jDcgc] \}}|d|c}}zdz|_ycc}}w)Nr+[,:r\)typesr6fieldsjoinzipr2)r0rrkvs r"rZztuple_type.__init__sf 1s5z 1 #((3t{{DJJ;W#X41aqc1#J#XYY\__ #XsA3 c|jSr.r1r^s r"rztuple_type.__str__rr$c,t|jSr.)rFrr^s r"rztuple_type.__iter__rr$c^|jDcgc]}|j|c}Scc}wr.)rr)r0rtys r"rztuple_type.to_irs#,0JJ7b!777s*c |j|Sr.)r)r0indexs r" __getitem__ztuple_type.__getitem__szz%  r$cyrr/r^s r"rhztuple_type.is_tuplerr$ct|t|uxr4|j|jk(xr|j|jk(Sr.)rYrrrfs r"rztuple_type.__eq__s<DzT%[(fTZZ5;;-Ff4;;Z_ZfZfKffr$cg}|jD](}|j||\}}|j|*t|||fSr.)rrappendr)r0rrr;rrXs r"rztuple_type._unflatten_irsP** !B,,Wf=ME6 MM%  !VT"F**r$r.rr)rrrr)rrrrrzTuple[tuple, int]) r9rRrSrZrrrrrhrrr/r$r"rrs,`  8!g+r$rceZdZdZy) slice_typecd|_y)Nrr1r^s r"rZzslice_type.__init__s   r$Nr9rRrSrZr/r$r"rrs!r$rrrrrrrrrrrrrrrrrrrrc |dk(rtS|dk(r|rtS|dk(r|stS|dk(r|rtS|dk(r|stS|dk(r|rt S|dk(r|st S|dk(r|rtS|dk(r|stStd|d|)Nrrrr r zUnsupported bitwidth z and signedness ) rrrrrrrrrr)rsigneds r" get_int_dtypers1} Q6 Qv RF R RF R RF R 0 :J6(STTr$ceZdZdZd[fd Zd\dZd]dZed^dZed^dZ ed^dZ ed^d Z ed^d Z ed^d Z ed^d Zed^d Zed^dZed^dZed^dZed^dZed^dZed^dZed^dZed^dZed^dZed^dZed^dZed^dZed^dZed^dZed^dZed^dZed^dZ ed^dZ!ed^d Z"ed^d!Z#ed^d"Z$ed^d#Z%ed^d$Z&ed^d%Z'ed^d&Z(ed^d'Z)ed^d(Z*ed^d)Z+ed^d*Z,ed^d+Z-ed^d,Z.ed^d-Z/e0d.Z1ed_d`d/Z2dad0Z3dad1Z4dad2Z5dbd3Z6dad4Z7dad5Z8dad6Z9dcdad7Z:dddad8Z;dad9Zdfdad<Z?dfdad=Z@dfdad>ZAdfdad?ZBdfdad@ZCdfdadAZDdadBZEdadCZFdadDZGdadEZHdadFZIdadGZJdadHZKdgdadIZLdgdadJZMdadKZNdadLZOdadMZPdadNZQdgdadOZRdadPZSdhdadQZTdidadRZUdhdadSZVdidadTZWd_dadUZXdcdadVZYdjdadWZZdjdadXZ[de\fdkdYZ]d^dadZZ^xZ_S)lr?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. :) ct|||_|jr |jnd|_d|_|jD]}|xj |zc_t |j |_||_|j|_ t|jDcgc] }t |c}|_ycc}w)Not called by user code.r/rN) rrZhandlerarrrUrYrrr)r0rrYsrs r"rZztensor.__init__s  #'==?TZZ   A JJ!OJ tzz*  [[ $**=QIaL=> =s,C c:|j|jyr.rrrs r"rztensor._flatten_ir#t{{#r$c~t|jdzdjd|jDzdzS)Nrrc32K|]}t|ywr.)r.0rs r" z!tensor.__str__..(s0LAQ0Lsr\)rrrrr^s r"rztensor.__str__&s24::$tyy0L0L'LLsRRr$Nc t||d|SNT)sanitize_overflowraddr0rgrs r"rhztensor.__add__*4$JJr$c t||d|Srrrs r"rlztensor.__radd__.5$$JJr$c t||d|Srsubrs r"roztensor.__sub__2rr$c t||d|Srrrs r"rqztensor.__rsub__6rr$c t||d|Srmulrs r"rsztensor.__mul__:rr$c t||d|Srrrs r"rwztensor.__rmul__>rr$cFt|}tj|||Sr.rrtruedivrs r"ryztensor.__truediv__Bs!$U+eX66r$cFt|}tj|||Sr.rrs r"r{ztensor.__rtruediv__Gs!$U+tX66r$cFt|}tj|||Sr.rrfloordivrs r"r}ztensor.__floordiv__Ls!$U+  uh77r$cFt|}tj|||Sr.rrs r"rztensor.__rfloordiv__Qs!$U+  h77r$cFt|}tj|||Sr.rrmodrs r"ruztensor.__mod__V$U+||D%22r$cFt|}tj|||Sr.r rs r"__rmod__ztensor.__rmod__[$U+||E422r$c.tj||Sr.)rminusr0rs r"rztensor.__neg__as~~dH--r$c.tj||Sr.)rinvertrs r"rztensor.__invert__estX..r$cFt|}tj|||Sr.rrand_rs r"rztensor.__and__k$U+}}T5(33r$cFt|}tj|||Sr.rrs r"__rand__ztensor.__rand__p$U+}}UD(33r$cFt|}tj|||Sr.rror_rs r"rz tensor.__or__ur r$cFt|}tj|||Sr.rrs r"__ror__ztensor.__ror__zrr$cFt|}tj|||Sr.rrxor_rs r"rztensor.__xor__rr$cFt|}tj|||Sr.r!rs r"__rxor__ztensor.__rxor__rr$c^t||t|}tj|||Sr.rrrshlrs r"rztensor.__lshift__s)e$$U+||D%22r$c^t||t|}tj|||Sr.r&rs r" __rlshift__ztensor.__rlshift__s)t$$U+||E422r$ct||t|}|jjrt j |||St j |||Sr.rrrrIrashrlshrrs r"rztensor.__rshift__sLe$$U+ :: # # %==uh7 7==uh7 7r$ct||t|}|jjrt j |||St j |||Sr.r+rs r" __rrshift__ztensor.__rrshift__sLt$$U+ :: # # %==h7 7==h7 7r$c\tj||}tj|||Sr.rrO greater_thanrs r"rz tensor.__gt__s)""5(3$$T5(;;r$c\tj||}tj|||Sr.r1rs r"rztensor.__rgt__s)""5(3$$UD(;;r$c\tj||}tj|||Sr.rrO greater_equalrs r"rz tensor.__ge__s)""5(3%%dE8<B 78RXX=MRTRYRYRafhfmfmfu #=bT!BCC  D r$cJd)zTransposes a 2D tensor.z0Transposition must be created by the AST Visitorr/r^s r"rztensor.T s IHHur$c"t|||||S)z3 Alias for :py:func:`tensor.cast`. r)cast)r0rfp_downcast_roundingbitcastrs r"toz tensor.tos D%!5wRRr$cyr.r/r0rs r" broadcast_toztensor.broadcast_to r$cyr.r/r0dimss r"transz tensor.transrgr$cyr.r/ris r"permuteztensor.permute rgr$cyr.r/r^s r"rz tensor.split#rgr$cyr.r/res r"viewz tensor.view&rgr$cyr.r/res r"reshapeztensor.reshape)rgr$cyr.r/)r0axiss r"rUztensor.expand_dims,rgr$cyr.r/)r0rrarbs r"r`z tensor.cast/rgr$cyr.r/)r0rXmaskboundary_checkcache_modifiereviction_policys r"storez tensor.store2rgr$cyr.r/)r0offsetss r"advanceztensor.advance5rgr$cyr.r/)r0cmpvalsemscopes r" atomic_casztensor.atomic_cas8rgr$cyr.r/r0rrwrrs r" atomic_xchgztensor.atomic_xchg;rgr$cyr.r/rs r" atomic_addztensor.atomic_add>rgr$cyr.r/rs r" atomic_maxztensor.atomic_maxArgr$cyr.r/rs r" atomic_minztensor.atomic_minDrgr$cyr.r/rs r" atomic_andztensor.atomic_andGrgr$cyr.r/rs r" atomic_orztensor.atomic_orJrgr$cyr.r/rs r" atomic_xorztensor.atomic_xorMrgr$cyr.r/r^s r"expz tensor.expPrgr$cyr.r/r^s r"logz tensor.logSrgr$cyr.r/r^s r"cosz tensor.cosVrgr$cyr.r/r^s r"sinz tensor.sinYrgr$cyr.r/r^s r"sqrtz tensor.sqrt\rgr$cyr.r/r^s r"rsqrtz tensor.rsqrt_rgr$cyr.r/r^s r"absz tensor.absbrgr$cyr.r/)r0rt combine_fn keep_dimss r"reducez tensor.reduceergr$cyr.r/)r0rtrreverses r"associative_scanztensor.associative_scanhrgr$cyr.r/)r0indicesrts r"gatherz tensor.gatherkrgr$cyr.r/)r0num_binss r" histogramztensor.histogramnrgr$cyr.r/)r0divs r"cdivz tensor.cdivqrgr$cyr.r/r^s r"sigmoidztensor.sigmoidtrgr$cyr.r/)r0 ieee_roundings r"softmaxztensor.softmaxwrgr$cyr.r/r^s r"ravelz tensor.ravelzrgr$cyr.r/r0rtreturn_indicesreturn_indices_tie_break_leftrs r"maxz tensor.max}rgr$cyr.r/r0rttie_break_leftrs r"argmaxz tensor.argmaxrgr$cyr.r/rs r"minz tensor.minrgr$cyr.r/rs r"argminz tensor.argminrgr$cyr.r/)r0rtrrs r"sumz tensor.sumrgr$cyr.r/)r0rtrs r"xor_sumztensor.xor_sumrgr$cyr.r/r0rtrs r"cumsumz tensor.cumsumrgr$cyr.r/rs r"cumprodztensor.cumprodrgr$cyr.r/)r0r[ descendings r"sortz tensor.sortrgr$cyr.r/)r0r[s r"flipz tensor.fliprgr$)rYrrrr.NFNrrraz Optional[str]rbrrr?rztuple[tensor, tensor]r`)Nr/r+r+NNNNNF)NFTF)TF)rF)r[rUrrUrr?)`r9rRrSr8rZrrr(rhrlrorqrsrwryr{r}rrur rrrrrrrr$rr)rr/rrrrrrrrrrGrrLrrrrrrrcrfrkrmrrprrrUr`r{r~rrrrrrrrrrrrrrrrrrrrrrrrrrrrrrr CONSTEXPR_0rrrrs@r"r?r?sS&? $S K K K K K K K K K K K K 7 7 7 7 8 8 8 8 3 3 3 3  . . / /  4 4 4 4 3 3 3 3 4 4 4 4 3 3  3 3  8 8 8 8 < < < <  = = = =  9 9 9 9  : : : :  5 5 5 5 9 9 9 9 ; ; : :  - -   II S S                                         %)+   r$r?c^eZdZdddZddZdZddZdZdZdZ d Z d Z d Z d Z dd Zy)rNc|Dcgc]}|c}|_d}|xs)t|jDcgc] }|| c}|_ycc}wcc}w)Ncrt|trtSt|trtS|jSr.)rWrrrUrYrHs r"get_typez tuple.__init__..get_types*!U# !S!  66Mr$)r;rrY)r0rrYirrIs r"rZztuple.__init__sJ"&'Qq'  JJT[['I 'IJ ((Js A Ac,t|tr t|}t|tr|j|Sddl}t|t |j fsJt |j|j|j|jSNr) rWrrUr;rRrSrrVrWrX)r0idxrRs r"rztuple.__getitem__sq c3 C.C c9 %;;s# # cE8>>#:; ;;SYYsxx%@AB Br$cf|j|jjj|Sr.)r;rYrrrs r" __getattr__ztuple.__getattr__s&{{499++11$788r$c|t|tr t|}t|tsJ||j|<yr.)rWrrUr;)r0rrXs r" __setitem__ztuple.__setitem__s3 c3 C.C#y)))  Cr$c|t|tr t|}t|j|jzSr.)rWr:rr;rfs r"rhz tuple.__add__s- eT "%LET[[5<</00r$cjt|tsJt|j|jzSr.)rWrUrr;rXrfs r"rsz tuple.__mul__s*%+++T[[5;;.//r$cddl}t|t|jfr t|}t |j |j k(Sr)rRrWr:rrUr;)r0rgrRs r"rz tuple.__eq__s9 edHNN3 4%LE 455r$cTddl}t|j|jSr)rRrkrr;)r0rRs r"rlztuple.__hash__sNHNN4;;/00r$cdt|jDcgc] }t|c}Scc}wr.)rr;)r0rIs r"rz tuple.__str__s#DKK0qCF0110s-c,t|jSr.)rFr;r^s r"rztuple.__iter__sDKK  r$c,t|jSr.)r6r;r^s r"__len__z tuple.__len__s4;;r$cH|jD]}|j|yr.)r;r)r0rrs r"rztuple._flatten_irs! #A MM' " #r$r.)rr:rYr)rrU)rr)r9rRrSrZrrrrhrsrrlrrrrr/r$r"rrs@ KC9! 1 06 12! #r$rceZdZdZy)rScL||_||_||_t|_yr.)rVrWrXrrY)r0rVrWrXs r"rZzslice.__init__s    L r$Nrr/r$r"rSrSs!r$rSc<eZdZddZd dZd dZd dZd dZd dZy) tensor_descriptor_base_typec||_yr.r)r0rs r"rZz$tensor_descriptor_base_type.__init__s $r$cBt|||j}||dzfSr)$_experimental_tensor_descriptor_baser)r0rrrXs r"rz)tensor_descriptor_base_type._unflatten_irs%4WV_dooVfqj  r$cV|j|jj|Sr.)create_tensor_descriptor_typerrrs r"rz!tensor_descriptor_base_type.to_irs"44T__5J5J75STTr$c"d|jdS)Nztensor_descriptor)>w)Gk$J[J[JaJabiJjkkr$ct||xr4|j|jk(xr|j|jk(Sr.)rrrr)r0rgrs r"rztensor_descriptor_type.__eq__YsKw~e$k$//U=M=M*MkTXTeTeW\WiWiUj kr$)rrrrrrrr)r9rRrSrZrrrrrs@r"rrEs) lkkr$rc,eZdZdZdfd ZddZxZS)rz9A descriptor representing a tensor in global memory. c t|||t|t|Dcgc]}|jc}t|Dcgc]}|jc}|_||_||_ycc}wcc}w)r)rrN)rrZrrrYrr)r0rrrrrrs r"rZz(_experimental_tensor_descriptor.__init__bsi ,* !5"9a166"9:#W$=QVV$=>    #:$=s A6A;c|j|j|jd|jD|jd|jDy)Nc34K|]}|jywr.rrs r"rz>_experimental_tensor_descriptor._flatten_ir..qs4Aqxx4c34K|]}|jywr.r%rs r"rz>_experimental_tensor_descriptor._flatten_ir..rs6Aqxx6r&)rrextendrrrs r"rz+_experimental_tensor_descriptor._flatten_iros>t{{#444666r$)r List[tensor]rr)rrr)r9rRrSr8rZrrrs@r"rr^s 7r$rcTtj|d}|dk(xs |dk(xs|dk(S)N01trueon)osgetenv)var_namers r"get_bool_env_varr2us. (C A 8 /qF{ /a4i/r$c>t|tr |jS|Sr.r)rs r"rere}s!Yww Hr$cDt|}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 )rer program_idrtrs r"r5r5s! t $D   tX ..r$cDt|}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 )rer num_programsr6s r"r8r8s! t $D  x 00r$c\t|}t|}tj|||Sr.)rerarange)rVendrs r"r:r:s)  &E c "C ??5#x 00r$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 cTt|}|Dcgc] }t|c}Scc}wr.)re)rrs r"rrs&  &E,1 2q  " 22 2s%c2t|}t||Sr.)rrrs r"_shape_check_implr>s % E Lr$ctt|}t|}t|}tj||||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 )r>rerfull)rrXrrs r"r@r@s7 e $E  &E  &E ==uh 77r$c0tj|||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)inputrgrs r" broadcastrDs  ( (x @@r$r_cXtt|}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) )r>rJrbroadcast_impl_shaperCrrs r"rfrfs)" .u5 6E  ( (x @@r$cNt|}|sd}tj|||S)a Permutes the dimensions of a tensor. If the parameter :code:`dims` is not specified, the function defaults to a (1,0) permutation, effectively transposing 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. )rrrJrrmrCrrjs r"rkrks+, D !D    E4 22r$cFt|}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. rIrJs r"rmrms#( D !D   E4 22r$Fc2tj||||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 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). Current implementation of `cat` supports only can_reorder=True. )rcat)rCrg can_reorderrs r"rMrM3s <<uk8 <rJrrrrGs r"rprps4" NO .u5 6E   E5dH MMr$)rNrcZtt|}tj||||S)af Returns a tensor with the same number of elements as input but with the provided shape. :param input: The input tensor. :type input: Block :param shape: The new shape. :code:`shape` can be passed as a tuple or as individual parameters: :: # These are equivalent reshape(x, (32, 32)) reshape(x, 32, 32) )r>rJrrr)rCrNrrs r"rrrrs+" .u5 6E   E5+x @@r$c`| |cxkr|ksntd|d| d||dk\r|S||zS)Nz invalid axis z . Expected z <= axis < rr)rtndims r" _wrap_axisr^sF ET D =k4% D6RSS194-$+-r$ctj||}t|}t|tt fr t |n|g}t|jt|z}|Dcgc]}tt||}}tt|t|k7rtd||}t|D]}tj|||}|Scc}w)aR Expand the shape of a tensor, by inserting new length-1 dimensions. Axis indices are with respect to the resulting tensor, so ``result.shape[axis]`` will be 1 for each axis. :param input: The input tensor. :type input: tl.tensor :param axis: The indices to add new axes :type axis: int | Sequence[int] z7expand_dims received duplicate axes, normalized axes = )rrOrerWr rr:r6rr^setrsortedrU)rCrtraxesnew_ndimdrZrPs r"rUrUs   uh /E t $D#D8U*;<4:4&D5;;#d)+HBF GQJ*1-x 8 GD G 3t9~T"RSWRXYZZ C D\5""3845 J Hs+C"ctj||}t|}t|}t|}|rtj|||Stj||||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 )rrOrerbr`)rCrrarbrs r"r`r`sd$   uh /E  &E./CD!'*Guh77 ==x1E FFr$c  | |Jd|<|xrd|jjv}|r|s|dnd} tjd| }t |}t |}t |}t j |||||||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. :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`). z;Only one of input_precision and allow_tf32 can be specifiedtf32ieeeTRITON_F32_DEFAULT)rqallowed_dot_input_precisionsr/r0rerdot) rCrgaccinput_precision allow_tf32max_num_imprecise_acc out_dtyper supports_tf32default_precisions r"rkrks2  "j&8w:ww 8 \Vx/?/?/\/\%\ '4* HZFbh))$8:KL)/:O#I.I/0EF <<uc?`_ 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 )rerrOr ) pointerrwrgrxpadding_optionryrzvolatilers r"r r Bs\ t $D  &E !!$1 ""5(3(8N(8N)/:O"8,H ==$~~~_n!8 --r$cZtt||}tj|||S)zQ Reinterpret a generic pointer as a TMA-backed tensor descriptor object. )rrerreinterpret_tensor_descriptor)desc_ptrrrrblock_tys r"+_experimental_reinterpret_tensor_descriptorr~s+ -e4kBH  1 1(Hh OOr$cFt||||}|j||S)a Experimental feature to access TMA descriptors loads. This is an escape hatch to easily exercise TTGIR operations. This will be removed in the future and shouldn't be used in production code. This loads a tensor of data based on the descriptor and offsets. r_)rr ) desc_pointerr}rrrdescs r"_experimental_descriptor_loadrs( 7|UE\d eD 99Wx9 00r$cpt||j|j|}|j|||S)a Experimental feature to access TMA descriptors stores. This is an escape hatch to easily exercise TTGIR operations. This will be removed in the future and shouldn't be used in production code. This stores a tensor of data based on the descriptor and offsets. r_)rrrr{)rrXr}rrs r"_experimental_descriptor_storers5 7|U[[RWR]R]hp qD ::gux: 88r$c tj||}t|}|tj||}t|}t|}tj|||||||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"} )rrOrer{)r~rXrwrxryrzrs r"r{r{sjT   uh /E t $D !!$1(8N)/:O >>'5$P_ai jjr$c 8tj|||||||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)baserrr}rorderrs r"rrs"  " "4+uV^ __r$c0tj|||S)z Advance a block pointer :param base: the block pointer to advance :param offsets: the offsets to advance, a tuple by dimension )rr~)rr}rs r"r~r~s   D'8 44r$c4tj|||||S)aMake an experimental 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._experimental_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)rrrrrs r"$_experimental_make_tensor_descriptorrsn  * *4h WWr$cdfd }|S)Nc: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 )r8)funcdocstrhas_cmpr2s r" _decoratorz&_add_atomic_docstr.._decorator/sLv69  : :F     r$rrrrr/)r2rrs`` r"_add_atomic_docstrr-s4 r$zcompare-and-swapT)rctj||}tj||}t|}t|}tj||||||Sr.)rrOrer)r~rrrrrs r"rrLsW   S( +C   S( +C c "C  &E   wS#uh GGr$exchangectj||}t|}t|}t|}tj||||||Sr.)rrOrerr~rrwrrrs r"rrWsO   S( +C c "C  &E t $D   dC IIr$rctj||}t|}t|}t|}tj||||||Sr.)rrOrerrs r"rrbO   S( +C c "C  &E t $D   wT3x HHr$rctj||}t|}t|}t|}tj||||||Sr.)rrOrerrs r"rrmrr$rctj||}t|}t|}t|}tj||||||Sr.)rrOrerrs r"rrxrr$z logical andctj||}t|}t|}t|}tj||||||Sr.)rrOrerrs r"rrrr$z logical orctj||}t|}t|}t|}tj||||||Sr.)rrOrerrs r"rrsO   S( +C c "C  &E t $D   gsD#uh GGr$z logical xorctj||}t|}t|}t|}tj||||||Sr.)rrOrerrs r"rrrr$ctj||}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. )rrOrwhere) conditionrIyrs r"rrs?"""9h7IQAQA >>)Q8 44r$c^t|}t|}tj||||Sr.)rrrrIrrrs r"rr,QAQA <<1/ ::r$c^t|}t|}tj||||Sr.)rrrrs r"rrrr$c^t|}t|}tj||||Sr.)rrrrs r"rrrr$ctj||}tj||}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_)rrO_promote_bfloat16_to_float32reminimumrIr propagate_nanrs r"rra 1h'A1h'A$Q:A$Q:A' 6M   Aq- ::r$ctj||}tj||}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_)rrOrremaximumrs r"rrrr$ctj||}tj||}tj||}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_)rrOrreclamp)rIrrrrs r"rr s$ 1h'A   S( +C   S( +C$Q:A &sX >C &sX >C' 6M >>!S#}h ??r$c dfd }|S)Ncd}|ddddz }|ddddz }|ddd d z }|j |_|S) Na~ 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 reduction should be done. If None, reduce all dimensions :type axis: int :param keep_dims: if true, keep the reduced dimensions with length 1 :type keep_dims: boolz :param z-: if true, return index corresponding to the z value :type z: boolzC: if true, in case of a tie (i.e., multiple elements have the same zI value), return the left-most index for values that aren't NaN :type zc: the desired data type of the returned tensor. If specified, the input tensor is casted to :code:`z` before the operation is performed. This is useful for preventing data overflows. If not specified, integer and bool dtypes are upcasted to :code:`tl.int32` and float dtypes are upcasted to at least :code:`tl.float32`. :type z : tl.dtyper1formatr8)rr dtype_argr2return_indices_arg tie_break_args r"rz)_add_reduction_docstr.._decorator, s  )   LTFS  f( (F  $  ?]^b]cd /# #F   ;y{DzEE +Z# #F}}$}/  r$rr/)r2rrrrs```` r"_add_reduction_docstrr) s4 r$c#VK|j}d|j|ywr.)get_insertion_pointrestore_insertion_point)rips r"_insertion_guardrI s%  $ $ &B  ##B's')czttrtf|dSfd}fdtt|}"t t dj tj|}|r.tfd|D}|Stfd|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 rrr*rc JDcgc]}|jjc}dz}|jd}t 5 fd} j |t t ||}t|Dcgc] \}}t|j||"}}} j|i} t| tr| jg} n| D cgc]} | j} }  j| dddycc}wcc}}wcc} w#1swYyxYw)Nrrc&|jSr.rrrs r"z5reduce..make_combine_region..f aggh/r$r )rYr get_regionrcreate_block_with_parentr:maprTr?argcall_JitFunctionrWrcreate_reduce_ret) reduce_opt param_typesregionrblockrrrresultsrrrr*rrCs r"make_combine_regionz#reduce..make_combine_regionb s.34qvv}}4q8 %%a( h ' 1/E55fd3ukCZ>[\E:CK:PQBF599Q<,QDQ 11*d21NG'6*">>*-4518855 &H & & 0 1 15 R 6 1 1/D 9D>%D#9DD/D DD"cXtj|D]}t|d}|S)Nrr_)rRrangerU)rndims_rs r" expand_ndimszreduce..expand_ndimsp s.& 5AAq84A 5r$c3:K|]}t|yw)r_N)rU)rrrrts r"rzreduce..| sMA Ath??Msc3\K|]#}|tdj%yw)rN)r6r)rrrrCs r"rzreduce..~ s%J QE!HNN(;<Js),) rWr?rrer^r6rr reductionr) rCrtrrrr*rrZrs ``` `` @r"rrP s% uizYQYfpqrstt 1 t $D#I.I $E!HNN 34   UD*=x HC  MMMC JJcJJC Jr$cp|jj}|tur|jt|S|S)Nr_)rYrbfloat16rcrt)rr scalar_tys r"rr s1 IHttGht// Hr$ct|}|j|}td||}t|jdkDrgt j t|jDcgc] }t |} }| |=t|| |}t||j|}t||f|||||\} } | | fScc}w)Nrr_rr) rerr:r6rRrrUrUrfr) rCrtrrrr*nrrdaxes_to_expandrvaluerindicess r"_reduce_with_indicesr s t $D DA 1a( +E 5;;!08s5;;?O0PQ1)A,QQ 4 E>HEUEKK(CuendJ)^f)35FH 8 Rs+Ccdfd }|S)Nc:d}|j|_|S)Nz 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: intr1r)rrr2s r"rz$_add_scan_docstr.._decorator s#}}$}/  r$rr/)r2rs` r"_add_scan_docstrr s  r$cttrtf||dSfd}t|}|"t |t dj }tj|||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 rUrc JDcgc]}|jjc}dz}|jd}t 5 fd} j |t t ||}t|Dcgc] \}}t|j||"}}} j|i} t| tr| jg} n| D cgc]} | j} }  j| dddycc}wcc}}wcc} w#1swYyxYw)Nrrc&|jSr.rrs r"rz?associative_scan..make_combine_region.. rr$r)rYrrrrr:rrTr?rrrWrcreate_scan_ret)scan_oprrrrrrrrrrrrr*rrCs r"rz-associative_scan..make_combine_region s.34qvv}}4q8 ##A& h ' //E55fd3ukCZ>[\E:CK:PQBF599Q<,QDQ 11*d21NG'6*">>*-4518855 $H $ $g . / /5 R 6 / /r)rWr?rrer^r6rr)rCrtrrrr*rs` ` `` r"rr sx%  4Wxdnopqrr / t $D $E!HNN 34  $ $UD2EwPX YYr$cFt|}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 :type input: Tensor :param num_bins: number of histogram bins :type num_bins: int )rerr)rCrrr*s r"rr s##8,H   eXx 88r$cHt|}tj||||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 )rerr)srcrrtrs r"rr s# t $D ??3tX 66r$c,tj|S)zA Insert a barrier to synchronize all threads in a block. )r debug_barrierr_s r"rr s  ! !( ++r$c rt|tr|g}t|D]c\}}t|tstd|dt|jt r@td|dt |jd|Dcgc]}|j}}tj||Scc}w)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[r\) rWrUrTrGrXrrYr multiple_ofrCr;rrrdrIs r"rr s &)$&!t1!Y'oaS0KLM M!''3'oaS0abfghgngnboappqrs s t & &!agg &F &   v ..' B4c rt|tr|g}t|D]c\}}t|tstd|dt|jt r@td|dt |jd|Dcgc]}|j}}tj||Scc}w)z^ Let the compiler know that the `value` first values in :code:`input` are contiguous. rrrr\) rWrUrTrGrXrrYrmax_contiguousrs r"rr s &)$&!t1!Y'oaS0KLM M!''3'oaS0abfghgngnboappqrs s t & &!agg &F &  " "5& 11'rc rt|tr|g}t|D]c\}}t|tstd|dt|jt r@td|dt |jd|Dcgc]}|j}}tj||Scc}w)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]. rrrr\) rWrUrTrGrXrrYr max_constancyrs r"rr& s&)$&!t1!Y'oaS0KLM M!''3'oaS0abfghgngnboappqrs s t & &!agg &F &  ! !% 00'rcVtjtj|||S)z< Allow compiler to assume the :code:`cond` is True. )rassumerO)condrs r"r r 9 s" ??8--dH=x HHr$  )sepr;fileflushrcy)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/)rr;rrrr;s r" static_printrF s r$cy)z Assert the condition at compile time. Does not require that the :code:`TRITON_DEBUG` environment variable is set. .. highlight:: python .. code-block:: python tl.static_assert(BLOCK_SIZE == 1024) Nr/r msgrs r" static_assertrV s r$)hexrc.ddl}t|}t|ts J|dd}|D]}||jvsd}n|s J|dg}|D]'}|j t j||)t j||||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) stringrerWr printablerrrO device_print) prefixrrrrb_asciichnew_argsrs r"rrd sD  (F fc "=vhn$== "G V%% %G  6vh5667H;**39:;  3 AAr$cnt|}tjtj||||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. )rer device_assertrOrs r"r!r! s0& c "C  ! !("4"4T8"Dc8 TTr$c |t|}t|}t|}t|} t|d}tjt t |}|}|D cgc]} tj| |c} x} rttj|ddd} | d} | D]} | | | \}} | jrEt| D]\}} | | | \| |<}|Dcgc]}t|| j}}| Dcgc]}|j}}|j||||Dcgc]}|j!|c}|||st#j%d|dSt'fdt|DS#t$r d}|f}YjwxYwcc} wcc}wcc}wcc}w)ac 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 :param _builder: the builder :return: one tensor or a tuple of tensors of the given dtypes TF)rarithmetic_check allow_lhs_ptr allow_rhs_ptrrc3ZK|]"\}}tj||$ywr.)r? get_result)rrrcalls r"rz)inline_asm_elementwise.., s%PEAr*B/Ps(+)rerFrGrVr`r _DtypeClassrrOrbinary_op_type_checking_implrrTrrcreate_inline_asmrr?r'r)asm constraintsrris_purepackrhas_multiple_outputsres_tysr dispatch_argsbin_op_type_checking broadcast_argitemrrdtrrrr(s @r"inline_asm_elementwiser7 st c "C%k2K t $D!'*G U # KK-u 5EGFJKs++C:KK}K&  1 1"  &a( ! ID3D-H A} I   $]3 P4&:4&O# a ! PEJKrz"m&9&9:KGK!./Aqxx/G/  % %c;_fAgY["((8BTAgiprv wD dooa('!*55 PYw=OP PP; $ L L/Ags) F#F*-F/F47F9F'&F'c$eZdZdZddZdZdZy) 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|ts J|d|td|_n!t|ts J|d||_|td|_||_yt|ts J|d||_||_y)Nz7 used as tl.static_range start value is not a constexprrz6 used as tl.static_range step value is not a constexprrz5 used as tl.static_range end value is not a constexpr)rWrUrXrVr;)r0arg1arg2rXs r"rZzstatic_range.__init__F s$ *ltf4k,ll* <!! DIdI. o4&8n0o o.DI <"1DJDHdI. n4&8m0n n.DJDHr$ctdNz8static_range can only be used in @triton.jit'd functionsr]r^s r"rzstatic_range.__iter__U UVVr$ctdr>r]r^s r"__next__zstatic_range.__next__X r?r$rr9rRrSr8rZrrAr/r$r"r9r94 s" WWr$r9c(eZdZdZ ddZdZdZy)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. Nc|td|_n||_|td|_||_n||_||_||_||_||_||_y)Nrr)rUrXrVr; num_stagesloop_unroll_factordisallow_acc_multi_bufferflatten)r0r;r<rXrErFrGrHs r"rZzrange.__init__} s\ <!! DIDI <"1DJDHDJDH$"4)B& r$ctdNz4tl.range can only be used in @triton.jit'd functionsr]r^s r"rzrange.__iter__ QRRr$ctdrJr]r^s r"rAzrange.__next__ rKr$)NNNNFFrBr/r$r"rr\ s#@X\:?"SSr$rc t|dk(r tdtt|jd}t||k7rtdt|d|g} g} |D]t} t | t r7| j | j| j | jJ| j t| | j | vt| } | |vrtd|jd| || d} || d} |r t| |} t |||| | | j||| 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) r6rr:r7rWr?rrrrYrrr)rlib_namelib_pathrarg_type_symbol_dict ret_shaper.rnum_args arg_typesarg_listrsymbolret_types r"dispatchrW sp  A%8994,1134Q78H 4yH##&t9+VH:?@ @IH! c6 "   SYY ' OOCJJ '   T#Y ' OOC !i I,,**>*C*C*E)FfYKYZ Z&i03' 215 !(I6Hd8XvxPXAY[bcemnnr$c |j}d}d}g} tjt|D]\} t j || ||| <| j || j|| jjs[d}^t| dkDrt| } d} | |vrd} |d} |D]} t j| | || \}} tjt|D]$} t j|| | || \|| <}&|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#)copyrRrr6rrOrrrYrarr*rcreate_extern_elementwiserW)rNrOrrPr.rr2 all_scalarrQrSrr#r4r5rrs r"extern_elementwiser\ syIIKMJII ^^C . /#--mA.>I aq)//0   ) ) +J   9~)$  , ,$ %a( ! hD'DDT=ZbVf h A} hM 23 kA"*"G"G VWHXZgiqYi#k M! a k%++I  - -D D(Hm=QS\^ego ppr$c0tj|||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*)rvryrs r"binary_op_type_legalizationr^ s  0 0c7 CCr$ct|S)z#A decorator for external functions.)r(rMs r"externr` s 2;r$)r!rrrrr.)rrrrrr)rCr?)FNrrrr)NNr/r+r+r+FN)rr)Nr/r+r+N)rr?) rr?rr)rr)rzList[constexpr]rrr)r2rrrrCallable[[T], T]r)NNNN)TN)rrU)rrU) r2rrrrrrrrra)FNN)r2rrra)rrr;r)r+N) r,rr-rrr rzUnion[dtype, Sequence[dtype]]r.rr/r) rNrrOrrr:rPdictrQrr.r) rNrrOrrr:rPrbr.r) __future__rwarningsr contextlibrenumr functoolsrrrVr r r r r rrrR runtime.jitrr4r/ _C.libtritonrr+r_utilsrrrr' PROPAGATE_NAN PropagateNanr(rDrJr>rOrQrUrrrrrrr)rrrrrrrrrrrrrrrfloat8e5 float8e5b16 float8e4nv float8e4b8 float8e4b15float16rrtfloat64pi32_trr?rrSrrrrr2rer5r8r:r8rr>r@rDrfrkrmrMrrSrrprrr^rUr`rkrur rrrr{rr~rrrrrrrrrrrrrrNONErrrrrrrrrrrrrrrrr rrrr!r7r9rrWr\r^r`r/r$r"rvs "%$ LLL A CL%   )X ,.  + +  C)C)Ll 6 " " " "*J9IJ9` !5!H'|'''T++B!! V} V} V} g g g g x x x >J 9  9  J - = - - e U8Z ZZ z A#JA#H!!#)#2@W:@WFk8k27&J7.0  / /$ 1 1  1 1 &'>%>? 3   8 8,  A  A)-A A&)-3 34#'3 3,  =  =  ) ).   B!%N N(',tA A&. 8G G< ]amt!h !hH KO[`ls )  )P rt"&8- 8-v 9=PBfP P 1 1 9 9.k .kb  `  `5 5  6X 6X 6X6X! 6X % 6X 6X|>&5H6 HJJ  JEI IEI IEI IM"I# IL!H" HM"I# I 5 56 ; ;  ; ;  ; ;  -9->->; ;* -9->->; ;* 2>2C2Cd@ @D[_+/%(4D@(( - -`      ,  !Z !ZH 9  9 7  7( , ,  /  /  2  2 1 1$ I I %(TE\`            $)D-B -B` U U, >BQ$(Q03Q QN%W%WP6S6S~&*(o(oV $'q 'qTDr$