
    >̄i1                    d    d dl mZ ddlmZ d ZddZd Zd Zd Z G d	 d
      Z	 G d d      Z
y)    )annotations   )driverc                    t        |      dkD  rt        |j                               dz   nd}|d   t        |d         D ]  \  }}d|||z   <    ||fS )Nr      ids_of_tensormapsz*CUtensorMap)lenmaxkeys	enumerate)	constants	signatureidsnum_regular_signaturesi_s         _/home/obispo/Crisostomo_bridge/mision_env/lib/python3.12/site-packages/triton/compiler/utils.pygenerate_cu_signaturer      so    :=i.1:LS!12Q6RS
+c"567 	CDAq4BI,q01	C,,,    c                `    g }t        |       D ]  }|j                  t        d              |S )NT)dummy)rangeappendInfoFromBackendForTensorMap)nretr   s      r   dummy_tensormaps_infor   $   s1    
C1X <

.T:;<Jr   c                `    g }| D ]&  }t        |      }||_        |j                  |       ( |S )N)infos)r   ids_of_folded_argsr   )r   r    r   infoes        r   parse_tma_infor#   +   s:    
C 'd31

1 Jr   c                z    i }| 4t        |       D ]$  \  }}|j                  |j                                & |S d }|S N)r   updateget_address_tma_mapping)tensormaps_infor   r   r"   s       r   get_tma_mappingr)   4   sM    
C"o. 	4DAqJJq0023	4 J Jr   c                R    d }| | D cg c]  }|j                          }}|S c c}w r%   )get_id_of_tensormap)r(   r   r"   s      r   get_ids_of_tensormapsr,   >   s4    
C"0?@1q$$&@@J As   $c                      e Zd ZdZdZdZddZd Zd ZddZ	d Z
d	 Zd
 Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zy)r   r   r   Nc                    || _         d| _        |s"t        |t              s| j	                  |       y |s"t        |t              r| j                  |       y |r| j                          y y )N )r   r    
isinstancedict_extract_info_from_backend_extract_info_from_dict_dummy)selfr   r   s      r   __init__z$InfoFromBackendForTensorMap.__init__N   sV    
"$Zt4++E2:eT2((/KKM r   c                   t         j                  t         j                  k  sJ t         j                  dk(  rt        j                  j
                  d   | _        d| _        d| _        g d| _	        g d| _
        g d| _        g d| _        t        j                  j                  d   | _        t        j                  j                  d	   | _        t        j                  j"                  d
   | _        d| _        t        j                  j(                  d   | _        t         xj                  dz  c_        y t         j                  dk(  rt        j                  j
                  d   | _        d| _        d| _        g d| _	        g d| _
        g d| _        g d| _        t        j                  j                  d   | _        t        j                  j                  d	   | _        t        j                  j"                  d
   | _        d| _        t        j                  j(                  d   | _        t         xj                  dz  c_        y y )Nr   CU_TENSOR_MAP_DATA_TYPE_FLOAT16   )      r<   )      r<   r<   )   @   r   r   )r   r   r   r   CU_TENSOR_MAP_INTERLEAVE_NONECU_TENSOR_MAP_SWIZZLE_32B"CU_TENSOR_MAP_L2_PROMOTION_L2_128B   !CU_TENSOR_MAP_FLOAT_OOB_FILL_NONEr      )r   r   Nr   utilsCUtensorMapDataTypetensorDataType
tensorRankglobalAddressArgIdxglobalStridesArgIdxglobalDimsArgIdxboxDimselementStridesCUtensorMapInterleave
interleaveCUtensorMapSwizzleswizzleCUtensorMapL2promotionl2PromotionTMADescArgIdxCUtensorMapFloatOOBfilloobFillr5   s    r   r4   z"InfoFromBackendForTensorMap._dummyX   s   *,,/J/L/LLLL&((A-"(,,"B"BCd"eDDO'(D$'5D$$2D!)DL".D$ll@@A`aDO!<<::;VWDL%||BBCghD!#D!<<??@cdDL'))Q.)&((A-"(,,"B"BCd"eDDO'(D$'5D$$2D!)DL".D$ll@@A`aDO!<<::;VWDL%||BBCghD!#D!<<??@cdDL'))Q.) .r   c                   |j                   | _         |j                  | _        |j                  | _        |j                  | _        |j                  | _        |j
                  | _        |j                  | _        |j                  | _        |j                  | _        |j                  | _	        |j                  | _
        |j                  | _        y r%   rJ   rK   rL   rM   rN   rO   rP   rR   rT   rV   rY   rW   r5   r   s     r   r2   z6InfoFromBackendForTensorMap._extract_info_from_backendy   s    #22**#(#<#< #(#<#<  % 6 6}}#22**}} ,,}}"00r   c                    |d   | _         |d   | _        |d   | _        |d   | _        |d   | _        |d   | _        |d   | _        |d   | _        |d	   | _        |d
   | _	        |d   | _
        |d   | _        y )NrJ   rK   rL   rM   rN   rO   rP   rR   rT   rV   rY   rW   r\   r]   s     r   r3   z3InfoFromBackendForTensorMap._extract_info_from_dict   s    #$45-#()>#? #()>#?  %&8 9Y'#$45-Y' /Y'"?3r   c                ^    | j                   | j                  t        | j                        z   iS r%   )rL   rW   r	   r    rZ   s    r   r'   z3InfoFromBackendForTensorMap.get_address_tma_mapping   s)    (($*<*<s4CZCZ?[*[\\r   c                F    | j                   t        | j                        z   S r%   )rW   r	   r    rZ   s    r   r+   z/InfoFromBackendForTensorMap.get_id_of_tensormap   s    !!C(?(?$@@@r   c                    | j                   S r%   )rW   rZ   s    r   getTMADescArgIdxz,InfoFromBackendForTensorMap.getTMADescArgIdx   s    !!!r   c                   t         j                  j                  d   dt         j                  j                  d   dt         j                  j                  d   dt         j                  j                  d   dt         j                  j                  d   d	t         j                  j                  d
   d	t         j                  j                  d   dt         j                  j                  d   dt         j                  j                  d   d	t         j                  j                  d   dt         j                  j                  d   dt         j                  j                  d   dt         j                  j                  d   di|   S )NCU_TENSOR_MAP_DATA_TYPE_UINT8r   CU_TENSOR_MAP_DATA_TYPE_UINT16r   CU_TENSOR_MAP_DATA_TYPE_UINT32r9   CU_TENSOR_MAP_DATA_TYPE_INT32CU_TENSOR_MAP_DATA_TYPE_UINT64   CU_TENSOR_MAP_DATA_TYPE_INT64r8   CU_TENSOR_MAP_DATA_TYPE_FLOAT32CU_TENSOR_MAP_DATA_TYPE_FLOAT64 CU_TENSOR_MAP_DATA_TYPE_BFLOAT16#CU_TENSOR_MAP_DATA_TYPE_FLOAT32_FTZ CU_TENSOR_MAP_DATA_TYPE_TFLOAT32$CU_TENSOR_MAP_DATA_TYPE_TFLOAT32_FTZ)r   rH   rI   )r5   dtypes     r   bytes_from_typez+InfoFromBackendForTensorMap.bytes_from_type   s1   LL,,-LMqLL,,-MNPQLL,,-MNPQLL,,-LMqLL,,-MNPQLL,,-LMqLL,,-NOQRLL,,-NOQRLL,,-NOQRLL,,-OPRSLL,,-RSUVLL,,-OPRSLL,,-STVW
  	r   c                    | j                   S r%   )rJ   rZ   s    r   getTensorMapDataTypez0InfoFromBackendForTensorMap.getTensorMapDataType       """r   c                    | j                   S r%   )rR   rZ   s    r   getInterleavez)InfoFromBackendForTensorMap.getInterleave       r   c                    | j                   S r%   )rT   rZ   s    r   
getSwizzlez&InfoFromBackendForTensorMap.getSwizzle       ||r   c                    | j                   S r%   )rV   rZ   s    r   getL2Promotionz*InfoFromBackendForTensorMap.getL2Promotion   s    r   c                    | j                   S r%   )rY   rZ   s    r   
getOobFillz&InfoFromBackendForTensorMap.getOobFill   r{   r   c                    | j                   S r%   )rK   rZ   s    r   getTensorRankz)InfoFromBackendForTensorMap.getTensorRank   rx   r   c                    | j                   S r%   )rO   rZ   s    r   
getBoxDimsz&InfoFromBackendForTensorMap.getBoxDims   r{   r   c                    | j                   S r%   )rP   rZ   s    r   getElementStridesz-InfoFromBackendForTensorMap.getElementStrides   ru   r   c                D    | j                  | j                  |      }||   S r%   )getOriginArgIdxrL   )r5   argsidxs      r   getGlobalAddressz,InfoFromBackendForTensorMap.getGlobalAddress   s$    ""4#;#;TBCyr   c                    g }| j                   D ]E  }d}|dk(  rd}n(|dk  r|dk7  r| dz
  }n| j                  ||      }||   }|j                  |       G |S Nr   r<   r   )rN   r   r   )r5   r   shaper"   tr   s         r   getGlobalDimsz)InfoFromBackendForTensorMap.getGlobalDims   st    && 	AA BwQ17BF**1d3ILLO	 r   c                   | j                  |      D cg c]  }t        |       }}| j                  j                         }g }t	        | j
                        D ]a  }d}||   dk(  rt	        |      D ]
  }|||   z  } n+||   dk  r	d||   z
  }n| j                  ||   |      }	||	   }|j                  |       c |dd  }|D cg c]   }|| j                  | j                        z  " }
}|
S c c}w c c}w r   )
r   intrM   copyr   rK   r   r   rr   rJ   )r5   r   r"   t_globalDimst_globalStridesArgIdxstrides_in_elementsr   r   iinew_idxstrides_in_bytess              r   getGlobalStridesz,InfoFromBackendForTensorMap.getGlobalStrides   s$   (,(:(:4(@A1AAA $ 8 8 = = ? t' 	*AA$Q'2-( *Bb))A* 'q)A-.q11../DQ/GNM&&q)	* 2!"5SfgaA 4 4T5H5H IIgg) B& hs   C/%C4c                    | j                   r5t        t        |            D cg c]  }|| j                   vs| }}||   S |S c c}w r%   )r    r   r	   )r5   r   r   r   ids_before_folding_args        r   r   z+InfoFromBackendForTensorMap.getOriginArgIdx   sK    ""16s4y1A%fAQdNeNeEea%f"%f)#..J &gs
   AAc                   t         j                  j                  | j                         | j	                         | j                  |      | j                  |      | j                  |      | j                         | j                         | j                         | j                         | j                         | j                               S r%   )r   rH   cuTensorMapEncodeTiledrt   r   r   r   r   r   r   rw   rz   r}   r   )r5   r   s     r   	tensormapz%InfoFromBackendForTensorMap.tensormap   s    ||22%%' !!$'t$!!$'OO""$ OO!OO
 	
r   c                h   t        | j                  | j                  t        | j                        t        | j
                        | j                  | j                  t        | j                        t        | j                        | j                  | j                  | j                  | j                  f      S r%   )hashr    rL   tuplerN   rM   rJ   rK   rO   rP   rR   rT   rV   rY   rZ   s    r   __hash__z$InfoFromBackendForTensorMap.__hash__  s    T,,d.F.FdNcNcHd4334d6I6I4??\abfbnbn\o4../$,,PTP`P`bfbnbnp q 	qr   c                J   t        || j                        sy| j                  | j                  | j                  | j
                  | j                  | j                  | j                  | j                  | j                  | j                  | j                  | j                  f|j                  |j                  |j                  |j
                  |j                  |j                  |j                  |j                  |j                  |j                  |j                  |j                  fk(  S NF)r0   	__class__r    rL   rN   rM   rJ   rK   rO   rP   rR   rT   rV   rY   )r5   others     r   __eq__z"InfoFromBackendForTensorMap.__eq__  s    %0'')A)A4CXCXZ^ZrZr##T__dllDDWDWY]YhYhjnjvjv   #(":":E<U<UW\WmWm"'";";U=Q=QSXScScejerer"'"6"68H8H%--Y^YjYj"'--"11 	1r   r   )r   r1   )__name__
__module____qualname__rG   r   ntmar6   r4   r2   r3   r'   r+   rb   rr   rt   rw   rz   r}   r   r   r   r   r   r   r   r   r   r   r   r/   r   r   r   r   I   s    	A	ADB14]A""# #
  .
 q
	1r   r   c                       e Zd Zd ZddZd Zy)TensorMapManagerc                    i | _         y r%   )tensormaps_devicerZ   s    r   r6   zTensorMapManager.__init__  s
    !#r   c                L   || j                   v rt        | j                   |         S |\  }}|j                  |      }d}t        j                  j                  |      }t        j                  j                  |||       || j                   |<   t        | j                   |         S )N   )r   r   r   r   rH   
cuMemAlloccuMemcpyHtoD)r5   keyr"   r   t_tensormapTENSORMAP_SIZE_IN_BYTESt_tensormap_devices          r   __getitem__zTensorMapManager.__getitem__"  s    $(((t--c233IQ++d+K&)#!'!8!89P!QLL%%&8+G^_*<D""3't--c233r   c                    | j                   j                         D ]$  \  }}t        j                  j	                  |       & y r%   )r   itemsr   rH   	cuMemFree)r5   r   vs      r   __del__zTensorMapManager.__del__.  s6    **002 	&DAqLL""1%	&r   N)r   r   )r   r   r   r6   r   r   r/   r   r   r   r     s    $
4&r   r   N)r   )
__future__r   runtimer   r   r   r#   r)   r,   r   r   r/   r   r   <module>r      s=   , # -Q1 Q1h& &r   