a zdP@sddlmZddlmZddlmZmZmZddlZddl m Z m Z ddl m Z edZd d ZGd d d ZGd ddeZGdddeZGdddeZedZedZedZedZedZedZedZedZedZedZedZedZedZ ed Z!ed!Z"eeZ#Gd"d#d#Z$Gd$d%d%Z%d&d'Z&e dd(d)Z'e dd*d+Z(e dd,d-Z)d.d/Z*e dd0d1Z+e dd2d3Z,e dd4d5Z-e dd6d7Z.e dd9d:Z/e dd;d<Z0e dd=d>Z1e dd@dAZ2e ddCdDZ3e ddEdFZ4dGdHdIdJdKZ5e e5dLddMdNZ6e e5dOddPdQZ7e e5dRddSdTZ8e e5dUddVdWZ9e e5dXddYdZZ:e e5d[dd\d]Z;e e5d^dd_d`Ze ddfdgZ?e ddhdiZ@dGdHdIdjdkZAe eAdlddmdnZBe eAdoddpdqZCe eAdrddsdtZDe eAduddvdwZEe eAdxddydzZFdGdHdId{d|ZGe eGd}dd~dUZHe eGddddZIe eGddddXZJe eGddddZKe eGddddZLe eGddddZMe dddZNe dddZOe dddZPejQddZRejQddZSejQddZTejQdd}ZUejQeAdddZVejQeAddddZWejQddZXejQddZYejQddZZejQddZ[e ddddZ\GdddZ]dS)) annotations)Enum)CallableListTypeVarN)builtinsemantic)irTcCs<t|trt||tSt|trd|kr8dkrLnnt||tSd|kr`dkrtnnt||tSd|krdkrnnt| |t Sd|krdkrnnt| |t St d|dnBt|t rt||tSt|trt|j|St|tr|Sd s8Jd |d t|d dS) NilllllzNonrepresentable integer .Fzcannot convert z of type z to tensor) isinstancebooltensorZget_int1int1intZ get_int32int32uint32Z get_int64int64uint64 RuntimeErrorfloatZget_fp32float32 constexpr _to_tensorvaluetype)xbuilderr]/var/www/html/stable-diffusion-webui/venv/lib/python3.9/site-packages/triton/language/core.pyr s&      rc@s~eZdZgdZgdZgdZdgZgdZdgZGddde Z d d Z d d Z d dZ ddZddZddZddZddZddZddZddZdd Zd!d"Zd#d$Zd%d&Zd'd(Zd)d*Zd+d,Zd-d.Zd/d0Zd1d2Zd3d4Z e!d5d6Z"e!d7d8Z#e!d9d:Z$dd;dd?Z&d@dAZ'e(dBdCZ)dDdEdFdGdHZ*dIdJZ+e(dKdLdMdNZ,dOdPZ-dQS)Rdtype)rint8int16rr)uint8uint16rr)fp8fp16bf16fp32fp64r&)r'r(r)r*voidc@seZdZdZdZdS)zdtype.SIGNEDNESSrrN)__name__ __module__ __qualname__SIGNEDUNSIGNEDrrrr SIGNEDNESS-sr1cCs||_|tjtjtjtjvs*J||tjvr\tjj|_t | dd|_ |j |_ n|tjvrtjj |_t | dd|_ |j |_ n|tjvr |dkrd|_d|_ nX|dkrd|_d|_ nB|d krd |_d|_ n,|d krd |_d |_ n|dkrd|_d|_ n|dkrd|_ dS)Nrr&r' r(r) r*5@r+r)namer! SINT_TYPES UINT_TYPESFP_TYPES OTHER_TYPESr1r/Zint_signednessrsplitZ int_bitwidthZprimitive_bitwidthr0Zfp_mantissa_width)selfr<rrr __init__1s8$         zdtype.__init__cCs |jdkS)Nr&r<rBrrr is_fp8Osz dtype.is_fp8cCs |jdkS)Nr'rDrErrr is_fp16Rsz dtype.is_fp16cCs |jdkS)Nr(rDrErrr is_bf16Usz dtype.is_bf16cCs |jdkS)Nr)rDrErrr is_fp32Xsz dtype.is_fp32cCs |jdkS)Nr*rDrErrr is_fp64[sz dtype.is_fp64cCs |jdkS)NrrDrErrr is_int1^sz dtype.is_int1cCs |jdkS)Nr"rDrErrr is_int8asz dtype.is_int8cCs |jdkS)Nr#rDrErrr is_int16dszdtype.is_int16cCs |jdkS)NrrDrErrr is_int32gszdtype.is_int32cCs |jdkS)NrrDrErrr is_int64jszdtype.is_int64cCs |jdkS)Nr$rDrErrr is_uint8mszdtype.is_uint8cCs |jdkS)Nr%rDrErrr is_uint16pszdtype.is_uint16cCs |jdkS)NrrDrErrr is_uint32sszdtype.is_uint32cCs |jdkS)NrrDrErrr is_uint64vszdtype.is_uint64cCs |jtjvSN)r<r!r?rErrr is_floatingyszdtype.is_floatingcCs |jtjvSrT)r<r!CUSTOMIZED_FP_TYPESrErrr is_customized_floating|szdtype.is_customized_floatingcCs |jtjvSrT)r<r!STANDARD_FP_TYPESrErrr is_standard_floatingszdtype.is_standard_floatingcCs |jtjvSrT)r<r!r=rErrr is_int_signedszdtype.is_int_signedcCs |jtjvSrT)r<r!r>rErrr is_int_unsignedszdtype.is_int_unsignedcCs|jtjtjvSrT)r<r!r=r>rErrr is_intsz dtype.is_intcCs|SrT)rKrErrr is_boolsz dtype.is_boolcCs tddS)NzNot implementedrrrrr is_voidsz dtype.is_voidcCsdSNFrrrrr is_blockszdtype.is_blockcCsdSr`rrrrr is_ptrsz dtype.is_ptr)othercCst|tsdS|j|jkSr`)r r!r<rBrcrrr __eq__s z dtype.__eq__cCs || SrTrerdrrr __ne__sz dtype.__ne__cCs t|jfSrT)hashr<rErrr __hash__szdtype.__hash__cCs|SrTrrErrr scalarsz dtype.scalar ir.builderzir.typerreturncCs|jdkr|S|jdkr$|S|jdvr6|S|jdvrH|S|jdvrZ|S|jdvrl|S|jdkr~|S|jdkr|S|jd kr| S|jd kr| S|jd kr| St d |d dS)Nr+r)r"r$)r#r%)rr)rrr&r'r(r)r*zfail to convert z to ir type) r<Z get_void_tyZ get_int1_tyZ get_int8_tyZ get_int16_tyZ get_int32_tyZ get_int64_tyZ get_fp8_tyZ get_half_tyZ get_bf16_tyZ get_float_tyZ get_double_ty ValueErrorrBrrrr to_irs.           z dtype.to_ircCs|jSrTrDrErrr __str__sz dtype.__str__strrmcCs|jS)z"See cache_key_part() in triton.cc.rDrErrr cache_key_partszdtype.cache_key_partcCs d|jS)Nztriton.language.rDrErrr __repr__szdtype.__repr__N).r,r-r.r=r>r?rVrXr@rr1rCrFrGrHrIrJrKrLrMrNrOrPrQrRrSrUrWrYrZr[r\r] staticmethodr_rarbrergripropertyrjrprqrtrurrrr r!%sZ    r!c@sreZdZddddddZddd d d Zd d ZddZddZdddddZdddddZ e ddZ dS) pointer_typerr!r) element_ty address_spacecCs,t|tstd||_||_||_dS)Nz,element_ty is a {type(element_ty).__name__}.)r r! TypeErrorryrzrqr<)rBryrzrrr rCs  zpointer_type.__init__rkzir.pointer_typerlcCs||j|dSNr)Z get_ptr_tyryrprorrr rpszpointer_type.to_ircCsd|jdS)Nzpointer<>ryrErrr rqszpointer_type.__str__cCs|SrTrqrErrr ruszpointer_type.__repr__cCsdSNTrrErrr rbszpointer_type.is_ptrrrcrmcCs&t|tsdS|j|jko$|j|jkSr`)r rxryrzrdrrr res zpointer_type.__eq__cCs || SrTrfrdrrr rgszpointer_type.__ne__cCs|SrTrrErrr rjszpointer_type.scalarN)r) r,r-r.rCrprqrurbrergrwrjrrrr rxsrxc@s~eZdZdddddZdddd d Zd d Zd dZddZddddZdddddZ dddddZ e ddZ dS) block_typer!r)ryshapecCsb||_|stdt|dtr.dd|D}||_d|_|jD]}|j|9_q@||_dS)Nz0d block_type is forbiddenrcSsg|] }|jqSrr.0srrr z'block_type.__init__..r)ryr{r rrnumelrqr<)rBryrrrrr rCs zblock_type.__init__rkz ir.block_typerlcCs||j||jSrT)Z get_block_tyryrprrorrr rpszblock_type.to_ircCsd|jd|jdS)N )rrrErrr rq%szfunction_type.__str__rkrcs4fdd|jD}fdd|jD}||S)Ncsg|]}|qSrrp)rtyrrr r)rz'function_type.to_ir..csg|]}|qSrr)rZret_typerrr r*r)rrZget_function_ty)rBrZir_param_typesrrrr rp(szfunction_type.to_irN)r,r-r.rCrqrprrrr r srr+rr"r#rrr$r%rrr&r'r(r)r*c@s6eZdZdZddZddddZdd Zd d Zd d ZddZ ddZ ddZ ddZ ddZ ddZddZddZddZd d!Zd"d#Zd$d%Zd&d'Zd(d)Zd*d+Zd,d-Zd.d/Zd0d1Zd2d3Zd4d5Zd6d7Zd8d9Zd:d;Zdd?Z!d@dAZ"dBdCZ#dDdEZ$dFdGZ%dHdIZ&dJdKZ'dLS)MrzL This class is used to store a value that is known at compile-time. cCst|tr|j|_n||_dSrTr rr)rBrrrr rCKs  zconstexpr.__init__rrrscCsd|jdS)Nz constexpr[]rrErrr ruQszconstexpr.__repr__cCst|j|jSrTrrrdrrr __add__Tszconstexpr.__add__cCst|j|jSrTrrdrrr __radd__Wszconstexpr.__radd__cCst|j|jSrTrrdrrr __sub__Zszconstexpr.__sub__cCst|j|jSrTrrdrrr __rsub__]szconstexpr.__rsub__cCst|j|jSrTrrdrrr __mul__`szconstexpr.__mul__cCst|j|jSrTrrdrrr __mod__cszconstexpr.__mod__cCst|j|jSrTrrdrrr __rmul__fszconstexpr.__rmul__cCst|j|jSrTrrdrrr __truediv__iszconstexpr.__truediv__cCst|j|jSrTrrdrrr __rtruediv__lszconstexpr.__rtruediv__cCst|j|jSrTrrdrrr __floordiv__oszconstexpr.__floordiv__cCst|j|jSrTrrdrrr __rfloordiv__rszconstexpr.__rfloordiv__cCst|j|jkSrTrrdrrr __gt__uszconstexpr.__gt__cCst|j|jkSrTrrdrrr __rgt__xszconstexpr.__rgt__cCst|j|jkSrTrrdrrr __ge__{szconstexpr.__ge__cCst|j|jkSrTrrdrrr __rge__~szconstexpr.__rge__cCst|j|jkSrTrrdrrr __lt__szconstexpr.__lt__cCst|j|jkSrTrrdrrr __rlt__szconstexpr.__rlt__cCst|j|jkSrTrrdrrr __le__szconstexpr.__le__cCst|j|jkSrTrrdrrr __rle__szconstexpr.__rle__cCst|j|jkSrTrrdrrr reszconstexpr.__eq__cCst|j|jkSrTrrdrrr rgszconstexpr.__ne__cCs t|jSrT)rrrErrr __bool__szconstexpr.__bool__cCs t|j SrTrrErrr __neg__szconstexpr.__neg__cCst|j|j@SrTrrdrrr __and__szconstexpr.__and__cCst|jo |jSrTrrdrrr logical_andszconstexpr.logical_andcCst|j|jBSrTrrdrrr __or__szconstexpr.__or__cCst|jp |jSrTrrdrrr logical_orszconstexpr.logical_orcCs t|j SrTrrErrr __pos__szconstexpr.__pos__cCs t|jSrTrrErrr __invert__szconstexpr.__invert__cCst|j|jSrTrrdrrr __pow__szconstexpr.__pow__cCst|j|j?SrTrrdrrr __rshift__szconstexpr.__rshift__cCst|j|j>SrTrrdrrr __lshift__szconstexpr.__lshift__cCs t|j SrTrrErrr __not__szconstexpr.__not__cOs|j|i|SrTr)rBargskwdsrrr __call__szconstexpr.__call__N)(r,r-r.__doc__rCrurrrrrrrrrrrrrrrrrrrrergrrrrrrrrrrrrrrrrr rFsJrc@seZdZddddZddddZedQd d ZdRd d ZedSddZdTddZ edUddZ dVddZ edWddZ dXddZ edYddZedZddZed[ddZed\d d!Zed]d"d#Zed^d$d%Zed_d&d'Zed`d(d)Zedad*d+Zedbd,d-Zedcd.d/Zeddd0d1Zeded2d3Zedfd4d5Zedgd6d7Zedhd8d9Zedid:d;Zedjdd?Z edld@dAZ!edmdBdCZ"edndDdEZ#edodFdGZ$edpdHdIZ%edqdJdKZ&e'dLdMZ(edrdOdPZ)d S)srr!)rcCsl||_d|_|r|j|_d|_|jD]}|j|9_q(t|j|_||_|j|_dd|jD|_dS)N)rrcSsg|] }t|qSr)rrrrr rrz#tensor.__init__..)handlerrarrrrjr!)rBrrrrrr rCs  ztensor.__init__rrrscCs(t|jdddd|jDdS)N[,css|]}t|VqdSrT)rrrrrr rz!tensor.__str__..r)rrr!joinrrErrr rqsztensor.__str__NcCst||}t|||SrT)rr addrBrc_builderrrr rs ztensor.__add__cCs|j||dSNr)rrrrr rsztensor.__radd__cCst||}t|||SrTrr subrrrr rs ztensor.__sub__cCst||}t|||SrTrrrrr rs ztensor.__rsub__cCst||}t|||SrT)rr mulrrrr rs ztensor.__mul__cCs|j||dSr)rrrrr rsztensor.__rmul__cCst||}t|||SrTrr truedivrrrr rs ztensor.__truediv__cCst||}t|||SrTrrrrr rs ztensor.__rtruediv__cCst||}t|||SrTrr floordivrrrr rs ztensor.__floordiv__cCst||}t|||SrTrrrrr rs ztensor.__rfloordiv__cCst||}t|||SrTrr modrrrr rs ztensor.__mod__cCst||}t|||SrTrrrrr __rmod__s ztensor.__rmod__cCs t||SrT)r minusrBrrrr rsztensor.__neg__cCs t||SrT)r invertrrrr r sztensor.__invert__cCst||}t|||SrT)rr and_rrrr rs ztensor.__and__cCst||}t|||SrT)rr or_rrrr rs z tensor.__or__cCst||}t|||SrT)rr Zxor_rrrr __xor__s ztensor.__xor__cCst||}t|||SrT)rr Zshlrrrr r s ztensor.__lshift__cCs4t||}|jr"t|||St|||SdSrT)rr!rZr ZashrZlshrrrrr r%s  ztensor.__rshift__cCst||}t|||SrTrr greater_thanrrrr r0s z tensor.__gt__cCst||}t|||SrTrrrrr r5s ztensor.__rgt__cCst||}t|||SrTrr Z greater_equalrrrr r;s z tensor.__ge__cCst||}t|||SrTrrrrr r@s ztensor.__rge__cCst||}t|||SrTrr less_thanrrrr rFs z tensor.__lt__cCst||}t|||SrTrrrrr rKs ztensor.__rlt__cCst||}t|||SrTrr Z less_equalrrrr rQs z tensor.__le__cCst||}t|||SrTrrrrr rVs ztensor.__rle__cCst||}t|||SrT)rr equalrrrr re\s z tensor.__eq__cCst||}t|||SrT)rr not_equalrrrr rgas z tensor.__ne__cCst||}t|||SrT)rr rrrrr rfs ztensor.logical_andcCst||}t|||SrT)rr rrrrr rks ztensor.logical_orcCs t||SrT)r not_rrrr rrsztensor.__not__cCslt|tr|g}|}t|D]J\}}t|trH|jdurHt|||}q|tdddkrZqdsJdq|S)NF unsupported)r slice enumeraterrr Z expand_dims)rBZslicesrretZdimslrrr __getitem__vs ztensor.__getitem__cCsds JddS)NFz0Transposition must be created by the AST VisitorrrErrr r sztensor.TFcCs0t|tr|j}|r"t|||St|||SrT)r rrr bitcastcast)rBr!rrrrr tos  z tensor.to)N)N)N)N)N)N)N)N)N)N)N)N)N)N)N)N)N)N)N)N)N)N)N)N)N)N)N)N)N)N)N)N)N)FN)*r,r-r.rCrqrrrrrrrrrrrrrrrrrrrrrrrrrrrrrergrrrrrwr rrrrr rs                                  rcCst|tr|jS|SrTr)vrrr _constexpr_to_values rcCst|}t||S)z Returns the id of the current program instance along the given :code:`axis`. :param axis: The axis of the 3D launch grid. Has to be either 0, 1 or 2. :type axis: int )rr program_idaxisrrrr rsrcCst|}t||S)z Returns the number of program instances launched along the given :code:`axis`. :param axis: The axis of the 3D launch grid. Has to be either 0, 1 or 2. :type axis: int )rr num_programsrrrr rsrcCst|}t|}t|||S)a  Returns contiguous values within the open interval [:code:`start`, :code:`end`). :param start: Start of the interval. Must be a power of two. :type start: int :param stop: End of the interval. Must be a power of two >= start. :type stop: int )rr arange)startendrrrr rs rcCsjt|}t|D]J\}}t|ts2td|dt|jtstd|dt|jdqdd|DS)NzShape element  must have type `constexpr`1 must have type `constexpr[int]`, got `constexpr[rcSsg|] }t|qSr)rrrrrr rrz%_shape_check_impl..)rrr rr{rrr)ridrrr _shape_check_impls  rcCs(t|}t|}t|}t||||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 )rrr full)rrr!rrrr rs rcCst|||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 )r Zbroadcast_impl_value)inputrcrrrr broadcasts rcCst|}t|||S)z 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: Tuple[int] )rr Zbroadcast_impl_shaperrrrrr broadcast_tos rcCs t||SrT)r trans)rrrrr rsrFcCst||||S)aq 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) )r cat)rrcZ can_reorderrrrr r srcCst|}t|||S)z Returns a tensor with the same elements as `input` but a different shape. The order of the elements may not be preserved. :param input: The input tensor. :type input: :param shape: The desired shape. :type shape: Tuple[int] rr viewrrrr rs rcCst|}t|||SrTrrrrr reshape.srTcCst|}t||||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:`float16`, :code:`bfloat16`, :code:`float32`} :param other: The second tensor to be multiplied. :type other: 2D tensor of scalar-type in {:code:`float16`, :code:`bfloat16`, :code:`float32`} )rr dot)rrcZ allow_tf32rrrr r9s rc CsZt|durt||}t|dur,t||}t|}t|}t|}t|||||||S)a Return a tensor of data whose values are, elementwise, loaded from memory at location defined by :code:`pointer`. :code:`mask` and :code:`other` are implicitly broadcast to :code:`pointer.shape`. :code:`other` is implicitly typecast to :code:`pointer.dtype.element_ty`. :param pointer: Pointers to the data to be loaded. :type pointer: Block of dtype=triton.PointerDType :param mask: if mask[idx] is false, do not load the data at address :code:`pointer[idx]`. :type mask: Block of triton.int1, optional :param other: if mask[idx] is false, return other[idx] :type other: Block, optional :param cache_modifier: changes cache option in nvidia ptx 'type cache_modifier: str, optional N)rrr load)pointermaskrcZcache_modifierZeviction_policyZvolatilerrrr rNs    rcCs0t||}t|dur t||}t||||S)ad Stores :code:`value` tensor of elements in memory, element-wise, at the memory locations specified by :code:`pointer`. :code:`value` is implicitly broadcast to :code:`pointer.shape` and typecast to :code:`pointer.dtype.element_ty`. :param pointer: The memory locations where the elements of :code:`value` are stored. :type pointer: Block of dtype=triton.PointerDType :param value: The tensor of elements to be stored. :type value: Block :param mask: If mask[idx] is false, do not store :code:`value[idx]` at :code:`pointer[idx]`. :type mask: Block of triton.int1, optional N)rrr store)rrrrrrr rks   rrrzCallable[[T], T])r<rmcsdddfdd }|S)Nr funcrmcsd}|jd|_|S)Na+ Performs an atomic {name} 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 compare-and-swap. :type pointer: Block of dtype=triton.PointerDType :param cmp: The values expected to be found in the atomic object :type cmp: Block of dtype=`pointer.dtype.element_ty` :param val: The values to copy in case the expected value matches the contained value. :type val: Block of dtype=`pointer.dtype.element_ty` rDformatrrZdocstrrDrr _decorators z&_add_atomic_docstr.._decoratorrr<r rrDr _add_atomic_docstrsr zcompare-and-swapcCs$t||}t||}t||||SrT)rr atomic_cas)rcmpvalrrrr r s  r ZexchangecCst||}t||||SrT)rr atomic_xchgrrrrrrr rs rrcCst||}t||||SrT)rr atomic_addrrrr rs rmaxcCst||}t||||SrT)rr atomic_maxrrrr rs rmincCst||}t||||SrT)rr atomic_minrrrr rs rz logical andcCst||}t||||SrT)rr atomic_andrrrr rs rz logical orcCst||}t||||SrT)rr atomic_orrrrr rs rz logical xorcCst||}t||||SrT)rr atomic_xorrrrr rs rcCs.t||}t||}t||}t||||S)a Returns a tensor of elements from either :code:`x` or :code:`y`, depending on :code:`condition`. Note that :code:`x` and :code:`y` are always evaluated regardless of the value of :code:`condition`. If you want to avoid unintended memory operations, use the :code:`mask` arguments in `triton.load` and `triton.store` instead. The shape of :code:`x` and :code:`y` are both broadcast to the shape of :code:`condition`. :code:`x` and :code:`y` must have the 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. )rr where) conditionryrrrr rs   rcCs"t||}t||}t|||SrT)rr umulhi)rrrrrr rs  rcCst|}t||||SrT)rr fdiv)rr ieee_roundingrrrr rsrcsdddfdd }|S)Nr rcsd}|jd|_|S)Nzj Computes the element-wise {name} of :code:`x` :param x: the input values :type x: Block rDrr rDrr r sz)_add_math_1arg_docstr.._decoratorrr rrDr _add_math_1arg_docstrs r Z exponentialcCs t||SrT)r exprrrrr r!sr!znatural logarithmcCs t||SrT)r logr"rrr r#sr#ZcosinecCs t||SrT)r cosr"rrr r$sr$ZsinecCs t||SrT)r sinr"rrr r% sr%z square rootcCs t||SrT)r sqrtr"rrr r&&sr&csdddfdd }|S)Nr rcsd}|jd|_|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 reduction should be done rDrr rDrr r 2sz)_add_reduction_docstr.._decoratorrr rrDr _add_reduction_docstr0s r'maximumcCst|}t|||SrT)rr rrrrrrr r?sz maximum indexcCst|}t|||SrT)rr argmaxr)rrr r*Fsr*minimumcCst|}t|||SrT)rr rr)rrr rMsz minimum indexcCst|}t|||SrT)rr argminr)rrr r,Tsr,sumcCst|}t|||SrT)rr r-r)rrr r-[szxor sumcCst|}t|||SrT)rr xor_sumr)rrr r.bsr.cCs t|SrT)r debug_barrierrrrr r/nsr/cCs~t|tr|g}t|D]J\}}t|ts:td|dt|jtstd|dt|jdqdd|D}t||S)ze Let the compiler knows that the values in :code:`input` are all multiples of :code:`value`. values element rrrcSsg|] }|jqSrrrrrr rrzmultiple_of..) r rrr{rrrr multiple_ofrvaluesrrrrrr r1ss   r1cCs~t|tr|g}t|D]J\}}t|ts:td|dt|jtstd|dt|jdqdd|D}t||S)z_ Let the compiler knows that the `value` first values in :code:`input` are contiguous. r0rrrcSsg|] }|jqSrrrrrr rrz"max_contiguous..) r rrr{rrrr max_contiguousr2rrr r4s   r4cCst|dk|| SNr)rrrrr abssr7cCs||d|S)z Computes the ceiling division of :code:`x` by :code:`div` :param x: the input number :type input: Block :param div: the divisor :param div: Block rr)rdivrrr cdivs r9cCstj||k||S)z Computes the element-wise minimum of :code:`x` and :code:`y`. :param input: the first input tensor :type input: Block :param other: the second input tensor :type other: Block tritonlanguagerrrrrr r+s cCstj||k||S)z Computes the element-wise maximum of :code:`x` and :code:`y`. :param input: the first input tensor :type input: Block :param other: the second input tensor :type other: Block r:r=rrr r(s sigmoidcCsddtj| Sr|)r;r<r!r6rrr r>ssoftmaxcCs8|tj|d}tj|}tj|d}t|||Sr5)r;r<rr!r-r)rrznumZdenrrr r?s cCstj||jgS)zm Returns a contiguous flattened view of :code:`x` :param x: the input tensor :type x: Block )r;r<rrr6rrr ravelsrBc CsR|||}||}||}||}t|||}|||} |||} | | fS)a Transforms indices of a row-major size_i*size_j matrix into those of one where indices are row major for each group of size_j rows. For example, for size_i = size_j = 4 and size_g = 2, it will transform [[0 , 1 , 2 , 3 ], [4 , 5 , 6 , 7 ], [8 , 9 , 10, 11], [12, 13, 14, 15]] into [[0, 2, 4 , 6 ], [1, 3, 5 , 7 ], [8, 10, 12, 14], [9, 11, 13, 15]] )r+) rjZsize_iZsize_jZsize_gZijZsize_gjZgroup_idZoff_iZnew_iZnew_jrrr swizzle2ds   rDcCs t|d|S)a' Returns a tensor filled with the scalar value 0 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 dtype: Data-type of the new array, e.g., :code:`tl.float16` :type dtype: DType r)r)rr!rrr zeross rEcCst|j|jSrT)rErr!)rrrr zeros_like srFrc Gsddl}|}t|tr|j}t|ts4J|dd}|D]}||jvrr?rBrDrErFrJrPrrrr s8   *"0uZ