g^(hc ddlZddlZddlZddlZddlZddlZddlZddlZddlZddl m Z ddl m Z ddl mZe dZgdZdLdZedZdMdZd Zd ZdNd Zd Zefd ZefdZefdZdZGddZdZdZdZdeedZ dZ!dZ"dZ#dZ$dOdZ%dZ&dNdZ'dNdefd Z(e)d!k(rddl*ZejVjYejVj[e.Z/e/ejVvrejVjae/ddl1Z1d"Z2d#e2Z3e1jhd$e2%Z5e5jmd&'Z7d(Z8d)Z9e7jud*e9%Z;e;jyd+e3,d-Z9e7jud.e9%Z=e=jyd+e3,d/Z9e7jud0e9%Z>e>jyd+e3,e8e>d1Z9e7jud2e9%Z?e?jyd+e3,e8e?d3Z9e7jud4e9%Z@e@jyd5e3,e@jyd6e3,e8e@d7ZAeAD]w\ZBZ9e7jueBe9%ZCeCjyd+e3,d8ZDeCjyd9d:eEdeD;dd?eD@eBdAk(sadBZDeCjydCdDdEeDFye5jZGdGZHdHZIeGjd0k(r,eHeGjZLeIeGjeeLyeGjd2k(r,eHeGjZLeIeGjeeLyeGjd*k(r!eHeGjZLeNe!eLyeGjd.k(r!eHeGjZLeNe"eLyeGjd4k(r?eHeGjZOeHeGjZPeIeGjeeOePyeGjdAk(rCeHeGjZLeIeGje%eLeGjeGj0IyeGjdJk(r8eHeGjZLeIeGje(eLeGjKyyy)PN) lru_cache)groupby)Any)format_flamegraphsegmentsmemorycomparec`|d}|d}|s|jdd}|d}|d|d|S)Nlinefilename/name:)split)f full_filenameifnamefuncs T/var/www/skyplay_api_hub/venv/lib/python3.12/site-packages/torch/cuda/_memory_viz.py _frame_fmtrsJ & A jME  C $ V9DWAaS$  cHgd}gd}|D]}||vsy|D]}||vsyy)N)zunwind::unwindzCapturedTraceback::gathergather_with_cpp_start__libc_start_mainPyEval_ PyObject_ PyFunction_) z core/boxingz /Registerz /Redispatchz pythonrun.czModules/main.czObjects/call.czObjects/methodobject.czpycore_ceval.hzceval.czcpython/abstract.hFT)rr omit_functionsomit_filenamesofs r _frame_filterr%sK N N : > rc~|r t|}|Dcgc] }t|d|dr t||"c}Scc}w)Nrr )reversedr%r)framesrreversers r _frames_fmtr*?sJ&!  6AjM 2 1m$  s%:cd|vr'|ddjdg}|ddd}||fS|jd|d}g}||fS)Nhistoryrr( real_sizerequested_sizesize)get)br(r-s r_block_extra_legacyr2IsiA~9a$$Xr2iLOK0  9 EE*AfI6  9 rc4d|vr t|S|d|dfS)Nr(r.)r2)r1s r _block_extrar4Ss*q"1%% X;*+ ++rc|dtjd}tjj|sFddl}t d||j jd|tjdd|g|dd g}tj|tjtjd }|jJ|jJ|jj||jj|jj!}|jj|j#|j#dk(sJ|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_callPopenPIPEr9r:writeclosereadwait)flamegraph_linesflamegraph_scripturllibargspresults rrrZs #BIIK=? 77>>+ , ./@.ABC"" \   w.?@A }g 6D JOOJOOg A 77   88  GGMM"#GGMMO XX]]_FHHNNFFH 668q== Mrc d}|D]}d|vr5t|\}}|j|d|dd||d|dncd}|dD]Y}|d}||z }d |vr,|d }|j|d|dd||d|d=|j|d|dd |d[|d |z } | s|j|d|dd | dy) Nc@|sydjt|dS)Nz ;T)r))joinr*)r(s rframes_fragmentz&_write_blocks..frames_fragmentvs!xx FD9::rr,rTstate  rr-r(z; r/z; )r4rH) rprefixblocksrVr1r(accounted_for_sizehszgapss r _write_blocksr`us/; > A )5a &F& GG(!AgJ"RHI Jy--  GGvha' |8D6< =%>rctj}|dD] }d|dd|d}t|||d"||jS)Nrstream_stream;seg_addressr[ioStringIOr`getvaluesnapshotrrsegrZs rrrs_ A #03x=/s9~.>?aX/0 QZZ\ **rctj}|dD]}d|d}t|||d||jS)Nrrbrcr[rfrjs rrrsT A #03x=/*aX/0 QZZ\ **rc d}d}tj}|Dchc] }|| }}|Dchc] }|| }}td||z D cgc]\} } |  c} } td||z D cgc]\} } |  c} } |D]&}|||vst|d|||d(|D]&}|||vst|d|||d(||j Scc}wcc}wcc} } wcc} } w)Nc|d|dfS)Nre total_sizer!rls r_seg_keyzcompare.._seg_keysIL 122rcd|dd|dS)Nrbrcrdrer!rqs r _seg_infozcompare.._seg_infosXuS^,<==rzonly_before = z only_after = z only_before;r[z only_after;)rgrhrAr`ri) beforeafterrrrrtrrl before_segs after_segsa_s rr r s$3> A,23S8C=3K3+01C(3-1J1 N;+CE$!QAEF GH M*{*BDA1DE FGM C= * !|IcN+;*:;S] KL QZZ\ **41EDsC9C> D 7 D cVdD]}t|dkr |d|dcS|dz}!|ddS)N)KiMiGiTiPiEiZig@z3.1fB.1fYiB)abs)numunits r _format_sizersI> s8f $ZvQ' ' v #Yc?rceZdZdZdZdZy)Bytesc||_yN)value)selfrs r__init__zBytes.__init__s  rc2t|j|zSr)rr)rrhss r__add__z Bytes.__add__sTZZ#%&&rc,t|jSr)rr)rs r__repr__zBytes.__repr__sDJJ''rN)__name__ __module__ __qualname__rrrr!rrrrs'(rrc,td|dDS)Nc38K|]}|ddk(s |dyw)rWactive_allocatedr/Nr!).0r1s r zcalc_active..s TQ1W:AS3SqyTs  r[)sumrqs r calc_activers T#h-T TTrcV||z}d}|dk7r||z dz}d|dd}t||S)Nr|rdz (rz % internal))r) free_external free_internaltotalsuffixpcts r _report_freersJ M )E F zu$+c#Yk*El^F8 $$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 >tj}|jdttdd}d}d}d}t |ddD]}||dz }d}d}d} g} d} |dD]O} | d d k(} | r3t | \}}| j| |d f| |z } || d |z z }n|| d z }| | d z } Q|| z }||z }||z }|dd z tzd z}t|Dcgc]}d}}t|Dcgc]}d}}d}t| D]\}\}}} ||z }||z}|tz}|d z tzd z}tt| rdnd|dzz}t||D]g}t||tz}t||d ztz}||xx||z tz z cc<||dk7rdt||dz||<c|||<i|ddk(rdnd|d}dj|}||z| z|dk(sJ|ddk7rd|dnd}|dtk\s|jd|dt|ddt!||d|d |jdt#|dd|jdt|d|jd t|d|jd!t!||d|jt$||z|z|k(sJ|j'Scc}wcc}w)"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 rrc"|dt|fS)Nrp)r)xs rzsegsum..s<+a.(Irkeyrpr[rWrTr/rXgryAz 0123456789* rcr|z , stream_z stream_[z] z allocated, z freerYz segments: ztotal_reserved: ztotal_allocated: z total_free: )rgrhrHr PAGE_SIZEsortedr4appendrange enumeratechrordmaxminintrUrlenlegendri) dataouttotal_reservedtotal_allocatedrrrlseg_free_externalseg_free_internal seg_allocated all_rangesboffsetr1activerzallocated_sizensegoccupiedfrac active_sizerstart_r/finish_startfinishmjsercbodys rsegsumrs ++-CII'i(8'9DENOMM ZI6 #l++  X !AwZ#55F$0O!>!!7ND"AB/ !QvY%??!!QvY.! qy G ! =(** ** L!A%)3a7!&t-AC--"4[))) )2:)> $ %A%f 4 KtmGi'Eki/!3FCv3/1r6:;A5&) $I .!a%9!45QAEY..A;#%"/DGbL0A"BHQK"#HQK $ $8})3x=//Jwwx   1 1M ASEV V V/28}/A8CM?+r |  ) IID6E#l"345\ 13DEFeF8SUW g6nII 3tJ/0145II ~!6 7r:;II!%"8!9<=II \-GHKLIIf = (? :n LL L <<>E.)s + L Lctjfd}tdD]%\}}|s jd|d||'j S)Nc Rgi}i}gdfd}fd}d}jt|dd}dD] }||dz } t|D].\}}|ddk(r\|d |d } } |} || \} } | d } | }n| | z }j| d | d |dt| d| | |f|| <|| z }k|ddk(rE|d |d } } |j | | ddf\}}}jd|dt| d|ddk(re|d |d } } || z}|j | | ddf\}}}jd|dt| d||vsj |||=%|ddk(rO|d |d } } |}j|d| dt| dj || | f||| <||ddk(rY|d |d } } |j | | }jd|dt| d||vsljj |||=|ddk(r6|d } |d}jd t| d!t|d"j|d1jd#t|y)$NrcrjSdzdz}}dz ttd|z|dk(rdS|S)Nrrryrr|)poprr)rr free_names next_names r_namez$trace..format.._nameFs\!~~''?INqA NI#c(Q,'(qAv(=> >1(=> >rcD]\}}}||k\s |||zks||fcStdD]&\}}|d}|d}||k\s|||zksd||fcSy)Nrrerseg_)NN)r)addrrsaddrr/rrlrsegment_intervalss r find_segmentz+trace..format..find_segmentNs%6 '!eT5=TEDL%8;& '$D$45 -3I+,5=TEDL%8!!:u,,  - rz entries rrpactionallocrr/MEMz = rrz] free_requestedzdel z # rYfree_completedz# free completed for rX segment_allocz = cudaMalloc(z, z) segment_freez cudaFree(z) # oom device_freezraise OutOfMemoryError # z requested, z free in CUDA z TOTAL MEM: )rHrrrr0r)entriessegment_addr_to_nameallocation_addr_to_namerrcountrrlrrr/nseg_nameseg_addroffsetrrzfreerrrrrs @@@rformatztrace..format>sC"$!"$   ?  S\N*-. # 0C c,/ /N 0"'*- $HE1{g%vY& dG%1$%7"(#$H!F!H_F QCs8*AfXQuT{m3GH12D%0@'- 8 00vY& d488dD?QR a Dc%+b9:8 00vY& d 488dD?QR a 1$qt RHI22%%d+/58/vY& dw TF.bt SIJ!(($d);<-1$T*8.vY& d+//d; IdV4d }B?@//%%d+,T28%y' /d }Lt Ude QCr(#[- $\ Ke ~./r device_traceszDevice z ---------------- )rgrhrrHri)rrrdrs` @rtracer;sb ++-CP0d$/01 IIs"56 7 1I <<>ra  cn|tjdtdtj|}|ddt |dzz zz }t j|jd}tjd|dg}tjdt|jd |S) Nz;device argument is deprecated, plots now contain all device) stacklevelr8zsnapshot.pickle)rbase64z $VIZ_KINDz $SNAPSHOT) warningswarn FutureWarningpickledumpsrr b64encodedecodejson_memory_viz_templatereplacerepr)rviz_kinddevicebufferencoded_buffer json_formats r _format_vizrs  I  \\$ F gS[1_,--F%%f-44W=N**'8NSTUK  ' ' T(^ D L L[ rc.t||sd|Sd|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_segmentss r trace_plotrs3   !   -  rc Xddl}ddlm}ddlm}m}|j ijjD]}|j|jk(s|j}g}|rD|j|j|jfvr|j||j}|rD|j|j }|s|j j"dkDs||<|j$j't)dzD cgc]} gc} t)dzD cgc] } | dddgd c} dfddfd } fd } i} j*D]\}}\}}}t-||s||j.k(r| |||| ||f<7||j0k(r+| | j3||f|j4q||j6k(r@| | j3||f|j4| |||dz| ||dzf<||j8k(s| |||d | ||f<| j;Dcgc]'\\}}}|j4|d |d |df)}}}}t=t?|tAjBdD]\} }d| }|d}|D]D\} }}}||kr|dj||z dd|dj|d||d||z}F||dkse|dj|d|z dddDcgc] }|ds | c}d<dD]2}|dxx|dzcc<|dr|dj|ddd4Scc} wcc} wcc}}}wcc}w)Nr) _EventType)Action TensorKeyr)rrerprcr[)rrc>|jdk(r |jSS)Ncuda)typeindex)r device_counts r to_devicez'_profile_to_snapshot..to_devices ;;& <<  rc|j}|jj} d|}|d|d|kDr||d<t|d||z|d< jj ||}||j jnd} j |d}|D cgc]} dd| j d}} d ||d||d } |r d |j| | Scc} w) Nrrerpunknownr!noner)r r rr)rrr/rcr(categoryr) rstorageptrr _categoriesr0rlowerr)r/ tensor_keyversion during_tracerrrlr%stackrPrallocation_stacksmemory_profilerkr!s rallocatez&_profile_to_snapshot..allocates:,,-!!%%z"6* y> !S^d%:!C N  td{ L"--11*gF,4,@8==&&(i!%%j"5JOPQfa@PP    _ %f - 4 4Q 7Qs)C&c `dD](}d|j||d|dd|dd*y)N)rrrrr/rr()rrr/rcr()r)rrrrks rrz"_profile_to_snapshot..freesJ5 A _ %f - 4 4!&M!&M#Ho   rF)r,rr/r(rrrer[inactive)r/rWr)r/rWr.r(rp)T)"torchtorch._C._profilerrtorch.profiler._memory_profilerrr_memory_profile_op_tree sorted_nodestag AllocationparentPyCallPyCCallrfrom_allocation extra_fields alloc_sizerr rtimeline isinstanceCREATEDESTROYrrINCREMENT_VERSION PREEXISTINGitemsrroperator itemgetter)profiler3rrreventr;python_parentsrrzrr0r kv_to_elem_timerr*r+r/ blocks_at_endr[rl last_addrrr(r.r r/rkr!s @@@@@r_profile_to_snapshotrQs!-A,,.N((55 8 99 -- -\\FN::*"3"3Z5G5G!HH"))&1++E,>,>?Cu))44q8)7!#& 8::**,L&+L1,<&=>"> q 01   !     &H 4 J7E6M6M2v, Gd*i0  V]] "08z70SJ G, - v~~ % W 56 *BSBS8T U v// / W 56 *BSBS8T U4<j'A+5J GaK0 1v)) )08j'1J G, -(-7,<,<,> ( !Z5 :$$ %uV}eFmU8_UM"&"7X=P=PQR=STz"6* N %+ $ !AtT64H $$dY.>%TU M  /&*$  t I $ s<( ( M \*Y6L !(,4J+?QC3x=CQHZ #S LS^+8} M #l*;j!Q RS OY?  Z0Rs N5N!,N  N'N'c2t|}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)rQr)rJrrks r profile_plotrSUs$G,H x!96 BBrrct|d|S)NzAllocator State Historyr)rrs r segment_plotrUcs t6 ??r__main__ztorch.cuda.memory._snapshot()zpickled memory statistics from z#Visualize memory dumps produced by ) descriptionr)destc.|jddddy)N-o--outputz output.svgz$flamegraph svg (default: output.svg)defaulthelp) add_argument)rPs r_outputr`xs    7  rzmPrints overall allocation statistics and a visualization of how the allocators segments are currently filled.statsinput)r^z`Prints buffer of the most recent allocation events embedded in the snapshot in a Pythonic style.rzaGenerate a flamegraph that visualizes what memory is stored in each allocator segment (aka block)rzNGenerate a flamegraph the program locations contributing to CUDA memory usage.rzGenerate a flamegraph that shows segments (aka blocks) that have been added or removed between two different memorys snapshots.r rurv))rz]Generate a visualization over time of the memory usage recorded by the trace as an html file.)rUzfVisualize 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)rr]r^z4path to save the visualization(default: output.html)rZr[z output.htmlr\rz?visualize change to segments rather than individual allocationsz-sz --segments store_true)rr^c|dk(rtjj}n t|d}t j |}t |tr|gd}|S)N-rb)rtraces)sysr9ropenrloadrBlist)rrrs r_readrlsJ 3;   AT4 A{{1~ dD ! $3D rcht|d5}|j|dddy#1swYyxYw)Nw)rirH)rrrs r_writeros- $_  GGDM   s(1)rrrU)r)F)FFr)NF)Rrrgr rHr<rrDrhr functoolsr itertoolsrtypingrcache__all__rr%r*r2r4rr`rrr rrrrrrrrr rrrQrSrUros.pathr>realpathdirname__file__thedirremoveargparsefn_namepickledArgumentParserparseradd_subparsers subparsersr`rW add_parserstats_ar_trace_a segments_amemory_a compare_aplotscmd trace_plot_ar^r parse_argsrOrlrorrbroutputrArurvrr!rrrs    $ @!@,6>6*;+(9+.?+4((U%  y!"# OdYx $*EP C@s@ z WW  bggooh7 8F  -G/y9G $X $ $9'CF&&H&5J BK##G#EG w/tK##G#EG w/uK&&z{&KJG'2 J Y$$X;$GH '0 H >%%i[%II 8'2 71 I E" [!,,Sk,J !!'!8n!!$ dQU!VE!!$ MPT!U , TD  % %l