Ë
    ©š¾j÷  ã                   óÎ  — d dl mZ d dlmZ d dlmc mZ d dlm	Z
 d dlZd dlmZ ej                  de
j                  de
j                  de
j                  fd„«       Zej                  de
j                  de
j                  d	e
j                  fd
„«       Zd„ Zd„ Z G d„ dej&                  «      ZdZi deej.                  ej0                  fi¥eD � ci c]  } d| › �eej.                  ej0                  f“Œ! c} ¥Zi e¥deej6                  ej8                  fi¥Zdededefd„Zde dedefd„Z!de dede"eef   fd„Z#de$e ef   de"e$e ef   e$e ej&                  f   f   fd„Z%yc c} w )é    )ÚAny)ÚprotonN)Ú
MockTensorÚsizeÚseq_idÚmetric_value_sizec                 ó¨  — d}|dz   }t        j                  ||d¬«      |z  }t        j                  | |z   |«       |dz   |z  }t        j                  ||«      }	t        j                  d|«      }
t        j
                  d|	«      D ]J  }|
||z  z   }||k  }t        j                  ||z   |¬«      }t        j                  | ||z   |z  z   ||¬«       ŒL y )Né€   é   Úrelaxed©Úsemr   )Úmask)ÚtlÚ
atomic_addÚstoreÚcdivÚarangeÚrangeÚload)Ú
device_ptrÚdevice_offset_ptrr   r   Úmetric_value_ptrr   Ú
BLOCK_SIZEÚrecord_sizeÚdevice_offsetÚ	num_itersÚoffsetsÚiÚcur_offsetsr   Úmetric_values                  ú`/home/spazio/kingdom_calls/work5/transenv/lib/python3.12/site-packages/triton/profiler/metric.pyÚtensor_metric_kernelr#   	   sÜ   € ð  #€JØ# aÑ'€Kô —M‘MÐ"3°[ÀiÔPÐSWÑW€MÜ‡H�HˆZ˜-Ñ'¨Ô0Ø" QÑ&¨$Ñ.€MÜ—‘Ð)¨:Ó6€IÜ�i‰i˜˜:Ó&€GÜ�X‰X�a˜Ó#ò ]ˆØ  J¡Ñ.ˆØÐ.Ñ.ˆÜ—w‘wÐ/°+Ñ=ÀDÔIˆÜ
�‰�˜}¨{Ñ:¸dÑBÑBÀLÐW[Ö\ñ	]ó    r!   c                 ó®   — t        j                  |dd¬«      |z  }t        j                  | |z   |«       |dz   |z  }t        j                  | |z   |«       y )Né   r   r   r   )r   r   r   )r   r   r   r   r!   r   s         r"   Úscalar_metric_kernelr'      sQ   € ô —M‘MÐ"3°Q¸IÔFÈÑM€MÜ‡H�HˆZ˜-Ñ'¨Ô0Ø" QÑ&¨$Ñ.€MÜ‡H�HˆZ˜-Ñ'¨Õ6r$   c                 ó\  —  | j                   |dddœŽ}|j                  «        t        |j                  dd «      }t        |dd «      }|€(t        j
                  j                  «       j                  }|j                  j                  |z  }|j                  ||j                  j                  fS )N)r   r   )ÚgridÚ	num_warpsÚtargetÚ	warp_size)ÚwarmupÚ_init_handlesÚgetattrÚmetadataÚdriverÚactiveÚget_current_targetr,   r*   ÚfunctionÚshared)Ú	kernel_fnÚargsÚkernelr+   r,   Únum_threadss         r"   Ú_get_kernelr:   &   s‘   € ØˆY×Ñ˜t¨%¸1Ò=€FØ
×ÑÔÜ�V—_‘_ h°Ó5€FÜ˜ ¨TÓ2€IØÐÜ—M‘M×4Ñ4Ó6×@Ñ@ˆ	Ø—/‘/×+Ñ+¨iÑ7€KØ�?‰?˜K¨¯©×)?Ñ)?Ð?Ð?r$   c            	      óT  — t        t        j                  «      } d}d}d}t        t        | | ||| |«      \  }}}t        t
        | | |||«      \  }}}	t        j                  j                  «       }
t        j                  j                  |
«      }t        j                  |||||||	«       y )Nr   r   )r   r   Úuint64r:   r#   r'   r1   r2   Úget_current_deviceÚget_current_streamÚ	libprotonÚset_metric_kernels)Úmock_ptrÚmock_seq_idÚ	mock_sizeÚmock_metric_value_sizeÚtensor_metric_kernel_fnÚ tensor_metric_kernel_num_threadsÚtensor_metric_kernel_sharedÚscalar_metric_kernel_fnÚ scalar_metric_kernel_num_threadsÚscalar_metric_kernel_sharedÚdeviceÚstreams               r"   r@   r@   1   sÀ   € Üœ"Ÿ)™)Ó$€HØ€KØ€IØÐÜ]hÜØØØØØØó^ÑZÐÐ=Ð?Zô ^iÜØØØØØó^ÑZÐÐ=Ð?Zô �]‰]×-Ñ-Ó/€FÜ�]‰]×-Ñ-¨fÓ5€FÜ× Ñ ØØØØ(Ø#Ø(Ø#õr$   c                   ó   ‡ — e Zd Zˆ fd„Zˆ xZS )Ú_TensorMetricc                 ón   •— t         ‰| �  |j                  «       ||j                  «       «       || _        y ©N)ÚsuperÚ__init__Údata_ptrÚnumelÚ_value)ÚselfÚvalueÚmetric_type_indexÚ	__class__s      €r"   rR   z_TensorMetric.__init__V   s)   ø€ Ü‰Ñ˜Ÿ™Ó)Ð+<¸e¿k¹k»mÔLàˆ�r$   )Ú__name__Ú
__module__Ú__qualname__rR   Ú__classcell__)rY   s   @r"   rN   rN   T   s   ø„ ÷ð r$   rN   )é   é   é    é@   ÚflopsÚbytesrW   ÚconvertÚreturnc                 óò   — t        | d«      r3| j                  «       dkD  r| j                  «       n| j                  «       } t	        | t
        t        f«      r| D �cg c]
  } ||«      ‘Œ c}S  || «      S c c}w )NrS   r   )ÚhasattrrT   ÚtolistÚitemÚ
isinstanceÚlistÚtuple)rW   rd   Úvs      r"   Ú_scalar_or_vector_valuern   j   s^   € Üˆu�jÔ!Ø"'§+¡+£-°!Ò"3�—‘”¸¿¹»ˆÜ�%œ$¤˜Ô'Ø$)Ö*˜q‘˜•
Ò*Ð*Ù�5‹>Ðùò +s   ÁA4Úkeyc                 óH   — | t         v rt         |    \  }}}t        ||«      S |S rP   )ÚROOFLINE_DTYPESrn   )ro   rW   rd   Ú_s       r"   Ú_scalar_metric_valuers   r   s.   € ð ŒoÑÜ'¨Ñ,‰ˆ��AÜ& u¨gÓ6Ð6Ø€Lr$   c                 óæ  — | t         v rLt         |    \  }}}| dk(  r|j                  «       n|j                  «       }||j                  «       dkD  r|fS |fS |j	                  «       rG|j                  «       }|j                  «       dkD  r|t
        j                  fS |t
        j                  fS |j                  «       }|j                  «       dkD  r|t
        j                  fS |t
        j                  fS )Nrc   r   )
rq   ÚlongÚdoublerT   Úis_floating_pointr?   Úmetric_type_vector_double_indexÚmetric_type_double_indexÚmetric_type_vector_int64_indexÚmetric_type_int64_index)ro   rW   rr   Úscalar_indexÚvector_indexs        r"   Ú_tensor_metric_value_and_indexr~   {   sØ   € Ø
ŒoÑÜ(7¸Ñ(<Ñ%ˆˆ<˜Ø # w¢�—
‘
”°E·L±L³NˆØ e§k¡k£m°aÒ&7�lÐIÐI¸\ÐIÐIð ×ÑÔ Ø—‘“ˆØ�;‰;‹=˜1ÒØœ)×CÑCÐCÐCØ”i×8Ñ8Ð8Ð8à�J‰J‹L€EØ‡{�{ƒ}�qÒØ”i×>Ñ>Ð>Ð>Ø”)×3Ñ3Ð3Ð3r$   Úmetricsc                 ó  — i }i }| j                  «       D ]h  \  }}t        |d«      rH|j                  j                  dk(  rt	        ||«      ||<   Œ;t        ||«      \  }}t        ||«      ||<   ŒZt	        ||«      ||<   Œj ||fS )NrS   Úcpu)Úitemsrg   rK   Útypers   r~   rN   )r   Útensor_metricsÚscalar_metricsro   rW   Úmetric_indexs         r"   Útransform_tensor_metricsr‡   Ž   s”   € Ø€NØ%'€NØ—m‘m“oò C‰
ˆˆUÜ�5˜*Ô%Ø�|‰|× Ñ  EÒ)Ü&:¸3ÀÓ&F�˜sÒ#ä&DÀSÈ%Ó&PÑ#��|Ü&3°E¸<Ó&H�˜sÒ#ä"6°s¸EÓ"BˆN˜3ÒðCð ˜>Ð)Ð)r$   )&Útypingr   Útriton._C.libprotonr   r?   Útriton.runtime.driverÚruntimer1   Útriton.languageÚlanguager   Útritonr   Újitr<   r#   r'   r:   r@   ÚTensorMetricrN   ÚFLOPS_WIDTHSÚfloatry   rx   ÚFLOPS_DTYPESÚintr{   rz   rq   rƒ   rn   Ústrrs   rl   r~   Údictr‡   )Úwidths   0r"   ú<module>r˜      sé  ðÝ Ý 3ß &Ð &Ý Û Ý ð ‡�ð]¸b¿i¹ið ]ÐQS×QZÑQZð ]Ø,.¯I©Iò]ó ð]ð& ‡�ð7¸b¿i¹ið 7ÐQS×QZÑQZð 7Ðjl×jsÑjsò 7ó ð7ò@ò ôF�I×*Ñ*ô ð €ðØ�˜	×:Ñ:¸I×<eÑ<eÐfÐgðð "öàð �ˆwˆ˜% ×!CÑ!CÀY×EnÑEnÐoÑoòð€ðØðàˆc�9×4Ñ4°i×6^Ñ6^Ð_ñ€ð 3ð °ð ¸#ó ð˜cð ¨#ð °#ó ð4¨ð 4°Cð 4¸EÀ#ÀsÀ(¹Oó 4ð&* d¨3°¨8¡nð *¸¸tÀCÈÀH¹~ÈtÐTWÐYb×YoÑYoÐToÑOpÐ?pÑ9qô *ùò_s   Ã$E"