a
    öDf.                     @  sT   d dl mZ ddlmZ dd ZdddZdd	 Zd
d Zdd ZG dd dZ	dS )    )annotations   )driverc                 C  sV   t |dkrt| d nd}|d d urNt|d D ]\}}d||| < q8||fS )Nr      Zids_of_tensormapsz*CUtensorMap)lenmaxkeys	enumerate)	constants	signatureidsZnum_regular_signaturesi_ r   b/nfs/NAS7/SABIOD/METHODE/ermites/ermites_venv/lib/python3.9/site-packages/triton/compiler/utils.pygenerate_cu_signature   s
     r   c                 C  s&   g }t | D ]}|tdd q|S )NT)dummy)rangeappendInfoFromBackendForTensorMap)nretr   r   r   r   dummy_tensormaps_info$   s    r   c                 C  s,   g }| D ]}t |d}||_|| q|S )Ninfos)r   ids_of_folded_argsr   )r   r   r   infoer   r   r   parse_tma_info+   s    
r   c                 C  s6   i }| d ur.t | D ]\}}||  qnd }|S N)r	   updateget_address_tma_mapping)tensormaps_infor   r   r   r   r   r   get_tma_mapping4   s    r#   c                 C  s   d }| d urdd | D }|S )Nc                 S  s   g | ]}|  qS r   )get_id_of_tensormap.0r   r   r   r   
<listcomp>B       z)get_ids_of_tensormaps.<locals>.<listcomp>r   )r"   r   r   r   r   get_ids_of_tensormaps>   s    r)   c                   @  s   e Zd ZdZdZdZd5ddZdd Zd	d
 ZddddZ	dd Z
dd Zdd Zdd Zdd Zdd Zdd Zdd Zdd  Zd!d" Zd#d$ Zd%d& Zd'd( Zd)d* Zd+d, Zd-d. Zd/d0 Zd1d2 Zd3d4 ZdS )6r   r   r   NFc                 C  sP   || _ d| _|s&t|ts&| | n&|s@t|tr@| | n|rL|   d S )Nr   )r   r   
isinstancedict_extract_info_from_backend_extract_info_from_dict_dummy)selfr   r   r   r   r   __init__N   s    z$InfoFromBackendForTensorMap.__init__c                 C  sN  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  jd7  _d S t jdkrJt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  jd7  _d S d S )Nr   CU_TENSOR_MAP_DATA_TYPE_FLOAT16   )      r5   )      r5   r5   )   @   r   r   )r   r   r   r   ZCU_TENSOR_MAP_INTERLEAVE_NONEZCU_TENSOR_MAP_SWIZZLE_32BZ"CU_TENSOR_MAP_L2_PROMOTION_L2_128B   Z!CU_TENSOR_MAP_FLOAT_OOB_FILL_NONEr      )r   r   Nr   utilsCUtensorMapDataTypetensorDataType
tensorRankglobalAddressArgIdxglobalStridesArgIdxglobalDimsArgIdxboxDimselementStridesZCUtensorMapInterleave
interleaveZCUtensorMapSwizzleswizzleZCUtensorMapL2promotionl2PromotionTMADescArgIdxZCUtensorMapFloatOOBfilloobFillr/   r   r   r   r.   X   s>    








z"InfoFromBackendForTensorMap._dummyc                 C  sd   |j | _ |j| _|j| _|j| _|j| _|j| _|j| _|j| _|j| _|j	| _	|j
| _
|j| _d S r   r?   r@   rA   rB   rC   rD   rE   rF   rG   rH   rJ   rI   r/   r   r   r   r   r,   y   s    z6InfoFromBackendForTensorMap._extract_info_from_backendr+   r   c                 C  s|   |d | _ |d | _|d | _|d | _|d | _|d | _|d | _|d | _|d	 | _|d
 | _	|d | _
|d | _d S )Nr?   r@   rA   rB   rC   rD   rE   rF   rG   rH   rJ   rI   rL   rM   r   r   r   r-      s    










z3InfoFromBackendForTensorMap._extract_info_from_dictc                 C  s   | j | jt| j iS r   )rA   rI   r   r   rK   r   r   r   r!      s    z3InfoFromBackendForTensorMap.get_address_tma_mappingc                 C  s   | j t| j S r   )rI   r   r   rK   r   r   r   r$      s    z/InfoFromBackendForTensorMap.get_id_of_tensormapc                 C  s   | j S r   )rI   rK   r   r   r   getTMADescArgIdx   s    z,InfoFromBackendForTensorMap.getTMADescArgIdxc                 C  s   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 )NZCU_TENSOR_MAP_DATA_TYPE_UINT8r   ZCU_TENSOR_MAP_DATA_TYPE_UINT16r   ZCU_TENSOR_MAP_DATA_TYPE_UINT32r2   ZCU_TENSOR_MAP_DATA_TYPE_INT32ZCU_TENSOR_MAP_DATA_TYPE_UINT64   ZCU_TENSOR_MAP_DATA_TYPE_INT64r1   ZCU_TENSOR_MAP_DATA_TYPE_FLOAT32ZCU_TENSOR_MAP_DATA_TYPE_FLOAT64Z CU_TENSOR_MAP_DATA_TYPE_BFLOAT16Z#CU_TENSOR_MAP_DATA_TYPE_FLOAT32_FTZZ CU_TENSOR_MAP_DATA_TYPE_TFLOAT32Z$CU_TENSOR_MAP_DATA_TYPE_TFLOAT32_FTZ)r   r=   r>   )r/   Zdtyper   r   r   bytes_from_type   s     z+InfoFromBackendForTensorMap.bytes_from_typec                 C  s   | j S r   )r?   rK   r   r   r   getTensorMapDataType   s    z0InfoFromBackendForTensorMap.getTensorMapDataTypec                 C  s   | j S r   )rF   rK   r   r   r   getInterleave   s    z)InfoFromBackendForTensorMap.getInterleavec                 C  s   | j S r   )rG   rK   r   r   r   
getSwizzle   s    z&InfoFromBackendForTensorMap.getSwizzlec                 C  s   | j S r   )rH   rK   r   r   r   getL2Promotion   s    z*InfoFromBackendForTensorMap.getL2Promotionc                 C  s   | j S r   )rJ   rK   r   r   r   
getOobFill   s    z&InfoFromBackendForTensorMap.getOobFillc                 C  s   | j S r   )r@   rK   r   r   r   getTensorRank   s    z)InfoFromBackendForTensorMap.getTensorRankc                 C  s   | j S r   )rD   rK   r   r   r   
getBoxDims   s    z&InfoFromBackendForTensorMap.getBoxDimsc                 C  s   | j S r   )rE   rK   r   r   r   getElementStrides   s    z-InfoFromBackendForTensorMap.getElementStridesc                 C  s   |  | j|}|| S r   )getOriginArgIdxrA   )r/   argsidxr   r   r   getGlobalAddress   s    z,InfoFromBackendForTensorMap.getGlobalAddressc                 C  s`   g }| j D ]P}d}|dkr d}n0|dk r<|dkr<| d }n| ||}|| }|| q
|S )Nr   r5   r   )rC   rY   r   )r/   rZ   shaper   tr[   r   r   r   getGlobalDims   s    
z)InfoFromBackendForTensorMap.getGlobalDimsc           
        s   dd   |D } j }g }t jD ]l}d}|| dkr\t|D ]}||| 9 }qHn2|| dk rvd||  }n || |}|| }|| q,|dd  } fdd|D }	|	S )Nc                 S  s   g | ]}t |qS r   )intr%   r   r   r   r'      r(   z@InfoFromBackendForTensorMap.getGlobalStrides.<locals>.<listcomp>r   r5   r   c                   s   g | ]}|   j qS r   )rP   r?   r%   rK   r   r   r'      r(   )r_   rB   copyr   r@   rY   r   )
r/   rZ   Zt_globalDimsZt_globalStridesArgIdxZstrides_in_elementsr   r^   iiZnew_idxZstrides_in_bytesr   rK   r   getGlobalStrides   s     
z,InfoFromBackendForTensorMap.getGlobalStridesc                   s0    j r( fddtt|D }|| S |S d S )Nc                   s   g | ]}| j vr|qS r   )r   )r&   r   rK   r   r   r'      r(   z?InfoFromBackendForTensorMap.getOriginArgIdx.<locals>.<listcomp>)r   r   r   )r/   r[   rZ   Zids_before_folding_argr   rK   r   rY      s    z+InfoFromBackendForTensorMap.getOriginArgIdxc                 C  sR   t j|  |  | || || ||  | 	 | 
 |  |  |  S r   )r   r=   ZcuTensorMapEncodeTiledrQ   rV   r\   r_   rc   rW   rX   rR   rS   rT   rU   )r/   rZ   r   r   r   	tensormap   s    z%InfoFromBackendForTensorMap.tensormapc                 C  sH   t | j| jt| jt| j| j| jt| jt| j	| j
| j| j| jfS r   )hashr   rA   tuplerC   rB   r?   r@   rD   rE   rF   rG   rH   rJ   rK   r   r   r   __hash__  s    z$InfoFromBackendForTensorMap.__hash__c                 C  sx   t || jsdS | 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)r*   	__class__r   rA   rC   rB   r?   r@   rD   rE   rF   rG   rH   rJ   )r/   otherr   r   r   __eq__  s    z"InfoFromBackendForTensorMap.__eq__)NF)__name__
__module____qualname__r<   r   Zntmar0   r.   r,   r-   r!   r$   rN   rP   rQ   rR   rS   rT   rU   rV   rW   rX   r\   r_   rc   rY   rd   rg   rj   r   r   r   r   r   I   s4   

!r   N)r   )

__future__r   Zruntimer   r   r   r   r#   r)   r   r   r   r   r   <module>   s   	
	
