widdlZddlZddlZddlZddlZddlZddlZddlmZm Z m Z m Z m Z m Z mZddlmZddlmZddlmZmZmZddlmZmZmZddlmZmZdd lmZd d lm Z m!Z!m"Z"dd l#m$Z$d Z%dZ&dede'fdZ(dede'fdZ)dede'fdZ*dede'fdZ+dede'fdZ,dZ-e'e.e/dhZ0GddZ1Gddej2Z3Gddej2Z4dZ5dZ6dS)N)AnyCallableDictOptionalTupleTypeUnion)language)ir) constexprtensor str_to_ty)_unwrap_if_constexprnv_tma_desc_type_value) _normalize_tyget_jit_fn_file_line) JITFunction)CompilationErrorCompileTimeAssertionFailureUnsupportedLanguageConstruct) ModuleTypecH|rdt|jzS|rzmangle_fn..(s!B!B!BB)B--!B!B!Br8cDg|]}|dt|S)c)repr)r<r constantss r7r=zmangle_fn..)s3!Y!Y!Y!Q"="=il););"="=!Y!Y!Yr8._d_'_sq_[]__)r/sortedreplace)namearg_tysrAmangled_arg_namesmangled_constantsrets ` r7 mangle_fnrP&s!B!B'!B!B!BCC!Y!Y!Y!YviGXGX!Y!Y!YZZ)11#u==)11#v>>)11#s;;CCCMM  = =& = =*; = =C Jr8oreturnc,t|tSN) isinstancerrQs r7_is_triton_valuerW2 a  r8c,t|tSrT)rUrrVs r7_is_triton_tensorrZ6rXr8c,t|tSrT)rUr rVs r7 _is_constexprr\:s a # ##r8ctt|o)|j p|jjdkS)Nr)rZtyper-numelrVs r7_is_triton_scalarr`>s2 Q   P):):%:%Oafla>OPr8c:t|ttfSrT)rUlisttuplerVs r7 _is_list_likerdBs a$ ' ''r8c |jrbt|D]T\}}t|s>t|s/t |j|d|jd|j|d|SdSdS)Nz Function z= is marked noinline, but was called with non-scalar argument :)noinline enumerater\r`rsrc__name__ arg_names)nodefnargsidxargs r7_check_fn_argsrqFs {!$  HC %% .?.D.D 2FDD DDjljvwzj{DDBDD  r8c eZdZdZdZdZdS)enter_sub_regionc||_dSrT) generator)selfrus r7__init__zenter_sub_region.__init__Us "r8cN|jj|_|jj|_i|j_|jj|_|jj |_ |j|jfSrT) rulscopecopyliveins local_defs prev_defsbuilderget_insertion_block insert_blockget_insertion_point insert_point)rvs r7 __enter__zenter_sub_region.__enter__Xs~~,1133 27799$&! N2FFHH N2FFHH|T...r8c|jj|j|j|j_|j|j_dSrT)rur~restore_insertion_pointrr{ryr}r|)rvrnkwargss r7__exit__zenter_sub_region.__exit__as< 66t7HIII $ $(N!!!r8N)rj __module__ __qualname__rwrrr;r8r7rsrsSsA###///33333r8rsc<eZdZdZdefdZdefdZdefdZdej defdZ dej defdZ dej defd Zdejdefd Zdejdefd Zdejdefd Zdejdefd ZdejdefdZdejdefdZdejdefdZdS)ContainsReturnCheckerc||_dSrT)gscope)rvrs r7rwzContainsReturnChecker.__init__js  r8rRc:tfd|DS)Nc3BK|]}|VdSrTvisit)r<srvs r7 z5ContainsReturnChecker._visit_stmts..ns-//Q4::a==//////r8)any)rvbodys` r7 _visit_stmtsz"ContainsReturnChecker._visit_stmtsms&////$//////r8ct|trB|js;|}t |j|SdSNF)rUrrgparserrr)rvrmfn_nodes r7_visit_functionz%ContainsReturnChecker._visit_functionpsO b+ & & Er{ EhhjjG(55;;GDD Dur8c8d}tj|D]\}}t|tr7|D]3}t|tjr|p||}4Qt|tjr|p||}|Sr)ast iter_fieldsrUrbASTr)rvrlrOrvalueitems r7 generic_visitz#ContainsReturnChecker.generic_visitws-- / /HAu%&& /!66D!$006!5TZZ%5%56E37++ /.TZZ.. r8rlc t|jtjrV|jj|jvrA|j|jj}t ||j}||SdS| |jSr) rUrrNameidrgetattrattrrr)rvrlrrms r7visit_Attributez%ContainsReturnChecker.visit_Attributesy dj#( + + z} ++ DJM2UDI..++B///5zz$*%%%r8ct|jtjurdS|j|jvr'|j|j}||SdSr)r^ctxrStorerrr)rvrlrms r7 visit_Namez ContainsReturnChecker.visit_NamesR >>SY & &5 7dk ! !TW%B''++ +ur8cdS)NTr;rvrls r7 visit_Returnz"ContainsReturnChecker.visit_Returntr8cdSrr;rs r7 visit_Assignz"ContainsReturnChecker.visit_Assign ur8cdSrr;rs r7visit_AugAssignz%ContainsReturnChecker.visit_AugAssignrr8c6||jSrTrrrs r7 visit_Modulez"ContainsReturnChecker.visit_Module  +++r8c6||jSrTrrs r7visit_FunctionDefz'ContainsReturnChecker.visit_FunctionDefrr8c||j}|jr|p||j}|SrT)rrorelse)rvrlrOs r7visit_IfzContainsReturnChecker.visit_IfsB ** ; 87**4;77C r8cj||jp||jSrT)rrrrs r7 visit_IfExpz!ContainsReturnChecker.visit_IfExps)zz$)$$? 4;(?(??r8c6||jSrT)rfuncrs r7 visit_Callz ContainsReturnChecker.visit_Callszz$)$$$r8N)rjrrrwboolrrrr AttributerrrReturnrAssignr AugAssignrModuler FunctionDefrIfrIfExprCallrr;r8r7rrhs0D0000T T     &CM &d & & & &sxD CMd ,,,,,,,co,$,,,, SV     @ @d@@@@%sx%D%%%%%%r8rcveZdZU dfdedeedeefdZdee e e e e efDZeeefed <ed ejjfd ejfd ejffd ZdZdZdedeeefddfdZdZdZ dZ!dZ"dZ#dZ$dZ%dZ&dZ'dZ(dZ)dZ*d Z+d!Z,d"Z-d#Z.d$Z/d%Z0e1j2d&e1j3d'e1j4d(e1j5d)e1j6d*e1j7d+e1j8d,e1j9d-e1j:d.e1j;d/e1j<d0e1j=d1i Z>ee?e1j@efed2<d3ZAd4ZBd5ZCd6ZDd7ZEd8ZFd9ZGe1jHd:e1jId;e1jJde1jMd?iZNee?e1jOefed@<dAZPe1jQdBe1jRdCe1jSdDe1jTdEiZUee?e1jVefedF<dGZWdHZXdIZYdJZZdKZ[dLZ\dMZ]de^eeffdNZ_defdOZ`dPefdQZadRZbdSZcdTe1jdfdUZee1jfdVe1jgdWiZhee?e1jiefedX<ejjkdYkr dZZld[Zmd\Znd]Zod^Zpd_Zqd`ZrfdaZsdbZtdTe1juddfdcZvddZwejjxevejjyeweze ewe eeweiZ{ee|e}e1jugeffede<xZ~S)g CodeGeneratorNFrjit_fnfunction_types file_namecz||_tj||_||_|dz |_|j||d||j_| |j_| in| |j_| |j n| |_ | in| |_ ||_ i|_ |D]\}}t|t r$| |j||j |<>t'|dd}|| vr$t'| ||j|j |<w||j |<i|_||_||_||_||_| |_d|_||_g|_d|_i|_||_ d|_!d|_"dS)NrrrF)#contextr r~r begin_lineset_locoptions codegen_fns module_map create_modulemodulefunction_ret_types prototyperitemsrUrgetrjrry attributesrAr function_name is_kernelcur_noderg scf_stackret_typer|_define_name_lookupdereference_namermvisiting_arg_default_value)rvrrrrrArrrrrrrrrgrrkv module_names r7rwzCodeGenerator.__init__s z'** "$q. Y A666& $/  (2(:""  6""N" LLNN # #DAq!Z(( !+ A!>!> A!!\266Kj((!(K)@!*!M!M A!" A $" *"    .06:6N6N6P6P+0'''r8ci|] }|j| Sr;rjr<rs r7 zCodeGenerator.s(t(t(t1Q(t(t(tr8builtin_namespaceprintminmaxc8t|jj||SrT)rrri)rvrlmessages r7 _unsupportedzCodeGenerator._unsupporteds+DKOT7KKKr8c t}|j||}||urdSt|rdS|jdi|x}rt |dkSdS)NFT__annotations__r )objectrrr\r)rvrK absent_markervalas r7_is_constexpr_globalz"CodeGenerator._is_constexpr_globals koodM22 -  5    4  1266::4@@ @1 3 ##{2 2ur8cdtffd dtffd tdtdtffd }|S)NrKc:j||SrT)ryr)rKabsentrvs r7 local_lookupz7CodeGenerator._define_name_lookup..local_lookups;??400 0r8cHj||}t||u|jvt |t ut |tt|ddt|dd dt |tj  |j tjdddkg r|St!t#jd |d d d ) N__triton_builtin__Frrztriton.language"TRITON_ALLOW_NON_CONSTEXPR_GLOBALS01z. Cannot access global variable a from within @jit'ed function. Triton kernels can only access global variables that are annotated as constexpr (`x: triton.language.constexpr = 42` or `x = triton.language.constexpr(42)`). Alternatively, set the envvar TRITON_ALLOW_NON_CONSTEXPR_GLOBALS=1, but we do not promise to support this forever.  )rrrrr^rrUrr startswithr r&rrosenviron NameErrortextwrapdedentrJ)rKrrrvs r7 global_lookupz8CodeGenerator._define_name_lookup..global_lookups +//$//C6M44+A#AII+sK00C!5u==Cr22==>OPPsHN33--d333JNN#GMMQTT     !4/3!4!4!455 6=WT35G5GII Ir8rRcx}jjfD]}|||}||ur|cSt|d)Nz is not defined)rrr )rKrlookup_functionrrrrrvs r7 name_lookupz6CodeGenerator._define_name_lookup..name_lookup1se"F#/@V@Z#Z ! !'f55&& LLL't44455 5r8)r*rr)rvrrrrs` @@@r7rz!CodeGenerator._define_name_lookup s 1s 1 1 1 1 1 1 I I I I I I I:  6c 6c 6 6 6 6 6 6 6 6 6r8rKrrRc.||j|<||j|<dS)z This function: called by visit_Assign() & visit_FunctionDef() to store left value (lvalue) 1. record local defined name (FIXME: should consider control flow) 2. store tensor in self.lvalue N)ryr|)rvrKrs r7 set_valuezCodeGenerator.set_value;s" " D %r8cn|j}|j}||fSrT)r~get_locr)rvlocips r7_get_insertion_point_and_locz*CodeGenerator._get_insertion_point_and_locDs3l""$$ \ - - / /3wr8cn|j||j|dSrT)r~rr)rvrrs r7_set_insertion_point_and_locz*CodeGenerator._set_insertion_point_and_locLs4 ,,R000 S!!!!!r8ct|s|g}|D]4}||t|tjrdS5dSrT)rdrrUrr)rvstmtsstmts r7visit_compound_statementz&CodeGenerator.visit_compound_statementSsfU## GE  D JJt   $ ++     r8cFtj||dSrTr NodeVisitorrrs r7rzCodeGenerator.visit_Module_  %%dD11111r8cj|j}|Jfd|jD}|S)Nc:g|]}|Sr;r)r<r6rvs r7r=z,CodeGenerator.visit_List..e#555C 3555r8)rrelts)rvrlrr(s` r7 visit_ListzCodeGenerator.visit_Listbs?jj""{{{555549555 r8c|j}|'jgtj}nt |trNfd|D}d|D}jd|Dt|}nLtj |j}j|j g|j }j |_ n%j |krtdj d|j}j|dS)NcZg|]'}tj|j(Sr;)r semantic to_tensorr~)r<rrvs r7r=z.CodeGenerator.visit_Return..os.ZZZ1(+55aFFZZZr8cg|] }|j Sr;r^r<rs r7r=z.CodeGenerator.visit_Return..ps444A444r8cg|] }|j Sr;handler0s r7r=z.CodeGenerator.visit_Return..qs;;;1ah;;;r8zInconsistent return types:  and )rrr~rOr voidrUrcr,r-r3r^rr3 create_blockset_insertion_point_to_end)rvrl ret_valueret_ty ret_values ret_typesrOpost_ret_blocks` r7rzCodeGenerator.visit_ReturnisYJJtz**   L  R ]FF  5 ) ) ZZZZPYZZZJ44444I L  ;; ;;; < < <9%%FF#--iFFC L  cj\ * * *XF = "DMM ]f $ $V$-VVfVVWW W2244 //?????r8cx |j\}}jr|dt |jjdddD]\}}|jj| dz }|j}|j}tj |tj } |tj | g|} ntj | ||} j rJd_ | d_ #d_ wxYwjrd nd } jjjjj| j_jjj} g} d }t1t3|D]!}|jvrLj|}t7|st9j|}| |X|jvr/j|D]!\}}j|||"tAjj!|tDrj|d d| tGj|jj!||dz }#j$}tK|| D]\}}&||j'| (|j)j$*rJj+j+tXj-kr,tXj-_+j.gntAj+t^rtaj+nj+gj_1j2jjj.fd jj1Dj3|rj4|dSdS)Nz,nested function definition is not supported.rrrtargetsr)targetr annotationTFpublicprivaterztt.nv_tma_desccg|];}j j|j.sI=, **288DL+A+ABB,,,r8)5rrnrmrrhdefaultsrCrprrrr AnnAssignrrr~get_or_insert_functionrrrrHrg push_backadd_entry_blockrangelenrAr\r appendr set_arg_attrrU param_typesrrrziprset_insertion_point_to_startr rhas_terminatorrr r5rOrcrbr; reset_typefinalizer7)rvrlrk kwarg_namesr default_valuearg_noderCrK st_target init_node visibilityentry arg_valuesrocstr insert_ptarg_name arg_values` r7rzCodeGenerator.visit_FunctionDefs!%DI!6!6 ; 7 Z##D*XYY Y )$)*XXY ,55dk4CU6:n6J6J4<6X6XZdfjfsuu dg&&&'')) s9~~&&  ADN""nQ'$S))7#DN1$566C!!#&&&'''+q'9?? e,,S$>>>>dn8=?OPPCG((.>BBB!!&c):):DNHX]>^>^'stDM':':':eierdsDN $ G  t~33DLAA B B B L  .2      ? L 3 3I > > > > > ? ?s %D  Dcg}|jD]}|||gz }||j}||fSrT)rnrkwarg)rvrlrkrprYs r7visit_argumentszCodeGenerator.visit_argumentssO 9 + +C $**S//* *IIjj,, +%%r8cPtj|||jSrT)rr#rrprs r7 visit_argzCodeGenerator.visit_args! %%dD111xr8c~||j}||j}||j}|tkrP||jvrt |dt|st |}||j|<|j|S||S)Nz4 is already defined. constexpr cannot be reassigned.) rrCrBrr ry ValueErrorr\r)rvrlrCrBrs r7visit_AnnAssignzCodeGenerator.visit_AnnAssignsZZ00 DK(( 4:&&  " "$$ F"E"E"EFFF '' )!%(("'DK ;v& &  &&&r8cg}t|tjr|||jgz }n#|jD]}|||gz }t |dkr||d|d}||j}t|s|g}t|s|g}tj f}t||D]p\}}t|}|Dt|s5t||s%tj||j}|||qdS)Nrz2simultaneous multiple assignment is not supported.r)rUrrKrrBrArPrrrdr r&rTrrWr,r-r~r) rvrl_namesrBnamesvaluesnative_nontensor_typesrKrs r7rzCodeGenerator.visit_Assignso dCM * * / tzz$+../ /FF, / /4::f--.. v;;??##D*^__ _q DJ''U## GEV$$ XF"*.!3uf-- ( (KD%(//E #E**!e%;<<!!)33E4<HH NN4 ' ' ' ' ( (r8c4|jj}tj|tj}tj||j|j}tj|jg|}| || |S)Nr?r@) rBrrrLoadBinOpoprrrr)rvrlrKlhsrhsassigns r7rzCodeGenerator.visit_AugAssignsz{~h$CHJJ///iTWdj11T[M=== 6$$T***r8ct|jtjur|jS||jSrT)r^rrrrrrs r7rzCodeGenerator.visit_Name s4 >>SY & &7N$$TW---r8cFtj||dSrTr"rs r7 visit_StorezCodeGenerator.visit_Storer$r8cFtj||dSrTr"rs r7 visit_LoadzCodeGenerator.visit_Loadr$r8cHfd|jD}t|S)Nc:g|]}|Sr;r)r<xrvs r7r=z-CodeGenerator.visit_Tuple..s#111! 1 111r8)r(rc)rvrlrns` r7 visit_TuplezCodeGenerator.visit_Tuples*1111ty111T{{r8ct|r t||||jSt|r6tjdd|}t||||jSt|||S)N_builderz__(.*)__z__r\1__)rZrr~resub)rv method_namervrwreverse_method_names r7_apply_binary_methodz"CodeGenerator._apply_binary_methods S ! ! I,73 ,,S4<HHH H S ! ! Q"$&j+"N"N 473 344S4<PPP P(wsK((---r8cZ||j}||j}|jt |j}|3||d|jj | |||S)Nz8AST binary operator '{}' is not (currently) implemented.) rleftright_method_name_for_bin_oprr^rurformatrjrrvrlrvrwrs r7 visit_BinOpzCodeGenerator.visit_BinOp!sjj##jj$$266tDG}}EE  ##D$^$e$efjfmfv$w$wyy y((c3???r8__add____sub____mul__ __truediv__ __floordiv____mod____pow__ __lshift__ __rshift____and____or____xor__rc |j|||j|j}|j}i}|jr|j|||_i|_||j|j}|j}g}g}g} |D]&} |df|dffD]Q\} } | | vrH| | j || j ks,Jd| d|| j d| d| | j R| |vs| |vr| | | | |vr || j n || j | | |vr|| j n|| j | |vr| |vr || || <| |vr| |vr || || <(t||zD]} | |vr|| j } || j }| |ksJd| d| d |d | | | | | || j ||||||| fS) Nthenelsezinitial value for `z ` is of type z , but the z block redefines it as zMismatched type for z between then block (z) and else block ())r~rUr rrr|rzrryr^rQr3get_typerIkeys)rvrlr{ then_block else_block then_defs else_defsror; ir_ret_typesrKdefs block_namethen_tyelse_tys r7visit_then_else_blocksz$CodeGenerator.visit_then_else_blocks9s 11*=== %%di000\5577 O((**  ; < L 5 5j A A A!,,..DK DO  ) )$+ 6 6 6,,..I99;;J   0 0D&/%89f:M$N X X j4<<:?gdm.@@@@XdXXASXX#-XXFJ4joXXA@@y  DI$5$5 T"""  9J9J4!5!5PYZ^P_Pdeee##$-J.J.IdO$:$C$C$E$E$E3>++inn.>.>>?? C CDu}}o*Go*Gg%%%.t..'..#*...&%% LL      W % % %    $ 6 ? ? A A B B B B)ZUI|[[r8c`t|5}|\}}|j}|j}|j||j|j|||||||\}}}} } |j} |j||r J||j| fd|D|j||r J||j| fd|D| D]} | |  dddn #1swxYwY|j | t|D]T\} }tj | | | | }|||UdS)Nc*g|]}|jSr;r2r<nrs r7r=z4CodeGenerator.visit_if_top_level.. 4X4X4XQYq\5H4X4X4Xr8c*g|]}|jSr;r2r<rrs r7r=z4CodeGenerator.visit_if_top_level..rr8)rsr~r6r7create_cond_branchr3rrV create_branch add_argumentrUrhr corerrpr)rvcondrlsrr{ip_blockrrror;r endif_blockr4rrK new_tensorrrs @@r7visit_if_top_levelz CodeGenerator.visit_if_top_levelpsz d # # -r " GX2244J2244J L 3 3H = = = L + +DKZ P P P++D':zRR YIy*j%L,3355K L 3 3J ? ? ?!0022 C CzO C C2 L & &{4X4X4X4XRW4X4X4X Y Y Y L 3 3J ? ? ?!0022 C CzO C C2 L & &{4X4X4X4XRW4X4X4X Y Y Y" - -((,,,, -) - - - - - - - - - - - - - - -0 11+>>> '' - -GAt!--kooa.@.@)A,OOJ NN4 , , , , - -sFF%%F),F)ct5}|\}}\}}j}|jrjnd} |||| \}} } } }||jfd| D|jd} | | j | t| dkr&j fd| D|js| } n'| | j | t| dkr&j fd| Ddddn #1swxYwYt| D]T\} }t j| | | | }||UdS)NcDg|]}|jSr;rHr~rIs r7r=z.CodeGenerator.visit_if_scf..s'.Z.Z.Z"rxx /E/E.Z.Z.Zr8Trc*g|]}|jSr;r2rs r7r=z.CodeGenerator.visit_if_scf.. -Q-Q-Qail.A-Q-Q-Qr8c*g|]}|jSr;r2rs r7r=z.CodeGenerator.visit_if_scf..rr8)rsrr~r6rrr create_if_opr3merge_block_beforeget_then_blockr7rPcreate_yield_opget_else_blockrhr rr get_resultr)rvrrlrr{rrlast_locrrror;if_oprrKrrrs` @@r7 visit_if_scfzCodeGenerator.visit_if_scfs d # # SrJGQ<<>>LB2244J8< M22444J++D':zRR NIy*j%A  - -b( ; ; ;L--.Z.Z.Z.ZPY.Z.Z.Z\`\gimnnE  ) )%*>*>*@*@ A A A L 3 3E4H4H4J4J K K K5zzA~~ ,,-Q-Q-Q-Q5-Q-Q-QRRR; F"1133 --e.B.B.D.DEEE L 3 3E4H4H4J4J K K K5zzA~~ ,,-Q-Q-Q-Q5-Q-Q-QRRR) S S S S S S S S S S S S S S S,!'' - -GAt!--e.>.>q.A.A9Q<PPJ NN4 , , , , - -sG/HHHc ||j}t|r|tj|j}t|j|}|r5|j r| |d| ||dS| ||dSt|}t|tvr^| |dddtDt|j|r|jn|j}||dS)NrzCannot have `return` statements inside `while` or `for` statements in triton (note that this also applies to `return` statements that are inside functions transitively called from within `while`/`for` statements)O`if` conditionals can only accept values of type {{{}}}, not objects of type {}, c3$K|] }|jV dSrTrrs r7rz)CodeGenerator.visit_If..$!G!G!*!G!G!G!G!G!Gr8)rtestrZtor int1r~rrrrrrrr^_condition_typesrr/rjrrr )rvrlrcontains_return active_blocks r7rzCodeGenerator.visit_Ifsrzz$)$$ T " " 8778=4<7@@D3DK@@FFtLLO .>U++TUUU''d33333!!$-----'--DDzz!111''krr !G!G6F!G!G!GGGT +--... )-=499$+L  ) ), 7 7 7 7 7r8c |||j}t|rZ|tj|j}t|5|\}}|j }|j |tj ||j |j}|j}|j }|j |tj ||j|j}|j}||||j|jksJd|jd|j|j} | tjkr| |jgng} |j| |jd} || | rL|j| |j|jg|j| || | rL|j| |j|jg| r3tj| d| ndcdddS#1swxYwYdSt=|}t#|t>vr^| |d!d"dt>Dt#|j#|r||j S||jS) NrzATernary expression with dynamic condition has inconsistent types r4Trrrc3$K|] }|jV dSrTrrs r7rz,CodeGenerator.visit_IfExp..rr8)$rrrZrr rr~rsrr6rUr,r-rrrrr^r5rHrr3rrr7rrrrrrrrrr/rj) rvrlrrrrthen_valrelse_valr ret_type_irrs r7rzCodeGenerator.visit_IfExps zz$)$$ T " "1 /778=4<7@@D!$''! d! d#@@BB H!\6688  99*EEE#,66tzz$)7L7Ldl[[!\==?? !\6688  99*EEE$,66tzz$+7N7NPTP\]]!\==?? 11"h???} 555{X`Xe{{ltly{{655#=@HHM@Y@Yx~~dl;;<<_a  11+t{DQQ--e.B.B.D.DEEEDL;;E > :##D*tAQ*t*t*tuu u W % % ?'77B''>>> > t'77B'')) ) t t t##rRrrZ^_fZgZgZprrtt t ts B449C-__neg____pos____not__ __invert__rcTt|sJd|dt|sJd|dt|t|ksJd|dt|r.|j|jks Jd|d|jd|jd dSdS) Nzcannot reassign constxpr z in the loopzcannot reasign constexpr zLoop carried variable z changed typezLoop-carried variable z has initial type z but is re-assigned to z: in loop! Please make sure that the type stays consistent.)rWr^rZ)rvrKloop_vallive_vals r7_verify_loop_carried_variablez+CodeGenerator._verify_loop_carried_variable$s))YY+Yt+Y+Y+YYY)))YY+Yt+Y+Y+YYY)H~~h///1]$1]1]1]///$X.. @(-8=2P2P2P @T @ @X] @ @%-] @ @ @3Q2PP @ @2P2Pr8c Z t5}|\}}\}}j}j|j||jj j }| g} g} g} |D]q} | |vrk|| } || } | | || | | | j | |r||jfd| Dd| D}j|fd| Djt%| D][\}} t&j|| |j| <j| j | <\|j}jj|jfdt;t=| Dj|fd| D}j|t%| D][\}} t&j||| |j| <j| j | <\j||jj j }g}|D]!} | |vr||| "j d|Ddddn #1swxYwYt%| D]R\}} t&j|!|| |}|j| <|j | <S|j"D] }JddS) NcDg|]}|jSr;rrIs r7r=z-CodeGenerator.visit_While..Ms(4`4`4`PRRXXdl5K5K4`4`4`r8cg|] }|j Sr;r2r<rps r7r=z-CodeGenerator.visit_While..Ns4U4U4UCSZ4U4U4Ur8cDg|]}|jSr;rrIs r7r=z-CodeGenerator.visit_While..Qs(AmAmAm]_"((4.Ys(:n:n:nST<;K;KA;N;N:n:n:nr8cDg|]}|jSr;rrIs r7r=z-CodeGenerator.visit_While..\s(@l@l@l\^$,AWAW@l@l@lr8cg|] }|j Sr;r2r<ys r7r=z-CodeGenerator.visit_While..ks)C)C)Cq!()C)C)Cr8FzNot implemented)&rsrr~r6rUrrQr rpopr|eraserr^rcreate_while_opcreate_block_with_parent get_beforerhr rrrpryrrr7create_condition_opr3rOrP get_afterrrrrr#r)rvrlrr{rrrdummy loop_defsror; init_argsrKrrwhile_oprr after_blockyieldsnew_defrrs` @r7 visit_WhilezCodeGenerator.visit_While-s d # #= Er$& !G\<<>>LBL--//E L 5 5e < < < N ! !$ ' ' '  ) )$) 4 4 4 N   I KKMMMEII! / /7??(H&t}H66tXxPPPLL&&&$$X]333$$X...  - -b( ; ; ;|334`4`4`4`V_4`4`4`4U4U94U4U4UWWH <@@ATATAVAVAmAmAmAmclAmAmAmooL L 5 5l C C C$U++ : :4$,M$8$89I9I!9L9LiXYl$[$[ D!(, D(9%%::di((D L 3 3L A A A L , ,T[:n:n:n:nX]^abk^l^lXmXm:n:n:n o o o,??@R@R@T@T@l@l@l@lbk@l@l@lnnK L 5 5k B B B$U++ : :4$,M$8$89K9KYWX\$Z$Z D!(, D(9%% N ! !$ ' ' '  ) )$) 4 4 4 N   IF! 3 37??MM)D/222 L ( ()C)CF)C)C)C D D D{= E= E= E= E= E= E= E= E= E= E= E= E= E= E= E@!'' , ,GAtm**8+>+>q+A+A9Q<PPG 'DK $+DOD ! !K 6 6D ++ + +5 6 6sP P++P/2P/c|jjjdksJ||j}||j}t |r|||jS||S)Nrsr) r __class__rjrrslicerZ __getitem__r~)rvrlrvslicess r7visit_SubscriptzCodeGenerator.visit_Subscriptwstx!*f4444jj$$DJ'' S ! ! B??6DL?AA A6{r8c*fd|jDS)Nc:g|]}|Sr;r)r<dimrvs r7r=z0CodeGenerator.visit_ExtSlice..r'r8)dimsrs` r7visit_ExtSlicezCodeGenerator.visit_ExtSlices5555495555r8c |jj}fd|jjD}t fd|jjD}|t jkr||i|}t|j j |j j |j j }|D]g}t|j|jj<|j|jD]"}t(j|#hdSd} d} |t jur,||i|}|j } |j } |j } |j} |j} n|turt3|dkr|dn&t)jd} t3|dkr|dn$|jjd} t3|dkr|dn&t)jd} nt7dd}t9| r&| j dkrt| j } d}| | } } t j| j} t j| j} t j| j} | j !r2| j !r| j !s(tEd | j d | j d | j d t j#| j | j }t j#|| j }|$j}|j%t j&j j'j(k}| j)} | j)} | j)} j*| ||} j*| ||} j*| ||} j+|},|jjt j&-||t]5}|\}}/\}}j0}j1|j23||jj24|5g}g}g}j6D]q}||vrkj6|}||}7||||3||3||3|r8||j9| | | d |D}| .|:d j;| | .|:dj;| j23|j1|<d|=_i_6t}|D]m\}},|t j&-|<d?|dz||j@n|jj24g}j6D]I}||vrC|3t jj6|jJt3|dkr$jAd|D|<dB} | Cdks Jdj1|<d|D}|r6jE| |}jF|| }j|jjj)G|,|jjt j&-||dddn #1swxYwYt}|D]W\}},|t j&-|H|||j@X|jD] }JddS)Nc:g|]}|Sr;rr<rprvs r7r=z+CodeGenerator.visit_For..s#???TZZ__???r8c3BK|]}|VdSrTrr<keywordrvs r7rz*CodeGenerator.visit_For..s/QQ74::g..QQQQQQr8rrr zAOnly `range` and `static_range` iterators are currently supportedFTz0For loop bounds and step must all be ints, are (rrcg|] }|j Sr;r2rs r7r=z+CodeGenerator.visit_For..s>_>_>_csz>_>_>_r8z tt.num_stagesztt.loop_unroll_factorcg|] }|j Sr;r2rs r7r=z+CodeGenerator.visit_For..s-G-G-G1ah-G-G-Gr8z7We use SCF, so the loop body should only have one blockz)Don't know what to do with else after for)Iriterrrndictkeywordsr static_rangerOstartrendstepr ryrBrr rrrr#r num_stagesloop_unroll_factorrPNum RuntimeErrorr\r,r-r~r&r%r3integer_promote_implrHr)rr'r(r3create_int_castrGrrrsrr6rUrrQrrr|rr create_for_opset_attrget_int32_attrget_bodyrzrhrpr^r get_parentsizeget_induction_var create_sub create_addreplace_all_uses_withr)!rvrl IteratorClass iter_args iter_kwargsiteratorr+rrr/r0lbubr. negative_stepiv_type iv_ir_type iv_is_signedivrr{rrrblockrrrorKrrfor_op for_op_regions!` r7 visit_ForzCodeGenerator.visit_Fors<  49>22 ???? ??? QQQQdi>PQQQQQ H1 1 1$}i?;??H !5x|7I8=K^__L! > >.7ll DKN+--di888 K>>DO11$====> F ! HN * *$}i?;??HBB=D!,J!)!<   e # #"%Y!!3!31CGAJJ9O9OB!$Y!!3!31DINSTDU9V9VB#&y>>A#5#59Q<<4::cgajj;Q;QDDbcc c    4:>>dj[))D MB   ( (T\ : :   ( (T\ : : **4>>x   v(9(9 vARARATAT vtrxtt[][cttgkgqtttuu u#8828LL#88$*MM]]4<00 -1D1O1VV Y Y{ \ ) )"j, G G \ ) )"j, G G|++D*lKK \ ' ' 3 3 t{~x}';';B'H'HIII d # #? Nr$& !G\<<>>LBL--//E L 5 5e < < < N ! !$ ' ' '  ) )$) 4 4 4 N   KKMMMIFE , ,7??#t4H&t}H66tXxPPPLL&&&$$X...MM(+++  - -b( ; ; ;\//B>_>_U^>_>_>_``F%1L1LZ1X1XYYY!- 79T9TUg9h9hiii N ! !$ ' ' ' L 5 5fooa6H6H I I I!,,..DK DO$U++ j j4tX]%9%9&//!:L:L:P:PQRUVQV:W:WY_`aYbYg%h%hiiii  ) )$) 4 4 4 N   F d d7??MM("3"="=dod>SUYUa"b"bccc6{{Q ,,-G-G-G-G-GHHH"OOA..99;;M %%''1,,,.g,,, L 5 5fooa6H6H I I I))++B 5\,,R44\,,R44 K  ' . D DR H H H NN4;>8=+?+?G+L+L M M M? N? N? N? N? N? N? N? N? N? N? N? N? N? N? ND!'' ] ]GAt NN4!5!5f6G6G6J6JFSTIN![![ \ \ \ \K 6 6D EE E E5 6 6s"Re  eec||j}||j}||j}t |||SrT)rlowerupperr.r)rvrlrOrPr.s r7 visit_SlicezCodeGenerator.visit_Slice sM 4:&& 4:&&zz$)$$UE4(((r8c6||jSrT)rrrs r7 visit_IndexzCodeGenerator.visit_Indexszz$*%%%r8cD|j||jfSrT)rprrrs r7 visit_keywordzCodeGenerator.visit_keywordsxDJ////r8c||j}|j||jnd}tj|||jS)Nrr)rrmsgr r device_assertr~)rvrlrrWs r7 visit_AssertzCodeGenerator.visit_AssertsSzz$)$$&*h&:djj"""}**4t|*LLLr8rmc|tj|jgRi|fd|jDdDi}dt DfdD}fdt DdD}dD}t |j||}|j|stj g|} |j } t|\} } t|j| | |||j|||j|j| | |jj|jj|jj} | |n3#t.$r&}t1|jj|jd|d}~wwxYw| j}||j|<n |j|}|j|}|j||}|d ks|dS|d kr#tA|!d |Sg}tE|D]>}|#tA|!|||?tI|S) Nc g|] }| Sr;r;)r<rKrns r7r=z2CodeGenerator.call_JitFunction..s444tT 444r8cNg|]"}t|r|nt|#Sr;)rWr rs r7r=z2CodeGenerator.call_JitFunction..s0QQQS',,@)C..QQQr8c6g|]\}}t||Sr;)r\)r<rrps r7r=z2CodeGenerator.call_JitFunction.. s)LLLFAss9K9KLaLLLr8c"i|] }|| Sr;r;)r<rrns r7rz2CodeGenerator.call_JitFunction..!s444AQQ444r8c&g|] \}}|vrdn|SrTr;)r<rrp constexprss r7r=z2CodeGenerator.call_JitFunction..#s)OOOVQZSOOOr8c g|] }||j SrTr2rs r7r=z2CodeGenerator.call_JitFunction..$sBBB3#/CJ///r8c g|] }||j SrTr/rs r7r=z2CodeGenerator.call_JitFunction..%sAAA#SXr8) rrrrrgrrrrrrr)%inspect getcallargsrmrkrhrPrjr has_functionr function_type __globals__rrrrrgr~rrrrr Exceptionrrrirr get_functioncallget_num_resultsrrrOrQrc)rvrmrnrrrAarg_vals arg_typesfn_namerrrrruecallee_ret_typesymbolcall_opresultsrr`s ` @r7call_JitFunctionzCodeGenerator.call_JitFunctions"25:4:::6::4444r|444QQDQQQ LLiooLLL 4444444 OOOOyOOOBB$BBBAAAAA BKI>>{''00 ? .r9==I^F$8$<$< !Iz%dlIvzS\eiep-/wW[Wn/1{i\f.2l.BPTP\Ph151H JJJI  T ++++ T T T&t{ tLLRSS T(0O/>D #G , ,"5g>O))'22,##FH55  " " $ $ ) )_-D4  $ $ & &! + +',,Q//AA AG7224455 R Rvg&8&8&;&;_Q=OPPQQQQ>> !s'E// F9!FFc@t|j}j|}| ||St fd|jD}fd|jD}t|tr(t||| |||St|drt|jstj|rddji}t'j|}d|jvr|d< ||i||S#t,$r!}t/jj|d|d}~wwxYw|jvrt9t|}||i|S)Nc3BK|]}|VdSrTrr$s r7rz+CodeGenerator.visit_Call..Os/DD74::g&&DDDDDDr8c:g|]}|Sr;rr"s r7r=z,CodeGenerator.visit_Call..Pr'r8__self__r _generator)rrr statically_implemented_functionsrr)r*rnrUrrqrthasattrrWrxr r is_builtinr~rc signature parametersrhrrrirrpr0) rvrlrmstatic_implementationkwsrn extra_kwargssigros ` r7rzCodeGenerator.visit_CallIs !$**TY"7"7 8 8 $ E I I" M M ,((t44 4DDDDdmDDDDD555549555 b+ & & 8 4T * * *((T377 7 B # # K(8(E(E K(-JbJbceJfJf K& 5L#B''Cs~---1 \* Kr47<73777 K K K't{dCCJ K '..00 0 0+T22Dr43s0 D:: E%E  E%c*t|jSrTr rrs r7visit_ConstantzCodeGenerator.visit_Constanths$$$r8rlct|jdkr||d||jd}||jd}|jt |j}|3||d|jj | |||S)Nr z^chained boolean operators (A or B or C) are not supported; use parentheses to split the chain.rrz9AST boolean operator '{}' is not (currently) implemented.) rPrprr_method_name_for_bool_oprr^rurrjrrs r7 visit_BoolOpzCodeGenerator.visit_BoolOpks t{  q ##vxx xjjQ((jjQ((377TW FF  ##QXXY]Y`Yijjll l((c3???r8 logical_and logical_orr)c*t|jSrTrrs r7visit_NameConstantz CodeGenerator.visit_NameConstant{sTZ(( (r8c*t|jSrT)r rrs r7 visit_NumzCodeGenerator.visit_Num~sTV$$ $r8cDttj|SrT)r r literal_evalrs r7 visit_StrzCodeGenerator.visit_StrsS-d3344 4r8c||j}t|r2|jdkr'tj|d|jSt||jS)NT)rr)r~) rrrZrr r,permuter~r)rvrlrvs r7rzCodeGenerator.visit_Attributesdjj$$ S ! ! Pdi3&6&6$,,S&$,,OO OsDI&&&r8cFtj||dSrTr"rs r7 visit_ExprzCodeGenerator.visit_Exprr$r8cdSrTr;rs r7visit_NoneTypezCodeGenerator.visit_NoneTyperr8c t|j}t|D]\}}t|tjrt |j||<8t|tjr|j }| |j}t|s3| |dt t|z|dkrdndt|zdz|j||<t!dt|d|S)Nz^Cannot evaluate f-string containing non-constexpr conversion values, found conversion of type rz{}z{!}z:encountered unexpected node of type {} in a JoinedStr noder)rbrprhrUrConstantr*rFormattedValue conversionrr\rr^chrrAssertionErrorr/)rvrlrprrconversion_code evaluateds r7visit_JoinedStrzCodeGenerator.visit_JoinedStrsDdk""!&)) w wHAu%.. w ,,q E3#566 w"'"2 JJu{33 $Y//0++xd9oo../000&5q%8%8TTdSEYEY>Y\_>_gghqhwxxq $%a%h%himnsitit%u%uvvvwwvr8c |dStj5tjdttjdt|j}|j}||_t|dr\t|drL|j |j |j |j z|j |j} t|}nJ#t $rt"$r3}t!|jj|jt)|dd}~wwxYw|r!||_|j ||cdddS#1swxYwYdS)Nignorelineno col_offset)warningscatch_warnings simplefilterDeprecationWarningPendingDeprecationWarningrr~rr{rrrrrsuperrrrhrrir@)rvrl last_noderrOrors r7rzCodeGenerator.visits < F  $ & &    !(,> ? ? ?  !(,E F F F I|++--H DMtX&& 274+F+F 2 $$T^T_t{5RTXTcddd<//11 ZggmmD))#    Z Z Z't{ tAwwOOUYY Z  / )  $$X...1                  s6CE?"!DE?E .EE  'E??FFcx||dt|j)Nzunsupported AST node type: {})rrr^rjrs r7rzCodeGenerator.generic_visits1&E&L<RVZZM`&a&abbbr8c8t|j}d|cxkrdkrnnt|jrtdt ||jd}t |tstd|s|dkrd}nN ||jd}n,#t$r}dt|zdz}Yd}~nd}~wwxYwt|j j |t |dS) Nrr z=`static_assert` requires one or two positional arguments onlyzqAssertion condition could not be determined at compile-time. Make sure that it depends only on `constexpr` valuesrrz')rPrnr*r3rrrUrNotImplementedErrorrhr@rrri)rvrl arg_countpassedrros r7execute_static_assertz#CodeGenerator.execute_static_asserts< NN I"""""""""s4='9'9"[\\ \%djj1&>&>??&$'' %D  dA~~X"jj166GG XXXG$q''QTWWGGGGGGX.dkotEYZaEbEbcc cts% C C/C**C/c,dtjffd }|S)Nrlcdfd|jDD}fd|jD}t|i|S)Nc4i|]\}}|t|Sr;)r)r<rKrs r7rz>CodeGenerator.static_executor..ret..s7D%*511r8c3BK|]}|VdSrTrr$s r7rz=CodeGenerator.static_executor..ret..s/#U#UGDJJw$7$7#U#U#U#U#U#Ur8cTg|]$}t|%Sr;)rrr"s r7r=z>CodeGenerator.static_executor..ret..s,OOOc(C99OOOr8)r*rnr )rvrlrrn python_fns` r7rOz*CodeGenerator.static_executor..retss#U#U#U#Ut}#U#U#UCPOOOTYOOODYY44455 5r8)rr)rrOs` r7static_executorzCodeGenerator.static_executors1 6CH 6 6 6 6 6 6 r8rz)NFNFNr)rjrrrrrr*rwrPrbrOfloatintrUrrrrupdater r device_printminimummaximumrrrr rr rrrr rr)rrrgrirlrrrr{r}rrrrAddSubMultDivFloorDivModPowLShiftRShiftBitAndBitOrBitXorrroperatorrrrrrrrEqNotEqLtLtEGtGtErcmpoprUSubUAddNotInvertrunaryoprrrrrMrQrSrrUrYrtrrBoolOprAndOrrboolopsys version_inforrrrrrrrrrrr static_assert static_printrrzrr __classcell__)rs@r7rrsjnMN0000al00X`aeXf00,4SM00000000d)u(tdESXZ]_ikr@s(t(t(ttCH~ttt (-,-  !  ! LLL   ---^&c&% 0A*B&t&&&&"""   222@@@2H?H?H?T&&&'''"(((2+++... 222222...@@@   )  n   L L I 8 I >T$s|"4c"9:   5\5\5\n---@---688863/3/3/j   @@@$ #)Xsvx(TWTZ\dfifmow<d4 ?C#78 t t t )SXy#'9cjR^?tD$5s$:;@@@H6H6H6T666E6E6E6N))) &&&0U38_0000MCMMMM -";-"-"-"-"^   >%%% @ @ @ @ @>AWmUXU[]izast_to_ttir..s+As1C1CJ **1---r8c.i|]\}}||Sr;r;)r<keyrcst_keys r7rzast_to_ttir..s'XXXeuXXXr8i1rTcFg|]\}}|jvt|Sr;)rAr)r<rrrs r7r=zast_to_ttir.. s1ppp$!QaWeWoNoNo1NoNoNor8) rrArrrrrrrrr)attrsrArrgrzr@rbr}rp get_constantsfilter_out_constants get_fn_attrsrrr rfrrrrr)rmrrrrrrrArrtys new_constantsr new_attrsfn_attrs all_constantsrmrrrrurOrs`` @r7 ast_to_ttirrs  EJJJJGXXXX~7O7U7U7W7WXXXI ^ " "FGGN++M ~'..00 1 1C''))M $$ 88A$=+;q+@+@#M! **,,I%%''HNN$$M'''pppp.*B*H*H*J*JpppI044Iz&r955Igy=hu%'HXa)3WR]jtvvvIOOBHHJJ  CCK Jr8)7rrcrrrr rtypingrrrrrrr rr _C.libtritonr r rr language.corerrr runtime.jitrrruntimererrorsrrrtypesrr#rPrrWrZr\r`rdrqrr^rrsr#rrrrr;r8r7rs   DDDDDDDDDDDDDDDDDD3333333333JJJJJJJJJJ========!!!!!!aaaaaaaaaa...$   !!!!!!!!!!!!$S$T$$$$QQQQQQ(S(T((((#ttDzz*33333333*U%U%U%U%U%COU%U%U%piiiiiCOiiiX!   r8