U yhca@snddlZddlZddlZddlZddlZddlZddlmZddlm Z ddl m Z ddl Z ddl Z ddlZedZddddgZd]d d Zed d Zd^ddZddZddZd_ddZddZefddZefddZefddZddZGdddZddZd d!Zd"Zd#eed$Z d%d&Z!d'd(Z"d)Z#d*d+Z$d`d,d-Z%d.d/Z&dad0d1Z'dbe d2d3d4Z(e)d5krjddl*Zej+,ej+-e.Z/e/ej+krej+0e/ddl1Z1d6Z2d7e2Z3e1j4d8e2d9Z5e5j6d:d;Z7dZ9e7j:d?e9d9Z;e;je>jdDZ9e7j:de9d9Z?e?jeHeGjKZLeIeGjMe%eLeGjQeGjd[n,eGjJd4krjeHeGjKZLeIeGjMe(eLeGjQd\dS)cN) lru_cache)Any)groupbyformat_flamegraphsegmentsmemorycompareFcCs>|d}|d}|s"|dd}|d}|d|d|S)Nlinefilename/name:)split)f full_filenameifnamefuncrH/var/www/html/venv/lib/python3.8/site-packages/torch/cuda/_memory_viz.py _frame_fmts rc Cs`ddddddddg}d d d d d dddddg }|D]}||kr0dSq0|D]}||krHdSqHdS)Nzunwind::unwindzCapturedTraceback::gatherZgather_with_cpp_startZ__libc_start_mainZPyEval_Z PyObject_Z PyFunction_z core/boxingz /Registerz /Redispatchz pythonrun.czModules/main.czObjects/call.czObjects/methodobject.czpycore_ceval.hzceval.czcpython/abstract.hFTr)r r Zomit_functionsZomit_filenamesZofrrr _frame_filters6  rcs|r t|}fdd|DS)Ncs(g|] }t|d|drt|qS)r r )rr).0rrrr >sz_frames_fmt..)reversed)framesrreverserrr _frames_fmt;sr cCsJd|kr.|dddg}|ddd}n|d|d}g}||fS)Nhistoryrr real_sizerequested_sizesize)get)brr"rrr_block_extra_legacy@s r'cCs d|krt|S|d|dfS)Nrr#)r')r&rrr _block_extraIsr(cCs|dkrdtd}tj|sXddl}td||jd|t dd|g|dd g}tj |tj tj d d }|j dk st |jdk st |j ||j |j}|j||dkst |S) Nz/tmp/z_flamegraph.plrzDownloading flamegraph.pl to: zNhttps://raw.githubusercontent.com/brendangregg/FlameGraph/master/flamegraph.plchmodz+xz --countnamebytesutf-8)stdinstdoutencoding)osgetuidpathexistsurllib.requestprintrequest urlretrieve subprocess check_callPopenPIPEr,AssertionErrorr-writeclosereadwait)Zflamegraph_linesZflamegraph_scripturllibargspresultrrrrOs*      c Csdd}|D]}d|krPt|\}}||d|dd||d|dn~d}|dD]p}|d }||7}d |kr|d }||d|dd||d|dq\||d|dd |dq\|d |} | r ||d|dd | dq dS)NcSs|sdSdt|ddS)Nz ;T)r)joinr )rrrrframes_fragmentesz&_write_blocks..frames_fragmentr!rDstate  rr"rz; r$z; )r(r<) rprefixblocksrFr&rZaccounted_for_sizehszZgapsrrr _write_blocksds  , ," rNcCsJt}|dD],}d|dd|d}t|||dq||S)Nrstream_stream;seg_addressrKioStringIOrNgetvaluesnapshotrrsegrJrrrr{s  cCs@t}|dD]"}d|d}t|||dq||S)NrrOrPrKrSrWrrrrs  csdddd}t}fdd|D}fdd|D}tdd d ||Dtd d d ||D|D]*}||krxt|d |||dqx|D]*}||krt|d|||dq||S)NcSs|d|dfS)NrR total_sizerrYrrr_seg_keyszcompare.._seg_keycSsd|dd|dS)NrOrPrQrRrr[rrr _seg_infoszcompare.._seg_infocsh|] }|qSrrrrYr\rr szcompare..csh|] }|qSrrr^r_rrr`szonly_before = cSsg|] \}}|qSrrra_rrrrszcompare..z only_after = cSsg|] \}}|qSrrrarrrrsz only_before;rKz only_after;)rTrUr4rNrV)beforeafterrr]rZ before_segsZ after_segsrYrr_rrs  cCs>dD],}t|dkr(|d|dS|d}q|ddS)N)ZKiZMiZGiZTiPiZEiZZig@z3.1fB.1fYiB)abs)numunitrrr _format_sizes   rnc@s$eZdZddZddZddZdS)BytescCs ||_dSN)value)selfrqrrr__init__szBytes.__init__cCst|j|Srp)rorq)rrrhsrrr__add__sz Bytes.__add__cCs t|jSrp)rnrq)rrrrr__repr__szBytes.__repr__N)__name__ __module__ __qualname__rsrurvrrrrrosrocCstdd|dDS)Ncss"|]}|ddkr|dVqdS)rGactive_allocatedr$Nr)rr&rrr s zcalc_active..rK)sumr[rrr calc_activesr}cCs>||}d}|dkr.||d}d|dd}t||S)Nrfrd (riz % internal))ro) free_external free_internaltotalsuffixZpctrrr _report_frees  ri@zC Legend: [a ] - a segment in the allocator ^-- a page aU of memory in the segment a-z: pages filled with a single block's content ' ': page is completely free *: page if completely full with multiple blocks 0-9: page is partially full with tensors of multiple blocks (9 == 90% full) (X% internal) - of the free memory, X% is free because we rounded the size of the allocation. c" CsTg}t}|dttdd}d}d}d}t|ddddD]`}||d7}d}d} d} g} d} |d D]b} | d d k}|rt| \}}| | |d f| |7} | | d |7} n || d 7}| | d 7} qt|| 7}||7}|| 7}|ddtd}ddt|D}ddt|D}d}t | D]\}\}}}||7}||}|t}|dtd}t t |rxdnd|d}t||D]n}t ||t}t ||dt}||||t7<||dkrdt||d||<n|||<qq4|ddkrdn d|d}d|} || | |dksJt|ddkrfd|dnd}|dtkrF|d| dt|ddt|| d|d qF|d!t|dd |d"t|d |d#t|d |rd$t|d%t|d&nd}!|d't||d |t||||ksLt|S)(aVisually reports how the allocator has filled its segments. This printout can help debug fragmentation issues since free fragments will appear as gaps in this printout. The amount of free space is reported for each segment. We distinguish between internal free memory which occurs because the allocator rounds the allocation size, and external free memory, which are the gaps between allocations in a segment. Args: data: snapshot dictionary created from _snapshot() zSummary of segments >= z in size rrcSs|dt|fS)NrZ)r})xrrrzsegsum..keyrZrKrGrzTr$cSsg|]}dqS)rHrrrcrrrrszsegsum..cSsg|]}dqS)grrrrrrsrbArHz 0123456789* rPrfz , stream_z stream_[z] z allocated, z freerIz segments: ztotal_reserved: ztotal_allocated: rz internal + z external)z total_free: )rTrUr<ro PAGE_SIZEsortedr(appendrange enumeratechrordmaxminintrEr;rlenlegendrV)"datarouttotal_reservedZtotal_allocatedrrrYZseg_free_externalZseg_free_internalZ seg_allocatedZ all_rangesZboffsetr&Zactivercallocated_sizeZnsegZoccupiedfracZ active_sizerZstart_r$Zfinish_startfinishmjserPbodyZinternal_externalrrrsegsumsr          4$ rcsRtfdd}tdD]&\}}|r"d|d||q"S)Nc sgi}i}gdfdd}fdd}d}t|dd}dD]}||d7}qTt|D]t\}}|d d kr|d |d } } |} || \} } | dkrd } | }n| | }| d| d|dt| d| | |f|| <|| 7}qn|d dkrX|d |d } } || | ddf\}}}d|dt| dqn|d dkr|d |d } } || 8}|| | ddf\}}}d|dt| d||kr|||=qn|d dkr0|d |d } } |}|d| dt| d|| | f||| <qn|d dkr|d |d } } || | }d|dt| d||kr|||=qn|d d kr|d } |d!}d"t| d#t|d$qn|dqnd%t|dS)&NrcsJr Sdd}}d7ttd||dkrBdn|S)Nrrrbrrf)poprr)rr) free_names next_namerr_name$s z$trace..format.._namecsD]*\}}}||kr|||kr||fSqtdD]>\}}|d}|d}||kr<|||kr.format..find_segmentz entries rrZactionallocrr$ZMEMz = rrz] free_requestedzdel z # rIfree_completedz# free completed for rHZ segment_allocz = cudaMalloc(z, z) Z segment_freez cudaFree(z) # ZoomZ device_freezraise OutOfMemoryError # z requested, z free in CUDA z TOTAL MEM: )r<rrror%r)entriesZsegment_addr_to_nameZallocation_addr_to_namerrcountrrYrrr$nZseg_nameZseg_addroffsetr rcfreerr)rrrrformatsn   &        "ztrace..format device_traceszDevice z ---------------- )rTrUrr<rV)rrrdrrrtracesO ra  cCst|dk rtjdtddt|}|ddt|d7}t|d}t d|dg}t dt | d |S) Nz;device argument is deprecated, plots now contain all device) stacklevelr+zsnapshot.pickle)r base64z $VIZ_KINDz $SNAPSHOT) warningswarn FutureWarningpickledumpsrr b64encodedecodejson_memory_viz_templatereplacerepr)rZviz_kinddevicebufferZencoded_bufferZ json_formatrrr _format_vizs rcCst||s dnd|S)aGenerate a visualization over time of the memory usage recorded by the trace as an html file. Args: data: Memory snapshot as generated from torch.cuda.memory._snapshot() device (torch.device, optional): Generate the trace for this device, needed if multiple devices have allocations. plot_segments (bool, optional): Plots memory returned from cudaMalloc, rather than individual allocations. Defaults to False. Returns: str: HTML of visualization Active Memory TimelinezActive Cached Memory Timeliner)rr plot_segmentsrrr trace_plots rc sddl}ddlm}m}ddlm}|ijjD]f}|j |j kr8|j }g}|rz|j |j |j fkrr|||j }qR||j}|r8|jjdkr8||<q8|jddtdDddtdDdfd d dfd d } fdd} i} jD]\} } \}}}t||s.q| |jkrP| |||| ||f<n| |jkrx| | ||f|jnh| |jkr| | ||f|j| |||d| ||df<n$| |jkr| |||dd| ||f<qfdd| D}tt|t !ddD]\}}d|}|d}|D]P\}}}}||kr^|d||dd|d|d||d||}q0||dkr|d|d|ddqdddDd<dD]<}|d|d8<|ds|d|dddqΈS)Nr)Action TensorKey) _EventTypecSsg|]}gqSrrrrrrrsz(_profile_to_snapshot..rcSsg|]}|dddgdqS)Nr)rrRrZrPrKr)rrrrrrs)rrcs|jdkr|jSSdS)Ncuda)typeindexr) device_countrr to_devices z'_profile_to_snapshot..to_deviceTc s|j}|jj}d|}|ddks6|d|kr>||d<t|d|||d<j||}|dk rt|jnd}|d}dd|D}d||d ||d } |rd || | S) NrrRrZunknownrcSsg|]}dd|jdqS)noner)r r r )r )rrBrrrrsz:_profile_to_snapshot..allocate..rr)rrr$rPrcategoryr) rZstorageZptrrZ _categoriesr%r lowerr) r$ tensor_keyversion during_tracerrrYrstackr)allocation_stacksmemory_profilerXrrrallocates   z&_profile_to_snapshot..allocatec s8dD].}d|||d|dd|ddqdS)N)rrrrr$rr)rrr$rPr)r)rrr)rXrrrsz"_profile_to_snapshot..freeF)rcs2g|]*\\}}}|j|d|d|dfqS)rr$rr)rrrevent)rrrrs rrrRrKZinactive)r$rGrz)r$rGr#rrZcSsg|]}|dr|qS)rKrr^rrrrs)T)"torchZtorch.profiler._memory_profilerrrZtorch._C._profilerrZ_memory_profileZ_op_treeZ sorted_nodestagZ AllocationparentZPyCallZPyCCallrZfrom_allocationZ extra_fieldsZ alloc_sizerrrZtimeline isinstanceZCREATEZDESTROYrrZINCREMENT_VERSIONZ PREEXISTINGitemsrroperator itemgetter)profilerrrrrrZpython_parentsrrrZ kv_to_elemtimerrrr$Z blocks_at_endrrKrYZ last_addrrcrrr)rrrrXrr_profile_to_snapshotsr                      rcCst|}t|d|S)a}Generate a visualization over time of the memory usage recorded by kineto memory profiling as an html file. Args: profile: profile as generated by `torch.profiler.profile(profile_memory=True)` device (torch.device, optional): Generate the trace for this device, needed if multiple devices have allocations. Returns: str: HTML of visualization r)rr)rrrXrrr profile_plots r)rcCs t|d|S)NzAllocator State Historyr)rrrrr segment_plotsr__main__ztorch.cuda.memory._snapshot()zpickled memory statistics from z#Visualize memory dumps produced by ) descriptionr)destcCs|jddddddS)N-o--outputz output.svgz$flamegraph svg (default: output.svg)defaulthelp) add_argument)rBrrr_output'srzmPrints overall allocation statistics and a visualization of how the allocators segments are currently filled.statsinput)rz`Prints buffer of the most recent allocation events embedded in the snapshot in a Pythonic style.zaGenerate a flamegraph that visualizes what memory is stored in each allocator segment (aka block)zNGenerate a flamegraph the program locations contributing to CUDA memory usage.zGenerate a flamegraph that shows segments (aka blocks) that have been added or removed between two different memorys snapshots.rdre))rz]Generate a visualization over time of the memory usage recorded by the trace as an html file.)rzfVisualize how allocations are packed into allocator segments at each point in a trace as an html file.z]visualize trace from this device (default: chooses the only device with trace info or errors)z-dz--device)rrrz4path to save the visualization(default: output.html)rrz output.htmlrz?visualize change to segments rather than individual allocationsz-sz --segments store_true)rrcCs>|dkrtjj}n t|d}t|}t|tr:|gd}|S)N-rb)rZtraces)sysr,ropenrloadrlist)r rrrrr_readUs     r c Cs$t|d}||W5QRXdS)Nw)rr<)r rrrrr_write_s r )rrr)F)FF)N)NF)N)N)Rrrr/rTr7r functoolsrtypingr itertoolsrrrrcache__all__rrr r'r(rrNrrrrnror}rrrrrrrrrrrrwos.pathr1realpathdirname__file__Zthedirremoveargparsefn_nameZpickledArgumentParserparseradd_subparsersZ subparsersrr add_parserZstats_arZtrace_aZ segments_aZmemory_aZ compare_aZplotscmdZ trace_plot_arr parse_argsrAr r rrroutputr4rdrerrrrrs              MY i