B
     \[                 @   s$  d Z ddlmZmZmZ ddlZddlZddlZddlZddl	Z	ddl
Z
ddlZddlZddlmZ ddlmZmZ ddlmZ ddlmZ ddlmZ dd	lmZ d
dlmZmZmZ d
dlmZmZmZ ddlm Z! ddl"Z#ej$dkrddl%m&Z& nddlm&Z& e'e(Z)G dd deZ*dd Z+dZ,dd Z-dZ.dd Z/dZ0dd Z1dZ2dd Z3d Z4G d!d" d"e5Z6G d#d$ d$e5Z7e7 Z8G d%d& d&e5Z9eG d'd( d(e9Z:G d)d* d*e&Z;G d+d, d,e9Z<G d-d. d.e9Z=G d/d0 d0e5Z>G d1d2 d2e5Z?G d3d4 d4e5Z@G d5d6 d6e5ZAG d7d8 d8e5ZBG d9d: d:e5ZCG d;d< d<e5ZDG d=d> d>e9ZEG d?d@ d@e5ZFG dAdB dBejGZHG dCdD dDe5ZIG dEdF dFe5ZJG dGdH dHe5ZKdIdJ ZLdKdL ZMdMdN ZNdOdP ZOdQdR ZPdSdT ZQdUdV ZRdWdX ZSdYdZ ZTd[d\ ZUd]d^ ZVd_d` ZWdadb ZXdcdd ZYeY dkZZdS )ez"
HSA driver bridge implementation
    )absolute_importprint_functiondivisionN)contextmanager)defaultdictdeque)total_ordering)mviewbuf)utils)config   )HsaSupportErrorHsaDriverErrorHsaApiError)enums	enums_extdrvapi)longint)   r   )Sequencec               @   s   e Zd ZdS )HsaKernelTimedOutN)__name__
__module____qualname__ r   r   6lib/python3.7/site-packages/numba/roc/hsadrv/driver.pyr   %   s   r   c             C   s(   ydddg|  S  t k
r"   dS X d S )NCPUGPUZDSPZUnknown)
IndexError)devicer   r   r   _device_type_to_string)   s    r    z!/opt/rocm/lib/libhsa-runtime64.soc        
         sf  t jdt} | dkrt  tddks>tjdks>tjdkrFt	  nt
j}ddg}d	 | d k	ryt j| } W n  tk
r   td
|  Y nX t j| std|  | g}n g fdd|D  }g }g }x`|D ]X}y||}W nB tk
r* } z"|t j|  || W d d }~X Y qX |S qW t|rFt  nddd |D }	t|	 d S )NZNUMBA_HSA_DRIVER0P   Zwin32darwinz/usr/libz
/usr/lib64zlibhsa-runtime64.soz'NUMBA_HSA_DRIVER %s is not a valid pathznNUMBA_HSA_DRIVER %s is not a valid file path.  Note it must be a filepath of the .so/.dll/.dylib or the driverc                s   g | ]}t j| qS r   )ospathjoin).0x)dlnamer   r   
<listcomp>R   s    z _find_driver.<locals>.<listcomp>
c             s   s   | ]}t |V  qd S )N)str)r(   er   r   r   	<genexpr>f   s    z_find_driver.<locals>.<genexpr>)r%   environgetDEFAULT_HSA_DRIVER_raise_driver_not_foundstructZcalcsizesysplatform_raise_platform_not_supportedctypesZCDLLr&   abspath
ValueErrorr   isfileOSErrorappendallr'   _raise_driver_error)
ZenvpathZdlloaderZdldirZ
candidatesZpath_not_existZdriver_load_errorr&   Zdllr.   errmsgr   )r*   r   _find_driver3   sD    



rA   z8
HSA is not currently supported on this platform ({0}).
c               C   s   t ttjd S )N)r   PLATFORM_NOT_SUPPORTED_ERRORformatr5   r6   r   r   r   r   r7   o   s    r7   z
The HSA runtime library cannot be found.

If you are sure that the HSA is installed, try setting environment
variable NUMBA_HSA_DRIVER with the file path of the HSA runtime shared
library.
c               C   s   t td S )N)r   DRIVER_NOT_FOUND_MSGr   r   r   r   r3   |   s    r3   zD
A HSA runtime library was found, but failed to load with error:
%s
c             C   s   t t|  d S )N)r   DRIVER_LOAD_ERROR_MSG)r.   r   r   r   r?      s    r?   zdriver missing function: %s.
c               @   s4   e Zd Zdd Zdd Zdd Zdd Zd	d
 ZdS )Recyclerc             C   s   g | _ d| _d S )NT)_garbageenabled)selfr   r   r   __init__   s    zRecycler.__init__c             C   s   | j | |   d S )N)rG   r=   service)rI   objr   r   r   free   s    zRecycler.freec             C   s*   x| j D ]}|| qW | j d d = d S )N)rG   
_finalizer)rI   rL   r   r   r   _cleanup   s    zRecycler._cleanupc             C   s    | j rt| jdkr|   d S )N
   )rH   lenrG   rO   )rI   r   r   r   rK      s    zRecycler.servicec             C   s   |    d| _d S )NF)rO   rH   )rI   r   r   r   drain   s    zRecycler.drainN)r   r   r   rJ   rM   rO   rK   rR   r   r   r   r   rF      s
   rF   c               @   s   e Zd ZdZdZdZejZe	j
ejfe	jejfe	jejfe	jejfe	jejfdZdd Zdd Zdd	 Zd
d Zedd Zedd Ze	je	je	jdfddZdddZdd Zdd Z edd Z!dd Z"dd Z#dS )Driverz0
    Driver API functions are lazily bound.
    N)version_majorversion_minorZ	timestamptimestamp_frequencyZsignal_max_waitc             C   s&   | j }|d k	r|S t| }|| _ |S )N)
_singletonobject__new__)clsrL   r   r   r   rY      s    
zDriver.__new__c          
   C   sx   y&t jrtdt | _d| _d | _W n. tk
rT } zd| _|| _W d d }~X Y nX d | _i | _t	 | _
t | _d S )NzHSA disabled by userFT)r   ZDISABLE_HSAr   rA   libis_initializedinitialization_error
_agent_mapZ	_programsrF   	_recyclerweakrefWeakSet_active_streams)rI   r.   r   r   r   rJ      s    
zDriver.__init__c          
      sh    j r
d S d _ y   W n4 tk
rP } z| _td| W d d }~X Y nX tj fdd}d S )NTzError at driver init: 
%s:c                 s@   yx j D ]} |   q
W W n tk
r0   Y nX  j  d S )N)agentsreleaseAttributeErrorr_   rR   )agent)rI   r   r   shutdown   s    z(Driver._initialize_api.<locals>.shutdown)r\   Zhsa_initr   r]   r   atexitregister)rI   r.   rg   r   )rI   r   _initialize_api   s    zDriver._initialize_apic                sX   | j d k	rd S |   g   fdd}t|}| |d  tdd  D }|| _ d S )Nc                s     |  tjS )N)r=   r   HSA_STATUS_SUCCESS)agent_idctxt)	agent_idsr   r   on_agent   s    
z+Driver._initialize_agents.<locals>.on_agentc             s   s   | ]}|t |fV  qd S )N)Agent)r(   rl   r   r   r   r/      s    z,Driver._initialize_agents.<locals>.<genexpr>)r^   rj   r   ZHSA_ITER_AGENT_CALLBACK_FUNCZhsa_iterate_agentsdict)rI   ro   callbackZ	agent_mapr   )rn   r   _initialize_agents   s    

zDriver._initialize_agentsc             C   s   |    | jd kS )N)rj   r]   )rI   r   r   r   is_available   s    zDriver.is_availablec             C   s   |    | j S )N)rs   r^   values)rI   r   r   r   rc      s    zDriver.agentsc          	   C   s4   t  }|d kst| ||||t| t|S )N)r   hsa_ext_program_tAssertionErrorhsa_ext_program_creater8   byrefProgram)rI   modelprofilerounding_modeoptionsZprogramr   r   r   create_program  s
    
zDriver.create_programc             C   s^   |d krt | j}t|}tj| }|dd |D  }t }| |||t| t	|j
S )Nc             S   s   g | ]
}|j qS r   )_id)r(   cr   r   r   r+     s    z(Driver.create_signal.<locals>.<listcomp>)tuplerc   rQ   r   hsa_agent_thsa_signal_tZhsa_signal_creater8   ry   Signalvalue)rI   Zinitial_valueZ	consumersZconsumers_lenZconsumers_typeresultr   r   r   create_signal  s    


zDriver.create_signalc                s   |    y,| j  \}}| }| |t| |jS  tk
rH   Y nX y| j  }W n tk
rt   t Y nX | j	d k	rt
d| j	 |  }x | D ]\}}t||| qW  fdd}	|	|}
t|  |
 |
S )NzError at driver init: 
%s:c                s    fdd}|S )Nc                 s   t d  | |S )Nzcall driver api: %s)_loggerdebug)argskwargs)fnfnamer   r   wrapped:  s    z;Driver.__getattr__.<locals>.driver_wrapper.<locals>.wrappedr   )r   r   )r   )r   r   driver_wrapper9  s    z*Driver.__getattr__.<locals>.driver_wrapper)rj   _hsa_propertiesZhsa_system_get_infor8   ry   r   KeyError_api_prototypesre   r]   r   	_find_apiitemssetattr)rI   r   enumtypr   protoZlibfnkeyvalr   Zretvalr   )r   r   __getattr__  s,    


zDriver.__getattr__c                s>   yt | j S  tk
r    Y nX  fdd}t|  | |S )Nc                 s   t t  d S )N)r   MISSING_FUNCTION_ERRMSG)r   Zkws)r   r   r   absent_functionL  s    z)Driver._find_api.<locals>.absent_function)getattrr[   re   r   )rI   r   r   r   )r   r   r   C  s    zDriver._find_apic             C   s   t tdd tt| jS )z^Returns a ordered list of components

        The first device should be picked first
        c             S   s   | j S )N)is_component)ar   r   r   <lambda>X  s    z#Driver.components.<locals>.<lambda>)listfilterreversedsortedrc   )rI   r   r   r   
componentsR  s    zDriver.componentsc             C   s   t  }| j| |S )N)Streamrb   add)rI   str   r   r   create_stream[  s    zDriver.create_streamc             C   s&   t d x| jD ]}|  qW dS )zc
        Implicit synchronization for all asynchronous streams
        across all devices.
        zimplicit syncN)r   inforb   synchronize)rI   r   r   r   r   implicit_sync`  s    
zDriver.implicit_sync)N)$r   r   r   __doc__rW   r^   r   ZAPI_PROTOTYPESr   r   ZHSA_SYSTEM_INFO_VERSION_MAJORr8   c_uint16ZHSA_SYSTEM_INFO_VERSION_MINORZHSA_SYSTEM_INFO_TIMESTAMPc_uint64Z#HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCYZHSA_SYSTEM_INFO_SIGNAL_MAX_WAITr   rY   rJ   rj   rs   propertyrt   rc   HSA_MACHINE_MODEL_LARGEHSA_PROFILE_FULL'HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULTr   r   r   r   r   r   r   r   r   r   r   rS      s2   



	
'	rS   c               @   s   e Zd Zdd Zdd ZdS )
HsaWrapperc             C   s   y| j | \}}W n& tk
r8   td| j|f Y nX tt| j}| }t|d}|r^|nt	|}|| j
|| |r|jtjkr|jS t|S d S )Nz%r object has no attribute %rZ_length_)r   r   re   	__class__r   hsa_hsa_info_functionhasattrr8   ry   r   Z_type_c_charr   r   )rI   r   r   r   funcr   Zis_array_typeZresult_buffr   r   r   r   m  s    
zHsaWrapper.__getattr__c             C   s(   t ttt| | j  | j  S )N)r   setdirtype__dict__keysr   )rI   r   r   r   __dir__  s    zHsaWrapper.__dir__N)r   r   r   r   r   r   r   r   r   r   l  s   r   c               @   sn  e Zd ZdZdZejejd fej	ejd fej
ejfejejfejejd fejejfejejfejejfejejfejejfejejfejejfejejfejejd fejejfdZdd Z e!d	d
 Z"e!dd Z#e!dd Z$e!dd Z%e!dd Z&dd Z'dd Z(d.ddZ)dd Z*dd Z+dd Z,d d! Z-d"d# Z.d$d% Z/d&d' Z0d(d) Z1d*d+ Z2d,d- Z3dS )/rp   zrAbstracts a HSA compute agent.

    This will wrap and provide an OO interface for hsa_agent_t C-API elements
    Zhsa_agent_get_info@   r      )namevendor_namefeaturewavefront_sizeworkgroup_max_dimZgrid_max_dimgrid_max_sizeZfbarrier_max_sizeZ
queues_maxqueue_max_size
queue_typeZnode_deviceZ
cache_sizeisac             C   s*   || _ tj| _t | _|   |   d S )N)r   r   r_   r   _queues_initialize_regions_initialize_mempools)rI   rl   r   r   r   rJ     s
    zAgent.__init__c             C   s
   t | jS )N)r    r   )rI   r   r   r   r     s    zAgent.devicec             C   s   | j tj@ dkS )Nr   )r   r   Z!HSA_AGENT_FEATURE_KERNEL_DISPATCH)rI   r   r   r   r     s    zAgent.is_componentc             C   s   | j S )N)_regions)rI   r   r   r   regions  s    zAgent.regionsc             C   s   | j S )N)	_mempools)rI   r   r   r   mempools  s    zAgent.mempoolsc             C   s   t | jddd dS )z&
        log2(wavefront_size)
        N1)binr   index)rI   r   r   r   wavebits  s    zAgent.wavebitsc                sF   g   fdd}t |}tj|d  tfdd D _d S )Nc                s     |  tjS )N)r=   r   rk   )	region_idrm   )
region_idsr   r   	on_region  s    
z,Agent._initialize_regions.<locals>.on_regionc                s   g | ]}t  |qS r   )	MemRegioninstance_for)r(   r   )rI   r   r   r+     s   z-Agent._initialize_regions.<locals>.<listcomp>)r   Z'HSA_AGENT_ITERATE_REGIONS_CALLBACK_FUNCr   Zhsa_agent_iterate_regionsr   _RegionListr   )rI   r   rr   r   )r   rI   r   r     s    
zAgent._initialize_regionsc                sH   g  d fdd	}t |}tj|d  tfdd D _d S )Nc                s     |  tjS )N)r=   r   rk   )r   rm   )mempool_idsr   r   r     s    
z-Agent._initialize_mempools.<locals>.on_regionc                s   g | ]}t  |qS r   )MemPoolr   )r(   Z
mempool_id)rI   r   r   r+     s   z.Agent._initialize_mempools.<locals>.<listcomp>)N)r   Z+HSA_AMD_AGENT_ITERATE_MEMORY_POOLS_CALLBACKr   Z"hsa_amd_agent_iterate_memory_poolsr   r   r   )rI   r   rr   r   )r   rI   r   r     s    
zAgent._initialize_mempoolsNc             C   s   |d k	st || jkst tj}|d kr4td |n||}ttj }	|d kr\tdn|}|d krrtdn|}t	
| j||||||t|	 t| |	}
| j|
 t|
S )Nr   )rw   r   r   ZHSA_QUEUE_CALLBACK_FUNCr8   castZPOINTERZhsa_queue_tc_uint32r   Zhsa_queue_creater   ry   Queuer   r   r`   proxy)rI   sizerr   dataprivate_segment_sizegroup_segment_sizer   Zcb_typcbr   qr   r   r   _create_queue  s    
zAgent._create_queuec             O   s   t j|d< | j||S )Nr   )r   ZHSA_QUEUE_TYPE_SINGLEr   )rI   r   r   r   r   r   create_queue_single  s    
zAgent.create_queue_singlec             O   s   t j|d< | j||S )Nr   )r   ZHSA_QUEUE_TYPE_MULTIr   )rI   r   r   r   r   r   create_queue_multi  s    
zAgent.create_queue_multic             C   s    xt | jD ]}|  qW dS )zJ
        Release all resources

        Called at system teardown
        N)r   r   rd   )rI   r   r   r   r   rd     s    zAgent.releasec             C   s   | j | | j| d S )N)r   remover_   rM   )rI   queuer   r   r   release_queue  s    zAgent.release_queuec             C   s$   d | j| j| j| j| jrdndS )Nz#<HSA agent ({0}): {1} {2} '{3}'{4}>z (component) )rC   r   r   r   r   r   )rI   r   r   r   __repr__  s
    zAgent.__repr__c             C   s   | j | j| jfS )N)r   r   r   )rI   r   r   r   _rank  s    zAgent._rankc             C   s"   t | tr|  | k S tS d S )N)
isinstancerp   r   NotImplemented)rI   otherr   r   r   __lt__  s    
zAgent.__lt__c             C   s"   t | tr|  | kS tS d S )N)r   rp   r   r   )rI   r   r   r   r   __eq__#  s    
zAgent.__eq__c             C   s   t |  S )N)hashr   )rI   r   r   r   __hash__)  s    zAgent.__hash__c             C   s   t | S )N)Context)rI   r   r   r   create_context,  s    zAgent.create_context)NNNNN)4r   r   r   r   r   r   ZHSA_AGENT_INFO_NAMEr8   r   ZHSA_AGENT_INFO_VENDOR_NAMEZHSA_AGENT_INFO_FEATUREr   Zhsa_agent_feature_tZHSA_AGENT_INFO_WAVEFRONT_SIZEr   Z HSA_AGENT_INFO_WORKGROUP_MAX_DIMr   ZHSA_AGENT_INFO_GRID_MAX_DIMZ
hsa_dim3_tZHSA_AGENT_INFO_GRID_MAX_SIZEZ HSA_AGENT_INFO_FBARRIER_MAX_SIZEZHSA_AGENT_INFO_QUEUES_MAXZHSA_AGENT_INFO_QUEUE_MAX_SIZEZHSA_AGENT_INFO_QUEUE_TYPEZhsa_queue_type_tZHSA_AGENT_INFO_NODEZHSA_AGENT_INFO_DEVICEZhsa_device_type_tZHSA_AGENT_INFO_CACHE_SIZEZHSA_AGENT_INFO_ISAZ	hsa_isa_tr   rJ   r   r   r   r   r   r   r   r   r   r   r   rd   r   r   r   r   r   r   r   r   r   r   r   rp     sL   










  
	rp   c               @   s8   e Zd ZdZdd Zdd Zdd Zdd	 Zd
d ZdS )r   )_allglobals	readonlysprivatesgroupsc             C   s^   t || _t dd |D | _t dd |D | _t dd |D | _t dd |D | _d S )Nc             s   s   | ]}|j d kr|V  qdS )globalN)kind)r(   r)   r   r   r   r/   5  s    z'_RegionList.__init__.<locals>.<genexpr>c             s   s   | ]}|j d kr|V  qdS )readonlyN)r  )r(   r)   r   r   r   r/   6  s    c             s   s   | ]}|j d kr|V  qdS )privateN)r  )r(   r)   r   r   r   r/   7  s    c             s   s   | ]}|j d kr|V  qdS )groupN)r  )r(   r)   r   r   r   r/   8  s    )r   r   r   r  r  r  )rI   Zlstr   r   r   rJ   3  s
    
z_RegionList.__init__c             C   s
   t | jS )N)rQ   r   )rI   r   r   r   __len__:  s    z_RegionList.__len__c             C   s
   || j kS )N)r   )rI   itemr   r   r   __contains__=  s    z_RegionList.__contains__c             C   s
   t | jS )N)r   r   )rI   r   r   r   __reversed__@  s    z_RegionList.__reversed__c             C   s
   | j | S )N)r   )rI   idxr   r   r   __getitem__C  s    z_RegionList.__getitem__N)	r   r   r   	__slots__rJ   r	  r  r  r  r   r   r   r   r   0  s   r   c               @   s   e Zd ZdZdZejejfej	e
jfeje
jfeje
jfeje
jfeje
jfeje
jfdZejdejdejdejdiZdd	 Zed
d Zedd Zdd Zdd Zi Ze dd Z!dS )r   z{Abstracts a HSA mem pool.

    This will wrap and provide an OO interface for hsa_amd_memory_pool_t
    C-API elements
    Zhsa_amd_memory_pool_get_info)segment_flagsr   alloc_allowedalloc_granulealloc_alignmentZaccessible_by_allr  r  r  r  c             C   s   || _ || _| j | _dS )z{Do not instantiate MemPool objects directly, use the factory class
        method 'instance_for' to ensure MemPool identityN)r   _owner_agent_as_parameter_)rI   rf   Zpoolr   r   r   rJ   k  s    zMemPool.__init__c             C   s   | j | j S )N)_segment_name_mapr  )rI   r   r   r   r  r  s    zMemPool.kindc             C   s   | j S )N)r  )rI   r   r   r   rf   v  s    zMemPool.agentc             C   s   | j dkr| j|@ S dS dS )a  
            Determines if a given feature is supported by this MemRegion.
            Feature flags are found in "./enums_exp.py" under:
                * hsa_amd_memory_pool_global_flag_t
                Params:
                check_flag: Feature flag to test
        r  FN)r  r  )rI   
check_flagr   r   r   supportsz  s    

zMemPool.supportsc             C   s\   | j s
t|dkstt }td}t| j||t| |j	d krXt
d| |S )Nr   zFailed to allocate from {})r  rw   r8   c_void_pr   r   Zhsa_amd_memory_pool_allocater   ry   r   r   rC   )rI   nbytesbuffflagsr   r   r   allocate  s    


zMemPool.allocatec             C   s8   y
| j | S  tk
r2   | ||}|| j |< |S X d S )N)_instance_dictr   )rZ   ownerr   new_instancer   r   r   r     s    


zMemPool.instance_forN)"r   r   r   r   r   r   Z HSA_AMD_MEMORY_POOL_INFO_SEGMENTr   Zhsa_amd_segment_tZ%HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGSr8   r   ZHSA_AMD_MEMORY_POOL_INFO_SIZEc_size_tZ.HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWEDc_boolZ.HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULEZ0HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALIGNMENTZ*HSA_AMD_MEMORY_POOL_INFO_ACCESSIBLE_BY_ALLr   ZHSA_AMD_SEGMENT_GLOBALZHSA_AMD_SEGMENT_READONLYZHSA_AMD_SEGMENT_PRIVATEZHSA_AMD_SEGMENT_GROUPr  rJ   r   r  rf   r  r  r  classmethodr   r   r   r   r   r   G  s6   

r   c            	   @   s   e Zd ZdZdZejejfej	ej
fejejfejejfejejfejejfejejfejejfdZejdejdejdejdiZdd	 Zed
d Zedd Zdd Zdd Z dd Z!i Z"e#dd Z$dS )r   zsAbstracts a HSA memory region.

    This will wrap and provide an OO interface for hsa_region_t C-API elements
    Zhsa_region_get_info)r  r  host_accessibler   alloc_max_sizer  r  r  r  r  r  r  c             C   s   || _ || _| j | _dS )zDo not instantiate MemRegion objects directly, use the factory class
        method 'instance_for' to ensure MemRegion identityN)r   r  r  )rI   rf   r   r   r   r   rJ     s    zMemRegion.__init__c             C   s   | j | j S )N)r  r  )rI   r   r   r   r    s    zMemRegion.kindc             C   s   | j S )N)r  )rI   r   r   r   rf     s    zMemRegion.agentc             C   s   | j dkr| j|@ S dS dS )a  
            Determines if a given feature is supported by this MemRegion.
            Feature flags are found in "./enums.py" under:
                * hsa_region_global_flag_t
                Params:
                check_flag: Feature flag to test
        r  FN)r  r  )rI   r  r   r   r   r    s    

zMemRegion.supportsc             C   sF   | j s
t|| jkst|dks$tt }t| j|t| |S )Nr   )	r  rw   r&  r8   r  r   Zhsa_memory_allocater   ry   )rI   r  r  r   r   r   r    s    
zMemRegion.allocatec             C   s   t | d S )N)r   hsa_memory_free)rI   Zptrr   r   r   rM     s    zMemRegion.freec             C   s8   y
| j | S  tk
r2   | ||}|| j |< |S X d S )N)r  r   )rZ   r   r   r!  r   r   r   r     s    


zMemRegion.instance_forN)%r   r   r   r   r   r   ZHSA_REGION_INFO_SEGMENTr   Zhsa_region_segment_tZHSA_REGION_INFO_GLOBAL_FLAGSZhsa_region_global_flag_tr   Z#HSA_AMD_REGION_INFO_HOST_ACCESSIBLEr8   r#  ZHSA_REGION_INFO_SIZEr"  ZHSA_REGION_INFO_ALLOC_MAX_SIZEZ'HSA_REGION_INFO_RUNTIME_ALLOC_ALIGNMENTZ%HSA_REGION_INFO_RUNTIME_ALLOC_GRANULEZ%HSA_REGION_INFO_RUNTIME_ALLOC_ALLOWEDr   ZHSA_REGION_SEGMENT_GLOBALZHSA_REGION_SEGMENT_READONLYZHSA_REGION_SEGMENT_PRIVATEZHSA_REGION_SEGMENT_GROUPr  rJ   r   r  rf   r  r  rM   r  r$  r   r   r   r   r   r     s<   
r   c               @   sR   e Zd Zdd Zdd Zdd Zedd Zd	d
 ZdddZ	dd Z
dd ZdS )r   c             C   s&   t || _|| _| j| _tj| _dS )zThe id in a queue is a pointer to the queue object returned by hsa_queue_create.
        The Queue object has ownership on that queue objectN)r`   r   _agentr   r  r   Zhsa_queue_destroyrN   )rI   rf   Z	queue_ptrr   r   r   rJ     s    zQueue.__init__c             C   s   | j |  d S )N)r(  r   )rI   r   r   r   rd     s    zQueue.releasec             C   s   t | jj|S )N)r   r   contents)rI   r   r   r   r   r      s    zQueue.__getattr__c       
      c   s   | j j}|jd }t|ttjks,t||j }t	| j d}x0t
| j }||  krl||j k rFn qFP qFW ||@ }||j}|| }	tt|	dt| |	V  t| j jj| d S )Nr   r   )r   r)  r   r8   sizeofr   hsa_kernel_dispatch_packet_trw   r   Z!hsa_queue_add_write_index_acq_relZ!hsa_queue_load_read_index_acquireZfrom_addressZbase_addressmemset	addressofZhsa_signal_store_releaseZdoorbell_signal)
rI   Zpacket_typeZqueue_structZ
queue_maskZpacket_array_tr   Zread_offsetZqueue_offsetr   packetr   r   r   _get_packet  s     

zQueue._get_packetc          	   C   sl   |  tjV}|j|_d}|tjtj> O }|tjtj> O }|tj	tj
> O }|dtj> O }||_W d Q R X d S )Nr   r   )r/  r   Zhsa_barrier_and_packet_tr   Zdep_signal0r   HSA_FENCE_SCOPE_SYSTEM%HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE%HSA_PACKET_HEADER_RELEASE_FENCE_SCOPEZHSA_PACKET_TYPE_BARRIER_ANDHSA_PACKET_HEADER_TYPEZHSA_PACKET_HEADER_BARRIERheader)rI   Z
dep_signalr.  r4  r   r   r   insert_barrier  s    zQueue.insert_barrierNc          	   C   s  t d|j t|}|t|ks&td|  k r:dks@n t||ksLt|t| jjd | krd}t|	|t| jjd | |d k	r|nt
d}| tj}	|	 j|tj> O  _|d |	_|dkr|d nd|	_|dkr|d nd|	_|d |	_|dkr|d nd|	_|dkr*|d nd|	_|j|	_|j|	_|d krNdn|j|	_|j|	_|j|	_d}
|
tjtj > O }
|
tjtj!> O }
|
tj"tj#> O }
|
|	_$W d Q R X |d krt d d}|j%|d	sd
}t&|j	|d	d S )Nzdispatch %sr   r   z"workgroupsize is too big {0} > {1}r      z&wait for sychronous kernel to completerP   )timeoutz'Kernel timed out after {timeout} second)'r   r   r   rQ   rw   r   r(  r   r   rC   r   r   r/  r   r+  Zsetupr   Z+HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONSZworkgroup_size_xZworkgroup_size_yZworkgroup_size_zZgrid_size_xZgrid_size_yZgrid_size_zr   completion_signalkernel_objectr   Zkernarg_addressr   r   r0  r1  r2  ZHSA_PACKET_TYPE_KERNEL_DISPATCHr3  r4  wait_until_ne_oner   )rI   symbolZkernargsZworkgroup_sizeZ	grid_sizesignalZdimsmsgsr.  r4  r7  r   r   r   dispatch0  sF    



zQueue.dispatchc             C   s   t tt| jj| j  S )N)r   r   r   r   r)  r   r   )rI   r   r   r   r   l  s    zQueue.__dir__c             C   s   t | S )N)ManagedQueueProxy)rI   r   r   r   ownedp  s    zQueue.owned)NNN)r   r   r   rJ   rd   r   r   r/  r5  r?  r   rA  r   r   r   r   r     s     
9r   c               @   s   e Zd Zdd Zdd ZdS )r@  c             C   s   t || _d S )N)r`   ref_queue)rI   r   r   r   r   rJ   u  s    zManagedQueueProxy.__init__c             C   s   t |  |S )N)r   rC  )rI   r
  r   r   r   r   x  s    zManagedQueueProxy.__getattr__N)r   r   r   rJ   r   r   r   r   r   r@  t  s   r@  c               @   s2   e Zd ZdZdd Zdd Zdd Zdd	d
ZdS )r   zThe id for the signal is going to be the hsa_signal_t returned by create_signal.
    Lifetime of the underlying signal will be tied with this object".
    Note that it is likely signals will have lifetime issues.c             C   s$   || _ | j | _t| tj| j  d S )N)r   r  r
   finalizer   Zhsa_signal_destroy)rI   Z	signal_idr   r   r   rJ     s    zSignal.__init__c             C   s   t | jS )N)r   Zhsa_signal_load_relaxedr   )rI   r   r   r   load_relaxed  s    zSignal.load_relaxedc             C   s   t | jS )N)r   Zhsa_signal_load_acquirer   )rI   r   r   r   load_acquire  s    zSignal.load_acquireNc             C   sH   d}d}|dkrd}n|t j | }t | jtj||tj |  |kS )zL
        Returns a boolean to indicate whether the wait has timeout
        r   i@B Nr   )r   rV   Zhsa_signal_wait_acquirer   r   ZHSA_SIGNAL_CONDITION_NEZHSA_WAIT_STATE_ACTIVErE  )rI   r7  ZoneZmhzZexpirer   r   r   r:    s    zSignal.wait_until_ne_one)N)r   r   r   r   rJ   rE  rF  r:  r   r   r   r   r   |  s
   r   c               @   s0   e Zd Zdd Zedd Zdd Zdd Zd	S )

BrigModulec             C   s*   t |}|| _t t |tj| _dS )z5
        Take a byte buffer of a Brig module
        N)r8   create_string_buffer_bufferr   r-  r   Zhsa_ext_module_tr   )rI   Zbrig_bufferbufr   r   r   rJ     s    
zBrigModule.__init__c          	   C   s&   t |d}| }W d Q R X t|S )Nrb)openreadrG  )rZ   	file_namefinrJ  r   r   r   	from_file  s    zBrigModule.from_filec             C   s
   t | jS )N)rQ   rI  )rI   r   r   r   r	    s    zBrigModule.__len__c             C   s   d tt| t| S )Nz!<BrigModule id={0} size={1}bytes>)rC   hexidrQ   )rI   r   r   r   r     s    zBrigModule.__repr__N)r   r   r   rJ   r$  rP  r	  r   r   r   r   r   rG    s   	rG  c               @   s:   e Zd ZejejejdddfddZdd Zd
dd	Z	dS )rz   Nr   r   c       
   	   C   s   t  | _|d kstdd }td}ttj	||t
| |jsVtd||f t  | _ttj	||t
| j | j||||t
| j}	||	 | j| _t| | jj| j d S )Nc             S   sD   | t jk	r@t }t| t| t|j	
d t|   d S )Nzutf-8)r   rk   r8   Zc_char_pr   Zhsa_status_stringry   r   r   r   decodeexit)Z
hsa_statusr=  r   r   r   check_fptr_return  s
    
z+Program.__init__.<locals>.check_fptr_returnr   z(HSA system extension %s.%s not supported)r   rv   r   rw   r8   r#  r   Zhsa_system_extension_supportedr   ZHSA_EXTENSION_FINALIZERry   r   Zhsa_ext_finalizer_1_00_pfn_t_ftablZhsa_system_get_extension_tablerx   r  r
   rD  Zhsa_ext_program_destroy)
rI   r{   r|   r}   r~   rT   rU   rU  Zsupportretr   r   r   rJ     s,    




zProgram.__init__c             C   s   | j | j|j d S )N)rV  Zhsa_ext_program_add_moduler   )rI   moduler   r   r   
add_module  s    zProgram.add_modulec             C   sT   t  }t  }tt|dt| | j| j	||||t
jt| t|S )zN
        The program object is safe to be deleted after ``finalize``.
        r   )r   Zhsa_code_object_tZhsa_ext_control_directives_tr8   r,  ry   r*  rV  Zhsa_ext_program_finalizer   r   ZHSA_CODE_OBJECT_TYPE_PROGRAM
CodeObject)rI   r   Zcallconvr~   code_objectZcontrol_directivesr   r   r   rD    s    
zProgram.finalize)r   N)
r   r   r   r   r   r   r   rJ   rY  rD  r   r   r   r   rz     s   'rz   c               @   s   e Zd Zdd ZdS )rZ  c             C   s$   || _ | j | _t| tj| j  d S )N)r   r  r
   rD  r   Zhsa_code_object_destroy)rI   r[  r   r   r   rJ     s    zCodeObject.__init__N)r   r   r   rJ   r   r   r   r   rZ    s   rZ  c               @   s,   e Zd Zdd Zdd Zdd Zdd Zd	S )

Executablec             C   sF   t  }ttjtjd t| || _	| j	| _
t| tj| j	 d S )N)r   Zhsa_executable_tr   Zhsa_executable_creater   r   ZHSA_EXECUTABLE_STATE_UNFROZENr8   ry   r   r  r
   rD  Zhsa_executable_destroy)rI   exr   r   r   rJ     s    zExecutable.__init__c             C   s   t | j|j|jd  d S )N)r   Zhsa_executable_load_code_objectr   )rI   rf   r[  r   r   r   load  s    zExecutable.loadc             C   s   t | jd dS )z0Freeze executable before we can query for symbolN)r   Zhsa_executable_freezer   )rI   r   r   r   freeze  s    zExecutable.freezec          
   C   s<   t  }t| jd t|d|jdt| t	||S )Nasciir   )
r   Zhsa_executable_symbol_tr   Zhsa_executable_get_symbolr   r8   rH  encodery   Symbol)rI   rf   r   r;  r   r   r   
get_symbol  s    

zExecutable.get_symbolN)r   r   r   rJ   r^  r_  rc  r   r   r   r   r\     s   
r\  c               @   sF   e Zd ZdZejejfejej	fej
ej	fejej	fdZdd ZdS )rb  Zhsa_executable_symbol_get_info)r9  Zkernarg_segment_sizer   r   c             C   s   || _ || _d S )N)r   r   )rI   r   Z	symbol_idr   r   r   rJ   2  s    zSymbol.__init__N)r   r   r   r   r   Z(HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECTr8   r   Z6HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZEr   Z4HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZEZ6HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZEr   rJ   r   r   r   r   rb    s   rb  c               @   sN   e Zd ZdZdddZdd Zdd Zd	d
 Zdd Ze	dd Z
dd ZdS )MemoryPointerTNc             C   sF   t |tst|| _|| _|| _|| _|| _|d k	| _d| _	d| _
d S )NTr   )r   r   rw   contextdevice_pointerr   Z_hsa_memsize_	finalizer
is_managedis_aliverefct)rI   re  pointerr   rg  r   r   r   rJ   :  s    
zMemoryPointer.__init__c             C   s2   y| j r| jr|   W n   t  Y nX d S )N)rh  ri  rg  	traceback	print_exc)rI   r   r   r   __del__E  s
    zMemoryPointer.__del__c             C   s   t t| S )N)OwnedPointerr`   r   )rI   r   r   r   ownL  s    zMemoryPointer.ownc             C   s&   | j r"| jstd|   d| _dS )z8
        Forces the device memory to the trash.
        zFreeing dead memoryFN)rh  ri  RuntimeErrorrg  )rI   r   r   r   rM   O  s
    zMemoryPointer.freec             C   s(   | j j}t| j|| j}tt| |S )N)rf  r   rd  re  r   ro  r`   r   )rI   rk  viewr   r   r   rr  Y  s    zMemoryPointer.viewc             C   s   | j S )N)rf  )rI   r   r   r   device_ctypes_pointer^  s    z#MemoryPointer.device_ctypes_pointerc             G   sB   t |}|dkrdS |tj dd |D  }t||d| j dS )z
        Grant access to given *agents*.
        Upon return, only the listed-agents and the owner agent have direct
        access to this pointer.
        r   Nc             S   s   g | ]
}|j qS r   )r   )r(   r   r   r   r   r+   k  s    z1MemoryPointer.allow_access_to.<locals>.<listcomp>)rQ   r   r   r   Zhsa_amd_agents_allow_accessrf  )rI   rc   ctZagent_arrayr   r   r   allow_access_tob  s    
zMemoryPointer.allow_access_to)N)r   r   r   __hsa_memory__rJ   rn  rp  rM   rr  r   rs  ru  r   r   r   r   rd  7  s   

rd  c               @   s   e Zd Zdd Zdd ZdS )
HostMemoryc             C   s6   || _ || _|| _|| _| j| _| j| _| jj| _d S )N)re  rA  r   host_pointerhandleZ_buflen_r   Z_bufptr_)rI   re  r   rk  r   r   r   r   rJ   q  s    zHostMemory.__init__c             C   s   | S )Nr   )rI   r   r   r   rp  |  s    zHostMemory.ownN)r   r   r   rJ   rp  r   r   r   r   rw  p  s   rw  c               @   s&   e Zd ZdddZdd Zdd ZdS )	ro  Nc             C   s<   || _ | j  jd7  _|d kr(| j | _n|jr2t|| _d S )Nr   )_memrj  _viewrh  rw   )rI   Zmemptrrr  r   r   r   rJ     s    

zOwnedPointer.__init__c             C   sf   y:| j  jd8  _| j jdks"t| j jdkr8| j   W n& tk
rN   Y n   t  Y nX d S )Nr   r   )rz  rj  rw   rM   ReferenceErrorrl  rm  )rI   r   r   r   rn    s    zOwnedPointer.__del__c             C   s   t | j|S )z$Proxy MemoryPointer methods
        )r   r{  )rI   r   r   r   r   r     s    zOwnedPointer.__getattr__)N)r   r   r   rJ   rn  r   r   r   r   r   ro    s   
	ro  c               @   s   e Zd ZdZe Zdd Zdd Ze	dd Z
e	dd	 Ze	d
d Ze	dd Ze	dd ZdddZdddZdd ZdS )r   z2
    A context is associated with a component
    c             C   s   t || _| jjr6|j}| jj|| jd}| | _t	
 | _tj}tj}dd |jjD }d | _d | _x,|D ]$}||r|| _||rp|| _qpW d S )N)rr   c             S   s   g | ]}|j r|qS r   )r  )r(   mpr   r   r   r+     s    z$Context.__init__.<locals>.<listcomp>)r`   r   r(  r   r   r   	_callbackrA  _defaultqueuer
   Z
UniqueDictallocationsr   Z.HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINEDZ,HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINEDr   r   _coarsegrain_mempool_finegrain_mempoolr  )rI   rf   ZqsZdefqZcoarse_flagZ	fine_flagZ	alloc_mpsr}  r   r   r   rJ     s     




zContext.__init__c             C   s   t || td d S )Nr   )r   Z_check_errorr5   rT  )rI   Zstatusr   r   r   r   r~    s    zContext._callbackc             C   s   | S )Nr   )rI   r   r   r   unproxy  s    zContext.unproxyc             C   s   | j S )N)r  )rI   r   r   r   default_queue  s    zContext.default_queuec             C   s   | j S )N)r(  )rI   r   r   r   rf     s    zContext.agentc             C   s$   | j d krd| j}t|| j S )Nz*coarsegrain mempool is not available in {})r  rC   r(  r:   )rI   r=  r   r   r   coarsegrain_mempool  s    
zContext.coarsegrain_mempoolc             C   s$   | j d krd| j}t|| j S )Nz(finegrain mempool is not available in {})r  rC   r(  r:   )rI   r=  r   r   r   finegrain_mempool  s    
zContext.finegrain_mempoolNTc          	   C   s  | j j}| j j}t }t }|dks,|dkr|dk	r~xL|D ]>}d}	x|D ]}
||
rH|	d7 }	qHW |	t|kr:|| q:W n|}xV|D ]@}|dkr|tjsq|r|j	r|| q|j	s|| qW nt
d| t|dkstdd}x>|D ]6}yt| j ||}W n tk
r*   Y qX P qW |dkrNt
d|||f ttj}tt| |||| |d	}|jdkrt
d
|| j|j< | S )a  
        Allocates memory.
        Parameters:
        nbytes the number of bytes to allocate.
        memTypeFlags the flags for which the memory region must have support,                     due to the inherent rawness of the underlying call, the                     validity of the flag is not checked, cf. C language.
        hostAccessible boolean as to whether the region in which the                       allocation takes place should be host accessible
        r   r   Nr   r   zUnknown device type string "%s"z!No suitable memory regions found.zMemory allocation failed. No agent/region               combination could meet allocation restraints               (hardware = %s, size = %s, flags = %s).)rg  zMemoryPointer has no value)r(  r   r   r   r  rQ   r=   r   Z%HSA_REGION_GLOBAL_FLAG_COARSE_GRAINEDr%  rq  rw   r   r   r  r   _make_mem_finalizerr   r'  rd  r`   r   r   r  rp  )rI   r  ZmemTypeFlagsZhostAccessibleZhwZall_regZ	flag_ok_rr   rcountr  memr   rO  rW  r   r   r   memalloc  sX    







zContext.memallocr   Fc             C   s^   |r
| j n| j}||}ttj}tt| |||| |d}|j	|  || j
|j< | S )z
        Allocates memory in a memory pool.
        Parameters:
        *nbytes* the number of bytes to allocate.
        *allow_acces_to*
        *finegrain*
        )rg  )r  r  r  r  r   Zhsa_amd_memory_pool_freerd  r`   r   ru  r  r   rp  )rI   r  ru  	finegrainZmempoolr  rO  r}  r   r   r   mempoolalloc*  s    	



zContext.mempoolallocc             C   s*   | j |||d}tt| ||j|jdS )N)ru  r  )r   rk  r   )r  rw  r`   r   rf  r   )rI   r   r  ru  r  r   r   r   memhostalloc>  s    zContext.memhostalloc)NT)r   F)r   r   r   r   r`   ra   rb   rJ   r~  r   r  r  rf   r  r  r  r  r  r   r   r   r   r     s   
K
r   c               @   sL   e Zd ZdZdd Zdd Zdd Zdd	 Zd
d Zdd Z	e
dd ZdS )r   z.
    An asynchronous stream for async API
    c             C   s   t  | _tt| _d S )N)r   _signalsr   r   
_callbacks)rI   r   r   r   rJ   I  s    zStream.__init__c             C   s(   t | jdkr| d | j| dS )zA
        Add a signal that corresponds to an async task.
        d   2   N)rQ   r  _syncr=   )rI   r<  r   r   r   _add_signalM  s    
zStream._add_signalc             C   s$   t |st| j|   | d S )N)callablerw   r  _get_last_signalr=   )rI   rr   r   r   r   _add_callbackV  s    zStream._add_callbackc             C   s   | j r| j d S dS )z&
        Get the last signal.
        r   N)r  )rI   r   r   r   r  Z  s    zStream._get_last_signalc             C   s   |  t| j dS )z)
        Synchronize the stream.
        N)r  rQ   r  )rI   r   r   r   r   `  s    zStream.synchronizec             C   sf   d}x\| j r`||krP | j  }| dkr4|  x| j| D ]
}|  q@W | j|= |d7 }qW d S )Nr   r   )r  popleftrE  r:  r  )rI   limitrt  Zsigr   r   r   r   r  f  s    

zStream._syncc             c   s   | V  |    dS )z
        A context manager that waits for all commands in this stream to execute
        and commits any pending memory transfers upon exiting the context.
        N)r   )rI   r   r   r   auto_synchronizes  s    zStream.auto_synchronizeN)r   r   r   r   rJ   r  r  r  r   r  r   r  r   r   r   r   r   E  s   	r   c                s    fdd}|S )z
    finalises memory
    Parameters:
    dtor a function that will delete/free held memory from a reference

    Returns:
    Finalising function
    c                s"   | j  tj fdd}|S )Nc                  s:   t d   r(t dj   j=    d S )NzCurrent allocations: %szAttempting delete on %s)r   r   r   r   )r  dtorry  syncr   r   core  s    z7_make_mem_finalizer.<locals>.mem_finalize.<locals>.core)r  r   r   )re  ry  r  )r  )r  ry  r  r   mem_finalize  s    z)_make_mem_finalizer.<locals>.mem_finalizer   )r  r  r   )r  r   r  }  s    	r  c             C   s
   t | jS )z$Get the device pointer as an integer)rs  r   )rL   r   r   r   rf    s    rf  c             C   s   | dkrt dS t|  | jS )z,Get the ctypes object for the device pointerNr   )r  require_device_memoryrs  )rL   r   r   r   rs    s    rs  c             C   s   t | ddS )aj  All HSA dGPU memory object is recognized as an instance with the
    attribute "__hsa_memory__" defined and its value evaluated to True.

    All HSA memory object should also define an attribute named
    "device_pointer" which value is an int(or long) object carrying the pointer
    value of the device memory address.  This is not tested in this method.
    rv  F)r   )rL   r   r   r   is_device_memory  s    r  c             C   s   t | stddS )z8A sentry for methods that accept HSA memory object.
    zNot a HSA memory object.N)r  	Exception)rL   r   r   r   r    s    r  c             C   s*   t | ttfr| S t | tj}t| |S )z
    NOTE: The underlying data pointer from the host data buffer is used and
    it should not be changed until the operation which can be asynchronous
    completes.
    )r   intlongnpZvoidr	   Zmemoryview_get_buffer)rL   Zforcewritabler   r   r   rx    s    rx  c             C   s8   t d |dk rtd| tt|t|| dS )a  
    Copy data from a host memory region to a dGPU.
    Parameters:
    context the dGPU context
    dst a pointer to the destination location in dGPU memory
    src a pointer to the source location in host memory
    size the size (in bytes) of data to transfer
    z	CPU->dGPUr   zInvalid size given: %sN)r   r   r:   r   hsa_memory_copyrf  rx  )re  dstsrcr   r   r   r   host_to_dGPU  s    	
r  c             C   s8   t d |dk rtd| tt|t|| dS )a  
    Copy data from a host memory region to a dGPU.
    Parameters:
    context the dGPU context
    dst a pointer to the destination location in dGPU memory
    src a pointer to the source location in host memory
    size the size (in bytes) of data to transfer
    z	dGPU->CPUr   zInvalid size given: %sN)r   r   r:   r   r  rx  rf  )re  r  r  r   r   r   r   dGPU_to_host  s    	
r  c             C   s8   t d |dk rtd| tt|t|| d S )Nz
dGPU->dGPUr   zInvalid size given: %s)r   r   r:   r   r  rf  )re  r  r  r   r   r   r   dGPU_to_dGPU  s    
r  c             C   s*   t d t| |t|t|||d d S )NzAsync CPU->dGPU)dst_ctxsrc_ctxr  r  r   stream)r   r   async_copy_dgpurx  rf  )r  r  r  r  r   r  r   r   r   async_host_to_dGPU  s    
r  c             C   s*   t d t| |t|t|||d d S )NzAsync dGPU->CPU)r  r  r  r  r   r  )r   r   r  rx  rf  )r  r  r  r  r   r  r   r   r   async_dGPU_to_host  s    
r  c             C   s*   t d t| |t|t|||d d S )NzAsync dGPU->dGPU)r  r  r  r  r   r  )r   r   r  rf  )r  r  r  r  r   r  r   r   r   async_dGPU_to_dGPU  s    
r  c       
      C   s   |dk rt d| td}| }|d k	rLt|j}dt||f}	n
dd |f}	tj	|| j
j||j
j|f|	  || d S )Nr   zInvalid size given: %sr   )r:   r   r   r  r   r   r   r8   ry   Zhsa_amd_memory_async_copyr(  r  )
r  r  r  r  r   r  r8  Zdependent_signalZdsignalZsignalsr   r   r   r    s    

r  c              C   sB   d} y,x&t jD ]}|jr|jdkr| d7 } qW W n   Y nX | S )zM
    Returns the number of discrete GPUs present on the current machine.
    r   r   r   )r   rc   r   r   )Zngpusr   r   r   r   
dgpu_count  s    r  )[r   Z
__future__r   r   r   r5   rh   r%   r8   r4   rl  r`   Zlogging
contextlibr   collectionsr   r   Znumba.utilsr   Znumbar	   r
   r   errorr   r   r   r   r   r   r   r   r  Znumpyr  Z	PYVERSIONZcollections.abcr   Z	getLoggerr   r   r   r    r2   rA   rB   r7   rD   r3   rE   r?   r   rX   rF   rS   r   r   rp   r   r   r   r   r@  r   rG  rz   rZ  r\  rb  rd  ZMemAllocrw  ro  r   r   r  rf  rs  r  r  rx  r  r  r  r  r  r  r  r  Zdgpu_presentr   r   r   r   <module>   s   

9
 B +VW $@9 +8