tj,UddlmZddlZddlZddlZddlZddlZddlZddlZddl Z ddl Z ddl Z ddl Z ddl Z ddlZddlZddlZddlZddlZddlZddlZddlZddlZddlZddlZddlmZmZmZmZmZmZm Z ddl!m!Z!ddlm"Z"ddl m#Z#ddl$m%Z%m&Z&m'Z'm(Z(m)Z)m*Z*m+Z+m,Z,m-Z-m.Z.m/Z/ddl0m1Z1m2Z2m3Z3dd lm4Z4ddl5Z5ddl6Z6ddl7m8cm9Z:dd l;mZ>dd l?m@Z@dd lAmBZBddlCmDZDddl7mEZEmFZFddlGmHZHddgZIddlJmKZKddlLmMZMmNZNmOZOmPZPmQZQe,rhddlmRZRmSZSmTZTddlUmVZVddl6mWZWmXZXmYZYddlZm[Z[ddl\m]Z]ddl^m_Z_ddl`maZaddlbmcZcddldmeZeddlfmgZgdd lhmiZidd!ljmkZkmlZlmmZmmnZnmoZompZpdd"lqmrZrdd#lsmtZtmuZugd$Zve/d%Zwejxdd(Zydd)lzm{Z{dd*l|m}Z}dd+l~mZdd,lmZdd-lmZdd.lmZdd/lmZmZmZmZmZmZdd0lmZmZdd1lmZmZdd2lmZdd3lmZejd4kZe jeZd5Ze/d6Zee5je5jfZe6jeze6jYzdzZerd7nejd8d9Zd:d;dZd>Zd>Zd?ZeDe6je6je6je6je6je6je6je6je6je6je6je6je6je6je6jgZd@edA<dBZeedz zdkredCks JdDddGZddKZGdLdMe5jZejdNOGdPdQZdddYZ ddd\ZĐddcZ ddddZejxddeZǐddiZȐddlZɐddpZʐddsZːddvZ̐ddyZ͐dd}ZddZΐddZϐddZАddZѐddZҐddZdfddZԐddZՐdddZ dddZ dddZؐddZِddZڐddZېddZܐddZe2dZe/ddNZee'e%efefZGdde+e(eefZddĄZddńZddȄZdd̄ZddЄZddԄZ dƐddڄZdd݄Zdd߄ZddZddZddZddZddZddZddZddZddZegde6jjdgngzZddZddZddZddlZddZgZded<ddZddZejdאdZej dؐdِdZeZeZeZdZd dڐd ZdZd dېdZej"dCdܐdZGdde*Z ejGddZ GddZ Gdde Z Gdde Z dݐdZGd d!e Zejdސd"ZGd#d$ZGd%d&eZejxdߐdd)Zej"dd*Zej"dd+Zdd,Zdd-Z dƐdd2Zddd4Zdd9Zdd;Zdd<Zdd=Zdd>ZdZdZdNd?ddCZddZdDddJZ dddKZ!dZdLddMZ"dZdLddNZ#ddTZ$ej"dUddVZ%ej"dUddWZ&dXZ'ej"dUddYZ(ddbZ)ddfZ* dddkZ+ddmZ,ee5jzZ-dnedg<ejx dddpZ.ejxddqZ/ejxddrZ0ejxddsZ1ejxdduZ2ddvZ3ddwZ4ddxZ5ddyZ6ddzZ7dd~Z8 dddZ9ddZ:GddZ;ddZ<ddZ=ddZ>ddZ?ddZ@ddZAddZBejddZC dƐddZDddZEddZFddZGddZHddZIddZJejddZKddZLejxddZMejxd dZNejxddZOddZPddZQd dZRd dZSddZTddZUd dZVdѐdZWGddejXZYd dZZddZ[ddZ\ dƐddÄZ]ddńZ^ dƐddʄZ_dd˄Z`dd΄Zadd҄ZbdӄfddڄZcdۄfdd܄Zddd߄ZeddZfddZgejGddZhejddZiddZjddZkddZlej"dddZm dddddddZnddZoddZpd dZqd!dZrd"dZsd#dZtd$dZud%dZvd&d Zwd'dZxd(dZy dƐd)dZzd*dZ{d+dZ|d,dZ}ddZ~dd Zd!d"d#d$d%d&d&d'Zd(eDZejd)Zd-d*Zd.d+Zd/d-Zd/d.Zejxd0d0ZejGd1d2ZiZd3ed4<d1d8ZeDZd9ed:<d2d;ZdƐd3d=Zd4d?Zd5d@Ze/dAZe/dBZGdCdDeeefZe1dNEddNdOd6dIZd7dKZGdLdMejXZejxd8dNZddOZd9dQZd:dTZd:dUZd;dWZddXZd<dYZddZZd=d[Zd\Zd>d]Zd>d^Zd?dcZ dd@dkZdAdmZdBdpZddqZdCduZdDdwZejdNOGdxdyZedze%fZeeegefZGd{d|ZeZdEd~ZdFdZdGdZdHdZdIdZeDgdZddZe"dJdZdKdZdLdZdMdZ ddNdZdS(O) annotationsN)Callable Collection GeneratorIteratorMappingMutableMapping MutableSet)datetime) lru_cache)StringIO) Anycast ConcatenateGenericLiteral NamedTupleProtocol TYPE_CHECKING TypeAlias TypeGuardTypeVar)dataclass_transform ParamSpecSelf)mock)datasheet_tops)DeviceProperties)_needs_inductor_compile) dtype_abbrs) OrderedSet) tree_flatten tree_map_only)has_triton_package!activation_quantization_aten_passinductor_autotune_lookup_table _sympy_subs) free_symbolsfree_unbacked_symbolsGuardOnDataDependentSymNode IterateExprsShapeEnv)IterableSequence ValuesView)Path)SymBoolSymFloatSymInt)ELEMENTWISE_TYPE_PROMOTION_KIND) GraphModule)Node) ScalingType) WorkspaceArgPythonWrapperCodegen)Dep GraphLowering)Buffer ExternKernelIRNodeLayout OperationReinterpretViewCompiledFxGraph)BaseSchedulerNodeSchedulerBuffer)cudampsxpumtiaTreturnstrcdtD}t|dksJt|dkrdn|}|S)Nc`g|]+}tt|)|,S)getattrtorch is_available.0xs S/var/www/html/CA-Chatbot/venv/lib/python3.11/site-packages/torch/_inductor/utils.py z get_gpu_type..ms3KKK'%*;*;*H*H*J*JK!KKKr9rrJ) GPU_TYPESlenpop) avail_gpusgpu_types rZ get_gpu_typerbksTKKYKKKJ z??a    ZA--vv:>>3C3CH Or\)get_interface_for_device)detect_fake_mode) DeviceType) EventList)GraphTransformObserver) ShapeProp)CeilDivCleanDivFloorDivIdentityMaxModularIndexing) make_symbolSymT) bound_sympy ValueRangesconfig)ceildivwin32!inductor_do_bench_using_profiling_TspvTORCHINDUCTOR_XPU_KERNEL_FORMATzebinz.cubinz.hsaco.)rJhiprLzOrderedSet[torch.dtype]_TMA_SUPPORTED_DTYPES@zmust be power of 2nbytesintc.|tzdz t zS)z/Round up to the nearest multiple of ALIGN_BYTESr9) ALIGN_BYTES)rs rZ_alignrs [ 1 $ 44r\v sympy.Exprboolct|tjtjtfr't t t |jSt|tp"tj |ttkS)z:v can be statically proven to be a multiple of ALIGN_BYTES) isinstancesympyAddrmallmap _is_alignedargsaligngcdr)rs rZrrs_!eiC011-3{AF++,,, a   K59Q #<#< #KKr\c2eZdZdZdZdZed dZd S) rz. s+QQQA5:##$#77QQQr\cNg|]"}tjd#Srrrs rZr[zfp8_bench.. s+OOO!!!!55OOOr\ activities RunCudaModuleNc>g|]\}}||SrS) elapsed_time)rXses rZr[zfp8_bench..s( G G G41aQ^^A   G G Gr\ raw eventsself_device_time_totalsort_by row_limitcpg|]3}|jtjkrtjd|j1|4S)zfused_abs_max_\d) device_typereCUDArematchnamerXevents rZr[zfp8_bench.."sJ   !Z_44H0%*==I JIIr\c3$K|] }|jV dSr)device_time_totalrs rZ zfp8_bench..-s%QQE3QQQQQQr\@@profiling results: %s ms)rUrJ synchronizeemptyrfloat16rrecordrangezero_rmaxprofilerprofileProfilerActivityrnvtxtensorzipmeanitemlogdebug key_averagestablerfevents statistics)rrrcache start_event end_eventr estimate_msn_warmupn_repeatpitimesresfiltered_eventss rZ fp8_benchrsBDDD J KJu}V L L LE*"""66K   t 44I 1XX      J**9559K1c&;.//00H1c# +,,--H8__   QQxQQQKOOuXOOOI    N + 0          x " "A KKMMM N ! ! # # #&&77                  aL   ! ! ! !     G G3{I+F+F G G G                 " *U   " "CIIlIIann$$-EQS$TTUUU       O  OQQQQQ Q Q   II(#... Js8(BK=* J5 K=J K=J A(K==LLFis_vetted_benchmarkingc`ddlm}m}|t}||||||S)Nr)gpu_benchmark_lockmay_distort_benchmarking_result)$torch._inductor.runtime.benchmarkingrr_do_bench_using_profiling)rrrrrr locked_benchs rZdo_bench_using_profilingr5sg" &%&?@@L 8 * *< 8 8 FC/  r\ kineto_events Iterable[Any]profiler_eventsrexpected_device_typerec  t d fd d|D}t||kr+td|dtd t|d |D] } |d }|D]}|}|}|} ||kr[| d krU| vs |d krK| vrG|dkr/|| | z dz z }|d krtdtd|dz |z S)NrrrONonec|jtjkrdS|j|jD] }|dSr)rreCPUaddr cpu_children)rchildbenchmark_event_idscollect_cpu_event_idss rZrz;_get_do_bench_profile_result..collect_cpu_event_idsYs^    . . F)))' ) )E ! !% ( ( ( ( ) )r\c\g|])}|jtkr|jtjk'|*SrS)r_DO_BENCH_PROFILE_EVENT_NAMErrerrs rZr[z0_get_do_bench_profile_result..asB  :5 5 5   / /  0 / /r\z Expected  z profiling events. Found z events.gpu_user_annotationrz Context Syncrz$Failed to capture device events for r|)rrrOr ) r!r^ RuntimeErrorrlinked_correlation_idcorrelation_id activity_typerrend_nsstart_ns) rr rr benchmark_eventsrdevice_time_usrrrrrs @@rZ_get_do_bench_profile_resultr"Qs ,6<<)))))))$  (( 5 5 5#? 5 5)** 5 5 5   "%%e$$$$NKK % ; ; = =--//++--     #7 7 7!666%)<<<)Q..>EX3X3X .. u||~~0@0@@FJ JN R3O R R R    F "X --r\c L|sddlm}|t}|}t |}||t jtdt j|}| d} | d} | tdD] } | |!| || | dz } tdt|| z } tdt|| z }t| D] } | |tt jj|}t jt jjj|g 5}t|D]\} | t jt*5|d d d n #1swxYwY]|d d d n #1swxYwYt,d t,t2jrAt,|d d t;|jj||tt@|}t,dt,d||S)aB Returns benchmark results by examining torch profiler events. This could be more accurate as it doesn't count CPU side overhead. The benchmarked function is wrapped in a profiler record_function so cache-clearing kernels can be excluded without relying on raw CUDA event grouping or kernel names. r)may_ban_benchmarkingrrTrrr9rNrrrrzprofiling time breakdownr)!rr$rbupperrcrrUrrrrrrrrrTrrrrrecord_functionrrr isEnabledForloggingDEBUGrrr"kineto_resultsrre)rrrrr$rdevice_type_upperdevice_interfacerrrrrrrprofile_activityrresults rZrrs "MMMMMM..K#))++/ <<BDDD  """ KJuy M M ME#((t(<@QRR    N + /     ' x  A KKMMM//0LMM                  $$&&&''''''''''''''' IIl  &&  NN   " "+Cr " R R   * !((**   -.. FII()))II(&111 Ms7A J I& J&I* *J-I* .JJJc ddlm}tjdd|duo(t t tjdddS#t$rYdSt$r}dt|vsJYd}~dSd}~wwxYw) Nr) roi_alignztorchvision::nmsMeta torchvisionr0Fztorchvision::nms does not exist) torchvision.opsr0rU_C%_dispatch_has_kernel_for_dispatch_keyhasattrrTops ImportErrorrrP)r0rs rZhas_torchvision_roi_alignr9s ------ 667I6RRR$  EI}d 3 3[* *   uu 0CFF::::uuuuusAA B " B +BB rtorch.device | None | str torch.devicec:|tjdjSt|trtj|}|jdvrM|jFt|j}tj|j|j S|S)Nr)cpumeta)index) rUrrrrPtyper?rcWorkercurrent_devicerr,s rZ decode_devicerDs ~|C  ''&#&f%% {/))fl.B3FK@@|FK/?/F/U/U/W/WXXXX Mr\itIterable[sympy.Expr]c`tjtj|tjjSr) functoolsreduceoperatormulrSOnerEs rZ sympy_productrOs  HL"egk : ::r\seq1Sequence[sympy.Expr]seq2c t|t|ksJtjtdt ||DS)Nc3&K|] \}}||zV dSrrS)rXabs rZrzsympy_dot..s*>>daAE>>>>>>r\)r^rexpandsumr)rPrRs rZ sympy_dotrYsN t99D ! ! ! ! <>>c$oo>>>>> ? ??r\indicessizesct|t|ksJtjj}t ||D] \}}||z|z}|S)z7Row-major flatten: per-dimension indices -> flat index.)r^rrLZeror)rZr[flatr?sizes rZ flatten_indexr`s\ w<<3u:: % % % % 7fdtDS)z=Row-major decomposition: flat index -> per-dimension indices.c hg|].\}}tt|dzd|/Sr9N)rnrO)rXrr_r?r[s rZr[z#decompose_index..sM    At }U1q577^<d|DS)Nc.i|]}t||SrS)rrWs rZ zunique..s ! ! !BqEE1 ! ! !r\)valuesrNs rZuniquerms" ! !b ! ! ! ( ( * **r\numberint | sympy.Exprdenomc t|tjst|tjr4ttj|tj|St|t rt|t s/J|dt |d|dt |t||S)Nz: , )rrExprrisympifyrr@runtime_ceildiv)rnrps rZrurus&%*%%DE5:)F)FDu}V,,emE.B.BCCC fc " "z%'='= ;;T&\\;;U;;d5kk;; = 65 ) ))r\keytorch.dtype | Nonec|dSt|dd}idddddd d d d d dddd dddddddddddddddddd d!d"d#d$dd%d&d'd(}|d)t|Dt |tr|n d*||S)+Nz*i8r|rri1 float8e4nvfp8e4nvfloat8e5fp8e5 float8e4b15fp8e4b15 float8e4b15x4 fp8e4b15x4 float8_e4m3fn float8_e5m2float8_e4m3fnuzfp8e4b8float8_e5m2fnuzfp8e5b16float8_e8m0fnuu8float4_e2m1fn_x2rfp16bfloat16bf16float32fp32float64fp64int8i8int16i16i32i64u16u32u64)int32int64uint8uint16uint32uint64ci|]}||SrSrS)rXrs rZrkz_type_of..Js1111111r\*)rPsplitupdatelistrlr)rv dtype_strtyss rZ_type_ofr&su  {uCs##B'I  i  G  z         w  9  :  $  D  6  F  6! " 6# $ % & ' (3   C8JJ11d3::<<00111222S#&& @33,@I,@,@@r\lstIterable[int | torch.SymInt]cd|DS)z Gets the shape and stride of a tensor. For non-symbolic tensors, this is trivial. But for symbolic tensors, we need to map from SymIntNode into sympy.Expr. c6g|]}tj|SrS)rrtrXrs rZr[z-convert_shape_to_inductor..Vs" * * *EM!   * * *r\rSrs rZconvert_shape_to_inductorrNs + *c * * **r\valint | torch.SymIntcRt|tjr |jjS|S)z Convert SymInt to sympy.Expr, leave int as is. Unlike sympy.sympify() which converts int to sympy.Integer, this function preserves int as int and only converts SymInt to Expr. rrUr4nodeexprrs rZconvert_symint_to_exprrYs'#u|$$x} Jr\rcddlm}t|tr|nNt|tjrt|n%|jjj |dS)zL Like convert_shape_to_symint, but operates on a single expression. r9VN)hint) virtualizedrrrrrgraphsizevars shape_envcreate_symintnode)rrs rZconvert_to_symintresr a   !U]++ LCFFF!+==ad=KK r\Iterable[int | sympy.Expr]list[int | torch.SymInt]cd|DS)zz Takes a list of shapes from Inductor and converts them into symints (or just ints if all shapes are static). c,g|]}t|SrS)rrs rZr[z+convert_shape_to_symint..}s! . . .Q a . . .r\rSrs rZconvert_shape_to_symintrvs / .# . . ..r\optorch._ops.OpOverloadcHtd|jjDS)z- Does this op overload have aliasing c3(K|] }|jduVdSr) alias_inforXrUs rZrzis_view..s)FFAq|4'FFFFFFr\)any_schema argumentsrs rZis_viewrs& FF1EFFF F FFr\cdSNFrS)rs rZrsr\user7is_pointwise_fn'Callable[[torch._ops.OpOverload], bool]c|jdkrdSt|jtjjs|jt jusdSttjj|j}|t just|r tfd|j DStj j |jvp |S)z Do all uses of this op have torch.Tag.pointwise or return True for optional `is_pointwise_fn` Uses in views ops will follow the views uses call_functionFc38K|]}t|VdSr)is_pointwise_use)rXurs rZrz#is_pointwise_use..s.KKA#A77KKKKKKr\)rrtargetrU_ops OpOverloadrJgetitemrrrusersTag pointwisetags)rrrs ` rZrrs v  u3:uz4559<xGW9W9Wu %*' 4 4F !!!WV__!KKKKKKKKKK 9 &+ - H1H1HHr\rrr list[Any]kwargsdict[str, Any]&tuple[GraphModule, list[torch.Tensor]]c tjgd fd }j|gt tj|||fR}t |jjdkr+t|jjdj dkr|f} |tj i}|fS) Narg torch.TensorrOr7cx|dtS)Nr)append placeholderr^)rg graph_argss rZadd_tensor_argz)gen_gm_and_inputs..add_tensor_args8#}}43z??44555r\r9rTensor)rrrOr7) rUfxGraphrr#rr^rreturnsrPr@outputr6)rrrrrgmrrs @@rZgen_gm_and_inputsrs A%'J6666666 1? u|^dF^LL   D FN "##q(( &q). / /8 ; ;wHHTNNN   b! $ $B z>r\rJr c|dkrdSt|}|r|dSdSNr=)rcrVrrCs rZrrsT /77$$&&'$$&&&&&''r\modelCallable[..., Any]example_inputs Sequence[Any]rct|tjdtj}t |D]}||}t|tj}|J||z S)Ni9)rrU manual_seedtime perf_counterr)rrrrt0rr.t1s rZtimedrs  d   B 5\\'F   B     7Nr\rS ?repeatbaselinectjfdt|D}tj|z }t ||z d|S)Nc4g|]}tSrS)r)rXrrrrrs rZr[z%print_performance..s'LLLuneV 4 4LLLr\z.6f)rUrrmedianprintr)rrrrr rtimingstooks``` ` rZprint_performancers{lLLLLLLLeFmmLLLG < 5 (D TH_ " "### 99;;r\objmethodc`t||t||fddS)zKReplace obj.method() with a new method that returns a precomputed constant.cSrrS)r.srZrz#precompute_method..sr\N)rTsetattr)rrr.s @rZprecompute_methodrs8 !WS& ! ! # #F C(((((r\methodsrc0|D]}t||dS)zFReplace methods with new methods that returns a precomputed constants.N)r)rrrs rZprecompute_methodsrs.''#v&&&&''r\rUrVcPt||kt||kz Sr)rrUrVs rZcmprs! q1u::AE ""r\rYint | Sequence[int]r_ Sequence[int]ct|tr|g|zSt|dkr"t||dg|zS|S)Nr9r)rrr^r@)rYr_s rZ pad_listliker sS!SsTz 1vv{{tAww!v%% Hr\tuple[_T, ...]list[_T]cVt|dkrgSdd}t||S) NrelemrxrOrPct|tr|Sddlm}t||sJ|S)Nr9)rH)rrP schedulerrHget_name)r$rHs rZ sort_funcztuple_sorted..sort_funcsP dC  K000000$ 122222}}r\rv)r$rxrOrP)r^sorted)rYr(s rZ tuple_sortedr+s? 1vv{{  ! # # ##r\PRV) covariantc.eZdZed dZdd Zd S) CachedMethodrrrOr cdSrrS)rs rZ clear_cachezCachedMethod.clear_caches),r\rP.argsrP.kwargsr-cdSrrSselfrrs rZ__call__zCachedMethod.__call__r\N)rrrOr )rr3rr4rOr-)rrr staticmethodr2r8rSr\rZr0r0s2,,,\,DDDDDDr\r0!Callable[Concatenate[Any, P], RV]CachedMethod[P, RV]c |j}d|dd|i}td|ddd|tj|||d}dfd }||_|S)N___cacherz def zC_cache_on_self(self): try: return self.zy except AttributeError: pass rv = fn(self) object.__setattr__(self, "z%", rv) return rv _cache_on_selfr7rrOr cLt|rt|dSdSrr6delattrr7rvs rZr2z"cache_on_self..clear_cache,s5 4    D#       r\r7rrOr rexeclstriprHwrapsr2)rrctxwrapperr2rvs @rZ cache_on_selfrLs ;D t   C *C  ! (+  FHH    "iob!!#&=&=&=">??G&G Nr\c t|S)z] Variant of cache_on_self for properties. The only difference is the type signature. )rL)rs rZcache_property_on_selfrN4s   r\ class_name*Callable[[FN_TYPE[P, RV]], FN_TYPE[P, RV]]cdfd }|S)NrFN_TYPE[P, RV]rOc dd|jdd|i}tdddd|tj||d }dfd }||_|S)Nr>rr?rz def inner(self: Any, *args: P.args, **kwargs: P.kwargs) -> RV: args_kwargs = (args, tuple(sorted(kwargs.items()))) if not hasattr(self, "z2"): object.__setattr__(self, "z%", {}) cache = self.z try: return cache[args_kwargs] except KeyError: pass rv = fn(self, *args, **kwargs) cache[args_kwargs] = rv return rv innerr7rrOr cLt|rt|dSdSrrBrDs rZr2z.wrapper..clear_cachebs5tS!! #c""""" # #r\rErF)rrJrTr2rvrOs @rZrKz'cache_on_self_and_args..wrapperCs4:33 333Rj  (+    03   "   $ )   ,$ ##CL11 # # # # # #( r\)rrRrOrRrS)rOrKs` rZcache_on_self_and_argsrV>s) $$$$$$L Nr\ node_schedule*Sequence[BaseSchedulerNode] | ExternKernelOrderedSet[Node]cddlm}t|tr6t jt jd|DtSt||j r|j StS)Nr9ircTg|]%}t|d|j|jj&S)r)r6roriginsrXrs rZr[z%aggregate_origins..tsK   4((  .2Y   !   r\) r\rrrHrIrJor_r!rAr^)rWr\s rZaggregate_originsrbls-&& L  *    LL    M2? 3 3$$||r\Sequence[BaseSchedulerNode]descriptive_names8Literal[True, 'torch', 'original_aten', 'inductor_node']ct|}|dkr.dfd|D}tt|}n|dkrg}|D]}|jdkrd}d}d|jvr|jdd}nd |jvr|jd d}d }|sOt |d t r||d |z||d j|ztt|}n|d kr d |D}ntd dg|zS)N original_atenc|jd}d}t|tjjr |jj}n@t|tjjr!t| }|S)Nrgr`) r>rrUrr_overloadpacketrHigherOrderOperatorrPr)originrgrvs rZget_origin_meta_strz2get_fused_kernel_name..get_origin_meta_strsl"K8MC-)>?? 0#3<M5:+IJJ 0-,,..//Jr\chg|].}|jdkr!d|jvr|jd#|/S)rrg)rr>)rXrkrls rZr[z)get_fused_kernel_name..sW   yO++6;.. O,8  ' '988r\rUrr`source_fn_stackrfwd_source_fn_stackbackwardr9 inductor_nodec2g|]}|jdk |jSr)rrrXrks rZr[z)get_fused_kernel_name..s-   "VY/5Q5QFK5Q5Q5Qr\rfused) rbr*r!rr>rrPrrNotImplementedErrorjoin)rWrd all_originssourcesrk source_fnsuffixrls @rZget_fused_kernel_namer|s$M22KO++       %   G,,-- g % %! C CFyO++ $ 33 & ,= >r BII*fk99 & ,A B2 FI'F ilC00CNN9Q<9999NN9Q<#86#ABBBG,,-- o - -  &1   "! 88WI' ( ((r\rKr<tuple[str, str]c   t|}d|D}tjt}tjt}d|rt d|D}t |dkr[|djtds%dtj D}|_ | fd |D]D}d |j vr|j d |j d } d} t| tjjrt#| j} n@t| tjjr!t#| } | r || |jd |j vr:|j d dj} || |j|j d d kr%||j|jFdnd} jd| dd|dd|d} jdg} t5|D]G\}}| jd|ddt5|Hhddlm| jdt }g}t|jslddlm }d5fd! }d6d$ d7 fd' }|D]R}t|d(r|j!t|j!d)r|j!j"|j!j"D]}|j|vr |#|j|j$|j}|H|||j\}}| jd*|d+||d,|dt|j!d-ri|j!j%]|j!j%D]P}|j$|j}|$|||j\}}|d.|zQT|D]~}|&d/0}|Etj'j(r4| )fd1|*D_| jd|| jd2d3|| d4| fS)8aH Retrieves metadata information for a kernel. Args: node_schedule (Union[Sequence[BaseSchedulerNode], ExternKernel]): Either a sequence of BaseSchedulerNode objects or an ExternKernel instance. wrapper (PythonWrapperCodegen): An instance of PythonWrapperCodegen, used to define the code comment format. Returns: tuple[str, str]: A tuple containing two strings: - The first string represents the kernel's metadata. - The second string represent the kernel's detailed metadata. c(g|]}|jdk |Srsrrts rZr[z'get_kernel_metadata..s$WWW&):V:Vf:V:V:Vr\Nc3$K|] }|jV dSr)r)rXns rZrz&get_kernel_metadata..s$"C"Cq17"C"C"C"C"C"Cr\r9r)_inductor_kernel_metadata_node_to_idx_mapci|]\}}|| SrSrS)rXidxrs rZrkz'get_kernel_metadata..s"V"V"Vfc11c"V"V"Vr\cj|Sr)r)r single_graphs rZrz%get_kernel_metadata..slTUVWr\r)rg from_nodepartitioner_tag is_backwardzTopologically SortedUnsortedrz Source Nodes: [rrz], Original ATen: []z" Source node to ATen node mapping: z => r[z Graph fragment:rbuffer-ir.TensorBox | ir.Buffer | ir.TorchBindObjectrw_namerPrOtuple[str, ir.Layout | None]ct|jr,t|jjr|jjj}n|j}||}n|j} |}n#t$rd}YnwxYw||fSr)r TensorBoxdata StorageBox origin_noder get_layoutrv)rrrrlayoutr\s rZget_buffer_infoz,get_kernel_metadata..get_buffer_infosfbl335 K995#)+"2">KK"("4K&"DD&+D"#..00FF*"""!FFF"V|#sA,, A;:A;shape Iterable[int]cHddd|DdS)N[rrc,g|]}t|SrS)rPrWs rZr[z@get_kernel_metadata..stringify_shape..s%<%<%<c!ff%<%<%.stringify_shapes.@499%<%.stringfy_layouts|>2&5ofk&B&B#D '6v}'E'E$G!'-}$6!? FL1?3C?(?*;???r\ read_writesreadsz %z : Tensor z = PlaceHolder[target=writes%T)include_tensor_metadatac30K|]}jd|VdS)rN)comment)rXlinerKs rZrz&get_kernel_metadata..OsG))11411))))))r\z return , )rrrrPrOr)rrrOrP)rrrOrP)+rb collections defaultdictrr!r^rr6renodesrsortr>rrUrrrPrirjrrgetrrwkeysr*itemsr`r\rArrrrrtry_get_bufferr format_nodeversionr}extend splitlines)!rWrKrxinductor_nodesfrom_node_dictoriginal_aten_dict unique_graphsnode_to_idx_maprrgrvsort_strmetadatadetailed_metadata original_noder all_reads all_writesrrrrrr input_namerw output_namerformatted_noder\rrs! ` @@@rZget_kernel_metadatars$$M22KWW;WWWN ,T22N$066 L ""C"CN"C"C"CCC }   " ")!,2L<)TUU Y"V"V ,BT8U8U"V"V"VIX F   WWWW    88 di ' 'DIo,F,R Io6MC-)>?? 0-788M5:+IJJ 0-,,..// :"3'..ty999 $) # #)K(+0C 3  & &ty 1 1 1 1 Y]], - - > > 49 % , ,TY 7 7 7)5)A%%zH ? C CX C CtyyATATAVAV7W7W C C99%7%<%<%>%>?? C C C $OOOOP &~';';'='= > >   u   P P= P Pdiiu 6N6N P P      GO!E!E!EFFF%/\\ " -99B = & & & & & & $ $ $ $ $ $( A A A A      # = =q-00AM4I1='22 q}7J7V]0  6Y..$! af---!"!7!7!?!?!>$-<_VQV-L-L* F)00&\\J\\.v66\\NX\\\ AM844 = ,8]1==!"!7!7!?!?!>$)8)H)H Q"))# *;<<<<" R RD!--d-KKN)em.?)"(()))) . 9 9 ; ;))) "((GO)P)P)P)PQQQQ  GO!U!Usxx ?S?S!U!UVVV TYY011 11r\ initial_queueIterable[torch.fx.Node] skip_filterCallable[[Any], bool] | NoneOrderedSet[torch.fx.Node]ct|}t|}|r\|}|jD]>}|r ||r||vr*||||?|\|S)zJReturns the set of nodes whose values depend on those within initial_queue)rr!r_rrr)rr dominated_setrusers rZdominated_nodesr[s ''M}--M +  ""J + +D {{400 =((!!$'''$$T*** + r\Sequence[IRNode]dict[str, IRNode]cddlmd fd t|\}}fd|D}t|\}}fd |D}tt jg||RS) Nr9r[rrBrOrct|jr|jSt|jr|jSt|jo(t|jjjjf Sr) rrrrrBComputedBuffer InputsKernel InputBufferTemplateBuffer)rr\is_unrealized_nodes rZrz*gather_origins..is_unrealized_nodets a & & .%%af-- - a ' ' .%%af-- -!RY'' !!  1 1 -  r\c4g|]}| |jSrSr^rXrrs rZr[z"gather_origins..s-WWWc?Q?QRU?V?VWckWWWr\c4g|]}| |jSrSrrs rZr[z"gather_origins..s,SSSC;M;Mc;R;RSCKSSSr\)rrBrOr)r`r\r"r! itertoolschain) rrkwargs_flattenrkwargs_origins args_flatten args_originsr\rs @@rZgather_originsros        "%V,,NAWWWW^WWWN"4((OL!SSSS<SSSL ioE|EnEEE F FFr\rcPd dd fd d fd d fd |S) z Normal sympy str is very slow, this is a lot faster. The result are somewhat worse, as it doesn't do as much simplification. So don't use this for final codegen. rrrOrct|tjo(t|jdko|jddkS)Nrr)rrMulr^r)rs rZ is_neg_leadzsympy_str..is_neg_leads9 tUY ' ' VC NNa,? VDIaLTVDV r\rPcnt|tjrt|jdkrP|jdr:|jdd|jdjdSdt |jS|S)Nrr9rz - z + )rrrr^rrwr)rr sympy_str_muls rZ sympy_str_addz sympy_str..sympy_str_adds dEI & & '49~~""{{49Q<'@'@"'- ! 55__-- RS HYZ[H\:]:]___zz#mTY"?"?@@@ =&& &r\ct|tjrL|rd|jdSdt |jS|S)N-r9z * )rrrrrwr)rrsympy_str_atoms rZrz sympy_str..sympy_str_mulsx dEI & & ({4   B:>>$)A,77999zz#ndi"@"@AAA!>$'' 'r\ct|tjr|jSt|tjtjfrd|dSt|t tttfr=|j j dd tt|jdSt!|S)N()rr)rrSymbolrrrrnrjrkrlfuncrrwr sympy_strrrP)rrs rZrz!sympy_str..sympy_str_atoms dEL ) ) 9  uy%)4 5 5 -}}T**--- - (HM N N i(RR499SDI5N5N+O+ORRR Rt99 r\)rrrOrrrrOrPrS)rrrrrs @@@@rZrrs    ' ' ' ' ' ' ' ( ( ( ( ( ( ( =  r\ValueRanges[Any]cddlm}tjr2t |jddx}r|jdkrt|Stj S)Nr9r current_node index_expr) rrrtcompute_all_boundsrT interpreterrrqrrunknown)r?rfx_nodes rZget_bounds_index_exprrsh !% ~tDD DW% Nl * *5!!!"$$$r\prefixc|ddkS)NrrrSrs rZprefix_is_reductionr s !9 r\c|ddvS)Nr)rYyzrSrs rZprefix_is_pointwiser s !9 ''r\rpr sympy.SymbolcL|tjksJt||ddS)9 Used to generate an integer-nonnegative symbol. Tinteger nonnegative)rpSIZEro)rrs rZsympy_index_symbol_with_prefixrs0 TY     vsDd C C CCr\checkc6|s tjo tjSr)rtdebug_index_assertsassert_indirect_indexing)rs rZgenerate_assertrs  /V/ TV5TTr\rcL|ddksJtj|ddS)rrrTr)rrrs rZsympy_index_symbolrs. 7c>>>> <d = = ==r\ replacementsdict[sympy.Expr, Any]c"t||S)z When the passed replacement symbol v is a string, it is converted to a symbol with name v that have the same replaced expression integer and nonnegative properties. r')rrs rZ sympy_subsr!s t\ * **r\&TypeGuard[torch.SymInt | torch.Tensor]cxt|tjp t|tjo|jSr)rrUr4r_has_symbolic_sizes_strides)rUs rZ is_symbolicr%s2 a & & 1el##E(Er\c4td|DS)Nc34K|]}t|VdSr)r%rs rZrz"any_is_symbolic..s(,,!{1~~,,,,,,r\r)rs rZany_is_symbolicr)s ,,t,,, , ,,r\)z,aten._fused_moving_avg_obs_fq_helper.defaultz7aten._fused_moving_avg_obs_fq_helper_functional.defaultzfbgemm.dense_to_jagged.defaultz%fbgemm.jagged_to_padded_dense.defaultrun_and_save_rng_staterun_with_rng_statezaten._local_scalar_densezaten._assert_scalarzaten.topk.defaultrtorch.fx.GraphModuletorch.fx.Node | Nonecddlm}|jjD]@}t |r|cS|jdx}||r|cSAdS)Nr)r*r)%torch.fx.experimental.symbolic_shapesr*rris_cudagraph_unsafe_fx_noder>r)rr*rrs rZ%get_first_incompatible_cudagraph_noder1s|LKKKKK &t , , KKK9=='' 'C 49N9Ns9S9S 4KKK 4r\cttt|jj}|jdksJ|S)z$Get the output node from an FX graphr)nextiterreversedrrr)r last_nodes rZ output_noder7+s<T(28>223344I <8 # # # # r\OrderedSet[torch.device]c|jd}td|D}t|jd}t |t r|n|f}td|D}||zS)Nrrc3K|]H}t|jdtj4|jdjVIdSrN)rr>rrUrrr_s rZrz"get_all_devices..4s`99  dimmE**EL 9 99 %999999r\rc3K|]g}t|tjj!t|jdtjS|jdjVhdSr;)rrUrr7r>rrrrXrs rZrz"get_all_devices..<sw77  c58= ) )7 sx||E**EL 9 9 7 777777r\)r find_nodesr!r7rrtuple)rplaceholder_nodes input_devicesout_argout_args out_devicess rZget_all_devicesrE2s++}+==.899%999//M "oo"1%G$We44Dww7*H,677777--K ; &&r\cttjD]}|dstj|}|jD]}|drt ||}t|tj j j j rV|j D]N}t|tj j j jr#|jjjOtj|=dtjvr.sP$%#*!#3#3!"27??27<<@PRS3T3T#U#U#3#3#3r\c>td|S)Nz*Failed to remove temporary cache dir at %s)exc_info)rwarning)rrzrinductor_cache_dirs rZrzfresh_cache..s$S[[@&%6A66r\) ignore_errorsonerrorz(on error, temporary cache dir kept at %s)rhtorch._inductor.cpp_builderrutempfilemkdtemprnrrrkrzrwrdictr^existslistdirr is_windowsrUrLrVr_shutilrmtree Exceptionr)rorqrsrufilesrr}s @@rZ fresh_cachersNNNDDDDDD11(2Bs2K2K2KLL' /1C D D   II35G H H H77 /::   ,.>??  mT22 }--2224W222w~~&677 " +; < <%,,).                               $  || ) 6 6 8 8 )&((( M")ll      >@RSSS   shF9AE B D4( E 4D8 8E ;D8 <E ? F9 EF9EA%F98G39'G  G33H)reverseseqr list[int]c|j}tt|}tt ||d}|stt |S|S)NTrvr) __getitem__rr^rr*r5)rrgettera_rsort_idxs rZargsortrs[ _F C//CF3FD999::H (HX&&''' Or\rr-)Sequence[int | torch.SymInt | sympy.Expr]cd fd }dt|D}t|tj||}d |D}|S) z Return a symbolic sort order for optimization-only layout heuristics. Data-dependent comparisons can fall back to non-guarding optimization hints, so callers must not use this order to make correctness decisions. rUtuple[int, sympy.Expr]rVrOrc|\}|\}dfd }dfd |kfd rd S|kfdrdS||krdS||krd SdS)Nr bool | torch.SymInt | sympy.ExprfallbackCallable[[], bool]rOrct|tr|S |S#t$r |cYSwxYwr)rr evaluate_exprr+)rrrs rZevaluatez*argsort_sym..cmp..evaluatesb$%%   " ..t444. " " "xzz!!! "s/AAlhsrrhsc\||kSr)optimization_hint)rrrs rZhint_ltz)argsort_sym..cmp..hint_lts-..s33i6Q6QRU6V6VV Vr\cSrrSa_valb_valrsrZrz*argsort_sym..cmp..775%+@+@r\rcSrrSrsrZrz*argsort_sym..cmp..rr\r9r)rrrrrOr)rrrrrOrrS) rUrVa_idxb_idxrrrrrs @@@rZrzargsort_sym..cmps u u " " " " " " W W W W W W 8EEM#@#@#@#@#@#@ A A 2 8EEM#@#@#@#@#@#@ A A 1 5==1 5==2qr\chg|]/\}}|t|tjr |jjn|f0SrSr)rXrrs rZr[zargsort_sym..sJ    C Z5<88?afkka@   r\rcg|]\}}|SrSrS)rXrrs rZr[zargsort_sym..s & & &fc1c & & &r\)rUrrVrrOr)rer*rH cmp_to_key)rrrrexprsr.s` rZ argsort_symrs!!!!!!H  nn   E 5i2377 I I IE & & & & &F Mr\r torch.dtypecv|tjkrdStjd|S)NrrSr)rUrr element_sizers rZget_dtype_sizers7  q ;r ' ' ' 4 4 6 66r\ceZdZUded<dS) LineContextrcontextNrrrrrSr\rZrr&sLLLLLr\rc$eZdZUded<ded<dS)ValueWithLineMaprPrzlist[tuple[int, LineContext]]line_mapNrrSr\rZrr*s'JJJ++++++r\rceZdZdZd6d7dZejd8d Zd9d Zd:dZ d:dZ dZ d;dZ dd?dZd>d?d Zd@d#ZdAdBd'ZdAdCd(ZdAdCd)Zd>d?d*ZdDd-Zd:d.ZdEd1ZdFd4Zd5S)GIndentedBufferrinitial_indentrrOr c"g|_||_dSr)_lines_indent)r7rs rZ__init__zIndentedBuffer.__init__3sBD % r\tabwidthric#VK|j} ||_dV||_dS#||_wxYwr)r)r7rprevs rZ set_tabwidthzIndentedBuffer.set_tabwidth7s@} !$DM EEE DMMMDDM s  (rct}d}g}|jD]}t|tr|}|$n4t|tr|||jfW|}t|tsJ|||d|d| dzz }t| |S)Nr9r) r rrDeferredLineBaserrrrPwritecountrgetvalue)r7bufrlinemaplirs rZgetvaluewithlinemapz"IndentedBuffer.getvaluewithlinemap@sjj 13+ & &B".// rtt< B ,, 2:///dC(( ( (( IIdOOO IIdOOO TZZ%%% %AA 888r\rPc4|jSr)rrr7s rZrzIndentedBuffer.getvalueTs''))//r\ct}|jD]}t|tr|}|$nt|tr;|}t|t sJ|dr||dd|||d|S)N\rr) r rrrrrPendswithrr)r7rrrs rZ getrawvaluezIndentedBuffer.getrawvalueWsjj+  B".// rtt< B ,, dC(( ( ((}}T""  $ss)$$$$ $ $||~~r\c|jSrrrs rZ get_lines_refzIndentedBuffer.get_lines_refks {r\c8|jdSr)rclearrs rZrzIndentedBuffer.clearns r\rc*t|jSr)rrrs rZ__bool__zIndentedBuffer.__bool__qsDK   r\c&d|j|jzzS)Nr)rrrs rZrzIndentedBuffer.prefixtsdlT]233r\c0|ddS)Nr writeliners rZnewlinezIndentedBuffer.newlinew tr\r$LineContext | DeferredLineBase | strct|tr|j|dSt|trA|j||dS|r2|j||dS|jddSNr`)rrrrr with_prefixrstripr7rs rZrzIndentedBuffer.writelinezs dK ( ( # K  t $ $ $ $ $ . / / # K  t// >> ? ? ? ? ? ZZ\\ # K  $++--777 8 8 8 8 8 K  r " " " " "r\c0||dS)zGWrite to JIT buffer only. On a plain IndentedBuffer, same as writeline.Nrrs rZ writeline_jitzIndentedBuffer.writeline_jits tr\cdS)z;Write to AOTI buffer only. No-op on a plain IndentedBuffer.NrSrs rZ writeline_aotzIndentedBuffer.writeline_aotr9r\F other_codeIndentedBuffer | strrc4|||dS)zESplice to JIT buffer only. On a plain IndentedBuffer, same as splice.rNsplicer7rrs rZ splice_jitzIndentedBuffer.splice_jits Je ,,,,,r\cdS)z.ctxsQ LLF "LL ' &  & s+=rOri contextlibcontextmanagerr7rrJs`` rZindentzIndentedBuffer.indentsB  " ' ' ' ' ' ' # " 'suu r\c&|xj|z c_dSrrr7rs rZ do_indentzIndentedBuffer.do_indent  r\c&|xj|zc_dSrrrs rZ do_unindentzIndentedBuffer.do_unindentrr\c t|trtd}|jD]X}t|tsA|r?t |t |t |z }Ytj |rd}|jD]b}t|tr|j |2t ||t|dcdStj|}|r|}|sdS|}|dD]}t ||dS)Ninfrr)rrrrrminr^rHmathisinfrrrtextwrapdedentrstripr)r7rrrrrs rZrzIndentedBuffer.splices j. 1 1 25\\F") I I!$ 44II TS5G5G)GHHFz&!! ") H HdK00HK&&t,,,,",,T4F 3FGGGG  H H "44J 1'..00  #**,,J%%d++ 2 2((q1111 2 2r\rCallable[[Any], Any]cbt|j}fd|jD|_|S)Nrc&g|] }|SrSrS)rXrrs rZr[z&IndentedBuffer.map..s!999Tdd4jj999r\)rrr)r7rrs ` rZrzIndentedBuffer.maps7DL9999999T[999  r\cPt|d|dS)Nrr)r@rrs rZ__repr__zIndentedBuffer.__repr__s't**11t}}1111r\otherrc|j|jksJt|j}||j||j|S)Nr )rrrr)r7r$rs rZ__add__zIndentedBuffer.__add__sV|u},,,,DL999 t{### u|$$$ r\new_line$DeferredLineBase | LineContext | strc||jvSrr)r7r's rZcontainszIndentedBuffer.containss4;&&r\NrrrrOr )rrrOri)rOrrOrPrOr rOr)rrrOr F)rrrrrOr )rrrOr r)rrrOrrrrOr )rrrOr)r$rrOr)r'r(rOr)rrrrrr r rrrrrrrrrrrrrrrrrrrrr#r&r*rSr\rZrr0s0H&&&&&!!!!9999(0000(!!!!4444####JJJJ-----KKKKK!!!!     222220 2222''''''r\rceZdZdZddfd Zedd Zd Zdd Zdd Z d Z dddZ dddZ dddZ d d!dZd d"dZd d"dZdfd ZxZS)#DualIndentedBufferzIndentedBuffer that simultaneously accumulates JIT and AOTI output. The base class (self._lines) holds JIT content. self.aot holds AOTI content. By default, writeline/splice write to both. Use _jit/_aot variants for mode-specific writes. rrrrOr crt|t||_dSr)superrraot)r7r __class__s rZrzDualIndentedBuffer.__init__s. (((!.11r\rc|S)uRead-only accessor for JIT buffer (for splicing into result). WARNING: Do NOT use jit for writing — writes would go to both buffers because self.writeline is overridden. Use writeline_jit/splice_jit instead. rSrs rZjitzDualIndentedBuffer.jits  r\cpt|||j|dSr)rrr6rs rZrzDualIndentedBuffer.writelines4  t,,, 4     r\c<t||dSr)rrrs rZrz DualIndentedBuffer.writeline_jits  t,,,,,r\c:|j|dSr)r6rrs rZrz DualIndentedBuffer.writeline_aots 4     r\c:|D]}||dSrrrs rZrzDualIndentedBuffer.writelinesrr\Frrct|||t|tr|jn|}|j||dSNr)rrrr3r6r7rr aot_others rZrzDualIndentedBuffer.splices] dJe<<<(5GHH XJNNj   /////r\c@t|||dSr?)rrrs rZrzDualIndentedBuffer.splice_jits#dJe<<<<.ctx s LLF "LL H   &   + &   F*     &   F*      s A'A<r rs`` rZrzDualIndentedBuffer.indents=  " + + + + + # " +suu r\cP|xj|z c_|jxj|z c_dSrrFrs rZrzDualIndentedBuffer.do_indent.   F"r\cP|xj|zc_|jxj|zc_dSrrFrs rZrzDualIndentedBuffer.do_unindentrHr\czt|jdSr)r5rr6r7r7s rZrzDualIndentedBuffer.clears,   r\r+r,rOrr.r0rrrOr r)rrr1)rrrrrpropertyr9rrrrrrrrrrr __classcell__r7s@rZr3r3sf2222222X!!!----!!!!!!!00000=====00000     ##########r\r3c6eZdZdZd dZd dZd dd Zd dd Zd S) AotOnlyBuffera*IndentedBuffer for pure-AOTI codegen. Mirror of the base class's pure-JIT defaults: writeline_aot/splice_aot write to the buffer; writeline_jit/splice_jit are no-ops. Lets call sites use writeline_jit/writeline_aot uniformly across pure-JIT, pure-AOTI, and dual-wrapper modes. rOr cdSrrSrs rZrzAotOnlyBuffer.writeline_jit+ r\c0||dSrrrs rZrzAotOnlyBuffer.writeline_aot.rr\FrrcdSrrSrs rZrzAotOnlyBuffer.splice_jit1rTr\c4|||dSr?rrs rZrzAotOnlyBuffer.splice_aot4s Je ,,,,,r\Nr.r0rM)rrrrrrrrrSr\rZrRrR"sx         -------r\rRc^ddlm}|jjrt St S)zConstruct the IndentedBuffer subclass matching the current codegen mode. Pure AOTI -> AotOnlyBuffer (writeline_aot writes; writeline_jit drops). Pure JIT -> IndentedBuffer (writeline_jit writes; writeline_aot drops). r9r)rrraot_moderRrrs rZmake_codegen_bufferrZ8s9 w   r\c(eZdZdfd Zd dZxZS) FakeIndentedBufferrOr cHtdSr)r5rrKs rZrzFakeIndentedBuffer.__init__Fs r\rrPrcj|dkrt||Std|d)Nr7zTried to call self.z on FakeIndentedBuffer. This bufferis currently used on TritonTemplateKernel to prevent actualwrites to the body without explicitly specifying the body with`TritonTemplateKernel.set_subgraph_body(name)`)object__getattribute__r)r7rs rZr`z#FakeIndentedBuffer.__getattribute__IsK ;  **466 6 =$ = = =   r\r.)rrPrOr)rrrrr`rOrPs@rZr\r\EsQ        r\r\c#Ktjtj}} dV||ct_t_dS#||ct_t_wxYwr)rHstdoutstderr)initial_stdoutinitial_stderrs rZrestore_stdout_stderrrfTsR%(ZNN@ !/ CJJJ CJ????s ;AcReZdZdZddZddZdd Zdd Zdd ZddZ ddZ ddZ dS)rz.A line that can be 'unwritten' at a later timerrPc@|sd}||_dSr)rrrs rZrzDeferredLineBase.__init__`s"zz|| D r\rOrrct)zJReturns either self.line or None to indicate the line has been 'unwritten'rvrs rZr8zDeferredLineBase.__call__e!!r\rct)z3Returns a new deferred line with the same conditionrjrs rZ _new_linezDeferredLineBase._new_lineirkr\rc>|||jSrrmr)r7rs rZrzDeferredLineBase.with_prefixms!~~444555r\cZ||jSr)rmrrHrs rZrHzDeferredLineBase.lstripps"~~di..00111r\r? int | slicecB||j|Srro)r7r?s rZrzDeferredLineBase.__getitem__ss~~di.///r\rc*t|jSr)rrrs rZrzDeferredLineBase.__bool__vsDIr\rc*t|jSr)r^rrs rZ__len__zDeferredLineBase.__len__ys49~~r\N)rrPrOrr)rrPrOr)rrPrOr)rOr)r?rqrOrr/rOr) rrrrrr8rmrrHrrrurSr\rZrr]s88 """"""""666622220000r\rc4eZdZdZd fd Zd d Zd d ZxZS)DelayReplaceLinez6At end of codegen call `line.replace(key, value_fn())`rvrPvalue_fnCallable[[], str]rcft|||_||_dSr)r5rrvrz)r7rvrzrr7s rZrzDelayReplaceLine.__init__s-   r\rOcf|j|j|Sr)rreplacervrzrs rZr8zDelayReplaceLine.__call__s$y  4==??;;;r\c8t|j|j|Sr)ryrvrzrs rZrmzDelayReplaceLine._new_lines$->>>r\)rvrPrzr{rrPr-)rrPrOry)rrrrrr8rmrOrPs@rZryry}sk@@!!!!!! <<<<????????r\ryindex_or_deviceint | torch.devicect|tjr|}n!tjt|}t j|}tjjr=|jJ|jdks |jdkrt ddSdS|j dkrdnd}|j }||kr!t d ||d  dSdS) N rz6GPU arch does not support max_autotune_gemm mode usageFTrLr~Dz,Not enough SMs to use max_autotune_gemm mode)min_sms avail_sms)extra) rrUrrbrcreaterr}majorrrr@multi_processor_count)rrproprrs rZ is_big_gpurs/5<00? lnno>>  "6 * *D }z%%% :>>TZ2-- KKP Q Q Q5tK5((bbbG*I7 :%I>>    u 4r\ctjr#tjjStjdjS)NrJ)rUrLrVget_device_propertiesgpu_subslice_countrJrrSr\rZget_max_num_smsrsI yDy..00CC : + +F 3 3 IIr\ctjsdStjtj}|jdkS)zEReturns true if the device is a NVIDIA B200, otherwise returns false.Fr)rUrJrVrrBr)device_propertiess rZ using_b200rsQ : " " $ $u 889R9R9T9TUU  "b ((r\ctjo4tjjduo!tjdkS)zCReturn true for NVIDIA CUDA devices with compute capability SM100+.N)rr)rUrJrVrr}get_device_capabilityrSr\rZis_nvidia_sm100_or_laterrsG  !! : M  % : J , , . .' 9r\ctjrtStj}t||ndz S)zFHandle experimental carveout if set otherwise return hardware SM countNr)rUrLrVrr4_get_sm_carveout_experimental)carveouts rZ get_num_smsrsS y!   x5577H   H,@a HHr\num_tma_descriptors num_programs int | Noner:cddlm}m}|t}|d}||zt z}|||||jS)zKBuilds and returns a WorkspaceArg for the device side TMA workspace buffer.r9)r:WorkspaceZeroModeNF)r zero_moder outer_name)codegen.commonr:rr from_boolTMA_DESCRIPTOR_SIZE unique_name)rrrr:rrr_s rZget_tma_workspace_argrs @???????"}} !++E22I - -0C CD <+<+--    r\block_kctjjsdSdtjdjvr|dkrdSdS)Nrgfx942r~r9r)rUrr}rJr gcnArchName)rs rZget_default_kpackrsE = q5:33A66BBBwRT}}q 1r\rrCallowed_layout_dtypeslist[torch.dtype]c|j|vr!td|j|t|jjo|j|vot |jS)NzDNot using template since dtype %s is not in allowed layout dtypes %s)rrris_gpurr@r)rrs rZ_use_template_for_gpursl|000 R L !   v}!"" & L1 1 & v} % %r\backendc|dtjdDvS)Nc6g|]}|SrSrrWs rZr[z)_use_autotune_backend..- r\r)r%rtmax_autotune_gemm_backendsrrs rZ_use_autotune_backendrP ==??!<BBDDJJ3OO r\c|dtjdDvS)Nc6g|]}|SrSrrWs rZr[z._use_conv_autotune_backend..rr\r)r%rtmax_autotune_conv_backendsrrs rZ_use_conv_autotune_backendrrr\c|dtjdDvS)Nc6g|]}|SrSrrWs rZr[z9_use_conv_bwd_weight_autotune_backend..0   r\r)r%rt%max_autotune_conv_bwd_weight_backendsrrs rZ%_use_conv_bwd_weight_autotune_backendrsP ==??=CCEEKKCPP r\c|dtjdDvS)Nc6g|]}|SrSrrWs rZr[z8_use_conv_bwd_input_autotune_backend.. rr\r)r%rt$max_autotune_conv_bwd_input_backendsrrs rZ$_use_conv_bwd_input_autotune_backendr sP ==??<BBDDJJ3OO r\) enable_int32 enable_float8check_max_autotunerrrcddlm}m}tjtjtjg}|r.tjtjtjtjg}|r+|tj tj gt|j j ot||p|j j dko|j|vo?t jpt jp| o$t'do||j |jS)Nr9)BackendFeaturehas_backend_featurer=TRITON)rrrrUrrrrrrrrrr@rrrt max_autotunemax_autotune_gemmrTRITON_TEMPLATES)rrrrrr layout_dtypess rZuse_triton_templatersDCCCCCCC]ENEMBMT u{S Ge153DEFFFv})**A)&-@@ O "e+M 0M P VF$< VDV@V P "( + + P  ~/N O O r\ output_layout add_guardsmatricesrBr Layout | Nonerc ddlm}ddlmdfd dfd }d fd dfd dfd |o%t fd|Do ||S)u^ Return True iff *all* supplied tensors satisfy the CUDA TMA constraints that Triton relies on today. * https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__TENSOR__MEMORY.html A tensor is accepted when: * 1 ≤ rank ≤ 5 (cuTensorMapEncodeTiled) * dtype in _TMA_SUPPORTED_DTYPES (CUtensorMapDataType enum) * Base pointer 16-byte aligned * Exactly one contiguous ("inner") dim with stride 1 * All "outer" dims have 16-byte aligned strides * Inner dim size × itemsize is a multiple of 16 * For 1-byte dtypes (e.g. FP8), inner dim ≥ 32 r)has_triton_tma_devicer9r expr_bytesrorOrcNjj|tSr)rrstatically_known_multiple_of TMA_ALIGNMENT)rrs rZ_alignedzcan_use_tma.._alignedDsw<._is_tma_compatible_layoutGsR >4 - x && 5!!%%888r\r]rBc6|}|}|}|jjvrdS|x}|jdkr |||S|||S)NFrL)get_size get_stride get_dtyper'runaligned_buffers get_devicer@)r]r[rrm_devicerr_is_tma_compatible_xpus rZ_is_tma_compatible_matrixz.can_use_tma.._is_tma_compatible_matrixTs ,,..  ::<<174 4 45  &H 3 8N8N))%%@@ @!!%%888r\r[rQrSequence[_IntLike]rrcjt|}|j}|dks|dkrdS|tvrdSr? jj|} jj|}n fd|D} fd|D} fdt |D}t|dkrdS|d}t |D]\} } | |kr  | |zsdS||}  | |zsdS|dkr" jj| dsdSd S) Nr9rFcNg|]!}jj|"SrSrr!replace_backed_symbols_with_hints)rXrrs rZr[z;can_use_tma.._is_tma_compatible..ts9JK BB1EEr\cNg|]!}jj|"SrSr)rXstrs rZr[z;can_use_tma.._is_tma_compatible..ws9KM BB2FFr\cZg|]'\}}jj|d%|(Srrrstatically_known_equals)rXrrrs rZr[z;can_use_tma.._is_tma_compatible..|sH   2w77A>>    r\r T)r^itemsizerrr guard_int_seqrestatically_known_geq)r[rrrankrsizes_i strides_irT inner_idxrr inner_dimrrrs rZrz'can_use_tma.._is_tma_compatiblebs 5zz> !88taxx5 - - -5  g&44U;;G(66w??IIOTGQXI     "9--    u::??5!H y))  EArI~~8BM** uu I& x H,-- 5 q==!1!F!FyRT!U!U=5tr\c,|d}jj|}jj|dsdSd}|D]D}jj|}jj||rdSEdS)Nrr9FlT)rrrrstatically_known_gt) r[rr last_stridelast_stride_hint MAX_UINT32r_ size_hintrs rZrz+can_use_tma.._is_tma_compatible_xpus bk 7+MM   w778H!LL 5   D(JJ4PPIw33IzJJ uu tr\c3.K|]}|VdSrrS)rXr]rs rZrzcan_use_tma..s/??))!,,??????r\)rrorOr)rrrOr)r]rBrOr)r[rQrrrrrOr)torch.utils._tritonrrrr) rrrrrrrrrrs ` @@@@@rZ can_use_tmar/s/":99999XXXXXX 9 9 9 9 9 9 9 9 9 9 9 9 9 9 933333333j.  5 ????h??? ? ? 5 % %m 4 4r\ctjtjj}g}|D]U}t |t t jfr ||krdS-|t j ||V|sdSddl m }t|dkr|dn t j |}|r|jj|n|jj|S)NFTr9rr)rUiinforrrrrrrLerrr^Andrrguard_or_falsestatically_known_true)r[r int32_max conditionsr_r conditions rZ_descriptor_shape_fits_in_int32rs EK((,IJ99 dS%-0 1 1 9iuu    ehtY77 8 8 8 8 t!$ZA!5!5 1 59j;QI  ?'' 222 W  3 3I > >r\rcVtjjsdStd|DsdStfd|DsdStjjrt |jsdStjj dStjjr|nd}t||dS)NFc3`K|])}t|dkV*dS)rN)r^r)rXr]s rZrz*use_triton_tma_template..s788!s1::<<  A%888888r\c3^K|]'}t|V(dS)rN)rr)rXr]rs rZrz*use_triton_tma_template..sL  ( LLLr\rTr) rtrwenable_persistent_tma_matmulrenable_template_tma_storerr_rUrr}r)rrrrs ` rZuse_triton_tma_templaters = 5u 88x888 8 8u   u }.7Vz888u }$t$mE O]]4F : N N NNr\c`t|||dsdSddlm}ddlm}|o |S)NrFr)%has_triton_tensor_descriptor_host_tmar9is_datacenter_blackwell_arch)rrrcodegen.cuda.cuda_envr)rrrrrs rZ!use_triton_blackwell_tma_templatersm # :   uIIIIIICCCCCC 1 0 2 2 U7S7S7U7UUr\scale_option_ar8scale_option_b scaling_typeslist[ScalingType]c||vo||vSrrS)rrr s rZuse_triton_scaling_templater#s ] * N~/NNr\)maxsizech tjdduS#t$rYdSwxYw)zCheck if CuTeDSL is importable; cache the result for reuse. Call ensure_cute_available.cache_clear() after installing CuTeDSL in the same interpreter to retry the import. cutlassNF importlibutil find_specr8rSr\rZensure_cute_availabler+sD~'' 22$>> uu # 11c tjddu}n#t$rYdSwxYw|rt |S)zCheck if NVIDIA Universal GEMM (cutlass_api) is importable; cache the result for reuse. Call ensure_nv_universal_gemm_available.cache_clear() after installing cutlass_api in the same interpreter to retry the import. cutlass_apiNF)r(r)r*r8_ensure_fp4_dtype_registered) availables rZ"ensure_nv_universal_gemm_availabler1 s_N,,];;4G uu'$&&& s !$ 22cddl} |jtjdS#t t f$r&ddl|jjfd}||j_YdSwxYw)a$Patch cutlass_api to handle torch.float4_e2m1fn_x2 -> cutlass.Float4E2M1FN. NOTE: cutlass_api doesn't natively map this dtype. We patch the lookup function in-place so all callers (including TensorWrapper) pick up the change. Remove once cutlass_api adds native FP4 support. rNcH|tjkrjS|Sr)rUr Float4E2M1FN)r_origr&s rZ_patchedz._ensure_fp4_dtype_registered.._patched* s(...++5<< r\)cutlass_api.utilsrXcutlass_type_from_torch_typerUrKeyErrorrdr&)r.r6r5r&s @@rZr/r/ s B66u7MNNNNN n % B B B!>      :B 6666 Bs$.3A%$A%ch tjdduS#t$rYdSwxYw)aGCheck if nvMatmulHeuristics is importable; cache the result for reuse. nvMatmulHeuristics provides performance model-based kernel selection for NVIDIA GEMM operations. Call ensure_nvmatmul_heuristics_available.cache_clear() after installing nvMatmulHeuristics in the same interpreter to retry the import. nvMatmulHeuristicsNFr'rSr\rZ$ensure_nvmatmul_heuristics_availabler<2 sE~''(<==TII uur,mat_amat_ba_is_2db_is_2doffs Any | Nonebias scale_resultctsdStdsdSddlm}t |jjsdS|sdStjg} t|| sdStj stj sdSt|||sdStd||fDrdS|r|rdS|dS||dSdS) a Returns True if we can use the blackwell kernel for grouped mm. Required conditions: 1. CuTeDSL backend is enabled 2. CuTeDSL is available 3. We are on a blackwell arch 4. The dtype is bf16 5. Max autotune or max autotune gemm is enabled 6. A, B, and the output are 16B aligned 7. We are not using dynamic shapes 8. A is 2d 9. B is 3d 10. Offsets are provided 11. Bias and Scale are not provided FCUTEDSLr9r)rc34K|]}t|VdSr) is_dynamicrWs rZrz3use_blackwell_cutedsl_grouped_mm..t s( 1 1Q:a== 1 1 1 1 1 1r\NT)r+rrrrrr@rUrrrtrrrr) r=r>rr?r@rArCrDrrs rZ use_blackwell_cutedsl_grouped_mmrIB s"2 ! " "u  + +uCCCCCC &-$ % %u ' ' ) )u^$M  7 7u  6#;u ue6 : : :u 1 15%. 1 1 111u gu |u <3u 4r\r]rkcddlm}|jj||z|zd}|dks|t jjkrdSddlm }tj j rdStj tjtjg}t!||o&t jp t jot'd}|r6|s,t(d t jjdS|S) Nr9rrrrF)try_import_cutlassCUTLASSzFailed to import CUTLASS lib. Please check whether _inductor.config.cutlass.cutlass_dir %s is set correctly. Skipping CUTLASS backend for now.)rrrrrrtr&cutlass_backend_min_gemm_sizecodegen.cutlass.utilsrMrUrr}rrrrrrrrr cutlass_dir) rr]rrJr gemm_sizerMrrs rZuse_cutlass_templaterS s 221q519r2JJIA~~V^%QQQu999999 }u]ENEK@Mfm44 -  . s/ C Cc $ $S ) ) C C C C C Cr\c3XK|]$}|jjvV%dSr)r'rr)rXtrs rZrz1use_nv_universal_gemm_template.. s5 O O1::<<174 4 O O O O O Or\T)r/rXr+r1rrraot_compilationrr@rUrr}rtrrrr) rr]rrJr=r>rAr dims_to_checktensors_to_checkrrXs @@rZuse_nv_universal_gemm_templatera sU<POOOOO " "u - / /u  * *uu }V##u}'8#u  6#;u 1IM}Q C C C C] C C CCCuu~ %%% O O O O>N O O OOOu 4r\op_namectjj}|dkrdS|d|dDvS)z8Check if CUTLASS should be used for the given operation.ALLTc6g|]}|SrSrrWs rZr[z'_use_cutlass_for_op.. s IIIQqwwyyIIIr\r)rtr&cutlass_enabled_opsr%r)rb enabled_opss rZ_use_cutlass_for_oprh sX.4::<.< s8  k ! !g&<&< &<&<&) %%Q 3;;  EAI ! # #  $ $Q ' ' ' ' RZ1__  % %a ( ( ( (  ! !! $ $ $ $ , << #55FF#&88>IK  ''r\cJtj|jSr)rUrJrrrs rZ_rocm_native_device_arch_namer] s : + +F 3 3 ??r\Ntuple[str | None, Callable[[], list[Any]], Callable[[], list[Any]], type[Any]]c ddl}ddlm}m}ddlm}t j|j }n'#t$rd d}d d}Gdd }d}YnwxYw||||fS) Nr)gen_ops_librarygen_ops_preselected)CKGemmOperationrOrcgSrrSrSr\rZrz*try_import_ck_lib..gen_ops_librarys Ir\cgSrrSrSr\rZrz.try_import_ck_lib..gen_ops_preselectedv rr\ceZdZdS)*try_import_ck_lib..CKGemmOperationN)rrrrSr\rZrry s Dr\r)rOr) ck4inductor(ck4inductor.universal_gemm.gen_instancesrrck4inductor.universal_gemm.oprrkrzdirname__file__r8)rrrrpackage_dirnames rZtry_import_ck_librb s              '//+*>??                      O-@/ QQs69!AAcftjstjsdStjjsdS|jjdkrdSt|j}dtj j Dp| dd|ifd tj j zD}|sdS|jtjtjtjfvrdSt%\}}}}|st&ddS|tj _dS) NFrJcFi|]}|dd|S):r)r)rXrJs rZrkz#use_ck_template.. s(DDDaqwws||ADDDr\rrc g|] }| SrSrS)rXrJrequested_archss rZr[z#use_ck_template.. s.!!!  !!!r\z,Please pip install Composable Kernel packageT)rtrrrUrr}rr@rrrarchrrck_supported_archrrrrrrrck_dir)r native_archrequested_supported_archsck_package_dirnamerrs @rZuse_ck_templater sA  6#;u = u }V##u0 >>KDD6;3CDDD#q!;IO!!!! %%''&+*GG!!! %u |EM5>5=IIIu"3"5"51a  BCCCu+FK 4r\cddlm}tdo9t|o*|jj||z|zddkS)Nr9rCKrrLrrrrrrrrrr]rrJrs rZuse_ck_gemm_templater se d## K F # # K G  . .q1uqy2 . F F Jr\cddlm}tdo9t|o*|jj||z|zddkS)Nr9rCKTILErrLrrrs rZuse_ck_tile_gemm_templater se h'' K F # # K G  . .q1uqy2 . F F Jr\c>tdot|S)Nr)rrrs rZuse_ck_conv_templater s %d + + G0G0GGr\cRtjp tjo|jjdkSr)rtrrrr@rs rZ_use_template_for_cpur s'7v7 & -  %&r\mat1ReinterpretView | Buffermat2cddlm}t|j|sJ|jj}|jj}t |oe|tj koCt|dko0t|dko|d|dko |ddk}t|||do|j p|S)Nr9)rCrF)require_constant_mat2) r\rCrrr_rrrrUrr^use_cpp_gemm_template is_contiguous)rrrrC mat1_size mat1_stridemat1_each_batch_is_contiguouss rZuse_cpp_bmm_templater s dk6 * *** *   I+$Kf%% " NN   - " ^^q  "   " "^y| +  " ^q " !t5 Q Q Q  !!##D'Dr\mat2_transposedr is_woq_int4 q_group_sizec vddlm}ddlm}ddlm} ddlm} t|rtdsdStj j sdS| tjtjfv} tjtjtjtjtjg} | ||| r|jnd||\} }}}}}t+||frdSt-||jr|}| | \}}|d | ||| | |t3| | }dd}|j| vo:|duo6||o+t-||jo|p| S)Nr9r[)create_micro_gemm)*get_gemm_template_output_and_compute_dtype)mm_argsCPPF) out_dtyper use_4x2_dim micro_gemm) input_dtype input2_dtype output_dtype num_threadsuse_refrrYrBrOrcf||ddkS)Nrr9) freeze_layoutrrYs rZis_last_dim_stride1z2use_cpp_gemm_template..is_last_dim_stride1 s* ||~~b!Q&&r\)rYrBrOr)r`r\codegen.cpp_micro_gemmrcodegen.cpp_utilsrkernel.mm_commonrrrrtcppweight_prepackrrUrrrrhalfrhas_free_symbolsrBaseView unwrap_viewparallel_num_threadsris_module_buffer)rrrrrrrr\rrr int8_gemmrr]rrJrrrrs rZrr s$999999MMMMMM))))))  ( (0Ee0L0Lu : $u  U[%*$==I]ENEJ UZXM")'  "+5&,,' ###Aq!VT4Au$ $$"!!@@AQAQRROL!""   NN$$^^%%!(**!   J''''  % C d " C   % % C tR] + + C " " $ $ A,A(A r\cRtjp tj ptdS)NATEN)rtrrrrSr\rZuse_aten_gemm_kernelsr s*7v7  ' v & &'r\cReZdZUejdZded<d dZd dZd d Z d S)DebugDirManagerrrPprev_debug_namerOr cBttj|_dSr)r3rcounterrrs rZrzDebugDirManager.__init__& s.//r\ctjjj|_|jd|j|_|jtjj_dS)N_tmp_)rU_dynamortdebug_dir_rootrrnew_namers rZ __enter__zDebugDirManager.__enter__) sA$}3B/??dg?? .2m +++r\rr_cntj|j|jtjj_dSr)rrrrrUrrtr)r7rs rZ__exit__zDebugDirManager.__exit__. s* dm$$$.2.B +++r\Nr.)rr_rOr ) rrrrrrrrrrrSr\rZrr" suioa  G0000<<<< CCCCCCr\rCallable[P, _T]r3r4tuple[_T, list[str]]cddlm}td fd }tj|d|5t j||i|}dddn #1swxYwY|tfS) Nr9r>coderPrOr c2|dSr)rr source_codess rZsave_output_codez*run_and_get_code..save_output_code< sr\rrrPrOr ) rr?r!rpatchr_rUrresetr)rrrr?rr.rs @rZrun_and_get_coder3 s %$$$$$$.LLL   =*<>N O O%% T$V$$%%%%%%%%%%%%%%% 4 %% %%s'A00A47A4c|dd}t|g|Ri|\}}g}|D]}tjrGtjjdur4|tjd|tj n3|tjd|tj |r d|D}||fS)N remove_quoteFTzR"TRITON\((.*?)\)TRITON"z '''.*?'''c"g|] }|dd S)rrS)rXrRs rZr[z'run_and_get_kernels..T s :::vad|:::r\) r_rrtrnrwautotune_at_compile_timerrfindallDOTALL)rrrrr.rkernelsrs rZrun_and_get_kernelsrE s::ne44L+B@@@@@@FLG;;   F&-"HPT"T"T NN2:&A4SS T T T T NN2:lD")DD E E E  ;::':::G 7?r\tuple[Any, list[str]]c.dfd }t|S)NrOrch}||Sr)rXrp)r.rs rZrun_with_backwardz1run_fw_bw_and_get_code..run_with_backwardY s-  r\)rOr)r)rr s` rZrun_fw_bw_and_get_coder X s2 - . ..r\cbddlm}gdfd dfd }tj|d |5tj|d 5t j||i|}dddn #1swxYwYdddn #1swxYwYS)zLGet the inductor-generated code, but skip any actual compilation or running.r9r>rrPrOr c2|dSrrrs rZrz"get_code..save_output_codeg sD!!!!!r\r7r?rcGdd}|jr|n|\}}|j|r|j|S)Nc"eZdZdZd dZd dZd S) @get_code..patched_compile_to_module..DummyModulez4This is empty to replace the generated triton modulerOr cdSrrSrs rZrzIget_code..patched_compile_to_module..DummyModule.__init__n sr\rrrcdSrrSr6s rZcallzEget_code..patched_compile_to_module..DummyModule.callq sr\Nr.rrrrrOr )rrrrrrrSr\rZ DummyModulerk sB F F          r\r)rncodegen_with_cpp_wrappercodegenr)r7r wrapper_code kernel_coders rZpatched_compile_to_modulez+get_code..patched_compile_to_modulej s        04/? SD ) ) + + +T\\^^ " k +,,,  0  [. / / /{}}r\compile_to_modulerNr)r7r?rOr)rr?rrr_rUrr)rrrr?rrrrs @@rZget_codera st$$$$$$ L"""""",  .0I     -);=MNN     B                                  s5"B$'B  B$ B B$B B$$B(+B(ct|g|Ri|}dt|cxkrdksnJdt||dSNr9rz%expected one or two code outputs got r)rr^)rrrrs rZget_triton_coder ssB000000L L!! & & & &Q & & & & &CL0A0ACC ' & & ?r\ct|g|Ri|\}}dt|cxkrdksnJdt||dSr)rr^)rrrrrs rZrun_and_get_triton_coder! sy'r;D;;;F;;OA| L!! & & & &Q & & & & &CL0A0ACC ' & & ?r\tuple[Any, list[GraphLowering]]cddlmddlm}|jgd fd }t j|d |5||i|}dddn #1swxYwY|fS) Nrr>rFrrrrOr cv|i||d}t|sJ|dS)Nr)rr)rrrr?graph_lowerings real_inits rZ fake_initz-run_and_get_graph_lowering..fake_init sQ 4"6"""Q%/////u%%%%%r\rr)torch._inductor.graphr?torch._inductor.output_coderGrrrr_) rrrrGr'r.r?r%r&s @@@rZrun_and_get_graph_loweringr* s433333;;;;;;(IO&&&&&&&&   ?J B B%%T$V$$%%%%%%%%%%%%%%% ? ""s AAAaten_op override_fnc#Kddlm}|j|} tj|||j|<dV||j|<dS#||j|<wxYw)z Override the lowering of aten_op with override_fn. The first argument of override_fn is the original lowering fn. r)loweringN)torch._inductorr. loweringsrHpartial)r+r,r.orig_fns rZoverride_loweringr3 s{)((((( )G.&/&7 W&M&M7# &-7###g7#----s !A Apre_fnpost_fnCallable[..., Any] | Nonec~ddlm}|jd fd }tjj|d|S) zr Add hook functions to be called at the beginning and end of Scheduler.__init__. Used for unit tests. r) Schedulerr&rrrOcT||||}r |||SrrS)r&routr2r5r4s rZrKz(add_scheduler_init_hook..wrapper sEy%   gi''  & GIu % % % r\r)r&rrrrOr)torch._inductor.schedulerr8runittestrrr_)r4r5r8rKr2s`` @rZadd_scheduler_init_hookr= sh433333 G =  % %iW E EEr\msgctjrt|dSt|dS)z Warnings that will be actionable for PyTorch developers, but not end users. Allows us to easily disable them in stable releases but keep them on for nightly builds. N)rtdeveloper_warningsrrinfo)r>s rZdeveloper_warningrB s:   C  r\c tjd}|dzttjkrZttj|dzdkr4tj|dzddkrtj|dzSn#t$rYnwxYwtjD]0}|dr|tddcS1dS)a An experimental API used only when config.benchmark_kernel is true. The benchmark name is only available at codegen time. So we can not directly call it in benchmark_all_kernels which is run after codegen. The function assumes the argument after --only is the benchmark name. It works for torchbench.py/hugginface.py/timm_models.py. But for ad-hoc scripts, this function may return None. There are 2 flavors of --only argument we need to handle: 1. --only model_name 2. --only=model_name z--onlyr9rrz--only=N)rHargvr?r^ ValueErrorrJ)rrs rZget_benchmark_namerF s  hnnX&& !Gc#(mm # #CHS1W%&&**q!!$++8C!G$ $      x)) >>) $ $ )s9~~''( ( ( ( ) 4sBB B)(B)rc4td|DS)Nc3"K|] }|dkV dSrdrSrWs rZrzis_ones.. &%%!qAv%%%%%%r\rrs rZis_onesrL  %%u%%% % %%r\c4td|DS)Nc3"K|] }|dkV dS)rNrSrWs rZrzis_zeros.. rIr\rJrKs rZis_zerosrP rMr\inputsSequence[torch.Tensor]c4td|DS)Nc3K|];}t|tj|jtjdkV. sY  dEL ) ) u|E***r\rJ)rQs rZ is_cpu_devicerU s0   r\ct|tjs Jd|jr tjStjS)Nz8only support sympy.Expr as input to get_sympy_Expr_dtype)rrrsrrUrrrs rZget_sympy_Expr_dtyperW sG c5: & &B & ~{}r\should_profile Iterator[Any]c/K|r5tjj|i|5}|VddddS#1swxYwYdSdVdSr)rUrr)rXrrrs rZ maybe_profiler[% s ^ #T 4V 4 4 GGG                   s -11cZtjj}|dkrtj}|SNr9)rtrthreadsrUget_num_threads)r^s rZrr. s(j G{{')) Nr\ctddlm}|}|dtjjrdndS)Nr9)get_backend_options num_stagesrr)runtime.triton_helpersrarrUrr})raoptionss rZget_backend_num_stagesre5 sD;;;;;;!!##G ;;|%-*;%BQQ C CCr\ct|tjjjjdk}||StjstddSddl m }m }tjo!tj dk}|tj tjtjfvsJt!j|jd rdd lm}|}|tj tjfvr|r |||Stjjjjdkr|tj|S|tj|S|tj tjfvr |r ||Stjjjjdkr|tjS|tjS) z We don't want to throw errors in this function. First check to see if the device is in device_info.py, then fall back to the inaccurate triton estimation. tf32)is_tf32Nzget_device_tflops: no Triton fallback available for non-CUDA devices. Returning 0.0; roofline estimates will use memory bandwidth only.rr)get_max_simd_tflopsget_max_tensorcore_tflops)rr clock_rate)max_clock_rate)rrUbackendsrJmatmulfp32_precisionrVrrtriton.testingrirjrrrrinspect signature parametersrtorch._utils_internalrl)rds_topsrirj SM80OrLaterrlsm_clocks rZget_device_tflopsrx= s  u~*1@FJG : " " $ $  P   sMMMMMMMM*))++ 0P0P0R0RW1K U]ENEMB B B B B,--8<<\JJ6888888!>## U]EN3 3 3 3,,UH== = >  % 4 > >,,U]HEE E&&u}h?? ? U]EN3 3 3 3,,U33 3 >  % 4 > >,,U];; ;&&u}55 5r\c"ddlm}|S)Nr get_dram_gbps)rpr{rzs rZget_gpu_dram_gbpsr|o s ,,,,,, =??r\ctddlm}|jjdddS)NrrVmax_shared_mem)triton.runtimerVrWrXrrr~s rZget_gpu_shared_memoryrv s>%%%%%% =  4 4Q 7 7 ; ;. s7JJz!UZ((<_JJJJJJr\r()rs rZrr s JJcJJJ J JJr\c tddlm}|D]}t||j|j|j|j|jfrJt| pds#t| pdrdSzt||j stdt|dS)Nr9r[rSTzunexpected type for is_dynamic F)r`r\rrrrrr@rmaybe_get_sizemaybe_get_striderB TypeErrorr@)rr\r]s rZrHrH s  I I   bmR[":KRYW  I  0 0 2 2 8b99 =M""$$*>> tt Ary)) I Gd1ggGGHH H 5r\ceZdZdZdZdS) Placeholder KERNEL_NAMEDESCRIPTIVE_NAMEN)rrrrrrSr\rZrr s K*r\rrr6inpcddlm}tjdd5}t j}t j}t |t|j|td|j |t|j |tj }t||5||j dddn #1swxYwYtj |z } ||j |j |td |j |t|j |||k} t"d ||j| | ddddS#1swxYwYdS) Nr9)stable_topological_sortrzutf-8)modeencoding)r fake_modezBefore: )filezAfter: zZ%s, save before/after graph to %s, graph before/after are the same = %s, time elapsed = %s)pattern_matcherrrNamedTemporaryFileior rhrd propagater rr nowrglint recompilerrrAr) rrrr>rr| before_ioafter_io start_time time_elapsedr]s rZpass_execution_and_saver sA988888  $      KMM ;==C R#3C#8#8999CSII $"($$1---- bhY''''\^^ #B , ,   DNNN               |~~ 2 )))    ###!,,,, bhX&&&&    H$5$5$7$7 7  h  F     +                  s7BG<C GC G C !CGGG input_bufBuffer | Operation | Nonec`ddlm}t||jo|S)zB Check if input buffer is a multi-outputs template buffer r9r[)r`r\rris_multi_outputs_templaterr\s rZrr sA 9b/00 2  / / 1 1r\cddlm}t||jo1t |jdkot |jdS)zL Check if input buffer is a output of multi-outputs template buffer r9r[r)r`r\r MultiOutputr^rQrrs rZ#is_output_of_multi_outputs_templater s\  9bn-- ;   ! !Q & ; %i&6q&9 : :r\rNode | Operation | Nonetorch._ops.OperatorBase | NonecD|dSddlm}t||jo"t||j o |dup|j|upt ||juottj j do#|jtj j j j kpttj j do#|jtj j jj kpBttj j do#|jtj j jj kS)NFr9r[all_to_all_singleall_gather_into_tensorreduce_scatter_tensor)r`r\r_CollectiveKernel _WaitKernel op_overloadr@FallbackKernelr6rUr7torchrecrdefaultrrrrr\s rZ is_collectiver s% |u 4-.. 3400 0 3 4Z 14+r1  T b''  *,?@@U$ (:(L(TT   *,DEEE$9%<DE  *,CDDY$ (:(P(XX/r\IRNode | Operation | Nonec:ddlm}t||juSNr9r[)r`r\r@r)rr\s rZis_waitr) s' :: ''r\snoderH filter_fn*Callable[[BaseSchedulerNode], bool] | Nonecddlm}t||rtd|jDSt |jo|dup ||S)NrGroupedSchedulerNodec34K|]}t|VdSr)contains_collectiverWs rZrz&contains_collective..6 s+@@a&q))@@@@@@r\)r;rrrsnodesrr)rrrs rZrr/ sr?>>>>>%-..A@@5<@@@@@@  $ $ P)t*;*Oyy?O?OPr\cddlm}t||rtd|jDSt |jS)Nrrc34K|]}t|VdSr) contains_waitrWs rZrz contains_wait..? s*::=##::::::r\)r;rrrrrr)rrs rZrr; sV>>>>>>%-..#::U\::::::uz"""r\Operation | None9torch._ops.OpOverload | Collection[torch._ops.OpOverload]cddlm}t|tjjr|g}t||jo|j|vSr)r`r\rrUrrrrrs rZis_fallback_oprD sT"ej+,,T dB- . . I43Cr3IIr\buf_name name_to_bufname_to_fused_nodecL|||jSr) defining_opr')rrrs rZbuf_name_to_fused_snoderO s# k(3?HHJJ KKr\cdSrrSrs rZrrZ ur\collected_node_setMutableSet[BaseSchedulerNode]dict[str, SchedulerBuffer]dict[str, BaseSchedulerNode] criteria_cbCallable[[Any], bool]c||rdS|||jD]1}t|j||}||vrt |||||2dS)Nr)runmet_dependenciesrrfind_recursive_deps_of_node)rrrrrdepdefining_op_for_deps rZrrU s{55!!!'   5 Hk#5   "4 4 4 #    #        r\cdSrrSrs rZrrs rr\c ||rdS|||D]}|jD]}}|jJ|jdkr)|j|vrE||j}||vrit |||||~dS)NOUTPUTr)r get_outputsrrr'find_recursive_users_of_node)rrrrroruser_ops rZrrn s{55!!!    G  D9(((y!!##x//y!!##+===(););)=)=>G,,, (""'       r\dynamo_gm_num_inputsaot_fw_gm_num_inputscFtjjjrdnd}||z |z S)zaComputes the number of inputs to the aot fw graph which have fixed addresses (params and buffers)rr)rU _functorchrtfunctionalize_rng_ops)rrnum_rng_seed_offset_inputss rZnum_fw_fixed_argumentsr s2  $ :A "6 69S SSr\fx_gc d d}d}g}|jjD]2}|jdkr%||r|||dz }3|t t t |ksJt |S) z> Infers which inputs are static for a backwards graph rYr7rOrcJd|jvod|jvod|jvod|jvS)Ntangentsbwd_seedbwd_base_offset bwd_rng_staterrs rZis_saved_tensorz'count_tangents..is_saved_tensor s@ af $ .!&( .!/ . qv-  r\rrr9)rYr7rOr)rrrrrrr^)rr  arg_countstatic_arg_idxsrs rZcount_tangentsr  s     IO Z  4= q!! 2&&y111 NI d5_)=)=#>#>?? ? ? ? ?   r\cg}t|jjD]A\}}|jdkrn0|jdr||B|S)a Returns indices of backward graph inputs that are always at fixed addresses: primals (parameters/buffers/user inputs saved for backward). Excludes saved activations which may not be at fixed addresses when the forward is partitioned for CUDA graphs. rprimals_)rerrrrrJr)r static_idxsrrs rZget_static_bw_input_idxsr snKDJ,--$$Q 4= E 6  Z ( ( $   s # # # r\c:eZdZUded<d dZed dZd S) BoxedBoolrrrOc|jSr)rrs rZrzBoxedBool.__bool__ s zr\rrBoxedBool | boolcBt|tr d|_|SdSr)rrrres rZdisablezBoxedBool.disable s% c9 % % CIJur\Nr/)rrrOr)rrrrrr:rrSr\rZrr sSKKK\r\r kernel_listc#Kddlm}|j ddfd }tj|d|5dVddddS#1swxYwYdS)Nr9r;Tr7r< kernel_namerPrrrrgpurcpp_definitionrOrcN|||||||Srr )r7rrrrrrorig_define_kernels rZ define_kernelz.collect_defined_kernels..define_kernel s; ;'''!! +{Hc>   r\r)NTN)r7r<rrPrrPrrrrrrrrrOr)codegen.wrapperr<rrrr_)rr<rrs` @rZcollect_defined_kernelsr! s555555-; $%)            /- P P sAAAc |dzS)N __original__rSrs rZ get_cloned_parameter_buffer_namer$ s .  r\c|tvSr)r]rs rZrr s Y r\c(tjjduS)z,Check if we're running on ROCm/HIP platform.N)rUrr}rSr\rZis_rocmr' s = D ((r\ backend_namer int | strrtuple[str, ...] | Nonec$tsdS ddlm}ddlm}||||}||i}n$#t tttf$rYdSwxYwt|dd}|dSt|S)Nr) GPUTarget) make_backendsupported_fp8_dtypes) r$triton.backends.compilerr,triton.compiler.compilerr- parse_optionsrdr8r9rrTr?)r(rrr,r-rrdr.s rZ_triton_supported_fp8_dtypesr2 s   t666666999999<y99,v&&44R88 K9 =tt#7,BDII#t % & &&s7A A+*A+)triton_backend triton_archrtorch.device | str | Noner3r4int | str | Nonec|tjtjtjtjfvrdSt |d}||dStj|}|jdkrdStj sdS|j }|tj }tj |}tr#d}|jdd}nd}|jdz|jz}t'|dd}|t)d | |dkrd nd }t+|||}|dS||vS) NTrrJr}rrrrz0triton_arch must be provided with triton_backendrr)rUrrrrr removeprefixrr@rJrVr?rBrr'rrrminorrTrEr2) rrr3r4r triton_dtype device_index propertiesr.s rZis_triton_fp8_dtype_supportedr= s       tE??//44L >4f%% ;& 4z&&(( 4|   :4466LZ55lCC 99 C"N$066s;;A>KK#N$*R/*2BBKJ T:: KLLL(E11BBr 7 Y#t / //r\c,|dkot|S)NrK)rrs rZdevice_need_guardr?=s U? -vf~~-r\c6|tjkr@tjr"tjdkS|tjkr tjrdS|tjtjfvS)N)rrT)rUrrJrVrrLrrrs rZ,needs_fallback_due_to_atomic_add_limitationsrAAst 5:#:#:#<#<z//11F:: %. UY%;%;%=%= tej111r\r self_dtype src_dtypesrc_device_type src_is_tensorc*|jtjjjtjjjfvr|dS|jtjjjkrdnd}|d|fvp|ot|ot|p|jtjjjkoA|dko;|o9|dko3tj j o"tj j ptdkp2||ko|tjtjfvptjS)NFrrXr=r9)overloadpacketrUr7atenscatter_reduce_scatter_reducescatter_rrArtrfallback_scatter_reduce_sumdynamic_threadsrrr$are_deterministic_algorithms_enabled)rrrBrCrDrE reduce_tys rZuse_scatter_fallbackrPJs= " IN *EIN,I J K K  "u+uy~/FFFE tY// 8  H'' H&,#?#? ??RVWW  r\ example_inputcnt|jjsdStjpt |Sr)rrr@rtassume_aligned_inputsr`)ras rZshould_assume_input_alignedrds6 -&+ , ,u  ' K+<]+K+KKr\rctjj}|st jS|jr |jjst jS|jj}|Sr) rU_guardsTracingContexttry_getr  nullcontextrrsuppress_guards)tracing_contextrs rZ#maybe_get_suppress_shape_guards_ctxrlsw m2::<.runsZ#9 );$ $  [eJ {   ;  k : : : r\)rrrOr)r^)rrrrSs``` rZalign_inputs_from_check_idxsrsN  ?q            Jr\cd|vrd}nNtdt||Ddz}t j||fd}t j|||S)Nrc3,K|]\}}|dz |zVdSrdrS)rXrrs rZrz)clone_preserve_strides.. s/TTf$TTTTTTr\r9r)r_rXrrrU as_stridedclone)rY needed_sizers rZclone_preserve_stridesrsAFFHH}}  TT#affhh :S:STTT T TWX X  a+ 6 6 < < > >F  FAFFHHahhjj 9 99r\rrcheck_inputs_idxsreturn_pair_idxsOrderedSet[int] | None-tuple[list[torch.Tensor], list[torch.Tensor]]cdg}g}|du}|D]}||}t|tjsJdt||t zrHt |||<|r4||vr0|||||||fS)z Clones misaligned tensors which we inferred were aligned. Returns a tuple of [old_tensors], [new_tensors] for every cloned tensor which is in `return_pair_idxs`. Nz Expected tensors only, but got: )rrUrr@data_ptr ALIGNMENTrr)rrrrrret_pair_definedr_inps rZrrs')K&(K(t3  2 2!}$ --  ;tDzz ; ;  - ==??Y & 22488JqM 2A)9$9$9""4((("":a=111  ##r\static_input_idxscg}|D]Y}||}t|tjr5|tzdkr||Zt |t |kr|S|S)z[ We require all inputs to be aligned, so introduce a copy for any that aren't. r)rrUrrrrr^)rQraligned_static_input_idxsrr}s rZremove_unaligned_input_idxsr1s!# 22s  eU\ * * 20@0@90LQR/R/R % , ,S 1 1 1 $%%->)?)???(( r\rcddlm}tjtjj}|jjj}|jjj j }tj r"|jj ||dS|jj||krdS|jr%|jj|dkrdS||o|||kS)zCheck if an expression fits within 32-bit integer range. NOTE: This function intentionally does not install guards. Callers are responsible for guarding (e.g. via check_leq) when they decide to use 32-bit indexing based on this result. r9rTg@xDF)rrrUrrrrrguarding_hint_or_throwrhas_guarding_hintrtassume_32bit_indexing check_leqr r^)rrint_maxrrs rZexpr_fits_within_32bitrCsk%+&&*GW-D(2D # ""1g...t w--a7l;;t  7  1 1!d( ; ; 5  Q   H$:$:1$=$=$HHr\compiled_graphrGctjj}||jt |jdksJt ||jJ|jD]}||jddtjjx}r|jd fd |jtfd|DdSdSdS) NrFrrrO)float | int | SymInt | SymFloat | SymBoolc~t|Sr|S|Sr)rdeserialize_symexprevaluate_symexpr)rfakify_first_callrs rZmap_exprz4set_tracing_context_output_strides..map_exprsE ("1vv (@(<.s+55!((1++555555r\)rrrOr) rUrfrgrhoutput_stridesr^r~rrr?)rrrrrJrrrs @@@rZ"set_tracing_context_output_stridesrps>m*2244Gw5A7)**a////).99 ,888#2  E}&--d3333$)!-6>>@@@3>(+(=%9999999&--5555u55555'AA  r\ctj tjStjsdStjrdS ddlm}n#t$rYdSwxYw|tj dkS)NFrREMOTE_CACHE_VERSIONz.pytorch/remote_cache:fx_graph_memcache_version) rtfx_graph_remote_cache is_fbcoderU_utils_internalis_fb_unit_testtorch._inductor.fb.remote_cacherModuleNotFoundErrorjustknobs_getval_intrs rZ should_use_remote_fx_graph_cachers #/++    u ,,..uHHHHHHH uu 5#8#M#M8$$ sA A$#A$c.tjdd|S)Nz [^a-zA-Z0-9_]r)rsubrs rZnormalize_namers 6"C . ..r\ztl.int1z tl.float8e4nvz tl.float8e5z tl.float8e4b8ztl.float8e5b16ztl.uint8)ztl.boolztl.float8_e4m3fnztl.float8_e5m2ztl.float8_e4m3fnuzztl.float8_e5m2fnuzztl.float8_e8m0fnuztl.float4_e2m1fn_x2ci|]\}}|| SrSrS)rXrJrs rZrkrksGGG$!QAGGGr\z^.*[.]ctdt|}t||S)z"Convert torch.dtype to triton typetl.)_triton_type_rerrP_triton_type_mappingr)rtriton_type_names rZ triton_typers6&**5#e**==  # #$46F G GGr\ct||}|dd}tt|}t |tjsJ|S)Nrr`)_torch_triton_mappingrr~rTrUrr)r adjusted_type type_namers rZtriton_type_to_torchrsY)--eU;;M%%eR00Iy))I i - --- - r\rc|j o||ko||ko|j|jko|j|jkow||ko)||kSr) is_mkldnnr_rrruntyped_storagerr^rrs rZis_same_tensorrs N < IIKK5::<< ' < KKMMU\\^^ + < J%+ % < K5< '  <  " " + + - -1F1F1H1H1Q1Q1S1S S  <    ! !U%9%9%;%; ;r\c8|jo||koi|j|jkoY|j|jkoItjj|tjj|kSr)rr_rrrUr7mkldnnrrs rZis_same_mkldnn_tensorrs  P IIKK5::<< ' P J%+ % P K5< ' P I  % %d + +uy/?/H/H/O/O O r\tuple[str, ...]cdS)N)risnan logical_not logical_andsignbitand_leltgegteqneraxorrSrSr\rZ boolean_opsrs   r\c$eZdZUded<ded<dS) OpDtypeRuler5type_promotion_kindrwoverride_return_dtypeNrrSr\rZrrs*8888------r\rzdict[str, OpDtypeRule]op_dtype_propagation_rulesrr5rc6t||t|<dSr)rr)rrrs rZ#register_op_dtype_propagation_rulesrs% (32((t$$$r\zOrderedSet[str]op_requires_libdevice_fp64c:t|dSr)rrrs rZ#register_op_requires_libdevice_fp64rs""4(((((r\rcddlm}|s|jj}|dkr t jS|dkrdS|dkr t jS|dkr t jSt j S)Nrrr=rKrLtpu) rkrrget_current_device_or_throwr@rt cpu_backend xpu_backend tpu_backend cuda_backend)rrs rZget_current_backendrs------ Ag99;;@ e!!   u   !!   !!""r\torch.device | Nonech|/|jdkr$tj|jSdS)z+Check if the given device supports float64.NrLT)r@rUrLrhas_fp64rs rZdevice_supports_fp64rs2 fkU22y..v66?? 4r\c|tjtjfvr/tjjrt dkr tjS|S)z"Maybe upcast [b]float16 to float32rw)rUrrrtrwcodegen_upcast_to_fp32rrrs rZupcast_compute_typer sC %-000 M 0 1  ! !X - -} Lr\KeyTypeValTypec\eZdZdZddZdd Zdd ZddZdddZd dZ d!dZ d"dZ d#dZ dS)$ ScopedDictz A dictionary-like object that allows for scoped updates. It maintains an original dictionary and a set of new items that can override the original items within the scope. The original dictionary is unmodified. original_dictMapping[KeyType, ValType]c"||_i|_dSrr new_items)r7rs rZrzScopedDict.__init__7s*13r\rvrrOrcH||jvr |j|S|j|SrrrrDs rZrzScopedDict.__getitem__;s* $. >#& &!#&&r\rr c||j|<dSr)r)r7rvrs rZ __setitem__zScopedDict.__setitem__@s#sr\r_rc&||jvp||jvSrrrDs rZ __contains__zScopedDict.__contains__Csdn$At/A(AAr\NrValType | Nonecd||jvr |j|S|j||Sr)rrr)r7rvrs rZrzScopedDict.getFs5 $. >#& &!%%c7333r\rc^t|j}|jD]}||jvr|dz }|Sr])r^rr)r7rrJs rZruzScopedDict.__len__Ks@ " # #  A***Qr\Iterator[KeyType]c#RK|jEd{V|jD]}||jvr|VdSrr)r7rJs rZ__iter__zScopedDict.__iter__RsT%%%%%%%%  A***  r\c8t|jp|jSr)rrrrs rZrzScopedDict.__bool__XsD&8$.999r\ctrrjrDs rZ __delitem__zScopedDict.__delitem__[s!!r\)rr)rvrrOr)rvrrrrOr )rvr_rOrr)rvrrr rOr rw)rOr r/)rvrrOr ) rrrrrrrr rrurrrrSr\rZrr/s4444'''' $$$$BBBB44444  ::::""""""r\r)frozen_defaultrtype[Any] | Nonerc.dfd }||S||S)NrrxrOc2tj|dS)NT)kw_onlyr) dataclasses dataclass)rrs rZwrapzir_dataclass..wrapas$S$vFFFFr\)rrxrOrxrS)rrrs ` rZ ir_dataclassr_sAGGGGGG { 499r\list[int] | Nonecvtjj}||jr |jjSdSr)rUrfrgrh fw_metadatabw_donated_idxs)rks rZget_donated_idxsr is7m2::<rrr)rr6r normalizedrrrZrs rZ,_fx_node_is_input_dependent_cudagraph_unsafer=s=<<<<< ^F fej3 4 4u  ( !) (0 (' GL'.t     !"IAvY'G  ?sx'<JKA(( 44 5r\c|j}t|tvrdSt|tjjrtjjj |j vrdSt|rdS|j dx}Jt|ttfs|gn|}|D]&}t|tjr |jrdS'dS)a  Check if an FX node is cudagraph-unsafe. This includes: - Ops in FORBIDDEN_CUDAGRAPH_OPS (CPU sync, dynamic alloc, etc.) - Ops with the cudagraph_unsafe tag - Input-dependent unsafe ops (e.g., index_put with boolean indices) - Ops with sparse tensor outputs TrNF)rrPFORBIDDEN_CUDAGRAPH_OPSrrUrrr4rcudagraph_unsaferr=r>rrr?r is_sparse)rrrvalsrs rZr0r0s^F 6{{---t 65:011 HL )V[ 8 8t4G<<t|&&&3&sT5M::Cuu  A!U\** q{ tt 5r\rDcddlm}t||j|jfrdSt||j|jfsdSt|dd}|t|rdSdS)ah Returns True if the node is an op that is not cudagraphable. This includes: - Ops in FORBIDDEN_CUDAGRAPH_OPS (CPU sync, dynamic alloc, etc.) - Ops with the cudagraph_unsafe tag - index_put_ with boolean indices (triggers .nonzero() during capture) - Control flow nodes (Conditional, WhileLoop) - Ops with sparse tensor outputs r9r[TFrN) r`r\r Conditional WhileLooprrArTr0)rr\rs rZis_cudagraph_unsafe_oprFs$677t dR.@ A AudIt,,G:7CCt 5r\ctjdd}tjrXddlm}|}|rFtj|dd}|r!tj ||gn|}|S)NLD_LIBRARY_PATHr`r)get_runtime_pathrMlib) rkrlrrtrlibfb.py.parutilrIrzrwpathsep)rzrI runtime_pathlib_paths rZget_ld_library_pathrOs :>>+R 0 0D K555555''))  Kw||L)UCCH8<J2:??Hd#3444(D Kr\c@ddlm}t||o|jduS)Nr)SubgraphPythonWrapperCodegen)torch._inductor.codegen.wrapperrQrpartition_signatures)rKrQs rZ#is_codegen_graph_partition_subgraphrTs9LLLLLL 7899 5  ( 4r\ctjjjjp t jduotjjjSr)rUrLrtrw cudagraphs&_unstable_customized_partition_wrapperrKgraph_partitionrSr\rZis_using_cudagraph_partitionrYs8 %0 F 1 9 E 1 / 01r\cddlm}|jj|dr,|jj|dr t jSt jS)Nr9rli) rrrrstatically_known_ltrrUrr)r_rs rZdtype_from_sizer\"sew++ e '  / /h ? ?{{r\)r=rLcd|dkr#tjjSd|vrdSdS)z; Returns True if the device supports MKL-DNN BF16. r=rLTF)rUr7r_is_mkldnn_bf16_supportedrs rZis_mkldnn_bf16_supportedr`0;ey99;;; +  t 5r\cd|dkr#tjjSd|vrdSdS)z; Returns True if the device supports MKL-DNN FP16. r=rLTF)rUr7r_is_mkldnn_fp16_supportedr_s rZis_mkldnn_fp16_supportedrd<rar\elementsSequence[Sequence[T]]headers Sequence[T]c d|D}|D]l}t|t|ksJt|D]8\}}t||tt|||<9mg}|ddt ||Dt|t|dzzt|dz z}|d|z|D]B}|ddt ||DCd|S) NcFg|]}tt|SrS)r^rP)rXrs rZr[ztabulate_2d..Is$ + + +ac#a&&kk + + +r\|c3,K|]\}}d||dVdSrNrS)rXhrs rZrztabulate_2d..Os2HH41a,a,,,,HHHHHHr\rr9rc3,K|]\}}d||dVdSrmrS)rXrrs rZrztabulate_2d..Ts2HHtq!l!QllllHHHHHHr\r)r^rerrPrrwrrX)rergwidthsrowrrr total_widths rZ tabulate_2drsHs] + +7 + + +F443xx3w<<''''cNN 4 4DAqF1Is3q66{{33F1II 4 E LLHH3w3G3GHHHHHIIIf++Vq1S[[1_EK LL{"###JJ SXXHHs37G7GHHHHHIIII 99U  r\dict1rdict2 d1_defaultr  d2_defaultEGenerator[tuple[KeyType, ValType | None, ValType | None], None, None]c# Kt|t|z}|D];}||}||}|||n|||n|fV.patch_configs""; 0L0LMM =*6N; ' ' ' 5L00r;rrrrr   00r\c||tt|}||krtd|||||<dS)NzDOverriding: %s=%s when aot_inductor_mode.compile_standalone is True.)rrTrtrrrs rZforce_patch_configz8maybe_aoti_standalone_config..force_patch_configs_""; 0L0LMM L KKV    '3{###r\z$aot_inductor_mode.compile_standalonezaot_inductor.package_cpp_onlyTz aot_inductor.embed_kernel_binaryz#aot_inductor.emit_multi_arch_kernelz+aot_inductor.model_name_for_generated_files aoti_modelzaot_inductor.link_libtorchzaot_inductor.dynamic_linkageFz"aot_inductor.cross_target_platformz$aot_inductor.package_constants_in_sowindowszconfig.aot_inductor.package_constants_in_so is not supported for windows cross-compilation. Please use config.aot_inductor.package_constants_on_disk_format = binary_blob.)r~rrrPrrrOr )rrtaot_inductor_modecompile_standalonecopyrUrr} test_configs use_libtorch aot_inductorcross_target_platformpackage_constants_in_sor)r~rrrrrs rZmaybe_aoti_standalone_configr{s"     3 3 3 3(++. 3 $((**NR ^%DdKKK ^%GNNN Au}GXCX     I<      (   ,   >+I5QQQ*..,1 -00.3  )).E) ]   r\ consts_sizetuple[bool, bool]cdtjjr$tjjdkrt dtjjr,tjjdkrt dd}d}||fStjjdkrd}d}||fS|dkrdSd}tj }||fS) a  Decide whether we should mmap weights, and whether to store the weights with .so. If force_mmap_weights or package_constants_on_disk_format == "binary_blob" configs are set, respect the config. Returns tuple (use_external_weights, use_mmap_weights). binary_blobzconfig.aot_inductor.package_constants_on_disk_format = binary_blob and config.aot_inductor.force_mmap_weights cannot both be True.rzKwhen cross_target_platform is windows, use_mmap_weights should not be true.TFi5w)FF)rtrforce_mmap_weights package_constants_on_disk_formatrrr)ruse_mmap_weightsuse_external_weightss rZdetermine_aoti_mmap_flagsrs .   @M Q Q J   -6   4 A A]  $#%555 ;}LL# #%555m##| !+--- !1 11r\cddlm}|jj}|dSt |t st d|dkrdStjd|st ddS) zL Validates if a model name is suitable for use in code generation. rrsNTz4Invalid AOTI model name: Model name must be a stringr`z^[a-zA-Z_][a-zA-Z0-9_]*$zVInvalid AOTI model name: Model name can only contain letters, numbers, and underscores) r/rtrmodel_name_for_generated_filesrrPrErr)rt model_names rZis_valid_aoti_model_namers '&&&&&$CJt j# & &QOPPPRt 8/ < <  d    4r\r, unbacked_onlyOrderedSet[sympy.Symbol]cB|rt|St|Sr)r*r))rYrs rZget_free_symbolsrs$$Q'''Ar\dict[str, str]citjdtjdtjt ji}tjrtj d|d<|S)zA Get a base environment for running Python subprocesses. PYTHONPATHTORCH_CUSTOM_PYTHONPATHr PYTHONHOME) rkrlrrLrwrHrzrtr sysconfigget_path)envs rZpython_subprocess_envrsv   *  bjnn %rzsx'@'@    C 7%.v66L Jr\c(eZdZUdZded<ded<dS)CUDAGraphWrapperMetadataz Metadata for Customized CUDAGraphWrapper. Currently assumes there is 1 dynamo graph and will extend to multiple graphs in the future. rnum_partitionspartition_indexNrrSr\rZrr6s8r\r.ceZdZUdZded<dS)CUDAGraphWrapperNzCUDAGraphWrapperType | NonerK)rrrrKrrSr\rZrrMs#+/G//////r\rCUDAGraphWrapperTypec|t_dSr)rWrK)rKs rZ!set_customized_partition_wrappersr_s5<*222r\ tuple[list[Any], dict[str, Any]]c td|j}t|tjjjr$||j|j \}}ng|j|j }|j }| ||}t|}t|dd}t|tjjrFd|jjD}|dt%|D]}||dt)j||f\}}dd fd|D}dd  d fd fd |D}t)j||\}}||fS)NrArc*g|]}|j |jSrS) kwarg_onlyrr=s rZr[z%snode_args_kwargs..os1   S^ H   r\rOrct|tjjjo:t|tjjjtjjjf Sr)rrUrLr\rBGeneratorStateOpaqueObjectStaters rZ _is_tensor_irz(snode_args_kwargs.._is_tensor_irvsO!U_/677 _  .0B0T UA A =  r\cvg|]5}|r&tjj|dn|6S)T)replace_symbols_with_hints)rUrLr\ir_node_to_tensor)rXrUrs rZr[z%snode_args_kwargs..|sZ  =   ,,Q4,PPP r\rc0tj|||S)Nr)rUr)r_rrs rZ_tensorz"snode_args_kwargs.._tensors{4uV<<<.to_real_tensors@!U\** Hgaffhh22 r\c&g|] }|SrSrS)rXrUrs rZr[z%snode_args_kwargs..s#666q""666r\r/)rOr)rrrOr)rrrrUrLr\runflatten_argsrQ constant_argsrfill_non_provided_argsrrTrrrrr^r_pytreer"tree_unflatten) rrrrrpositional_arg_namesarg_name flat_argsflat_args_pytree_specrrrs @@@rZsnode_args_kwargsrcs  + +D$*9::**4;8JKK ff22t12  & &tV 4 4D &\\F$ t44K+uz455'  + 3 =   -[s4yy[9 ' 'H JJx & & & &'-':D&>'J'J$I$     I==== 7666I666I(4IJJLD& <r\rr=cddlm}|j}|jjr"||jjdz}|dS)Nr9rr)rr fwd_rng_stater r)rrrrr8rJ)rrdep_names rZis_nonfreeable_buffersrs`xH w|=(();<<   I  r\ template_dirr1ct||dz 5}|cdddS#1swxYwYdS)z,Load a template file and return its content.z .py.jinjaN)openread)rrr|s rZ load_templaters l//// 0 0Avvxxs 7;;c|j}t|tjjtjjfsJdt |tjsdSttj j j j tj j jj g}||vrdSttj jjg}t|tjjr||vSt#| S)zLDecide whether fallback for a node. This is only used in inductor lite mode.z6Expected OpOverload or HigherOrderOperator, but found F)rrrUrrrjr@rtfallback_by_defaultr!r7rH_assert_scalarrlift_fresh_copy higher_order triton_kernel_wrapper_functionalr)rr"skip_fallback_due_to_dynamic_shape fallback_hopss rZshould_fallback_by_defaultrs [F & (FG  OONV NNOO   %u*4 IN ) 1 IN * 2 **&333u    @AM&%*899'&&&t,, ,,r\) z-torch.ops._c10d_functional.all_reduce.defaultz.torch.ops._c10d_functional.all_reduce_.defaultz9torch.ops._c10d_functional.all_gather_into_tensor.defaultz8torch.ops._c10d_functional.reduce_scatter_tensor.defaultz4torch.ops._c10d_functional.all_to_all_single.defaultz6torch.ops._c10d_functional_autograd.all_reduce.defaultzBtorch.ops._c10d_functional_autograd.all_gather_into_tensor.defaultzAtorch.ops._c10d_functional_autograd.reduce_scatter_tensor.defaultz=torch.ops._c10d_functional_autograd.all_to_all_single.defaultz(torch.ops._c10d_functional.isend.defaultz(torch.ops._c10d_functional.irecv.defaultz0torch.ops._c10d_functional.batch_p2p_ops.defaultc|tvS)z0Check if an operation is a collective operation.)COLLECTIVE_OPS)rbs rZis_collective_oprs n $$r\cbtjr ddlm}|S#t$rgcYSwxYwgS)Nrtlx_only_cuda_options)rtr)torch._inductor.fb.tlx_templates.registryrr8rs rZrrs[    W W W W W W( (   III  s  ,,r c||zdz |z|zS)z(Round x up to the nearest multiple of y.r9rS)rYr s rZ _round_uprs UQY1  !!r\mat_sizetuple[Any, Any] scale_sizetuple[Any, ...] scale_numel mat_dtype scale_dtypeeq_fnCallable[[Any, Any], bool]tuple[Any | None, Any | None]cddlm}m}||dr|j|jfSt |dkrF||d|dr||dds*||ddr&||d|dr|j|jfS||d|dr&||dt|dds>||d|dr4||dt|ddr|j|jfS||dt|ddr4||dt|ddr|j |jfS|tj krdnd}|tj kr|tj krt|ddtt||dzddz} t|ddtt||dzddz} ||| s ||| r|j|jfS|tjkrtjjst|ddtt||dzddz} t|ddtt||dzddz} ||| s ||| r|j|jfSnjt|dd|z|dz} t||dzd|dz} ||| s ||| r|j|jfSd S) z: Core implementation for scale/swizzle inference. r)r8 SwizzleTyper9rrr~rrNN)torch.nn.functionalr8r TensorWise NO_SWIZZLEr^RowWiseruBlockWise1x128BlockWise128x128rUrrr BlockWise1x16SWIZZLE_32_4_4rrr} BlockWise1x32) rrrrrrr8r K_multiplierexpected_numel_aexpected_numel_bs rZ_infer_scale_swizzle_implr s>=<<<<<<< u[!>%{'=== :! E*Q-! - - ?%% 1 q2I2I ? E*Q- # # ?(-jmXa[(I(I ?& (>> > E*Q-! - - FjmWXa[#%>%>?? F E*Q-! - -  F jmWXa[#%>%>??  F-{/EE E 5A S 9 9 : : Huu qM78A;44@ @  H/1GG G"U%;;;11LE***{e>Q/Q/Q$Xa[#66 L8A;. 3 3Q: :  %Xa[#66 L8A;. 3 3Q: :   5. / / I55FV3W3W I,k.HH He***}  I(!c::Y x{2B77>>   )!c::Y x{2B77>>  u["233 Muu-88 M#0+2LLL M 'x{B77,FRST &|hqk'A2FFRST u["233 Iuu-88 I#0+2HHH :r\matscalect|jd|jdft|j||j|jdS)a Infer the scaling type and swizzle mode from matrix and scale tensor shapes/dtypes. This function determines how scale factors are laid out relative to the matrix: - TensorWise: Single scale for entire tensor - RowWise: One scale per row - BlockWise1x128/128x128: Block-scaled with float32 scales - BlockWise1x32: MXFP8 with float8_e8m0fnu scales (swizzled on NVIDIA) - BlockWise1x16: NVFP4 with float8_e4m3fn scales (swizzled) Args: mat: The matrix tensor (FP8 or FP4) scale: The scale factor tensor Returns: Tuple of (ScalingType, SwizzleType) or (None, None) if unrecognized rr9c||kSrrSrs rZrz%infer_scale_swizzle..`s 16r\rrrrrr)r rr?numelr)r r s rZinfer_scale_swizzler FsW( %)A, ! -%%KKMM)K!!    r\r@ transposecddlm|}|}|r|d|df}|r tjt j|dnd}d fd }tt|d kr|d|dfn |ddft|||j |j | S) z Infer the scaling type and swizzle mode for IR nodes (used during graph lowering). This is the IR-compatible version of infer_scale_swizzle, using symbolic size comparisons via V.graph.sizevars.statically_known_equals. rrr9rUrrVrOrcDjj||S)z5Compare values using symbolic equality when possible.r)rUrVrs rZ symbolic_eqz+infer_scale_swizzle_ir..symbolic_eq{sw771===r\rr )rUrrVrrOr) rkrrrHrIrJrKr r^r?r)r r r rrrr rs @rZinfer_scale_swizzle_irr ds.-----||~~H!!J.QK!-DNT)"8<Q???STK>>>>>> %/28}}/A/A(1+x{++QR UVGW$$)K    r\r-)rrrOr)rrrOr)rr)rrrrrrrOr)rrF) rrrrrrrrrOr) rr r r rrr rerOrr/)rr:rOr;)rErFrOr)rPrQrRrQrOr)rZrQr[rQrOr)r?rr[rQrOra)rErgrOrh)rnrorprorOro)rvrwrOrP)rrrOra)rrrOro)rrorOr)rrrOr)rrrOr)rr7rrrOr)rrrrrrrOr)rJ)rrPrOr )r9rJ) rrrrrrrrPrOr)rSrrrrJ)rrrrrrrrr rrrPrOr)rrrrPrOr )rrrrrOr )rUrrVrrOr)rYrr_rrOr)rYr!rOr")rr;rOr<)rOrPrOrP)rWrXrOrY)rWrcrdrerOrP)rWrXrKr<rOr}r)rrrrrOr)rrrrrOrr)r?rrOr)rrPrOr)rrprrrOr)rrrOr)rrPrOr)rrrrrOr)rUrrOr")rrrOr)rr,rOr-)rr,rOr7)rr,rOr8r.)rrrOr)rvrPrrPrOri)NNT)rorprqrrrsrrOri)rrrrrOr)rr-rrrrrOr)rrrOrrLr r+)rrrOrrw)rrrr;rrrOr:)r~)rrrOr)rrCrrrOr)rrPrOr) rrCrrrrrrrOr)rrBrrrrrOrr0)r[rQrrrOr)rrBrrCrrrOr)rr8rr8r r!rOr)r=rr>rrrCr?rr@rrArBrCrBrDrBrOr) rrCr]rrrrJrrOrr)rrCr]rTrrTrJrTr=rBr>rBrArUrrVrOr)rbrPrOrr) r]rTrrTrJrTrirrOr)r]rTrrTrJrTrOr)r]rTrrTrJrTrOr)rrPrOrP)rOr)rrCrOr)rrCrrrrBrOr)FTFN)rrCrrBrrBrrrrrrrrrOr)rrrr3rr4rOr)rrrOr)rrrr3rr4rOr)rrrr3rr4rOrP)rrrr3rr4rOr")r+rr,rrOri)r4rr5r6rOr)r>rPrOr rv)rrrOr)rQrRrOr)rrrOr)rXrrrrrrOrY)rrrOr)rrPrOr)rrPrOr)rr rOr) rrrr6rrr>rPrOr )rrrOr)rrrrrOr)rrrOr)rrHrrrOr)rrHrOr)rrrrrOr)rrPrrrrrOr) rrHrrrrrrrrrOr )rrrrrOr)rr,rOr)rr,rOr)rrrOri)rrPrOrP)rrrrOr)r(rPrr)rrrOr*) rrrr5r3rrr4r6rrrOr)rrPrOr)rrrOr)rrrrrrBrrCrrDrPrErrOr)rWrcrOr )rrrOr)rarrOr)rOr)rrrr3rr4rOrm)rQrzrOr{)rrrrrrrOr)rYrrOr)rrrrrrrOr)rQrzrrrOr)rrrOr)rrrrGrOr )rrrOrP)rrPrOr)rrrrrOr)rOr)rrPrr5rrwrOr )rrPrOr )rrrrOrP)rrrOr)rrrOr)rrrrrOr)rOr)rOr")rr.rOr})rr4rOr)rrDrOr)rKr<rOr)r_rrOr)rrPrOr)rerfrgrhrOrP) rtrrurrvr rwr rOrx)r~rrOr)rrrOr)rYr,rrrOr)rOr)rKrrOr )rrHrOr)rr=rOr)rrPrr1rOrP)rr4rOr)rOr)rYrr rrOr)rrrrrrrrrrrrrOr)r rr rrOr)r r@r r@r rrOr( __future__rrr renumrHr(rqrrr(rrJrkrrrrrHrrrrr<collections.abcrrrrrr r r r r typingrrrrrrrrrrrtyping_extensionsrrrrrrUtorch.utils._pytreerX_pytreer$torch._inductor.analysis.device_infortorch._inductor.runtime.hintsr!torch.fx.passes.regional_inductorrtorch.utils._dtype_abbrsr torch.utils._ordered_setr!r"r#rr$OPTIMUS_EXCLUDE_POST_GRAD#torch.fx.experimental._size_hintingr(r/r)r*r+r,r-r.r/r0pathlibr1r2r3r4torch._prims_commonr5torch.fxr6 torch.fx.noder7rr8rr:r r< dependenciesr=rr?r\r@rArBrCrDrE output_coderGr&rHrIr]rNrrbtorch._dynamo.device_interfacerctorch._dynamo.utilsrdtorch.autogradretorch.autograd.profiler_utilrf(torch.fx.passes.graph_transform_observerrgtorch.fx.passes.shape_proprhtorch.utils._sympy.functionsrirjrkrlrmrntorch.utils._sympy.symbolrorptorch.utils._sympy.value_rangesrqrrr`rtruntime.runtime_utilsruru _IS_WINDOWS getLoggerrrrrxrrs VarRangesrr InputTypegetenvXPU_KERNEL_FORMATGPU_KERNEL_BIN_EXTSr_rrrrrrrrrrrrrrrrrrrrrrrFunctionrrrrrr"rr9rDrOrYr`rfrmrrrrrrrrrrrrrrr r+r,r-FN_TYPEr0rLrNrVrbr|rrrrrr r rrrr!r%r) frozensetrr}r?r1r7rErZr_r`rfrhr rnrclear_on_fresh_inductor_cacheclear_inductor_cachesfresh_inductor_cacherrrrrrr3rRrZr\rfrryrrrrrrrrrrrrrrrrrr#r+r1r/r<rIrSrarhrTrprtrrrrrrrrrrrrrrr rrr!r*r3r=rBrFrLrPrUrWr[rrerxr|rrrrrrrrHEnumrrrrrrrrrrrrrr rrr!r$rr'r2r=r?rArPr\r`rdrlryr~rrrrrrrrrrrcompilerrrrrrrrrrrrrrrrrrr r"r+r-r3r=r0rFrOrTrYr\SUPPORTED_MKLDNN_DEVICESr`rdrsr}rrrrrrPartitionFnTyperrrWrrrrrrrrrr r r rSr\rZr@ s}!"""""""                                 CBBBBBBBBB $$$$$$$$$??????::::::EEEEEE000000//////;;;;;;;;222222($ <;;;;;>>>>>>>>>>>//////////CCCCCC$$$$$$""""""//////,,,,,,555555!!!!!!$$$$$$TTTTTTTTTTTTTTTT,,,,,,======== + * *  GCLL  DCCCCC000000%%%%%%222222KKKKKK00000087777777DDDDDDDD======lg% g!! C WT]] UZ' ( L3  - 4 SEEibi(I7SS   " " "     2<               22( {Q'A--+2B2B2BDX2B2BB5555 LLLL     EN   d###$#"GGGGGX#( 86.6.6.6.v#( NNNNNb     ;;;;@@@@    ++++ * * * *%A%A%A%AP++++    "////GGGG@OIIIII00''''' (%'      )))) '''' ####     $ $ $ $ IcNN WTT""" ;sAv&* +EEEEE8WQU^EEE:++++\,4)4)4)4)n_2_2_2_2H15(GGGG:,,,,^ % % % %(((( D D D DUUUU > > > >++++ ---- $)   ( =  (   +6    ''''& <!#""""     """"( +/77777v!5$"49      ( 777777tQ7777* ,,,,,,,, `'`'`'`'`'`'`'`'FLLLLLLLL^-----N---,               @@@@@ ? ? ? ? ?' ? ? ? 8 JJJJ  ))))IIII $(            # <7;uBBBBBBL5:4BGOOOOOO2BG V V V V V V OOOOQ     Q     BBB0Q     >>>>BPEEEEEPJJJJEJ&&&&& EF( . 5(5(5(5(p @@@@ RRRR:""""JHHHH&&&& :""&#=====@'''' CCCCCCCC"&&&&$&////((((V    ####* ....$FJFFFFF*    B&&&&&&&&  DDDD .6.6.6.6b  QQQQ . . . .0000    ((((####KKKK(*****$)***     F        "*.#####L((((=A Q Q Q Q Q####JJJJLLLL*=)<      <*=)< :TTTT    2               ,!!!!)))) T''''0)-10"&$( 101010101010h....2222$$$$NHHHHB    LLLL ' ' ' ' 2:0 : : : :04$$$$$<$*I*I*I*IZ:&//// '#)*$%  HG*>*D*D*F*FGGG"*Y''HHHH      & ........ 687777/9jll::::))))#####" ')   ')  -"-"-"-"-" 01-"-"-"`D)))D*)49 44442YYYY5555!!!!H""""J4    1111*            &"&!%      FRRRRj&2&2&2&2R66d###       $# 38$./@ 00000000*:)9););&====++++\     "-"-"-"-L   $%%%%       """" MMMM`B"""""""r\