U 9%e @s(ddlmZddlmZddlmZddlmZddlm Z m Z m Z m Z ddl mZddlmZd d lmZmZe d Zd Zd Zd d dddZddddZddZGdddZGdddeZGdddeZGdddeZedZed Zed!Z ed"Z!ed#Z"ed$Z#ed%Z$ed&Z%ed'Z&ed(Z'ed)Z(ed*Z)ed+Z*ed,Z+ed-Z,ed.Z-ed/Z.ee"Z/Gd0d1d1Z0Gd2d3d3Z1d4d5Z2edd7d8Z3edd9d:Z4edd;d<Z5d=d>Z6edd?d@Z7eddAdBZ8eddCdDZ9eddEdFZ:eddHdIZ;eddJdKZeddPdQZ?edRe-d6fdSdTZ@ed6d6eAdUdUdUdGd6fdVdWZBeddYdZZCedd3d[d\d]ZDedd3d[d^d_ZEd`dadbdcddZFeeFdeddfdgZGeeFdhddidjZHeeFdkddldmZIeeFdnd dodpZJeeFdqd drdsZKeeFdtd dudvZLeeFdwd dxdyZMeeFdzd d{d|ZNedd}d~ZOedddZPedddZQd`dadbddZReeRddddZSeeRddddZTeeRddddZUeeRddddZVeeRddddZWeeRddddZXdd`d`d`dadddZYeddZZedddZ[edddZ\edddZ]eddZ^eddZ_eddZ`eddZaeddZbeddZceeYdddddddnZdeeYddddddZeeddZfeddZgeddZheddZieeYdddddddqZjeeYdddddd„ZkeddĄZleeYdŃdddńZmeddȄZneeYdɃd dd˄Zod!d`d`d`daddd̈́Zped"ddτZqeepdЃd#ddЄZreddӄZseepdԃd$ddԄZted%ddׄZued&ddلZved'ddۄZwed(dd݄Zxeddd6dGd6dd`d`dddZyed)ddZzed6dddZ{ed*ddZ|GdddZ}d+d`d`dddddddZ~d,d`d`ddddddZddZd6S(-) annotations)contextmanager)Enum)wraps)CallableListSequenceTypeVar)ir)jit)mathsemanticTiZ__triton_builtin__)fnreturncs0ts ttfdd}t|td|S)zMark a function as a builtin.cs&d|ks|ddkrtd||S)N_builderzdDid you forget to add @triton.jit ? (`_builder` argument must be provided outside of JIT functions.) ValueError)argskwargsrS/var/www/html/Darija-Ai-API/env/lib/python3.8/site-packages/triton/language/core.pywrappers zbuiltin..wrapperT)callableAssertionErrorrsetattrTRITON_BUILTIN)rrrrrbuiltins   r boolrcCs t|tdS)z-Is this a registered triton builtin function?F)getattrrrrrr is_builtin%sr$cCst|trt||tSt|trd|kr8dkrLnnt||tSd|kr`dkrtnnt||tSd|krdkrnnt| |t Sd|krdkrnnt| |t St d|dnt|t r^d }d d d }td |}|t dks<|dks<||ks<||kr8|krLnnt||tSt||tSn(t|trvt|j|St|tr|Sdstd|dt|ddS)NilllllzNonrepresentable integer .g8g?r absinfgFzcannot convert z of type z to tensor) isinstancer!tensorZget_int1int1intZ get_int32int32uint32Z get_int64int64uint64 RuntimeErrorfloat __builtins__Zget_fp32float32Zget_fp64float64 constexpr _to_tensorvaluertype)xbuilderZ min_float32Z max_float32Zabs_xrrrr7*sB          r7c@seZdZddddgZddddd gZd d d d dddgZd dddgZdgZGdddeZ ddZ ddZ ddZ ddZ ddZddZd d!Zd"d#Zd$d%Zd&d'Zd(d)Zd*d+Zd,d-Zd.d/Zd0d1Zd2d3Zd4d5Zd6d7Zd8d9Zd:d;Zdd?Zd@dAZ e!dBdCZ"e!dDdEZ#e!dFdGZ$ddHdIdJZ%ddHdKdLZ&dMdNZ'e(dOdPZ)dQdRdSdTdUZ*dVdWZ+e(dXdYdZd[Z,d\d]Z-d^S)_dtypeint8int16r-r/r+uint8uint16r.r0fp8e4b15fp8e4fp8e5fp16bf16fp32fp64voidc@seZdZdZdZdS)zdtype.SIGNEDNESSrr N)__name__ __module__ __qualname__SIGNEDUNSIGNEDrrrr SIGNEDNESSSsrNcCs||_|tjtjtjtjks*t||tjkr^tjj|_ t | dd|_ |j |_ n*|tjkrtjj|_ t | dd|_ |j |_ n|tjkrx|dkrd|_d|_ d|_n|dkrd|_d|_ d|_n|d krd |_d|_ d|_n|d krd |_d |_ d|_nh|dkr,d|_d |_ d|_nJ|dkrJd|_d|_ d|_n,|dkrhd|_d|_ d|_ntd|n|dkrd|_ dS)Nr,rArBrCr rD rEr&rF rG5@iz Unsupported floating-point type rHr)namer< SINT_TYPES UINT_TYPESFP_TYPES OTHER_TYPESrrNrLZint_signednessr,split int_bitwidthprimitive_bitwidthrMZfp_mantissa_widthZ exponent_biasr1)selfrZrrr__init__WsT$            zdtype.__init__cCs d|jkS)NZfp8rZrbrrris_fp8sz dtype.is_fp8cCs |jdkS)NrBrdrerrris_fp8e4szdtype.is_fp8e4cCs |jdkS)NrArdrerrr is_fp8e4b15szdtype.is_fp8e4b15cCs |jdkS)NrDrdrerrris_fp16sz dtype.is_fp16cCs |jdkS)NrErdrerrris_bf16sz dtype.is_bf16cCs |jdkS)NrFrdrerrris_fp32sz dtype.is_fp32cCs |jdkS)NrGrdrerrris_fp64sz dtype.is_fp64cCs |jdkS)Nr+rdrerrris_int1sz dtype.is_int1cCs |jdkS)Nr=rdrerrris_int8sz dtype.is_int8cCs |jdkS)Nr>rdrerrris_int16szdtype.is_int16cCs |jdkS)Nr-rdrerrris_int32szdtype.is_int32cCs |jdkS)Nr/rdrerrris_int64szdtype.is_int64cCs |jdkS)Nr?rdrerrris_uint8szdtype.is_uint8cCs |jdkS)Nr@rdrerrr is_uint16szdtype.is_uint16cCs |jdkS)Nr.rdrerrr is_uint32szdtype.is_uint32cCs |jdkS)Nr0rdrerrr is_uint64szdtype.is_uint64cCs |jtjkSN)rZr<r]rerrr is_floatingszdtype.is_floatingcCs |jtjkSrv)rZr<STANDARD_FP_TYPESrerrris_standard_floatingszdtype.is_standard_floatingcCs |jtjkSrv)rZr<r[rerrr is_int_signedszdtype.is_int_signedcCs |jtjkSrv)rZr<r\rerrris_int_unsignedszdtype.is_int_unsignedcCs|jtjtjkSrv)rZr<r[r\rerrris_intsz dtype.is_intcCs|Srv)rmrerrris_boolsz dtype.is_boolcCs tddS)NzNot implementedr1rrrris_voidsz dtype.is_voidcCsdSNFrrrrris_blockszdtype.is_blockcCsdSrrrrrris_ptrsz dtype.is_ptr)othercCst|tsdS|j|jkSr)r)r<rZrbrrrr__eq__s z dtype.__eq__cCs || Srvrrrrr__ne__sz dtype.__ne__cCs t|jfSrv)hashrZrerrr__hash__szdtype.__hash__cCs|Srvrrerrrscalarsz dtype.scalar ir.builderzir.typer;rcCs|jdkr|S|jdkr$|S|jdkr6|S|jdkrH|S|jdkrZ|S|jdkrl|S|jdkr~|S|jdkr|S|jd kr| S|jd kr| S|jd kr| S|jd kr| S|jd kr| Std|ddS)NrHr+)r=r?)r>r@)r-r.)r/r0rCrBrArDrErFrGzfail to convert z to ir type)rZZ get_void_tyZ get_int1_tyZ get_int8_tyZ get_int16_tyZ get_int32_tyZ get_int64_tyZ get_fp8e5_tyZ get_fp8e4_tyZget_fp8e4b15_tyZ get_half_tyZ get_bf16_tyZ get_float_tyZ get_double_tyrrbr;rrrto_irs6             z dtype.to_ircCs|jSrvrdrerrr__str__sz dtype.__str__strr"cCs|jS)z"See cache_key_part() in triton.cc.rdrerrrcache_key_partszdtype.cache_key_partcCs d|jS)Nztriton.language.rdrerrr__repr__szdtype.__repr__N).rIrJrKr[r\r]rxr^rrNrcrfrgrhrirjrkrlrmrnrorprqrrrsrtrurwryrzr{r|r} staticmethodrrrrrrpropertyrrrrrrrrrr<LsZ  -    r<c@sreZdZddddddZddd d d Zd d ZddZddZdddddZdddddZ e ddZ dS) pointer_typer r<r,) element_ty address_spacecCs,t|tstd||_||_||_dS)Nz,element_ty is a {type(element_ty).__name__}.)r)r< TypeErrorrrrrZ)rbrrrrrrc s  zpointer_type.__init__rzir.pointer_typercCs||j|dS)Nr )Z get_ptr_tyrrrrrrrszpointer_type.to_ircCsd|jdS)Nzpointer<>rrerrrrszpointer_type.__str__cCs|Srvrrerrrrszpointer_type.__repr__cCsdSNTrrerrrrszpointer_type.is_ptrr!rrcCs&t|tsdS|j|jko$|j|jkSr)r)rrrrrrrrs zpointer_type.__eq__cCs || Srvrrrrrr$szpointer_type.__ne__cCs|Srvrrerrrr'szpointer_type.scalarN)r ) rIrJrKrcrrrrrrrrrrrrr src@s~eZdZdddddZdddd d Zd d Zd dZddZddddZdddddZ dddddZ e ddZ dS) block_typer<r)rshapecCs||_|stdt|dtr.dd|D}||_d|_|jD]}|j|9_q@|jtkrvtd|jdtd||_ dS) Nz0d block_type is forbiddenrcSsg|] }|jqSrr8.0srrr 7sz'block_type.__init__..r znumel (z') exceeds triton maximum tensor numel ()) rrr)r6rnumelTRITON_MAX_TENSOR_NUMELrrrZ)rbrrrrrrrc-s  zblock_type.__init__rz ir.block_typercCs||j||jSrv)Z get_block_tyrrrrrrrrBszblock_type.to_ircCsd|jd|jdS)N )rrrerrrrcszfunction_type.__str__rr;cs4fdd|jD}fdd|jD}||S)Ncsg|]}|qSrrrtyrrrrgsz'function_type.to_ir..csg|]}|qSrr)rret_typerrrrhs)rrZget_function_ty)rbr;Zir_param_typesrrrrrfszfunction_type.to_irN)rIrJrKrcrrrrrrr^srrHr+r=r>r-r/r?r@r.r0rCrBrArDrErFrGc@sFeZdZdZddZddddZdd 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'dLdMZ(dNdOZ)dPS)Qr6zL This class is used to store a value that is known at compile-time. cCst|tr|j|_n||_dSrvr)r6r8)rbr8rrrrcs  zconstexpr.__init__rr"cCsd|jdS)Nz constexpr[]rrerrrrszconstexpr.__repr__cCs|jSrvrrerrr __index__szconstexpr.__index__cCst|j|jSrvr6r8rrrr__add__szconstexpr.__add__cCst|j|jSrvrrrrr__radd__szconstexpr.__radd__cCst|j|jSrvrrrrr__sub__szconstexpr.__sub__cCst|j|jSrvrrrrr__rsub__szconstexpr.__rsub__cCst|j|jSrvrrrrr__mul__szconstexpr.__mul__cCst|j|jSrvrrrrr__mod__szconstexpr.__mod__cCst|j|jSrvrrrrr__rmul__szconstexpr.__rmul__cCst|j|jSrvrrrrr __truediv__szconstexpr.__truediv__cCst|j|jSrvrrrrr __rtruediv__szconstexpr.__rtruediv__cCst|j|jSrvrrrrr __floordiv__szconstexpr.__floordiv__cCst|j|jSrvrrrrr __rfloordiv__szconstexpr.__rfloordiv__cCst|j|jkSrvrrrrr__gt__szconstexpr.__gt__cCst|j|jkSrvrrrrr__rgt__szconstexpr.__rgt__cCst|j|jkSrvrrrrr__ge__szconstexpr.__ge__cCst|j|jkSrvrrrrr__rge__szconstexpr.__rge__cCst|j|jkSrvrrrrr__lt__szconstexpr.__lt__cCst|j|jkSrvrrrrr__rlt__szconstexpr.__rlt__cCst|j|jkSrvrrrrr__le__szconstexpr.__le__cCst|j|jkSrvrrrrr__rle__szconstexpr.__rle__cCst|j|jkSrvrrrrrrszconstexpr.__eq__cCst|j|jkSrvrrrrrrszconstexpr.__ne__cCs t|jSrv)r!r8rerrr__bool__szconstexpr.__bool__cCs t|j Srvrrerrr__neg__szconstexpr.__neg__cCst|j|j@Srvrrrrr__and__szconstexpr.__and__cCst|jo |jSrvrrrrr logical_andszconstexpr.logical_andcCst|j|jBSrvrrrrr__or__szconstexpr.__or__cCst|j|jASrvrrrrr__xor__szconstexpr.__xor__cCst|jp |jSrvrrrrr logical_orszconstexpr.logical_orcCs t|j Srvrrerrr__pos__szconstexpr.__pos__cCs t|jSrvrrerrr __invert__szconstexpr.__invert__cCst|j|jSrvrrrrr__pow__szconstexpr.__pow__cCst|j|j?Srvrrrrr __rshift__szconstexpr.__rshift__cCst|j|j>Srvrrrrr __lshift__szconstexpr.__lshift__cCs t|j Srvrrerrr__not__szconstexpr.__not__cOs |j||Srvr)rbrkwdsrrr__call__szconstexpr.__call__N)*rIrJrK__doc__rcrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrr6sNr6c@sFeZdZddddZddddZed[d d Zd\d d Zed]ddZd^ddZ ed_ddZ d`ddZ edaddZ dbddZ edcddZeddddZededdZedfd d!Zedgd"d#Zedhd$d%Zedid&d'Zedjd(d)Zedkd*d+Zedld,d-Zedmd.d/Zednd0d1Zedod2d3Zedpd4d5Zedqd6d7Zedrd8d9Zedsd:d;Zedtdd?Z edvd@dAZ!edwdBdCZ"edxdDdEZ#edydFdGZ$edzdHdIZ%ed{dJdKZ&ed|dLdMZ'ed}dNdOZ(ed~dPdQZ)eddRdSZ*eddTdUZ+e,dVdWZ-eddYdZZ.d S)r*r<)r9cCsl||_d|_|r|j|_d|_|jD]}|j|9_q(t|j|_||_|j|_dd|jD|_dS)N)r r cSsg|] }t|qSrr6rrrrrsz#tensor.__init__..)handlerrrr6r9rr<)rbrr9rrrrrcs  ztensor.__init__rr"cCs(t|jdddd|jDdS)N[,css|]}t|VqdSrv)rrrrr sz!tensor.__str__..r)rr<joinrrerrrrsztensor.__str__NcCst||}t|||Srv)r7raddrbrrrrrrs ztensor.__add__cCs|j||dSNr)rrrrrrsztensor.__radd__cCst||}t|||Srvr7rsubrrrrrs ztensor.__sub__cCst||}t|||Srvrrrrrr#s ztensor.__rsub__cCst||}t|||Srv)r7rmulrrrrr's ztensor.__mul__cCs|j||dSr)rrrrrr,sztensor.__rmul__cCst||}t|||Srvr7rtruedivrrrrr/s ztensor.__truediv__cCst||}t|||Srvrrrrrr4s ztensor.__rtruediv__cCst||}t|||Srvr7rfloordivrrrrr8s ztensor.__floordiv__cCst||}t|||Srvrrrrrr=s ztensor.__rfloordiv__cCst||}t|||Srvr7rmodrrrrrBs ztensor.__mod__cCst||}t|||Srvrrrrr__rmod__Gs ztensor.__rmod__cCs t||Srv)rminusrbrrrrrMsztensor.__neg__cCs t||Srv)rinvertrrrrrQsztensor.__invert__cCst||}t|||Srvr7rand_rrrrrWs ztensor.__and__cCst||}t|||Srvrrrrr__rand__\s ztensor.__rand__cCst||}t|||Srvr7ror_rrrrras z tensor.__or__cCst||}t|||Srvrrrrr__ror__fs ztensor.__ror__cCst||}t|||Srvr7rZxor_rrrrrks ztensor.__xor__cCst||}t|||Srvrrrrr__rxor__ps ztensor.__rxor__cCst||}t|||Srvr7rZshlrrrrrus ztensor.__lshift__cCst||}t|||Srvrrrrr __rlshift__zs ztensor.__rlshift__cCs4t||}|jr"t|||St|||SdSrvr7r<rzrZashrZlshrrrrrrs  ztensor.__rshift__cCs4t||}|jr"t|||St|||SdSrvrrrrr __rrshift__s  ztensor.__rrshift__cCst||}t|||Srvr7r greater_thanrrrrrs z tensor.__gt__cCst||}t|||Srvrrrrrrs ztensor.__rgt__cCst||}t|||Srvr7rZ greater_equalrrrrrs z tensor.__ge__cCst||}t|||Srvrrrrrrs ztensor.__rge__cCst||}t|||Srvr7r less_thanrrrrrs z tensor.__lt__cCst||}t|||Srvrrrrrrs ztensor.__rlt__cCst||}t|||Srvr7rZ less_equalrrrrrs z tensor.__le__cCst||}t|||Srvrrrrrrs ztensor.__rle__cCst||}t|||Srv)r7requalrrrrrs z tensor.__eq__cCst||}t|||Srv)r7r not_equalrrrrrs z tensor.__ne__cCst||}t|||Srv)r7rrrrrrrs ztensor.logical_andcCst||}t|||Srv)r7rrrrrrrs ztensor.logical_orcCs t||Srv)rnot_rrrrrsztensor.__not__cCst|tr|g}|}t|D]h\}}t|trH|jdkrHt|||}qt|trr|jdkrr|jdkrr|j dkrrqdst d|q|S)NFzunsupported tensor index: ) r)slice enumerater6r8r expand_dimsstartstopstepr)rbZslicesrretdimslrrr __getitem__s (ztensor.__getitem__cCsds tddS)NFz0Transposition must be created by the AST Visitor)rrerrrrsztensor.TFcCs0t|tr|j}|r"t|||St|||Srv)r)r6r8rbitcastcast)rbr<rrrrrtos  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)N)N)N)N)N)FN)/rIrJrKrcrr rrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrr*s                                       r*cCst|tr|jS|Srvr)vrrr_constexpr_to_values rNcCst|}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_idaxisrrrrrsrcCst|}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_programsrrrrrsrcCst|}t|}t|||S)ar Returns contiguous values within the left-closed and right-open interval [:code:`start`, :code:`end`). End - Start must be less than or equal to 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 > start. :type end: int32 )rrarange)r endrrrrr s 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)rrr:rrrr7sz%_shape_check_impl..)rrr)r6rr8r,r9)ridrrr_shape_check_impl0s  r cCs(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 )r rrfull)rr8r<rrrrr!:s r!cCst|||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 )rZbroadcast_impl_value)inputrrrrr broadcastPs r#cCst|}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] )r rZbroadcast_impl_shaper"rrrrr broadcast_to]s r%cCs t||Srv)rtrans)r"rrrrr&ksr&FcCst||||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) )rcat)r"rZ can_reorderrrrrr'psr'cCst|}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] )r rviewr$rrrr(s r(cCst|}t|||Srv)r rreshaper$rrrr)sr)cCsH| |kr|ks4ntd|d| d||dkr@|S||S)Nz invalid axis z . Expected z <= axis < rr)rndimrrr _wrap_axissr+cst|}t|trt|n|g}t|jt|fdd|D}tt|t|krftd||}t|D]}t |||}qr|S)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] csg|]}tt|qSr)r+rrrZnew_ndimrrrszexpand_dims..z7expand_dims recieved duplicate axes, normalized axes = ) rr)rlistlenrsetrsortedrr)r"rrZaxesr arr-rrs rTcCs"t|}t|}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`} )rrdot)r"rZ allow_tf32Z out_dtyperrrrr3s r3c Csft|dk rt||}t|dk r,t||}t|}t|}t|}t|}t||||||||| S)al Return a tensor of data whose values are loaded from memory at location defined by `pointer`: (1) `pointer` could be a single element pointer, then a scalar will be loaded - `mask` and `other` must be scalar too - `other` is implicitly typecast to `pointer.dtype.element_ty` - `boundary_check` and `padding_option` must be empty (2) `pointer` could be element-wise tensor of pointers, in which case: - `mask` and `other` are implicitly broadcast to `pointer.shape` - `other` is implicitly typecast to `pointer.dtype.element_ty` - `boundary_check` and `padding_option` must be empty (3) `pointer` could be a block pointer defined by `make_block_ptr`, in which case: - `mask` and `other` must be None - `boundary_check` and `padding_option` can be specified to control the behavior of out-of-bound access :param pointer: Pointer to the data to be loaded :type pointer: `triton.PointerType`, or block of `dtype=triton.PointerType` :param mask: if `mask[idx]` is false, do not load the data at address `pointer[idx]` (must be `None` with block pointers) :type mask: Block of `triton.int1`, optional :param other: if `mask[idx]` is false, return `other[idx]` :type other: Block, optional :param boundary_check: tuple of integers, indicating the dimensions which should do the boundary check :type boundary_check: tuple of ints, optional :param padding_option: should be one of {"", "zero", "nan"}, do padding while out of bound :param cache_modifier: changes cache option in NVIDIA PTX :type cache_modifier: str, optional :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 N)rr7rload) pointermaskrboundary_checkZpadding_optioncache_modifiereviction_policyZvolatilerrrrr5s#    r5rc CsFt||}t|dk r t||}t|}t|}t|||||||S)a Store a tensor of data into memory locations defined by `pointer`: (1) `pointer` could be a single element pointer, then a scalar will be stored - `mask` must be scalar too - `boundary_check` and `padding_option` must be empty (2) `pointer` could be element-wise tensor of pointers, in which case: - `mask` is implicitly broadcast to `pointer.shape` - `boundary_check` must be empty (3) or `pointer` could be a block pointer defined by `make_block_ptr`, in which case: - `mask` must be None - `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 :param eviction_policy: changes eviction policy in NVIDIA PTX :type eviction_policy: str, optional N)r7rrstore)r6r8r7r8r9r:rrrrr;s    r;)basec Cst|||||||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)r<rstridesoffsetsZ block_shapeorderrrrrr=(s r=cCst|||S)z Advance a block pointer :param base: the block pointer to advance :param offsets: the offsets to advance, a tuple by dimension )radvance)r<r?rrrrrA7srArzCallable[[T], T])rZrcsdddfdd }|S)Nrfuncrcsd}|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` rdformatrrCZdocstrrdrr _decoratorHs z&_add_atomic_docstr.._decoratorrrZrGrrdr_add_atomic_docstrFsrIzcompare-and-swapcCs.t||}t||}t|}t|||||Srv)r7rr atomic_cas)r6cmpvalsemrrrrrJ[s  rJZexchangecCs$t||}t|}t|||||Srv)r7rr atomic_xchgr6rLr7rMrrrrrNds rNrcCs$t||}t|}t|||||Srv)r7rr atomic_addrOrrrrPls rPmaxcCs$t||}t|}t|||||Srv)r7rr atomic_maxrOrrrrRts rRmincCs$t||}t|}t|||||Srv)r7rr atomic_minrOrrrrT|s rTz logical andcCs$t||}t|}t|||||Srv)r7rr atomic_andrOrrrrUs rUz logical orcCs$t||}t|}t|||||Srv)r7rr atomic_orrOrrrrVs rVz logical xorcCs$t||}t|}t|||||Srv)r7rr atomic_xorrOrrrrWs rWcCs.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 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. )r7rwhere) conditionr:yrrrrrXs   rXcCs"t||}t||}t|||Srv)r7rumulhi)r:rZrrrrr[s  r[cCst|}t||||Srv)rrfdiv)r:rZZ ieee_roundingrrrrr\sr\csdddfdd }|S)NrrBcsd}|jd|_|S)Nzk Computes the element-wise {name} of :code:`x`. :param x: the input values :type x: Block rdrDrFrdrrrGsz)_add_math_1arg_docstr.._decoratorrrHrrdr_add_math_1arg_docstrs r]Z exponentialcCs t||Srv)rexpr:rrrrr^sr^znatural logarithmcCs t||Srv)rlogr_rrrr`sr`ZcosinecCs t||Srv)rcosr_rrrrasraZsinecCs t||Srv)rsinr_rrrrbsrbz square rootcCs t||Srv)rsqrtr_rrrrcsrczabsolute valuecCs t||Srv)rr'r_rrrr'sr')rZreturn_indices_arg tie_break_argrcsdddfdd }|S)NrrBcsLd}dk r"|ddd7}dk r:|dd7}|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 donez :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 NaNrdrDrFrZrdrerrrGs  z)_add_reduction_docstr.._decoratorrrZrdrerGrrfr_add_reduction_docstrsrhccs|}dV||dSrv)Zget_insertion_pointZrestore_insertion_point)r;iprrr_insertion_guardsrjcsTttr"tf|ddSfdd}|dk rDt|}t||S)aRApplies 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 :param combine_fn: a function to combine two groups of scalar tensors (must be marked with @triton.jit) r _generatorrc sddD}t||d}|d}tzfdd|jD}||fddt|jD}j|id}t|tr|j g}ndd|D}j |W5QRXdS) NcSsg|] }|jjqSrr9rrtrrrr)sz7reduce..make_combine_region..r rcsg|]}|qSrrrrrrr.scs g|]\}}t||qSrr*argrrrblockrrr0srcSsg|] }|jqSrrrrrrrr6s) r get_regionrjrcreate_block_with_parentrcall_JitFunctionr)r*rZcreate_reduce_ret)Z reduce_op in_scalar_tys prototyperegionrrresultshandlesrrl combine_fnr"rsrmake_combine_region(s      z#reduce..make_combine_regionN)r)r*reducerrZ reductionr"rrrrlrrrrrs  rcCsB|jj}|r(|jdkr(|jt|dS|tkr>|jt|dS|S)NrWr)r9rr|r`rr-bfloat16r4)ror scalar_tyrrr_promote_reduction_input=s rc Cst|}|j|}td||d}t|jdkrjddtt|jD}||=t|||d}t||j|d}t||f||||d\}} || fS)Nrrr cSsg|] }t|qSrrr,rrrrUsz(_reduce_with_indices..rk)rrrr/rangerr%r) r"rrrrlnindexZaxes_to_expandZrvalueZrindicesrrr_reduce_with_indicesMs   rcCst||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 rXr:rZrrrminimum_s rcCst||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 rrrrrmaximumls rc CsF|r||ko||k}nd}||kp$|}t|||}t|||}||fSrr) value1index1value2index2tie_break_lefttiegtZv_retZi_retrrr_argmax_combine{s   rcCst||||dSrrrrrrrrr_argmax_combine_tie_break_leftsrcCst||||dSrrrrrr_argmax_combine_tie_break_fastsrcCs t||Srv)rrQrrrr _fast_maxsrreturn_indicesreturn_indices_tie_break_left)rdrecCs|t|}|r*|rt||tSt||tSnNt|jjdkrlt|jrT|t }n|j sbt |t }t ||tSdSNrW)rrrrr6r<rarwrr4is_integer_typerr-rrr"rrrrrrrQs   z maximum indexr)recCst||d|d\}}|SNT)rr)rQr"rr_r rrrargmaxsrc CsF|r||ko||k}nd}||kp$|}t|||}t|||}||fSrr) rrrrrrltZ value_retZ index_retrrr_argmin_combines   rcCst||||dSrrrrrr_argmin_combine_tie_break_leftsrcCst||||dSrrrrrr_argmin_combine_tie_break_fastsrcCs t||Srv)rrSrrrr _fast_minsrcCs|t|}|r*|rt||tSt||tSnNt|jjdkrlt|jrT|t }n|j sbt |t }t ||tSdSr)rrrrr6r<rarwrr4rrr-rrrrrrrSs   z minimum indexcCst||d|d\}}|Sr)rSrrrrargminsrcCs||Srvrr2brrr _sum_combinesrsumcCst|}t||tSrv)rrrr"rrrrrscCs||ASrvrrrrr _xor_combinesrzxor sumcCs6|jj}|stdt||d}t||t||dS)Nz#xor_sum only supported for integersrrk)r9rr|rrrr)r"rrrlrrrrxor_sums rcsdddfdd }|S)NrrBcsd}|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 scan should be donerdrDrFrdrrrGsz$_add_scan_docstr.._decoratorrrgrrdr_add_scan_docstrs rcsLttr"tf|ddSfdd}t|}t||S)auApplies 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) rkrc sddD}t||d}|d}tzfdd|jD}||fddt|jD}j|id}t|tr|j g}ndd|D}j |W5QRXdS) NcSsg|] }|jjqSrrmrnrrrr-szAassociative_scan..make_combine_region..r rcsg|]}|qSrrrrrrr2scs g|]\}}t||qSrrprrrsrrr4srucSsg|] }|jqSrrvrwrrrr:s) rryrjrrzrr{r)r*rZcreate_scan_ret)Zscan_opr|r}r~rrrrrrsrr,s      z-associative_scan..make_combine_region)r)r*associative_scanrrrrrrrs  rcumsumcCst|}t||tSrv)rrrrrrrrBscCs||Srvrrrrr _prod_combineLsrcumprodcCst|}t||tSrv)rrrrrrrrQscCs t|S)zA Insert a barrier to synchronize all threads in a block. )r debug_barrierrrrrr]srcCs~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|] }|jqSrrrrrrrqszmultiple_of..) r)r6rrr8r,r9r multiple_ofr"valuesrrrrrrres   rcCs~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. rrrrcSsg|] }|jqSrrrrrrrsz"max_contiguous..) r)r6rrr8r,r9rmax_contiguousrrrrrus   rcCs~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 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]. rrrrcSsg|] }|jqSrrrrrrrsz!max_constancy..) r)r6rrr8r,r9r max_constancyrrrrrs   r  )seprfileflushr)rrcGsdS)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)rrrrrrrrr static_prints rcCsdS)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)condmsgrrrr static_asserts rrcGsddl}t|}t|ts(t|dd}|D]}||jkr0d}qHq0|sZt|dg}|D]}|t||qbt |||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) :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. rNz is not stringTFz is not an ascii string) stringrr)rr printableappendr7r device_print)prefixrrrZb_asciichnew_argsrqrrrrs rc Cst|}ddl}|}||}t|dr>|j}||}q"d}d}d}|dk rl|jj}|jjj}|jj }t 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. rNrIunknown) rinspect currentframe getmodulehasattrf_backf_codeco_name co_filenamef_linenor device_assertr7) rrrrframemodulelineno func_name file_namerrrrs    rc@s*eZdZdZd ddZddZddZdS) 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. NcCsnt|tst|dkr"td|_nt|ts0t||_|dkrPtd|_||_nt|ts^t||_||_dS)Nr r)r)r6rr r r)rbZarg1Zarg2r rrrrcs  zstatic_range.__init__cCs tddSNz8static_range can only be used in @triton.jit'd functionsr~rerrr__iter__+szstatic_range.__iter__cCs tddSrr~rerrr__next__.szstatic_range.__next__)NN)rIrJrKrrcrrrrrrr s rr.dicttuple)lib_namelib_pathrarg_type_symbol_dict ret_shapeis_purec Cs t|dkrtdtt|d}t||krLtdt|d|g} g} |D]@} t| tr| | j| | jqX| t | | | qXt | } | |krtd|d| nD|| d} || d} |rt | |} t|||| | | ||| SdS)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 N) r/rr.keysr)r*rr<rr9rrr)rCrrrrrrrnum_args arg_typesarg_listrqsymbolrrrrdispatch6s*        r)rrrrrc Cs|}d}d}g} tt|D]8} t|| ||| <| || j|| jr d}q t| dkrt| } d} | |kr~d} |d} t |D]\} } t j | | || d\}} qtt|D]"} t j || | || d\|| <}q|s| j }t |d}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)arithmetic_checkZcreate_extern_elementwise)copyrr/r7rr<r9rrrrZbinary_op_type_checking_implrr#r)rrrrrrZ dispatch_argsZ all_scalarrrrrZ broadcast_argitemrrCrrrextern_elementwise`sB    rcCst|S)z#A decorator for external functions.)r rrrrexternsr)N)N)N)N)N)N)N)FN)N)N)N)Nrr4r4N)N)N)NN)NNN)NNN)NNN)NNN)NNN)NNN)NNN)N)N)FN)N)N)N)N)N)N)NN)NN)N)NN)NFT)T)NFT)T)N)NNN)NN)NN)r)r)N)N)N)N)r4N)r4N)N)N) __future__r contextlibrenumr functoolsrtypingrrrr Z_C.libtriton.tritonr Z runtime.jitr r4rrrrrr r$r7r<rrrrHr+r=r>r-r/r?r@r.r0Zfloat8e5Zfloat8e4Z float8e4b15Zfloat16rr4r5Zpi32_tr6r*rrrrr r!r#r%r&r'r(r)r+rr3rr5r;r=rArIrJrNrPrRrTrUrVrWrXr[r\r]r^r`rarbrcr'rhrjrrrrrrrrrrQrrrrrrSrrrrrrrrrrrrrrrrrrrrrrrrrrs      "?"2{v           .$  !             "    !--*)