B
     \                 @   s`  d dl mZmZmZ d dlZd dlZddlmZ ddlmZ ddl	m
Z
mZ ejZejZejZG dd	 d	ejZejZejZejZejZejZejZeZejZejZejZejZejZeZ ej!rejnej"Z#ejZ$ejZ%eZ&ejZ'ejZ(G d
d dejZ)ejZ*ejZ+ejZ,ejZ-ejZ.ejZ/G dd dejZ0G dd dejZ1G dd dejZ2e2Z3ejZ4ejZ5ejZ6ejZ7ejZ8ejZ9ejZ:eZ;ejZ<eZ=eZ>ejZ?ejZ@eZAejZBeZCejZDejZEeZFejZGej"ZHdZIG dd dejZJeeJZKejZLejZMejZNG dd dejZOeePejQe#ejRZSejZTeZUejZVejZWejZXejZYejZZejZ[G dd dejZ\eKZ]eZ^ejZ_ejZ`G dd dejZaePeeeeejbee^ZcePee^ZdePee^e]ZeePee^e]ejRZfePee^efejRZgePee^e_ejRZhePee^e;ejieaejbe?ee=ZjG dd dejZkeZlejZmejZnejZoG dd dejZpG dd  d ejZqejZrG d!d" d"ejZsG d#d$ d$ejZteZuejZvejZwejZxG d%d& d&ejZyePeeejzZ{ePdeee)ejzZ|ePee&ejzZ}ePee=eAejzZ~ePeeje>eejRZePdejRZePeeUejRZd'd( Ze \ZZd)d* Zeeeejbged+eg ed+eg ed+eeejRged+eejejejeejQged+eejejejejRged+eeeejRged+ee{ejzged+eeeeejged+eejeejejeejQged+ee#ej"eeee ged+ee ged+e#e gd,e#e gd,de e#gd,de e#gd,e#e e#gd,e#e e#gd,e#e e#gd,e#e e#gd,e#e e#e#gd,e#e e#e#gd,e#e e#e#gd,e#e e#e#gd,de e#gd,de e#gd,de e#gd,de e#gd,de e#gd,de e#gd,de e#gd,de e#gd,de e#gd,de e#gd,de e#gd,de e#gd,de e#gd,de e#gd,de e#gd,de e#gd,de e#gd,de e#gd,de e#gd,de e#gd,e#e e$e#eje%gd,e#e e$e#eje%gd,eeej"e'e|ejRej"ej"eee)ged+ee&ej"e'ej"e eee)ged+eee)ged+eee)ged+ejee)gd,ejee)gd,ejee)gd,ejee)gd,dee)ejgd,dee)ejgd,ejee)ejejgd,ejee)ejejgd,ejee)ejejgd,ejee)ejejgd,ejee)ejgd,ejee)ejgd,ejee)ejgd,ejee)ejgd,dee)ejgd,dee)ejgd,ee&e6ejRged+eee}ejzged+ee&ejeejRged+eejRged+eejRejRejged+eejReeged+eejRejged+eejRejged+eejbee;ged+ee;e<ejRged+ee;e;eejQged+eee>eejRe>ejbeejReejged+eejRejejbee=ged+ee=ged+ee=e@ejRged+ee=ejbeeAged+eeAeBejRged+ee=e~ejRged+eeeDejbeeCged+eeeCgd-eeeCee=ejbgd-eeeCejbgd-eeeCeEejRgd-eeCejbejRged+eeCeejbejRged+eeCeejbejRged+eeCeej"ged+eeeCejbejbeejieeFgd-eeeFeGejRgd-eeeFeGejRgd-eeeeeNgd-eeeeNgd-eeee)ejgd-eeee eeOgd-eeejeejgd-eee e$e#eSejRgd-eeeeejRgd-eej"ej"ee ee$ee#eje%ee#gd-eeeeejRgd-eeee)ej"eej"gd-eeeUeWejRgd-eeeeejRgd-eeeUejej"eejRgd-eeejRgd-eeejReejReejej"ee e gd-eeeeUe[ejRgd-eeej"eeeej"ejRgd-eeeUeUeejQgd-eeejReUej"gd-eeejRejeeejeejRgd-eeejRgd-eeejRgd-eeej"eeejej"eejeejReejeejRgd-eeeejRgd-eeeeeqee\ejReelgd-d.yZdS )/    )print_functionabsolute_importdivisionN   )utils   )enums)HsaApiError
HsaWarningc               @   s(   e Zd ZdejfdejfdejfgZdS )
hsa_dim3_txyzN)__name__
__module____qualname__ctypesc_uint32_fields_ r   r   6lib/python3.7/site-packages/numba/roc/hsadrv/drvapi.pyr      s   r   c               @   sH   e Zd ZdZdefdejfdejfdefdejfdejfdejfgZ	d	S )
hsa_queue_tzwIn theory, this should be aligned to 64 bytes. In any case, allocation
    of this structure is done by the hsa librarytypeZfeaturesZbase_addressZdoorbell_signalsize	reserved1idN)
r   r   r   __doc__hsa_queue_type_tr   r   c_void_phsa_signal_tr   r   r   r   r   r   .   s   r   c               @   s   e Zd Zdejfdejfdejfdejfdejfdejfdejfdejfd	ejfd
ejfdejfdejfdejfdejfdefgZdS )hsa_kernel_dispatch_packet_theaderZsetupZworkgroup_size_xZworkgroup_size_yZworkgroup_size_z	reserved0Zgrid_size_xZgrid_size_yZgrid_size_zZprivate_segment_sizeZgroup_segment_sizeZkernel_objectZkernarg_address	reserved2completion_signalN)	r   r   r   r   c_uint16r   c_uint64r   r   r   r   r   r   r    @   s   r    c               @   sN   e Zd ZdZdejfdejfdejfdejfdejd fdejfd	e	fgZ
d
S )hsa_agent_dispatch_packet_tz5This should be aligned to HSA_PACKET_ALIGN_BYTES (64)r!   r   r"   Zreturn_addressarg   r#   r$   N)r   r   r   r   r   r%   r   r   r&   r   r   r   r   r   r   r'   U   s   r'   c               @   sT   e Zd Zdejfdejfdejfdefdefdefdefdefd	ejfd
efg
ZdS )hsa_barrier_and_packet_tr!   r"   r   Zdep_signal0Zdep_signal1Zdep_signal2Zdep_signal3Zdep_signal4r#   r$   N)	r   r   r   r   r%   r   r   r&   r   r   r   r   r   r*   a   s   r*      c            	   @   sT   e Zd Zdeje fdefdefdejfdejd fdej	fdej	fd	ejfgZ
d
S )BrigModuleHeaderZidentificationZ	brigMajorZ	brigMinorZ	byteCounthash@   ZreservedZsectionCountZsectionIndexN)r   r   r   r   Zc_charMODULE_IDENTIFICATION_LENGTHBrigVersion32_tr&   c_uint8r   r   r   r   r   r   r,      s   r,   c               @   s    e Zd ZdejfdejfgZdS )!hsa_amd_profiling_dispatch_time_tstartendN)r   r   r   r   r&   r   r   r   r   r   r2      s   r2   c               @   s,   e Zd Zdejfdejfdejd fgZdS )hsa_amd_image_descriptor_tversionZdeviceIDdatar   N)r   r   r   r   r   r   r   r   r   r   r5      s   r5   c               @   sn   e Zd Zdejfdejfdejfdejfdejfdejfdejfdejd	 fd
efdejfdejd fgZ	dS )hsa_ext_control_directives_tZcontrol_directives_maskZbreak_exceptions_maskZdetect_exceptions_maskZmax_dynamic_group_sizeZmax_flat_grid_sizeZmax_flat_workgroup_sizer   Zrequired_grid_sizer   Zrequired_workgroup_sizeZrequired_dimr#   K   N)
r   r   r   r   r&   r%   r   r   r1   r   r   r   r   r   r8      s   r8   c               @   s4   e Zd ZdefdefdefdefdefdefgZ	dS )hsa_ext_finalizer_1_00_pfn_tZhsa_ext_program_createZhsa_ext_program_destroyZhsa_ext_program_add_moduleZhsa_ext_program_iterate_modulesZhsa_ext_program_get_infoZhsa_ext_program_finalizeN)
r   r   r   HSA_EXT_PROGRAM_CREATE_FPTRHSA_EXT_PROGRAM_DESTROY_FPTRHSA_EXT_PROGRAM_ADD_MODULE_FPTR$HSA_EXT_PROGRAM_ITERATE_MODULES_FPTRHSA_EXT_PROGRAM_GET_INFO_FPTRHSA_EXT_PROGRAM_FINALIZE_FPTRr   r   r   r   r   r:     s   r:   c               @   s   e Zd ZdefdefgZdS )hsa_ext_image_format_tZchannel_typeZchannel_orderN)r   r   r   hsa_ext_image_channel_type_thsa_ext_image_channel_order_tr   r   r   r   r   rA     s   rA   c               @   s<   e Zd ZdefdejfdejfdejfdejfdefgZdS )hsa_ext_image_descriptor_tZgeometrywidthZheightZdepthZ
array_sizeformatN)r   r   r   hsa_ext_image_geometry_tr   c_size_trA   r   r   r   r   r   rD   $  s   rD   c               @   s    e Zd ZdejfdejfgZdS )hsa_ext_image_data_info_tr   Z	alignmentN)r   r   r   r   rH   r   r   r   r   r   rI   0  s   rI   c               @   s   e Zd ZdefdefgZdS )hsa_ext_image_region_toffsetN)r   r   r   r   r   r   r   r   r   rJ   6  s   rJ   c               @   s"   e Zd ZdefdefdefgZdS )hsa_ext_sampler_descriptor_tZcoordinate_modeZfilter_modeZaddress_modeN)r   r   r   !hsa_ext_sampler_coordinate_mode_thsa_ext_sampler_filter_mode_t!hsa_ext_sampler_addressing_mode_tr   r   r   r   r   rL   A  s   rL   c              C   sb   t  } t  }xHdd ttD D ]2}tt|}d|krD|| |< q$d|kr$|||< q$q$W | |fS )Nc             S   s   g | ]}| d r|qS )ZHSA_)
startswith).0namer   r   r   
<listcomp>  s    z2_build_reverse_error_warn_maps.<locals>.<listcomp>ZSTATUS_ERRORZSTATUS_INFO)r   Z
UniqueDictdirr   getattr)Zerr_mapZwarn_maprR   coder   r   r   _build_reverse_error_warn_maps|  s    


rW   c             C   sd   | t jkr`| t jkr:t| d}d|j|}t| |n&t| d}d|j|}t	
|t d S )NZUNKNOWN_HSA_ERRORzCall to {0} returned {1}ZUNKNOWN_HSA_INFO)r   ZHSA_STATUS_SUCCESSZHSA_STATUS_ERROR	ERROR_MAPgetrF   r   r	   WARN_MAPwarningswarnr
   )resultfuncZ	argumentsZerrnamemsgZwarnnamer   r   r   _check_error  s    

r`   )restypeargtypeserrcheck)ra   rb   )rc   ra   rb   )yZhsa_status_stringZhsa_initZhsa_shut_downZhsa_system_get_infoZhsa_system_extension_supportedZhsa_system_get_extension_tableZhsa_agent_get_infoZhsa_iterate_agentsZ hsa_agent_get_exception_policiesZhsa_agent_extension_supportedZhsa_signal_createZhsa_signal_destroyZhsa_signal_load_acquireZhsa_signal_load_relaxedZhsa_signal_store_relaxedZhsa_signal_store_releaseZhsa_signal_exchange_acq_relZhsa_signal_exchange_acquireZhsa_signal_exchange_relaxedZhsa_signal_exchange_releaseZhsa_signal_cas_acq_relZhsa_signal_cas_acquireZhsa_signal_cas_relaxedZhsa_signal_cas_releaseZhsa_signal_add_acq_relZhsa_signal_add_acquireZhsa_signal_add_relaxedZhsa_signal_add_releaseZhsa_signal_subtract_acq_relZhsa_signal_subtract_acquireZhsa_signal_subtract_relaxedZhsa_signal_subtract_releaseZhsa_signal_and_acq_relZhsa_signal_and_acquireZhsa_signal_and_relaxedZhsa_signal_and_releaseZhsa_signal_or_acq_relZhsa_signal_or_acquireZhsa_signal_or_relaxedZhsa_signal_or_releaseZhsa_signal_xor_acq_relZhsa_signal_xor_acquireZhsa_signal_xor_relaxedZhsa_signal_xor_releaseZhsa_signal_wait_acquireZhsa_signal_wait_relaxedZhsa_queue_createZhsa_soft_queue_createZhsa_queue_destroyZhsa_queue_inactivateZ!hsa_queue_load_read_index_acquireZ!hsa_queue_load_read_index_relaxedZ"hsa_queue_load_write_index_acquireZ"hsa_queue_load_write_index_relaxedZ#hsa_queue_store_write_index_relaxedZ#hsa_queue_store_write_index_releaseZ!hsa_queue_cas_write_index_acq_relZ!hsa_queue_cas_write_index_acquireZ!hsa_queue_cas_write_index_relaxedZ!hsa_queue_cas_write_index_releaseZ!hsa_queue_add_write_index_acq_relZ!hsa_queue_add_write_index_acquireZ!hsa_queue_add_write_index_relaxedZ!hsa_queue_add_write_index_releaseZ"hsa_queue_store_read_index_relaxedZ"hsa_queue_store_read_index_releaseZhsa_region_get_infoZhsa_agent_iterate_regionsZhsa_memory_allocateZhsa_memory_freeZhsa_memory_copyZhsa_memory_assign_agentZhsa_memory_registerZhsa_memory_deregisterZhsa_isa_from_nameZhsa_isa_get_infoZhsa_isa_compatibleZhsa_code_object_serializeZhsa_code_object_deserializeZhsa_code_object_destroyZhsa_code_object_get_infoZhsa_code_object_get_symbolZhsa_code_symbol_get_infoZhsa_code_object_iterate_symbolsZhsa_executable_createZhsa_executable_destroyZhsa_executable_load_code_objectZhsa_executable_freezeZhsa_executable_get_infoZ%hsa_executable_global_variable_defineZ+hsa_executable_agent_global_variable_defineZ'hsa_executable_readonly_variable_defineZhsa_executable_validateZhsa_executable_get_symbolZhsa_executable_symbol_get_infoZhsa_executable_iterate_symbolshsa_amd_coherency_get_typerd   Z&hsa_amd_profiling_set_profiler_enabledZ#hsa_amd_profiling_get_dispatch_timeZ/hsa_amd_profiling_convert_tick_to_system_domainZhsa_amd_signal_async_handlerZhsa_amd_async_functionZhsa_amd_signal_wait_anyZhsa_amd_image_get_info_max_dimZhsa_amd_queue_cu_set_maskZhsa_amd_memory_pool_get_infoZ"hsa_amd_agent_iterate_memory_poolsZhsa_amd_memory_pool_allocateZhsa_amd_memory_pool_freeZhsa_amd_memory_async_copyZ"hsa_amd_agent_memory_pool_get_infoZhsa_amd_agents_allow_accessZhsa_amd_memory_pool_can_migrateZhsa_amd_memory_migrateZhsa_amd_memory_lockhsa_amd_memory_unlockre   hsa_amd_interop_map_bufferrf   Zhsa_amd_image_create)Z
__future__r   r   r   r   r[    r   r   errorr	   r
   ZPOINTERZ_PTRr&   Zhandle_structZc_intZhsa_status_tZ	Structurer   Zhsa_access_permission_tZhsa_endianness_tZhsa_machine_model_tZhsa_profile_tZhsa_system_info_tZhsa_extension_tZhsa_agent_tZhsa_agent_feature_tZhsa_device_type_tZ!hsa_default_float_rounding_mode_tZhsa_agent_info_tZhsa_exception_policy_tr   ZHSA_LARGE_MODELr   Zhsa_signal_value_tZhsa_signal_condition_tZhsa_wait_state_tZhsa_region_tr   Zhsa_queue_feature_tr   Zhsa_packet_type_tZhsa_fence_scope_tZhsa_packet_header_tZhsa_packet_header_width_tZ"hsa_kernel_dispatch_packet_setup_tZ(hsa_kernel_dispatch_packet_setup_width_tr    r'   r*   Zhsa_barrier_or_packet_tZhsa_region_segment_tZhsa_region_global_flag_tZhsa_region_info_tZhsa_symbol_kind_tZhsa_variable_allocation_tZhsa_symbol_linkage_tZhsa_variable_segment_tZ	hsa_isa_tZhsa_isa_info_tZhsa_code_object_tZhsa_callback_data_tZhsa_code_object_type_tZhsa_code_object_info_tZhsa_code_symbol_tZhsa_code_symbol_info_tZhsa_executable_tZhsa_executable_state_tZhsa_executable_info_tZhsa_executable_symbol_tZhsa_executable_symbol_info_tr0   r/   r,   ZBrigModule_tZhsa_amd_agent_info_tZhsa_amd_region_info_tZhsa_amd_coherency_type_tr2   Z	CFUNCTYPEZc_boolr   Zhsa_amd_signal_handlerZhsa_amd_segment_tZhsa_amd_memory_pool_tZ!hsa_amd_memory_pool_global_flag_tZhsa_amd_memory_pool_info_tZhsa_amd_memory_pool_access_tZhsa_amd_link_info_type_tZhsa_amd_memory_pool_link_info_tZ hsa_amd_agent_memory_pool_info_tr5   Zhsa_ext_module_tZhsa_ext_program_tZhsa_ext_program_info_tZ#hsa_ext_finalizer_call_convention_tr8   Zc_char_pr;   r<   r=   Z-HSA_EXT_PROGRAM_ITERATE_MODULES_CALLBACK_FUNCr>   r?   Zc_int32r@   r:   Zhsa_ext_image_trG   rB   rC   rA   rD   Zhsa_ext_image_capability_trI   rJ   Zhsa_ext_sampler_trO   rM   rN   rL   Z	py_objectZHSA_ITER_AGENT_CALLBACK_FUNCZHSA_QUEUE_CALLBACK_FUNCZ'HSA_AGENT_ITERATE_REGIONS_CALLBACK_FUNCZ(HSA_CODE_OBJECT_ITERATE_SYMBOLS_CALLBACKrH   ZHSA_ALLOC_CALLBACK_FUNCTIONZvoid_fn_ptrZ+HSA_AMD_AGENT_ITERATE_MEMORY_POOLS_CALLBACKrW   rX   rZ   r`   r%   ZAPI_PROTOTYPESr   r   r   r   <module>   s  






































	

	
	
			
	



	



	
	


	

	
	

