j 5UddlmZddlZddlZddlZddlZddlZddlmZddlm Z ddl m Z m Z m Z mZmZddlmZddlmZerddlmZdadadagad ed <gZd ed <dadaejGd dZgaded<da da!da"dZ#eGddZ$dZ%GddZ&iZ'ded<eGddZ(e()dgde()dgde()dgd e()d!gd"e()d#gd$d%Z*d&Z+d'Z,d(Z-d)Z.d*Z/d+Z0d,Z1d-Z2d.Z3d/Z4d8d2Z5e d9d5Z6d6Z7d7Z8dS):) annotationsN) dataclass) lru_cache)DictListSetTuple TYPE_CHECKING)config)get_benchmark_name)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__rHC:\Projects\Historical\venv\Lib\site-packages\torch/_inductor/metrics.pyrr(s0       rrzList[CppOuterLoopFusedCount]!cpp_outer_loop_fused_inner_countscdadadatt dadatda da da dS)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_reorderingrrrresetr'8sp%&"%++--- !01-rcPeZdZUdZded<ded<ded<ded<ded<ded<d S) CachedMetricsDeltasz] The subset of metrics we want update across cache hits, e.g., the FxGraphCache. rrrr"r#r r%N)rrr__doc__rrrrr)r)Psf  ''''222222rr)cHdtjtDS)Ncg|] }|j Srname).0fields r z%get_metric_fields..`s L L L5EJ L L Lr) dataclassesfieldsr)rrrget_metric_fieldsr4_s" 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. returnNonecni|_tD]}t||j|<dSN)cached_metricsr4globals)selfmetrics r__init__zCachedMetricsHelper.__init__jsB ')) < s r get_deltaszCachedMetricsHelper.get_deltasosU ')) T TF$+IIf$58KF8S$SM& ! !"33]333rdeltac|tD],}t|xxt||z cc<-dSr:)r4r<getattr)rCr>s r apply_deltasz CachedMetricsHelper.apply_deltasvsN')) 8 8F IIf   !7!7 7     8 8rN)r7r8)r7r))rCr))rrrr*r?rB staticmethodrFrrrr6r6csf <<<< 4444888\888rr6zDict[str, MetricTable]REGISTERED_METRIC_TABLESc`eZdZUded<ded<dZded<dZd Zd Zd Ze d Z d S) MetricTablestr table_namez List[str] column_namesrrnum_rows_addedcR|jtvrdS|t|jtks.Jt|jdtt |jt ks@Jt |jdt t g}|fd|jDz }||dS)Nz v.s. c g|] }| Srr)r/ column_namerow_dicts rr1z'MetricTable.add_row..sKKK+%KKKr)rLenabled_metric_tableslenrMsetkeysr _write_row)r=row_fnrowrRs @radd_rowzMetricTable.add_rowsC ?"7"9"9 9 9 F6884$%% * *    $#$$ ; ;CMM ; ;   4$%% MMOO* *    $#$$ B BC ,@,@ B B      KKKK9JKKKK rcd|jdS)N metric_table_z.csv)rL)r=s routput_filenamezMetricTable.output_filenames4t4444rc|}t|d5}tj|d}|dg|jzddddS#1swxYwYdS)Nw lineterminator model_name)r]opencsvwriterwriterowrM)r=filenamefdrfs r write_headerzMetricTable.write_headers'')) (C  @BZ4888F OO\NT->> ? ? ? @ @ @ @ @ @ @ @ @ @ @ @ @ @ @ @ @ @s5A''A+.A+c|}|jdkr3tj|s||xjdz c_t |D]+\}}t|tr|d}n|d}n|}|||<,t|d5}tj |d}| |ddddS#1swxYwYdS)Nrz.6far`ra) r]rNospathexistsrj enumerate isinstancefloatrdrerfrg)r=rYrhidxorig_valnew_valrirfs rrWzMetricTable._write_rowsF''))  ! # #BGNN8,D,D #       q &s^^  MC(E** #%++!"CHH (C  !BZ4888F OOC  ! ! ! ! ! ! ! ! ! ! ! ! ! ! ! ! ! !s.,C''C+.C+c:t||}|t|<dSr:)rJrH)r.rMtables rregister_tablezMetricTable.register_tables"D,//). &&&rN) rrrrrNrZr]rjrWrGrzrrrrJrJsOOON$555@@@ !!!(//\///rrJ 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_pathr|r~ kernel3_pathkernel0_latencyr}rkernel3_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_gbcddlm}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 rl) PyCodeCache)get_triton_kernel) codecacherwrapper_benchmarkrloadinspect getsourcefn)kernel_module_coderrmodkernels r_parse_kernel_fn_codersc '&&&&&444444   - . .C  s # #F  VY\ * **rcDt|S)zJ Return the line of code for the kernel excluding the decorators. )rT splitlines)proper_kernel_fn_codes r_parse_kernel_line_of_coders $//11 2 22rc~|dkrdStjd|}|s Jd|dS)Nforeachzsize_hints=(\[[0-9, ]*\]),zsize_hints missing!rlresearchgroup)rrms r_parse_size_hintsr&sI)##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!rlr)rrrs r_parse_reduction_hintr/sJCCCt 9;MNNA ?????1 771::rc,||Sr:)count)rpatterns r_count_patternr7s & &w / //rc|d}|dsJ|d}|d}||dz|}|d}t |S)Nrdef (z):rl,)r startswithindexsplitrT)rdef_line start_idxend_idxdecl_csvcompss r _count_argsr;s$//11!4H   v & &&& &s##InnT""G A /0H NN3  E u::rc@|d}||dS)z Skip decorators. rN)r)kernel_fn_code start_poss r_parse_proper_kernel_fn_coderEs&$$V,,I )** %%rc~tj|d|}|r"t|dSdS)Nz = ([\d]+)rl)rrrr)rnumel_arg_namers r _parse_numelrMs@ ^0002GHHA1771::trcztjd|}|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.]+)rlN)rrrtr)rrrs r_parse_kernel_args_num_gbrUsB 1>BBA QWWQZZ    trc * 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. rl)"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) rHitemsrSr]rorprqunlinkrj)r.ryrhs rpurge_old_log_filesrs 05577!! e (** * *,,..Hw~~h'' $ (###    !!rr7Set[str]c4ttjSr:)enabled_metric_tables_implr rSrrrrSrSs %f&B C CCr config_strrKct}|dD]E}|}|s|tvsJd|d||F|S)NrzMetric table name z is not registered)rUrstriprHadd)renabledr.s rrrseeG  %%zz||   , , , , 8 8 8 8 - , , D Nrc"|tvSr:)rSr-s ris_metric_table_enabledrs (** **rcJ|tvsJd|dt|S)Nz Metric table z is not defined)rHr-s rrrs3 + + + +-RT-R-R-R + + + #D ))r)r7r)rrKr7r)9 __future__rrer2rrorr functoolsrtypingrrrr r torch._inductorr torch._inductor.utilsr torch._inductor.schedulerr rrr rrrr"r#rrr$r%r&r'r)r4r6rHrJrzrrrrrrrrrrrrSrrrrrrrs5"""""""  !!!!!!88888888888888""""""444444<;;;;;;!"   8: 9999 !!!!!!!! CE!DDDD,-)0  3 3 3 3 3 3 3  3MMM888888882465555 8/8/8/8/8/8/8/ 8/v         " -       6 + + + 333000&&&&'''T ! ! !DDDD      +++*****r