Wj\9UddlmZddlZddlZddlZddlZddlZddlmZddlm Z ddl m Z m Z m Z ddlmZddlmZddlmZe rdd lmZdd lmZdd lmZdadadagad ed <gZded<da da!ejGddZ"ga#ded<da$da%da&da'ded<da(da)dSdZ*eGddZ+dTdZ,GddZ-iZ.ded <eGd!d"Z/e/0d#gd$e/0d%gd&e/0d'gd(e/0d)gd*e/0d+gd,dUd/Z1dVd1Z2dWd4Z3dXd5Z4dYd7Z5dVd8Z6dZd:Z7d[d=Z8d\d?Z9d]dBZ:dSdCZ;d^dEZdadKZ?e/0dLgdMdbdRZ@dS)c) annotationsN) dataclass) lru_cache)Optional TYPE_CHECKINGUnion)config)get_benchmark_name) OrderedSet)Callable)Config)BaseSchedulerNodez#list[tuple[BaseSchedulerNode, int]]nodes_num_elemz%list[tuple[BaseSchedulerNode, float]] node_runtimesc(eZdZUded<dZded<dS)CppOuterLoopFusedCountintinner_kernel_numberrlocal_buffer_numberN)__name__ __module__ __qualname____annotations__r\/home/longshao/multi-rider-rag/.venv/lib/python3.11/site-packages/torch/_inductor/metrics.pyrr+s0       rrzlist[CppOuterLoopFusedCount]!cpp_outer_loop_fused_inner_countsrnum_auto_chunkingreturnNonecdadadatt dadatda da da da da dadS)Nr)generated_kernel_countgenerated_cpp_vec_kernel_countnum_bytes_accessedrclearrir_nodes_pre_fusioncpp_to_dtype_countrnum_comprehensive_padding)num_matches_for_scatter_upon_const_tensornum_loop_reorderingparallel_reduction_countcodegen_mix_order_reductionrrrrresetr-As%&"%++--- !01- "#rcPeZdZUdZded<ded<ded<ded<ded<ded<d S) CachedMetricsDeltasz] The subset of metrics we want update across cache hits, e.g., the FxGraphCache. rr"r#r&r'r$r)N)rrr__doc__rrrrr/r/_sf  ''''222222rr/ list[str]cHdtjtDS)Ncg|] }|j Srname).0fields r z%get_metric_fields..os L L L5EJ L L Lr) dataclassesfieldsr/rrrget_metric_fieldsr;ns" L LK$67J$K$K L L LLrc:eZdZdZd dZd dZed dZd S) CachedMetricsHelperz A helper class to help calculate and apply counter deltas for those metrics we want to save with cache entries (e.g., FxGraphCache) and apply on a cache hit. rr cni|_tD]}t||j|<dSN)cached_metricsr;globals)selfmetrics r__init__zCachedMetricsHelper.__init__ysB ')) < .sQQQ h{+QQQrc 3jK|].}t|tttdfV/dSr?) isinstancerQfloattype)r6is r z&MetricTable.add_row..s9HHq:a#ud4jj!9::HHHHHHr) rRenabled_metric_tableslenrSr keysr all _write_row)rBrUbnrowrZs @radd_rowzMetricTable.add_rowsD ?"7"9"9 9 9 F6884$%%X6664$%% < > ? ? ? @ @ @ @ @ @ @ @ @ @ @ @ @ @ @ @ @ @s5A''A+.A+rglist[str | float | None]c|}|jdkr3tj|s||xjdz c_t |D]+\}}t|tr|d}n|d}n|}|||<,t|d5}tj |d}| |ddddS#1swxYwYdS)Nrz.6farnro) rkrTospathexistsrx enumerater\r]rrrsrtru)rBrgrvidxorig_valnew_valrwrts rrezMetricTable._write_rowsF''))  ! # #BGNN8,D,D #       q &s^^  MC(E** #%++!"CHH (C  !BZ4888F OOC  ! ! ! ! ! ! ! ! ! ! ! ! ! ! ! ! ! !s.,C''C+.C+r5c:t||}|t|<dSr?)rPrN)r5rStables rregister_tablezMetricTable.register_tables"D,//). &&&rN)rUrVrr )rrQrL)rgryrr )r5rQrSr1rr ) rrrrrTrhrkrxrerMrrrrrPrPsOOON(5555@@@@ !!!!(///\///rrP slow_fusion) kernel1_pathkernel1_latency kernel2_pathkernel2_latencyfused_kernel_pathfused_kernel_latencyslow_down_ratio graph_stats)graph_idnum_nodes_before_fusionnum_nodes_after_fusionpersistent_red_perf) kernel0_pathrr kernel3_pathkernel0_latencyrrkernel3_latency size_hintsreduction_hint'fusion_failure_due_to_indexing_mismatch)pre_grad_graph_idpost_grad_graph_id node1_name node2_namenode1_debug_strnode2_debug_strcommon_buffer_namesfailure_reasonkernel_metadata kernel_name kernel_pathkernel_categoryrr line_of_codenum_load num_store num_for_loopnum_atomic_addnum_argsxnumelynumelrnumelkernel_args_num_gbkernel_module_coderQcddlm}ddlm}||}||}t j|jjS)z The kernel_module_code is the python module that contains kernel function code. kernel function is the proper triton kernel function annotated with @triton.jit r{) PyCodeCache)get_triton_kernel) codecacherwrapper_benchmarkrloadinspect getsourcefn)rrrmodkernels r_parse_kernel_fn_coder sc '&&&&&444444   - . .C  s # #F  VY\ * **rproper_kernel_fn_codecDt|S)zJ Return the line of code for the kernel excluding the decorators. )rb splitlines)rs r_parse_kernel_line_of_coder0s $//11 2 22rr Optional[str]c~|dkrdStjd|}|s Jd|dS)Nforeachzsize_hints=(\[[0-9, ]*\]),zsize_hints missing!r{researchgroup)rrms r_parse_size_hintsr7sI)##t /1CDDA #####1 771::rcz|dvrdStjd|}|s Jd|dS)N) reductionpersistent_reductionz$reduction_hint=ReductionHint\.(\w*),z/reduction_hint not found in kernel source code!r{r)rrrs r_parse_reduction_hintr@sLCCCt 9;MNNA ?????1 771::rpatternc,||Sr?)count)rrs r_count_patternrJs & &w / //rc|d}|dsJ|d}|d}||dz|}|d}t |S)Nrdef (z):r{,)r startswithindexsplitrb)rdef_line start_idxend_idxdecl_csvcompss r _count_argsrNs$//11!4H   v & &&& &s##InnT""G A /0H NN3  E u::rkernel_fn_codec@|d}||dS)z Skip decorators. rN)r)r start_poss r_parse_proper_kernel_fn_coderXs&$$V,,I )** %%rnumel_arg_name Optional[int]c~tj|d|}|r"t|dSdS)Nz = ([\d]+)r{)rrrr)rrrs r _parse_numelr`s@ ^0002GHHA1771::trOptional[float]cztjd|}|r"t|dS dS)z inductor meta looks like: inductor_meta={... 'mutated_arg_names': [], 'no_x_dim': False, 'kernel_num_gb': 2.0}, z.kernel_num_gb.:\s*([0-9.]+)r{N)rrr]r)rrrs r_parse_kernel_args_num_gbrhsB 1>BBA QWWQZZ    trrrc * ddlm}||t|t| t |t t td fddS)z An utility to log kernel metadata. We may parse metadata from kernel source code here. It's fine to parse the generated kernel code here since the logging is disabled by default. It would hurt compilation time. r{)"get_kernel_category_by_source_coderc"tdtdtdtdttdtdtdtdS) Nztl.loadztl.storezfor z tl.atomic_addrrrr)rrrr)rrkernel_line_of_coderrrrrsrz%log_kernel_metadata..s&&.$,/&'>N8HH55JKK&''//           rctD]p\}}|tvr[|}tj|rt j||qdS)z Purge the old log file at the beginning when the benchmark script runs. Should do it in the parent process rather than the child processes running each individual model. N) rNitemsrarkr~rrunlinkrx)r5rrvs rpurge_old_log_filesrs 05577!! e (** * *,,..Hw~~h'' $ (###    !!rOrderedSet[str]c4ttjSr?)enabled_metric_tables_implr rarrrraras %f&B C CCr config_strct}|dD]E}|}|s|tvsJd|d||F|S)NrzMetric table name z is not registered)r rstriprNadd)renabledr5s rrrs)||G  %%zz||  //// 9 9 9 90//  D Nrr5boolc"|tvSr?)rar4s ris_metric_table_enabledrs (** **rcJ|tvsJd|dt|S)Nz Metric table z is not defined)rNr4s rrrs3 + + + +-RT-R-R-R + + + #D ))rkernel_autotunerr triton_config latency_msr r latencyr]c^tdfddS)Nrc,tdS)Nr)rQ)r rrrsrrz,log_kernel_autotune_result..s!&& [[!   r)rrh)rrr rs````rlog_kernel_autotune_resultr sV&''//       rrL)rr1)rrQrrQ)rrQrr)rrQrrQrr)rrQrrQrr)rrQrrQrr)rrQrrQ)rrQrrQrr)rrQrrQrr)rrQrrQrrQrr )rr)rrQrr)r5rQrr)r5rQrrP) rrQrrQr r rr]rr )A __future__rrsr9rr~rr functoolsrtypingrrrtorch._inductorr torch._inductor.utilsr torch.utils._ordered_setr collections.abcr %torch._inductor.runtime.triton_compatr torch._inductor.schedulerrr"r#r$rrrr&r'rrr(r)r*rr+r,r-r/r;r=rNrPrrrrrrrrrrrrrarrrr rrrrs """""""  !!!!!!1111111111""""""444444//////<((((((<<<<<<;;;;;;!"   8: 9999 !!!!!!!! CE!DDDD,-)<  3 3 3 3 3 3 3  3MMMM888888882465555 :/:/:/:/:/:/:/ :/z         " -       6 + + + + 33330000&&&&*))))X ! ! ! !DDDD      ++++****        r