DŽgca ddlZddlZddlZddlZddlZddlZddlmZddlm Z ddl m Z ddl Z ddl Z ddlZedZgdZdLdZedZdMdZd Zd ZdNd Zd Zefd ZefdZefdZdZGddZdZdZdZdeedZ dZ!dZ"dZ#dZ$dOdZ%dZ&dNdZ'dNde fd 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)Any)groupby)format_flamegraphsegmentsmemorycomparec`|d}|d}|s|jdd}|d}|d|d|S)Nlinefilename/name:)split)f full_filenameifnamefuncs ^/home/mcse/projects/flask_80/flask-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|dst||#c}Scc}w)Nrr )reversedr%r)framesrreversers r _frames_fmtr*;s>&!28 dQM!F)UVWaUb>+ , ./@.ABC"" \^o qw.?@A }g 6DZ__Z__W^_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_fragmentes!xx FD9::rr,rTstate  rr-r(z; r/z; )r4rH) rprefixblocksrVr1r(accounted_for_sizehszgapss r _write_blocksr`ds*;> A )5a &F& GGvha' |1_V-D,EQGYFZZ\] ^!" y\ J{^"b("q=x[FGGvha' |1_V5L4MQrdRTUVGGvha' |>"RHI Jy--  GGvha' |8D6< =!>rctj}|dD] }d|dd|d}t|||d"||jS)Nrstream_stream;seg_addressr[ioStringIOr`getvaluesnapshotrrsegrZs rrr{s_ 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|rd (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 g}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|rd!t|d"t|d#nd}!|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..sq TU6Wrkeyrpr[rWrTr/rXgryAz 0123456789* rcr|z , stream_z stream_[z] z allocated, z freerYz segments: ztotal_reserved: ztotal_allocated: rz internal + z external)z total_free: )rgrhrHr PAGE_SIZEsortedr4appendrange enumeratechrordmaxminintrUrlenlegendri)"datarouttotal_reservedtotal_allocatedrrrlseg_free_externalseg_free_internal seg_allocated all_rangesboffsetr1activerzallocated_sizensegoccupiedfrac active_sizerstart_r/finish_startfinishmjsercbodyinternal_externals" rsegsumrsH ++-CII'i(8'9DENOMMd:&,WX0^#l++  X !AwZ#55F$0O!>!!7ND"AB/ !QvY%??!!QvY.! qy G ! =(** ** L!A%)3a7!&t-AC--"4[))) )2:)> $ %A%f 4 K}Gi'Eki/!3FCv3/1r6:;A5&) $I .!a%9!45QAEY..A;#%"/DGbL0A"BHQK"#HQK $ $8})3x=//Jwwx  #44}DLHYYYY/28}/A8CM?+r |  ) II$r%L(9":!;<%&79JKLERXQYY[] ^_0^bII 3tJ/0145II ~!6 7r:;II!%"8!9<=cp"U=12,u]?S>TT^_vxII \-GHKLIIf = (? :n LL L <<>?.)s - L6 L;ctjfd}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.._name$s\!~~''?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_segment,s%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..formatsB#%!"$   ?  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' 5eDk],uUY{m[jkl QCr(#W+ $X Ke ~./r device_traceszDevice z ---------------- )rgrhrrHri)rrrdrs` @rtracersb ++-CN0^$/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 ' [ ABrc.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_plotrs# t]5qw xxPoqw xxrc Xddl}ddlm}m}ddlm}|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)Action TensorKey) _EventTyper)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 NL 14$;?L!--11*gF,4,@8==&&(i!%%j"5JOPQfa@PPdaSXfn o  _ %f - 4 4Q 7 Qs)C&c `dD](}d|j||d|dd|dd*y)N)rrrrr/rr()rrr/rcr()r)rrrrks rrz"_profile_to_snapshot..freesM5 RA _ %f - 4 4>CFm>CFm@A@Eh 6Q R RrF)r.rr/r(rrrer[inactive)r/rWr)r/rWr.r(rp)T)"torchtorch.profiler._memory_profilerrrtorch._C._profilerr_memory_profile_op_tree sorted_nodestag AllocationparentPyCallPyCCallrfrom_allocation extra_fields alloc_sizerr"rtimeline isinstanceCREATEDESTROYrrINCREMENT_VERSION PREEXISTINGitemsrroperator itemgetter)profiler5rrreventr=python_parentsrrzrr2r kv_to_elemtimerr,r-r/ blocks_at_endr[rl last_addrrr(r0r"r1rkr#s @@@@@r_profile_to_snapshotrSs!A-,,.N((55 8 99 -- -\\FN::*"3"3Z5G5G!HH"))&1++E,>,>?Cu))44q8)7!#& 8 ::**,L&+L1,<&=>"> 38 q8H2I K)/ !'!%$% ! " $KH "RJ 6D5L5L h1f+z7T*i0  V]] "08z70SJ G, - v~~ % W 56 *BSBS8T U v// / W 56 *BSBS8T U4%TU M $9K_cou!v wt I  $ s<( ( M #l*;i*GR\!] ^ _,4J+?QC3x=CQHZ #S LS^+8} M #l*;j!Q RS O]?KtMRs 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)rSr)rLrrks r profile_plotrUs$G,H x!96 BBrrct|d|S)NzAllocator State Historyr)rrs r segment_plotrWs 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_outputrb's tZDjkrzmPrints 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.)rWzfVisualize 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)r r_r`z4path to save the visualization(default: output.html)r\r]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)sysr9ropenrloadrDlist)rrrs r_readrnUsJ 3;   AT4 A{{1~ dD ! $3D rcht|d5}|j|dddy#1swYyxYw)Nw)rkrH)rrrs r_writerq_s- $_  GGDM   s(1)rrrW)r)F)FFr)NF)Rrrjr<rgrDr  functoolsrtypingr itertoolsrrrrJcache__all__rr%r*r2r4rr`rrr rrrrrrrrr rrrSrUrWros.pathr>realpathdirname__file__thedirremoveargparsefn_namepickledArgumentParserparseradd_subparsers subparsersrbrY add_parserstats_aratrace_a segments_amemory_a compare_aplotscmd trace_plot_ar`r parse_argsrOrnrqrrdroutputrArurvrr!rrrs   $ @!>e , *>.*;+(9+.?+2((U%  y!"# KZVr B ygR C@s@ z WW  bggooh7 8F  -G/y9G $X $ $3VW^V_1` aF&&H&5JlBK##G#EG w/tK##G#EG w/uK&&z{&KJG'2 JbK$$X;$GH '0 H>K%%i[%II 8'2 71 I E" Z[!,,Sk,J !!'!8n!!$ dQU!VE!!$ MPT!U , TD  % %dLTX % Y Z    D {{j TZZ t{{HTN+  TZZ t{{F4L)  TZZ  fTl  TZZ  eDk  !t{{#djj!t{{GFE23  $TZZ t{{JtDKKt}}]^  &TZZ t{{LdkkBC '{r