a
    öDf!                     @   s   d dl Z d dlZd dlZd dlZd dlmZ d dlZddlmZ ddl	m
Z
 ddlmZ G dd	 d	e jd
ZG dd deZG dd dZG dd deZG dd deZG dd deZG dd deZG dd dZdd ZeeZdS )    N)Path   )_build   )get_cache_manager)driverc                   @   s.   e Zd ZdZdZedd ZddddZdS )	
DriverBaser   r   c                   C   s    t jt jt jtddS )Nz..Zthird_party)ospathjoindirnameabspath__file__ r   r   b/nfs/NAS7/SABIOD/METHODE/ermites/ermites_venv/lib/python3.9/site-packages/triton/runtime/driver.pythird_party_dir   s    zDriverBase.third_party_dirN)returnc                 C   s   d S Nr   selfr   r   r   __init__   s    zDriverBase.__init__)__name__
__module____qualname__CUDAHIPstaticmethodr   r   r   r   r   r   r      s
   
r   )	metaclassc                       s$   e Zd Z fddZdd Z  ZS )	CudaUtilsc                    s"   t | dstt| | | _| jS Ninstance)hasattrsuperr   __new__r    cls	__class__r   r   r#       s    
zCudaUtils.__new__c              	   C   s  t jt jt}tt j|dd }t	|
d }t|}d}||}|d u rt }t j|d}t|d}	|	| W d    n1 s0    Y  td||}
t|
d$}	|j|	 |d	d
}W d    n1 s0    Y  W d    n1 s0    Y  dd l}|jd|}|j|}|j| |j| _|j| _|j| _|j| _|j| _|j | _ |j!| _!|j"| _"|j#| _#|j$| _$|j%| _%|j&| _&d S )Nbackendszcuda.cutf-8zcuda_utils.somain.cwZ
cuda_utilsrbTbinaryr   )'r	   r
   r   realpathr   r   r   	read_texthashlibmd5encode	hexdigestr   get_filetempfileTemporaryDirectoryopenwriter   putreadimportlib.utilutilspec_from_file_locationmodule_from_specloaderexec_moduleload_binaryget_device_propertiesZCUtensorMapDataTypeZCUtensorMapInterleaveZCUtensorMapSwizzleZCUtensorMapL2promotionZCUtensorMapFloatOOBfillZcuTensorMapEncodeTiled
cuMemAlloccuMemcpyHtoD	cuMemFreeZcuOccupancyMaxActiveClustersr   r   srckeycachefname
cache_pathtmpdirZsrc_pathfso	importlibspecmodr   r   r   r   %   s<    


(RzCudaUtils.__init__r   r   r   r#   r   __classcell__r   r   r&   r   r      s   r   c                   @   s*   e Zd Zdd ZedddZdd ZdS )	TensorMapManagerc                 C   s
   i | _ d S r   )tensormaps_devicer   r   r   r   r   I   s    zTensorMapManager.__init__)rI   c                 C   sf   || j v rt| j | S |\}}||}d}tj|}tj||| || j |< t| j | S d S )N   )rV   intZ	tensormapr   utilsrD   rE   )r   rI   eargsZt_tensormapZTENSORMAP_SIZE_IN_BYTESZt_tensormap_devicer   r   r   __getitem__L   s    


zTensorMapManager.__getitem__c                 C   s$   | j  D ]\}}tj| q
d S r   )rV   itemsr   rY   rF   )r   _vr   r   r   __del__X   s    zTensorMapManager.__del__N)r   r   r   r   tupler\   r`   r   r   r   r   rU   G   s   rU   c                       sB   e Zd Ze Z fddZdd Ze dd Z	dd Z
  ZS )	
CudaDriverc                    s"   t | dstt| | | _| jS r   )r!   r"   rb   r#   r    r$   r&   r   r   r#   `   s    
zCudaDriver.__new__c                    sx   t  | _| j| _d| _dd l  jj| _zddlm	} || _
W n  ty^    fdd| _
Y n0  jj| _ jj| _d S )NZcubinr   )_cuda_getCurrentRawStreamc                    s    j | jS r   )cudaZcurrent_streamZcuda_stream)idxtorchr   r   <lambda>p       z%CudaDriver.__init__.<locals>.<lambda>)r   rY   r   backendZ
binary_extrg   rd   get_device_capabilityZtorch._Crc   Zget_current_streamImportErrorZcurrent_deviceget_current_deviceZ
set_deviceZset_current_device)r   rc   r   rf   r   r   e   s    


zCudaDriver.__init__c                 C   s.   |   }| |}|d d |d  }d|fS )Nr   
   r   rd   )rm   rk   )r   ZdeviceZ
capabilityr   r   r   get_current_targett   s    
zCudaDriver.get_current_targetc                 C   sL   t |}|d urHtdd |D }t|D ]\}}|tj||f  q*|S )Nc                 S   s"   g | ]}t |d r| n|qS )data_ptr)r!   rp   ).0argr   r   r   
<listcomp>   ri   z8CudaDriver.assemble_tensormap_to_arg.<locals>.<listcomp>)listra   	enumerateappendrb   tensormap_manager)r   Ztensormaps_infor[   Zargs_with_tmaZargs_ptrirZ   r   r   r   assemble_tensormap_to_arg{   s    z$CudaDriver.assemble_tensormap_to_arg)r   r   r   rU   rw   r#   r   	functools	lru_cachero   ry   rT   r   r   r&   r   rb   ]   s   
rb   c                       s$   e Zd Z fddZdd Z  ZS )HIPUtilsc                    s"   t | dstt| | | _| jS r   )r!   r"   r|   r#   r    r$   r&   r   r   r#      s    
zHIPUtils.__new__c              	   C   sX  t jt jt}tt j|dd }t	|
d }t|}d}||}|d u rt }t j|d}t|d}	|	| W d    n1 s0    Y  td||}
t|
d$}	|j|	 |d	d
}W d    n1 s0    Y  W d    n1 s0    Y  dd l}|jd|}|j|}|j| |j| _|j| _d S )Nr(   zhip.cr)   zhip_utils.sor*   r+   Z	hip_utilsr,   Tr-   r   )r	   r
   r   r/   r   r   r   r0   r1   r2   r3   r4   r   r5   r6   r7   r8   r9   r   r:   r;   r<   r=   r>   r?   r@   rA   rB   rC   rG   r   r   r   r      s(    


(RzHIPUtils.__init__rS   r   r   r&   r   r|      s   r|   c                       s$   e Zd Z fddZdd Z  ZS )	HIPDriverc                    s"   t | dstt| | | _| jS r   )r!   r"   r}   r#   r    r$   r&   r   r   r#      s    
zHIPDriver.__new__c                 C   s   t  | _| j| _d S r   )r|   rY   r   rj   r   r   r   r   r      s    zHIPDriver.__init__rS   r   r   r&   r   r}      s   r}   c                       s$   e Zd Z fddZdd Z  ZS )UnsupportedDriverc                    s"   t | dstt| | | _| jS r   )r!   r"   r~   r#   r    r$   r&   r   r   r#      s    
zUnsupportedDriver.__new__c                 C   s   d | _ d | _d S r   )rY   rj   r   r   r   r   r      s    zUnsupportedDriver.__init__rS   r   r   r&   r   r~      s   r~   c                       sL   e Zd Zdd Zdd Zdd Z fddZd	d
 Zdd Zdd Z	  Z
S )	LazyProxyc                 C   s   || _ d | _d S r   _init_fn_obj)r   Zinit_fnr   r   r   r      s    zLazyProxy.__init__c                 C   s   | j d u r|  | _ d S r   )r   r   r   r   r   r   _initialize_obj   s    
zLazyProxy._initialize_objc                 C   s   |    t| j|S r   )r   getattrr   r   namer   r   r   __getattr__   s    zLazyProxy.__getattr__c                    s2   |dv rt  || n|   t| j|| d S )Nr   )r"   __setattr__r   setattrr   )r   r   valuer&   r   r   r      s    zLazyProxy.__setattr__c                 C   s   |    t| j| d S r   )r   delattrr   r   r   r   r   __delattr__   s    zLazyProxy.__delattr__c                 C   s,   | j d u r"d| jj d| j dS t| j S )N<z for z not yet initialized>)r   r'   r   r   reprr   r   r   r   __repr__   s    
zLazyProxy.__repr__c                 C   s   |    t| jS r   )r   strr   r   r   r   r   __str__   s    zLazyProxy.__str__)r   r   r   r   r   r   r   r   r   r   rT   r   r   r&   r   r      s   r   c                  C   s4   dd l } | jjd urt S | j r*t S t S d S )Nr   )rg   versionZhipr}   rd   Zis_availablerb   r~   rf   r   r   r   initialize_driver   s    
r   )abcr1   r	   r6   pathlibr   rz   Zcommon.buildr   rJ   r   Zruntimer   ABCMetar   objectr   rU   rb   r|   r}   r~   r   r   r   r   r   r   <module>   s$   )-#