K i^ ddlmZmZddlZddlZddlZddlZddlZddlZddl Z ddl Z ddl m Z ddl mZddlmZddlmZmZmZmZmZmZmZmZmZmZddlmZddlmZd d l m!Z!d d l"m"Z"d d l m#Z#d dl$m%Z%m&Z&m'Z'm(Z(d dl)m*Z*ddl+m,Z,dZ-dZ.edZ/Gddej`Z1d6dZ2GddZ3iZ4gZ5dZ6d7dZ7Gddee/Z8dZ9dZ:dZ;Gd d!ZGd%d&e intlenmaxminlistfloatprintrangegetattr isinstancecopymathF)super__init__namehashlibsha256encodehasherglobals nonlocalssupported_python_builtins GLUON_MODULE TRITON_MODULEsupported_modulesused_global_valsvisiting_arg_default_value)selfr3r8r9src __class__s X/mnt/ssd/data/python-lab/Trading/venv/lib/python3.12/site-packages/triton/runtime/jit.pyr2zDependenciesFinder.__init__.st  nnSZZ%89  " * &     " "TV*/'c6|jjSN)r7 hexdigestr@s rCretzDependenciesFinder.retYs{{$$&&rDctj|jryt|dd}|j t S)NT __module__)inspect isbuiltinfuncr- startswithr<)r@noderOmodules rC_is_triton_builtinz%DependenciesFinder._is_triton_builtin]s6   TYY '|R0  //rDc>t|tsJ|jj|jjzD]_}|\}}|j|\}}|j|\}}||k7s2t d|d|d|j d|j d|d |jj|j|j}|tt|ddz }|jj|jd y) NGlobal variable z has value z when compiling z, but inner kernel z has conflicting value z7 from when it was first compiled. This is not allowed.noinlineFr$) r. JITCallabler>keys RuntimeErrorr3__name__update cache_keystrr-r7r6)r@rOkvar_name_v1v2func_keys rC _update_hashzDependenciesFinder._update_hashcs/$ ,,,&&++-0E0E0J0J0LL AKHa))!,EB))!,EBRx"&xj B4?OPTPYPY{Zmnrn{n{m|}TUWTXXOP   $$T%:%:;>>Cj%899 8??734rDcddlm}|t|turyt |ddryt |dddk(ryt |t r|j|yt|r*t |tst ||std||jry|/tj||f|j|t|f<y) Nr constexpr__triton_builtin__FrKrLztriton.language.extra.libdevicez!Unsupported function referenced: ) language.corergtyperr-r.rWrdcallablerYr?r/deepcopyr>id)r@valvar_dictr3rgs rCrecord_referencez#DependenciesFinder.record_referenceus- ;$s)z1  3,e 4  3 b )-N N  c; '   c "  C=C!6z#y?Y!B3%HI I  * *   ;?==;Mx:XD ! !4H"6 7rDc>t|jtjur |jS|jj vryfd}||j\}}|jj vr|Sj|||j|S)Ncjj|d}||jfSjj|d}||jfSy)NNN)r8getr9)r3rnr@s rC name_lookupz2DependenciesFinder.visit_Name..name_lookupsZ,,""4.CDLL((..$$T40CDNN**rD)rjctxastStorerm local_namesr:rp)r@rQrurnros` rC visit_NamezDependenciesFinder.visit_Names >SYY &77N 77d&& & $DGG, X 77d44 4J c8TWW5 rDc^|jDcgc]}|j|c}Scc}wrF)eltsvisit)r@rQelts rC visit_TuplezDependenciesFinder.visit_Tuples$,0995C 3555s*cf|j|j}t|tjr6|j|j}t|tjr6t |dd}|||j vryt ||j}|j||S)NrZrL) r}valuer.rw Attributer-r=attrrp)r@rQlhslhs_namerIs rCvisit_Attributez"DependenciesFinder.visit_Attributesjj$cmm,**SYY'Ccmm,3 B/ ;(d&<&<<c499% c" rDc|jjDchc]}|jc}|_|j|ycc}wrF)argsargry generic_visit)r@rQrs rCvisit_FunctionDefz$DependenciesFinder.visit_FunctionDefs4/3yy~~>CGG> 4 ?sAcpfd}tj|j|j|jr |jgng|j D]}j |||j|jj |j||jy)Nc jrJd_|D]}|j| d_y#d_wxYw)NTF)r?r})defaultsexprr@s rCvisit_defaultsz:DependenciesFinder.visit_arguments..visit_defaultssS 8::::26/$)D' 4()38/%/s << A) itertoolschain posonlyargsrvararg kwonlyargsr} kw_defaultskwargr)r@rQrrs` rCvisit_argumentsz"DependenciesFinder.visit_argumentss 8??4#3#3TYYQUQ\Q\ bdfjfufuv C JJsO  t''( :: ! JJtzz "t}}%rDc|j|}t|tr|xjt |zc_y|jj |yrF)r}r.r)rysetadd)r@rQtargets rCvisitAssnTargetz"DependenciesFinder.visitAssnTargetsED! fd #   F +      (rDct|jdk7r td|j|jd|j |y)Nrz2Simultaneous multiple assignment is not supported.r)r&targets TypeErrorrrr@rQs rC visit_AssignzDependenciesFinder.visit_AssignsG t||  ! PQ Q T\\!_- 4 rDc\|j|j|j|yrFrrrrs rCvisit_AnnAssignz"DependenciesFinder.visit_AnnAssign$ T[[) 4 rDc\|j|j|j|yrFrrs rC visit_ForzDependenciesFinder.visit_ForrrD)returnNoners)rZrK __qualname____doc__r2propertyrIrSrdrprzrrrrrrrr __classcell__rBs@rCr"r""s` )0V''0 5$ D06 ! &@) !! !rDr"cddlmcm}t|tr|j }|j dr7|jd}t|}|j dsJd|ddzS|jdrdt|ddzS|j drdt|ddzS|j drt|jdSt||jrdt|jSt||jr |j}n(t|tr |j}n t |}t!j"|j%d d |S) Nrzconst const**krztl._trL)triton.language.corelanguagecorer.r]striprP removeprefix _normalize_tyendswith pointer_type element_tydtyper3rjrZrrtreplace)tyrs rCrr s?''"c XXZ == ")Br"B==% %%"QR&= ;;s r#2w// / == r!"v.. . ==  !78 8 B)) *=/011 B # WW B  [[ W % ) )"**T2*> CCrDceZdZdZ d dZedZed dZed dZedZ edZ e dZ e d Z y ) KernelParamzBRepresents a parameter (name plus metadata) to a @jit'ed function.c<||_||_||_||_yrF)num_paramdo_not_specializedo_not_specialize_on_alignment)r@rparamrrs rCr2zKernelParam.__init__)s  !2.L+rDc.|jjSrF)rr3rHs rCr3zKernelParam.name0s{{rDc|jjr1|jjtjjk(ryt |jjS)NrL)r annotationrM ParameteremptyrrHs rCrzKernelParam.annotation4sD{{%%)?)?7CTCTCZCZ)ZT[[3344rDc|j}|jdr|dd}n|jdr|dd}|ttjvr |jSy)NrrrrrL)rrPrrvalues)r@as rCannotation_typezKernelParam.annotation_type:s] OO << !"A \\# !"A .5578 8?? "rDcd|jvSNrg)rrHs rC is_constexprzKernelParam.is_constexprEsdoo--rDcr|jryd|jvxs|jjdS)NFrr)rrrPrHs rCis_constzKernelParam.is_constIs1   $//)MT__-G-G-MMrDc.|jjSrF)rdefaultrHs rCrzKernelParam.defaultOs{{"""rDcd|jjtjjk7SrF)rrrMrrrHs rC has_defaultzKernelParam.has_defaultSs#{{""g&7&7&=&===rDN)rr%rzinspect.Parameterrboolrrrr])rZrKrrr2rr3rrrrrrrrDrCrr&sLM15M  55 ..NN ##>>rDrc8ddlmddlmdfd S)Nrrfrrcyttryttr7|r d|nd}dk(r|rydkr dkrd |fSd kr d krd |fSd |fSttryt drZj |f}t j|d}|!|drdndt|dz}|t |<|r d|nd}||fSttrdjfStrdfSttrPDcgc] }| }}fd} | |Dcgc]}|d c}} | |Dcgc]}|d c}} | | fSttrTt jdsJtjj } d| tjddfSt rat jdsJtjj } d| tjdj ddfSt#dt%zcc}wcc}wcc}w)N)rgN)u1Nr%)alignr)rgriii32llu64i64)fp32Ndata_ptrrrrtensorrgcNtdrt|St|S)N_fields)hasattrrjtuple)valsrs rCzAcreate_specialize_impl..specialize_impl..s&'#y:Qid3i&6W\]aWbrDz tensordesc<>,zUnsupported type: %s)r.rr%r*rr dtype2strrtrrWr\rrbaser) block_shapelayoutrrj)rrspecialize_valuerkeydskresxspec make_tupletysrXinnerGluonTensorDescriptorrgspecialize_extraspecialize_impls` rCrz/create_specialize_impl..specialize_implas\ ;& T " S !?O"3U;UYCax,'SSI%5s|###"2s|#s|# U #! S* %99h'C--T*C{"1vt32DSV2LL!$ #BR"3>X\C:  [ )/ / Y '% % U #0341OA&4D4bJD1qad12CT2qt23D;  - .388Z0 00&sxx~~6E!%coo)>(?qA4H H 2 3388Z0 00&sxx~~6E!%coo)>(?qaPRVW W2T#Y>? ?512sI 3 I I)FTT)rrg'triton.experimental.gluon.nvidia.hopperr)rrrgrs`@@@rCcreate_specialize_implr\s$a-@-@^ rDcttdk(rtjtdtd}|||dS)NrcyrFr)r`kwargss rCrzmangle_type..srD)r)r&specialize_impl_cacheappendr)r specializers rC mangle_typer sA !Q&$$%;.sxtxx$T%'YRX'YrDr)r@rs``rC __getitem__zKernelInterface.__getitem__s ZYrDN)rr )rZrKr__annotations__rrrDrCr r s FZrDr c |jDcic],\}}||jjdk(r t|n|.}}}ddl}|||j Dcgc] }t |c}t |j|j Dcgc] }t |c}t |j|jd} |j| } | Scc}}wcc}wcc}w)Nrr)r3 signature constant_keys constant_vals attrs_keys attrs_valsoptionsr) itemsrBrZr]jsonrXr)r__dict__dumps) r3r constantsattrsrrrrrobjserialized_objs rCserialize_specialization_datar#senetetevwWaWZ\aEOO$<$<$Gc%jURwIw9QZQ_Q_Qa?bAQ?b Y    %**,0OQa0O_cdidpdpdr_s##C C ZZ_N x@c0Os1C, C2C7 c t|jt|k(sJg}t|jj|D]\}}|jr|j d|d)|j rdnd}|jrdnd}|jrdnd}d|d|d|d|d } |jrt|jtr"|jdk(s|jdd d vrd }|r#|j d |jd | d|j d |jd|j |  d} ddjtt| |jjdgzddj|jjDcgc] }d|d| c}ddj|d} |jjD cic];\}} | j t"j$j&urd|| j =} }} t(| d<t+|j,| d<t/| | | dScc}wcc} }w)a2 Equivalent to sig.bind followed by apply_defaults. This generates a native Python function (using exec) which can be memoized on a per-kernel basis to avoid having to run these expensive functions -- which constitute much of the kernel launch overhead -- every time we run the kernel. z("constexpr", )TrueFalsezspecialize_impl(, rNr)fpbfFz("z",) + z[1:]z", None)cx|djtjjur|dS|dd|dS)Nrrz =default_)rrMrr)rs rCrz0create_function_from_signature..sAAaDLLG,=,=,C,CCAaDAaD6QZ[\]^[_Z`IarDz def dynamic_func(z **optionsz): params = {'z': z} specialization = [rz-] return params, specialization, options default_rWr dynamic_func)r& parametersziprXrrrrrrr.r]joinr)maprrrMrrrWrget_arg_specializationexec)sigkparamsbackendspecializationr3kprrrrIr func_bodyrfunc_namespaces rCcreate_function_from_signaturer<s s~~ #g, .. .N++-w70b ??  ! !N4&": ;!#v'H$&$8$8fJ!@@GfE$TF"XJb BugQOC!!b00#6))T1R5G5G5K|5[%* "))Br/A/A.B&T*RS#))Br/A/A.B(*KL%%/'0, bC))DS#..*>*>*@!ABk]RSTU 3>>;N;N;PQ4QtfCv.QRSTxx/01I>>//1 D% == 1 1 7 7 7 4&5==(N %0N=!(>w?]?](^N$% N# . ))%R s I=AJc8|jd|jS)N.)rKrfns rC get_full_namerAsmm_Aboo. //rDcdeZdZdZdZedZdZedZdZ dZ dZ ee e Z y ) rWc||_tj||_ tj|\|_|_t||_ tj|_ tjdj|j}|t!j"d|t j$j'd}||_d|_i|_|j.|_|j0|_|j2|_|j4|_|j6|_y#t $r}td|d}~wwxYw)Nz1@jit functions should be defined in a Python filerLz^def\s+\w+\s*\()r@rMrgetsourcelinesraw_srcstarting_line_numberOSError ValueErrorrA_fn_name threadingRLock _hash_locktextwrapdedentr1research MULTILINEstart_srchashr>rrZr __globals__rK)r@r@erAs rCr2zJITCallable.__init__s **2. Y6=6L6LR6P 3DL$3&b) #//+oobggdll34")).R\\BHHJKL  TVzz   OO>>--7 YPQWX X Ys"D== E EEcn|jtj|jjzSrF)rUrMgetclosurevarsr@r9rHs rCget_capture_scopezJITCallable.get_capture_scopes('"8"8"A"K"KKKrDc|j5|j|jcdddSd|j|_tj|j j }t|j|j||j}|j|j|jt|jz|_tt!|j"j%|_ddlm}|xjt|j"j%Dcgc]\\}}\}}t+||r||fc}}}z c_t-j.|jj1dj3|_ddd|jScc}}}w#1swY|jSxYw)Nz recursion:)r3r8r9rArrfr$)rLrTrIrMrXr@r9r"rUrAr}parserIr]rFdictsortedr>rrrgr.r4r5r6rG)r@r9dependencies_finderrgr3r`rns rCr\zJITCallable.cache_keys__ Nyy$yy N N %T]]O4DI..tww7AAI"4$--QUQaQamv9=#C   % %djjl 3+//#d6O6O2PPDI$(0C0T0T0Z0Z0\)]$^D ! 6 II9=9N9N9T9T9V=="5)4Xc1!+C!; $Sk=> >I tyy'7'7'@AKKMDI# N$yy = N$yys$GDG G %AG GG%ctj|j}t|tjsJt |j dk(sJt|j dtjsJ|S)Nrr)rwr[rSr.Moduler&body FunctionDef)r@trees rCr[zJITCallable.parse2s_yy#$ +++499~"""$))A,888 rDcddlm}||S)Nr)constexpr_type)rre)r@res rCrjzJITCallable.type9s7d##rDc d|_||_y)a" The only method allowed to modify src. Bypasses the __setattr__ restriction by calling super().__setattr__ directly. Note that it is the callers responsibility to make sure any triton functions that call this function have the `.hash` value reset to None. N)rTrS)r@new_srcs rC_unsafe_update_srczJITCallable._unsafe_update_src>s  rDctd)NzqCannot set attribute 'src' directly. Use '_unsafe_update_src()' and manually clear `.hash` of all callersinstead.)AttributeErrorrHs rC_set_srczJITCallable._set_srcHs() )rDc|jSrF)rSrHs rC_get_srczJITCallable._get_srcMs yyrD)fgetfsetN) rZrKrr2rYrr\r[rjrhrkrmrArrDrCrWrWsX (DL2$$)  x 0CrDrWc,eZdZUded<ded<ded<y)JitFunctionInforrRr]r3 JITFunction jit_functionN)rZrKrrrrDrCrqrqSs  IrDrqct|t|f}|j|d}||St|t|z}|||<|SrF)rr]rt)kernel_key_cacher8rrr\s rCcompute_cache_keyrvZsW  #g, /C $$S$/IN#c'l2I%S rDcreZdZdZ ddZdZdZdZdZdZ dfd Z d Z d Z d Z d Zd ZxZS)rrcy)NFrrHs rCis_gluonzJITFunction.is_gluongsrDc |sy|jj} |jj} djt |j |dD cgc]\} } | j d| c} } } | d|jd|jd|jd|jd|jd | d }t|j}t||||d ||}||||j|j|j|j|j|j|||d }|||t| | |d |i||dScc} } w)Nr(rz: z [num_warps=z , num_ctas=z , num_stages=z, enable_fp_fusion=z, launch_cooperative_grid=](r%r) rdevicer num_warpsnum_ctas num_stagesenable_fp_fusionlaunch_cooperative_grid extern_libsconfigsspecialization_data is_warmuprF)rreprr@compileis_manual_warmupalready_compiled)r@rrKr1r0paramsr3r}r~rrrrAr#rrq)r@hookrrr|rrrrr3rRrr arg_reprsr full_namerrs rC _call_hookzJITFunction._call_hookjsww####IIc$++WZ[\W]F^_%**Rt4_` {7#4#4"5[AQAQ@RR_`g`r`r_stGHOH`H`Gaa{|C|[|[{\\^_h^iijk!$''* ;IyR[]def]gipruv#" **((!,, ' 8 8'.'F'F"..#6"  vtT2C*6*&"   +`sE cTt|sJ|jj|y)z Add a hook that will be executed prior to the execution of run function with args and kwargs passed into the kernel N)rk pre_run_hooksr)r@rs rCadd_pre_run_hookzJITFunction.add_pre_run_hooks$ ~~ !!$'rDcddlm}m}m}m}t j j}||}||_||_||_t|j|j|}ii|||fS)z1 Precompute as much as possible. r)CompiledKernelr ASTSource make_backend) compilerrrrrractiveget_current_targetr<rr)r@rrrrrr7binders rC create_binderzJITFunction.create_bindersd PO113v&, "/ WU2vw..rDc |j|}|jDcgc]}|j}}|Dcgc]}|d }}t||D cic]\} } | |  } } } d|vsJdd|vsJdd|vsJd|D]#} | |jvs| |vst d| zt |d } | D cic]&} | tt|j| (} } |Dcgc]}|d  }}t |d }|D cic]} | |jt||  }} || | |fScc}wcc}wcc} } wcc} wcc}wcc} w) Nr device_typez=device_type option is deprecated; current target will be usedr|z8device option is deprecated; current device will be usedstreamz8stream option is deprecated; current stream will be usedz2Keyword argument %s was specified but unrecognisedc |dk(Srr)r`rns rCrz(JITFunction._pack_args..s 3+;MrDrc"t|tSrF)r.r])r`rs rCrz(JITFunction._pack_args..s Z35GrD) parse_optionsrr3r0rKeyErrorrrr)r parse_attr)r@r7r bound_argsr8rrsigkeyssigvalsr^vr constexprspathattrvalsr s rC _pack_argszJITFunction._pack_argss''/#';;/a166//!/0A1Q400(+GW(=>fq!QT> >F*k,kk*v%a'aa%v%a'aa% YA(((Qg-=SVWWXX Y#7,MN [efSWd-d:3D3D3F.GNNf f"01QAaD11h(GHPUV1G&&'81'EFFVV :u44%00>g1Vs"D; E E+E 1 E#Ec |jd|jxstjj|d<tj j }tj j|}|jD] }||i| |j|\}} } } } | |i|\} }}t| ||}|j|d}|4|j| || ||\}}}}|j|||||||}|yt}|jjD]6\\}}\}}|j||x}|k7s$t!d|d|d||s|Jt#|r|| }t%|}|d}|dkDr|dnd}|dkDr|dnd}t'|dr|j)}|j*||g| j-}|j.|||||j0|j2|tjj4tjj6g | j-|S) NdebugrUz1 has changed since we compiled this kernel, from z to rrrresult)rtrrruntimerrget_current_deviceget_current_streamr device_cachesrvr _do_compileobjectr>rrYrkr&rrlaunch_metadatarr functionpacked_metadatalaunch_enter_hooklaunch_exit_hook) r@rrrrr|rr kernel_cacherurr7rrr8rrkernelrrr  not_presentr3r`rn globals_dictnewVal grid_sizegrid_0grid_1grid_2rs rCr zJITFunction.runs **Wdjj9PU]]=P=Pw11311&9&& "D $ !& ! "CGBTBTU[B\? &/5d.Ef.E+ NG 0.'J!!#t, >48OOGVU_aoDK5M 1GY E%%c9fj'SXZ`aF~h .2.C.C.I.I.K q *IT1*\&**4==#E"&tf,]^a]bbfgmfnoqq q # ##~J'D I!WF )A T!W1F )A T!W1Fvx(4f44T6XJDUDUDWXO FJJvvvvvH^H^`o}}66 8V8V nYcYjYjYl n rDcT|j |jS|j|SrF)_reprrI)r@r`s rCrzJITFunction.reprs" $ 2t}}E 1 ErDc |r|ng}|r|ng}t|||j|_||_||_||_||_||_g|_ t|jjjD]T\} } | |vxs| j|v} | |vxs| j|v} |jjt!| | | | Vt#|j$|_d|_||_||_|jD cgc]} | jc} |_|jD cgc]} | j0s| j2c} |_g|_ycc} wcc} wrF)r1r2rKrRversionrrrrr enumeraterr/rr3rrrrrrrrV arg_namesrrrr)r@r@rrrrrVrrirdnsdns_oaprBs rCr2zJITFunction.__init__sW1B-Ki)Goq& mm  !2.L+ . !$..";";"B"B"DE CHAu((KEJJ:K,KC88hEJJJhv.F-GGbcgcpcpbqrt tE#3O#DE (9 "-? U BHH$5$5e$<%% G   0 >? %l3 SZ01)+6<<>? /y9??A U E4!8ue C  u%"0081a!''0           s 9F,(F c Bj\}} }jtjjgryj t jj} | Ctt| } fd} f d} | j| | | }|Sjj}|<jtjjg|S)NcBjjS)N)rr _env_vars)rr)env_varsrr@rArsrC async_compilez.JITFunction._do_compile..async_compileSs!||C@P@P\d|eerDc r |<jtjjg yrF)rrrjit_post_compile_hook) rr rr|rrrr@rrs rCfinalize_compilez1JITFunction._do_compile..finalize_compileVs9$* S! C CS)U[]gip!&1rD)rr)rrrrjit_cache_hookrr active_modertrrsubmitrrr)r@rrr|rrr rr`r7 async_moder\rrrrrrArs```````` @@@@rCrzJITFunction._do_compileFs .2.@.@.H+ a! ??5==77iQ[]dglfmou vnnT9j%@#//335  !68H%c7GXFI f f 1 1  &&y-AQRF  \\#fg>N>N\OF &L  OOEMM??iQWYcelotnu" $ rDctd)Nz:Cannot call @triton.jit'd outside of the scope of a kernel)rYr@rrs rC__call__zJITFunction.__call__csWXXrDcPd|jd|jjdS)Nz JITFunction(:r%)rRr@rrHs rC__repr__zJITFunction.__repr__fs&dkk]!DGG,@,@+ACCrD)rz bool | None)NNNNNNN)rZrKrryrrrrr rr2rrrrrrrs@rCrrrres`,  , \( /502hFmq;?" H]! F:YDrDrrcyrFrr?s rCjitrosrDrrrrrrrVcyrFrrs rCrrtsrDc@dfd }|||S|S)a< Decorator for JIT-compiling a function using the Triton compiler. :note: When a jit'd function is called, arguments are implicitly converted to pointers if they have a :code:`.data_ptr()` method and a `.dtype` attribute. :note: This function will be compiled and run on the GPU. It will only have access to: * python primitives, * builtins within the triton package, * arguments to this function, * other jit'd functions :param fn: the function to be jit-compiled :type fn: Callable c t|sJtjjrddlm}||St |S)Nr)InterpretedFunction)rrrrrVrr)rkrr interpret interpreterrrr) r@rrrrrrVrrs rC decoratorzjit..decoratorsj|| == " " 8&r7N_Fdlq08tUdf f"3/M! /  rDr@r rzJITFunction[T]r) r@rrrrrrrVrs ``````` rCrrs&:& ~}rDcNeZdZdZedZddZdZedZedZ y) rzr Can be used in place of real tensors when calling: kernel.warmup(MockTensor(torch.float32), ...) cl|jjdk(r|jdk(r t|S|S)Nrtorch)rBrZrKr)rs rCrzMockTensor.wrap_dtypes. == ! !W ,71Jc? " rDNc*|dg}||_||_y)Nr)rshape)r@rrs rCr2zMockTensor.__init__s =CE  rDcdg}|jddD]}|j|d|ztt|S)Nrr)rrrreversed)r@stridessizes rCstridezMockTensor.stridesG#JJqrN /D NN72;- . /Xg&''rDcyNrrrrDrCrzMockTensor.data_ptrrDcyrrrrDrC ptr_rangezMockTensor.ptr_rangerrDrF) rZrKrr staticmethodrr2rrrrrDrCrrsM   ( rDrcJeZdZdZdZdZd dZdZdZdZ dZ d Z d Z y ) TensorWrapperc||_||_|j|_|j|_|jj|_yrF)rrdatar|r)r@rrs rCr2zTensorWrapper.__init__s5  II kk YY__ rDc6|jjSrF)rrrHs rCrzTensorWrapper.data_ptrsyy!!##rDc4|jj|SrF)rr)r@rs rCrzTensorWrapper.stridestyy&&rDc<d|jd|jdS)NzTensorWrapper[r{r%)rrrHs rC__str__zTensorWrapper.__str__s |2dii[::rDc6|jjSrF)r element_sizerHs rCr zTensorWrapper.element_sizesyy%%''rDc^t|jj|jSrF)rrcpurrHs rCrzTensorWrapper.cpusTYY]]_djj99rDcN|jj|jyrF)rcopy_)r@others rCrzTensorWrapper.copy_s  #rDc^t|jj|jSrF)rrclonerrHs rCrzTensorWrapper.clonesTYY__. ;;rDc`t|jj||jSrF)rrtor)r@r|s rCrzTensorWrapper.tos TYY\\&14::>>rDc`t|jj||jSrF)rr new_emptyr)r@sizess rCrzTensorWrapper.new_emptys"TYY007DDrDNr) rZrKrr2rrr r rrrrrrrDrCrrs5%$';(:$<?ErDrct|tr;||jjk(r |jSt|j|St |dr t||St dt |d)NrzCannot reinterpret a r>)r.rrrrrrj)rrs rC reinterpretrsk&-( FKK%% %;; !e4 4  $VU++/V ~Q?@@rDcL|}t|ts|j}t|ts|jjj}|j }t |jD].\}}|jjds&||z }||fS||fS)Nzdef ) r.rWr@__code__ co_filenamerFrrErrP)r@base_fn file_name begin_lineidxlines rCget_jit_fn_file_liner#sG+.**+. ##//I--J w/ T ::< " "6 * # J  j   j  rDceZdZdZdZy)BoundConstexprFunctionc ||_||_yrF)__self____func__)r@instancer@s rCr2zBoundConstexprFunction.__init__'s   rDcB|j|jg|i|SrF)r(r'rs rCrzBoundConstexprFunction.__call__+s!t}}T]]J!Q!)!,,JJdggt&v&  J == " "J~7Js BB)rZrKrr2r0rrrs@rCr,r,/s )-rDr,ct|S)z Wraps an arbitrary Python function so that it can be called at compile-time on constexpr arguments in a Triton function and returns a constexpr result. )r,r?s rCconstexpr_functionr5Ms R  rDr)Fr)rOptional[Callable]rr6rOptional[Iterable[int | str]]rr7rOptional[bool]rVr8rzCallable[[T], JITFunction[T]]rF)r@z Optional[T]rr6rr6rr7rr7rr8rVr8rz4Union[JITFunction[T], Callable[[T], JITFunction[T]]])H __future__rrrwr/r4rMrrJrOrM collectionsr dataclassesr functoolsrtypingrr r r r r rrrrtriton.tools.tensor_descriptorrtypesrrLrrr_utilsrrrrcachertriton._C.libtritonrr<r;r NodeVisitorr"rrrrrr r r#r<rArWrqrvrrrrrrr#r%r,r5rrDrCrDs>,  #!%ddd;ee ?! 3  CLb!b!TD4/>/>d  4n@ Zgaj Z 7*t0_1_1D   BD+q1BDT    #*.7;DH #   (  5  %B    #   4 #*.7;DH #44  4 ( 4 5 4%B4 44:4xB"E"EJ A!$=[= <!rD