o ۗi7@s4ddlmZddlZddlmZmZmZmZmZddl Z ddl m Z ddl m Zddl mZed ZGd d d eZd6ddZd6ddZd7ddZd8dd Zd9d:d#d$Zd;d)d*Z ! +d<d=d0d1Zd>d4d5Zd?d9d:Zd?d;d<Zd?d=d>Zd@d?d@Zd@dAdBZdAdDdEZ d@dFdGZ!dBdLdMZ"dBdNdOZ#dCdRdSZ$dDdTdUZ%dEdVdWZ&dEdXdYZ'dEdZd[Z(dEd\d]Z)dEd^d_Z*dFd`daZ+dEdbdcZ,dEdddeZ-dEdfdgZ.dGdhdiZ/dHdjdkZ0dIdldmZ1dJdpdqZ2dEdrdsZ3dEdtduZ4dEdvdwZ5dEdxdyZ6dEdzd{Z7dEd|d}Z8dKddZ9dLddZ:dMddZ;dNddZ<dOddZ=dPddZ>dQddZ?dRddZ@dSddZAdTddZBdUddZCdVddZDdWddZE dXdYddZFddZGddZHddZIddZJddZKddZLddZMddZNddZOdZddDŽZPd[ddʄZQd\dd̄ZRd]dd؄ZSd^ddڄZTdd܄ZUddބZVd_ddZWd`ddZXdaddZYdbddZZdbddZ[dbddZ\dbddZ]dbddZ^dbddZ_dbddZ`ddZadcddZbddddZcdeddZddfd d Zed d ZfdgddZgdhddZhdiddZidjddZjdjddZkdjddZldkd d!Zmdld%d&Zndmd)d*Zodkd+d,Zpd-d.Zqd9d/d0Zrdnd2d3Zsdnd4d5ZtdS(o) annotationsN)ListOptionalSequenceTupleTypeVar)ir)core)mathTcseZdZfddZZS)IncompatibleTypeErrorImplcs@||_||_d|jd|j|_tt||jdS)Nzinvalid operands of type  and )type_atype_b__repr__messagesuperr__init__)selfrr __class__V/home/ubuntu/transcripts/venv/lib/python3.10/site-packages/triton/language/semantic.pyrsz"IncompatibleTypeErrorImpl.__init__)__name__ __module__ __qualname__r __classcell__rrrrrsraxisintbuilder 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.r0a_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)r1r3a_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| sW| rj|r\tjS| rg| rgtjStjS| 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_intr9r5rr>)r1r?r3rArB scalar_ty tensor_tyrrrcomputation_type_impl=s: rST 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 .rr6r!g8g?rabsinfzcannot convert z of type z to tensor) isinstancer@r(r)get_int1int1r r+uint32int64uint64r'fullfloat __builtins__rIrK constexpr to_tensorrFr9type)xr!rTr6 min_float32 max_float32abs_xrrrrios>         rirr allow_ptr_aNonecCsJ|r!|s t|||r||krt|||r#t||dSdSN)is_ptrr is_floating)rrrorrrcheck_ptr_type_impls   rtFlhstl.tensor | numbers.NumberrhsTuple[tl.tensor, tl.tensor]cCst|tj}t|tj}|r|} t||}|r|} t||}|jj} |jj} t| | |t| | ||r|| s|| s|t| || ||} |rN| dkrN| sX|r\| dkr\| r\t d|rft d| | |dnt || |}|rvt d| | |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.rrZ)r_numbersNumberrirjscalarrtrrrSis_int_unsignedr'recastbroadcast_impl_value)rurwr! allow_lhs_ptr allow_rhs_ptrarithmetic_checkrB lhs_is_scalar rhs_is_scalar lhs_scalar rhs_scalar lhs_sca_ty rhs_sca_ty ret_sca_tyrrrbinary_op_type_checking_implsB        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@Fr z! overflow detected for operation )rjr{r4optionssanitize_overflowrPr}r(rcget_int_max_valuer) get_int64get_int_min_valueand_ less_equal greater_equalr device_assert) rurwr!rrrret max_value min_valuecondmsgrrr binary_op_sanitize_overflow_impls   rinputotherrcCst|||dd\}}|jj}|jj}|r|rtd|r3|s3||}}|jj}|jj}|rDt||j|j|jS| rUt| |j|j|jS| ro|rbt |||t t||j|j|jStd|)NTzcannot add pointers togetherrC)rrjr{rrr9r(r) create_addptrhandlers create_faddrPradd create_add)rrrr!input_scalar_tyother_scalar_tyrrrrs$ rcCst|||dd\}}|jj}|r"t||jt||j|jS| r3t| |j|j|jS| rM|r@t |||t t||j|j|jStd|)NTFrC)rrjr{rrr(r)rrminusrs create_fsubrPrsub create_subr9rrrr!rQrrrrs rcCs|t|||\}}|jj}|rt||j|j|jS|r7|r*t |||t t| |j|j|jSt d|NrC) rrjr{rsr(r) create_fmulrrPrmul create_mulr9rrrrrsrcCst|||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 |jSNFTrC) rrjr{rsrPr}r(rIfp_mantissa_widthr9r) create_fdivr)rrr!rrrrrtruedivs  rcCst|||dddd\}}|jj}|jj}|rK|rKt||}t|||}t|||}|r>t| |j |j |jSt| |j |j |jSt d|r) rrjr{rPr>r} is_int_signedr(r) create_sdivr create_udivr9)rrr!rrret_tyrrrfloordiv0s   r ieee_roundingcCs^|jj}|jj}|r|stdt|||dddd\}}||j|j}t||jS)Nz4both operands of fdiv must have floating scalar typeFT) rjr{rsr9rrrr(r))rrrr!rrrrrrfdiv?srcCst|||dddd\}}|jj}|jj}|r2tjt||d||d}t|t||d|d|}|S| rj|j |j krLt d| d| d| r]t||j|j|jSt||j|j|jSt d|)NFT_builderz Cannot mod z by rDrC)rrjr{rsr floorrrrrPr5r9rrr(r) create_sremr create_urem)rrr!rQrrrrrrmodJs  rrky 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 )rr6rsr( PropagateNanALLr)create_minimumfrrjNONEcreate_minnumfr'r create_minsir| create_minuir9rkrrr!r6rrrminimumf  rcCrr)rr6rsr(rrr)create_maximumfrrjrcreate_maxnumfr'r create_maxsir| create_maxuir9rrrrmaximumxrrminmaxcCsnt|||\}}t|||\}}t|||\}}|j}|r/t||j|j|j||jStd|d)Nrz(. Only floating point clamp is supported) rr6rsr(r) create_clampfrrjr9)rkrrrr!r6rrrclamps rcCstt|||\}}|jj}|jj}|r|st||t||}||kr,t|||}||kr6t|||}||fSrq)rrjr{rPrr>r})rrr! input_sca_ty other_sca_tyrrrrbitwise_op_type_checking_impls    rcC*t|||\}}t||j|j|jSrq)rr(r) create_andrrjrrr!rrrrrcCrrq)rr(r) create_orrrjrrrror_rrcCrrq)rr(r) create_xorrrjrrrrxor_rrcCD|jst|td|}|jst|td|}t|||SNra)rjis_int1bitcastr(r6rrrrr logical_and   rcCrr)rjrrr(r6rrrrr logical_orrrcCs&|jst|td|}t||Sr)rjrrr(r6invert)rr!rrrnot_s  rcCrrq)rr(r) create_lshrrrjrrrrlshrrrcCrrq)rr(r) create_ashrrrjrrrrashrrrcCrrq)rr(r) create_shlrrjrrrrshlrrcCs|SrqrrrrrplussrcCsJ|jj}|rtd|dt||||}t ||d|S)Nz$wrong type argument to unary minus ()T) rjr{rrr'rr(r)get_null_valueto_irr)rr!r_0rrrrs rcCsP|jj}|s |rtd|dt|| ||}t |||S)Nz%wrong type argument to unary invert (r) rjr{rrrsr'rr(r)get_all_ones_valuerr)rr!r_1rrrrs  rv tl.block_typecCs&|jstjS|jj}ttj|Srq)rjis_blockr(rashape block_type)rrrrr _bool_likes rcCt|||\}}|jj}|rt||j|jt|S| rB| r4t| |j|jt|St| |j|jt|St d|r)rrjr{rsr(r)create_fcmpOGTrrrPrcreate_icmpSGTcreate_icmpUGTr9rrr!rQrrr greater_thanrcCrr)rrjr{rsr(r)create_fcmpOGErrrPrcreate_icmpSGEcreate_icmpUGEr9rrrrrrrcCrr)rrjr{rsr(r)create_fcmpOLTrrrPrcreate_icmpSLTcreate_icmpULTr9rrrr less_thanrr cCrr)rrjr{rsr(r)create_fcmpOLErrrPrcreate_icmpSLEcreate_icmpULEr9rrrrr.rrcCnt|||\}}|jj}|rt||j|jt|S| r0t| |j|jt|St d|r) rrjr{rsr(r)create_fcmpOEQrrrP create_icmpEQr9rrrrequal=rcCrr) rrjr{rsr(r)create_fcmpUNErrrP create_icmpNEr9rrrr not_equalIrrstartendcCst|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_r r'r@r(rr+r)create_make_range)rrr!is_start_int64 is_end_int64rangerrrrrarangeZs  r r List[int]r6cCst|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)numelrFr}r'rrgetattrnamesplat)rrFr6r! get_value_fnrrrreks   rerFcCsF|jr Jdt|dkr|St|j|}t||j||S)NzCannot splat a block tensorr) rjrlenr(rr6r) create_splatr)rFrr!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) rjr#r'r(rr{r)create_reshaper)rr*r+r!r#srrrrreshapes  r.cCsZdd|jD}||d|jst|||dSt|jj|}t| |j ||S)NcSsg|]}t|qSrr(_constexpr_to_value.0rkrrr zexpand_dims..r )rr!) rinsertrjrr&r(rr{r)create_expand_dimsr)rrr!r*rrrr expand_dimss   r7cCsX|sJdt|jdksJt|jj|jd|jdg}t||j|j|S)Nz;current implementation of `cat` always may reorder elementsr r) r(rr(rrjr{r) create_catr)rurwr+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!) r~rr7r_r(rhrrjr{r) create_joinrr.)r;r<r! was_rank_1two new_shaper9rrrrjoins     rCcCspt|jdks Jt|jddksJ|jdd}t|jj|}||j\}}t ||t ||fS)Nrr=r) r(rr(r0rrjr{ create_splitrr))r;r!rBr9outLHSoutRHSrrrsplits  rGdims 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|VqdSrqr/r2drrr szpermute..z?permute dims must be a permutation of 0, 1, ..., n-1, but were csg|]}j|qSrrrJrrrr3r4zpermute..) r(rr'sortedlistrr(rrjr{r) create_transr)rrHr!r9rrrpermutes "rQc 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 : ) rjrr(rr)r)rget_block_shapesr(r' enumerater{create_broadcast)rrr!r src_shapeiitemrrrbroadcast_impl_shapes,  rYc CsZ|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}| }qkn.t |t |krt t |t |D]}t| |jdt|jdg|}|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 rRr)rjrr(rr{rr)r)rrSr(rr6rTappendr'strrU) rurwr!lhs_tyrhs_ty lhs_shape rhs_shape_ ret_shaperWleftrightrrrrr~sl+'      r~ rounding_mode Optional[str]cCs<|durdS|dkrtjjS|dkrtjjStd|d)NrtnertzzInvalid rounding mode: z0. Supported rounding modes are 'rtne' and 'rtz'.)r ROUNDING_MODERTNERTZr')rdrrr_str_to_rounding_mode"srkdst_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 )rjrr(rr{rSrrr}primitive_bitwidthr'r[r)create_bitcastrr)rrlr!src_ty src_sca_ty dst_sca_tysrc_bitsdst_bitsrrrr,s  rfp_downcast_roundingcCsd|j}t|tjr |j}t|tjr|j}|r#t|j|j}||kr)|S|j}|j}t |}d}| rU| rU|j |j krU|durLt j j}n|t j jkrTd}n|durgtdt|dt||so|r|jddus{Jd|jd||||dS|r| s| r|s|rt||j||||S|r|r|r|stt|tj|||S| o| o|j |j k}|rt||j|||S| o| o|j |j k} | rt||j|||S| rJ| rJ|j!|j!ks|j"|j"krJ|#o|$ } |$r;|j%|} t|&| |j%} t'|| |St|(|j||| |S|)r| r|$ro|j%|} t|&| |j%} t'|| |S|#rt|*|j|||St|+|j|||S| r|)r|$s|#st|,|j|||St|-|j|||S|.r| r|j!} | dkrt|/|j|||S| d krt't|tj0|t|1d tj0|S| r|.rt|2|j|||S|.r(|.r(t|3|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.rrr rz cannot cast z to )4rjr_r(rhrFrrr{rSrkrsrmr rhrir'r[ is_fp8e4b15 codegen_fnsgetrOr)create_fp_to_fprrrMrLrNr}rIcreate_fp_trunc create_fp_extrPr4r5ris_boolr6rrcreate_int_castis_standard_floatingcreate_fp_to_sicreate_fp_to_uicreate_ui_to_fpcreate_si_to_fprrcreate_ptr_to_intrcrcreate_int_to_ptrrn)rrlr!rtrorprquse_custom_rounding truncate_fpext_fp sign_extendtyrbitwidthrrrr}?s              &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(rhrFr2elemrrrr3 "z0_canonicalize_boundary_check..rz'Duplicate dimension in `boundary_check`r)hasattrr_r r(setrN)boundary_check block_shapedimrrr_canonicalize_boundary_checks ,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 pointers3`tl.int1` should be rewrited in `tl.make_block_ptr`z@Padding option `nan` is not supported for integer block pointers)r'rj element_tyr(rarPr rrrrSr)create_tensor_pointer_loadr) ptrmaskrrrrr is_volatiler!elt_tyrlrrr_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)rjr{rrr'rrrYrSrr(raint8 pointer_type address_spacer}rr) create_loadrcreate_masked_load)rrrrrrrrr!ptr_tyrr|rrlrrrr _load_legacy*sP        rrrOptional[tl.tensor]rrrr[rrrc Cs^t|} t|} t|} |jr#|jjr#t||||| | | || St||||| | | || Srq) rrrrjrrrrrr) rrrrrrrrr!rrrrrrloadhs rdesc_ptrcCs<t||dd}||j|||t|t|}t||SNF require_i64)_convert_to_ir_valuescreate_descriptor_loadrrrrr(r))roffsetsrrrjr!rkrrrdescriptor_loadxs  rcCs*t||dd}t||j|j|tjSr)rr(r)create_descriptor_storervoid)rrFrr!rrrdescriptor_storesrglobal_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|]}|jqSrrr1rrrr3z$tensormap_create..cSrrrr1rrrr3rcSrrrr1rrrr3rcSrrrr1rrrr3r)r6r(rcr)create_tensormap_createrr) rrrrrrrrrrr!rrrtensormap_creates       rcCt||jtjSrq)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'rjrrSrrYr(rarr}r)create_tensor_pointer_storerr) rvalrrrrr!rrrrr_store_block_pointers"   2   rc Cs2|jjstd|jd|rtd|js0|jr%td|r0|jr0td|jrKt||j|}|durKt||j|}|jj}|j}|t j krgt j }t ||j }t|||}t|||}|s~t ||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)rjr{rrr'rrrYrSrr(rarrrr}r) create_storerrr|create_masked_store) rrrrrrr!rrrrr _store_legacys2        "rrc Cslt|}t|}|js|jjrtd|jr,|jjr,t |||||||St |||||||S)N"Cannot store to a constant pointer) rrrjis_constr{r'rrrrrr) rrrrrrr!rrrrrstoresr cmprrc CsNt|}t|}|jjj}|jdvrtdt| |j |j |j |||jS)N)rrz9atomic_cas only supports elements with width {16, 32, 64}) rrrjr{rrmr'r(r)create_atomic_casr)rr rrrr!rrrr atomic_cass   "rop&Tuple[tl.tensor, tl.tensor, tl.tensor]cCsB|jjstd|j|js|jjrtd|jjj}|tjur4|dkr4td|d|tj tj tj tj fvrLtd|dt ||jrk|dur^t||j|}|durkt||j|}t||jjj|}|s|d}tj }|jr|||j}ttj |j}t||}|||fS)Nz)Pointer argument of store instruction is rratomic_z does not support fp16z does not support T)rjr{rrr'rr rr(rGrarint16rHr[rrYrSr}r`r)rr))rrrrr!rmask_irmask_tyrrratom_red_typechecking_impls.       rc 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 )rrrrjr{rPrr(r)create_atomic_rmwr ATOMIC_OPMAXrUMAXrIrKr9rer+rcrrrbrdrr rUMINwhererrrrrr!sca_tyri_typei_vali_ptrui_typeui_valui_ptrposnegpos_retneg_retrrrr atomic_max,J       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_min not supported for dtype r^r )rrrrjr{rPrr(r)rr rMINrrrIrKr9rer+rcrrrbrdrr rrrrrrr atomic_minSr)r+c Csjt|||d|\}}}t|}t|}|jj}|rtjjntjj }t | ||j |j |j |||jS)Nr)rrrrjr{rsr rFADDADDr(r)rr)rrrrrr!rrrrr atomic_addzs $r.c CNt|||d|\}}}t|}t|}t|tjj|j |j |j |||j S)Nand) rrrr(r)rr rANDrrjrrrrrr!rrr atomic_and  r3c Cr/)Nor) rrrr(r)rr rORrrjr2rrr atomic_orr4r7c Cr/)Nxor) rrrr(r)rr rXORrrjr2rrr atomic_xorr4r:c Cr/)Nxchg) rrrr(r)rr rXCHGrrjr2rrr atomic_xchgs  r=cCsH||jjvsJd|jjd||}|dkrd}ttj|S)Nzinput_precision must be one of z. Got TF32X3TF32x3)lowerrallowed_dot_input_precisionsupperr$r INPUT_PRECISION)input_precisionr!rrr_str_to_dot_input_precisions rEaccrDmax_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 ()$rjrr6rOr(ruint8rGrHrIrvr}rdefault_dot_input_precisionrEr(rrFrwrxr{rP get_int32r+rNr'rLget_fp32rMget_fp16rr)rmax_num_imprecise_acc_defaultr) create_dot)rurwrFrDrGrHr!lhs_rankrhs_rankrMr ret_scalar_tyMNKBr acc_handlerrrdots  $   F0,         "  "r] float_formatcCs`|dkrtjjS|dkrtjjS|dkrtjjS|dkr tjjS|dkr(tjjStd|d)Ne4m3e5m2e2m3e3m2e2m1zInvalid float format: rY)r F8F6F4TYE4M3E5M2E2M3E3M2E2M1r')r^rrr_str_to_fp_typesrj lhs_scale rhs_scaletl.tensor | Nonec  Cs|jr |js Jt|j} t|j} | | kr dks9n| | kr+dks9nJd|jd|jdt|} t|} |dvsLJd||dvsWJd |t|tjoa|jdu} | shJd |jjd }|jjd d\}}|d kr~dnd }|||jjdksJd|jd|jd|dksJd|| dkr|jjdnd}t ||r|||gn||g}| d}|dur| ||r|||gn||g}n |j }|j|ksJ| rdn|j }t ||j |j | |j || ||S)NrrIrJrKr)rcr_r`zNYI: lhs_format )r_r`zNYI: rhs_format zNYI: rhs_scale not supportedrLrcr r=zCReduction dimension should pack the same number of elements; (lhs: rz!scaled_dot NYI for K < 64. Got K=r)rjrr(rrjr_r(rhrFrrQr)rr)create_dot_scaled)rurk lhs_formatrwrl rhs_formatrFrHr!rUrVlhs_format_enumrhs_format_enumrhs_scale_is_nonerXrZrYPACKEDr[rrr\rhs_scale_handlerrr dot_scaledsB  F      rv 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) r6r(rawarningswarnr}rrjrr~r) create_selectr)rwrkrr!r`rrrrr0s   rcCs"|r t||}n|}t||Srq)r(rr))rkrQrares_tyrrr wrap_tensorFs r|inputsSequence[tl.tensor]Tuple[tl.tensor, ...]csdurtfddDddjjt}|ks'Jd|dfddtDtfddDsAJd d dD|tfd dttDS) Nc3s&|]}t||jjgddVqdS)Tr>N)r.r#rFr2tr!rrrLQs$zreduction..rz&reduction axis must be < inputs rank (rcsg|] \}}|kr|qSrr)r2rWr-)rrrr3Wszreduction..c3s|] }|jjkVqdSrq)rjrrrMrrrLXsz-all reduction inputs must have the same shapecSrrrrrrrr3Zrc3*|]}t||jjVqdSrqr| get_resultrjr{r2rW)r} reduce_oprarrrL^() tuplerjrr(rTall create_reduceverifyr)r}rregion_builder_fnr!rankr)rr!r}rrarr reductionOs "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 shapecSrrrrrrrr3srz$associative_scan..c3rrqrrr}scan_oprrrrLwrz#associative_scan..)rjrr( create_scanrrr)r}rrrr!rrrrrassociative_scanfs ."rnum_binscCsJt|jdks Jd|jsJdt||j|ttj |fS)Nr z histogram only supports 1D inputz%histogram only supports integer input) r(rr6rPr(r)create_histogramrrr+)rrr!rrr histograms"rvaluescCs@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_contextrkrrrr 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_contiguousrcCr)NzCShape of input to max_constancy does not match the length of valuesz tt.constancyrrrrr max_constancyrrcCst|tjSrq)r(r)create_barrierrrrrr debug_barriersrprefixargshexcCs|ds |r |d7}|ds|r|ddd}t|dkr)|ds)d|}dd|D}dd|D}t|||||tjS)N rRr=rcSrrrr2argrrrr3rz device_print..cSs*g|]}|jtjtjtjtjtjfvqSr)r6r(rarrr+rcrrrrr3s*)endswithr( startswithr(r) create_printr)rrrr!new_args is_signedrrr device_printsrrrcCs$|jjsdSt||j|tjSrq)rdebugr(r) create_assertrr)rrr!rrrrsrcCrrq)r(r) create_assumerr)rr!rrrassumerrcCst|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|)NrWrXz@Block pointers only support 64 bit `shape/strides`, got a value z which is out of the rangerUrVzFBlock 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_r r(rhrFrrPr)r#r6rPrcr}r get_int64_tyrr+rj)r!rrrrr_convert_elem_to_ir_values*      rcs,t|drfdd|DSt|gS)Nrcsg|]}t|qSr)rrr!rrrr3sz)_convert_to_ir_values..)rr)r! list_likerrrrrs 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)rcSrrrrrrrr3rz"make_block_ptr..css2|]}t|tod|kodknVqdS)rUrVN)r_r rrrrrLs0z!make_block_ptr..zGExpected a list of constant integers (`int32_t` range) in `block_shape`cSrrrrrrrr3rzrSrirtrrrrrrrrrrrrrrrrrrrrrrrrrrrrr rrrr rer&r.r7r:rCrGrQrYr~rkrr}rrrrrrrrrrrrrrrrr rrr(r+r.r3r7r:r=rEr]rjrvrr|rrrrrrrrrrrrrrrrrrs       2 +  !                                  : t    >    ,   ' '    H(    '