Ë
    ÄÓ^jÙy  ã                  ó@  — 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   € ÜÜ�‰Ü�E‰E‹GØØ'ØØ!Ø ô	
ó	ð 	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   € ÜÜ�‰Ü�E‰E‹GØØ#ØØ!Ø ô	
ó	ð 	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   € Ü�=‰=Ü�‰Ü�E‰E‹GØØ!ØØ!Ø ô	
ó	ð 	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\  rçt        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 )z¹Return 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Ó:¼d¿j¹jÒHÜ×Ñœg×2Ñ2×8Ñ8Ó9¸RÒ?Ü—?‘?¤4§?¡?´1·6±6³8Ó#<Ó=Ð=Ü×Ñ¤× 4Ñ 4× :Ñ :Ó;¸rÒAÜ—?‘?¤4§?¡?Ó#4Ó5Ð5äÐEÄg×FZÑFZ×F`ÑF`ÐEaÐbÓcÐcÜ�9‰9×ÑÓ 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Ñ1°aÒ7ÐZÐ9ZÓZÐ7ÜÔ Ø&.Ð&:”5˜”?ÀˆäŒyœJ w°$ Ô7ÙØ)Ø!ØØØôð ô   ,´Ñ"<Ñ<ˆ
äŸ	™	×+Ñ+Ó-‰ˆˆS�#Ø ™9 s™?¬T¯Y©Y×-@Ñ-@Ñ@À1ÑDÌÏÉ×I\ÑI\Ñ\ˆÜ—9‘9×.Ñ.¬t¯y©y×/AÑ/AÓ/CÓDˆÜ‹|ˆô  Ÿ9™9×.Ñ.Ó0Ñˆˆd�DÜŸ	™	×*Ñ*Ó,‰ˆˆS�#Ø˜d S™jÑ(¨4°#©:¸Ñ+;Ñ;ˆØ˜oÑ-°Ñ8ˆô œœu W›~Ó.ˆØ˜ 
Ñ*Ñ*ˆô $Ó%ˆ	ØÐÜŸ™¨Ó.ˆIØò :�Ø%¨°S©Ñ9‘	ð:à! IÑ-‰Ià!ˆIð ˜T‘/ t¨e¡|¸Ñ&9Ñ:ˆØ˜v¨™|Ñ,ˆô Ó ˆÜ“ˆô 	ˆh˜˜I¨4Õ0Üˆh˜‰l˜E 9°$Õ7Üˆh˜‰l˜D )°Õ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ˆÜˆd�m‰m˜aÑ  T§^¡^¸Õ>Üˆd�m‰m˜aÑ  d§n¡n¸$Õ?Üˆd�m‰m˜aÑ §¡¨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Ñ BÄkÑQÑQˆÜ�IÓ¤%¨
Ó"3°rÑ"9Ñ:ˆÜ�—‘ Ñ(¨#¨s°D·N±NÈÕ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�#ÓØ!€HˆlÓ!ð
 "€HˆlÓ!Ø€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Ð  ó   € ä˜DŸN™NÓ+Ð+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   € à�y‰y˜6Ñ!Ð!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ÐIÀcÐ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Ÿ™Ñ.°*Ñ<ˆØ�|‰|˜e d§j¡jÑ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)ÚTraceWriteOptionsgü©ñÒMbP?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   … à€Eˆ5ÓØ!%Ð˜$Ó%Ø€M�4ÓØ"Ð�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Ñ1°aÒ7ÐZÐ9ZÓZÐ7ØˆŒ	Ø(ˆÔØ%ˆÔØ *¨R¡°BÑ6ˆÔØ+3Ð+?œ˜hœÀTˆŒØˆŒäŒyØˆDŒJØà×'Ñ'¨$×*>Ñ*>Ñ>ˆÜ—[‘[Ü˜[¨,Ó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-   € õ 	.à/3¯z©zÐ/E‰u�T—Z‘Z×(Ñ(Ó*Ó+ÐOÈ4ÐOrB   c                óR   — | j                   �| j                   j                  «        y y rW   )r  Úzero_ræ   s    rC   ÚresetzTraceSession.reset8  s!   € Ø�:‰:Ð!Ø�J‰J×ÑÕð "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   € Ø�9Š9œœÜ�J‰J×"Ñ"Ô$Ø×Ñ˜TŸY™YÔ'ØrB   c                óR   — | j                   j                  «       j                  «       S rW   )r  ÚcpuÚnumpyræ   s    rC   Ú
_raw_byteszTraceSession._raw_bytesG  s   € Ø�z‰z�~‰~Ó×%Ñ%Ó'Ð'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ˆ
äRWÐXcÓRdö
ØMNŒI”×2Ñ2°3¸¸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Ÿ\™\Ó*Ø—W‘W ’\�Ø—‘ð€ ð
 .0¯W©W¸ª\¨Ø—‘ð€ ð
 @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˜×*Ñ*Ó+ˆAØ7:¸T×=NÑ=NÒ7N�S˜D×-Ñ-°Ñ1Ò2ÐTUˆEà *™n¬}Ñ<ˆOØ˜‘8ˆDÜ˜1“Xò �Ø˜q‘y T×%6Ñ%6¸Ñ%:Ñ;�Ü&2×&>Ñ&>ØØ# c¬JÑ&6Ñ6ó'Ñ#�˜˜eð —‘ÜÜ˜t×/Ñ/°Ó8Ó9ØØ"Ø"Ÿj™jØ"&§.¡.ôõñð	ð, 	�‰ñð 	ô 		
ð ˆ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ØØÒ(Ô&Ó(ˆà�o‰oÓˆØ×#Ñ# CÓ(ˆØ×"Ñ" 3¨Ó.ˆÙÜÐ:Ô;Øä
�‰”B—G‘G—O‘O DÓ)Ò0¨S¸4Õ@Ø˜‘—‘ˆØ#%ˆô "›eˆ	Ü-0«UˆØ(*ˆØò 	)ˆAØ—v‘v˜qŸw™w¨¯©�1�ˆBØ×$Ñ$ Q¨Ô+Ø×)Ò)‘"¨qˆCØ$'×$5Ò$5�A˜‘F˜a’<¸AÀ¹FˆCØ�M‰M˜#ÔØ×Ñ˜c 3˜ZÕ(ð	)ô ˜)Ó$ò 	ˆCð ×$Ò$ð �c˜#�Y‘à˜=×,Ñ,¨S°!Ó4°SÐ9¸ÀÀSÀ	ÐJð ð
 ×ÑØ N¸3ÀqÐSYÐ[`ÐRaÑbôð ×ÑàØ0ØØØ)¨3Ð/ñõð	ô$ ˜|Ó,ò 	‰HˆC�Ø× Ò Ø  ¨¡¨# ¨f°S¸4±ZÀÐ4DÐE‘à  r¡	¨#˜Ð/�Ø×ÑØ M¸#ÀcÐTZÐ\aÐSbÑcôð ×ÑàØ/ØØØ)¨3Ð/ñõð	ð& ×#Ò#Ø×-Ñ-¨f°cÓ:ˆJØ:Dò (Ñ6��C˜˜c 3¨¨a°°Bà ×-Ñ-¨cÓ2ØØÜ(¨¬s´;Ó/?Ñ)?Ñ@Ø#%°¸1Ñ=ñ�ð ˜1’9Ø—I‘I ¨"¨v©+¸¿¹Ñ)BÈÈcÏiÉiÉ�IÕXà—I‘I ¨°°f±ÀÇ	Á	Ñ0I�IÔJØ×#Ñ# BÕ'ò(ð ˆJØò (�ØŸ6™6 1§7¡7¨A¯F©F�q�A�Ø×-Ò-‘b°1�Ø(+×(9Ò(9˜˜Q™ !’|ÀÀBÁ�ØŸF™F ašK‘S°Q·V±V¸q²[©cÀc�à ×-Ñ-¨a¯d©dÓ3ØØŸ4™4 &™=¨C¯I©IÑ5ØØÜ(¨¯©´´KÓ0@Ñ)@ÑAØ#%°¸1Ñ=ñ�ð —6‘6œZÒ'Ø!�B�s‘GØ×#Ñ# BÖ'ð!(ô$ �$˜‹_ð 	Q Ü�I‰I¨$¸|ÑLÈaÔP÷	QäÐ.¬s°6«{¨m¸;ÀtÀfÐ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Ø—v‘v˜qŸw™w¨¯©�1�ˆBØ×)Ò)‘"¨qˆCØ$'×$5Ò$5�A˜‘F˜a’<¸AÀ¹FˆCØ˜�*ˆCØ�v‰vœÒ$Ø˜cÑ" 1§4¡4Ñ(×/Ñ/°·±Õ5Ø—‘œ9Ò$Ø% cÑ*¨1¯4©4Ñ0�ÚØŸ™›�BØ—t‘t˜r“zØŸ
™
 B¨¯©¨r©	°3¸¸Q¿T¹TÀ1ÀaÈÈBÐ#OÖPà—
‘
˜AŸD™D ! S¨#¨q¯t©t°Q¸¸1¸bÐ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   øè ø€ Ò8¨a˜1˜t™8¨�/Ñ8ùó   ƒc              3  ó.   •K  — | ]  }|‰z
  d z  –— Œ y­wr‹  rA   rŒ  s     €rC   r�  z.TraceSession._write_summary.<locals>.<genexpr>  s   øè ø€ Ò;°˜a $™h¨1�_Ñ;ù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=×3FÑ3FÀwÔ3Oˆz× Ñ  ¨!Ó,¨QÒ/ÐU_ˆØ˜oÑ-ˆä/:¼4Ó/@ˆØ6@ò 	:Ñ2ˆB��S˜#˜s D¨!¨Q°Ø�q‹yØ˜SÑ!×(Ñ(¨¨s¯y©y©Õ9ð	:ð ˆÜ˜,Ó'ó 6	ˆCØ Ñ$ˆDÜ�D“	ˆAØ�AŠvØÜ�t“9˜q‘=ˆDÜ˜t›9¤c¨$£i�5ˆEÜÓ8°4Ô8Ó8¸1Ñ<ˆGØIJÈQÊœÓ;°dÔ;Ó;¸qÀ1¹uÒEÐTUˆJÜ69¸$³iÀ!²mÈÈAÊ”—‘˜:Ó&¬¨T«Ò2ÐSWˆBà×(Ñ(Ò/¨CˆDØ�3˜‘:ˆDØ�uŠ}Øò �AØÜœC¤ C¬¨S°1°u±9ÀÈÁÑ2OÓ)PÓ QÐTXÑ XÓYÐ[_ÐbcÑ[cÓdóàñô ñð
 ��Q‘à.3°eªm�U˜U‘] dÒ*ÀˆEØˆDØ4ò 
$�Ø %™i¨¨e˜�3�Ü% d›Oò ‘D�A�qØ�DØ˜1˜q™5‘L�CØ˜a“xØ  1™u˜ØHLÈqÊœs 3¬¨C°!°d±(¸dÑ1BÓ(CÔDÐVW˜Ø# e¨a¡iÑ/°$¸±,Ñ>˜Ùðð !$��q˜˜�W’ð
$ð �N‰Nà!Ø ×-Ñ-¨cÓ2ØØ $Ø Ø$Ø$Ø#*Ø&0Ø#'à $Ø$Ø$Ø04Ö 5¨1  Q£Ò 5ñ	ñöðG6	ôp �, Ó$ð 	¨Ü�I‰Ià'Ø'+Ø ŸY™YØ"×/Ñ/Ø'+×';Ñ';Ø$(×$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7ðV ñó ðó.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Ñ€
Ø˜A‰o€ó5ó
2ò
ò
ò
ò;òEò
%ð* ôO7�:ó O7ó ðO7ónóð ÷,ð ,ó ð,ð$ ÷ 1ð  1ó ð 1ðF ÷!ð !ó ð!÷IJò IJrB   