o Zh#6@s:UddlmZddlZddlZddlZddlZddlZddlmZddlm Z ddl m Z m Z m Z mZmZddlmZddlmZddlmZerOdd lmZdadadagad ed <gZd ed <dadaejGdddZ ga!ded<da"da#da$da%dZddZ&eGdddZ'd[ddZ(GdddZ)iZ*ded<eGdd d Z+e+,d!gd"e+,d#gd$e+,d%gd&e+,d'gd(e+,d)gd*d\d-d.Z-d]d1d2Z.d^d5d6Z/d_d7d8Z0d`d:d;Z1d]desz%get_metric_fields..) dataclassesfieldsr'rrrrget_metric_fieldsdsr1c@s2eZdZdZd ddZdddZedd d Zd 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. rrcCs&i|_tD] }t||j|<qdSN)cached_metricsr1globals)selfmetricrrr__init__os zCachedMetricsHelper.__init__r'cCs6i}tD]}t||j|||<qtdi|S)Nr)r1r5r4r')r6Z delta_metricsr7rrr get_deltasts zCachedMetricsHelper.get_deltasdeltacCs(tD]}t|t||7<qdSr3)r1r5getattr)r:r7rrr apply_deltas{s z CachedMetricsHelper.apply_deltasNrr)rr')r:r'rr)rrrr(r8r9 staticmethodr<rrrrr2hs   r2zdict[str, MetricTable]REGISTERED_METRIC_TABLESc@s`eZdZUded<ded<dZded<dd d ZdddZdddZdddZe dddZ dS) MetricTablestr table_namer) column_namesrrnum_rows_addedrow_fn4Callable[[], dict[str, Optional[Union[str, float]]]]rrcs|jtvrdS|t|jtks"Jt|jdtt|jtks=Jt|jdtt}|gfdd|jD}tdd|DsXJ|t t t |dS)Nz v.s. csg|]}|qSrr)r,Z column_nameZrow_dictrrr.sz'MetricTable.add_row..css|]}t|tVqdSr3) isinstancerA)r,irrr sz&MetricTable.add_row..) rBenabled_metric_tableslenrCr keysr all _write_rowrlistrA)r6rEZbnrowrrGradd_rows zMetricTable.add_rowcCsd|jdS)NZ metric_table_z.csv)rB)r6rrroutput_filenameszMetricTable.output_filenamecCsX|}t|d}tj|dd}|dg|jWddS1s%wYdS)Nw lineterminatorZ model_name)rSopencsvwriterwriterowrC)r6filenamefdrZrrr write_headers  "zMetricTable.write_headerrQcCs|}|jdkrtj|s||jd7_t|D]\}}t|tr,|d}n |dur3d}n|}|||<qt |d}t j |dd}| |WddS1sWwYdS)Nrz.6farUrV) rSrDospathexistsr^ enumeraterHfloatrXrYrZr[)r6rQr\idxZorig_valnew_valr]rZrrrrOs     "zMetricTable._write_rowr+cCst||}|t|<dSr3)r@r?)r+rCtablerrrregister_tables  zMetricTable.register_tableN)rErFrr)rrAr=)rQr)rr)r+rArCr)rr) rrrrrDrRrSr^rOr>rjrrrrr@s      r@Z slow_fusion) kernel1_pathkernel1_latency kernel2_pathkernel2_latencyZfused_kernel_pathZfused_kernel_latencyZslow_down_ratioZ graph_stats)Zgraph_idZnum_nodes_before_fusionZnum_nodes_after_fusionZpersistent_red_perf) Z kernel0_pathrkrmZ kernel3_pathZkernel0_latencyrlrnZkernel3_latency size_hintsreduction_hintZ'fusion_failure_due_to_indexing_mismatch)Zpre_grad_graph_idZpost_grad_graph_idZ node1_nameZ node2_nameZnode1_debug_strZnode2_debug_strZcommon_buffer_namesZfailure_reasonkernel_metadata kernel_name kernel_pathkernel_categoryrorpZ line_of_codeZnum_loadZ num_storeZ num_for_loopZnum_atomic_addnum_argsxnumelynumelrnumelZkernel_args_num_gbkernel_module_coderAcCs8ddlm}ddlm}||}||}t|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)Z codecacher{wrapper_benchmarkr|loadinspect getsourcefn)rzr{r|modZkernelrrr_parse_kernel_fn_codes   rproper_kernel_fn_codercCs t|S)zJ Return the line of code for the kernel excluding the decorators. )rL splitlines)rrrr_parse_kernel_line_of_code&s rru Optional[str]cCs.|dkrdStd|}|sJd|dS)NZforeachzsize_hints=(\[[0-9, ]*\]),zsize_hints missing!r_researchgroup)rzrumrrr_parse_size_hints-s    rcCs.|dvrdStd|}|sJd|dS)N)Z reductionZpersistent_reductionz$reduction_hint=ReductionHint\.(\w*),z/reduction_hint not found in kernel source code!r_r)rurzrrrr_parse_reduction_hint6s    rpatterncCs ||Sr3)count)rrrrr_count_pattern@ rcCsP|d}|ds J|d}|d}||d|}|d}t|S)Nrdef (z):r_,)r startswithindexsplitrL)rZdef_lineZ start_idxZend_idxZdecl_csvcompsrrr _count_argsDs    rkernel_fn_codecCs|d}||dS)z Skip decorators. rN)r)r start_posrrr_parse_proper_kernel_fn_codeNs  rnumel_arg_name Optional[int]cCs(t|d|}|rt|dSdS)Nz = ([\d]+)r_)rrrr)rrrrrr _parse_numelVsrOptional[float]cCs$td|}|rt|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)rrrfr)rrurrrr_parse_kernel_args_num_gb^s rrsrtc shddlm}||t|t|t|tttdfdddS)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_coderqcsXtdtdtdtdttdtdtdtdS) Nztl.loadztl.storezfor z tl.atomic_addrwrxryrr)rrrrrrurZkernel_line_of_codersrtrrprorrs$z%log_kernel_metadata..N) r}rrrrrrget_metric_tablerR)rsrtrzrrrrlog_kernel_metadatass   rcCsFtD]\}}|tvr |}tj|rt||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) r?itemsrKrSrbrcrdunlinkr^)r+rir\rrrpurge_old_log_filess   rOrderedSet[str]cCs ttjSr3)enabled_metric_tables_implr rKrrrrrKrrK config_strcCsNtt}|dD]}|}|sq |tvsJd|d||q |S)NrzMetric table name z is not registered)r rArstripr?add)renabledr+rrrrs    rr+boolcCs |tvSr3)rKr*rrris_metric_table_enabledrrcCs |tvs Jd|dt|S)Nz Metric table z is not defined)r?r*rrrrsrr=)rr))rzrArrA)rrArr)rzrArurArr)rurArzrArr)rrArrArr)rrArrA)rrArrArr)rrArurArr)rsrArtrArzrArr)rr)rrArr)r+rArr)r+rArr@)< __future__rrYr/rrbrr functoolsrtypingrrrrr Ztorch._inductorr Ztorch._inductor.utilsr Ztorch.utils._ordered_setr Ztorch._inductor.schedulerr rrrrrrr r!rrr"r#r$r%r&r'r1r2r?r@rjrrrrrrrrrrrrKrrrrrrrs           =         ,