
    ^jy                    @   U d Z ddlmZ ddlZddlZddlZddlZddlmZ ddl	m
Z
 ddlmZ ddlZddlZddlmZ ddlmZmZmZ ddlmZ dd	lmZmZ dd
lmZ ddlmZmZ ddlmZ dZ dZ!dZ"dZ#dZ$ ejJ                  d      Z&dZ' ejJ                  d      Z(g dZ)d/dZ*i Z+de,d<   d0dZ-d Z.e'dz  Z/e$dz  Z0d1dZ1d2dZ2d Z3d Z4d Z5d  Z6d! Z7d" Z8e G d# d$e             Z9d3d%Z:d3d&Z;e G d' d(             Z<e G d) d*             Z=e G d+ d,             Z> G d- d.      Z?y)4u  Intra-kernel trace profiler for CuTe-DSL kernels.

Emits Chrome Trace JSON (compatible with Perfetto / chrome://tracing) from
per-warp instrumentation inserted directly into CuTe-DSL kernels.

Toggle with QUACK_TRACE=1 env var.  When disabled (the default) every trace
call is a compile-time no-op — the JIT never emits any profiling PTX.

Design decisions
----------------
**Two-timer approach (inspired by Triton Proton).**
NVIDIA GPUs expose two timers accessible from PTX:
  - %globaltimer — device-wide, ~1 GHz, synchronized across all SMs.
  - %clock64    — per-SM cycle counter, ~2.1 GHz on H100, *not* synchronized
                  across SMs (confirmed empirically: cross-SM spread > 400M ticks
                  vs ~200 ticks for globaltimer on the same launch).
We read %globaltimer once at init and once at flush (per warp) to anchor each
warp's timeline to a device-wide epoch, then read %clock64 for every event.
This gives us low-overhead per-event timestamps (local SM register read) while
retaining cross-SM comparability.  During post-processing the per-slot pair
    (init_globaltimer, init_clock64) and (final_globaltimer, final_clock64)
auto-calibrates the clock64-to-nanosecond conversion:
    ratio = (final_gt - init_gt) / (final_clk - init_clk)
    event_ns = init_gt + (event_clk - init_clk) * ratio

**Compact events (inspired by ThunderKittens).**
Each event is 8 bytes: a raw 32-bit %clock value and a packed (region_id,
event_type) tag, stored with a single v2.u32 streaming store.  The device
writes the raw clock — no subtraction needed.  The host computes deltas
during post-processing using init_clock from metadata with proper u32
wraparound.  Block and warp identity (constant per slot) are stored once
in per-slot metadata instead of per event.

**Minimal live registers.**
The TraceContext dataclass carries only 3 DSL values across loop iterations:
  - slot_ptr (64-bit)  — base of this warp's interleaved [metadata|events]
  - cnt     (32-bit)  — circular buffer write index
  - is_active (1-bit) — predicate for stores (warp leader AND warp sampling)
init_clk is NOT stored — the device writes raw clock values and the host
subtracts init_clk during post-processing.  This saves one register vs
computing deltas on device.

**Interleaved per-slot layout.**
Each warp's metadata and events are contiguous in memory:
    [meta₀ events₀ | meta₁ events₁ | ...]
This means the device needs only ONE pointer (slot_ptr) instead of separate
metadata and event pointers, saving another register.

**Warp sampling.**
An optional warp_ids parameter restricts profiling to specific warps.
Non-selected warps execute predicated stores that the GPU evaluates to
hardware no-ops — zero store bandwidth and no branch divergence.

**Auto-interned region names.**
ctx.b("mma") / ctx.e("mma") auto-assign integer IDs via a module-level
registry at JIT time.  The host reads the same registry at write_trace time.
No region_names parameter needed on either side.

Usage
-----
Host:
    with TraceSession("trace.json", grid_size=G, block_size=B) as sess:
        my_kernel[grid, block](..., sess.ptr)

Device (safe to call from all lanes):
    ctx = TraceContext.create(trace_ptr)
    ctx.b("load"); ctx.e("load")
    ctx.flush()
    )annotationsN)Optional)defaultdict)	dataclass)Int32Int64
const_expr)Arch)llvmnvvm)T   )storestore_v2)
ParamsBaseQUACK_TRACE      z<IHH(   z<QQQQII)$thread_state_uninterruptiblethread_state_iowaitthread_state_runningthread_state_runnablethread_state_sleepingthread_state_unknownbackground_memory_dumplight_memory_dumpdetailed_memory_dumpvsync_highlight_colorgeneric_workgoodbadterribleblackgreywhiteyellowoliverail_responserail_animation	rail_idle	rail_loadstartupheap_dump_stack_frameheap_dump_object_typeheap_dump_child_node_arrowcq_build_runningcq_build_passedcq_build_failedcq_build_abandonedcq_build_attempt_runnigcq_build_attempt_passedcq_build_attempt_failedrail_animatecq_build_attempt_runningc                 P    t         j                  j                  t        d      dk(  S )zECheck QUACK_TRACE=1.  Evaluated at JIT time so disabled = no codegen. 1)osenvirongetQUACK_TRACE_ENV     d/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/torch/_vendor/quack/trace.pyenabledrD      s    ::>>/2.#55rB   zdict[int, str]_REGION_REGISTRYc                    t         j                         D ]  \  }}|| k(  s|c S  t        t               }| t         |<   |S )z<Assign a stable integer ID to a region name.  JIT-time only.)rE   itemslen)nameridns      rC   _intern_regionrL      sG    "((* Q9J 
C SJrB   c                 ,    t         j                          y)zOClear the registry.  Called by TraceContext.create so each kernel starts fresh.N)rE   clearrA   rB   rC   _reset_region_registryrO      s    rB      c                "    t         | t        z  z   S )z+Per-slot size in bytes (metadata + events).)METADATA_SIZE
EVENT_SIZE)per_warp_caps    rC   
_slot_sizerU      s    <*444rB   c                    | t        |      z  S N)rU   )total_slotsrT   s     rC   _buf_total_bytesrY      s    L111rB   c            
     n    t        t        j                  t        j                         g dddd            S )Nzmov.u64 $0, %globaltimer;=lTFhas_side_effectsis_align_stackr   r   
inline_asmr   i64rA   rB   rC   _read_globaltimerrb      s3    EEG'! 	
	 	rB   c            
     n    t        t        j                  t        j                         g dddd            S )Nzmov.u64 $0, %clock64;r[   TFr\   r_   rA   rB   rC   _read_clock64rd      s3    EEG#! 	
	 	rB   c            
         t        j                  t        j                  t	        j
                         g dddd            S )Nzmov.u32 $0, %clock;z=rTFr\   )cutlassr   r   r`   r   i32rA   rB   rC   _read_clockrh      s7    ==EEG!! 	
	 	rB   c                 v    t        j                  t        j                  t	        j
                                     S rW   )rf   r   r   read_ptx_sreg_smidr   rg   rA   rB   rC   
_read_smidrk     s"    ==009::rB   c                r    t        j                  | t        |      t         j                  j                        S )z:Create a cute global-memory pointer from an Int64 address.)cutemake_ptrr   AddressSpacegmem)dtypeaddrs     rC   	_gmem_ptrrs     s%    ==dT->->-C-CDDrB   c                    t         j                  j                  j                         j	                         t
        j                  k\  rt        j                  t         j                  j                        dk(  r:t        j                  t        j                  t        j                                     S t        j                  t         j                  j                        dk(  r't        j                  t        j                               S t        dt         j                  j                         t         j"                  j%                         dk(  S )zReturn a DSL predicate for the warp leader thread.

    Uses nvvm.elect_sync() on SM90+ (hardware single-thread election),
    falls back to lane_idx() == 0 on older architectures.
          z&CUDA_VERSION.major must be >= 12, got r   )rf   base_dslBaseDSL_get_dslget_arch_enumr
   sm_90r	   CUDA_VERSIONmajorBooleanr   
elect_syncr   bool
ValueErrorrm   archlane_idxrA   rB   rC   _is_warp_leaderr     s     ((*88:djjHg22889R???4??1668#<== 4 4 : :;rA??4??#455EgFZFZF`F`Eabcc991$$rB   c                      e Zd ZU dZdZded<   dZded<   dZded	<   dZd
ed<   dZ	ded<   e
	 	 d	 	 	 	 	 dd       Zd ZddZddZddZddZddZddZddZy)TraceContexta  Per-warp trace recorder for use inside CuTe-DSL kernels.

    Use the ``create`` classmethod (not ``__init__``) to construct.  Named
    regions (ctx.b("mma") / ctx.e("mma")) are resolved to integer IDs at JIT
    time.  Optional warp_ids restricts profiling to specific warps.

    Usage::

        ctx = TraceContext.create(trace_ptr)
        ctx.b("load"); ctx.e("load")
        ctx.flush()
    r   intrT   Nztuple | Nonewarp_idszcute.Pointerslot_ptrzcutlass.Int32cntzcutlass.Boolean	is_activec                b   ||dz
  z  dk(  sJ d       t                |t        |      nd}t               rt        |du       r | ||ddd      S t        |t
        z  z   }t        j                  j                         \  }}}||z  |z  t        j                  j                  z   dz
  t        j                  j                  z  }t        j                  j                  t        j                  j                               }	t               }
t        j                  j                         \  }}}t        j                  j                         \  }}}|||z  z   ||z  |z  z   }||z  |	z   }t        t         t#        |            }|||z  z   }t%               }|*t'        j(                  d      }|D ]
  }||	|k(  z  } ||z  }n|}|	dz  |
dz  d	z  z  }||d
z  z  }t+               }t-               }t/        |||d       t/        |dz   ||d       t/        |dz   ||d        | |||t!        d      |      S )u  Create and initialize a TraceContext.  Safe to call from all lanes.

        Only lane 0 (warp leader) performs stores; all other lanes execute
        the arithmetic but skip the writes via predication.  The caller does
        NOT need an ``if is_warp_leader():`` guard.

        Region names are auto-interned by ctx.b("name") / ctx.e("name") via a
        module-level registry — no explicit region_names list needed.
        r   r   per_warp_cap must be power of 2N)rT   r   r   r   r   F?   i        cscopr   r   )rO   tuplerD   r	   
META_ELEMSEVENT_ELEMSrm   r   	block_dim	WARP_SIZEmake_warp_uniformwarp_idxrk   	block_idxgrid_dimrs   r   r   r   rf   r~   rb   rd   r   )clsbuf_ptrrT   r   
SLOT_ELEMSbdxbdybdzwarps_per_blockr   smidbidxbidybidzgdxgdygdzlinear_blockslotbufr   	is_leaderr   widpackedinfogtclk64s                               rC   createzTraceContext.createB  s:     q 01a7Z9ZZ7 &.&:5?yJw$7)!   ,"<<
		++-S#9s?TYY-@-@@1DI\I\\99..tyy/A/A/CD|  99..0dD		**,S#dSj(4#:+;;o-8 uW~.
** $%	.I :%S9	:!I-I!I T/te|&9:v|,   	hI40hlE9$7hlD)6%a
 	
rB   c                >   | j                   yt               }t               }t        | j                   dz   || j                  d       t        | j                   dz   || j                  d       t        | j                   dz   | j
                  | j                  d       y)zEWrite final timer pair and event count.  Safe to call from all lanes.NrP   r   r   r   	   )r   rb   rd   r   r   r   )selfr   clks      rC   flushzTraceContext.flush  ss    ==  odmmaT^^>dmmadnn$?dmma4>>tDrB   c                2   | j                   y t               }t        | j                  | j                  dz
  z  t
        z  z   }t        |      t        |      dz  z  }t        | j                   |z   ||| j                  d       | xj                  dz  c_        y )Nr   r   r   r   )	r   rh   r   r   rT   r   r   r   r   )r   	region_id
event_typer   evt_offtags         rC   _recordzTraceContext._record  s|    == mD,=,=,A BkQQI%
"3r"9:(#sDNNMArB   c                0    | j                  |t               y rW   )r   EVENT_BEGINr   r   s     rC   record_bzTraceContext.record_b  s    Y,rB   c                0    | j                  |t               y rW   )r   	EVENT_ENDr   s     rC   record_ezTraceContext.record_e  s    Y	*rB   c                0    | j                  |t               y rW   )r   
EVENT_MARKr   s     rC   record_mzTraceContext.record_m  s    Y
+rB   c                B    | j                  t        |      t               y rW   )r   rL   r   r   rI   s     rC   bzTraceContext.b  s    ^D);7rB   c                B    | j                  t        |      t               y rW   )r   rL   r   r   s     rC   ezTraceContext.e  s    ^D)95rB   c                B    | j                  t        |      t               y rW   )r   rL   r   r   s     rC   mzTraceContext.m  s    ^D):6rB   )   N)r   zOptional[Int64]rT   r   r   z"tuple[int, ...] | list[int] | None)r   r   r   r   )r   r   )rI   str)__name__
__module____qualname____doc__rT   __annotations__r   r   r   r   classmethodr   r   r   r   r   r   r   r   r   rA   rB   rC   r   r   &  s     L#!Hl!
 "Hl!C!%I%  !7;	J
 J
 J
 5	J
 J
XE-+,867rB   r   c                    | dz  S )Nr   rA   r   s    rC   _unpack_warpr     s    D=rB   c                    | dz	  S )Nr   rA   r   s    rC   _unpack_smidr     s    Q;rB   c                  h    e Zd ZU dZded<   ded<   ded<   ded<   ded<   edd       Zedd	       Zy
)_Eventz:Reconstructed event with absolute timestamp (nanoseconds).r   tsidtypeblock	warp_smidc                ,    t        | j                        S rW   )r   r   r   s    rC   warpz_Event.warp      DNN++rB   c                ,    t        | j                        S rW   )r   r   r   s    rC   r   z_Event.smid  r   rB   Nreturnr   )r   r   r   r   r   propertyr   r   rA   rB   rC   r   r     sD    DGG
IJN, , , ,rB   r   c                      e Zd ZU dZded<   ded<   ded<   ded<   ded<   ded<   edd	       Zedd
       Zedd       Zedd       Z	ddZ
y)	_SlotMetaz(Per-slot metadata read back from device.r   init_gtinit_clkfinal_gt	final_clkr   r   c                     | j                   dz  S )N  r   r   s    rC   r   z_SlotMeta.block  s    yy6!!rB   c                &    | j                   dz	  dz  S )Nr   r   r   r   s    rC   r   z_SlotMeta.warp_smid  s    		R6))rB   c                     | j                   dz  S )z2Low 32 bits of init_clock64 (%clock at init time).    )r   r   s    rC   
init_clk32z_SlotMeta.init_clk32  s     }}z))rB   c                z    | j                   | j                  z
  }|dkD  r| j                  | j                  z
  |z  S dS )u>   clock64 ticks → nanoseconds conversion factor for this slot.r         ?)r   r   r   r   )r   dclks     rC   ratioz_SlotMeta.ratio  s9     ~~-8<q,4IcIrB   c                ^    || j                   z
  dz  }| j                  || j                  z  z   S )z9Convert a raw 32-bit clock value to absolute nanoseconds.r   )r   r   r   )r   raw_clock32deltas      rC   clock_to_nsz_SlotMeta.clock_to_ns  s-    t.*<||edjj000rB   Nr   )r   float)r   r   r   r  )r   r   r   r   r   r   r   r   r   r   r  rA   rB   rC   r   r     sx    2LMMN
I	H" " * * * * J J
1rB   r   c                  T    e Zd ZU dZded<   dZded<   dZded<   dZded	<   d
Zded<   y)TraceWriteOptionsgMbP?r  scaleTr   emit_complete_eventsFgroup_by_smidemit_summary_json   r   summary_hist_binsN)	r   r   r   r  r   r  r  r  r
  rA   rB   rC   r  r    s7    E5!%$%M4"t" s rB   r  c                      e Zd ZdZ	 ddddddd	 	 	 	 	 	 	 	 	 	 	 ddZed	        Zd
 Zd Zd Z	d Z
ddZddZddZdddZed        Zd Zy)TraceSessiona  Host-side profiling session.

    Allocates a single contiguous device buffer, provides one pointer (sess.ptr)
    to pass to the kernel, and writes Chrome Trace JSON on exit.

    Can be used as a context manager for automatic sync + write:

        with TraceSession("trace.json", grid_size=G, block_size=B) as sess:
            my_kernel[grid, block](..., sess.ptr)
        # trace.json written here
    Nr   r   r	  cuda)rT   	grid_size
block_sizer   devicec               b   ||dz
  z  dk(  sJ d       || _         || _        || _        |dz   dz  | _        |t	        |      nd | _        || _        t               sd | _        y | j                  | j                  z  }t        j                  t        ||      t        j                  |      | _        y )Nr   r   r          )rq   r  )pathrT   total_blocksr   r   r   r  rD   d_buftorchzerosrY   uint8)r   r  rT   r  r  r   r  rX   s           rC   __init__zTraceSession.__init__  s     q 01a7Z9ZZ7	(% *RB6+3+?hTyDJ''$*>*>>[[[,7++

rB   c                j    ddl m} | j                    || j                  j                               S dS )zwDevice pointer as Int64, or None when tracing is disabled.
        Pass directly as an Optional[Int64] kernel argument.r   )r   N)cutlass.cutlass_dslr   r  data_ptr)r   r   s     rC   ptrzTraceSession.ptr0  s-     	./3zz/EuTZZ((*+O4OrB   c                R    | j                   | j                   j                          y y rW   )r  zero_r   s    rC   resetzTraceSession.reset8  s!    ::!JJ "rB   c                    | S rW   rA   r   s    rC   	__enter__zTraceSession.__enter__<  s    rB   c                    | j                   rCt               r9t        j                  j	                          | j                  | j                          y)NF)r  rD   r  r  synchronizewrite_trace)r   exc_typeexc_valexc_tbs       rC   __exit__zTraceSession.__exit__?  s2    99JJ""$TYY'rB   c                R    | j                   j                         j                         S rW   )r  cpunumpyr   s    rC   
_raw_byteszTraceSession._raw_bytesG  s    zz~~%%''rB   c           
         | j                   | j                  z  }t        | j                        }t	        |      D cg c]!  }t        t        j                  |||z         # c}S c c}w rW   )r  r   rU   rT   ranger   METADATA_STRUCTunpack_from)r   rawrX   
slot_bytesss        rC   _read_metadatazTraceSession._read_metadataJ  sa    ''$*>*>> 1 12
RWXcRd
MNI223JGH
 	
 
s   &A%c                   | j                   | j                  z  }t        | j                        }g }t	        |      D ]  }||   j
                  }t        || j                        }|| j                  kD  r|| j                  dz
  z  nd}	||z  t        z   }
||   }t	        |      D ]  }|	|z   | j                  dz
  z  }t        j                  ||
|t        z  z         \  }}}|j                  t        t        |j                  |            |||j                  |j                                  |j#                  d        |S )Nr   r   )r   r   r   r   r   c                   | j                   t        | j                        | j                  t	        | j                        | j
                  dk(  rd| j                  fS | j
                  dk(  rd| j                  fS d| j                  fS )Nr   r   r   )r   r   r   r   r   r   r   )evs    rC   <lambda>z+TraceSession._read_events.<locals>.<lambda>l  su    R\\*R\\*WW\ 
 .0WW\ 
 @A rB   )key)r  r   rU   rT   r0  r   minrR   EVENT_STRUCTr2  rS   appendr   r   r  r   r   sort)r   r3  metasrX   r4  eventsr5  r   rK   startslot_events_offmetaiidxraw_clkeidetypes                    rC   _read_eventszTraceSession._read_eventsQ  sK   ''$*>*>> 1 12
{# 	A(,,CC**+A7:T=N=N7NSD--12TUE*n}<O8D1X qyT%6%6%:;&2&>&>#cJ&66'#e t//89""jj"&..	, 	 	 		
 rB   c                @    t         j                  |t        |            S rW   )rE   r?   r   )r   rJ   s     rC   _region_namezTraceSession._region_namew  s    ##CS22rB   c                   t               sy |xs
 t               }| j                         }| j                  |      }| j	                  ||      }|st        d       y t        j                  t        j                  j                  |      xs dd       |d   j                  }g }t               }t               }	i }
|D ]  }|j                  |j                  |j                  }}}|
j                  ||       |j                   r|n|}|j                   r|dz  |z  n|dz  }|j#                  |       |	j#                  ||f        t%        |      D ]`  }|j                   rd|d	nd|
j'                  |d      d	d
|d}|j)                  dd|dd|id       |j)                  dd|dd|id       b t%        |	      D ]]  \  }}|j                   rd|dz	  dd|dz  d}n	d|dz  d}|j)                  dd||d|id       |j)                  dd||d|id       _ |j*                  r| j-                  ||      }|D ]  \	  }}}}}}}}}| j/                  |      ||t0        |t3        t0              z     |||dd}|dk(  r2|j5                  d||z
  |j6                  z  ||j6                  z         n$|j5                  dd||z
  |j6                  z         |j)                  |        ng }|D ]  }|j                  |j                  |j                  }}}|j                   r|n|}|j                   r|dz  |z  n|dz  }|j8                  dk(  rd n|j8                  d!k(  rd"nd}| j/                  |j:                        ||j                  |z
  |j6                  z  ||t0        |j:                  t3        t0              z     |||dd#}|j8                  t<        k(  rd|d$<   |j)                  |        t?        |d%      5 }tA        jB                  d&|d'|       d d d        t        d(t3        |       d)|        |j*                  r!|jD                  r| jG                  |||       y y y # 1 sw Y   QxY w)*Nz&intra_kernel_profiler::trace: 0 events.T)exist_okr   r   r  zSM 03dz Block 04dMprocess_namerI   )phrI   pidtidargsprocess_sort_index
sort_indexzBlock z Warp r   02dzWarp thread_namethread_sort_index)smr   r   )rI   rU  rV  cnamerW  X)rT  r   durrE  t)rT  r5  r   Br   E)rI   rT  r   rU  rV  r^  rW  r5  wns)displayTimeUnittraceEventszintra_kernel_profiler::trace: z events -> )$rD   r  r.  r6  rJ  printr=   makedirsr  dirnamer   setr   r   r   
setdefaultr  addsortedr?   r>  r  _pair_begin_endrL  _CNAME_LISTrH   updater  r   r   r   openjsondumpr  _write_summary)r   r  optr3  r@  rA  min_tstrace_events	used_pidsused_threadsblock_to_smidr   r]  r   rd  rU  rV  pnametname
out_eventsr   r`  rJ   kindr9  rT  fs                              rC   r&  zTraceSession.write_trace|  s   y(&(oo##C(""3.:;
BGGOOD)0S4@#% "e	-0U(* 	)Avvqww1B$$Q+))"qC$'$5$5AFa<AFCMM#c3Z(	) )$ 	C $$ c#Y=,,S!4S9S	J 
 N3qSY[`Rab 0)3/	$ |, 	HC   #fS4Z4DEr	#/M#cTZ\aSbc /)3/	& ##--fc:J:D (6Cc3aB --c2(s;/?)?@#%1= 19II"v+)BciiIXIIf		0IIJ##B'( J (66177AFFqA--b1(+(9(9Q!|BFFaKSQVVq[cc --add344&=CII5(K0@)@A#%1= 66Z'!BsG##B'!($ $_ 	QII$|LaP	Q.s6{m;tfMN##(=(=j#6 )>#		Q 	Qs   >Q((Q1c                   t        d       }g }| D ]<  }|j                  |j                  |j                  }}}|j                  r|n|}|j                  r|dz  |z  n|dz  }	||	f}
|j
                  t        k(  r,||
   |j                     j                  |j                         |j
                  t        k(  rh||
   |j                     }|s|j                         }|j                  |k\  s|j                  ||j                  |z
  ||	|j                  d|||f	       |j                  |j                  d||	|j                  d|||f	       ? |S )zMMatch B/E events into (ts, dur, pid, tid, rid, kind, block, warp, sm) tuples.c                      t        t              S rW   )r   listrA   rB   rC   r:  z.TraceSession._pair_begin_end.<locals>.<lambda>  s    {[_O` rB   r   r  r   r   )r   r   r   r   r  r   r   r   r>  r   r   pop)rA  rv  thread_statesoutr   r]  r   rd  rU  rV  r;  stackt0s                rC   ro  zTraceSession._pair_begin_end  s/    <GG`;a 	CAvvqww1B))"qC$'$5$5AFa<AFC*Cvv$c"144(//59$%c*1440Bttrz

Br	3QTT1aB#OP

ADD!S#qttQ1bAB	C 
rB   c                  ' |j                  d      r|j                  dd      d   n|}|dz   }t        t              }|D ]3  \	  }}}	}
}}}}}|dk(  s||   j	                  ||j
                  z         5 g }t        |      D ]  }||   }t        |      }|dk(  rt        |      |z  't        |      t        |      }}t        'fd|D              |z  }|dkD  rt        'fd|D              |dz
  z  nd}t        '      dkD  r&|dkD  r!t        j                  |      t        '      z  nd }|j                  xs d}dg|z  }||kD  rK|D ]E  }|t        t        t        dt        d	||z
  ||z
  z              |z        |dz
        xx   dz  cc<   G n||d<   ||kD  r||z
  |z  nd}i }d
D ]m  }|dz  d|} }}t!        |      D ]L  \  }!}"|}#||"|z  z  }||k\  s|"|z  }$|$dkD  rt        dt        d	||#z
  |$z              nd}%|||!z  z   |%|z  z   }  n | |d| <   o |j	                  || j#                  |      |'||||||||||D "cg c]  }"|"|z  	 c}"dd        t%        |d      5 }&t'        j(                  |d|j
                  | j*                  | j,                  | j.                  |d|&d       d d d        t1        d|        y c c}"w # 1 sw Y   xY w)Nz.jsonr   r   z_summary.jsonc              3  .   K   | ]  }|z
  d z    ywr   NrA   .0dmeans     rC   	<genexpr>z.TraceSession._write_summary.<locals>.<genexpr>  s     8a1t8/8   c              3  .   K   | ]  }|z
  d z    ywr  rA   r  s     rC   r  z.TraceSession._write_summary.<locals>.<genexpr>  s     ;a$h1_;r  r	  g        r   )   
      2   K   Z   _   c   g      Y@p)binsr<  maxprob)regionrI   countmean_durcv_durmin_durmax_durvar_dur_popvar_dur_samplepercentileshistrd  re  )tracerf  r  blocksr   rT   regionsr   )indentz)intra_kernel_profiler::trace: summary -> )endswithrsplitr   r  r>  r  rn  rH   sumr<  r  absmathsqrtr
  r   	enumeraterL  rr  rs  rt  r  r   rT   rh  )(r   
trace_pathr~  rv  basesummary_pathregion_statsr   r`  rU  rV  rJ   r  r   rd  r]  r  dursrK   min_dmax_dvar_pop
var_samplecvr  r  r  w_binpctsr  qcumvalrE  cprevr  fracr  r  s(                                          @rC   ru  zTraceSession._write_summary  sa   3=3F3Fw3Oz  !,Q/U_o-/:4/@6@ 	:2BS#sD!QqyS!((syy9	: ,' 6	C$DD	AAvt9q=Dt9c$i5E84881<GIJQ;d;;q1uETUJ69$i!mA:&T2SWB((/CD3:Du} ACCS1u92O)P QTX XY[_bc[cd 
 Q.3emUU]d*ED4 
$%ie3%dO DAqD1q5LCax 1uHLqs3C!d(d1B(CDVW#eai/$,> !$qW
$ NN! --c2 $ $$#*&0#' $$$04 51Q 5	G6	p ,$ 	II''+ YY"//'+';';$($5$5& 	 	9,HI' !6
	 	s   K$.A	KKrW   )r  z
str | NonerT   r   r  r   r  r   r   z"list[int] | tuple[int, ...] | Noner  zstr | torch.device)r   zlist[_SlotMeta])r   zlist[_Event])rJ   r   r   r   )r  r   rv  zTraceWriteOptions | None)r   r   r   r   r  r   r  r!  r#  r*  r.  r6  rJ  rL  r&  staticmethodro  ru  rA   rB   rC   r  r    s    
  
 !7;%+

 	

 
 
 5
 #
: P P(
$L3
i7V  .PJrB   r  )r   r   )rI   r   r   r   )rT   r   r   r   )rX   r   rT   r   r   r   )r   r   r   r   )@r   
__future__r   rs  r  r=   structtypingr   collectionsr   dataclassesr   r  rf   cutlass.cuterm   r   r   r	   cutlass.base_dsl.archr
   cutlass._mlir.dialectsr   r   r  r   
copy_utilsr   r   cute_dsl_utilsr   r@   r   r   r   rS   Structr=  rR   r1  rp  rD   rE   r   rL   rO   r   r   rU   rY   rb   rd   rh   rk   rs   r   r   r   r   r   r   r  r  rA   rB   rC   <module>r     sm  DL #   	   # !    , , & - ! ' &
  	

 
v}}V$ &--	*%P6 $& . %* a
Ao5
2


;E
%* O7: O7 O7n , , ,$  1  1  1F ! ! !IJ IJrB   