1import ctypes 2import warnings 3 4from numba.core import utils 5 6from numba.roc.hsadrv import enums 7from .error import HsaApiError, HsaWarning 8 9_PTR = ctypes.POINTER 10 11# This deals with types which are defined as 12# typedef struct { uint64_t handle;}; 13handle_struct = ctypes.c_uint64 14 15#------------------------------------------------------------------------------ 16# HSA types from hsa.h, ordered as per header file 17 18hsa_status_t = ctypes.c_int # enum 19class hsa_dim3_t(ctypes.Structure): 20 _fields_ = [ 21 ('x', ctypes.c_uint32), 22 ('y', ctypes.c_uint32), 23 ('z', ctypes.c_uint32) 24 ] 25hsa_access_permission_t = ctypes.c_int # enum 26hsa_endianness_t = ctypes.c_int # enum 27hsa_machine_model_t = ctypes.c_int # enum 28hsa_profile_t = ctypes.c_int # enum 29hsa_system_info_t = ctypes.c_int # enum 30hsa_extension_t = ctypes.c_int # enum 31hsa_agent_t = handle_struct 32hsa_agent_feature_t = ctypes.c_int # enum 33hsa_device_type_t = ctypes.c_int # enum 34hsa_default_float_rounding_mode_t = ctypes.c_int # enum 35hsa_agent_info_t = ctypes.c_int # enum 36hsa_exception_policy_t = ctypes.c_int # enum 37hsa_signal_t = handle_struct 38hsa_signal_value_t = ctypes.c_uint64 if enums.HSA_LARGE_MODEL else ctypes.c_uint32 39hsa_signal_condition_t = ctypes.c_int # enum 40hsa_wait_state_t = ctypes.c_int # enum 41hsa_region_t = handle_struct 42hsa_queue_type_t = ctypes.c_int # enum 43hsa_queue_feature_t = ctypes.c_int # enum 44class hsa_queue_t(ctypes.Structure): 45 """In theory, this should be aligned to 64 bytes. In any case, allocation 46 of this structure is done by the hsa library""" 47 _fields_ = [ 48 ('type', hsa_queue_type_t), 49 ('features', ctypes.c_uint32), 50 ('base_address', ctypes.c_void_p), # if LARGE MODEL 51 ('doorbell_signal', hsa_signal_t), 52 ('size', ctypes.c_uint32), 53 ('reserved1', ctypes.c_uint32), 54 ('id', ctypes.c_uint32), 55 ] 56hsa_packet_type_t = ctypes.c_int # enum 57hsa_fence_scope_t = ctypes.c_int # enum 58hsa_packet_header_t = ctypes.c_int # enum 59hsa_packet_header_width_t = ctypes.c_int # enum 60hsa_kernel_dispatch_packet_setup_t = ctypes.c_int # enum 61hsa_kernel_dispatch_packet_setup_width_t = ctypes.c_int # enum 62class hsa_kernel_dispatch_packet_t(ctypes.Structure): 63 _fields_ = [ 64 ('header', ctypes.c_uint16), 65 ('setup', ctypes.c_uint16), 66 ('workgroup_size_x', ctypes.c_uint16), 67 ('workgroup_size_y', ctypes.c_uint16), 68 ('workgroup_size_z', ctypes.c_uint16), 69 ('reserved0', ctypes.c_uint16), # Must be zero 70 ('grid_size_x', ctypes.c_uint32), 71 ('grid_size_y', ctypes.c_uint32), 72 ('grid_size_z', ctypes.c_uint32), 73 ('private_segment_size', ctypes.c_uint32), 74 ('group_segment_size', ctypes.c_uint32), 75 ('kernel_object', ctypes.c_uint64), 76 # NOTE: Small model not dealt with properly...! 77 # ifdef HSA_LARGE_MODEL 78 ('kernarg_address', ctypes.c_uint64), 79 # SMALL Machine has a reserved uint32 80 ('reserved2', ctypes.c_uint64), # Must be zero 81 ('completion_signal', hsa_signal_t), 82 ] 83class hsa_agent_dispatch_packet_t(ctypes.Structure): 84 """This should be aligned to HSA_PACKET_ALIGN_BYTES (64)""" 85 _fields_ = [ 86 ('header', ctypes.c_uint16), 87 ('type', ctypes.c_uint16), 88 ('reserved0', ctypes.c_uint32), 89 # NOTE: Small model not dealt with properly...! 90 ('return_address', ctypes.c_void_p), 91 ('arg', ctypes.c_uint64 * 4), 92 ('reserved2', ctypes.c_uint64), 93 ('completion_signal', hsa_signal_t), 94 ] 95class hsa_barrier_and_packet_t(ctypes.Structure): 96 _fields_ = [ 97 ('header', ctypes.c_uint16), 98 ('reserved0', ctypes.c_uint16), 99 ('reserved1', ctypes.c_uint32), 100 ('dep_signal0', hsa_signal_t), 101 ('dep_signal1', hsa_signal_t), 102 ('dep_signal2', hsa_signal_t), 103 ('dep_signal3', hsa_signal_t), 104 ('dep_signal4', hsa_signal_t), 105 ('reserved2', ctypes.c_uint64), 106 ('completion_signal', hsa_signal_t), 107 ] 108 109hsa_barrier_or_packet_t = hsa_barrier_and_packet_t 110 111hsa_region_segment_t = ctypes.c_int # enum 112hsa_region_global_flag_t = ctypes.c_int # enum 113hsa_region_info_t = ctypes.c_int # enum 114hsa_symbol_kind_t = ctypes.c_int # enum 115hsa_variable_allocation_t = ctypes.c_int # enum 116hsa_symbol_linkage_t = ctypes.c_int # enum 117hsa_variable_segment_t = ctypes.c_int # enum 118hsa_isa_t = handle_struct 119hsa_isa_info_t = ctypes.c_int # enum 120hsa_code_object_t = handle_struct 121hsa_callback_data_t = handle_struct 122hsa_code_object_type_t = ctypes.c_int # enum 123hsa_code_object_info_t = ctypes.c_int # enum 124hsa_code_symbol_t = handle_struct 125hsa_code_symbol_info_t = ctypes.c_int # enum 126hsa_executable_t = handle_struct 127hsa_executable_state_t = ctypes.c_int # enum 128hsa_executable_info_t = ctypes.c_int # enum 129hsa_executable_symbol_t = handle_struct 130hsa_executable_symbol_info_t = ctypes.c_int # enum 131#------------------------------------------------------------------------------ 132 133 134#------------------------------------------------------------------------------ 135# HSA types from Brig.h, ordered as per header file 136# NOTE: not all of the definitions are needed 137BrigVersion32_t = ctypes.c_uint32 138MODULE_IDENTIFICATION_LENGTH=8 139class BrigModuleHeader(ctypes.Structure): 140 _fields_ = [ 141 ('identification', ctypes.c_char*MODULE_IDENTIFICATION_LENGTH), 142 ('brigMajor', BrigVersion32_t), 143 ('brigMinor', BrigVersion32_t), 144 ('byteCount', ctypes.c_uint64), 145 ('hash', ctypes.c_uint8*64), 146 ('reserved', ctypes.c_uint32), 147 ('sectionCount', ctypes.c_uint32), 148 ('sectionIndex', ctypes.c_uint64), 149 ] 150 151BrigModule_t = _PTR(BrigModuleHeader) 152 153#------------------------------------------------------------------------------ 154 155 156#------------------------------------------------------------------------------ 157# HSA types from hsa_ext_amd.h, ordered as per header file 158hsa_amd_agent_info_t = ctypes.c_int # enum 159hsa_amd_region_info_t = ctypes.c_int # enum 160hsa_amd_coherency_type_t = ctypes.c_int # enum 161class hsa_amd_profiling_dispatch_time_t(ctypes.Structure): 162 _fields_ = [ 163 ('start', ctypes.c_uint64), 164 ('end', ctypes.c_uint64), 165 ] 166 167# typedef bool (*hsa_amd_signal_handler)(hsa_signal_value_t value, void* arg); 168hsa_amd_signal_handler = _PTR( 169 ctypes.CFUNCTYPE(ctypes.c_bool, 170 hsa_signal_value_t, 171 ctypes.c_void_p) 172 ) 173 174hsa_amd_segment_t = ctypes.c_int # enum 175hsa_amd_memory_pool_t = handle_struct 176hsa_amd_memory_pool_global_flag_t = ctypes.c_int # enum 177hsa_amd_memory_pool_info_t = ctypes.c_int # enum 178hsa_amd_memory_pool_access_t = ctypes.c_int # enum 179hsa_amd_link_info_type_t = ctypes.c_int # enum 180hsa_amd_memory_pool_link_info_t = ctypes.c_int # enum 181hsa_amd_agent_memory_pool_info_t = ctypes.c_int # enum 182class hsa_amd_image_descriptor_t(ctypes.Structure): 183 _fields_ = [ 184 ('version', ctypes.c_uint32), 185 ('deviceID', ctypes.c_uint32), 186 ('data', ctypes.c_uint32*1), 187 ] 188#------------------------------------------------------------------------------ 189 190 191#------------------------------------------------------------------------------ 192# HSA types from hsa_ext_finalize.h, ordered as per header file 193hsa_ext_module_t = BrigModule_t 194 195hsa_ext_program_t = handle_struct 196hsa_ext_program_info_t = ctypes.c_int # enum 197hsa_ext_finalizer_call_convention_t = ctypes.c_int # enum 198class hsa_ext_control_directives_t(ctypes.Structure): 199 _fields_ = [ 200 ('control_directives_mask', ctypes.c_uint64), 201 ('break_exceptions_mask', ctypes.c_uint16), 202 ('detect_exceptions_mask', ctypes.c_uint16), 203 ('max_dynamic_group_size', ctypes.c_uint32), 204 ('max_flat_grid_size', ctypes.c_uint64), 205 ('max_flat_workgroup_size', ctypes.c_uint32), 206 ('reserved1', ctypes.c_uint32), 207 ('required_grid_size', ctypes.c_uint64*3), 208 ('required_workgroup_size', hsa_dim3_t), 209 ('required_dim', ctypes.c_uint8), 210 ('reserved2', ctypes.c_uint8*75), 211 ] 212 213# function pointers, that are used in the 214# "hsa_ext_finalizer_1_00_pfn_t" struct of pointers 215HSA_EXT_PROGRAM_CREATE_FPTR = ctypes.CFUNCTYPE( 216 hsa_status_t, # return value 217 hsa_machine_model_t, # machine_model 218 hsa_profile_t, # profile 219 hsa_default_float_rounding_mode_t, # default_float_rounding_mode 220 ctypes.c_char_p, # options 221 _PTR(hsa_ext_program_t)) # program 222 223HSA_EXT_PROGRAM_DESTROY_FPTR = ctypes.CFUNCTYPE( 224 hsa_status_t, # return value 225 hsa_ext_program_t) # program 226 227HSA_EXT_PROGRAM_ADD_MODULE_FPTR = ctypes.CFUNCTYPE( 228 hsa_status_t, # return value 229 hsa_ext_program_t, # program 230 hsa_ext_module_t) # module 231 232HSA_EXT_PROGRAM_ITERATE_MODULES_CALLBACK_FUNC = ctypes.CFUNCTYPE( 233 hsa_status_t, # return 234 hsa_ext_program_t, # program 235 hsa_ext_module_t, # module 236 ctypes.c_void_p) # data 237 238HSA_EXT_PROGRAM_ITERATE_MODULES_FPTR = ctypes.CFUNCTYPE( 239 hsa_status_t, # return value 240 hsa_ext_program_t, # program 241 HSA_EXT_PROGRAM_ITERATE_MODULES_CALLBACK_FUNC, # callback 242 ctypes.c_void_p) # data 243 244HSA_EXT_PROGRAM_GET_INFO_FPTR = ctypes.CFUNCTYPE( 245 hsa_status_t, # return value 246 hsa_ext_program_t, # program 247 hsa_ext_program_info_t, # attribute 248 ctypes.c_void_p) # value 249 250HSA_EXT_PROGRAM_FINALIZE_FPTR = ctypes.CFUNCTYPE( 251 hsa_status_t, # return value 252 hsa_ext_program_t, # program 253 hsa_isa_t, # isa 254 ctypes.c_int32, # call_convention 255 hsa_ext_control_directives_t, # control_directives 256 ctypes.c_char_p, #options 257 hsa_code_object_type_t, #code_object_type 258 _PTR(hsa_code_object_t)) # code_object 259 260# this struct holds function pointers 261class hsa_ext_finalizer_1_00_pfn_t(ctypes.Structure): 262 _fields_ = [ 263 ('hsa_ext_program_create', HSA_EXT_PROGRAM_CREATE_FPTR), 264 ('hsa_ext_program_destroy', HSA_EXT_PROGRAM_DESTROY_FPTR), 265 ('hsa_ext_program_add_module', HSA_EXT_PROGRAM_ADD_MODULE_FPTR), 266 ('hsa_ext_program_iterate_modules', 267 HSA_EXT_PROGRAM_ITERATE_MODULES_FPTR), 268 ('hsa_ext_program_get_info', HSA_EXT_PROGRAM_GET_INFO_FPTR), 269 ('hsa_ext_program_finalize', HSA_EXT_PROGRAM_FINALIZE_FPTR) 270 ] 271 272#------------------------------------------------------------------------------ 273 274 275 276#------------------------------------------------------------------------------ 277# HSA types from hsa_ext_image.h (NOTE: support incomplete) 278 279hsa_ext_image_t = handle_struct 280hsa_ext_image_geometry_t = ctypes.c_int # enum 281hsa_ext_image_channel_type_t = ctypes.c_int # enum 282hsa_ext_image_channel_order_t = ctypes.c_int # enum 283 284class hsa_ext_image_format_t(ctypes.Structure): 285 _fields_ = [ 286 ("channel_type", hsa_ext_image_channel_type_t), 287 ("channel_order", hsa_ext_image_channel_order_t) 288 ] 289 290class hsa_ext_image_descriptor_t(ctypes.Structure): 291 _fields_ = [ 292 ("geometry", hsa_ext_image_geometry_t), 293 ("width", ctypes.c_size_t), 294 ("height", ctypes.c_size_t), 295 ("depth", ctypes.c_size_t), 296 ("array_size", ctypes.c_size_t), 297 ("format", hsa_ext_image_format_t) 298 ] 299 300hsa_ext_image_capability_t = ctypes.c_int # enum 301 302class hsa_ext_image_data_info_t(ctypes.Structure): 303 _fields_ = [ 304 ("size", ctypes.c_size_t), 305 ("alignment", ctypes.c_size_t), 306 ] 307 308class hsa_ext_image_region_t(ctypes.Structure): 309 _fields_ = [ 310 ("offset", hsa_dim3_t), 311 ("offset", hsa_dim3_t), 312 ] 313 314hsa_ext_sampler_t = handle_struct 315hsa_ext_sampler_addressing_mode_t = ctypes.c_int # enum 316hsa_ext_sampler_coordinate_mode_t = ctypes.c_int # enum 317hsa_ext_sampler_filter_mode_t = ctypes.c_int # enum 318 319class hsa_ext_sampler_descriptor_t(ctypes.Structure): 320 _fields_ = [ 321 ("coordinate_mode", hsa_ext_sampler_coordinate_mode_t), 322 ("filter_mode", hsa_ext_sampler_filter_mode_t), 323 ("address_mode", hsa_ext_sampler_addressing_mode_t) 324 ] 325 326#NOTE: Not implemented yet: hsa_ext_images_1_00_pfn_t 327#------------------------------------------------------------------------------ 328 329#------------------------------------------------------------------------------ 330# callbacks that have no related typedef in the hsa include files 331 332HSA_ITER_AGENT_CALLBACK_FUNC = ctypes.CFUNCTYPE( 333 hsa_status_t, # return value 334 hsa_agent_t, # agent 335 ctypes.py_object) # this is a c_void_p used to wrap a python object 336 337HSA_QUEUE_CALLBACK_FUNC = ctypes.CFUNCTYPE( 338 None, # return value 339 hsa_status_t, 340 _PTR(hsa_queue_t), 341 ctypes.py_object) # this is a c_void_p used to wrap a python object 342 343HSA_AGENT_ITERATE_REGIONS_CALLBACK_FUNC = ctypes.CFUNCTYPE( 344 hsa_status_t, # return value 345 hsa_region_t, # region 346 ctypes.py_object) # this is a c_void_p used to wrap a python object 347 348# hsa_status_t (*callback)(hsa_code_object_t code_object, hsa_code_symbol_t symbol, void* data), 349HSA_CODE_OBJECT_ITERATE_SYMBOLS_CALLBACK = ctypes.CFUNCTYPE( 350 hsa_status_t, # return value 351 hsa_code_object_t, 352 hsa_code_symbol_t, 353 ctypes.py_object) # this is a c_void_p used to wrap a python object 354 355# hsa_status_t (*alloc_callback)(size_t size, hsa_callback_data_t data, void **address), 356HSA_ALLOC_CALLBACK_FUNCTION = ctypes.CFUNCTYPE( 357 hsa_status_t, # return value 358 ctypes.c_size_t, 359 hsa_callback_data_t, 360 _PTR(ctypes.c_void_p) # this might need to be a ptr to a py_object 361 ) 362 363void_fn_ptr = ctypes.CFUNCTYPE( 364 None, 365 ctypes.c_void_p) # this might need to be a ptr to a py_object 366 367# hsa_status_t (*callback)(hsa_amd_memory_pool_t memory_pool, void* data) 368HSA_AMD_AGENT_ITERATE_MEMORY_POOLS_CALLBACK = ctypes.CFUNCTYPE( 369 hsa_status_t, 370 hsa_amd_memory_pool_t, 371 ctypes.c_void_p) # this is a c_void_p used to wrap a python object 372 373 374#------------------------------------------------------------------------------ 375 376# Functions used by API calls returning hsa_status_t to check for errors ###### 377 378def _build_reverse_error_warn_maps(): 379 err_map = utils.UniqueDict() 380 warn_map = utils.UniqueDict() 381 382 for name in [name for name in dir(enums) if name.startswith('HSA_')]: 383 code = getattr(enums, name) 384 if 'STATUS_ERROR' in name: 385 err_map[code] = name 386 elif 'STATUS_INFO' in name: 387 warn_map[code] = name 388 else: 389 pass # should we warn here? 390 return err_map, warn_map 391 392ERROR_MAP, WARN_MAP = _build_reverse_error_warn_maps() 393 394 395def _check_error(result, func, arguments): 396 if result != enums.HSA_STATUS_SUCCESS: 397 if result >= enums.HSA_STATUS_ERROR: 398 errname = ERROR_MAP.get(result, "UNKNOWN_HSA_ERROR") 399 msg = "Call to {0} returned {1}".format(func.__name__, errname) 400 raise HsaApiError(result, msg) 401 else: 402 warnname = WARN_MAP.get(result, "UNKNOWN_HSA_INFO") 403 msg = "Call to {0} returned {1}".format(func.__name__, warnname) 404 warnings.warn(msg, HsaWarning) 405 406 407# The API prototypes 408# These are order based on header files. 409API_PROTOTYPES = { 410 411#------------------------------------------------------------------------------ 412# HSA functions from hsa.h, ordered as per header file. 413 414 # hsa_status_t hsa_status_string( 415 # hsa_status_t status, 416 # const char **status_string); 417 'hsa_status_string': { 418 'restype': hsa_status_t, 419 'argtypes': [hsa_status_t, _PTR(ctypes.c_char_p)], 420 'errcheck': _check_error 421 }, 422 423 # hsa_status_t hsa_init(void) 424 'hsa_init': { 425 'restype': hsa_status_t, 426 'argtypes': [], 427 'errcheck': _check_error 428 }, 429 430 # hsa_status_t hsa_shut_down(void) 431 'hsa_shut_down': { 432 'restype': hsa_status_t, 433 'argtypes': [], 434 'errcheck': _check_error 435 }, 436 437 # hsa_status_t hsa_system_get_info(hsa_system_info_t, void*) 438 'hsa_system_get_info': { 439 'restype': hsa_status_t, 440 'argtypes': [hsa_system_info_t, ctypes.c_void_p], 441 'errcheck': _check_error 442 }, 443 444 # hsa_status_t HSA_API hsa_system_extension_supported(uint16_t, uint16_t, 445 # uint16_t, bool *); 446 'hsa_system_extension_supported': { 447 'restype': hsa_status_t, 448 'argtypes': [ctypes.c_uint16, # extension 449 ctypes.c_uint16, # version_major 450 ctypes.c_uint16, # version_minor 451 _PTR(ctypes.c_bool)], # result 452 'errcheck': _check_error 453 }, 454 455 # hsa_status_t hsa_system_get_extension_table(uint16_t, uint16_t, 456 # uint16_t, void *); 457 'hsa_system_get_extension_table': { 458 'restype': hsa_status_t, 459 'argtypes': [ctypes.c_uint16, # extension 460 ctypes.c_uint16, # version_major 461 ctypes.c_uint16, # version_minor 462 ctypes.c_void_p], # result 463 'errcheck': _check_error 464 }, 465 466 # hsa_status_t hsa_agent_get_info(hsa_agent_t, hsa_agent_info_t, void*) 467 'hsa_agent_get_info': { 468 'restype': hsa_status_t, 469 'argtypes': [hsa_agent_t, hsa_agent_info_t, ctypes.c_void_p], 470 'errcheck': _check_error 471 }, 472 473 # hsa_status_t hsa_iterate_agents(hsa_status_t(*)(hsa_agent_t, void*), 474 # void*) 475 'hsa_iterate_agents': { 476 'restype': hsa_status_t, 477 'argtypes': [HSA_ITER_AGENT_CALLBACK_FUNC, ctypes.py_object], 478 'errcheck': _check_error 479 }, 480 481 # hsa_status_t hsa_agent_get_exception_policies(hsa_agent_t agent, 482 # hsa_profile_t profile, 483 # uint16_t *mask); 484 'hsa_agent_get_exception_policies': { 485 'restype': hsa_status_t, 486 'argtypes': [hsa_agent_t, hsa_profile_t, _PTR(ctypes.c_uint16)], 487 'errcheck': _check_error 488 }, 489 490 # hsa_status_t hsa_agent_extension_supported(uint16_t extension, hsa_agent_t agent, 491 # uint16_t version_major, 492 # uint16_t version_minor, bool *result); 493 'hsa_agent_extension_supported': { 494 'restype': hsa_status_t, 495 'argtypes': [ctypes.c_uint16, hsa_agent_t, ctypes.c_uint16, ctypes.c_uint16, 496 _PTR(ctypes.c_bool)], 497 'errcheck': _check_error 498 }, 499 500 #-------------------------------------------------------------------------- 501 # Signals 502 #-------------------------------------------------------------------------- 503 504 # hsa_status_t hsa_signal_create( 505 # hsa_signal_value_t initial_value, 506 # uint32_t agent_count, 507 # const hsa_agent_t *agents, 508 # hsa_signal_t *signal) 509 'hsa_signal_create': { 510 'restype': hsa_status_t, 511 'argtypes': [hsa_signal_value_t, 512 ctypes.c_uint32, 513 _PTR(hsa_agent_t), 514 _PTR(hsa_signal_t)], 515 'errcheck': _check_error 516 }, 517 518 # hsa_status_t hsa_signal_destroy( 519 # hsa_signal_t signal) 520 'hsa_signal_destroy': { 521 'restype': hsa_status_t, 522 'argtypes': [hsa_signal_t], 523 'errcheck': _check_error 524 }, 525 526 # hsa_signal_value_t hsa_signal_load_acquire( 527 # hsa_signal_t signal); 528 'hsa_signal_load_acquire': { 529 'restype': hsa_signal_value_t, 530 'argtypes': [hsa_signal_t], 531 }, 532 533 # hsa_signal_value_t hsa_signal_load_relaxed( 534 # hsa_signal_t signal); 535 'hsa_signal_load_relaxed': { 536 'restype': hsa_signal_value_t, 537 'argtypes': [hsa_signal_t], 538 }, 539 540 # void hsa_signal_store_relaxed( 541 # hsa_signal_t signal, 542 # hsa_signal_value_t value); 543 'hsa_signal_store_relaxed': { 544 'restype': None, 545 'argtypes': [hsa_signal_t, hsa_signal_value_t] 546 }, 547 548 # void hsa_signal_store_release( 549 # hsa_signal_t signal, 550 # hsa_signal_value_t value); 551 'hsa_signal_store_release': { 552 'restype': None, 553 'argtypes': [hsa_signal_t, hsa_signal_value_t], 554 }, 555 556 # hsa_signal_value_t hsa_signal_exchange_acq_rel( 557 # hsa_signal_t signal, 558 # hsa_signal_value_t value); 559 'hsa_signal_exchange_acq_rel': { 560 'restype': hsa_signal_value_t, 561 'argtypes': [hsa_signal_t, hsa_signal_value_t] 562 }, 563 564 # hsa_signal_value_t hsa_signal_exchange_acquire( 565 # hsa_signal_t signal, 566 # hsa_signal_value_t value); 567 'hsa_signal_exchange_acquire': { 568 'restype': hsa_signal_value_t, 569 'argtypes': [hsa_signal_t, hsa_signal_value_t] 570 }, 571 572 # hsa_signal_value_t hsa_signal_exchange_relaxed( 573 # hsa_signal_t signal, 574 # hsa_signal_value_t value); 575 'hsa_signal_exchange_relaxed': { 576 'restype': hsa_signal_value_t, 577 'argtypes': [hsa_signal_t, hsa_signal_value_t] 578 }, 579 580 # hsa_signal_value_t hsa_signal_exchange_release( 581 # hsa_signal_t signal, 582 # hsa_signal_value_t value); 583 'hsa_signal_exchange_release': { 584 'restype': hsa_signal_value_t, 585 'argtypes': [hsa_signal_t, hsa_signal_value_t] 586 }, 587 588 # hsa_signal_value_t hsa_signal_cas_acq_rel( 589 # hsa_signal_t signal, 590 # hsa_signal_value_t expected, 591 # hsa_signal_value_t value); 592 'hsa_signal_cas_acq_rel': { 593 'restype': hsa_signal_value_t, 594 'argtypes': [hsa_signal_t, hsa_signal_value_t, hsa_signal_value_t] 595 }, 596 597 # hsa_signal_value_t hsa_signal_cas_acquire( 598 # hsa_signal_t signal, 599 # hsa_signal_value_t expected, 600 # hsa_signal_value_t value); 601 'hsa_signal_cas_acquire': { 602 'restype': hsa_signal_value_t, 603 'argtypes': [hsa_signal_t, hsa_signal_value_t, hsa_signal_value_t] 604 }, 605 606 # hsa_signal_value_t hsa_signal_cas_relaxed( 607 # hsa_signal_t signal, 608 # hsa_signal_value_t expected, 609 # hsa_signal_value_t value); 610 'hsa_signal_cas_relaxed': { 611 'restype': hsa_signal_value_t, 612 'argtypes': [hsa_signal_t, hsa_signal_value_t, hsa_signal_value_t] 613 }, 614 615 # hsa_signal_value_t hsa_signal_cas_release( 616 # hsa_signal_t signal, 617 # hsa_signal_value_t expected, 618 # hsa_signal_value_t value); 619 'hsa_signal_cas_release': { 620 'restype': hsa_signal_value_t, 621 'argtypes': [hsa_signal_t, hsa_signal_value_t, hsa_signal_value_t] 622 }, 623 624 # void hsa_signal_add_acq_rel( 625 # hsa_signal_t signal, 626 # hsa_signal_value_t value); 627 'hsa_signal_add_acq_rel': { 628 'restype': None, 629 'argtypes': [hsa_signal_t, hsa_signal_value_t] 630 }, 631 632 # void hsa_signal_add_acquire( 633 # hsa_signal_t signal, 634 # hsa_signal_value_t value); 635 'hsa_signal_add_acquire': { 636 'restype': None, 637 'argtypes': [hsa_signal_t, hsa_signal_value_t] 638 }, 639 640 # void hsa_signal_add_relaxed( 641 # hsa_signal_t signal, 642 # hsa_signal_value_t value); 643 'hsa_signal_add_relaxed': { 644 'restype': None, 645 'argtypes': [hsa_signal_t, hsa_signal_value_t] 646 }, 647 648 # void hsa_signal_add_release( 649 # hsa_signal_t signal, 650 # hsa_signal_value_t value); 651 'hsa_signal_add_release': { 652 'restype': None, 653 'argtypes': [hsa_signal_t, hsa_signal_value_t] 654 }, 655 656 # void hsa_signal_subtract_acq_rel( 657 # hsa_signal_t signal, 658 # hsa_signal_value_t value); 659 'hsa_signal_subtract_acq_rel': { 660 'restype': None, 661 'argtypes': [hsa_signal_t, hsa_signal_value_t] 662 }, 663 664 # void hsa_signal_subtract_acquire( 665 # hsa_signal_t signal, 666 # hsa_signal_value_t value); 667 'hsa_signal_subtract_acquire': { 668 'restype': None, 669 'argtypes': [hsa_signal_t, hsa_signal_value_t] 670 }, 671 672 # void hsa_signal_subtract_relaxed( 673 # hsa_signal_t signal, 674 # hsa_signal_value_t value); 675 'hsa_signal_subtract_relaxed': { 676 'restype': None, 677 'argtypes': [hsa_signal_t, hsa_signal_value_t] 678 }, 679 680 # void hsa_signal_subtract_release( 681 # hsa_signal_t signal, 682 # hsa_signal_value_t value); 683 'hsa_signal_subtract_release': { 684 'restype': None, 685 'argtypes': [hsa_signal_t, hsa_signal_value_t] 686 }, 687 688 # void hsa_signal_and_acq_rel( 689 # hsa_signal_t signal, 690 # hsa_signal_value_t value); 691 'hsa_signal_and_acq_rel': { 692 'restype': None, 693 'argtypes': [hsa_signal_t, hsa_signal_value_t] 694 }, 695 696 # void hsa_signal_and_acquire( 697 # hsa_signal_t signal, 698 # hsa_signal_value_t value); 699 'hsa_signal_and_acquire': { 700 'restype': None, 701 'argtypes': [hsa_signal_t, hsa_signal_value_t] 702 }, 703 704 # void hsa_signal_and_relaxed( 705 # hsa_signal_t signal, 706 # hsa_signal_value_t value); 707 'hsa_signal_and_relaxed': { 708 'restype': None, 709 'argtypes': [hsa_signal_t, hsa_signal_value_t] 710 }, 711 712 # void hsa_signal_and_release( 713 # hsa_signal_t signal, 714 # hsa_signal_value_t value); 715 'hsa_signal_and_release': { 716 'restype': None, 717 'argtypes': [hsa_signal_t, hsa_signal_value_t] 718 }, 719 720 # void hsa_signal_or_acq_rel( 721 # hsa_signal_t signal, 722 # hsa_signal_value_t value); 723 'hsa_signal_or_acq_rel': { 724 'restype': None, 725 'argtypes': [hsa_signal_t, 726 hsa_signal_value_t] 727 }, 728 729 # void hsa_signal_or_acquire( 730 # hsa_signal_t signal, 731 # hsa_signal_value_t value); 732 'hsa_signal_or_acquire': { 733 'restype': None, 734 'argtypes': [hsa_signal_t, 735 hsa_signal_value_t] 736 }, 737 738 # void hsa_signal_or_relaxed( 739 # hsa_signal_t signal, 740 # hsa_signal_value_t value); 741 'hsa_signal_or_relaxed': { 742 'restype': None, 743 'argtypes': [hsa_signal_t, 744 hsa_signal_value_t] 745 }, 746 747 # void hsa_signal_or_release( 748 # hsa_signal_t signal, 749 # hsa_signal_value_t value); 750 'hsa_signal_or_release': { 751 'restype': None, 752 'argtypes': [hsa_signal_t, 753 hsa_signal_value_t] 754 }, 755 756 # void hsa_signal_xor_acq_rel( 757 # hsa_signal_t signal, 758 # hsa_signal_value_t value); 759 'hsa_signal_xor_acq_rel': { 760 'restype': None, 761 'argtypes': [hsa_signal_t, hsa_signal_value_t] 762 }, 763 764 # void hsa_signal_xor_acquire( 765 # hsa_signal_t signal, 766 # hsa_signal_value_t value); 767 'hsa_signal_xor_acquire': { 768 'restype': None, 769 'argtypes': [hsa_signal_t, hsa_signal_value_t] 770 }, 771 772 # void hsa_signal_xor_relaxed( 773 # hsa_signal_t signal, 774 # hsa_signal_value_t value); 775 'hsa_signal_xor_relaxed': { 776 'restype': None, 777 'argtypes': [hsa_signal_t, hsa_signal_value_t] 778 }, 779 780 # void hsa_signal_xor_release( 781 # hsa_signal_t signal, 782 # hsa_signal_value_t value); 783 'hsa_signal_xor_release': { 784 'restype': None, 785 'argtypes': [hsa_signal_t, hsa_signal_value_t] 786 }, 787 788 # hsa_signal_value_t HSA_API 789 # hsa_signal_wait_acquire(hsa_signal_t signal, 790 # hsa_signal_condition_t condition, 791 # hsa_signal_value_t compare_value, 792 # uint64_t timeout_hint, 793 # hsa_wait_state_t wait_state_hint); 794 'hsa_signal_wait_acquire': { 795 'restype': hsa_signal_value_t, 796 'argtypes': [hsa_signal_t, 797 hsa_signal_condition_t, 798 hsa_signal_value_t, 799 ctypes.c_uint64, 800 hsa_wait_state_t] 801 }, 802 803 # hsa_signal_value_t hsa_signal_wait_relaxed( 804 # hsa_signal_t signal, 805 # hsa_signal_condition_t condition, 806 # hsa_signal_value_t compare_value, 807 # uint64_t timeout_hint, 808 # hsa_wait_state_t wait_state_hint); 809 'hsa_signal_wait_relaxed': { 810 'restype': hsa_signal_value_t, 811 'argtypes': [hsa_signal_t, 812 hsa_signal_condition_t, 813 hsa_signal_value_t, 814 ctypes.c_uint64, 815 hsa_wait_state_t], 816 }, 817 818 #-------------------------------------------------------------------------- 819 # Queues 820 #-------------------------------------------------------------------------- 821 822 # hsa_status_t HSA_API 823 # hsa_queue_create(hsa_agent_t agent, uint32_t size, hsa_queue_type_t type, 824 # void (*callback)(hsa_status_t status, hsa_queue_t *source, 825 # void *data), 826 # void *data, uint32_t private_segment_size, 827 # uint32_t group_segment_size, hsa_queue_t **queue); 828 'hsa_queue_create': { 829 'restype': hsa_status_t, 830 'argtypes': [hsa_agent_t, 831 ctypes.c_uint32, 832 hsa_queue_type_t, 833 HSA_QUEUE_CALLBACK_FUNC, 834 ctypes.c_void_p, # data 835 ctypes.c_uint32, # private segment size 836 ctypes.c_uint32, # group segment size 837 _PTR(_PTR(hsa_queue_t))], 838 'errcheck': _check_error 839 }, 840 841 # hsa_status_t 842 # hsa_soft_queue_create(hsa_region_t region, uint32_t size, 843 # hsa_queue_type_t type, uint32_t features, 844 # hsa_signal_t doorbell_signal, hsa_queue_t **queue); 845 'hsa_soft_queue_create': { 846 'restype': hsa_status_t, 847 'argtypes': [hsa_region_t, 848 ctypes.c_uint32, 849 hsa_queue_type_t, 850 ctypes.c_uint32, 851 hsa_signal_t, 852 _PTR(_PTR(hsa_queue_t))], 853 'errcheck': _check_error 854 }, 855 856 # hsa_status_t hsa_queue_destroy( 857 # hsa_queue_t *queue) 858 'hsa_queue_destroy': { 859 'restype': hsa_status_t, 860 'argtypes': [_PTR(hsa_queue_t)], 861 'errcheck': _check_error 862 }, 863 864 # hsa_status_t hsa_queue_inactivate(hsa_queue_t *queue); 865 'hsa_queue_inactivate': { 866 'restype': hsa_status_t, 867 'argtypes': [_PTR(hsa_queue_t)], 868 'errcheck': _check_error 869 }, 870 871 # uint64_t hsa_queue_load_read_index_acquire(hsa_queue_t *queue); 872 'hsa_queue_load_read_index_acquire': { 873 'restype': ctypes.c_uint64, 874 'argtypes': [_PTR(hsa_queue_t)] 875 }, 876 877 # uint64_t hsa_queue_load_read_index_relaxed(hsa_queue_t *queue); 878 'hsa_queue_load_read_index_relaxed': { 879 'restype': ctypes.c_uint64, 880 'argtypes': [_PTR(hsa_queue_t)] 881 }, 882 883 # uint64_t hsa_queue_load_write_index_acquire(hsa_queue_t *queue); 884 'hsa_queue_load_write_index_acquire': { 885 'restype': ctypes.c_uint64, 886 'argtypes': [_PTR(hsa_queue_t)] 887 }, 888 889 # uint64_t hsa_queue_load_write_index_relaxed(hsa_queue_t *queue); 890 'hsa_queue_load_write_index_relaxed': { 891 'restype': ctypes.c_uint64, 892 'argtypes': [_PTR(hsa_queue_t)] 893 }, 894 895 # void hsa_queue_store_write_index_relaxed(hsa_queue_t *queue, uint64_t value); 896 'hsa_queue_store_write_index_relaxed': { 897 'restype': None, 898 'argtypes': [_PTR(hsa_queue_t), ctypes.c_uint64] 899 }, 900 901 # void hsa_queue_store_write_index_release(hsa_queue_t *queue, uint64_t value); 902 'hsa_queue_store_write_index_release': { 903 'restype': None, 904 'argtypes': [_PTR(hsa_queue_t), ctypes.c_uint64] 905 }, 906 907 # uint64_t hsa_queue_cas_write_index_acq_rel( 908 # hsa_queue_t *queue, 909 # uint64_t expected, 910 # uint64_t value); 911 'hsa_queue_cas_write_index_acq_rel': { 912 'restype': ctypes.c_uint64, 913 'argtypes': [_PTR(hsa_queue_t), ctypes.c_uint64, ctypes.c_uint64] 914 }, 915 916 # uint64_t hsa_queue_cas_write_index_acquire( 917 # hsa_queue_t *queue, 918 # uint64_t expected, 919 # uint64_t value); 920 'hsa_queue_cas_write_index_acquire': { 921 'restype': ctypes.c_uint64, 922 'argtypes': [_PTR(hsa_queue_t), ctypes.c_uint64, ctypes.c_uint64] 923 }, 924 925 # uint64_t hsa_queue_cas_write_index_relaxed( 926 # hsa_queue_t *queue, 927 # uint64_t expected, 928 # uint64_t value); 929 'hsa_queue_cas_write_index_relaxed': { 930 'restype': ctypes.c_uint64, 931 'argtypes': [_PTR(hsa_queue_t), ctypes.c_uint64, ctypes.c_uint64] 932 }, 933 934 # uint64_t hsa_queue_cas_write_index_release( 935 # hsa_queue_t *queue, 936 # uint64_t expected, 937 # uint64_t value); 938 'hsa_queue_cas_write_index_release': { 939 'restype': ctypes.c_uint64, 940 'argtypes': [_PTR(hsa_queue_t), ctypes.c_uint64, ctypes.c_uint64] 941 }, 942 943 # uint64_t hsa_queue_add_write_index_acq_rel( 944 # hsa_queue_t *queue, 945 # uint64_t value); 946 'hsa_queue_add_write_index_acq_rel': { 947 'restype': ctypes.c_uint64, 948 'argtypes': [_PTR(hsa_queue_t), ctypes.c_uint64] 949 }, 950 951 # uint64_t hsa_queue_add_write_index_acquire( 952 # hsa_queue_t *queue, 953 # uint64_t value); 954 'hsa_queue_add_write_index_acquire': { 955 'restype': ctypes.c_uint64, 956 'argtypes': [_PTR(hsa_queue_t), ctypes.c_uint64] 957 }, 958 959 # uint64_t hsa_queue_add_write_index_relaxed( 960 # hsa_queue_t *queue, 961 # uint64_t value); 962 'hsa_queue_add_write_index_relaxed': { 963 'restype': ctypes.c_uint64, 964 'argtypes': [_PTR(hsa_queue_t), ctypes.c_uint64] 965 }, 966 967 # uint64_t hsa_queue_add_write_index_release( 968 # hsa_queue_t *queue, 969 # uint64_t value); 970 'hsa_queue_add_write_index_release': { 971 'restype': ctypes.c_uint64, 972 'argtypes': [_PTR(hsa_queue_t), ctypes.c_uint64] 973 }, 974 975 # void hsa_queue_store_read_index_relaxed( 976 # hsa_queue_t *queue, 977 # uint64_t value); 978 'hsa_queue_store_read_index_relaxed': { 979 'restype': None, 980 'argtypes': [_PTR(hsa_queue_t), ctypes.c_uint64] 981 }, 982 983 # void hsa_queue_store_read_index_release( 984 # hsa_queue_t *queue, 985 # uint64_t value); 986 'hsa_queue_store_read_index_release': { 987 'restype': None, 988 'argtypes': [_PTR(hsa_queue_t), ctypes.c_uint64] 989 }, 990 991 #-------------------------------------------------------------------------- 992 # Memory 993 #-------------------------------------------------------------------------- 994 995 # hsa_status_t hsa_region_get_info( 996 # hsa_region_t region, 997 # hsa_region_info_t attribute, 998 # void *value); 999 'hsa_region_get_info': { 1000 'restype': hsa_status_t, 1001 'argtypes': [hsa_region_t, hsa_region_info_t, ctypes.c_void_p], 1002 'errcheck': _check_error, 1003 }, 1004 1005 # hsa_status_t hsa_agent_iterate_regions( 1006 # hsa_agent_t agent, 1007 # hsa_status_t (*callback)(hsa_region_t region, void *data), 1008 # void *data); 1009 'hsa_agent_iterate_regions': { 1010 'restype': hsa_status_t, 1011 'argtypes': [hsa_agent_t, 1012 HSA_AGENT_ITERATE_REGIONS_CALLBACK_FUNC, 1013 ctypes.py_object], 1014 'errcheck': _check_error 1015 }, 1016 1017 # hsa_status_t hsa_memory_allocate( 1018 # hsa_region_t region, 1019 # size_t size, 1020 # void **ptr); 1021 'hsa_memory_allocate': { 1022 'restype': hsa_status_t, 1023 'argtypes': [hsa_region_t, ctypes.c_size_t, _PTR(ctypes.c_void_p)], 1024 'errcheck': _check_error 1025 }, 1026 1027 # hsa_status_t hsa_memory_free( 1028 # void *ptr); 1029 'hsa_memory_free': { 1030 'restype': hsa_status_t, 1031 'argtypes': [ctypes.c_void_p], 1032 'errcheck': _check_error 1033 }, 1034 1035 # hsa_status_t HSA_API hsa_memory_copy( 1036 # void * dst, 1037 # const void * src, 1038 # size_t size); 1039 'hsa_memory_copy': { 1040 'restype': hsa_status_t, 1041 'argtypes': [ctypes.c_void_p, ctypes.c_void_p, ctypes.c_size_t], 1042 'errcheck': _check_error 1043 }, 1044 1045 # hsa_status_t HSA_API hsa_memory_assign_agent(void *ptr, 1046 # hsa_agent_t agent, 1047 # hsa_access_permission_t access); 1048 'hsa_memory_assign_agent': { 1049 'restype': hsa_status_t, 1050 'argtypes': [ctypes.c_void_p, hsa_agent_t, hsa_access_permission_t], 1051 'errcheck': _check_error 1052 }, 1053 1054 # hsa_status_t hsa_memory_register( 1055 # void *address, 1056 # size_t size); 1057 'hsa_memory_register': { 1058 'restype': hsa_status_t, 1059 'argtypes': [ctypes.c_void_p, ctypes.c_size_t], 1060 'errcheck': _check_error 1061 }, 1062 1063 # hsa_status_t hsa_memory_deregister( 1064 # void *address, 1065 # size_t size); 1066 'hsa_memory_deregister': { 1067 'restype': hsa_status_t, 1068 'argtypes': [ctypes.c_void_p, ctypes.c_size_t], 1069 'errcheck': _check_error 1070 }, 1071 1072 #-------------------------------------------------------------------------- 1073 # Code Object functions 1074 #-------------------------------------------------------------------------- 1075 1076 # hsa_status_t HSA_API hsa_isa_from_name(const char* name, 1077 # hsa_isa_t* isa); 1078 'hsa_isa_from_name': { 1079 'restype': hsa_status_t, 1080 'argtypes': [ctypes.c_char_p, _PTR(hsa_isa_t)], 1081 'errcheck': _check_error 1082 }, 1083 1084 # hsa_status_t HSA_API hsa_isa_get_info(hsa_isa_t isa, 1085 # hsa_isa_info_t attribute, 1086 # uint32_t index, 1087 # void* value); 1088 'hsa_isa_get_info': { 1089 'restype': hsa_status_t, 1090 'argtypes': [hsa_isa_t, hsa_isa_info_t, ctypes.c_void_p], 1091 'errcheck': _check_error 1092 }, 1093 1094 # hsa_status_t HSA_API hsa_isa_compatible(hsa_isa_t code_object_isa, 1095 # hsa_isa_t agent_isa, 1096 # bool* result); 1097 'hsa_isa_compatible': { 1098 'restype': hsa_status_t, 1099 'argtypes': [hsa_isa_t, hsa_isa_t, _PTR(ctypes.c_bool)], 1100 'errcheck': _check_error 1101 }, 1102 1103 # hsa_status_t HSA_API hsa_code_object_serialize( 1104 # hsa_code_object_t code_object, 1105 # hsa_status_t (*alloc_callback)(size_t size, 1106 # hsa_callback_data_t data, void **address), 1107 # hsa_callback_data_t callback_data, 1108 # const char *options, 1109 # void **serialized_code_object, 1110 # size_t *serialized_code_object_size); 1111 'hsa_code_object_serialize': { 1112 'restype': hsa_status_t, 1113 'argtypes': [HSA_ALLOC_CALLBACK_FUNCTION, 1114 hsa_callback_data_t, 1115 _PTR(ctypes.c_void_p), 1116 hsa_callback_data_t, 1117 ctypes.c_char_p, 1118 _PTR(ctypes.c_void_p), 1119 _PTR(ctypes.c_size_t)], 1120 'errcheck': _check_error 1121 }, 1122 1123 # hsa_status_t HSA_API hsa_code_object_deserialize( 1124 # void *serialized_code_object, 1125 # size_t serialized_code_object_size, 1126 # const char *options, 1127 # hsa_code_object_t *code_object); 1128 'hsa_code_object_deserialize': { 1129 'restype': hsa_status_t, 1130 'argtypes': [ctypes.c_void_p, 1131 ctypes.c_size_t, 1132 ctypes.c_char_p, 1133 _PTR(hsa_code_object_t)], 1134 'errcheck': _check_error 1135 }, 1136 1137 # hsa_status_t HSA_API hsa_code_object_destroy( 1138 # hsa_code_object_t code_object); 1139 'hsa_code_object_destroy': { 1140 'restype': hsa_status_t, 1141 'argtypes': [hsa_code_object_t], 1142 'errcheck': _check_error 1143 }, 1144 1145 # hsa_status_t HSA_API hsa_code_object_get_info( 1146 # hsa_code_object_t code_object, 1147 # hsa_code_object_info_t attribute, 1148 # void *value); 1149 'hsa_code_object_get_info': { 1150 'restype': hsa_status_t, 1151 'argtypes': [hsa_code_object_t, 1152 hsa_code_object_info_t, 1153 ctypes.c_void_p 1154 ], 1155 'errcheck': _check_error 1156 }, 1157 1158 # hsa_status_t HSA_API hsa_code_object_get_symbol( 1159 # hsa_code_object_t code_object, 1160 # const char *symbol_name, 1161 # hsa_code_symbol_t *symbol); 1162 'hsa_code_object_get_symbol': { 1163 'restype': hsa_status_t, 1164 'argtypes': [hsa_code_object_t, 1165 ctypes.c_char_p, 1166 _PTR(hsa_code_symbol_t) 1167 ], 1168 'errcheck': _check_error 1169 }, 1170 1171 # hsa_status_t HSA_API hsa_code_symbol_get_info( 1172 # hsa_code_symbol_t code_symbol, 1173 # hsa_code_symbol_info_t attribute, 1174 # void *value); 1175 'hsa_code_symbol_get_info': { 1176 'restype': hsa_status_t, 1177 'argtypes': [hsa_code_symbol_t, 1178 hsa_code_symbol_info_t, 1179 ctypes.c_void_p 1180 ], 1181 'errcheck': _check_error 1182 }, 1183 1184 # hsa_status_t HSA_API hsa_code_object_iterate_symbols( 1185 # hsa_code_object_t code_object, 1186 # hsa_status_t (*callback)(hsa_code_object_t code_object, hsa_code_symbol_t symbol, void* data), 1187 # void* data); 1188 'hsa_code_object_iterate_symbols': { 1189 'restype': hsa_status_t, 1190 'argtypes': [hsa_code_object_t, 1191 HSA_CODE_OBJECT_ITERATE_SYMBOLS_CALLBACK, 1192 ctypes.c_void_p 1193 ], 1194 'errcheck': _check_error 1195 }, 1196 1197 #-------------------------------------------------------------------------- 1198 # Executable functions 1199 #-------------------------------------------------------------------------- 1200 1201 # hsa_status_t HSA_API hsa_executable_create( 1202 # hsa_profile_t profile, 1203 # hsa_executable_state_t executable_state, 1204 # const char *options, 1205 # hsa_executable_t *executable); 1206 1207 "hsa_executable_create": { 1208 'restype': hsa_status_t, 1209 'argtypes': [hsa_profile_t, 1210 hsa_executable_state_t, 1211 ctypes.c_char_p, 1212 ctypes.POINTER(hsa_executable_t)], 1213 'errcheck': _check_error, 1214 }, 1215 1216 # hsa_status_t HSA_API hsa_executable_destroy( 1217 # hsa_executable_t executable); 1218 1219 "hsa_executable_destroy": { 1220 'errcheck': _check_error, 1221 'restype': hsa_status_t, 1222 'argtypes': [ 1223 hsa_executable_t, 1224 ], 1225 }, 1226 1227 # hsa_status_t HSA_API hsa_executable_load_code_object( 1228 # hsa_executable_t executable, 1229 # hsa_agent_t agent, 1230 # hsa_code_object_t code_object, 1231 # const char *options); 1232 1233 "hsa_executable_load_code_object": { 1234 'errcheck': _check_error, 1235 'restype': hsa_status_t, 1236 'argtypes': [ 1237 hsa_executable_t, 1238 hsa_agent_t, 1239 hsa_code_object_t, 1240 ctypes.c_char_p, 1241 ], 1242 }, 1243 1244 # hsa_status_t HSA_API hsa_executable_freeze( 1245 # hsa_executable_t executable, 1246 # const char *options); 1247 1248 "hsa_executable_freeze": { 1249 'errcheck': _check_error, 1250 'restype': hsa_status_t, 1251 'argtypes': [ 1252 hsa_executable_t, 1253 ctypes.c_char_p, 1254 ], 1255 }, 1256 1257 # hsa_status_t HSA_API hsa_executable_get_info( 1258 # hsa_executable_t executable, 1259 # hsa_executable_info_t attribute, 1260 # void *value); 1261 "hsa_executable_get_info": { 1262 'errcheck': _check_error, 1263 'restype': hsa_status_t, 1264 'argtypes': [ 1265 hsa_executable_t, 1266 hsa_executable_info_t, 1267 ctypes.c_void_p 1268 ], 1269 }, 1270 1271 # hsa_status_t HSA_API hsa_executable_global_variable_define( 1272 # hsa_executable_t executable, 1273 # const char *variable_name, 1274 # void *address); 1275 "hsa_executable_global_variable_define": { 1276 'restype': hsa_status_t, 1277 'argtypes': [hsa_executable_t, 1278 ctypes.c_char_p, 1279 ctypes.c_void_p], 1280 'errcheck': _check_error, 1281 }, 1282 1283 # hsa_status_t HSA_API hsa_executable_agent_global_variable_define( 1284 # hsa_executable_t executable, 1285 # hsa_agent_t agent, 1286 # const char *variable_name, 1287 # void *address); 1288 "hsa_executable_agent_global_variable_define": { 1289 'restype': hsa_status_t, 1290 'argtypes': [hsa_executable_t, 1291 hsa_agent_t, 1292 ctypes.c_char_p, 1293 ctypes.c_void_p], 1294 'errcheck': _check_error, 1295 }, 1296 1297 # hsa_status_t HSA_API hsa_executable_readonly_variable_define( 1298 # hsa_executable_t executable, 1299 # hsa_agent_t agent, 1300 # const char *variable_name, 1301 # void *address); 1302 "hsa_executable_readonly_variable_define": { 1303 'restype': hsa_status_t, 1304 'argtypes': [hsa_executable_t, 1305 hsa_agent_t, 1306 ctypes.c_char_p, 1307 ctypes.c_void_p], 1308 'errcheck': _check_error, 1309 }, 1310 1311 # hsa_status_t HSA_API hsa_executable_validate( 1312 # hsa_executable_t executable, 1313 # uint32_t* result); 1314 "hsa_executable_validate": { 1315 'restype': hsa_status_t, 1316 'argtypes': [hsa_executable_t, 1317 _PTR(ctypes.c_uint32)], 1318 'errcheck': _check_error, 1319 }, 1320 1321 # hsa_status_t HSA_API hsa_executable_get_symbol( 1322 # hsa_executable_t executable, 1323 # const char *module_name, 1324 # const char *symbol_name, 1325 # hsa_agent_t agent, 1326 # int32_t call_convention, 1327 # hsa_executable_symbol_t *symbol); 1328 "hsa_executable_get_symbol": { 1329 'errcheck': _check_error, 1330 'restype': hsa_status_t, 1331 'argtypes': [ 1332 hsa_executable_t, 1333 ctypes.c_char_p, # module_name (must be NULL for program linkage) 1334 ctypes.c_char_p, # symbol_name 1335 hsa_agent_t, 1336 ctypes.c_int32, 1337 ctypes.POINTER(hsa_executable_symbol_t), 1338 ], 1339 }, 1340 1341 # hsa_status_t HSA_API hsa_executable_symbol_get_info( 1342 # hsa_executable_symbol_t executable_symbol, 1343 # hsa_executable_symbol_info_t attribute, 1344 # void *value); 1345 "hsa_executable_symbol_get_info": { 1346 'errcheck': _check_error, 1347 'restype': hsa_status_t, 1348 'argtypes': [ 1349 hsa_executable_symbol_t, 1350 hsa_executable_symbol_info_t, 1351 ctypes.c_void_p, 1352 ], 1353 }, 1354 1355 1356 #hsa_status_t HSA_API hsa_executable_iterate_symbols( 1357 # hsa_executable_t executable, 1358 # hsa_status_t (*callback)(hsa_executable_t executable, hsa_executable_symbol_t symbol, void* data), 1359 # void* data); 1360 "hsa_executable_iterate_symbols": { 1361 'errcheck': _check_error, 1362 'restype': hsa_status_t, 1363 'argtypes': [ 1364 hsa_executable_symbol_t, 1365 hsa_executable_symbol_info_t, 1366 ctypes.c_void_p, 1367 ], 1368 }, 1369 1370 1371 #-------------------------------------------------------------------------- 1372 # AMD extensions from hsa_ext_amd.h 1373 #-------------------------------------------------------------------------- 1374 1375 # hsa_status_t HSA_API hsa_amd_coherency_get_type(hsa_agent_t agent, 1376 # hsa_amd_coherency_type_t* type); 1377 1378 "hsa_amd_coherency_get_type": { 1379 'errcheck': _check_error, 1380 'restype': hsa_status_t, 1381 'argtypes': [ 1382 hsa_agent_t, 1383 _PTR(hsa_amd_coherency_type_t), 1384 ], 1385 }, 1386 1387 # hsa_status_t HSA_API hsa_amd_coherency_set_type(hsa_agent_t agent, 1388 # hsa_amd_coherency_type_t type); 1389 "hsa_amd_coherency_get_type": { 1390 'errcheck': _check_error, 1391 'restype': hsa_status_t, 1392 'argtypes': [ 1393 hsa_agent_t, 1394 hsa_amd_coherency_type_t, 1395 ], 1396 }, 1397 1398 # hsa_status_t HSA_API 1399 # hsa_amd_profiling_set_profiler_enabled(hsa_queue_t* queue, int enable); 1400 "hsa_amd_profiling_set_profiler_enabled": { 1401 'errcheck': _check_error, 1402 'restype': hsa_status_t, 1403 'argtypes': [ 1404 _PTR(hsa_queue_t), 1405 ctypes.c_int, 1406 ], 1407 }, 1408 1409 # hsa_status_t HSA_API hsa_amd_profiling_get_dispatch_time( 1410 # hsa_agent_t agent, hsa_signal_t signal, 1411 # hsa_amd_profiling_dispatch_time_t* time); 1412 "hsa_amd_profiling_get_dispatch_time": { 1413 'errcheck': _check_error, 1414 'restype': hsa_status_t, 1415 'argtypes': [ 1416 hsa_agent_t, 1417 hsa_signal_t, 1418 _PTR(hsa_amd_profiling_dispatch_time_t) 1419 ], 1420 }, 1421 1422 # hsa_status_t HSA_API 1423 # hsa_amd_profiling_convert_tick_to_system_domain(hsa_agent_t agent, 1424 # uint64_t agent_tick, 1425 # uint64_t* system_tick); 1426 "hsa_amd_profiling_convert_tick_to_system_domain": { 1427 'errcheck': _check_error, 1428 'restype': hsa_status_t, 1429 'argtypes': [ 1430 ctypes.c_uint64, 1431 _PTR(ctypes.c_uint64) 1432 ], 1433 }, 1434 1435 # hsa_status_t HSA_API 1436 # hsa_amd_signal_async_handler(hsa_signal_t signal, 1437 # hsa_signal_condition_t cond, 1438 # hsa_signal_value_t value, 1439 # hsa_amd_signal_handler handler, void* arg); 1440 "hsa_amd_signal_async_handler": { 1441 'errcheck': _check_error, 1442 'restype': hsa_status_t, 1443 'argtypes': [ 1444 hsa_signal_t, 1445 hsa_signal_condition_t, 1446 hsa_signal_value_t, 1447 hsa_amd_signal_handler, 1448 ctypes.c_void_p, 1449 ], 1450 }, 1451 1452 #hsa_amd_async_function(void (*callback)(void* arg), void* arg); 1453 "hsa_amd_async_function": { 1454 'errcheck': _check_error, 1455 'restype': hsa_status_t, 1456 'argtypes': [ 1457 ctypes.POINTER(void_fn_ptr), 1458 ctypes.c_void_p, 1459 ], 1460 }, 1461 1462 #uint32_t HSA_API 1463 #hsa_amd_signal_wait_any(uint32_t signal_count, hsa_signal_t* signals, 1464 # hsa_signal_condition_t* conds, 1465 # hsa_signal_value_t* values, uint64_t timeout_hint, 1466 # hsa_wait_state_t wait_hint, 1467 # hsa_signal_value_t* satisfying_value); 1468 "hsa_amd_signal_wait_any": { 1469 'errcheck': _check_error, 1470 'restype': ctypes.c_uint32, 1471 'argtypes': [ 1472 ctypes.c_uint32, 1473 _PTR(hsa_signal_t), 1474 _PTR(hsa_signal_condition_t), 1475 _PTR(hsa_signal_value_t), 1476 ctypes.c_uint64, 1477 hsa_wait_state_t, 1478 _PTR(hsa_signal_value_t), 1479 ], 1480 }, 1481 1482 # hsa_status_t HSA_API hsa_amd_image_get_info_max_dim(hsa_agent_t agent, 1483 # hsa_agent_info_t attribute, 1484 # void* value); 1485 "hsa_amd_image_get_info_max_dim": { 1486 'errcheck': _check_error, 1487 'restype': hsa_status_t, 1488 'argtypes': [ 1489 hsa_agent_t, 1490 hsa_agent_info_t, 1491 ctypes.c_void_p, 1492 ], 1493 }, 1494 1495 # hsa_status_t HSA_API hsa_amd_queue_cu_set_mask(const hsa_queue_t* queue, 1496 # uint32_t num_cu_mask_count, 1497 # const uint32_t* cu_mask); 1498 "hsa_amd_queue_cu_set_mask": { 1499 'errcheck': _check_error, 1500 'restype': hsa_status_t, 1501 'argtypes': [ 1502 _PTR(hsa_queue_t), 1503 ctypes.c_uint32, 1504 _PTR(ctypes.c_uint32) 1505 ], 1506 }, 1507 1508 # hsa_status_t HSA_API 1509 # hsa_amd_memory_pool_get_info(hsa_amd_memory_pool_t memory_pool, 1510 # hsa_amd_memory_pool_info_t attribute, 1511 # void* value); 1512 "hsa_amd_memory_pool_get_info": { 1513 'errcheck': _check_error, 1514 'restype': hsa_status_t, 1515 'argtypes': [ 1516 hsa_amd_memory_pool_t, 1517 hsa_amd_memory_pool_info_t, 1518 ctypes.c_void_p 1519 ], 1520 }, 1521 1522 # hsa_status_t HSA_API hsa_amd_agent_iterate_memory_pools( 1523 # hsa_agent_t agent, 1524 # hsa_status_t (*callback)(hsa_amd_memory_pool_t memory_pool, void* data), 1525 # void* data); 1526 "hsa_amd_agent_iterate_memory_pools": { 1527 'errcheck': _check_error, 1528 'restype': hsa_status_t, 1529 'argtypes': [ 1530 hsa_agent_t, 1531 HSA_AMD_AGENT_ITERATE_MEMORY_POOLS_CALLBACK, 1532 ctypes.c_void_p 1533 ], 1534 }, 1535 1536 # hsa_status_t HSA_API hsa_amd_memory_pool_allocate 1537 # (hsa_amd_memory_pool_t memory_pool, size_t size, 1538 # uint32_t flags, void** ptr); 1539 "hsa_amd_memory_pool_allocate": { 1540 'errcheck': _check_error, 1541 'restype': hsa_status_t, 1542 'argtypes': [ 1543 hsa_amd_memory_pool_t, 1544 ctypes.c_size_t, 1545 ctypes.c_uint32, 1546 _PTR(ctypes.c_void_p) 1547 ], 1548 }, 1549 1550 # hsa_status_t HSA_API hsa_amd_memory_pool_free(void* ptr); 1551 "hsa_amd_memory_pool_free": { 1552 'errcheck': _check_error, 1553 'restype': hsa_status_t, 1554 'argtypes': [ 1555 ctypes.c_void_p 1556 ], 1557 }, 1558 1559 # hsa_status_t HSA_API hsa_amd_memory_async_copy(void* dst, 1560 # hsa_agent_t dst_agent, const void* src, 1561 # hsa_agent_t src_agent, size_t size, 1562 # uint32_t num_dep_signals, 1563 # const hsa_signal_t* dep_signals, 1564 # hsa_signal_t completion_signal); 1565 "hsa_amd_memory_async_copy": { 1566 'errcheck': _check_error, 1567 'restype': hsa_status_t, 1568 'argtypes': [ 1569 ctypes.c_void_p, 1570 hsa_agent_t, 1571 ctypes.c_void_p, 1572 hsa_agent_t, 1573 ctypes.c_size_t, 1574 ctypes.c_uint32, 1575 _PTR(hsa_signal_t), 1576 hsa_signal_t 1577 ], 1578 }, 1579 1580 # hsa_status_t HSA_API hsa_amd_agent_memory_pool_get_info( 1581 # hsa_agent_t agent, hsa_amd_memory_pool_t memory_pool, 1582 # hsa_amd_agent_memory_pool_info_t attribute, void* value); 1583 "hsa_amd_agent_memory_pool_get_info": { 1584 'errcheck': _check_error, 1585 'restype': hsa_status_t, 1586 'argtypes': [ 1587 hsa_agent_t, 1588 hsa_amd_memory_pool_t, 1589 hsa_amd_agent_memory_pool_info_t, 1590 ctypes.c_void_p 1591 ], 1592 }, 1593 1594 1595 # hsa_status_t HSA_API 1596 # hsa_amd_agents_allow_access(uint32_t num_agents, const hsa_agent_t* agents, 1597 # const uint32_t* flags, const void* ptr); 1598 "hsa_amd_agents_allow_access": { 1599 'errcheck': _check_error, 1600 'restype': hsa_status_t, 1601 'argtypes': [ 1602 ctypes.c_uint32, 1603 _PTR(hsa_agent_t), 1604 _PTR(ctypes.c_uint32), 1605 ctypes.c_void_p 1606 ], 1607 }, 1608 1609 1610 # hsa_status_t HSA_API 1611 # hsa_amd_memory_pool_can_migrate(hsa_amd_memory_pool_t src_memory_pool, 1612 # hsa_amd_memory_pool_t dst_memory_pool, 1613 # bool* result); 1614 "hsa_amd_memory_pool_can_migrate": { 1615 'errcheck': _check_error, 1616 'restype': hsa_status_t, 1617 'argtypes': [ 1618 hsa_amd_memory_pool_t, 1619 hsa_amd_memory_pool_t, 1620 _PTR(ctypes.c_bool) 1621 ], 1622 }, 1623 1624 1625 # hsa_status_t HSA_API hsa_amd_memory_migrate(const void* ptr, 1626 # hsa_amd_memory_pool_t memory_pool, 1627 # uint32_t flags); 1628 "hsa_amd_memory_migrate": { 1629 'errcheck': _check_error, 1630 'restype': hsa_status_t, 1631 'argtypes': [ 1632 ctypes.c_void_p, 1633 hsa_amd_memory_pool_t, 1634 ctypes.c_uint32 1635 ], 1636 }, 1637 1638 1639 # hsa_status_t HSA_API hsa_amd_memory_lock(void* host_ptr, size_t size, 1640 # hsa_agent_t* agents, int num_agent, 1641 # void** agent_ptr); 1642 "hsa_amd_memory_lock": { 1643 'errcheck': _check_error, 1644 'restype': hsa_status_t, 1645 'argtypes': [ 1646 ctypes.c_void_p, 1647 ctypes.c_size_t, 1648 _PTR(hsa_agent_t), 1649 ctypes.c_int, 1650 _PTR(ctypes.c_void_p) 1651 ], 1652 }, 1653 1654 1655 # hsa_status_t HSA_API hsa_amd_memory_unlock(void* host_ptr); 1656 "hsa_amd_memory_unlock": { 1657 'errcheck': _check_error, 1658 'restype': hsa_status_t, 1659 'argtypes': [ 1660 ctypes.c_void_p 1661 ], 1662 }, 1663 1664 1665 # hsa_status_t HSA_API 1666 # hsa_amd_memory_fill(void* ptr, uint32_t value, size_t count); 1667 "hsa_amd_memory_unlock": { 1668 'errcheck': _check_error, 1669 'restype': hsa_status_t, 1670 'argtypes': [ 1671 ctypes.c_void_p 1672 ], 1673 }, 1674 1675 # hsa_status_t HSA_API hsa_amd_interop_map_buffer(uint32_t num_agents, 1676 # hsa_agent_t* agents, 1677 # int interop_handle, 1678 # uint32_t flags, 1679 # size_t* size, 1680 # void** ptr, 1681 # size_t* metadata_size, 1682 # const void** metadata); 1683 "hsa_amd_interop_map_buffer": { 1684 'errcheck': _check_error, 1685 'restype': hsa_status_t, 1686 'argtypes': [ 1687 ctypes.c_uint32, 1688 _PTR(hsa_agent_t), 1689 ctypes.c_int, 1690 ctypes.c_uint32, 1691 _PTR(ctypes.c_size_t), 1692 _PTR(ctypes.c_void_p), 1693 _PTR(ctypes.c_size_t), 1694 _PTR(ctypes.c_void_p), 1695 ], 1696 }, 1697 1698 1699 # hsa_status_t HSA_API hsa_amd_interop_unmap_buffer(void* ptr); 1700 "hsa_amd_interop_map_buffer": { 1701 'errcheck': _check_error, 1702 'restype': hsa_status_t, 1703 'argtypes': [ 1704 _PTR(ctypes.c_void_p), 1705 ], 1706 }, 1707 1708 1709 # hsa_status_t HSA_API hsa_amd_image_create( 1710 # hsa_agent_t agent, 1711 # const hsa_ext_image_descriptor_t *image_descriptor, 1712 # const hsa_amd_image_descriptor_t *image_layout, 1713 # const void *image_data, 1714 # hsa_access_permission_t access_permission, 1715 # hsa_ext_image_t *image 1716 # ); 1717 "hsa_amd_image_create": { 1718 'errcheck': _check_error, 1719 'restype': hsa_status_t, 1720 'argtypes': [ 1721 hsa_agent_t, 1722 _PTR(hsa_ext_image_descriptor_t), 1723 _PTR(hsa_amd_image_descriptor_t), 1724 ctypes.c_void_p, 1725 hsa_access_permission_t, 1726 hsa_ext_image_t 1727 ], 1728 }, 1729 1730 #-------------------------------------------------------------------------- 1731 # Functions from hsa_ext_finalize.h 1732 # NOTE: To access these functions use the hsa_ext_finalizer_1_00_pfn_t 1733 # struct. 1734 #-------------------------------------------------------------------------- 1735 1736} 1737