
    Foj                        S SK Jr  S SKJr  S SKJs  Jr  S SKJ	r
  S SKrS SKJr  SSKJrJrJr  \R"                  S\
R$                  S\
R$                  S	\
R$                  4S
 j5       r\R"                  S\
R$                  S\
R$                  S\
R$                  4S j5       rS rS r " S S\R.                  5      rS\\\4   S\\\\4   \\\R.                  4   4   4S jrg)    )Any)protonN)
MockTensor   )
exit_stateenter_stateCOMPUTE_METADATA_SCOPE_NAMEsize	metric_idmetric_value_sizec                    Sn[         R                  " U5      n[         R                  " X-   U5        US-   U-  n[         R                  " XV5      n[         R                  " SU5      n	[         R
                  " SU5       HB  n
XU-  -   nX:  n[         R                  " XK-   US9n[         R                  " XU-   U-  -   XS9  MD     [         R                  " 5         Xu-   U-  n[         R                  " X5        g )N   r   r   )mask)tlloadstorecdivarangerangedebug_barrier)
device_ptrdevice_offset_ptrr
   r   metric_value_ptrr   
BLOCK_SIZEdevice_offset	num_itersoffsetsicur_offsetsr   metric_values                 P/mnt/workspace/venv-train/lib/python3.13/site-packages/triton/profiler/metric.pytensor_metric_kernelr"   
   s      #JGG-.MHHZ'3"Q&$.M)6Iii:&GXXa#J..ww/=DI
{:dBBL\	 $
 "6$>MHH.    r    c                 
   [         R                  " U5      n[         R                  " X-   U5        US-   U-  n[         R                  " X-   U5        US-   U-  n[         R                  " 5         [         R                  " X5        g )Nr   )r   r   r   r   )r   r   r
   r   r    r   s         r!   scalar_metric_kernelr%      sj    GG-.MHHZ'3"Q&$.MHHZ'6"Q&$.MHH.r#   c                 `    U R                   " USSS.6nUR                  5         UR                  $ )N)r   r   )grid	num_warps)warmup_init_handlesfunction)	kernel_fnargskernels      r!   _get_kernelr/   (   s,    t%1=F
??r#   c            	      <   [        [        R                  5      n SnSnSn[        [        U U UUU U5      n[        [
        U U UUU5      n[        R                  R                  5       n[        R                  R                  U5      n[        R                  " XEU5        g )Nr   r   )r   r   uint64r/   r"   r%   driveractiveget_current_deviceget_current_stream	libprotonset_metric_kernels)mock_ptrmock_metric_id	mock_sizemock_metric_value_sizetensor_metric_kernel_fnscalar_metric_kernel_fndevicestreams           r!   r7   r7   .   s    "))$HNI) * ]]--/F]]--f5F  !8SYZr#   c                   (   ^  \ rS rSrU 4S jrSrU =r$ )_TensorMetricI   c                 j   > [         TU ]  UR                  5       X!R                  5       5        Xl        g )N)super__init__data_ptrnumel_value)selfvaluemetric_type_index	__class__s      r!   rE   _TensorMetric.__init__K   s%    )+<kkmLr#   )rH   )__name__
__module____qualname____firstlineno__rE   __static_attributes____classcell__)rL   s   @r!   rA   rA   I   s     r#   rA   metricsreturnc                 @   0 n0 nU R                  5        GH  u  p4[        US5      (       a  UR                  R                  S:X  a  XBU'   M7  [	        [
        5        UR                  5       (       aF  UR                  5       nUR                  5       S:  a  [        R                  nOV[        R                  nOEUR                  5       nUR                  5       S:  a  [        R                  nO[        R                  n[        5         [!        XE5      X'   M  XBU'   GM     X!4$ )NrF   cpur   )itemshasattrr>   typer   r	   is_floating_pointdoublerG   r6   metric_type_vector_double_indexmetric_type_double_indexlongmetric_type_vector_int64_indexmetric_type_int64_indexr   rA   )rT   tensor_metricsscalar_metricskeyrJ   metric_indexs         r!   transform_tensor_metricsrf   Q   s    N%'Nmmo
5*%%||  E)&+s#78**,,!LLNE{{}q('0'P'P'0'I'I!JJLE{{}q('0'O'O'0'H'H&3E&H#"'3- &. ))r#   )typingr   triton._C.libprotonr   r6   triton.runtime.driverruntimer2   triton.languagelanguager   tritonr   stater   r   r	   jitr1   r"   r%   r/   r7   TensorMetricrA   dictstrtuplerf    r#   r!   <module>ru      s     3 & &    G G /bii /TVT]T] /,.II/ /$ /bii /TVT]T] /momvmv / /[6I** *d38n *tCH~tTWYbYoYoToOp?p9q *r#   