Ϧi7SSKJr SSKrSSKJrJrJrJrJr SSK r SSK J r SSK J r SSK Jr \"S 5r"S S \5rSnS jrSnS jrSoSjrSpSjrSqSrSjjrSsSjrStSuSjjrSvSjrSwSjrSwSjrSwSjrSxSjrSxSjrSySjr SxSjr!SzSjr"SzSjr#S{Sjr$S|Sjr%S}Sjr&S}S jr'S}S!jr(S}S"jr)S}S#jr*S~S$jr+S}S%jr,S}S&jr-S}S'jr.SS(jr/SS)jr0SS*jr1SS+jr2S}S,jr3S}S-jr4S}S.jr5S}S/jr6S}S0jr7S}S1jr8SS2jr9SS3jr:SS4jr;SS5jrSS8jr?SS9jr@SS:jrASS;jrBSS<jrCSS=jrDSS>jrESSS?jjrFS@rGSArHSBrISCrJSDrKSErLSFrMSGrNSHrOSSIjrPSSJjrQSSKjrRSSLjrSSSMjrTSNrUSOrVSSPjrWSSQjrXSSRjrYSSSjrZSSTjr[SSUjr\SSVjr]SSWjr^SSXjr_SSYjr`SZraSS[jrbSS\jrcSS]jrdSS^jreS_rfSS`jrgSSajrhSSbjriSScjrjSSdjrkSSejrlSSfjrmSSgjrnSShjroSSijrpSjrqSqSkjrrSSljrsSSmjrtg)) annotationsN)ListOptionalSequenceTupleTypeVar)ir)core)mathTc(^\rSrSrU4SjrSrU=r$)IncompatibleTypeErrorImplc>XlX lSURR5-S-URR5-Ul[[ U]UR5 g)Nzinvalid operands of type  and )type_atype_b__repr__messagesuperr__init__)selfrr __class__s W/var/www/html/ai-image-ml/venv/lib/python3.13/site-packages/triton/language/semantic.pyr"IncompatibleTypeErrorImpl.__init__sT  2T[[5I5I5KKgUX\XcXcXlXlXnn  '7 E)rrr)__name__ __module__ __qualname____firstlineno__r__static_attributes__ __classcell__)rs@rrrsFFrrcUS;a[SU35e[R"URU5[R5$)Nrr r z+program_id axis must be 0, 1, or 2 but got ) ValueErrortltensorcreate_get_program_idint32axisbuilders r program_idr/s= 9FtfMNN 99W2248"(( CCrcUS;a[SU35e[R"URU5[R5$)Nr&z-num_programs axis must be 0, 1, or 2 but got )r'r(r)create_get_num_programsr+r,s r num_programsr2"s= 9HOPP 99W44T:BHH EErcdURnURnURnURnXE:Xa X#:aU$U$U[RRR :Xa X#:aU$U$U[RRR :Xa X2:aU$U$[ SUSU35e)Nzunexpected signedness r) int_bitwidthint_signednessr(dtype SIGNEDNESSUNSIGNED TypeError)a_tyb_tya_rankb_ranka_snb_sns rinteger_promote_implr@-s   F   F   D   D |t0D0 $$-- -'t1T1 $$-- -'t1T1 ,TF%v> ??rcfX:waU(aX4OX 4upVUR5RUR5R::a=U(a4U[R[R4;a[R $U$UR 5(dUR 5(a[R$UR5(dUR5(a[R $UR5(dUR5(a'U(a[R $[R$UR5(dUR5(aaU(a[R $UR5(a%UR5(a[R$[R $UR5(a,UR5(aX:XaU$[R$UR5(aUR5(d[SUSU35eU(aMURUR:wa3[SUR5-S-UR5-S-5e[!X5$)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@)r: a_is_scalarr; b_is_scalar div_or_mod scalar_ty tensor_tys rcomputation_type_implrU=s !/:|  >>  ! !Y^^%5%;%; ;yRZZ,EEzz!  ||~~zz ||~~zz ||~~ :: ::  ||~~ ::  <<>>dllnn;; zz {{}}|t33 ;;== *4&dV<==d))T-@-@@5 G'QTXTaTaTcckkl l  ++rc[U[5(a4[R"UR U5[R 5$[U[ 5(aSUs=::aS:aO O[RnOrSUs=::aS:aO O[RnOQSUs=::aS:aO O[RnO0SUs=::aS:aO O[RnO[SUS35e[S XUS 9$[U[5(anS nS S S--n[S"U5nU[S5:XdUS:XdX:wdXFs=::aU::aO O[RnO[R n[S XUS 9$[U[R"5(a[%UR&U5$[U[R5(aU$U(a[)SUS[+U5S35eU$)NllzNonrepresentable integer .r6r.g8g?r absinfzcannot convert z of type z to tensor) isinstanceboolr(r)get_int1int1intr+uint32int64uint64r'fullfloat __builtins__rHrJ constexpr to_tensorrEr9type)xr. check_typer6 min_float32 max_float32abs_xs rrnrnos!Tyy))!,bgg66 As   Q  HHE a % IIE q 5 HHE a % IIE81=> >B88 Au   !QV+ U#A& E%L C< 6  .; .JJEJJEB88 Ar|| $ $'** Aryy ! !/!Id1gYjIJJ HrcUR5(aXU(d [X5eUR5(aX:wa [X5eUR5(a [X5eggN)is_ptrr is_floating)rr allow_ptr_as rcheck_ptr_type_implrzsZ }}+F; ; ==?? 0+F; ;     +F; ; rc[U[R5n[U[R5nU(a Un [X5nU(a Un [X5nURR n URR n [ XU5 [ XU5 U(aU R5(dU R5(d[XXU5n U(aW S:aU R5(d"U(a&W S:a U R5(a [S5eU(a [SW XS9O [X U5nU(a [SW XS9O [XU5n[XU5upX4$)Nrz{Cannot perform a binary operation between an unsigned tensor and a negative scalar. Perform a explicit cast on one of them.r\r])rbnumbersNumberrnroscalarrzrwrUis_int_unsignedr'rjcastbroadcast_impl_value)lhsrhsr. allow_lhs_ptr allow_rhs_ptrarithmetic_checkrR lhs_is_scalar rhs_is_scalar lhs_scalar rhs_scalar lhs_sca_ty rhs_sca_ty ret_sca_tys rbinary_op_type_checking_implrsRsGNN3MsGNN3M % %JJ > > 1 1 3 3J(>(@(@u**++**++yy..u||U\\JEJJWW  $ $ & &yy,,U\\5<<H%**UU    ! !  ,U7C Hyy++ELL%,,GTT &&78 99rc [XUSS5upURRnUR5(aN[R "UR UR[X5R5UR5$UR5(aE[R "URURUR5UR5$UR5(a]U(a[XU[5 [R "URURUR5UR5$[SU35e)NTFrB)rror~rwr(r)rrminusrx create_fsubrOrsub create_subr9rrrr.rSs rrrs/gtUSLE !!Iyy..u||U5=R=Y=YZ\a\f\fggyy,,U\\5<<H%**UU       ,U7C Hyy++ELL%,,GTT &yk2 33rc[XU5upURRnUR5(aE[R "UR URUR5UR5$UR5(a]U(a[XU[5 [R "URURUR5UR5$[SU35eNrB) rror~rxr(r) create_fmulrrOrmul create_mulr9rs rrrs/gFLE !!Iyy,,U\\5<<H%**UU       ,U7C Hyy++ELL%,,GTT &yk2 33rc [XUSSSS5upURRnURRnUR5(a#UR 5(a[ XU5nGOUR 5(a"UR5(a [ XU5nOUR 5(aLUR 5(a7[ U[ RU5n[ U[ RU5nOlUR5(aIUR5(a4URUR:a [ XU5nO[ XU5nO[SU35e[ R"URURUR5UR5$NFTrB) rror~rxrOrr(rHfp_mantissa_widthr9r) create_fdivr)rrr.rrs rtruedivrsT/gueUY[_`LEjj''Ojj''O""$$)?)?)A)AUW5    ! !o&A&A&C&CUW5    ! !o&<&<&>&>UBJJ0UBJJ0  $ $ & &?+F+F+H+H  , ,/P/P P9E9E*?*;<== 99W((u||Dejj QQrc r[XUSSSS5upURRnURRnUR5(aUR5(a[ X45n[ XU5n[ XU5nUR 5(aE[R"URURUR5UR5$[R"URURUR5UR5$[SU35er) rror~rOr@r is_int_signedr(r) create_sdivr create_udivr9)rrr.rrret_tys rfloordivr0s/gueUY[_`LEjj''Ojj''OO$:$:$<$<%oGUG,UG,    ! !99W00u||LejjY Y99W00u||LejjY Y &&78 99rc tURRnURRnUR5(aUR5(d [S5e[ XUSSSS5upUR UR UR 5n[R"X`R5$)Nz4both operands of fdiv must have floating scalar typeFT) ror~rxr9rrrr(r))rr ieee_roundingr.rrrs rfdivr?sjj''Ojj''O  & & ( (0K0K0M0MNOO/gueUZ\`aLE   ellELL 9C 99S** %%rc <[XUSSSS5upURRnURRnUR5(a;[R "[ XSU5US9n[U[XQSU5SU5nU$UR5(aURUR:wa3[SUR5-S-UR5-S-5eUR5(aE[R"UR!UR"UR"5UR5$[R"UR%UR"UR"5UR5$[SU35e)NFT_builderz Cannot mod z by rCrB)rror~rxr floorrrrrOr5r9rrr(r) create_sremr create_urem)rrr.rSrrrs rmodrJsX/gueUY[_`LE !!Ijj''O 4eW=P%U494I       # #'E'E EMI,>,>,@@6IOLdLdLffjoop p  " " $ $99W00u||LejjY Y99W00u||LejjY Y &yk2 33rc[XU5upURnUR5(aU[RR :XaE[R "URURUR5UR5$U[RR:XaE[R "URURUR5UR5$[SU35eUR5(aE[R "URURUR5UR5$UR5(aE[R "UR!URUR5UR5$[#SU35eNzUnexpected propagate_nan Unexpected dtype )rr6rxr( PropagateNanALLr)create_minimumfrroNONEcreate_minnumfr'r create_minsir create_minuir9rpy propagate_nanr.r6s rminimumrf8 'g 6DA GGE  BOO// /99W44QXXqxxH!&&Q Q boo22 299W33AHHahhGP P8HI I     yy--ahhA166JJ    yy--ahhA166JJ+E7344rc[XU5upURnUR5(aU[RR :XaE[R "URURUR5UR5$U[RR:XaE[R "URURUR5UR5$[SU35eUR5(aE[R "URURUR5UR5$UR5(aE[R "UR!URUR5UR5$[#SU35er)rr6rxr(rrr)create_maximumfrrorcreate_maxnumfr'r create_maxsir create_maxuir9rs rmaximumrxrrcX[XU5up[XU5up[XU5upURnUR5(aQ[R"UR UR UR UR U5UR5$[SUS35e)Nrz(. Only floating point clamp is supported) rr6rxr(r) create_clampfrror9)rpminmaxrr.r6s rclamprs+Cg>HC )!' :FA )!' :FA GGE yy..qxxSZZQ^_abagaghh+E72Z[\\rc@[XU5upURRnURRnUR5(aUR5(d [ X45e[ X45nXS:wa [ XU5nXT:wa [ XU5nX4$rv)rror~rOrr@r)rrr. input_sca_ty other_sca_tyrs rbitwise_op_type_checking_implrs/gFLE::$$L::$$L    (;(;(=(=' CC%lAJ!U0!U0 <rc[XU5up[R"URURUR5UR 5$rv)rr(r) create_andrrorrr.s rrr:0wGLE 99W'' ellCUZZ PPrc[XU5up[R"URURUR5UR 5$rv)rr(r) create_orrrors ror_rs:0wGLE 99W&&u||U\\BEJJ OOrc[XU5up[R"URURUR5UR 5$rv)rr(r) create_xorrrors rxor_rrrcURR5(d![U[R"S5U5nURR5(d![U[R"S5U5n[ XU5$Nre)rois_int1bitcastr(r6rrs r logical_andrsc ::    rxx/9 ::    rxx/9 g &&rcURR5(d![U[R"S5U5nURR5(d![U[R"S5U5n[ XU5$r)rorrr(r6rrs r logical_orrsc ::    rxx/9 ::    rxx/9 uW %%rcURR5(d![U[R"S5U5n[ X5$r)rorrr(r6invert)rr.s rnot_rs7 ::    rxx/9 % !!rc[XU5up[R"URURUR5UR 5$rv)rr(r) create_lshrrrors rlshrr:0wGLE 99W((u||Dejj QQrc[XU5up[R"URURUR5UR 5$rv)rr(r) create_ashrrrors rashrrrrc[XU5up[R"URURUR5UR 5$rv)rr(r) create_shlrrors rshlrrrcU$rvr\)rs rplusrs LrcURRnUR5(a[SUR 5-S-5e[ R "URURU55U5n[X0SU5$)Nz$wrong type argument to unary minus ()T) ror~rwr'rr(r)get_null_valueto_irr)rr.r_0s rrrsu::$$L?,BWBWBYY\__`` 7)),*<* 66??  ww FFLLE ==% ((rc|[XU5upURRnUR5(aD[R "UR URUR5[U55$UR5(aUR5(aD[R "URURUR5[U55$[R "URURUR5[U55$[SU35er)rror~rxr(r)create_fcmpOGTrrrOrcreate_icmpSGTcreate_icmpUGTr9rrr.rSs r greater_thanr/gFLE !!Iyy// ellKZX]M^__       " " $ $99W33ELL%,,OQ[\aQbc c99W33ELL%,,OQ[\aQbc c &yk2 33rc|[XU5upURRnUR5(aD[R "UR URUR5[U55$UR5(aUR5(aD[R "URURUR5[U55$[R "URURUR5[U55$[SU35er)rror~rxr(r)create_fcmpOGErrrOrcreate_icmpSGEcreate_icmpUGEr9rs rrrrrc|[XU5upURRnUR5(aD[R "UR URUR5[U55$UR5(aUR5(aD[R "URURUR5[U55$[R "URURUR5[U55$[SU35er)rror~rxr(r)create_fcmpOLTrrrOrcreate_icmpSLTcreate_icmpULTr9rs r less_thanr&rrc|[XU5upURRnUR5(aD[R "UR URUR5[U55$UR5(aUR5(aD[R "URURUR5[U55$[R "URURUR5[U55$[SU35er)rror~rxr(r)create_fcmpOLErrrOrcreate_icmpSLEcreate_icmpULEr9rs rrr.rrc[XU5upURRnUR5(aD[R "UR URUR5[U55$UR5(aD[R "URURUR5[U55$[SU35er) rror~rxr(r)create_fcmpOEQrrrO create_icmpEQr9rs requalr.=/gFLE !!Iyy// ellKZX]M^__     yy..u||U\\JJW\L]^^ &yk2 33rc[XU5upURRnUR5(aD[R "UR URUR5[U55$UR5(aD[R "URURUR5[U55$[SU35er) rror~rxr(r)create_fcmpUNErrrO create_icmpNEr9rs r not_equalr3Ir/rc[U[5(a[U[5(d [S5e[US- 5n[US- 5nU(dU(a [S5eX::a [S5eX- nXUS- -S:wa [S5eU/n[R "[R U5n[R"URX5U5$)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) rbrfr'rcr(rr+r)create_make_range)startendr.is_start_int64 is_end_int64rangerrs raranger<Zs eS ! !C)=)=JKK%2+&Nr ?L344 |XYY KE !>?? GE ]]288U +F 99W..u:F CCrc[U[R5(a.URRS:XdS5e[ XU5nOlUc [ S5eUS:Xa!URURU55nO![USUR35nU"U5n[R"X5n[XU5$)Nr zonly accepts size-1 tensorz2dtype must be specified when value is not a tensorrget_) rbr(r)numelrErr'r r getattrnamesplat)rrEr6r. get_value_fns rrjrjks%##{{  A%C'CC%U7+ =QR R A:**5;;w+?@E"7d5::,,?@L 'E %' w ''rcURR5(aS5e[U5S:XaU$[R"UR U5n[R "URURU5U5$)NzCannot splat a block tensorr) rorlenr(rr6r) create_splatr)rErr.rs rrBrBsgzz""$$C&CC $ 5zQ ]]5;; .F 99W))%,,> GGrc"SnUHnXE-nM URRU:wa [S5e[R"URR U5n[R "URURX5U5$)Nr z:reshape() cannot change total number of elements in tensor) ror?r'r(rr~r)create_reshaper)r dst_shape can_reorderr.r?srs rreshaperLss E   zz5 UVV ]]5::,,i 8F 99W++ELL)QSY ZZrcURVs/sHn[R"U5PM nnURUS5 URR 5(d [ XUS9$[R"URRU5n[R"URURU5U5$s snf)Nr )rr.) rr(_constexpr_to_valueinsertrorrBrr~r)create_expand_dimsr)rr-r.rprIrs r expand_dimsrQs49KK@Kq''*KI@ T1 ::   UW== ]]5::,,i 8F 99W// dCV LLAs Cc\U(dS5e[UR5S:Xde[R"URR URSURS-/5n[R "URURUR5U5$)Nz;current implementation of `cat` always may reorder elementsr r) rErr(rror~r) create_catr)rrrJr.ret_types rcatrUs| UUU; syy>Q  }}SXX__syy|ciil/J.KLH 99W'' CJJ? JJrc [XU5upUR/:HnU(a[USU5n[USU5n[URS[R 5(a[R "S5nOSnURU/-n[R "URRU5n[R"URURUR5U5nU(a [US/SUS9nU$)Nrr FrJr.) rrrQrbr(rmrror~r) create_joinrrL)abr. was_rank_1two new_shaperTrs rjoinr_s g .DABJ 1g & 1g &!''"+r||,,ll1o3%I}}QVV]]I6H ))G''!((;X FCcA3E7C Jrc[UR5S:de[R"URS5S:XdeURSSn[R"UR R U5nURUR5upE[R"XC5[R"XS54$)NrrWr ) rErr(rNrror~ create_splitrr))rZr.r^rToutLHSoutRHSs rsplitrds L1    " "1772; /1 45 4 I}}QVV]]I6H))!((3NF &# &# rc[UR5[U5:wa [S5e[SU55[ [ [U555:wa[SU35e[ R"URRUVs/sHo0RUPM sn5n[ R"URURU5U5$s snf)Nz5permute dims must have the same length as input shapec3N# UHn[R"U5v M g7frv)r(rN).0ds r permute..s 6Ab$$Q''s#%z?permute dims must be a permutation of 0, 1, ..., n-1, but were ) rErr'sortedlistr;r(rror~r) create_transr)rdimsr.rhrTs rpermuteros 5;;3t9$PQQ 6 66$uSY?O:PPZ[_Z`abb}}UZZ..0NAQ0NOH 99W))%,,=x HH1OsC& c URR5(dR[R"URU5n[R"UR UR U5U5$URR5n[U5[U5:wa[SUSU35eX:XaU$[U5H1upVXU:wdMUS:wdM[SXSUSUSUSU3 5e [R"URRU5n[R"URUR U5U5$)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 : ) rorr(rr)rFrget_block_shapesrEr' enumerater~create_broadcast)rrr.r src_shapeiitems rbroadcast_impl_shaperxs0 ::   uzz51yy--ellEBFKK ++-I 9~U#C N *3y>3y>:ii : :3::q I " fmmaS9_ MO"335 ; ^c)n ,3y>3y>:ii : :3::q I " fmmaS9_ MO"335 ; 9~Y///  +GALEqy  '1*%-  & "-/21v"68<"=?B4y"IKR"SUXY^U_"`aa,  !]]6==)? ? 99W++ELL&,,w:OPRX YYrc"URn[U[R5(a URn[U[R5(a URnUR 5(a9[R "URURR55nXA:XaU$URnURn[U5nSnUR5(anUR5(aYURUR:a?Uc[RRnOJU[RR:waSnO)Ub&[S[!U5-S-[!U5-5eUR#5(dUR#5(a8UR$R'S5cS5eUR$S"XX2S9$UR)5(aUR5(d1UR5(aUR)5(dU(aA[R*"UR-UR.UR1U5U5U5$UR35(aUR55(a*UR75(a:UR55(d%[9[9U[R:U5Xb5$UR5=(a0 UR5=(a URUR:nU(a@[R*"UR=UR.UR1U55U5$UR5=(a0 UR5=(a URUR:n U (a@[R*"UR?UR.UR1U55U5$URA5(Ga#URA5(Ga URBURB:wdURDURD:waURG5=(a URI5(+n URI5(aWURJR1U5n [R*"URMU 5URJ5n [OX U5$[R*"URQUR.UR1U5U 5U5$URS5(GaURA5(GaURI5(aWURJR1U5n [R*"URMU 5URJ5n [OX U5$URG5(a@[R*"URUUR.UR1U55U5$[R*"URWUR.UR1U55U5$URA5(aURS5(aURI5(dURG5(d@[R*"URYUR.UR1U55U5$[R*"UR[UR.UR1U55U5$UR]5(aURA5(aURBn U S:Xa@[R*"UR_UR.UR1U55U5$U S :XaX[O[9U[R`U5[R*"URcS 5[R`5U5$URA5(aUUR]5(a@[R*"UReUR.UR1U55U5$UR]5(aUUR]5(a@[R*"URgUR.UR1U55U5$S US U35e) 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 )4rorbr(rmrErrr~rrrrxrr rrr'r{ is_fp8e4b15 codegen_fnsgetrNr)create_fp_to_fprr rLrKrMrrHcreate_fp_trunc create_fp_extrOr4r5ris_boolr6r r3create_int_castis_standard_floatingcreate_fp_to_sicreate_fp_to_uicreate_ui_to_fpcreate_si_to_fprwcreate_ptr_to_intrhrcreate_int_to_ptrr)rrr.fp_downcast_roundingrrruse_custom_rounding truncate_fpext_fp sign_extendtyr bitwidths rrr?s ZZF&",,''& 55399 v}}ejj.I.I.KL  JJ11EFJ$:$:%%  ' '**G*G G  '@P@P@U@U)= !R%5%5%:%: :RVW>W[e[t[t>t ..0M9K9K9M5M     ""7+B711"5u{{CBU0 099W44U\\6<->-@-@     ""7+B711"5u{{CBU0 0  % % ' '99W44U\\6<>@@     z'?'?'A'A99W44U\\6<99W66u||V\\RYEZ[]cd d q=T%7;RYYwGXGXYZG[]_]e]e=fhop pz0022yy225<<gAVWY_``z0022yy// fll7>STV\]]4LtF8445rc2[RRnU(auUS:Xa[RRnU$US:Xa[RRnU$US:Xa[RR nU$[ SUS35eU$)Nz.ca.cgz.cvCache modifier  not supported)r CACHE_MODIFIERrCACGCVr'cache_modifiercaches r_str_to_load_cache_modifierrs    " "E U "%%((E L u $%%((E L u $%%((E L~.>nMN N Lrcv[RRnU(aUS:Xa[RRnU$US:Xa[RRnU$US:Xa[RR nU$US:Xa[RR nU$[SUS35eU$)Nz.wbrz.csz.wtrr)r rrWBrCSWTr'rs r_str_to_store_cache_modifierrs    " "E U "%%((E Lu $%%((E L u $%%((E L u $%%((E L~.>nMN N Lrc[RRnU(aSUS:Xa[RRnU$US:Xa[RRnU$[ SUS35eU$)N evict_last evict_firstzEviction policy r)r EVICTION_POLICYNORMAL EVICT_LAST EVICT_FIRSTr')eviction_policyevictions r_str_to_eviction_policyrsu!!((H l *))44H O  -))55H O//@OP P OrcSnU(aSUS:Xa[RRnU$US:Xa[RRnU$[ SUS35eU$)NzeronanzPadding option r)r PADDING_OPTIONPAD_ZEROPAD_NANr')padding_optionpaddings r_str_to_padding_optionrshG V #''00G N u $''//G N~.>nMN N Nrcv[RRnU(aUS:Xa[RRnU$US:Xa[RRnU$US:Xa[RRnU$US:Xa[RR nU$[ SUS35eU$)Nacquirereleaseacq_relrelaxedMemory semantic r)r MEM_SEMANTICACQUIRE_RELEASEACQUIRERELEASERELAXEDr') sem_optionsems r _str_to_semrs // ) )C  "//))C J9 $//))C J 9 $//11C J 9 $//))C J/ |>JK K Jrc2[RRnU(auUS:Xa[RRnU$US:Xa[RRnU$US:Xa[RRnU$[ SUS35eU$)Ngpuctasysrr)r MEM_SYNC_SCOPEGPUCTASYSTEMr') scope_optionscopes r _str_to_scopers    ! !E 5 %%))E L U "%%))E L U "%%,,E L/ ~^LM M LrcU(a[US5(dU/nUVs/sH0n[U[R5(a UROUPM2 nnUH3n[U[ 5(aSUs=::a[ U5:aM0e e [ U5S:de[ U5[ [U55:XdS5e[U5$gs snf)N__iter__rz'Duplicate dimension in `boundary_check`r\) hasattrrbr(rmrErfrEsetrk)boundary_check block_shapeelemdims r_canonicalize_boundary_checkrs~z22,-N]kl]kUY 4(F(F$**DP]kl!Cc3''A,Gs;7G,G GG,G GG">"Q&&&>"c#n*=&>>i@ii>n%% ms7Cc UcUb [S5eURRRn U [R:wdS5eU R 5(a)U[ RR:Xa [S5eURRn [X:R55n[R"URURX4XVU5U 5$)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'ro element_tyr(rerOr rrrrrr)create_tensor_pointer_loadr) ptrmaskrrrrr is_volatiler.elt_tyrs r_load_block_pointerrs 5,fgg XX + +F RWW SSS  }}7b&7&7&?&??[\\XX F2.BYBYB[\N 99**3::~PUalmou wwrc URRR5(d'[SURR 5S35eUcUb [S5eU(dU(a [S5eURR 5(dbU(a*URR 5(a [S5eU(a*URR 5(a [S5eURR 5(aNUb$[ XRR5U5nUb$[ X RR5U5nURRn U Rn U [R:Hn U (a<[Rn [R"XR5n [X U5nUb [X*U5nURR 5(a1URR5n [R"X5n OU n Uc3[R "UR#UR$XVU5U 5nOQ[R "UR'UR$UR$U(a UR$OSXVU5U 5nU (a[U[RU5nU$)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)ror~rwr'rrrxrrrr(reint8 pointer_type address_spacerrr) create_loadrcreate_masked_load)rrrrrrrrr.ptr_tyrrrrrs r _load_legacyr*sB 88?? ! ! # #01B1B1D0E]STT |)DEE.TU U 88     DII&&((de e UZZ((**ef f xx  'hh.G.G.I7SD  (0I0I0KWUEXX__F   FG)=)=>3( UG, xx))+v- |ii++CJJUW]^ii  & &szz4;;PU [_af'2 45;=3) Jrc [U5n [U5n [U5n URR 5(a8URR R 5(a[XX#XXU5 $[XX#XXU5 $rv) rrrrorwrrrr) rrrrrrrrr.rrrs rloadrhsv ( 7E&7H$^4G xxSXX0099;;"3eWU]lsttCughelmmrc [XQSS9nURURXRU5[ U5[ U55n[ R"Xd5$NF require_i64)_convert_to_ir_valuescreate_descriptor_loadrr rrr(r))desc_ptroffsetsrrror.rps rdescriptor_loadrxsQ#G%HG&&xGAT'B>'R'>'O QA 99Q rc[X2SS9n[R"URURURU5[R 5$r )r r(r)create_descriptor_storervoid)rrErr.s rdescriptor_storers>#G%HG 99W44X__ellT[\^`^e^e ffrc U(a#USR[R:Xde[R"U R UR UR UV s/sHoR PM sn UV s/sHoR PM sn UV s/sHoR PM sn UV s/sHoR PM sn UUUU 5 [R 5$s sn fs sn fs sn fs sn f)Nr)r6r(rhr)create_tensormap_createrr) rglobal_addressbox_dim global_dim global_strideelement_stride elem_typeinterleave_layout swizzle_mode fill_moder.rps rtensormap_creater!s a 0 6 6"(( BB B 99'' OO  ! !& 'w!XXw ') *z!XXz *, -}!XX} -- .~!XX~ .        ( * - .s C$;C)C.1C3c~[R"URUR5[R5$rv)r(r)#create_tensormap_fenceproxy_acquirerr)rr.s rtensormap_fenceproxy_acquirer$s) 99W@@QSUSZSZ [[rc Ub [S5eURRR5nURR 5(d [ XU5nURR 5(dS5eXqRR5:Xd&SUSURR5S35eURRRURR:Xd@SURRRSURRS35eURRRnU[ R:wdS5e[X75n[XU5n[ R"URURURX4U5[ R5$) 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'rorrrrrxr(rerrr)create_tensor_pointer_storerr) rvalrrrrr.rrs r_store_block_pointerr(s fgg((%%668K 88    "3W= 88    O OO  ((33 ] k]"4SXX5N5N5P4QQ[\]  88   ) )SXX-@-@ @qDWX[X`X`XkXkXvXvWwxQRURZRZReReQffpCqq @ XX + +F RWW SSS 2.NN sG $C 99W88SZZQ_hpqWW rc >URRR5(d'[SURR 5S35eU(a [S5eURR 5(d[URR 5(a [S5eU(a*URR 5(a [S5eURR 5(aK[ XRR5U5nUb$[ X RR5U5nURRnURnU[R:Xa<[Rn[R"XR5n[XU5n[XU5nU(dJ[R"UR!UR"UR"XE5[R$5$URRR'5(d [S5e[R"UR)UR"UR"UR"XE5[R$5$)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)ror~rwr'rrrxrrrr(rerrrrr) create_storerrrcreate_masked_store) rr'rrrrr.rrs r _store_legacyr,s 88?? ! ! # #01B1B1D0E^TUUAB B 88     88    ef f DII&&((de e xx"3(A(A(CWM  'hh.G.G.I7SD XX__F   F)=)=>3( sG $C yy--cjj#**eVXZX_X_`` 99   # # % %=>> 99W00SZZV[fhjhoho pprc [U5n[U5nURR5(d)URRR5(a [ S5eURR 5(a7URRR5(a[XX#XxU5$[XX#XxU5$)N"Cannot store to a constant pointer) rrrois_constr~r'rwrrr(r,) rr'rrrrr.rrs rstorer0s ) 8E&7H xxchhoo6688=>> xxSXX0099;;#CdEU\]]StUgVVrc F[U5n[U5nURRRnUR S;a [ S5e[R"URURURURX45UR5$)N)r5rz9atomic_cas only supports elements with width {16, 32, 64}) rrror~rrr'r(r)create_atomic_casr)rcmpr'rrr.rs r atomic_casr5sy c C % E++J$$L8TUU 99W..szz3::szzSV^`c`h`h iircTURRR5(d&[SURR 5-5eURR 5(d)URR R 5(a [S5eURRR nU[RLaUS:wa[SU-S-5eU[R[R[R[R4;a[SU-S-[U5-5eURR5(aNUb$[X RR!5U5nUb$[XRR!5U5n[#XRRR U5nU(dUR%S5n[RnURR5(afUR'X`RR!55n[R("[RURR!55n[R*"Xg5nXU4$)Nz)Pointer argument of store instruction is r.ratomic_z does not support fp16z does not support T)ror~rwr'rr/rr(rFrerint16rGr{rrxrrrrdrFrr))rr'ropr.rmask_irmask_tys ratom_red_typechecking_implr<s 88?? ! ! # #DsxxGXGXGZZ[[ xxchh11::<<=>>++JRZZB%KR*BBCCbggrww"++>>R*>>ZPQQ xx  'hh.G.G.I7SD ?&sHH,E,E,GQC sHHOO.. 8C ""4('' 88    **7HH4M4M4OPGmmBGGSXX-F-F-HIGyy* T>rc [XUSU5upn[U5n[U5nURRnUR 5(aUR 5(aj[R"UR[RRURURURX45UR5$[R"UR[RRURURURX45UR5$U[R[R 1;a[#SU35e[%/SXe5nU[R:Xa[R&O[R(n[+XU5n [+U[R,"US5U5n U[R:Xa[R.O[R0n [+XU5n [+U[R,"U S5U5n [3XU5n[5XU5n[R"UR[RRU RU R[7X.U5RX45U R5n[R"UR[RR8U RU R[7X/U5RX45U R5n[;UUUU5n[+UXe5$)Nrz#atomic_max not supported for dtype rar )r<rrror~rOrr(r)create_atomic_rmwr ATOMIC_OPMAXrUMAXrHrJr9rjr+rhrrrgrirr&rUMINwhererr'rrrr.sca_tyri_typei_vali_ptrui_typeui_valui_ptrposnegpos_retneg_retrs r atomic_maxrP,|/$wONCd c C % E XX__F }}    ! !99))",,*:*:CJJ TXT_T_adlnqnvnvx x99))",,*;*;SZZUYU`U`bemorowowy y bjj"**--=fXFGG C )D2::-RXX288F C )E C3W =E!RZZ/biiRYYG S7 +F S"//'15w ?F 7 +C Cw 'Cii!!",,"2"2ELL%,,"&t'":"A"A3 OPUPZPZ\Gii!!",,"3"3V]]FMM"&t'":"A"A3 OPVP[P[]G Wgw /C 3 ((rc [XUSU5upn[U5n[U5nURRnUR 5(aUR 5(aj[R"UR[RRURURURX45UR5$[R"UR[RRURURURX45UR5$U[R[R 1;a[#SU35e[%/SXe5nU[R:Xa[R&O[R(n[+XU5n [+U[R,"US5U5n U[R:Xa[R.O[R0n [+XU5n [+U[R,"U S5U5n [3XU5n[5XU5n[R"UR[RRU RU R[7X.U5RX45U R5n[R"UR[RR8U RU R[7X/U5RX45U R5n[;UUUU5n[+UXe5$)Nrz#atomic_min not supported for dtype rar )r<rrror~rOrr(r)r>r r?MINrrBrHrJr9rjr+rhrrrgrirr&rrArCrDs r atomic_minrTSrQrc [XUSU5upn[U5n[U5nURRnUR 5(a[ RRO[ RRn[R"URXpRURURX45UR5$)Nr)r<rrror~rxr r?FADDADDr(r)r>r)rr'rrrr.rEr9s r atomic_addrXzs/$wONCd c C % E XX__F$0022   8H8HB 99W..r::szz4;;X[cehemem nnrc $[XUSU5upn[U5n[U5n[R"UR [ RRURURURX45UR5$)Nand) r<rrr(r)r>r r?ANDrrorr'rrrr.s r atomic_andr]q/$wONCd c C % E 99W..r||/?/?SZZY]YdYdfiqXX rc $[XUSU5upn[U5n[U5n[R"UR [ RRURURURX45UR5$)Nor) r<rrr(r)r>r r?ORrror\s r atomic_orrbso/$gNNCd c C % E 99W..r|| CJJX\XcXcehpXX rc $[XUSU5upn[U5n[U5n[R"UR [ RRURURURX45UR5$)Nxor) r<rrr(r)r>r r?XORrror\s r atomic_xorrfr^rc $[XUSU5upn[U5n[U5n[R"UR [ RRURURURX45UR5$)Nxchg) r<rrr(r)r>r r?XCHGrror\s r atomic_xchgrjsq/$PNCd c C % E 99W..r||/@/@#**cjjZ^ZeZegjrXX rcUR5URR;d!SURRSU35eUR5nUS:XaSn[ [ R U5$)Nzinput_precision must be one of z. Got TF32X3TF32x3)lowerrallowed_dot_input_precisionsupperr@r INPUT_PRECISION)input_precisionr.s r_str_to_dot_input_precisionrssx  "goo&R&R Rp )'//*V*V)WW]^m]nop R%++-O("" 2%% 77rc  URR5(aURR5(deURR5(a!URR5(aGOUR[R [R [R[R[R4;dSUR35eUR[R [R [R[R[R4;dSUR35eURUR:Xd!SURSUR35eURR5(dURR5(a6[U[RU5n[U[RU5nUcURRn[X65n[UR 5n[UR 5nXxs=:XaS:Xd2O Xxs=:XaS:Xd$O SUR SUR S 35eUR S R"UR S R":XdVS UR S UR SUR S R"SUR S R"S 3 5eUR$R'S5cS5eUR$S"URUR5n UR S R"U S:a@UR S R"U S:a UR S R"U S:dSU SSU SSU S35eURR(R+5(aQURR([R :XdS5eUR-S5n [R.n OUR15(a [3S5eURR(R55(d)URR(R15(a"UR7S5n [Rn O9UR95(aUR;S5OUR7S5n Un URR S n URR S n URR S nUS:XaURR SOSn[R<"X(aXU /OX/5nUcUR?X(aXU /OX/5nOUR@nURU:XdeUcXURR5(a6URR5(aURRBnOXSnOUURR5(a6URR5(aXN:a[3SUSUS 35e[RD"URGUR@UR@UX45U5$)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: r rWzFirst 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 ()$rorr6rNr(ruint8rFrGrHrrrdefault_dot_input_precisionrsrErrErrr~rO get_int32r+rMr'rKget_fp32rLget_fp16rrFrmax_num_imprecise_acc_defaultr) create_dot)rraccrrmax_num_imprecise_acc out_dtyper.lhs_rankrhs_rankryr  ret_scalar_tyMNKBr acc_handles rdotrs 88    388#4#4#6#66 6 yycii..00 yyRWWbhh BKKZZ)) O,B399+*N O)yyRWWbhh BKKZZ)) O,B399+*N O)yyCII%k)OPSPYPY{Z_`c`i`i_j'kk% yy#))"7"7"9"93 G,3 G,!//EE1/KO399~H399~H  $1 $(A(AREpqtqzqzp{|EFIFOFOEPPQDRR A 99R=  #)) # E q( 3LSYYKXUVYV_V_`bVcViViUjjZ[^[d[deg[h[n[nZoopqq     " "> 2 >t@tt >&&~6sxxJL 99R=  ,q/ 1ciim6I6I\Z[_6\ IIbM  <? 2r,\!_,=W\RS_DUU_`lmn`o_p qr 3 xxxx"'')A+AA)   q !      vx x  " "chhoo&=&=&?&?   a  $-$5$5$7$7W  a W=M=Ma=P!  rA rA rA%]qA ]]=q1)qf EF {))"1qQi1&I ZZ xx6!!!$ 99    #))"2"2"4"4$+OO$Q$Q !$% ! 99    #))"2"2"4"49N9R67L6MM]^_]``abc c 99W'' CJJ Os rc`US:Xa[RR$US:Xa[RR$US:Xa[RR$US:Xa[RR $US:Xa[RR $[SUS35e)Ne4m3e5m2e2m3e3m2e2m1zInvalid float format: r[)r F8F6F4TYE4M3E5M2E2M3E3M2E2M1r') float_formats r_str_to_fp_typersv{{v{{v{{v{{v{{ -l^1= >>rc JURR5(aURR5(de[UR5n [UR5n Xs=:XaS:Xd2O Xs=:XaS:Xd$O SURSURS35e[ U5n [ U5n US;d SU35eUS;d S U35e[ U[ R5=(a URSLn U (dS 5eURRS nURRS SunnUS :XaSOS nUUURRS-:Xd"SURSURS35eUS:d SU<35eU S:XaURRSOSn[ R"UU(aUUU/OUU/5nURS5nUc!URUU(aUUU/OUU/5nOURnURU:XdeU (aSO URn[ R"URURURXRUU U5U5$)Nr rurvrwr )rrrzNYI: lhs_format )rrzNYI: rhs_format zNYI: rhs_scale not supportedrxrr rWzCReduction dimension should pack the same number of elements; (lhs: rz!scaled_dot NYI for K < 64. Got K=r)rorrErrrbr(rmrErr}rFrr)create_dot_scaled)r lhs_scale lhs_formatr rhs_scale rhs_formatrrr.rrlhs_format_enumrhs_format_enumrhs_scale_is_nonerrrPACKEDrrr rrhs_scale_handles r dot_scaledrs 88    388#4#4#6#66 6399~H399~H  $1 $(A(AREpqtqzqzp{|EFIFOFOEPPQDRR A%j1O%j1O 1 1R5Ej\3RR 1 ) )J-=j\+JJ )"9bll;W SW@W <<<  rA 88>>"# DAq&QAF     tRSVS\S\R]]fgjgpgpfqqr st 7:8aT::7%]qA ]]91q!Qi1a& AF   ! B {))"1q!Qi1a&I ZZ xx6!!!0ti6F6F 99!!#**i.>.>Q[Q[]m"1: ?@F HHrcUR[R:wa#[R"SUR35 [ U[RU5n[ XUSS5upURR5(a[XU5up[XU5upO[XU5upURn[R"URURURUR5U5$)Nzgtl.where with a non-boolean condition is deprecated and will error out in a future triton release. Got T) r6r(rewarningswarnrrrorrr) create_selectr) conditionrprr.rrs rrCrC0s"''! uvwFwFvG H Y1I 'gtT BDA~~  +I'B #A'21+I'B VVF 99W**9+;+;QXXqxxPRX YYrcnU(a[R"X5nOUn[R"X5$rv)r(rr))rprSrres_tys r wrap_tensorrFs)y4 99Q rc(^^^^ ^ Uc[U4SjT55mSnTSRRm [T 5nX:d SUS35e[ T 5VVs/sHupVXQ:wdM UPM snnm [ U 4SjT55(dS5eTR TVs/sHowRPM snU5mU"T5 TR5 [UUU 4Sj[[T5555$s snnfs snf)Nc3d># UH%n[XRR/STS9v M' g7f)TrXN)rLr?rE)rgtr.s rrireduction..Qs'f_eZ[wq77==/tWU_es-0rz&reduction axis must be < inputs rank (r c3T># UHoRRT:Hv M g7frv)ror)rgrrs rrirXs5fvv||u$fs%(z-all reduction inputs must have the same shapec3># UH8n[TRU5TURRT5v M: g7frvr get_resultror~)rgrvinputs reduce_oprs rrir^s9tas\]Y11!4fQinn6K6KYWWasAA) tuplerorrErsall create_reducerverifyr;) rr-region_builder_fnr.rankrvrKrrrrs ` ` @@@r reductionrOs |f_eff 1INN E u:D ;H@aHH;(/=/tq19/=I 5f5 5 5f7ff 5%%&@Axx&@$GIi   tafgjkqgrast tt>'As" D 1D +Dc^^^TSRRm[T5nU*Us=::aU:dO SUSUS35eUS:aX- nTH$nURRT:XaMS5e URTVs/sHofRPM snX5mU"T5 TR 5 [ UUU4Sj[[T5555$s snf)Nrz scan axis z must be < inputs rank (r z(all scan inputs must have the same shapec3># UH8n[TRU5TURRT5v M: g7frvr)rgrvrscan_oprs rri#associative_scan..ws9n[mVWW//2F1INN4I4I5QQ[mr)rorrE create_scanrrrr;) rr-rreverser.rrrrs ` @@rassociative_scanrfs 1INN E u:D 5D 4 S:dV3KD6QR!SS  ax  vv||u$P&PP$!!V"R>R>T)UV Hrc[UR5[U5:wa [S5eURR S[ R "XRR555 U$)NzDShape of input to max_contiguous does not match the length of valuesz tt.contiguityrErr'rrr rrrs rmax_contiguousrsQ 177|s6{"_``HHor||FHH.s,XKDz$$?4)?)?%)??Ks9;zGExpected a list of constant integers (`int32_t` range) in `block_shape`z# UHn[T5[U5:Hv M g7frv)rE)rgrrs rrirs!dCcis;3y>1Ccs$'zBExpected shape/strides/offsets/block_shape to have the same length)r rorwrrr'r(rerrrrrrbrmrErrkrlr;rEcreate_make_block_ptrrr)r) baserstridesrrorderr.rrs ` rmake_block_ptrrs "' 1E#G5G#G%HG 99    !5!5!>!>!@!@hii yyrww&D"//"''4993J3JKWU ; + +"m VabVadD",,!?!?4::TIVabK XKX X XRQR X 5* % %PU VPU:dBLL99TZZt CPUE V %=Ds5z!23 3s5ss 3 dET[Ccd d dMLM d  * *4;;Q\^c dF 99VR__R]]499;O;OQ\-]^ __%c Ws 97I))7I.c[X!SS9n[R"URURU5UR 5$r )r r(r)create_advancerro)rrr.s radvancers6#G%HG 99W++DKKA499 MMr)r-rfr. ir.builderreturn tl.tensor)r:tl.dtyper;rrr) r:rrPrcr;rrQrcrRrcrr)T)rqrc)rrrrryrcrNone)FFTF)rtl.tensor | numbers.Numberrrr.rrTuple[tl.tensor, tl.tensor])rrrrr.rrcallable) rrrrrrcr.rrr)rrrrr.rrr) rrrrrrcr.rrr)rprrrrtl.PropagateNanr.r) rprrrrrrrr.r)rrrrr.rrr)rrrrr.rrr)rrr.r)rrrr)rrr.rrr)rrr.rrr)rrrz tl.block_type)r7rfr8rfr.rrr)r List[int]r6rr.rrr)rErrrr.rrr) rrrIrrJrcr.rrr)rrr-rfr.rrr) rrrrrJrcr.rrr)rZrr[rr.rrr)rZrr.rrr)rrrnz Tuple[int]r.rrr)rrrrr.rrr)rrrrr.rrr)r Optional[str])rrrrr.rrrrv) rrrrr.rrrrr)rrrOptional[tl.tensor]rr rrrr{rr{rr{rrcr.rrr) rrrr{rr{r.rrr)rrrErr.rrr)rrrrrList[tl.tensor]rr rr rr rrfrrfrrfr rfr.rrr)rrr.rrr)rrr'rrr rr{rr{r.rrr)rrr4rr'rrr{rr{r.rrr) rrr'rrrr9r{r.rrz&Tuple[tl.tensor, tl.tensor, tl.tensor])rrr'rrrrr{rr{r.rrr)rrrrrrrrrrrfrrr.rrr)rr)rrrrrrrr rztl.tensor | Nonerrr.rrr) rrrprrrr.rrr)rSequence[tl.tensor]r-rfr.rrTuple[tl.tensor, ...]) rr r-rfrrcr.rrr )rrrrfr.rrr)rprrrrr)r.rrr) rr{rr rrcr.rrr)rrrr{r.rrr)rrr.rrr)u __future__rrtypingrrrrrr| _C.libtritonr r r(r r Exceptionrr/r2r@rUrnrzrrrrrrrrrrrrrrrrrrrrrrrrrrrrr&rr.r3r<rjrBrLrQrUr_rdrorxrrrrrrrrrrrrrrrrr!r$r(r,r0r5r<rPrTrXr]rbrfrjrsrrrrCrrrrrrrrrrrrr rrr\rrrs#";; CLF FD F @ /,&*/,/7/,d# V <]a,16QB&&:: ):644 )4$ 4 4 ) 4R4 :&&!*&485$5$ ]" +5 :U Q P Q '&" R R Q)$) 4 4 4 4 4 4"D"(.H[MK2 IL$2tpZ(04l5,l58Al5h       w,;| n n.1 nDG nVZ n n!* n ',5g   #  $ >\:)qXWW)3W8AW,j(27]6$)N$)No#(18EE&0E5>EP ? H 0 H=E HPZ H_h HPZ, u.o(o-Bo2n   8 VG BU0H $`NNr