//////////////////////////////////////////////////////////////////////////////// // // The University of Illinois/NCSA // Open Source License (NCSA) // // Copyright (c) 2014-2020, Advanced Micro Devices, Inc. All rights reserved. // // Developed by: // // AMD Research and AMD HSA Software Development // // Advanced Micro Devices, Inc. // // www.amd.com // // Permission is hereby granted, free of charge, to any person obtaining a copy // of this software and associated documentation files (the "Software"), to // deal with the Software without restriction, including without limitation // the rights to use, copy, modify, merge, publish, distribute, sublicense, // and/or sell copies of the Software, and to permit persons to whom the // Software is furnished to do so, subject to the following conditions: // // - Redistributions of source code must retain the above copyright notice, // this list of conditions and the following disclaimers. // - Redistributions in binary form must reproduce the above copyright // notice, this list of conditions and the following disclaimers in // the documentation and/or other materials provided with the distribution. // - Neither the names of Advanced Micro Devices, Inc, // nor the names of its contributors may be used to endorse or promote // products derived from this Software without specific prior written // permission. // // THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR // IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, // FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL // THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR // OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, // ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER // DEALINGS WITH THE SOFTWARE. // //////////////////////////////////////////////////////////////////////////////// // HSA C to C++ interface implementation. // This file does argument checking and conversion to C++. #include #include #include #include #include "core/inc/runtime.h" #include "core/inc/agent.h" #include "core/inc/host_queue.h" #include "core/inc/isa.h" #include "core/inc/memory_region.h" #include "core/inc/queue.h" #include "core/inc/signal.h" #include "core/inc/cache.h" #include "core/inc/amd_elf_image.hpp" #include "core/inc/amd_hsa_loader.hpp" #include "core/inc/amd_loader_context.hpp" #include "core/inc/hsa_ven_amd_loader_impl.h" #include "inc/hsa_ven_amd_aqlprofile.h" #include "core/inc/hsa_ext_amd_impl.h" namespace rocr { using namespace amd::hsa; template struct ValidityError; template <> struct ValidityError { enum { kValue = HSA_STATUS_ERROR_INVALID_SIGNAL }; }; template <> struct ValidityError { enum { kValue = HSA_STATUS_ERROR_INVALID_SIGNAL_GROUP }; }; template <> struct ValidityError { enum { kValue = HSA_STATUS_ERROR_INVALID_AGENT }; }; template <> struct ValidityError { enum { kValue = HSA_STATUS_ERROR_INVALID_REGION }; }; template <> struct ValidityError { enum { kValue = HSA_STATUS_ERROR_INVALID_QUEUE }; }; template <> struct ValidityError { enum { kValue = HSA_STATUS_ERROR_INVALID_CACHE }; }; template <> struct ValidityError { enum { kValue = HSA_STATUS_ERROR_INVALID_ISA }; }; template struct ValidityError { enum { kValue = ValidityError::kValue }; }; #define IS_BAD_PTR(ptr) \ do { \ if ((ptr) == nullptr) return HSA_STATUS_ERROR_INVALID_ARGUMENT; \ } while (false) #define IS_BAD_PROFILE(profile) \ do { \ if (profile != HSA_PROFILE_BASE && \ profile != HSA_PROFILE_FULL) { \ return HSA_STATUS_ERROR_INVALID_ARGUMENT; \ } \ } while (false) #define IS_BAD_EXECUTABLE_STATE(executable_state) \ do { \ if (executable_state != HSA_EXECUTABLE_STATE_FROZEN && \ executable_state != HSA_EXECUTABLE_STATE_UNFROZEN) { \ return HSA_STATUS_ERROR_INVALID_ARGUMENT; \ } \ } while (false) #define IS_BAD_ROUNDING_MODE(rounding_mode) \ do { \ if (rounding_mode != HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT && \ rounding_mode != HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO && \ rounding_mode != HSA_DEFAULT_FLOAT_ROUNDING_MODE_NEAR) { \ return HSA_STATUS_ERROR_INVALID_ARGUMENT; \ } \ } while (false) #define IS_BAD_FP_TYPE(fp_type) \ do { \ if (fp_type != HSA_FP_TYPE_16 && \ fp_type != HSA_FP_TYPE_32 && \ fp_type != HSA_FP_TYPE_64) { \ return HSA_STATUS_ERROR_INVALID_ARGUMENT; \ } \ } while (false) #define IS_BAD_FLUSH_MODE(flush_mode) \ do { \ if (flush_mode != HSA_FLUSH_MODE_FTZ && \ flush_mode != HSA_FLUSH_MODE_NON_FTZ) { \ return HSA_STATUS_ERROR_INVALID_ARGUMENT; \ } \ } while (false) #define IS_VALID(ptr) \ do { \ if (((ptr) == NULL) || !((ptr)->IsValid())) \ return hsa_status_t(ValidityError::kValue); \ } while (false) #define CHECK_STATUS(status) \ do { \ if ((status) != HSA_STATUS_SUCCESS) return status; \ } while (false) #define CHECK_ALLOC(ptr) \ do { \ if ((ptr) == nullptr) return HSA_STATUS_ERROR_OUT_OF_RESOURCES; \ } while (false) #define IS_OPEN() \ do { \ if (!core::Runtime::runtime_singleton_->IsOpen()) \ return HSA_STATUS_ERROR_NOT_INITIALIZED; \ } while (false) template static __forceinline bool IsValid(T* ptr) { return (ptr == NULL) ? NULL : ptr->IsValid(); } namespace AMD { hsa_status_t handleException(); template static __forceinline T handleExceptionT() { handleException(); abort(); return T(); } } // namespace amd #define TRY try { #define CATCH } catch(...) { return AMD::handleException(); } #define CATCHRET(RETURN_TYPE) } catch(...) { return AMD::handleExceptionT(); } //----------------------------------------------------------------------------- // Basic Checks //----------------------------------------------------------------------------- static_assert(sizeof(hsa_barrier_and_packet_t) == sizeof(hsa_kernel_dispatch_packet_t), "AQL packet definitions have wrong sizes!"); static_assert(sizeof(hsa_barrier_and_packet_t) == sizeof(hsa_agent_dispatch_packet_t), "AQL packet definitions have wrong sizes!"); static_assert(sizeof(hsa_barrier_and_packet_t) == 64, "AQL packet definitions have wrong sizes!"); static_assert(sizeof(hsa_barrier_and_packet_t) == sizeof(hsa_barrier_or_packet_t), "AQL packet definitions have wrong sizes!"); #ifdef HSA_LARGE_MODEL static_assert(sizeof(void*) == 8, "HSA_LARGE_MODEL is set incorrectly!"); #else static_assert(sizeof(void*) == 4, "HSA_LARGE_MODEL is set incorrectly!"); #endif #if !defined(HSA_LARGE_MODEL) || !defined(__linux__) // static_assert(false, "Only HSA_LARGE_MODEL (64bit mode) and Linux supported."); #endif namespace HSA { //---------------------------------------------------------------------------// // Init/Shutdown routines //---------------------------------------------------------------------------// hsa_status_t hsa_init() { TRY; return core::Runtime::runtime_singleton_->Acquire(); CATCH; } hsa_status_t hsa_shut_down() { TRY; IS_OPEN(); return core::Runtime::runtime_singleton_->Release(); CATCH; } //---------------------------------------------------------------------------// // System //---------------------------------------------------------------------------// hsa_status_t hsa_system_get_info(hsa_system_info_t attribute, void* value) { TRY; IS_OPEN(); IS_BAD_PTR(value); return core::Runtime::runtime_singleton_->GetSystemInfo(attribute, value); CATCH; } hsa_status_t hsa_extension_get_name(uint16_t extension, const char** name) { TRY; IS_OPEN(); IS_BAD_PTR(name); switch (extension) { case HSA_EXTENSION_FINALIZER: *name = "HSA_EXTENSION_FINALIZER"; break; case HSA_EXTENSION_IMAGES: *name = "HSA_EXTENSION_IMAGES"; break; case HSA_EXTENSION_PERFORMANCE_COUNTERS: *name = "HSA_EXTENSION_PERFORMANCE_COUNTERS"; break; case HSA_EXTENSION_PROFILING_EVENTS: *name = "HSA_EXTENSION_PROFILING_EVENTS"; break; case HSA_EXTENSION_AMD_PROFILER: *name = "HSA_EXTENSION_AMD_PROFILER"; break; case HSA_EXTENSION_AMD_LOADER: *name = "HSA_EXTENSION_AMD_LOADER"; break; case HSA_EXTENSION_AMD_AQLPROFILE: *name = "HSA_EXTENSION_AMD_AQLPROFILE"; break; default: *name = "HSA_EXTENSION_INVALID"; return HSA_STATUS_ERROR_INVALID_ARGUMENT; } return HSA_STATUS_SUCCESS; CATCH; } hsa_status_t hsa_system_extension_supported(uint16_t extension, uint16_t version_major, uint16_t version_minor, bool* result) { TRY; IS_OPEN(); if ((extension > HSA_EXTENSION_STD_LAST && (extension < HSA_AMD_FIRST_EXTENSION || extension > HSA_AMD_LAST_EXTENSION)) || result == NULL) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; } *result = false; if (extension == HSA_EXTENSION_PERFORMANCE_COUNTERS || extension == HSA_EXTENSION_PROFILING_EVENTS) return HSA_STATUS_SUCCESS; uint16_t system_version_major = 0; hsa_status_t status = core::Runtime::runtime_singleton_->GetSystemInfo( HSA_SYSTEM_INFO_VERSION_MAJOR, &system_version_major); assert(status == HSA_STATUS_SUCCESS); if (version_major <= system_version_major) { uint16_t system_version_minor = 0; if (version_minor <= system_version_minor) { *result = true; } } return HSA_STATUS_SUCCESS; CATCH; } hsa_status_t hsa_system_major_extension_supported(uint16_t extension, uint16_t version_major, uint16_t* version_minor, bool* result) { TRY; IS_OPEN(); IS_BAD_PTR(version_minor); IS_BAD_PTR(result); if ((extension == HSA_EXTENSION_IMAGES) && (version_major == 1)) { *version_minor = 0; *result = true; return HSA_STATUS_SUCCESS; } if ((extension == HSA_EXTENSION_FINALIZER) && (version_major == 1)) { *version_minor = 0; *result = true; return HSA_STATUS_SUCCESS; } if ((extension == HSA_EXTENSION_AMD_LOADER) && (version_major == 1)) { *version_minor = 0; *result = true; return HSA_STATUS_SUCCESS; } if ((extension == HSA_EXTENSION_AMD_AQLPROFILE) && (version_major == 1)) { *version_minor = 0; *result = true; return HSA_STATUS_SUCCESS; } *result = false; return HSA_STATUS_SUCCESS; CATCH; } static size_t get_extension_table_length(uint16_t extension, uint16_t major, uint16_t minor) { // Table to convert from major/minor to major/length struct sizes_t { std::string name; size_t size; }; static sizes_t sizes[] = { {"hsa_ext_images_1_00_pfn_t", sizeof(hsa_ext_images_1_00_pfn_t)}, {"hsa_ext_finalizer_1_00_pfn_t", sizeof(hsa_ext_finalizer_1_00_pfn_t)}, {"hsa_ven_amd_loader_1_00_pfn_t", sizeof(hsa_ven_amd_loader_1_00_pfn_t)}, {"hsa_ven_amd_loader_1_01_pfn_t", sizeof(hsa_ven_amd_loader_1_01_pfn_t)}, {"hsa_ven_amd_loader_1_02_pfn_t", sizeof(hsa_ven_amd_loader_1_02_pfn_t)}, {"hsa_ven_amd_loader_1_03_pfn_t", sizeof(hsa_ven_amd_loader_1_03_pfn_t)}, {"hsa_ven_amd_aqlprofile_1_00_pfn_t", sizeof(hsa_ven_amd_aqlprofile_1_00_pfn_t)}}; static const size_t num_tables = sizeof(sizes) / sizeof(sizes_t); if (minor > 99) return 0; std::string name; switch (extension) { case HSA_EXTENSION_FINALIZER: name = "hsa_ext_finalizer_"; break; case HSA_EXTENSION_IMAGES: name = "hsa_ext_images_"; break; // case HSA_EXTENSION_PERFORMANCE_COUNTERS: // name = "hsa_ext_perf_counter_"; // break; // case HSA_EXTENSION_PROFILING_EVENTS: // name = "hsa_ext_profiling_event_"; // break; // case HSA_EXTENSION_AMD_PROFILER: // name = "hsa_ven_amd_profiler_"; // break; case HSA_EXTENSION_AMD_LOADER: name = "hsa_ven_amd_loader_"; break; case HSA_EXTENSION_AMD_AQLPROFILE: name = "hsa_ven_amd_aqlprofile_"; break; default: return 0; } char buff[6]; sprintf(buff, "%02u", minor); name += std::to_string(major) + "_" + buff + "_pfn_t"; for (size_t i = 0; i < num_tables; i++) { if (sizes[i].name == name) return sizes[i].size; } return 0; } hsa_status_t hsa_system_get_extension_table(uint16_t extension, uint16_t version_major, uint16_t version_minor, void* table) { TRY; return HSA::hsa_system_get_major_extension_table( extension, version_major, get_extension_table_length(extension, version_major, version_minor), table); CATCH; } hsa_status_t hsa_system_get_major_extension_table(uint16_t extension, uint16_t version_major, size_t table_length, void* table) { TRY; IS_OPEN(); IS_BAD_PTR(table); if (table_length == 0) return HSA_STATUS_ERROR_INVALID_ARGUMENT; if (extension == HSA_EXTENSION_IMAGES) { if (version_major != core::Runtime::runtime_singleton_->extensions_.image_api.version.major_id) { return HSA_STATUS_ERROR; } hsa_ext_images_1_pfn_t ext_table; ext_table.hsa_ext_image_clear = hsa_ext_image_clear; ext_table.hsa_ext_image_copy = hsa_ext_image_copy; ext_table.hsa_ext_image_create = hsa_ext_image_create; ext_table.hsa_ext_image_data_get_info = hsa_ext_image_data_get_info; ext_table.hsa_ext_image_destroy = hsa_ext_image_destroy; ext_table.hsa_ext_image_export = hsa_ext_image_export; ext_table.hsa_ext_image_get_capability = hsa_ext_image_get_capability; ext_table.hsa_ext_image_import = hsa_ext_image_import; ext_table.hsa_ext_sampler_create = hsa_ext_sampler_create; ext_table.hsa_ext_sampler_destroy = hsa_ext_sampler_destroy; ext_table.hsa_ext_image_get_capability_with_layout = hsa_ext_image_get_capability_with_layout; ext_table.hsa_ext_image_data_get_info_with_layout = hsa_ext_image_data_get_info_with_layout; ext_table.hsa_ext_image_create_with_layout = hsa_ext_image_create_with_layout; memcpy(table, &ext_table, Min(sizeof(ext_table), table_length)); return HSA_STATUS_SUCCESS; } if (extension == HSA_EXTENSION_FINALIZER) { if (version_major != core::Runtime::runtime_singleton_->extensions_.finalizer_api.version.major_id) { return HSA_STATUS_ERROR; } hsa_ext_finalizer_1_00_pfn_t ext_table; ext_table.hsa_ext_program_add_module = hsa_ext_program_add_module; ext_table.hsa_ext_program_create = hsa_ext_program_create; ext_table.hsa_ext_program_destroy = hsa_ext_program_destroy; ext_table.hsa_ext_program_finalize = hsa_ext_program_finalize; ext_table.hsa_ext_program_get_info = hsa_ext_program_get_info; ext_table.hsa_ext_program_iterate_modules = hsa_ext_program_iterate_modules; memcpy(table, &ext_table, Min(sizeof(ext_table), table_length)); return HSA_STATUS_SUCCESS; } if (extension == HSA_EXTENSION_AMD_LOADER) { if (version_major != 1) return HSA_STATUS_ERROR; hsa_ven_amd_loader_1_03_pfn_t ext_table; ext_table.hsa_ven_amd_loader_query_host_address = hsa_ven_amd_loader_query_host_address; ext_table.hsa_ven_amd_loader_query_segment_descriptors = hsa_ven_amd_loader_query_segment_descriptors; ext_table.hsa_ven_amd_loader_query_executable = hsa_ven_amd_loader_query_executable; ext_table.hsa_ven_amd_loader_executable_iterate_loaded_code_objects = hsa_ven_amd_loader_executable_iterate_loaded_code_objects; ext_table.hsa_ven_amd_loader_loaded_code_object_get_info = hsa_ven_amd_loader_loaded_code_object_get_info; ext_table.hsa_ven_amd_loader_code_object_reader_create_from_file_with_offset_size = hsa_ven_amd_loader_code_object_reader_create_from_file_with_offset_size; ext_table.hsa_ven_amd_loader_iterate_executables = hsa_ven_amd_loader_iterate_executables; memcpy(table, &ext_table, Min(sizeof(ext_table), table_length)); return HSA_STATUS_SUCCESS; } if (extension == HSA_EXTENSION_AMD_AQLPROFILE) { if (version_major != hsa_ven_amd_aqlprofile_VERSION_MAJOR) { debug_print("aqlprofile API incompatible ver %d, current ver %d\n", version_major, hsa_ven_amd_aqlprofile_VERSION_MAJOR); return HSA_STATUS_ERROR; } os::LibHandle lib = os::LoadLib(kAqlProfileLib); if (lib == NULL) { debug_print("Loading '%s' failed\n", kAqlProfileLib); return HSA_STATUS_ERROR; } hsa_ven_amd_aqlprofile_pfn_t ext_table; ext_table.hsa_ven_amd_aqlprofile_version_major = (decltype(::hsa_ven_amd_aqlprofile_version_major)*) os::GetExportAddress(lib, "hsa_ven_amd_aqlprofile_version_major"); ext_table.hsa_ven_amd_aqlprofile_version_minor = (decltype(::hsa_ven_amd_aqlprofile_version_minor)*) os::GetExportAddress(lib, "hsa_ven_amd_aqlprofile_version_minor"); ext_table.hsa_ven_amd_aqlprofile_error_string = (decltype(::hsa_ven_amd_aqlprofile_error_string)*) os::GetExportAddress(lib, "hsa_ven_amd_aqlprofile_error_string"); ext_table.hsa_ven_amd_aqlprofile_validate_event = (decltype(::hsa_ven_amd_aqlprofile_validate_event)*) os::GetExportAddress(lib, "hsa_ven_amd_aqlprofile_validate_event"); ext_table.hsa_ven_amd_aqlprofile_start = (decltype(::hsa_ven_amd_aqlprofile_start)*) os::GetExportAddress(lib, "hsa_ven_amd_aqlprofile_start"); ext_table.hsa_ven_amd_aqlprofile_stop = (decltype(::hsa_ven_amd_aqlprofile_stop)*) os::GetExportAddress(lib, "hsa_ven_amd_aqlprofile_stop"); ext_table.hsa_ven_amd_aqlprofile_read = (decltype(::hsa_ven_amd_aqlprofile_read)*) os::GetExportAddress(lib, "hsa_ven_amd_aqlprofile_read"); ext_table.hsa_ven_amd_aqlprofile_legacy_get_pm4 = (decltype(::hsa_ven_amd_aqlprofile_legacy_get_pm4)*) os::GetExportAddress(lib, "hsa_ven_amd_aqlprofile_legacy_get_pm4"); ext_table.hsa_ven_amd_aqlprofile_get_info = (decltype(::hsa_ven_amd_aqlprofile_get_info)*) os::GetExportAddress(lib, "hsa_ven_amd_aqlprofile_get_info"); ext_table.hsa_ven_amd_aqlprofile_iterate_data = (decltype(::hsa_ven_amd_aqlprofile_iterate_data)*) os::GetExportAddress(lib, "hsa_ven_amd_aqlprofile_iterate_data"); bool version_incompatible = true; uint32_t version_curr = 0; version_major = HSA_AQLPROFILE_VERSION_MAJOR; if (ext_table.hsa_ven_amd_aqlprofile_version_major != NULL) { version_curr = ext_table.hsa_ven_amd_aqlprofile_version_major(); version_incompatible = (version_major != version_curr); } if (version_incompatible == true) { debug_print("Loading '%s' failed, incompatible ver %d, current ver %d\n", kAqlProfileLib, version_major, version_curr); return HSA_STATUS_ERROR; } memcpy(table, &ext_table, Min(sizeof(ext_table), table_length)); return HSA_STATUS_SUCCESS; } return HSA_STATUS_ERROR; CATCH; } //---------------------------------------------------------------------------// // Agent //---------------------------------------------------------------------------// hsa_status_t hsa_iterate_agents(hsa_status_t (*callback)(hsa_agent_t agent, void* data), void* data) { TRY; IS_OPEN(); IS_BAD_PTR(callback); return core::Runtime::runtime_singleton_->IterateAgent(callback, data); CATCH; } hsa_status_t hsa_agent_get_info(hsa_agent_t agent_handle, hsa_agent_info_t attribute, void* value) { TRY; IS_OPEN(); IS_BAD_PTR(value); const core::Agent* agent = core::Agent::Convert(agent_handle); IS_VALID(agent); return agent->GetInfo(attribute, value); CATCH; } hsa_status_t hsa_agent_get_exception_policies(hsa_agent_t agent_handle, hsa_profile_t profile, uint16_t* mask) { TRY; IS_OPEN(); IS_BAD_PTR(mask); IS_BAD_PROFILE(profile); const core::Agent* agent = core::Agent::Convert(agent_handle); IS_VALID(agent); *mask = 0; return HSA_STATUS_SUCCESS; CATCH; } hsa_status_t hsa_cache_get_info(hsa_cache_t cache, hsa_cache_info_t attribute, void* value) { TRY; IS_OPEN(); core::Cache* Cache = core::Cache::Convert(cache); IS_VALID(Cache); IS_BAD_PTR(value); return Cache->GetInfo(attribute, value); CATCH; } hsa_status_t hsa_agent_iterate_caches(hsa_agent_t agent_handle, hsa_status_t (*callback)(hsa_cache_t cache, void* data), void* data) { TRY; IS_OPEN(); const core::Agent* agent = core::Agent::Convert(agent_handle); IS_VALID(agent); IS_BAD_PTR(callback); return agent->IterateCache(callback, data); CATCH; } hsa_status_t hsa_agent_extension_supported(uint16_t extension, hsa_agent_t agent_handle, uint16_t version_major, uint16_t version_minor, bool* result) { TRY; IS_OPEN(); if ((extension > HSA_EXTENSION_STD_LAST && (extension < HSA_AMD_FIRST_EXTENSION || extension > HSA_AMD_LAST_EXTENSION)) || result == NULL) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; } *result = false; const core::Agent* agent = core::Agent::Convert(agent_handle); IS_VALID(agent); if (agent->device_type() == core::Agent::kAmdGpuDevice) { uint16_t agent_version_major = 0; hsa_status_t status = agent->GetInfo(HSA_AGENT_INFO_VERSION_MAJOR, &agent_version_major); assert(status == HSA_STATUS_SUCCESS); if (version_major <= agent_version_major) { uint16_t agent_version_minor = 0; if (version_minor <= agent_version_minor) { *result = true; } } } return HSA_STATUS_SUCCESS; CATCH; } hsa_status_t hsa_agent_major_extension_supported(uint16_t extension, hsa_agent_t agent_handle, uint16_t version_major, uint16_t* version_minor, bool* result) { TRY; IS_OPEN(); if ((extension > HSA_EXTENSION_STD_LAST && (extension < HSA_AMD_FIRST_EXTENSION || extension > HSA_AMD_LAST_EXTENSION)) || result == NULL) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; } *result = false; const core::Agent* agent = core::Agent::Convert(agent_handle); IS_VALID(agent); if (agent->device_type() == core::Agent::kAmdGpuDevice) { uint16_t agent_version_major = 0; hsa_status_t status = agent->GetInfo(HSA_AGENT_INFO_VERSION_MAJOR, &agent_version_major); assert(status == HSA_STATUS_SUCCESS); if (version_major <= agent_version_major) { *version_minor = 0; *result = true; } } return HSA_STATUS_SUCCESS; CATCH; } /// @brief Api to create a user mode queue. /// /// @param agent Hsa Agent which will execute Aql commands /// /// @param size Size of Queue in terms of Aql packet size /// /// @param type of Queue Single Writer or Multiple Writer /// /// @param callback Callback function to register in case Quee /// encounters an error /// /// @param service_queue Pointer to a service queue /// /// @param queue Output parameter updated with a pointer to the /// queue being created /// /// @return hsa_status hsa_status_t hsa_queue_create( hsa_agent_t agent_handle, uint32_t size, hsa_queue_type32_t type, void (*callback)(hsa_status_t status, hsa_queue_t* source, void* data), void* data, uint32_t private_segment_size, uint32_t group_segment_size, hsa_queue_t** queue) { TRY; IS_OPEN(); if ((queue == nullptr) || (size == 0) || (!IsPowerOfTwo(size)) || (type > HSA_QUEUE_TYPE_COOPERATIVE)) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; } core::Agent* agent = core::Agent::Convert(agent_handle); IS_VALID(agent); hsa_queue_type32_t agent_queue_type = HSA_QUEUE_TYPE_MULTI; hsa_status_t status = agent->GetInfo(HSA_AGENT_INFO_QUEUE_TYPE, &agent_queue_type); assert(HSA_STATUS_SUCCESS == status); if ((agent_queue_type == HSA_QUEUE_TYPE_SINGLE) && (type != HSA_QUEUE_TYPE_SINGLE)) { return HSA_STATUS_ERROR_INVALID_QUEUE_CREATION; } if (callback == nullptr) callback = core::Queue::DefaultErrorHandler; core::Queue* cmd_queue = nullptr; status = agent->QueueCreate(size, type, callback, data, private_segment_size, group_segment_size, &cmd_queue); if (status != HSA_STATUS_SUCCESS) return status; assert(cmd_queue != nullptr && "Queue not returned but status was success.\n"); *queue = core::Queue::Convert(cmd_queue); return status; CATCH; } hsa_status_t hsa_soft_queue_create(hsa_region_t region, uint32_t size, hsa_queue_type32_t type, uint32_t features, hsa_signal_t doorbell_signal, hsa_queue_t** queue) { TRY; IS_OPEN(); if ((queue == NULL) || (region.handle == 0) || (doorbell_signal.handle == 0) || (size == 0) || (!IsPowerOfTwo(size)) || (type < HSA_QUEUE_TYPE_MULTI) || (type > HSA_QUEUE_TYPE_SINGLE) || (features == 0)) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; } const core::MemoryRegion* mem_region = core::MemoryRegion::Convert(region); IS_VALID(mem_region); const core::Signal* signal = core::Signal::Convert(doorbell_signal); IS_VALID(signal); core::HostQueue* host_queue = new core::HostQueue(region, size, type, features, doorbell_signal); *queue = core::Queue::Convert(host_queue); return HSA_STATUS_SUCCESS; CATCH; } /// @brief Api to destroy a user mode queue /// /// @param queue Pointer to the queue being destroyed /// /// @return hsa_status hsa_status_t hsa_queue_destroy(hsa_queue_t* queue) { TRY; IS_OPEN(); IS_BAD_PTR(queue); core::Queue* cmd_queue = core::Queue::Convert(queue); IS_VALID(cmd_queue); cmd_queue->Destroy(); return HSA_STATUS_SUCCESS; CATCH; } /// @brief Api to inactivate a user mode queue /// /// @param queue Pointer to the queue being inactivated /// /// @return hsa_status hsa_status_t hsa_queue_inactivate(hsa_queue_t* queue) { TRY; IS_OPEN(); IS_BAD_PTR(queue); core::Queue* cmd_queue = core::Queue::Convert(queue); IS_VALID(cmd_queue); cmd_queue->Inactivate(); return HSA_STATUS_SUCCESS; CATCH; } /// @brief Api to read the Read Index of Queue using Acquire semantics /// /// @param queue Pointer to the queue whose read index is being read /// /// @return uint64_t Value of Read index uint64_t hsa_queue_load_read_index_scacquire(const hsa_queue_t* queue) { TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->LoadReadIndexAcquire(); CATCHRET(uint64_t); } /// @brief Api to read the Read Index of Queue using Relaxed semantics /// /// @param queue Pointer to the queue whose read index is being read /// /// @return uint64_t Value of Read index uint64_t hsa_queue_load_read_index_relaxed(const hsa_queue_t* queue) { TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->LoadReadIndexRelaxed(); CATCHRET(uint64_t); } /// @brief Api to read the Write Index of Queue using Acquire semantics /// /// @param queue Pointer to the queue whose write index is being read /// /// @return uint64_t Value of Write index uint64_t hsa_queue_load_write_index_scacquire(const hsa_queue_t* queue) { TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->LoadWriteIndexAcquire(); CATCHRET(uint64_t); } /// @brief Api to read the Write Index of Queue using Relaxed semantics /// /// @param queue Pointer to the queue whose write index is being read /// /// @return uint64_t Value of Write index uint64_t hsa_queue_load_write_index_relaxed(const hsa_queue_t* queue) { TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->LoadWriteIndexRelaxed(); CATCHRET(uint64_t); } /// @brief Api to store the Read Index of Queue using Relaxed semantics /// /// @param queue Pointer to the queue whose read index is being updated /// /// @param value Value of new read index void hsa_queue_store_read_index_relaxed(const hsa_queue_t* queue, uint64_t value) { TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); cmd_queue->StoreReadIndexRelaxed(value); CATCHRET(void); } /// @brief Api to store the Read Index of Queue using Release semantics /// /// @param queue Pointer to the queue whose read index is being updated /// /// @param value Value of new read index void hsa_queue_store_read_index_screlease(const hsa_queue_t* queue, uint64_t value) { TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); cmd_queue->StoreReadIndexRelease(value); CATCHRET(void); } /// @brief Api to store the Write Index of Queue using Relaxed semantics /// /// @param queue Pointer to the queue whose write index is being updated /// /// @param value Value of new write index void hsa_queue_store_write_index_relaxed(const hsa_queue_t* queue, uint64_t value) { TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); cmd_queue->StoreWriteIndexRelaxed(value); CATCHRET(void); } /// @brief Api to store the Write Index of Queue using Release semantics /// /// @param queue Pointer to the queue whose write index is being updated /// /// @param value Value of new write index void hsa_queue_store_write_index_screlease(const hsa_queue_t* queue, uint64_t value) { TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); cmd_queue->StoreWriteIndexRelease(value); CATCHRET(void); } /// @brief Api to compare and swap the Write Index of Queue using Acquire and /// Release semantics /// /// @param queue Pointer to the queue whose write index is being updated /// /// @param expected Current value of write index /// /// @param value Value of new write index /// /// @return uint64_t Value of write index before the update uint64_t hsa_queue_cas_write_index_scacq_screl(const hsa_queue_t* queue, uint64_t expected, uint64_t value) { TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->CasWriteIndexAcqRel(expected, value); CATCHRET(uint64_t); } /// @brief Api to compare and swap the Write Index of Queue using Acquire /// Semantics /// /// @param queue Pointer to the queue whose write index is being updated /// /// @param expected Current value of write index /// /// @param value Value of new write index /// /// @return uint64_t Value of write index before the update uint64_t hsa_queue_cas_write_index_scacquire(const hsa_queue_t* queue, uint64_t expected, uint64_t value) { TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->CasWriteIndexAcquire(expected, value); CATCHRET(uint64_t); } /// @brief Api to compare and swap the Write Index of Queue using Relaxed /// Semantics /// /// @param queue Pointer to the queue whose write index is being updated /// /// @param expected Current value of write index /// /// @param value Value of new write index /// /// @return uint64_t Value of write index before the update uint64_t hsa_queue_cas_write_index_relaxed(const hsa_queue_t* queue, uint64_t expected, uint64_t value) { TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->CasWriteIndexRelaxed(expected, value); CATCHRET(uint64_t); } /// @brief Api to compare and swap the Write Index of Queue using Release /// Semantics /// /// @param queue Pointer to the queue whose write index is being updated /// /// @param expected Current value of write index /// /// @param value Value of new write index /// /// @return uint64_t Value of write index before the update uint64_t hsa_queue_cas_write_index_screlease(const hsa_queue_t* queue, uint64_t expected, uint64_t value) { TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->CasWriteIndexRelease(expected, value); CATCHRET(uint64_t); } /// @brief Api to Add to the Write Index of Queue using Acquire and Release /// Semantics /// /// @param queue Pointer to the queue whose write index is being updated /// /// @param value Value to add to write index /// /// @return uint64_t Value of write index before the update uint64_t hsa_queue_add_write_index_scacq_screl(const hsa_queue_t* queue, uint64_t value) { TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->AddWriteIndexAcqRel(value); CATCHRET(uint64_t); } /// @brief Api to Add to the Write Index of Queue using Acquire Semantics /// /// @param queue Pointer to the queue whose write index is being updated /// /// @param value Value to add to write index /// /// @return uint64_t Value of write index before the update uint64_t hsa_queue_add_write_index_scacquire(const hsa_queue_t* queue, uint64_t value) { TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->AddWriteIndexAcquire(value); CATCHRET(uint64_t); } /// @brief Api to Add to the Write Index of Queue using Relaxed Semantics /// /// @param queue Pointer to the queue whose write index is being updated /// /// @param value Value to add to write index /// /// @return uint64_t Value of write index before the update uint64_t hsa_queue_add_write_index_relaxed(const hsa_queue_t* queue, uint64_t value) { TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->AddWriteIndexRelaxed(value); CATCHRET(uint64_t); } /// @brief Api to Add to the Write Index of Queue using Release Semantics /// /// @param queue Pointer to the queue whose write index is being updated /// /// @param value Value to add to write index /// /// @return uint64_t Value of write index before the update uint64_t hsa_queue_add_write_index_screlease(const hsa_queue_t* queue, uint64_t value) { TRY; core::Queue* cmd_queue = core::Queue::Convert(queue); assert(IsValid(cmd_queue)); return cmd_queue->AddWriteIndexRelease(value); CATCHRET(uint64_t); } //----------------------------------------------------------------------------- // Memory //----------------------------------------------------------------------------- hsa_status_t hsa_agent_iterate_regions( hsa_agent_t agent_handle, hsa_status_t (*callback)(hsa_region_t region, void* data), void* data) { TRY; IS_OPEN(); IS_BAD_PTR(callback); const core::Agent* agent = core::Agent::Convert(agent_handle); IS_VALID(agent); return agent->IterateRegion(callback, data); CATCH; } hsa_status_t hsa_region_get_info(hsa_region_t region, hsa_region_info_t attribute, void* value) { TRY; IS_OPEN(); IS_BAD_PTR(value); const core::MemoryRegion* mem_region = core::MemoryRegion::Convert(region); IS_VALID(mem_region); return mem_region->GetInfo(attribute, value); CATCH; } hsa_status_t hsa_memory_register(void* address, size_t size) { TRY; IS_OPEN(); if (size == 0 && address != NULL) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; } return HSA_STATUS_SUCCESS; CATCH; } hsa_status_t hsa_memory_deregister(void* address, size_t size) { TRY; IS_OPEN(); return HSA_STATUS_SUCCESS; CATCH; } hsa_status_t hsa_memory_allocate(hsa_region_t region, size_t size, void** ptr) { TRY; IS_OPEN(); if (size == 0 || ptr == NULL) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; } const core::MemoryRegion* mem_region = core::MemoryRegion::Convert(region); IS_VALID(mem_region); return core::Runtime::runtime_singleton_->AllocateMemory( mem_region, size, core::MemoryRegion::AllocateNoFlags, ptr); CATCH; } hsa_status_t hsa_memory_free(void* ptr) { TRY; IS_OPEN(); if (ptr == NULL) { return HSA_STATUS_SUCCESS; } return core::Runtime::runtime_singleton_->FreeMemory(ptr); CATCH; } hsa_status_t hsa_memory_assign_agent(void* ptr, hsa_agent_t agent_handle, hsa_access_permission_t access) { TRY; IS_OPEN(); if ((ptr == NULL) || (access < HSA_ACCESS_PERMISSION_RO) || (access > HSA_ACCESS_PERMISSION_RW)) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; } const core::Agent* agent = core::Agent::Convert(agent_handle); IS_VALID(agent); return HSA_STATUS_SUCCESS; CATCH; } hsa_status_t hsa_memory_copy(void* dst, const void* src, size_t size) { TRY; IS_OPEN(); if (dst == NULL || src == NULL) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; } if (size == 0) { return HSA_STATUS_SUCCESS; } return core::Runtime::runtime_singleton_->CopyMemory(dst, src, size); CATCH; } //----------------------------------------------------------------------------- // Signals //----------------------------------------------------------------------------- hsa_status_t hsa_signal_create(hsa_signal_value_t initial_value, uint32_t num_consumers, const hsa_agent_t* consumers, hsa_signal_t* hsa_signal) { return AMD::hsa_amd_signal_create(initial_value, num_consumers, consumers, 0, hsa_signal); } hsa_status_t hsa_signal_destroy(hsa_signal_t hsa_signal) { TRY; IS_OPEN(); core::Signal* signal = core::Signal::Convert(hsa_signal); signal->DestroySignal(); return HSA_STATUS_SUCCESS; CATCH; } hsa_signal_value_t hsa_signal_load_relaxed(hsa_signal_t hsa_signal) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->LoadRelaxed(); CATCHRET(hsa_signal_value_t); } hsa_signal_value_t hsa_signal_load_scacquire(hsa_signal_t hsa_signal) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->LoadAcquire(); CATCHRET(hsa_signal_value_t); } void hsa_signal_store_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->StoreRelaxed(value); CATCHRET(void); } void hsa_signal_store_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->StoreRelease(value); CATCHRET(void); } hsa_signal_value_t hsa_signal_wait_relaxed(hsa_signal_t hsa_signal, hsa_signal_condition_t condition, hsa_signal_value_t compare_value, uint64_t timeout_hint, hsa_wait_state_t wait_state_hint) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->WaitRelaxed(condition, compare_value, timeout_hint, wait_state_hint); CATCHRET(hsa_signal_value_t); } hsa_signal_value_t hsa_signal_wait_scacquire(hsa_signal_t hsa_signal, hsa_signal_condition_t condition, hsa_signal_value_t compare_value, uint64_t timeout_hint, hsa_wait_state_t wait_state_hint) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->WaitAcquire(condition, compare_value, timeout_hint, wait_state_hint); CATCHRET(hsa_signal_value_t); } hsa_status_t hsa_signal_group_create(uint32_t num_signals, const hsa_signal_t* signals, uint32_t num_consumers, const hsa_agent_t* consumers, hsa_signal_group_t* signal_group) { TRY; IS_OPEN(); if (num_signals == 0) return HSA_STATUS_ERROR_INVALID_ARGUMENT; for (uint i = 0; i < num_signals; i++) IS_VALID(core::Signal::Convert(signals[i])); for (uint i = 0; i < num_consumers; i++) IS_VALID(core::Agent::Convert(consumers[i])); core::SignalGroup* group = new core::SignalGroup(num_signals, signals); CHECK_ALLOC(group); if (!group->IsValid()) { delete group; return HSA_STATUS_ERROR_OUT_OF_RESOURCES; } *signal_group = core::SignalGroup::Convert(group); return HSA_STATUS_SUCCESS; CATCH; } hsa_status_t hsa_signal_group_destroy(hsa_signal_group_t signal_group) { TRY; IS_OPEN(); core::SignalGroup* group = core::SignalGroup::Convert(signal_group); IS_VALID(group); delete group; return HSA_STATUS_SUCCESS; CATCH; } hsa_status_t hsa_signal_group_wait_any_relaxed(hsa_signal_group_t signal_group, const hsa_signal_condition_t* conditions, const hsa_signal_value_t* compare_values, hsa_wait_state_t wait_state_hint, hsa_signal_t* signal, hsa_signal_value_t* value) { TRY; IS_OPEN(); const core::SignalGroup* group = core::SignalGroup::Convert(signal_group); IS_VALID(group); const uint32_t index = AMD::hsa_amd_signal_wait_any( group->Count(), const_cast(group->List()), const_cast(conditions), const_cast(compare_values), uint64_t(-1), wait_state_hint, value); if (index >= group->Count()) return HSA_STATUS_ERROR_INVALID_ARGUMENT; *signal = group->List()[index]; return HSA_STATUS_SUCCESS; CATCH; } hsa_status_t hsa_signal_group_wait_any_scacquire(hsa_signal_group_t signal_group, const hsa_signal_condition_t* conditions, const hsa_signal_value_t* compare_values, hsa_wait_state_t wait_state_hint, hsa_signal_t* signal, hsa_signal_value_t* value) { TRY; hsa_status_t ret = HSA::hsa_signal_group_wait_any_relaxed( signal_group, conditions, compare_values, wait_state_hint, signal, value); std::atomic_thread_fence(std::memory_order_acquire); return ret; CATCH; } void hsa_signal_and_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AndRelaxed(value); CATCHRET(void); } void hsa_signal_and_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AndAcquire(value); CATCHRET(void); } void hsa_signal_and_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AndRelease(value); CATCHRET(void); } void hsa_signal_and_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AndAcqRel(value); CATCHRET(void); } void hsa_signal_or_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->OrRelaxed(value); CATCHRET(void); } void hsa_signal_or_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->OrAcquire(value); CATCHRET(void); } void hsa_signal_or_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->OrRelease(value); CATCHRET(void); } void hsa_signal_or_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->OrAcqRel(value); CATCHRET(void); } void hsa_signal_xor_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->XorRelaxed(value); CATCHRET(void); } void hsa_signal_xor_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->XorAcquire(value); CATCHRET(void); } void hsa_signal_xor_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->XorRelease(value); CATCHRET(void); } void hsa_signal_xor_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->XorAcqRel(value); CATCHRET(void); } void hsa_signal_add_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AddRelaxed(value); CATCHRET(void); } void hsa_signal_add_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AddAcquire(value); CATCHRET(void); } void hsa_signal_add_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AddRelease(value); CATCHRET(void); } void hsa_signal_add_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->AddAcqRel(value); CATCHRET(void); } void hsa_signal_subtract_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->SubRelaxed(value); CATCHRET(void); } void hsa_signal_subtract_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->SubAcquire(value); CATCHRET(void); } void hsa_signal_subtract_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->SubRelease(value); CATCHRET(void); } void hsa_signal_subtract_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); signal->SubAcqRel(value); CATCHRET(void); } hsa_signal_value_t hsa_signal_exchange_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->ExchRelaxed(value); CATCHRET(hsa_signal_value_t); } hsa_signal_value_t hsa_signal_exchange_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->ExchAcquire(value); CATCHRET(hsa_signal_value_t); } hsa_signal_value_t hsa_signal_exchange_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->ExchRelease(value); CATCHRET(hsa_signal_value_t); } hsa_signal_value_t hsa_signal_exchange_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->ExchAcqRel(value); CATCHRET(hsa_signal_value_t); } hsa_signal_value_t hsa_signal_cas_relaxed(hsa_signal_t hsa_signal, hsa_signal_value_t expected, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->CasRelaxed(expected, value); CATCHRET(hsa_signal_value_t); } hsa_signal_value_t hsa_signal_cas_scacquire(hsa_signal_t hsa_signal, hsa_signal_value_t expected, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->CasAcquire(expected, value); CATCHRET(hsa_signal_value_t); } hsa_signal_value_t hsa_signal_cas_screlease(hsa_signal_t hsa_signal, hsa_signal_value_t expected, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->CasRelease(expected, value); CATCHRET(hsa_signal_value_t); } hsa_signal_value_t hsa_signal_cas_scacq_screl(hsa_signal_t hsa_signal, hsa_signal_value_t expected, hsa_signal_value_t value) { TRY; core::Signal* signal = core::Signal::Convert(hsa_signal); assert(IsValid(signal)); return signal->CasAcqRel(expected, value); CATCHRET(hsa_signal_value_t); } //===--- Instruction Set Architecture -------------------------------------===// using core::Isa; using core::IsaRegistry; using core::Wavefront; hsa_status_t hsa_isa_from_name( const char *name, hsa_isa_t *isa) { TRY; IS_OPEN(); IS_BAD_PTR(name); IS_BAD_PTR(isa); const Isa *isa_object = IsaRegistry::GetIsa(name); if (!isa_object) { return HSA_STATUS_ERROR_INVALID_ISA_NAME; } *isa = Isa::Handle(isa_object); return HSA_STATUS_SUCCESS; CATCH; } hsa_status_t hsa_agent_iterate_isas( hsa_agent_t agent, hsa_status_t (*callback)(hsa_isa_t isa, void *data), void *data) { TRY; IS_OPEN(); IS_BAD_PTR(callback); const core::Agent *agent_object = core::Agent::Convert(agent); IS_VALID(agent_object); const Isa *isa_object = agent_object->isa(); if (!isa_object) { return HSA_STATUS_SUCCESS; } return callback(Isa::Handle(isa_object), data); CATCH; } /* deprecated */ hsa_status_t hsa_isa_get_info( hsa_isa_t isa, hsa_isa_info_t attribute, uint32_t index, void *value) { TRY; IS_OPEN(); IS_BAD_PTR(value); if (index != 0) { return HSA_STATUS_ERROR_INVALID_INDEX; } const Isa *isa_object = Isa::Object(isa); IS_VALID(isa_object); return isa_object->GetInfo(attribute, value) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR_INVALID_ARGUMENT; CATCH; } hsa_status_t hsa_isa_get_info_alt( hsa_isa_t isa, hsa_isa_info_t attribute, void *value) { TRY; IS_OPEN(); IS_BAD_PTR(value); const Isa *isa_object = Isa::Object(isa); IS_VALID(isa_object); return isa_object->GetInfo(attribute, value) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR_INVALID_ARGUMENT; CATCH; } hsa_status_t hsa_isa_get_exception_policies( hsa_isa_t isa, hsa_profile_t profile, uint16_t *mask) { TRY; IS_OPEN(); IS_BAD_PROFILE(profile); IS_BAD_PTR(mask); const Isa *isa_object = Isa::Object(isa); IS_VALID(isa_object); // FIXME: update when exception policies are supported. *mask = 0; return HSA_STATUS_SUCCESS; CATCH; } hsa_status_t hsa_isa_get_round_method( hsa_isa_t isa, hsa_fp_type_t fp_type, hsa_flush_mode_t flush_mode, hsa_round_method_t *round_method) { TRY; IS_OPEN(); IS_BAD_FP_TYPE(fp_type); IS_BAD_FLUSH_MODE(flush_mode); IS_BAD_PTR(round_method); const Isa *isa_object = Isa::Object(isa); IS_VALID(isa_object); *round_method = isa_object->GetRoundMethod(fp_type, flush_mode); return HSA_STATUS_SUCCESS; CATCH; } hsa_status_t hsa_wavefront_get_info( hsa_wavefront_t wavefront, hsa_wavefront_info_t attribute, void *value) { TRY; IS_OPEN(); IS_BAD_PTR(value); const Wavefront *wavefront_object = Wavefront::Object(wavefront); if (!wavefront_object) { return HSA_STATUS_ERROR_INVALID_WAVEFRONT; } return wavefront_object->GetInfo(attribute, value) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR_INVALID_ARGUMENT; CATCH; } hsa_status_t hsa_isa_iterate_wavefronts( hsa_isa_t isa, hsa_status_t (*callback)(hsa_wavefront_t wavefront, void *data), void *data) { TRY; IS_OPEN(); IS_BAD_PTR(callback); const Isa *isa_object = Isa::Object(isa); IS_VALID(isa_object); const Wavefront &wavefront_object = isa_object->GetWavefront(); return callback(Wavefront::Handle(&wavefront_object), data); CATCH; } /* deprecated */ hsa_status_t hsa_isa_compatible( hsa_isa_t code_object_isa, hsa_isa_t agent_isa, bool *result) { TRY; IS_OPEN(); IS_BAD_PTR(result); const Isa *code_object_isa_object = Isa::Object(code_object_isa); IS_VALID(code_object_isa_object); const Isa *agent_isa_object = Isa::Object(agent_isa); IS_VALID(agent_isa_object); *result = Isa::IsCompatible(*code_object_isa_object, *agent_isa_object); return HSA_STATUS_SUCCESS; CATCH; } //===--- Code Objects (deprecated) ----------------------------------------===// namespace { hsa_status_t IsCodeObjectAllocRegion( hsa_region_t region, void *data) { assert(data); assert(((hsa_region_t*)data)->handle == 0); bool runtime_alloc_allowed = false; hsa_status_t status = HSA::hsa_region_get_info( region, HSA_REGION_INFO_RUNTIME_ALLOC_ALLOWED, &runtime_alloc_allowed); if (status != HSA_STATUS_SUCCESS) { return status; } if (runtime_alloc_allowed) { ((hsa_region_t*)data)->handle = region.handle; return HSA_STATUS_INFO_BREAK; } return HSA_STATUS_SUCCESS; } hsa_status_t FindCodeObjectAllocRegionForAgent( hsa_agent_t agent, void *data) { assert(data); assert(((hsa_region_t*)data)->handle == 0); hsa_device_type_t device = HSA_DEVICE_TYPE_CPU; hsa_status_t status = HSA::hsa_agent_get_info( agent, HSA_AGENT_INFO_DEVICE, &device); if (status != HSA_STATUS_SUCCESS) { return status; } if (device == HSA_DEVICE_TYPE_CPU) { return HSA::hsa_agent_iterate_regions(agent, IsCodeObjectAllocRegion, data); } return HSA_STATUS_SUCCESS; } hsa_status_t FindCodeObjectAllocRegion( void *data) { assert(data); assert(((hsa_region_t*)data)->handle == 0); return HSA::hsa_iterate_agents(FindCodeObjectAllocRegionForAgent, data); } amd::hsa::code::AmdHsaCodeManager *GetCodeManager() { return core::Runtime::runtime_singleton_->code_manager(); } } // namespace anonymous /* deprecated */ hsa_status_t hsa_code_object_serialize( hsa_code_object_t code_object, hsa_status_t (*alloc_callback)(size_t size, hsa_callback_data_t data, void **address), hsa_callback_data_t callback_data, const char *options, void **serialized_code_object, size_t *serialized_code_object_size) { TRY; IS_OPEN(); IS_BAD_PTR(alloc_callback); IS_BAD_PTR(serialized_code_object); IS_BAD_PTR(serialized_code_object_size); amd::hsa::code::AmdHsaCode *code = GetCodeManager()->FromHandle(code_object); if (!code) { return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; } hsa_status_t status = alloc_callback( code->ElfSize(), callback_data, serialized_code_object); if (status != HSA_STATUS_SUCCESS) { return status; } assert(*serialized_code_object); memcpy(*serialized_code_object, code->ElfData(), code->ElfSize()); *serialized_code_object_size = code->ElfSize(); return HSA_STATUS_SUCCESS; CATCH; } /* deprecated */ hsa_status_t hsa_code_object_deserialize( void *serialized_code_object, size_t serialized_code_object_size, const char *options, hsa_code_object_t *code_object) { TRY; IS_OPEN(); IS_BAD_PTR(serialized_code_object); IS_BAD_PTR(code_object); if (serialized_code_object_size == 0) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; } hsa_region_t code_object_alloc_region = {0}; hsa_status_t status = FindCodeObjectAllocRegion(&code_object_alloc_region); if (status != HSA_STATUS_SUCCESS && status != HSA_STATUS_INFO_BREAK) { return status; } assert(code_object_alloc_region.handle != 0); void *code_object_alloc_data = nullptr; status = HSA::hsa_memory_allocate( code_object_alloc_region, serialized_code_object_size, &code_object_alloc_data); if (status != HSA_STATUS_SUCCESS) { return status; } assert(code_object_alloc_data); memcpy( code_object_alloc_data, serialized_code_object, serialized_code_object_size); code_object->handle = reinterpret_cast(code_object_alloc_data); return HSA_STATUS_SUCCESS; CATCH; } /* deprecated */ hsa_status_t hsa_code_object_destroy( hsa_code_object_t code_object) { TRY; IS_OPEN(); void *code_object_data = reinterpret_cast(code_object.handle); if (!code_object_data) { return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; } if (!GetCodeManager()->Destroy(code_object)) { return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; } HSA::hsa_memory_free(code_object_data); return HSA_STATUS_SUCCESS; CATCH; } static std::string ConvertOldTargetNameToNew( const std::string &OldName, bool IsFinalizer, uint32_t EFlags) { std::string NewName = ""; bool xnack_supported = false; // FIXME #1: Should 9:0:3 be completely (loader, sc, etc.) removed? // FIXME #2: What does PAL do with respect to boltzmann/usual fiji/tonga? if (OldName == "AMD:AMDGPU:6:0:0") NewName = "amdgcn-amd-amdhsa--gfx600"; else if (OldName == "AMD:AMDGPU:6:0:1") NewName = "amdgcn-amd-amdhsa--gfx601"; else if (OldName == "AMD:AMDGPU:6:0:2") NewName = "amdgcn-amd-amdhsa--gfx602"; else if (OldName == "AMD:AMDGPU:7:0:0") NewName = "amdgcn-amd-amdhsa--gfx700"; else if (OldName == "AMD:AMDGPU:7:0:1") NewName = "amdgcn-amd-amdhsa--gfx701"; else if (OldName == "AMD:AMDGPU:7:0:2") NewName = "amdgcn-amd-amdhsa--gfx702"; else if (OldName == "AMD:AMDGPU:7:0:3") NewName = "amdgcn-amd-amdhsa--gfx703"; else if (OldName == "AMD:AMDGPU:7:0:4") NewName = "amdgcn-amd-amdhsa--gfx704"; else if (OldName == "AMD:AMDGPU:7:0:5") NewName = "amdgcn-amd-amdhsa--gfx705"; else if (OldName == "AMD:AMDGPU:8:0:1") { NewName = "amdgcn-amd-amdhsa--gfx801"; xnack_supported = true; } else if (OldName == "AMD:AMDGPU:8:0:0" || OldName == "AMD:AMDGPU:8:0:2") NewName = "amdgcn-amd-amdhsa--gfx802"; else if (OldName == "AMD:AMDGPU:8:0:3" || OldName == "AMD:AMDGPU:8:0:4") NewName = "amdgcn-amd-amdhsa--gfx803"; else if (OldName == "AMD:AMDGPU:8:0:5") NewName = "amdgcn-amd-amdhsa--gfx805"; else if (OldName == "AMD:AMDGPU:8:1:0") { NewName = "amdgcn-amd-amdhsa--gfx810"; xnack_supported = true; } else if (OldName == "AMD:AMDGPU:9:0:0" || OldName == "AMD:AMDGPU:9:0:1") { NewName = "amdgcn-amd-amdhsa--gfx900"; xnack_supported = true; } else if (OldName == "AMD:AMDGPU:9:0:2" || OldName == "AMD:AMDGPU:9:0:3") { NewName = "amdgcn-amd-amdhsa--gfx902"; xnack_supported = true; } else if (OldName == "AMD:AMDGPU:9:0:4" || OldName == "AMD:AMDGPU:9:0:5") { NewName = "amdgcn-amd-amdhsa--gfx904"; xnack_supported = true; } else if (OldName == "AMD:AMDGPU:9:0:6" || OldName == "AMD:AMDGPU:9:0:7") { NewName = "amdgcn-amd-amdhsa--gfx906"; xnack_supported = true; } else if (OldName == "AMD:AMDGPU:9:0:12") { NewName = "amdgcn-amd-amdhsa--gfx90c"; xnack_supported = true; } else { // Code object v2 only supports asics up to gfx906 plus gfx90c. Do NOT add // handling of new asics into this if-else-if* block. return ""; } if (IsFinalizer) { if (EFlags & ELF::EF_AMDGPU_FEATURE_XNACK_V2) NewName = NewName + ":xnack+"; else if (xnack_supported) NewName = NewName + ":xnack-"; } else { if (OldName == "AMD:AMDGPU:8:0:1") NewName = NewName + ":xnack+"; else if (OldName == "AMD:AMDGPU:8:1:0") NewName = NewName + ":xnack+"; else if (OldName == "AMD:AMDGPU:9:0:1") NewName = NewName + ":xnack+"; else if (OldName == "AMD:AMDGPU:9:0:3") NewName = NewName + ":xnack+"; else if (OldName == "AMD:AMDGPU:9:0:5") NewName = NewName + ":xnack+"; else if (OldName == "AMD:AMDGPU:9:0:7") NewName = NewName + ":xnack+"; else if (xnack_supported) NewName = NewName + ":xnack-"; } return NewName; } /* deprecated */ hsa_status_t hsa_code_object_get_info( hsa_code_object_t code_object, hsa_code_object_info_t attribute, void *value) { TRY; IS_OPEN(); IS_BAD_PTR(value); amd::hsa::code::AmdHsaCode *code = GetCodeManager()->FromHandle(code_object); if (!code) { return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; } switch (attribute) { case HSA_CODE_OBJECT_INFO_ISA: { char isa_name[64]; hsa_status_t status = code->GetInfo(attribute, &isa_name); if (status != HSA_STATUS_SUCCESS) { return status; } std::string isa_name_str(isa_name); bool IsFinalizer = true; uint32_t codeHsailMajor; uint32_t codeHsailMinor; hsa_profile_t codeProfile; hsa_machine_model_t codeMachineModel; hsa_default_float_rounding_mode_t codeRoundingMode; if (!code->GetNoteHsail(&codeHsailMajor, &codeHsailMinor, &codeProfile, &codeMachineModel, &codeRoundingMode)) { // Only finalizer generated the "HSAIL" note. IsFinalizer = false; } std::string new_isa_name_str = ConvertOldTargetNameToNew(isa_name_str, IsFinalizer, code->EFlags()); hsa_isa_t isa_handle = {0}; status = HSA::hsa_isa_from_name(new_isa_name_str.c_str(), &isa_handle); if (status != HSA_STATUS_SUCCESS) { return status; } *((hsa_isa_t*)value) = isa_handle; return HSA_STATUS_SUCCESS; } default: { return code->GetInfo(attribute, value); } } CATCH; } /* deprecated */ hsa_status_t hsa_code_object_get_symbol( hsa_code_object_t code_object, const char *symbol_name, hsa_code_symbol_t *symbol) { TRY; IS_OPEN(); IS_BAD_PTR(symbol_name); IS_BAD_PTR(symbol); amd::hsa::code::AmdHsaCode *code = GetCodeManager()->FromHandle(code_object); if (!code) { return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; } return code->GetSymbol(nullptr, symbol_name, symbol); CATCH; } /* deprecated */ hsa_status_t hsa_code_object_get_symbol_from_name( hsa_code_object_t code_object, const char *module_name, const char *symbol_name, hsa_code_symbol_t *symbol) { TRY; IS_OPEN(); IS_BAD_PTR(symbol_name); IS_BAD_PTR(symbol); amd::hsa::code::AmdHsaCode *code = GetCodeManager()->FromHandle(code_object); if (!code) { return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; } return code->GetSymbol(module_name, symbol_name, symbol); CATCH; } /* deprecated */ hsa_status_t hsa_code_symbol_get_info( hsa_code_symbol_t code_symbol, hsa_code_symbol_info_t attribute, void *value) { TRY; IS_OPEN(); IS_BAD_PTR(value); code::Symbol *symbol = code::Symbol::FromHandle(code_symbol); if (!symbol) { return HSA_STATUS_ERROR_INVALID_CODE_SYMBOL; } return symbol->GetInfo(attribute, value); CATCH; } /* deprecated */ hsa_status_t hsa_code_object_iterate_symbols( hsa_code_object_t code_object, hsa_status_t (*callback)(hsa_code_object_t code_object, hsa_code_symbol_t symbol, void *data), void *data) { TRY; IS_OPEN(); IS_BAD_PTR(callback); amd::hsa::code::AmdHsaCode *code = GetCodeManager()->FromHandle(code_object); if (!code) { return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; } return code->IterateSymbols(code_object, callback, data); CATCH; } //===--- Executable -------------------------------------------------------===// using amd::hsa::common::Signed; using amd::hsa::loader::Loader; using amd::hsa::loader::Executable; using amd::hsa::loader::CodeObjectReaderImpl; namespace { Loader *GetLoader() { return core::Runtime::runtime_singleton_->loader(); } } // namespace anonymous hsa_status_t hsa_code_object_reader_create_from_file( hsa_file_t file, hsa_code_object_reader_t *code_object_reader) { TRY; IS_OPEN(); IS_BAD_PTR(code_object_reader); std::unique_ptr reader( new (std::nothrow) CodeObjectReaderImpl()); CHECK_ALLOC(reader); hsa_status_t status = reader->SetFile(file); CHECK_STATUS(status); *code_object_reader = CodeObjectReaderImpl::Handle(reader.release()); return HSA_STATUS_SUCCESS; CATCH; } hsa_status_t hsa_code_object_reader_create_from_memory( const void *code_object, size_t size, hsa_code_object_reader_t *code_object_reader) { TRY; IS_OPEN(); IS_BAD_PTR(code_object); IS_BAD_PTR(code_object_reader); if (size == 0) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; } std::unique_ptr reader( new (std::nothrow) CodeObjectReaderImpl()); CHECK_ALLOC(reader); hsa_status_t status = reader->SetMemory(code_object, size); CHECK_STATUS(status); *code_object_reader = CodeObjectReaderImpl::Handle(reader.release()); return HSA_STATUS_SUCCESS; CATCH; } hsa_status_t hsa_code_object_reader_destroy( hsa_code_object_reader_t code_object_reader) { TRY; IS_OPEN(); CodeObjectReaderImpl *reader = CodeObjectReaderImpl::Object(code_object_reader); if (!reader) { return HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER; } delete reader; return HSA_STATUS_SUCCESS; CATCH; } /* deprecated */ hsa_status_t hsa_executable_create( hsa_profile_t profile, hsa_executable_state_t executable_state, const char *options, hsa_executable_t *executable) { TRY; IS_OPEN(); IS_BAD_PROFILE(profile); IS_BAD_EXECUTABLE_STATE(executable_state); IS_BAD_PTR(executable); // Invoke non-deprecated API. hsa_status_t status = HSA::hsa_executable_create_alt( profile, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, options, executable); if (status != HSA_STATUS_SUCCESS) { return status; } Executable *exec = Executable::Object(*executable); if (!exec) { return HSA_STATUS_ERROR_INVALID_EXECUTABLE; } if (executable_state == HSA_EXECUTABLE_STATE_FROZEN) { exec->Freeze(nullptr); } return HSA_STATUS_SUCCESS; CATCH; } hsa_status_t hsa_executable_create_alt( hsa_profile_t profile, hsa_default_float_rounding_mode_t default_float_rounding_mode, const char *options, hsa_executable_t *executable) { TRY; IS_OPEN(); IS_BAD_PROFILE(profile); IS_BAD_ROUNDING_MODE(default_float_rounding_mode); // NOTES: should we check // if default float // rounding mode is valid? // spec does not say so. IS_BAD_PTR(executable); Executable *exec = GetLoader()->CreateExecutable( profile, options, default_float_rounding_mode); CHECK_ALLOC(exec); *executable = Executable::Handle(exec); return HSA_STATUS_SUCCESS; CATCH; } hsa_status_t hsa_executable_destroy( hsa_executable_t executable) { TRY; IS_OPEN(); Executable *exec = Executable::Object(executable); if (!exec) { return HSA_STATUS_ERROR_INVALID_EXECUTABLE; } GetLoader()->DestroyExecutable(exec); return HSA_STATUS_SUCCESS; CATCH; } /* deprecated */ hsa_status_t hsa_executable_load_code_object( hsa_executable_t executable, hsa_agent_t agent, hsa_code_object_t code_object, const char *options) { TRY; IS_OPEN(); Executable *exec = Executable::Object(executable); if (!exec) { return HSA_STATUS_ERROR_INVALID_EXECUTABLE; } void *code_object_p = reinterpret_cast(code_object.handle); if (!code_object_p) { return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; } CodeObjectReaderImpl reader; reader.SetMemory(code_object_p, amd::elf::ElfSize(code_object_p)); return exec->LoadCodeObject(agent, code_object, options, reader.GetUri()); CATCH; } hsa_status_t hsa_executable_load_program_code_object( hsa_executable_t executable, hsa_code_object_reader_t code_object_reader, const char *options, hsa_loaded_code_object_t *loaded_code_object) { TRY; IS_OPEN(); Executable *exec = Executable::Object(executable); if (!exec) { return HSA_STATUS_ERROR_INVALID_EXECUTABLE; } CodeObjectReaderImpl *reader = CodeObjectReaderImpl::Object( code_object_reader); if (!reader) { return HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER; } hsa_code_object_t code_object = {reinterpret_cast(reader->GetCodeObjectMemory())}; return exec->LoadCodeObject( {0}, code_object, options, reader->GetUri(), loaded_code_object); CATCH; } hsa_status_t hsa_executable_load_agent_code_object( hsa_executable_t executable, hsa_agent_t agent, hsa_code_object_reader_t code_object_reader, const char *options, hsa_loaded_code_object_t *loaded_code_object) { TRY; IS_OPEN(); Executable *exec = Executable::Object(executable); if (!exec) { return HSA_STATUS_ERROR_INVALID_EXECUTABLE; } CodeObjectReaderImpl *reader = CodeObjectReaderImpl::Object( code_object_reader); if (!reader) { return HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER; } hsa_code_object_t code_object = {reinterpret_cast(reader->GetCodeObjectMemory())}; return exec->LoadCodeObject( agent, code_object, options, reader->GetUri(), loaded_code_object); CATCH; } hsa_status_t hsa_executable_freeze( hsa_executable_t executable, const char *options) { TRY; IS_OPEN(); Executable *exec = Executable::Object(executable); if (!exec) { return HSA_STATUS_ERROR_INVALID_EXECUTABLE; } return GetLoader()->FreezeExecutable(exec, options); CATCH; } hsa_status_t hsa_executable_get_info( hsa_executable_t executable, hsa_executable_info_t attribute, void *value) { TRY; IS_OPEN(); IS_BAD_PTR(value); Executable *exec = Executable::Object(executable); if (!exec) { return HSA_STATUS_ERROR_INVALID_EXECUTABLE; } return exec->GetInfo(attribute, value); CATCH; } hsa_status_t hsa_executable_global_variable_define( hsa_executable_t executable, const char *variable_name, void *address) { TRY; IS_OPEN(); IS_BAD_PTR(variable_name); Executable *exec = Executable::Object(executable); if (!exec) { return HSA_STATUS_ERROR_INVALID_EXECUTABLE; } return exec->DefineProgramExternalVariable(variable_name, address); CATCH; } hsa_status_t hsa_executable_agent_global_variable_define( hsa_executable_t executable, hsa_agent_t agent, const char *variable_name, void *address) { TRY; IS_OPEN(); IS_BAD_PTR(variable_name); Executable *exec = Executable::Object(executable); if (!exec) { return HSA_STATUS_ERROR_INVALID_EXECUTABLE; } return exec->DefineAgentExternalVariable( variable_name, agent, HSA_VARIABLE_SEGMENT_GLOBAL, address); CATCH; } hsa_status_t hsa_executable_readonly_variable_define( hsa_executable_t executable, hsa_agent_t agent, const char *variable_name, void *address) { TRY; IS_OPEN(); IS_BAD_PTR(variable_name); Executable *exec = Executable::Object(executable); if (!exec) { return HSA_STATUS_ERROR_INVALID_EXECUTABLE; } return exec->DefineAgentExternalVariable( variable_name, agent, HSA_VARIABLE_SEGMENT_READONLY, address); CATCH; } hsa_status_t hsa_executable_validate( hsa_executable_t executable, uint32_t *result) { TRY; IS_OPEN(); IS_BAD_PTR(result); Executable *exec = Executable::Object(executable); if (!exec) { return HSA_STATUS_ERROR_INVALID_EXECUTABLE; } return exec->Validate(result); CATCH; } hsa_status_t hsa_executable_validate_alt( hsa_executable_t executable, const char *options, uint32_t *result) { TRY; IS_OPEN(); IS_BAD_PTR(result); return HSA::hsa_executable_validate(executable, result); CATCH; } /* deprecated */ hsa_status_t hsa_executable_get_symbol( hsa_executable_t executable, const char *module_name, const char *symbol_name, hsa_agent_t agent, int32_t call_convention, hsa_executable_symbol_t *symbol) { TRY; IS_OPEN(); IS_BAD_PTR(symbol_name); IS_BAD_PTR(symbol); std::string mangled_name(symbol_name); if (mangled_name.empty()) { return HSA_STATUS_ERROR_INVALID_SYMBOL_NAME; } if (module_name && !std::string(module_name).empty()) { mangled_name.insert(0, "::"); mangled_name.insert(0, std::string(module_name)); } Executable *exec = Executable::Object(executable); if (!exec) { return HSA_STATUS_ERROR_INVALID_EXECUTABLE; } // Invoke non-deprecated API. return HSA::hsa_executable_get_symbol_by_name( executable, mangled_name.c_str(), exec->IsProgramSymbol(mangled_name.c_str()) ? nullptr : &agent, symbol); CATCH; } hsa_status_t hsa_executable_get_symbol_by_name( hsa_executable_t executable, const char *symbol_name, const hsa_agent_t *agent, // NOTES: this is not consistent with the rest of // of the specification, but seems like a better // approach to distinguish program/agent symbols. hsa_executable_symbol_t *symbol) { TRY; IS_OPEN(); IS_BAD_PTR(symbol_name); IS_BAD_PTR(symbol); Executable *exec = Executable::Object(executable); if (!exec) { return HSA_STATUS_ERROR_INVALID_EXECUTABLE; } loader::Symbol *sym = exec->GetSymbol(symbol_name, agent); if (!sym) { return HSA_STATUS_ERROR_INVALID_SYMBOL_NAME; } *symbol = loader::Symbol::Handle(sym); return HSA_STATUS_SUCCESS; CATCH; } hsa_status_t hsa_executable_symbol_get_info( hsa_executable_symbol_t executable_symbol, hsa_executable_symbol_info_t attribute, void *value) { TRY; IS_OPEN(); IS_BAD_PTR(value); loader::Symbol *sym = loader::Symbol::Object(executable_symbol); if (!sym) { return HSA_STATUS_ERROR_INVALID_EXECUTABLE_SYMBOL; } return sym->GetInfo(attribute, value) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR_INVALID_ARGUMENT; CATCH; } /* deprecated */ hsa_status_t hsa_executable_iterate_symbols( hsa_executable_t executable, hsa_status_t (*callback)(hsa_executable_t executable, hsa_executable_symbol_t symbol, void *data), void *data) { TRY; IS_OPEN(); IS_BAD_PTR(callback); Executable *exec = Executable::Object(executable); if (!exec) { return HSA_STATUS_ERROR_INVALID_EXECUTABLE; } return exec->IterateSymbols(callback, data); CATCH; } hsa_status_t hsa_executable_iterate_agent_symbols( hsa_executable_t executable, hsa_agent_t agent, hsa_status_t (*callback)(hsa_executable_t exec, hsa_agent_t agent, hsa_executable_symbol_t symbol, void *data), void *data) { TRY; IS_OPEN(); IS_BAD_PTR(callback); // NOTES: should we check if agent is valid? spec does not say so. const core::Agent *agent_object = core::Agent::Convert(agent); IS_VALID(agent_object); Executable *exec = Executable::Object(executable); if (!exec) { return HSA_STATUS_ERROR_INVALID_EXECUTABLE; } return exec->IterateAgentSymbols(agent, callback, data); CATCH; } hsa_status_t hsa_executable_iterate_program_symbols( hsa_executable_t executable, hsa_status_t (*callback)(hsa_executable_t exec, hsa_executable_symbol_t symbol, void *data), void *data) { TRY; IS_OPEN(); IS_BAD_PTR(callback); amd::hsa::loader::Executable *exec = amd::hsa::loader::Executable::Object(executable); if (!exec) { return HSA_STATUS_ERROR_INVALID_EXECUTABLE; } return exec->IterateProgramSymbols(callback, data); CATCH; } //===--- Runtime Notifications --------------------------------------------===// hsa_status_t hsa_status_string( hsa_status_t status, const char **status_string) { IS_BAD_PTR(status_string); const size_t status_u = static_cast(status); switch (status_u) { case HSA_STATUS_SUCCESS: *status_string = "HSA_STATUS_SUCCESS: The function has been executed successfully."; break; case HSA_STATUS_INFO_BREAK: *status_string = "HSA_STATUS_INFO_BREAK: A traversal over a list of elements has been interrupted by the " "application before completing."; break; case HSA_STATUS_ERROR: *status_string = "HSA_STATUS_ERROR: A generic error has occurred."; break; case HSA_STATUS_ERROR_INVALID_ARGUMENT: *status_string = "HSA_STATUS_ERROR_INVALID_ARGUMENT: One of the actual arguments does not meet a " "precondition stated in the documentation of the corresponding formal argument."; break; case HSA_STATUS_ERROR_INVALID_QUEUE_CREATION: *status_string = "HSA_STATUS_ERROR_INVALID_QUEUE_CREATION: The requested queue creation is not valid."; break; case HSA_STATUS_ERROR_INVALID_ALLOCATION: *status_string = "HSA_STATUS_ERROR_INVALID_ALLOCATION: The requested allocation is not valid."; break; case HSA_STATUS_ERROR_INVALID_AGENT: *status_string = "HSA_STATUS_ERROR_INVALID_AGENT: The agent is invalid."; break; case HSA_STATUS_ERROR_INVALID_REGION: *status_string = "HSA_STATUS_ERROR_INVALID_REGION: The memory region is invalid."; break; case HSA_STATUS_ERROR_INVALID_SIGNAL: *status_string = "HSA_STATUS_ERROR_INVALID_SIGNAL: The signal is invalid."; break; case HSA_STATUS_ERROR_INVALID_QUEUE: *status_string = "HSA_STATUS_ERROR_INVALID_QUEUE: The queue is invalid."; break; case HSA_STATUS_ERROR_OUT_OF_RESOURCES: *status_string = "HSA_STATUS_ERROR_OUT_OF_RESOURCES: The runtime failed to allocate the necessary " "resources. This error may also occur when the core runtime library needs to spawn " "threads or create internal OS-specific events."; break; case HSA_STATUS_ERROR_INVALID_PACKET_FORMAT: *status_string = "HSA_STATUS_ERROR_INVALID_PACKET_FORMAT: The AQL packet is malformed."; break; case HSA_STATUS_ERROR_RESOURCE_FREE: *status_string = "HSA_STATUS_ERROR_RESOURCE_FREE: An error has been detected while releasing a resource."; break; case HSA_STATUS_ERROR_NOT_INITIALIZED: *status_string = "HSA_STATUS_ERROR_NOT_INITIALIZED: An API other than hsa_init has been invoked while the " "reference count of the HSA runtime is zero."; break; case HSA_STATUS_ERROR_REFCOUNT_OVERFLOW: *status_string = "HSA_STATUS_ERROR_REFCOUNT_OVERFLOW: The maximum reference count for the object has been " "reached."; break; case HSA_STATUS_ERROR_INCOMPATIBLE_ARGUMENTS: *status_string = "HSA_STATUS_ERROR_INCOMPATIBLE_ARGUMENTS: The arguments passed to a functions are not " "compatible."; break; case HSA_STATUS_ERROR_INVALID_INDEX: *status_string = "HSA_STATUS_ERROR_INVALID_INDEX: The index is invalid."; break; case HSA_STATUS_ERROR_INVALID_ISA: *status_string = "HSA_STATUS_ERROR_INVALID_ISA: The instruction set architecture is invalid."; break; case HSA_STATUS_ERROR_INVALID_ISA_NAME: *status_string = "HSA_STATUS_ERROR_INVALID_ISA_NAME: The instruction set architecture name is invalid."; break; case HSA_STATUS_ERROR_INVALID_CODE_OBJECT: *status_string = "HSA_STATUS_ERROR_INVALID_CODE_OBJECT: The code object is invalid."; break; case HSA_STATUS_ERROR_INVALID_EXECUTABLE: *status_string = "HSA_STATUS_ERROR_INVALID_EXECUTABLE: The executable is invalid."; break; case HSA_STATUS_ERROR_FROZEN_EXECUTABLE: *status_string = "HSA_STATUS_ERROR_FROZEN_EXECUTABLE: The executable is frozen."; break; case HSA_STATUS_ERROR_INVALID_SYMBOL_NAME: *status_string = "HSA_STATUS_ERROR_INVALID_SYMBOL_NAME: There is no symbol with the given name."; break; case HSA_STATUS_ERROR_VARIABLE_ALREADY_DEFINED: *status_string = "HSA_STATUS_ERROR_VARIABLE_ALREADY_DEFINED: The variable is already defined."; break; case HSA_STATUS_ERROR_VARIABLE_UNDEFINED: *status_string = "HSA_STATUS_ERROR_VARIABLE_UNDEFINED: The variable is undefined."; break; case HSA_STATUS_ERROR_EXCEPTION: *status_string = "HSA_STATUS_ERROR_EXCEPTION: An HSAIL operation resulted in a hardware exception."; break; case HSA_STATUS_ERROR_INVALID_CODE_SYMBOL: *status_string = "HSA_STATUS_ERROR_INVALID_CODE_SYMBOL: The code object symbol is invalid."; break; case HSA_STATUS_ERROR_INVALID_EXECUTABLE_SYMBOL: *status_string = "HSA_STATUS_ERROR_INVALID_EXECUTABLE_SYMBOL: The executable symbol is invalid."; break; case HSA_STATUS_ERROR_INVALID_FILE: *status_string = "HSA_STATUS_ERROR_INVALID_FILE: The file descriptor is invalid."; break; case HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER: *status_string = "HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER: The code object reader is invalid."; break; case HSA_STATUS_ERROR_INVALID_CACHE: *status_string = "HSA_STATUS_ERROR_INVALID_CACHE: The cache is invalid."; break; case HSA_STATUS_ERROR_INVALID_WAVEFRONT: *status_string = "HSA_STATUS_ERROR_INVALID_WAVEFRONT: The wavefront is invalid."; break; case HSA_STATUS_ERROR_INVALID_SIGNAL_GROUP: *status_string = "HSA_STATUS_ERROR_INVALID_SIGNAL_GROUP: The signal group is invalid."; break; case HSA_STATUS_ERROR_INVALID_RUNTIME_STATE: *status_string = "HSA_STATUS_ERROR_INVALID_RUNTIME_STATE: The HSA runtime is not in the configuration " "state."; break; case HSA_STATUS_ERROR_FATAL: *status_string = "HSA_STATUS_ERROR_FATAL: The queue received an error that may require process " "termination."; break; case HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION: *status_string = "HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION: The agent attempted to access memory beyond " "the largest legal address."; break; case HSA_STATUS_ERROR_ILLEGAL_INSTRUCTION: *status_string = "HSA_STATUS_ERROR_ILLEGAL_INSTRUCTION: The agent attempted to execute an illegal shader " "instruction."; break; case HSA_EXT_STATUS_ERROR_IMAGE_FORMAT_UNSUPPORTED: *status_string = "HSA_EXT_STATUS_ERROR_IMAGE_FORMAT_UNSUPPORTED: Image format is not supported."; break; case HSA_EXT_STATUS_ERROR_IMAGE_SIZE_UNSUPPORTED: *status_string = "HSA_EXT_STATUS_ERROR_IMAGE_SIZE_UNSUPPORTED: Image size is not supported."; break; case HSA_EXT_STATUS_ERROR_IMAGE_PITCH_UNSUPPORTED: *status_string = "HSA_EXT_STATUS_ERROR_IMAGE_PITCH_UNSUPPORTED: Image pitch is not supported or invalid."; break; case HSA_EXT_STATUS_ERROR_SAMPLER_DESCRIPTOR_UNSUPPORTED: *status_string = "HSA_EXT_STATUS_ERROR_SAMPLER_DESCRIPTOR_UNSUPPORTED: Sampler descriptor is not " "supported or invalid."; break; case HSA_EXT_STATUS_ERROR_INVALID_PROGRAM: *status_string = "HSA_EXT_STATUS_ERROR_INVALID_PROGRAM: Invalid program"; break; case HSA_EXT_STATUS_ERROR_INVALID_MODULE: *status_string = "HSA_EXT_STATUS_ERROR_INVALID_MODULE: Invalid module"; break; case HSA_EXT_STATUS_ERROR_INCOMPATIBLE_MODULE: *status_string = "HSA_EXT_STATUS_ERROR_INCOMPATIBLE_MODULE: Incompatible module"; break; case HSA_EXT_STATUS_ERROR_MODULE_ALREADY_INCLUDED: *status_string = "HSA_EXT_STATUS_ERROR_MODULE_ALREADY_INCLUDED: Module already included"; break; case HSA_EXT_STATUS_ERROR_SYMBOL_MISMATCH: *status_string = "HSA_EXT_STATUS_ERROR_SYMBOL_MISMATCH: Symbol mismatch"; break; case HSA_EXT_STATUS_ERROR_FINALIZATION_FAILED: *status_string = "HSA_EXT_STATUS_ERROR_FINALIZATION_FAILED: Finalization failed"; break; case HSA_EXT_STATUS_ERROR_DIRECTIVE_MISMATCH: *status_string = "HSA_EXT_STATUS_ERROR_DIRECTIVE_MISMATCH: Directive mismatch"; break; case HSA_STATUS_ERROR_MEMORY_FAULT: *status_string = "HSA_STATUS_ERROR_MEMORY_FAULT: Agent attempted to access an inaccessible address."; break; case HSA_STATUS_ERROR_INVALID_MEMORY_POOL: *status_string = "HSA_STATUS_ERROR_INVALID_MEMORY_POOL: The memory pool is invalid."; break; case HSA_STATUS_CU_MASK_REDUCED: *status_string = "HSA_STATUS_CU_MASK_REDUCED: The CU mask was successfully set but the mask attempted to " "enable a CU which was disabled for the process. CUs disabled for the process remain " "disabled."; break; default: return HSA_STATUS_ERROR_INVALID_ARGUMENT; } return HSA_STATUS_SUCCESS; } } // namespace HSA } // namespace rocr