o o h2W@sddlmZddlZddlmZmZmZmZmZddl Z ddl m Z ddl m ZedZGd d d eZdPddZdPddZdQddZdRddZdSdTd"d#ZdUd(d)Z *dVdWd/d0ZdXd3d4ZdYd8d9ZdYd:d;ZdYdd?ZdZd@dAZd[dCdDZdZdEdFZ d\dKdLZ!d\dMdNZ"d]dQdRZ#d^dSdTZ$d_dUdVZ%d_dWdXZ&d_dYdZZ'd_d[d\Z(d_d]d^Z)d`d_d`Z*d_dadbZ+d_dcddZ,d_dedfZ-dadgdhZ.dbdidjZ/dcdkdlZ0dddodpZ1d_dqdrZ2d_dsdtZ3d_dudvZ4d_dwdxZ5d_dydzZ6d_d{d|Z7deddZ8dfddZ9dgddZ:dhddZ;diddZ<djddZ=dkddZ>dlddZ?dmddZ@dnddZAdoddZBdpddZCdqddZD drdsddZEddZFddZGddZHddZIddZJddZKddZLddZMddZNdtddƄZOduddʄZPdd̄ZQdvddЄZRdwddӄZSdxddՄZTdyddׄZUdzddZVd{ddZWddZXddZYd|ddZZd}ddZ[d~ddZ\dddZ]dddZ^dddZ_dddZ`dddZadddZbdddZcddZddd d Zedd dZfdddZgdddZhdddZiddZjdd!d"Zkdd$d%Zldd(d)Zmdd+d,Zndd.d/Zodd0d1Zpdd2d3Zqdd4d5Zrdd9d:Zsdd=d>Ztdd?d@ZudAdBZvdSdCdDZwddFdGZxddHdIZyddNdOZzdS() annotationsN)ListOptionalSequenceTupleTypeVar)ir)coreTcseZdZfddZZS)IncompatibleTypeErrorImplcs@||_||_d|jd|j|_tt||jdS)Nzinvalid operands of type  and )type_atype_b__repr__messagesuperr __init__)selfrr __class__l/var/www/html/construction_image-detection-poc/venv/lib/python3.10/site-packages/triton/language/semantic.pyrsz"IncompatibleTypeErrorImpl.__init__)__name__ __module__ __qualname__r __classcell__rrrrr sr axisintbuilder ir.builderreturn tl.tensorcC*|dvr td|t||tjS)Nrr rz+program_id axis must be 0, 1, or 2 but got ) ValueErrortltensorcreate_get_program_idint32rr rrr program_idr,cCr$)Nr%z-num_programs axis must be 0, 1, or 2 but got )r&r'r(create_get_num_programsr*r+rrr num_programs!r-r/a_tytl.dtypeb_tycCs|j}|j}|j}|j}||kr||kr|S|S|tjjjkr'||kr%|S|S|tjjjkr6||kr4|S|Std|d|)Nzunexpected signedness r) int_bitwidthint_signednessr'dtype SIGNEDNESSUNSIGNED TypeError)r0r2a_rankb_ranka_snb_snrrrinteger_promote_impl,sr= a_is_scalarbool b_is_scalar div_or_modcCsT||kr)|r ||fn||f\}}|j|jkr)|r'|tjtjfvr'tjS|S|s1|r4tjS|s<|r?tjS| sG| rO|rLtjStjS| r_| r_|r\tjStjS| sg| rjtjS| r{| r{||krx|StjS| r| st d|d||r|j|jkrt d|d|dt||S)Nunexpected type rzCannot use /, #, or % with x because they have different signedness;this is unlikely to result in a useful answer. Cast them to the same signedness.)kindvaluer'float16bfloat16float32is_fp64float64is_fp32is_fp16is_bf16is_fp8is_intr8r4rr=)r0r>r2r@rA scalar_ty tensor_tyrrrcomputation_type_impl<s: rRT check_typecCst|trt||tjSt|trdd|krdkr$nntj}n8d|kr.dkr4nntj}n(d|kr>dkrDnntj }nd|krNdkrTnntj }nt d|dt d |||d St|t rd }d d d}td|}|t dks|dks||ks||kr|krnntj}ntj}t d |||d St|tjrt|j|St|tjr|S|rtd|dt|d|S)NllzNonrepresentable integer .rr5r g8g?rabsinfzcannot convert z of type z to tensor) isinstancer?r'r(get_int1int1rr*uint32int64uint64r&fullfloat __builtins__rHrJ constexpr to_tensorrEr8type)xr rSr5 min_float32 max_float32abs_xrrrrhos>         rhrr allow_ptr_aNonecCsJ|r!|s t|||r||krt|||r#t||dSdSN)is_ptrr is_floating)rrrnrrrcheck_ptr_type_impls   rsFlhstl.tensor | numbers.NumberrhsTuple[tl.tensor, tl.tensor]cCst|tj}t|tj}|r|} t||}|r|} t||}|jj} |jj} t| | |t| | ||r| s| st| || ||} |rN| dkrN| sX|r\| dkr\| r\t d| r|r{| | krp| ks{nt d| d| |r| | kr| ksnt d| d| |rtd| | |dnt|| |}|rtd| | |dnt|| |}t|||\}}||fS)Nrz{Cannot perform a binary operation between an unsigned tensor and a negative scalar. Perform a explicit cast on one of them.zScalar z is out of range for type rrY)r^numbersNumberrhriscalarrsrqrRis_int_unsignedr&rOget_int_min_valueget_int_max_valuerdcastbroadcast_impl_value)rtrvr allow_lhs_ptr allow_rhs_ptrarithmetic_checkrA lhs_is_scalar rhs_is_scalar lhs_scalar rhs_scalar lhs_sca_ty rhs_sca_ty ret_sca_tyrrrbinary_op_type_checking_implsL      ""  r binary_opcallablec Cs|jjjdks |jjs dS|jj}|jj}||ksJ|s!Jt|tj|}t|tj|}|||d|}| }t | |tj}| }t | |tj}t t|||t||||} d|jd|j} t| | |dS)N@Frz! overflow detected for operation )rirzr3optionssanitize_overflowrOr~r'rbr}r( get_int64r|and_ less_equal greater_equalr device_assert) rtrvr rrrret max_value min_valuecondmsgrrr binary_op_sanitize_overflow_impls   rinputotherrcCsHt|||dd\}}|jj}|jj}|r|rtd|r3|s3||}}|jj}|jj}|rr|j}|jrf|jjdkrf|j rXt t j |j |}nt j |}||j|d}t ||j||jS|rt ||j|j|jS|r|rt|||tt ||j|j|jStd|)NTzcannot add pointers togetherrFrB)rrirzrqr8handler5r{r3is_blockr' block_typerbget_block_shapesto_ircreate_int_castr( create_addptrrr create_faddrOradd create_add)rrrr input_scalar_tyother_scalar_ty other_handlei64_tyrrrrs0   rcCst|||dd\}}|jj}|r"t||jt||j|jS| r3t| |j|j|jS| rM|r@t |||t t||j|j|jStd|)NTFrB)rrirzrqr'r(rrminusrr create_fsubrOrsub create_subr8rrrr rPrrrrs rcCs|t|||\}}|jj}|rt||j|j|jS|r7|r*t |||t t| |j|j|jSt d|NrB) rrirzrrr'r( create_fmulrrOrmul create_mulr8rrrrrsrcCst|||dddd\}}|jj}|jj}|r#|r#t|||}nI|r2|r2t|||}n:|rI|rIt|tj|}t|tj|}n#|re|re|j|jkr^t|||}nt|||}nt d|t | |j |j |jSNFTrB) rrirzrrrOr~r'rHfp_mantissa_widthr8r( create_fdivr)rrr rrrrrtruediv#s  rcCst|||dddd\}}|jj}|jj}|rK|rKt||}t|||}t|||}|r>t| |j |j |jSt| |j |j |jSt d|r) rrirzrOr=r~ is_int_signedr'r( create_sdivr create_udivr8)rrr rrret_tyrrrfloordiv=s   r ieee_roundingcCs^|jj}|jj}|r|stdt|||dddd\}}||j|j}t||jS)Nz4both operands of fdiv must have floating scalar typeFT) rirzrrr8rrrr'r()rrrr rrrrrrfdivLsrcCst|||dddd\}}|jj}|jj}|r%t||j|j|jS|r]|j |j kr?t d| d| d| rPt| |j|j|jSt||j|j|jSt d|)NFTz Cannot mod z by rCrB)rrirzrrr'r( create_fremrrOr4r8rr create_srem create_urem)rrr rPrrrrmodWs  rrjy propagate_nantl.PropagateNancCt|||\}}|j}|r<|tjjkr"t||j|j|j S|tjj kr5t| |j|j|j St d|| rMt||j|j|j S|r^t||j|j|j Std|NzUnexpected propagate_nan Unexpected dtype )rr5rrr' PropagateNanALLr(create_minimumfrriNONEcreate_minnumfr&r create_minsir{ create_minuir8rjrrr r5rrrminimump  rcCrr)rr5rrr'rrr(create_maximumfrrircreate_maxnumfr&r create_maxsir{ create_maxuir8rrrrmaximumrrminmaxcCsnt|||\}}t|||\}}t|||\}}|j}|r/t||j|j|j||jStd|d)Nrz(. Only floating point clamp is supported) rr5rrr'r( create_clampfrrir8)rjrrrr r5rrrclamps rcCstt|||\}}|jj}|jj}|r|st||t||}||kr,t|||}||kr6t|||}||fSrp)rrirzrOr r=r~)rrr input_sca_ty other_sca_tyrrrrbitwise_op_type_checking_impls    rcC*t|||\}}t||j|j|jSrp)rr'r( create_andrrirrr rrrrrcCrrp)rr'r( create_orrrirrrror_rrcCrrp)rr'r( create_xorrrirrrrxor_rrcCD|jst|td|}|jst|td|}t|||SNr`)riis_int1bitcastr'r5rrrrr logical_and   rcCrr)rirrr'r5rrrrr logical_orrrcCs&|jst|td|}t||Sr)rirrr'r5invert)rr rrrnot_s  rcCrrp)rr'r( create_lshrrrirrrrlshrrrcCrrp)rr'r( create_ashrrrirrrrashrrrcCrrp)rr'r( create_shlrrirrrrshlrrcCs|SrprrrrrplussrcCsJ|jj}|rtd|dt||||}t ||d|S)Nz$wrong type argument to unary minus ()T) rirzrqr&rr'r(get_null_valuerr)rr r_0rrrrs rcCsP|jj}|s |rtd|dt|| ||}t |||S)Nz%wrong type argument to unary invert (r) rirzrqrrr&rr'r(get_all_ones_valuerr)rr r_1rrrrs  rv tl.block_typecCs&|jstjS|jj}ttj|Srp)rirr'r`shaper)rrrrr _bool_likes rcCt|||\}}|jj}|rt||j|jt|S| rB| r4t| |j|jt|St| |j|jt|St d|r)rrirzrrr'r(create_fcmpOGTrrrOrcreate_icmpSGTcreate_icmpUGTr8rrr rPrrr greater_than rcCrr)rrirzrrr'r(create_fcmpOGErrrOrcreate_icmpSGEcreate_icmpUGEr8rrrrrrrcCrr)rrirzrrr'r(create_fcmpOLTrrrOrcreate_icmpSLTcreate_icmpULTr8rrrr less_than)rrcCrr)rrirzrrr'r(create_fcmpOLErrrOrcreate_icmpSLEcreate_icmpULEr8rrrrr8rrcCnt|||\}}|jj}|rt||j|jt|S| r0t| |j|jt|St d|r) rrirzrrr'r(create_fcmpOEQrrrO create_icmpEQr8rrrrequalGrcCrr) rrirzrrr'r(create_fcmpUNErrrO create_icmpNEr8rrrr not_equalSrrstartendcCst|tr t|tstdt|d?}t|d?}|s|r"td||kr*td||}||d@dkr:td|g}ttj|}t||||S)Nz/arange's arguments must be of type tl.constexpr zarange must fit in int32z=arange's end argument must be greater than the start argumentr rz#arange's range must be a power of 2) r^rr&r?r'rr*r(create_make_range)rrr is_start_int64 is_end_int64rangerrrrrarangeds  r!r List[int]r5cCst|tjr|jjdksJdt|||}n(|durtd|dkr,|||}n t |d|j }||}t||}t |||S)Nr zonly accepts size-1 tensorz2dtype must be specified when value is not a tensorrget_) r^r'r(numelrEr~r&rrgetattrnamesplat)rrEr5r get_value_fnrrrrdus   rdrEcCsF|jr Jdt|dkr|St|j|}t||j||S)NzCannot splat a block tensorr) rirlenr'rr5r( create_splatr)rErr rrrrr's  r' dst_shape can_reordercCsRd}|D]}||9}q|jj|krtdt|jj|}t||j|||S)Nr z:reshape() cannot change total number of elements in tensor) rir$r&r'rrzr(create_reshaper)rr+r,r r$srrrrreshapes  r/cCsZdd|jD}||d|jst|||dSt|jj|}t| |j ||S)NcSsg|]}t|qSrr'_constexpr_to_value.0rjrrr zexpand_dims..r )rr ) rinsertrirr'r'rrzr(create_expand_dimsr)rrr r+rrrr expand_dimss   r8cCsX|sJdt|jdksJt|jj|jd|jdg}t||j|j|S)Nz;current implementation of `cat` always may reorder elementsr r) r)rr'rrirzr( create_catr)rtrvr,r ret_typerrrcats "r;abcCst|||\}}|jgk}|rt|d|}t|d|}t|jdtjr*td}nd}|j|g}t|jj|}t | |j |j |}|rQt |dgd|d}|S)NrrFr,r ) rrr8r^r'rgrrirzr( create_joinrr/)r<r=r was_rank_1two new_shaper:rrrrjoins     rDcCspt|jdks Jt|jddksJ|jdd}t|jj|}||j\}}t ||t ||fS)Nrr>r) r)rr'r1rrirz create_splitrr()r<r rCr:outLHSoutRHSrrrsplits  rHdims Tuple[int]cs~tjt|kr tdtdd|Dttt|kr%td|tjj fdd|D}t | j ||S)Nz5permute dims must have the same length as input shapecss|]}t|VqdSrpr0r3drrr szpermute..z?permute dims must be a permutation of 0, 1, ..., n-1, but were csg|]}j|qSrrrKrrrr4r5zpermute..) r)rr&sortedlistr r'rrirzr( create_transr)rrIr r:rrrpermutes "rRc Cs|jst|j|}t||j||S|j}t|t|kr.t d|d|||kr4|St |D]#\}}|||kr[|dkr[t d||d|d|d|d| q8t|jj |}t| |j||S)Nz!Cannot broadcast, rank mismatch: z, r z3Cannot broadcast, the expanded size of the tensor (z ) must match the existing size (z) at non-singleton dimension : ) rirr'rr(r*rrr)r& enumeraterzcreate_broadcast)rrr r src_shapeiitemrrrbroadcast_impl_shapes,  rYc Cs^|j}|j}|r'|s't|j|j}t||j| |}||fS|sH|rHt|j|j}t||j| |}||fS|r+|r+| }| }t |t |krt t |t |D]}t| |jdt|jdg|j }|j}| }qkn/t |t |krt t |t |D]}t| |jdt|jdg|j }|j}| }qt |t |ksJg}t|D]3\} } || } | dkr|| q| dks| | kr|| qtdt| dt| dt| ||krt|j|} t||j|| }||kr+t|j|} t||j|| }||fS)Nrr z?Cannot make_shape_compatible: incompatible dimensions at index rSr)rirr'rrzrr(r*rrr)r r7valuesrTappendr&strrU) rtrvr lhs_tyrhs_ty lhs_shape rhs_shape_ ret_shaperWleftrightrrrrrsl+'      r rounding_mode Optional[str]cCs<|durdS|dkrtjjS|dkrtjjStd|d)NrtnertzzInvalid rounding mode: z0. Supported rounding modes are 'rtne' and 'rtz'.)r ROUNDING_MODERTNERTZr&)rerrr_str_to_rounding_mode,srldst_tycCs|j}|rt|j|j}||kr|S|j}|j}|s%|r+t|||S|j}|j}||krCt dt |dt |t | |j |||S)Nz!Cannot bitcast data-type of size z to data-type of size )rirr'rrzrrqr~primitive_bitwidthr&r\r(create_bitcastrr)rrmr src_ty src_sca_ty dst_sca_tysrc_bitsdst_bitsrrrr6s  rfp_downcast_roundingcCs@|j}|rt|j|j}||kr|S|j}|j}t|}d}|rC|rC|j|jkrC|dur:t j j }n|t j j krBd}n|durUt dt |dt ||s]|rt|jddusiJd|jd||||dS|r||s|r|s|rt||j||||S|r|r|r|stt|tj|||S|o|o|j|jk}|rt||j|||S|o|o|j|jk} | rt||j|||S|r8|r8|j|jks|j|jkr8| o |! } |!r)|j"|} t|#| |j"} t$|| |St|%|j||| |S|&r~|r~|!r]|j"|} t|#| |j"} t$|| |S| rpt|'|j|||St|(|j|||S|r|&r|!s| st|)|j|||St|*|j|||S|+r|r|j} | dkrt|,|j|||S| d krt$t|tj-|t|.d tj-|S|r|+rt|/|j|||S|+r|+rt|0|j|||SJd |d |) NFTz]fp_downcast_rounding should be set only for truncating fp conversions. Source scalar type is z and destination type is convert_custom_typesz0target doesn't provide conversion for this type._builderrr rz cannot cast z to )1rirr'rrzrrlrrrnr rirjr&r\ is_fp8e4b15 codegen_fnsgetrNr(create_fp_to_fprrrLrKrMr~rHcreate_fp_trunc create_fp_extrOr3r4ris_boolr5rrris_standard_floatingcreate_fp_to_sicreate_fp_to_uicreate_ui_to_fpcreate_si_to_fprqcreate_ptr_to_intrbrcreate_int_to_ptrro)rrmr rurprqrruse_custom_rounding truncate_fpext_fp sign_extendtyrbitwidthrrrr~Is            &r~cCs\tjj}|r,|dkrtjj}|S|dkrtjj}|S|dkr$tjj}|Std|d|S)Nz.ca.cgz.cvCache modifier  not supported)r CACHE_MODIFIERrCACGCVr&cache_modifiercacherrr_str_to_load_cache_modifierrcCsptjj}|r6|dkrtjj}|S|dkrtjj}|S|dkr$tjj}|S|dkr.tjj}|Std|d|S)Nz.wbrz.csz.wtrr)r rrWBrCSWTr&rrrr_str_to_store_cache_modifier  rcCsHtjj}|r"|dkrtjj}|S|dkrtjj}|Std|d|S)N evict_last evict_firstzEviction policy r)r EVICTION_POLICYNORMAL EVICT_LAST EVICT_FIRSTr&)eviction_policyevictionrrr_str_to_eviction_policysrcCsDd}|r |dkrtjj}|S|dkrtjj}|Std|d|S)NzeronanzPadding option r)r PADDING_OPTIONPAD_ZEROPAD_NANr&)padding_optionpaddingrrr_str_to_padding_optionsrcCsptjj}|r6|dkrtjj}|S|dkrtjj}|S|dkr$tjj}|S|dkr.tjj}|Std|d|S)Nacquirereleaseacq_relrelaxedMemory semantic r)r MEM_SEMANTICACQUIRE_RELEASEACQUIRERELEASERELAXEDr&) sem_optionsemrrr _str_to_semrrcCs\tjj}|r,|dkrtjj}|S|dkrtjj}|S|dkr$tjj}|Std|d|S)Ngpuctasysrr)r MEM_SYNC_SCOPEGPUCTASYSTEMr&) scope_optionscoperrr _str_to_scoperrcCs|rEt|ds |g}dd|D}|D]}t|tr(d|kr't|ks*JJqt|dks3Jt|tt|ksAJdt|SdS)N__iter__cS"g|] }t|tjr |jn|qSrr^r'rgrEr3elemrrrr4"z0_canonicalize_boundary_check..rz'Duplicate dimension in `boundary_check`r)hasattrr^rr)setrO)boundary_check block_shapedimrrr_canonicalize_boundary_check s ,rc Cs|dus|dur td|jjj} | tjksJd| r(|tjjkr(td|jj} t || }t | |j |||||| S)NK`mask` and `other` arguments cannot be specified for loading block pointers4`tl.int1` should be rewritten in `tl.make_block_ptr`z@Padding option `nan` is not supported for integer block pointers)r&ri element_tyr'r`rOr rrrrr(create_tensor_pointer_loadr) ptrmaskrrrrr is_volatiler elt_tyrmrrr_load_block_pointers rc  Cs|jjstd|jd|dur|durtd|s!|r%td|js@|r5|jr5td|r@|jr@td|jr_|durRt||j|}|dur_t||j|}|jj} | j} | t j k} | r}t j } t | | j } t|| |}|durt|| |}|jr|j} t | | } n| } |durt ||j|||| }nt ||j|j|r|jnd|||| }| rt|t j |}|S)NUnsupported ptr type z in `tl.load`z)`other` cannot be provided without `mask`z`padding_option` or `boundary_check` argument is not supported for loading a tensor ofpointers or loading a scalar. Because the compiler does not know the boundary; please use block pointers (defined by `make_block_ptr`) insteadEMask argument cannot be block type if pointer argument is not a blockzFOther argument cannot be block type if pointer argument is not a block)rirzrqr&rrrYrrr'r`int8 pointer_type address_spacer~rr( create_loadrcreate_masked_load)rrrrrrrrr ptr_tyrrrrmrrrr _load_legacy0sP        rrrOptional[tl.tensor]rrrr\rrrc Cs^t|} t|} t|} |jr#|jjr#t||||| | | || St||||| | | || Srp) rrrrirqrrrr) rrrrrrrrr rrrrrrloadns rdesc_ptrblock_tycCs ||j||}t||Srp)$create_reinterpret_tensor_descriptorrrr'$_experimental_tensor_descriptor_base)rrr rrrrreinterpret_tensor_descriptor~s rcCsht|dkrdS|ddksJd|dd|jd}|d|ks2J|d|d|ddS) NrrzAtensor descriptor block shape must have at least 8 rows, but got rr z2 tensor descriptor block shape must have at least  columns, but got )r)rn)rr5min_colsrrrvalidate_descriptor_blocks rdesc&tl._experimental_tensor_desciptor_basecCst|tjsJt|j|jt|j}t||ks&Jd|dt|t||dd}||j |t |t |}t ||j SNz expected z offsets, but got F require_i64)r^r'rrrr5r)_convert_to_ir_valuescreate_descriptor_loadrrrr(r)roffsetsrrr ndimrjrrrdescriptor_loads $r'tl._experimental_tensor_descriptor_basecCst|tjsJt|j|jt|j}t||ks&Jd|dt||j|jks.Jt||dd}t | |j |j |tj Sr) r^r'rrrr5r)rrr(create_descriptor_storervoid)rrErr rrrrdescriptor_stores $rc CsBt|tjsJ|dksJd|dksJdt|jdks'Jd|j|jddks6Jd|jt|jdksEJd |j|jdd ksTJd |j|j}d |jd }|jd|ksuJd |d|d|jdt|j|jd|jdg}t ||fddd}| |j |j || |} t | |S)Nz#cache modifier is not supported yetz$eviction policy is not supported yetrdescriptor must be 2D, but got rr *descriptor block must have 1 row, but got x offsets must be 1D, but got rz5descriptor gather must have at least 8 rows, but got rzdescriptor gather of  must have at least rFr)r^r'rr)rrr5rnrrcreate_descriptor_gatherrrr() r x_offsetsy_offsetrrr r5rrirjrrrdescriptor_gathers( rcCst|tjsJt|jdksJd|j|jddks&Jd|jt|jdks5Jd|j|jddksDJd|j|j}d |jd}|jd|kseJd |d |d |jdt ||fd dd}| |j |j |j |t dtj S)Nrrrr rrrz6descriptor scatter must have at least 8 rows, but got rzdescriptor scatter of rrFr)r^r'rr)rrshapaer5rnrcreate_descriptor_scatterrr(r)rrErrr r5rrrrdescriptor_scatters"rglobal_addressbox_dimList[tl.tensor] global_dim global_strideelement_stride elem_typeinterleave_layout swizzle_mode fill_modec Csj|r |djtjks Jt| |j|jdd|Ddd|Ddd|Ddd|D|||| tjS)NrcSg|]}|jqSrrr2rrrr4z$tensormap_create..cSrrrr2rrrr4rcSrrrr2rrrr4rcSrrrr2rrrr4r)r5r'rbr(create_tensormap_createrr) rr r r r rrrrrr rrrtensormap_creates       rcCt||jtjSrp)r'r(#create_tensormap_fenceproxy_acquirerr)rr rrrtensormap_fenceproxy_acquirerc Cs|durtd|jj}|jst|||}|js"Jd||jks7Jd|d|jd|jjj|jjksPJd|jjjd|jjd|jjj}|tjks^Jdt||}t |||}t | |j |j |||tj S) Nrz-Value argument must be block type or a scalarz Block shape(z) and value shape(z ) mismatchzBlock element type(z) and value element type(r)r&rirrrrYr'r`rr~r(create_tensor_pointer_storerr) rvalrrrrr rrrrr_store_block_pointers"   2   rc Cs6|jjstd|jd|rtd|js0|jr%td|r0|jr0td|jrKt||j|}|durKt||j|}|jj}|j}|t j krgt j }t ||j }t|||}t|||}|durt ||j|j||t jS|jjstdt ||j|j|j||t jS)Nrz in `tl.store`z`boundary_check` argument is not supported for storing a tensor of pointers or storing a scalar. Because the compiler does not know the boundary; please use block pointers (defined by `make_block_ptr`) insteadzFValue argument cannot be block type if pointer argument is not a blockrz"Mask must have boolean scalar type)rirzrqr&rrrYrrr'r`rrrr~r( create_storerrrcreate_masked_store) rrrrrrr rrrrr _store_legacys2        "r!rc Cslt|}t|}|js|jjrtd|jr,|jjr,t |||||||St |||||||S)N"Cannot store to a constant pointer) rrriis_constrzr&rqrrrr!) rrrrrrr rrrrrstoreCsr$cmprrc CsNt|}t|}|jjj}|jdvrtdt| |j |j |j |||jS)N)rrz9atomic_cas only supports elements with width {16, 32, 64}) rrrirzrrnr&r'r(create_atomic_casr)rr%rrrr rrrr atomic_casYs   "r(op&Tuple[tl.tensor, tl.tensor, tl.tensor]cCsF|jjstd|j|js|jjrtd|jjj}|tjur4|dkr4td|d|tj tj tj tj fvrLtd|dt ||jrk|dur^t||j|}|durkt||j|}t||jjj|}|dur|d}tj }|jr|||j}ttj |j}t||}|||fS)Nz)Pointer argument of store instruction is r"ratomic_z does not support fp16z does not support T)rirzrqr&rr#rr'rFr`rint16rGr\rrYrr~r_r*rr()rrrr)r rmask_irmask_tyrrratom_red_typechecking_implbs.       r/c Cst|||d|\}}}t|}t|}|jj}|rG|r3t| t j j |j |j |j |||jSt| t j j|j |j |j |||jS|tjtjhvrVtd|tgd||}|tjkretjntj}t|||} t|t|d|} |tjkrtjntj} t|| |} t|t| d|} t|||}t|||}t| t j j | j | j t|||j ||| j}t| t j j| j | j t|||j ||| j}t||||}t|||S)Nrz#atomic_max not supported for dtype r]r )r/rrrirzrOrr'r(create_atomic_rmwr ATOMIC_OPMAXrUMAXrHrJr8rdr*rbrrrarcrrrUMINwhererrrrrr sca_tyri_typei_vali_ptrui_typeui_valui_ptrposnegpos_retneg_retrrrr atomic_max}J       rBc Cst|||d|\}}}t|}t|}|jj}|rG|r3t| t j j |j |j |j |||jSt| t j j|j |j |j |||jS|tjtjhvrVtd|tgd||}|tjkretjntj}t|||} t|t|d|} |tjkrtjntj} t|| |} t|t| d|} t|||}t|||}t| t j j | j | j t|||j ||| j}t| t j j| j | j t|||j ||| j}t||||}t|||S)Nrz#atomic_min not supported for dtype r]r )r/rrrirzrOrr'r(r0r r1MINrr4rHrJr8rdr*rbrrrarcrrrr3r5r6rrr atomic_minrCrEc Csjt|||d|\}}}t|}t|}|jj}|rtjjntjj }t | ||j |j |j |||jS)Nr)r/rrrirzrrr r1FADDADDr'r(r0r)rrrrrr r7r)rrr atomic_adds $rHc CNt|||d|\}}}t|}t|}t|tjj|j |j |j |||j S)Nand) r/rrr'r(r0r r1ANDrrirrrrrr rrr atomic_and  rMc CrI)Nor) r/rrr'r(r0r r1ORrrirLrrr atomic_orrNrQc CrI)Nxor) r/rrr'r(r0r r1XORrrirLrrr atomic_xorrNrTc CrI)Nxchg) r/rrr'r(r0r r1XCHGrrirLrrr atomic_xchgs  rWcCsH||jjvsJd|jjd||}|dkrd}ttj|S)Nzinput_precision must be one of . Got TF32X3TF32x3)lowerrallowed_dot_input_precisionsupperr%r INPUT_PRECISION)input_precisionr rrr_str_to_dot_input_precisions r`accr_max_num_imprecise_acc out_dtypec Cs|jr |js J|jr|jrn@|jtjtjtjtjtj fvs.Jd|j|jtjtjtjtjtj fvsEJd|j|j|jksWJd|jd|j|j sa|j rot |tj|}t |tj|}|durw|j j }t||}t|j}t|j}||krdksn||krdksnJd|jd|jd |jd j|jd jksJd |jd |jd|jd jd|jd jd |jddusJd|jd|j|j} |jd j| dkr|jd j| dkr|jd j| dksJd| dd| dd| d|jjr6|jjtjks-Jd|d} tj} n1|r?td|jjsM|jjrV|d} tj } n|r`|dn|d} |} |jjd } |jjd } |jjd }|dkr|jjdnd}t| |r|| | gn| | g}|dur|| |r|| | gn| | g}n |j }|j|ksJ|dur|jr|jr|j j!}nd}n|jr|jr||krtd|d|d t"|#|j |j ||||S)NzUnsupported lhs dtype zUnsupported rhs dtype z&Both operands must be same dtype. Got rr+Both inputs must be either 2D or 3D; (lhs: vs rhs: rr>zFirst input shape (z) and second input shape z= are not compatible for matmul (second index of first shape (z0) must be equal to first index of second shape ( min_dot_sizez2target doesn't provide lower shape bounds for dot.rr zInput shapes should have M >= z, N >= z and K >= zonly int8 supported!zhout_dtype=bfloat16 is unsupported. Please use out_dtype=float32/float16 and cast with `.to(tl.bfloat16)`zmax_num_imprecise_acc (z) must be <= K ()$rirr5rNr'ruint8rFrGrHryr~rdefault_dot_input_precisionr`r)rrErzr{rzrO get_int32r*rMr&rKget_fp32rLget_fp16rr*rmax_num_imprecise_acc_defaultr( create_dot)rtrvrar_rbrcr lhs_rankrhs_rankrhr ret_scalar_tyMNKBr acc_handlerrrdots  $   F0,         "  "rx float_formatcCs.ttj|d}|durtd|d|S)NzInvalid float format: rX)r%r ScaleDotElemTypeTYr]r&)ryty_enumrrr_str_to_fp_typeLsr|cCstjtjtjtjd|}|dur-|dksJd||jtjks+Jd|j|S|j|kr4|Stjtjtjtjd|}|j|ksQJd|d|jt |||S)z If float_format is subbyte, make sure it's packed as uint8 and return it. Otherwise, return a tensor (perhaps bitcasting) of the specified float format. )e5m2e4m3bf16fp16Ne2m1z)Internal Error: Unexpected float format: z)e2m1 format must be packed as uint8. Got zUnexpected dtype for rX) r'float8e5 float8e4nvrGrFr{r5riuint16r)rryr triton_ty unsigned_tyrrr_bitcast_to_fp_typeSs   r lhs_scale lhs_format rhs_scale rhs_formattl.tensor | None fast_mathc  CsJ|jr |js Jt|j} t|j} | | kr dks9n| | kr+dks9nJd|jd|jd|j}|j}t|} t|} hd}||vsVJd|||vsaJd||dupot|tjoo|jdu}|dup~t|tjo~|jdu}t ||| }t ||| }|jjd }|jjd d\}}|d krdnd }|d krdnd }||||jjd ksJd |jd|jd| dkr|jjdnd}t ||r|||gn||g}| d}|dur| ||r|||gn||g}n |j }|j|ksJ|rdn|j }|rdn|j }t| |j || |j || |||S)Nrrdrerfr>rrr~r}rzNYI: lhs_format zNYI: rhs_format rgrr r>zCReduction dimension should pack the same number of elements; (lhs: r)rirr)rrEr|r^r'rgrrrlr*rr(create_dot_scaled)rtrrrvrrrarrcr rprqlhs_format_enumrhs_format_enumallowed_formatsrhs_scale_is_nonelhs_scale_is_nonersrurtPACKED_APACKED_Brvrrrwrhs_scale_handlelhs_scale_handlerrr dot_scaledesN  F      r conditioncCs|jtjkrtd|jt|tj|}t|||dd\}}|jr6t |||\}}t |||\}}nt |||\}}|j}t | |j |j |j |S)Nzgtl.where with a non-boolean condition is deprecated and will error out in a future triton release. Got T) r5r'r`warningswarnr~rrirrr( create_selectr)rrjrr rarrrrr5s   r5cCs"|r t||}n|}t||Srp)r'rr()rjrPrbres_tyrrr wrap_tensors rinputsSequence[tl.tensor]Tuple[tl.tensor, ...]csdurtfddDddjjt}|ks'Jd|dfddtDtfddDsAJd d dD|tfd dttDS) Nc3s&|]}t||jjgddVqdS)Tr?N)r/r$rEr3tr rrrMs$zreduction..rz&reduction axis must be < inputs rank (rcsg|] \}}|kr|qSrr)r3rWr.)rrrr4szreduction..c3s|] }|jjkVqdSrp)rirrrNrrrMsz-all reduction inputs must have the same shapecSrrrrrrrr4rc3*|]}t||jjVqdSrpr get_resultrirzr3rW)r reduce_oprbrrrM() tuplerirr)rTall create_reduceverifyr )rrregion_builder_fnr rankr)rr rrrbrr reductions "rreversecsdjjt}| |kr|ks!nJd|d|d|dkr)||7}D] }|jjks7Jdq+|ddD|||tfdd ttDS) Nrz scan axis z must be < inputs rank (rz(all scan inputs must have the same shapecSrrrrrrrr4rz$associative_scan..c3rrprrrscan_oprrrrMrz#associative_scan..)rirr) create_scanrrr )rrrrr rrrrrassociative_scans ."rsrcindexcCs|js Jdt|jj}t|jj|ksJd| |kr&|ks2nJd|d|d|dkr:||7}t|D]}||krEq>|jj||jj|ksYJd|dq>||j|j|}t||jj |jjS) Nzindex must be an integer tensorz0source and index tensors must have the same rankz gather axis z must be < source rank (rrz index dim z( must match the corresponding source dim) r5rOr)rirr  create_gatherrrrz)rrrr rrLgatherrrrrs . *rnum_binscCsJt|jdks Jd|jsJdt||j|ttj |gS)Nr z histogram only supports 1D inputz%histogram only supports integer input) r)rr5rOr'r(create_histogramrrr*)rrr rrr histograms"rrZcCs@tdt|jt|krtd|jdt||j|S)Nr zAShape of input to multiple_of does not match the length of valuesztt.divisibility) rr)rr&rset_attrr make_attr get_contextrjrZrrr multiple_ofsrcC:t|jt|kr td|jdt||j|S)NzDShape of input to max_contiguous does not match the length of valuesz tt.contiguityr)rr&rrr rrrrrrmax_contiguous rcCr)NzCShape of input to max_constancy does not match the length of valuesz tt.constancyrrrrr max_constancyrrcCst|tjSrp)r'r(create_barrierrrrrr debug_barriersrprefixargshexcCs|ds |r |d7}|ds|r|ddd}t|dkr)|ds)d|}dd|D}dd|D}t|||||tjS)N rSr>rcSrrrr3argrrrr4%rz device_print..cSs*g|]}|jtjtjtjtjtjfvqSr)r5r'r`rr,r*rbrrrrr4&s*)endswithr) startswithr'r( create_printr)rrrr new_args is_signedrrr device_printsrrrcCs$|jjsdSt||j|tjSrp)rdebugr'r( create_assertrr)rrr rrrr*srcCrrp)r'r( create_assumerr)rr rrrassume0rrcCst|tr t|}t|tjrH|r-d|jkrdks'nJd|jd||jSd|jkr8dksBnJd|jd||jSt|tjr|jjdksXJd |j saJd |j tj krv|rv| |j ||j S|j tjkr|sJd |j SJd t|)NrVrWz@Block pointers only support 64 bit `shape/strides`, got a value z which is out of the rangerTrUzFBlock pointers only support 32 bit `offsets/block_shape`, got a value r z*Expected a scalar in shape/strides/offsetsz8Expected an integer scalar type in shape/strides/offsetsFzzBlock pointers only support 32 bit `offsets/block_shape`, add a `.to(tl.int32)` or use regular indexing for 64 bit supportz3Unsupported element type in shape/strides/offsets: )r^rr'rgrErrkr(r$r5rOrbrr get_int64_tyrr*ri)r rrrrr_convert_elem_to_ir_value4s*      rcs,t|drfdd|DSt|gS)Nrcsg|]}t|qSr)rrr rrrr4Nsz)_convert_to_ir_values..)rr)r list_likerrrrrLs rbasec s8t||}t||}t||dd}|jr|jjr td|jjtjkr4t|t tj |jj |}t ds<gddDt ddDsPJd t |dsX|g}d d|D}t|ttt|ksoJd t fd d||||fDsJd ||j||||}t|t t|jjS)NFrzMExpected `base` to be a pointer type (but not a block pointer type or others)rcSrrrrrrrr4drz"make_block_ptr..css2|]}t|tod|kodknVqdS)rTrUN)r^rrrrrrMes0z!make_block_ptr..zGExpected a list of constant integers (`int32_t` range) in `block_shape`cSrrrrrrrr4krzr z-Tensor descriptor last dim must be 1 but got csg|]}t|qSr)rhr2rrrr4r5z*make_tensor_descriptor..cs"g|] }t|jtjdqS)rw)rhtor'rbr2rrrr4rcSrrr)r3r.rrrr4r) r)r&r'r1 _unwrap_shaper^rirrrcreate_make_tensor_descriptorr_experimental_tensor_descriptor)rrrrr rrirrrrmake_tensor_descriptors&     r)rrr r!r"r#)r0r1r2r1r"r1) r0r1r>r?r2r1r@r?rAr?r"r1)T)rSr?)rr1rr1rnr?r"ro)FFTF)rtrurvrur r!r"rw)rtr#rvr#r r!rr) rrurrurr?r r!r"r#)rrurrur r!r"r#) rrurrurr?r r!r"r#)rjr#rr#rrr r!) rjr#rr#rr#rrr r!)rr#rr#r r!r"rw)rr#rr#r r!r"r#)rr#r r!)rr#r"r#)rr#r r!r"r#)rr#r r#r"r#)rr#r"r)rrrrr r!r"r#)rr"r5r1r r!r"r#)rEr#rr"r r!r"r#) rr#r+r"r,r?r r!r"r#)rr#rrr r!r"r#) rtr#rvr#r,r?r r!r"r#)r<r#r=r#r r!r"r#)r<r#r r!r"rw)rr#rIrJr r!r"r#)rr#rr"r r!r"r#)rtr#rvr#r r!r"r#)rerf)rr#rmr1r r!r"r#rp) rr#rmr1r r!rurfr"r#)rr#rrrrrrrr\rr\rr\rr?r r!r"r#)rr#rrr r!) rrrr\rr\r r!r"r#)rrrEr#r r!r"r#)rr\rr\r r!r"r#)rEr#r r!r"r#)rr#r r#r r r r r r rr rrrrrrrrr r!r"r#)rr#r r!r"r#)rr#rr#rrrr\rr\r r!r"r#)rr#r%r#rr#rr\rr\r r!r"r#) rr#rr#rr#r)r\r r!r"r*)rr#rr#rr#rr\rr\r r!r"r#)rtr#rvr#rar#r_rfrbrrcr1r r!r"r#)ryr\)rr#ryr\r r!)rtr#rr#rr\rvr#rrrr\rarrr?rcr1r r!r"r#) rr#rjr#rr#r r!r"r#)rrrrr r!r"r) rrrrrr?r r!r"r) rr#rr#rrr r!r"r#)rr#rrr r!r"r#)rjr#rZr"r"r#)r r!r"r#) rr\rr rr?r r!r"r#)rr#rr\r r!r"r#)rr#r r!r"r#) rr#rr rr rrr r!r"r){ __future__rrtypingrrrrrrx _C.libtritonr rr r'r Exceptionr r,r/r=rRrhrsrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrr!rdr'r/r8r;rDrHrRrYrrlrr~rrrrrrrrrrrrrrrrrrrr!r$r(r/rBrErHrMrQrTrWr`rxr|rrr5rrrrrrrrrrrrrrrrrrrrrs      3 +  &  #                                : p    >     ,   ' '   I0   '