diff --git a/projects/clr/rocclr/cmake/ROCclrHSA.cmake b/projects/clr/rocclr/cmake/ROCclrHSA.cmake index 3ba360b24d..d28ca82e47 100644 --- a/projects/clr/rocclr/cmake/ROCclrHSA.cmake +++ b/projects/clr/rocclr/cmake/ROCclrHSA.cmake @@ -1,4 +1,4 @@ -# Copyright (c) 2020 - 2021 Advanced Micro Devices, Inc. All rights reserved. +# Copyright (c) 2020 - 2025 Advanced Micro Devices, Inc. All rights reserved. # # Permission is hereby granted, free of charge, to any person obtaining a copy # of this software and associated documentation files (the "Software"), to deal @@ -18,15 +18,54 @@ # OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN # THE SOFTWARE. -find_package(hsa-runtime64 1.11 REQUIRED CONFIG - PATHS - /opt/rocm/ - ${ROCM_INSTALL_PATH} - PATH_SUFFIXES - cmake/hsa-runtime64 - lib/cmake/hsa-runtime64 - lib64/cmake/hsa-runtime64) -target_link_libraries(rocclr PUBLIC hsa-runtime64::hsa-runtime64) +if(UNIX) + find_package(hsa-runtime64 1.11 REQUIRED CONFIG + PATHS + /opt/rocm/ + ${ROCM_INSTALL_PATH} + PATH_SUFFIXES + cmake/hsa-runtime64 + lib/cmake/hsa-runtime64 + lib64/cmake/hsa-runtime64) +else() + find_package(hsa-runtime64 CONFIG + PATHS + /opt/rocm/ + ${ROCM_INSTALL_PATH} + ${CMAKE_CURRENT_BINARY_DIR} + ${CMAKE_INSTALL_PREFIX} + ${CMAKE_INSTALL_PREFIX}/.. + PATH_SUFFIXES + rocr/lib/cmake/hsa-runtime64 + rocr/runtime/hsa-runtime + cmake/hsa-runtime64 + lib/cmake/hsa-runtime64 + lib64/cmake/hsa-runtime64) +endif() + +if (ROCR_DLL_LOAD) + find_path(AMD_HSA_INCLUDE_DIR hsa.h + HINTS + /opt/rocm + ${ROCM_INSTALL_PATH} + ${CMAKE_CURRENT_BINARY_DIR} + PATHS + ${CMAKE_CURRENT_BINARY_DIR}/.. + ${CMAKE_CURRENT_BINARY_DIR}/../.. + ${CMAKE_CURRENT_BINARY_DIR}/../../rocr + ${ROCCLR_SRC_DIR}/../../rocr-runtime/runtime/hsa-runtime + PATH_SUFFIXES + include + include/hsa + inc) + message("Roc CLR: " ${ROCCLR_SRC_DIR} "; HSA headers:" ${AMD_HSA_INCLUDE_DIR}) + target_compile_definitions(rocclr PUBLIC ROCR_DYN_DLL) + target_include_directories(rocclr PUBLIC ${AMD_HSA_INCLUDE_DIR}) +else() + target_link_libraries(rocclr PUBLIC hsa-runtime64::hsa-runtime64) +endif() + +#target_include_directories(rocclr PRIVATE ${AMD_HSA_INCLUDE_DIR}/..) find_package(NUMA) if(NUMA_FOUND) @@ -39,6 +78,7 @@ find_package(OpenGL REQUIRED) target_sources(rocclr PRIVATE ${ROCCLR_SRC_DIR}/device/rocm/rocappprofile.cpp + ${ROCCLR_SRC_DIR}/device/rocm/rocrctx.cpp ${ROCCLR_SRC_DIR}/device/rocm/rocblit.cpp ${ROCCLR_SRC_DIR}/device/rocm/rocblitcl.cpp ${ROCCLR_SRC_DIR}/device/rocm/roccounters.cpp diff --git a/projects/clr/rocclr/device/rocm/rocblit.cpp b/projects/clr/rocclr/device/rocm/rocblit.cpp index 7b8495eacb..a51a9163f1 100644 --- a/projects/clr/rocclr/device/rocm/rocblit.cpp +++ b/projects/clr/rocclr/device/rocm/rocblit.cpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2015 - 2025 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -317,8 +317,8 @@ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& d active.handle); hsa_status_t status = - hsa_amd_memory_async_copy_rect(&dstMem, &offset, &srcMem, &offset, &dim, agent, direction, - wait_events.size(), wait_events.data(), active); + Hsa::memory_async_copy_rect(&dstMem, &offset, &srcMem, &offset, &dim, agent, direction, + wait_events.size(), wait_events.data(), active); if (status != HSA_STATUS_SUCCESS) { gpu().Barriers().ResetCurrentSignal(); LogPrintfError("DMA buffer failed with code %d", status); @@ -338,7 +338,7 @@ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& d ClPrint(amd::LOG_DEBUG, amd::LOG_COPY2, "HSA Async Copy wait_event=0x%zx, completion_signal=0x%zx", (wait_events.size() != 0) ? wait_events[0].handle : 0, active.handle); - hsa_status_t status = hsa_amd_memory_async_copy( + hsa_status_t status = Hsa::memory_async_copy( (reinterpret_cast
(dst) + dstOffset), dstAgent, (reinterpret_cast(src) + srcOffset), srcAgent, size[0], wait_events.size(), wait_events.data(), active); @@ -385,8 +385,8 @@ bool DmaBlitManager::copyImageToBuffer(device::Memory& srcMemory, device::Memory image_region.range.y = size[1]; image_region.range.z = size[2]; - hsa_status_t status = hsa_ext_image_export(gpu().gpu_device(), srcImage.getHsaImageObject(), - dstHost, rowPitch, slicePitch, &image_region); + hsa_status_t status = Hsa::image_export(gpu().gpu_device(), srcImage.getHsaImageObject(), + dstHost, rowPitch, slicePitch, &image_region); result = (status == HSA_STATUS_SUCCESS) ? true : false; // hsa_ext_image_export need a system scope fence @@ -431,8 +431,8 @@ bool DmaBlitManager::copyBufferToImage(device::Memory& srcMemory, device::Memory image_region.range.y = size[1]; image_region.range.z = size[2]; - hsa_status_t status = hsa_ext_image_import(gpu().gpu_device(), srcHost, rowPitch, slicePitch, - dstImage.getHsaImageObject(), &image_region); + hsa_status_t status = Hsa::image_import(gpu().gpu_device(), srcHost, rowPitch, slicePitch, + dstImage.getHsaImageObject(), &image_region); result = (status == HSA_STATUS_SUCCESS) ? true : false; // hsa_ext_image_import need a system scope fence @@ -513,10 +513,10 @@ inline bool DmaBlitManager::rocrCopyBuffer(address dst, hsa_agent_t& dstAgent, c copyMask &= (engine == HwQueueEngine::SdmaRead ? sdmaEngineReadMask_ : sdmaEngineWriteMask_); if (copyMask == 0) { // Check SDMA engine status - status = hsa_amd_memory_copy_engine_status(dstAgent, srcAgent, &freeEngineMask); + status = Hsa::memory_copy_engine_status(dstAgent, srcAgent, &freeEngineMask); if (status == HSA_STATUS_SUCCESS) { - status = hsa_amd_memory_get_preferred_copy_engine(dstAgent, srcAgent, &recIdMask); + status = Hsa::memory_get_preferred_copy_engine(dstAgent, srcAgent, &recIdMask); } ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, @@ -552,9 +552,9 @@ inline bool DmaBlitManager::rocrCopyBuffer(address dst, hsa_agent_t& dstAgent, c copyEngine, dst, src, size, forceSDMA, engine, (wait_events.size() != 0) ? wait_events[0].handle : 0, active.handle); - status = hsa_amd_memory_async_copy_on_engine(dst, dstAgent, src, srcAgent, size, - wait_events.size(), wait_events.data(), active, - copyEngine, forceSDMA); + status = Hsa::memory_async_copy_on_engine(dst, dstAgent, src, srcAgent, size, + wait_events.size(), wait_events.data(), active, + copyEngine, forceSDMA); } else { kUseRegularCopyApi = true; } @@ -567,8 +567,8 @@ inline bool DmaBlitManager::rocrCopyBuffer(address dst, hsa_agent_t& dstAgent, c dst, src, size, (wait_events.size() != 0) ? wait_events[0].handle : 0, active.handle, engine); - status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent, size, wait_events.size(), - wait_events.data(), active); + status = Hsa::memory_async_copy(dst, dstAgent, src, srcAgent, size, wait_events.size(), + wait_events.data(), active); } if (status == HSA_STATUS_SUCCESS) { @@ -609,14 +609,14 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory, if (static_cast(srcMemory.owner())->ipcShared()) { hsa_amd_pointer_info_t info = {sizeof(hsa_amd_pointer_info_t)}; if (HSA_STATUS_SUCCESS == - hsa_amd_pointer_info(const_cast
(src), &info, nullptr, nullptr, nullptr)) { + Hsa::pointer_info(const_cast
(src), &info, nullptr, nullptr, nullptr)) { srcAgent = info.agentOwner; } } if (static_cast(dstMemory.owner())->ipcShared()) { hsa_amd_pointer_info_t info = {sizeof(hsa_amd_pointer_info_t)}; - if (HSA_STATUS_SUCCESS == hsa_amd_pointer_info(dst, &info, nullptr, nullptr, nullptr)) { + if (HSA_STATUS_SUCCESS == Hsa::pointer_info(dst, &info, nullptr, nullptr, nullptr)) { dstAgent = info.agentOwner; } } @@ -2737,7 +2737,7 @@ bool KernelBlitManager::runScheduler(uint64_t vqVM, hsa_queue_t* schedulerQueue, if (!dev().info().pcie_atomics_) { // Use a device side global atomics to workaround the reliance of PCIe 3 atomics - sp->write_index = hsa_queue_load_write_index_relaxed(schedulerQueue); + sp->write_index = Hsa::queue_load_write_index_relaxed(schedulerQueue); } else { sp->write_index = static_cast(-1ULL); } diff --git a/projects/clr/rocclr/device/rocm/roccounters.cpp b/projects/clr/rocclr/device/rocm/roccounters.cpp index 475732f7d1..46a9a1bd7c 100644 --- a/projects/clr/rocclr/device/rocm/roccounters.cpp +++ b/projects/clr/rocclr/device/rocm/roccounters.cpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2017 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2017 - 2025 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -565,7 +565,7 @@ bool PerfCounterProfile::initialize() { } // create the completion signal - if (hsa_signal_create(1, 0, nullptr, &completionSignal_) != HSA_STATUS_SUCCESS) { + if (Hsa::signal_create(1, 0, nullptr, &completionSignal_) != HSA_STATUS_SUCCESS) { LogError("Failed to create signal for profile counter"); return false; } @@ -604,7 +604,7 @@ hsa_ext_amd_aql_pm4_packet_t* PerfCounterProfile::createStopPacket() { PerfCounterProfile::~PerfCounterProfile() { if (completionSignal_.handle != 0) { - hsa_signal_destroy(completionSignal_); + Hsa::signal_destroy(completionSignal_); } if (profile_.command_buffer.ptr) { diff --git a/projects/clr/rocclr/device/rocm/roccounters.hpp b/projects/clr/rocclr/device/rocm/roccounters.hpp index 4ce2d70e41..e2d0f684dd 100644 --- a/projects/clr/rocclr/device/rocm/roccounters.hpp +++ b/projects/clr/rocclr/device/rocm/roccounters.hpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2017 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2017 - 2025 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -24,7 +24,6 @@ #include "top.hpp" #include "device/device.hpp" #include "device/rocm/rocdevice.hpp" -#include "hsa/hsa_ven_amd_aqlprofile.h" namespace amd::roc { @@ -108,19 +107,7 @@ class PerfCounterProfile : public amd::ReferenceCountedObject { //! Get the API tables bool Create() { hsa_agent_t agent = roc_device_.getBackendDevice(); - bool system_support, agent_support; - hsa_system_extension_supported(HSA_EXTENSION_AMD_AQLPROFILE, 1, 0, &system_support); - if (!system_support) { - LogError("HSA system does not support profile counter"); - return false; - } - hsa_agent_extension_supported(HSA_EXTENSION_AMD_AQLPROFILE, agent, 1, 0, &agent_support); - if (!agent_support) { - LogError("HSA agent does not support profile counter"); - return false; - } - - if (hsa_system_get_major_extension_table( + if (Hsa::system_get_major_extension_table( HSA_EXTENSION_AMD_AQLPROFILE, hsa_ven_amd_aqlprofile_VERSION_MAJOR, sizeof(hsa_ven_amd_aqlprofile_pfn_t), &api_) != HSA_STATUS_SUCCESS) { LogError("Failed to obtain aql profile extension function table"); diff --git a/projects/clr/rocclr/device/rocm/rocdevice.cpp b/projects/clr/rocclr/device/rocm/rocdevice.cpp index 7114688604..e7d1818f9e 100644 --- a/projects/clr/rocclr/device/rocm/rocdevice.cpp +++ b/projects/clr/rocclr/device/rocm/rocdevice.cpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2008 - 2024 Advanced Micro Devices, Inc. +/* Copyright (c) 2008 - 2025 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -243,7 +243,7 @@ Device::~Device() { ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_QUEUE, "Deleting hardware queue %p with refCount 0", queue->base_address); qIter = it.erase(qIter); - hsa_queue_destroy(queue); + Hsa::queue_destroy(queue); } } queuePool_.clear(); @@ -260,22 +260,13 @@ Device::~Device() { delete[] p2p_agents_list_; if (0 != prefetch_signal_.handle) { - hsa_signal_destroy(prefetch_signal_); + Hsa::signal_destroy(prefetch_signal_); } } -bool NullDevice::initCompiler(bool isOffline) { return true; } - -bool NullDevice::destroyCompiler() { return true; } - -void NullDevice::tearDown() { destroyCompiler(); } +void NullDevice::tearDown() {} bool NullDevice::init() { - // Initialize the compiler - if (!initCompiler(offlineDevice_)) { - return false; - } - // Create offline devices for all ISAs not already associated with an online // device. This allows code objects to be compiled for all supported ISAs. std::vector devices = getDevices(CL_DEVICE_TYPE_GPU, false); @@ -313,7 +304,7 @@ NullDevice::~NullDevice() {} hsa_status_t Device::iterateAgentCallback(hsa_agent_t agent, void* data) { hsa_device_type_t dev_type = HSA_DEVICE_TYPE_CPU; - hsa_status_t stat = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &dev_type); + hsa_status_t stat = Hsa::agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &dev_type); if (stat != HSA_STATUS_SUCCESS) { LogPrintfError("HSA_AGENT_INFO_DEVICE failed with %x", stat); @@ -322,8 +313,8 @@ hsa_status_t Device::iterateAgentCallback(hsa_agent_t agent, void* data) { if (dev_type == HSA_DEVICE_TYPE_CPU) { AgentInfo info = {agent, {0}, {0}, {0}}; - stat = hsa_amd_agent_iterate_memory_pools(agent, Device::iterateCpuMemoryPoolCallback, - reinterpret_cast(&info)); + stat = Hsa::agent_iterate_memory_pools(agent, Device::iterateCpuMemoryPoolCallback, + reinterpret_cast(&info)); if (stat == HSA_STATUS_SUCCESS) { cpu_agents_.push_back(info); } @@ -344,14 +335,12 @@ hsa_status_t Device::loaderQueryHostAddress(const void* device, const void** hos // ================================================================================================ bool Device::init() { - hsa_status_t status = HSA_STATUS_SUCCESS; - // Initialize the compiler - if (!initCompiler(offlineDevice_)) { - LogError("initCompiler failed."); + if (!Hsa::LoadLib()) { + LogPrintfWarning("Failed to load rocr library!"); return false; } - status = hsa_init(); + hsa_status_t status = Hsa::init(); // If there are no GPUs available, hsa_init will fail with HSA_STATUS_ERROR_OUT_OF_RESOURCES // but for NoGpu tests to pass, true needs to be returned @@ -366,10 +355,10 @@ bool Device::init() { return false; } - hsa_system_get_major_extension_table(HSA_EXTENSION_AMD_LOADER, 1, sizeof(amd_loader_ext_table), - &amd_loader_ext_table); + Hsa::system_get_major_extension_table(HSA_EXTENSION_AMD_LOADER, 1, sizeof(amd_loader_ext_table), + &amd_loader_ext_table); - status = hsa_iterate_agents(iterateAgentCallback, nullptr); + status = Hsa::iterate_agents(iterateAgentCallback, nullptr); if (status != HSA_STATUS_SUCCESS) { LogPrintfError("hsa_iterate_agents failed with %x", status); return false; @@ -398,8 +387,8 @@ bool Device::init() { auto agent = gpu_agents_[i]; char unique_id[32] = {0}; if (HSA_STATUS_SUCCESS == - hsa_agent_get_info(agent, static_cast(HSA_AMD_AGENT_INFO_UUID), - unique_id)) { + Hsa::agent_get_info(agent, static_cast(HSA_AMD_AGENT_INFO_UUID), + unique_id)) { if (std::string(unique_id).find(str_id) != std::string::npos) { str_id = std::to_string(i); break; @@ -446,7 +435,7 @@ bool Device::init() { // request via environment variable. By default the // System Memory is setup to be Coherent if (roc_device->settings().enableNCMode_) { - hsa_status_t err = hsa_amd_coherency_set_type(agent, HSA_AMD_COHERENCY_TYPE_NONCOHERENT); + hsa_status_t err = Hsa::coherency_set_type(agent, HSA_AMD_COHERENCY_TYPE_NONCOHERENT); if (err != HSA_STATUS_SUCCESS) { LogError("Unable to set NC memory policy!"); continue; @@ -524,20 +513,20 @@ extern const char* SchedulerSourceCode; void Device::tearDown() { NullDevice::tearDown(); - hsa_shut_down(); + Hsa::shut_down(); } // ================================================================================================ bool Device::create() { char agent_name[64] = {0}; - if (HSA_STATUS_SUCCESS != hsa_agent_get_info(bkendDevice_, HSA_AGENT_INFO_NAME, agent_name)) { + if (HSA_STATUS_SUCCESS != Hsa::agent_get_info(bkendDevice_, HSA_AGENT_INFO_NAME, agent_name)) { LogError("Unable to get HSA device name"); return false; } - if (HSA_STATUS_SUCCESS != hsa_agent_get_info(bkendDevice_, - (hsa_agent_info_t)HSA_AMD_AGENT_INFO_CHIP_ID, - &pciDeviceId_)) { + if (HSA_STATUS_SUCCESS != Hsa::agent_get_info(bkendDevice_, + (hsa_agent_info_t)HSA_AMD_AGENT_INFO_CHIP_ID, + &pciDeviceId_)) { LogPrintfError("Unable to get PCI ID of HSA device %s", agent_name); return false; } @@ -546,7 +535,7 @@ bool Device::create() { uint count; hsa_isa_t first_isa; } agent_isas = {0, {0}}; - if (HSA_STATUS_SUCCESS != hsa_agent_iterate_isas( + if (HSA_STATUS_SUCCESS != Hsa::agent_iterate_isas( bkendDevice_, [](hsa_isa_t isa, void* data) { agent_isas_t* agent_isas = static_cast(data); @@ -562,18 +551,18 @@ bool Device::create() { } uint32_t isa_name_length = 0; - if (HSA_STATUS_SUCCESS != hsa_isa_get_info_alt(agent_isas.first_isa, - (hsa_isa_info_t)HSA_ISA_INFO_NAME_LENGTH, - &isa_name_length)) { + if (HSA_STATUS_SUCCESS != Hsa::isa_get_info_alt(agent_isas.first_isa, + (hsa_isa_info_t)HSA_ISA_INFO_NAME_LENGTH, + &isa_name_length)) { LogPrintfError("Unable to get ISA name length for HSA device %s (PCI ID %x)", agent_name, pciDeviceId_); return false; } std::vector isa_name(isa_name_length + 1, '\0'); - if (HSA_STATUS_SUCCESS != hsa_isa_get_info_alt(agent_isas.first_isa, - (hsa_isa_info_t)HSA_ISA_INFO_NAME, - isa_name.data())) { + if (HSA_STATUS_SUCCESS != Hsa::isa_get_info_alt(agent_isas.first_isa, + (hsa_isa_info_t)HSA_ISA_INFO_NAME, + isa_name.data())) { LogPrintfError("Unable to get ISA name for HSA device %s (PCI ID %x)", agent_name, pciDeviceId_); return false; @@ -587,7 +576,7 @@ bool Device::create() { } if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, HSA_AGENT_INFO_PROFILE, &agent_profile_)) { + Hsa::agent_get_info(bkendDevice_, HSA_AGENT_INFO_PROFILE, &agent_profile_)) { LogPrintfError("Unable to get profile for HSA device %s (PCI ID %x)", agent_name, pciDeviceId_); return false; } @@ -596,9 +585,9 @@ bool Device::create() { // Check cooperative groups for HIP only if (amd::IS_HIP && (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, - static_cast(HSA_AMD_AGENT_INFO_COOPERATIVE_QUEUES), - &coop_groups))) { + Hsa::agent_get_info(bkendDevice_, + static_cast(HSA_AMD_AGENT_INFO_COOPERATIVE_QUEUES), + &coop_groups))) { LogPrintfError( "Unable to determine if cooperative queues are supported for HSA device %s (PCI ID %x)", agent_name, pciDeviceId_); @@ -610,8 +599,8 @@ bool Device::create() { // Get Agent HDP Flush Register Memory hsa_amd_hdp_flush_t hdpInfo; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, static_cast(HSA_AMD_AGENT_INFO_HDP_FLUSH), - &hdpInfo)) { + Hsa::agent_get_info(bkendDevice_, static_cast(HSA_AMD_AGENT_INFO_HDP_FLUSH), + &hdpInfo)) { LogPrintfError("Unable to determine HDP flush info for HSA device %s", agent_name); return false; } @@ -646,8 +635,8 @@ bool Device::create() { uint32_t hsa_bdf_id = 0; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, static_cast(HSA_AMD_AGENT_INFO_BDFID), - &hsa_bdf_id)) { + Hsa::agent_get_info(bkendDevice_, static_cast(HSA_AMD_AGENT_INFO_BDFID), + &hsa_bdf_id)) { LogPrintfError("Unable to determine BFD ID for HSA device %s (PCI ID %x)", agent_name, pciDeviceId_); return false; @@ -659,8 +648,8 @@ bool Device::create() { info_.deviceTopology_.pcie.function = (hsa_bdf_id & 0x07); uint32_t pci_domain_id = 0; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, static_cast(HSA_AMD_AGENT_INFO_DOMAIN), - &pci_domain_id)) { + Hsa::agent_get_info(bkendDevice_, static_cast(HSA_AMD_AGENT_INFO_DOMAIN), + &pci_domain_id)) { LogPrintfError("Unable to determine domain ID for HSA device %s (PCI ID %x)", agent_name, pciDeviceId_); return false; @@ -700,14 +689,14 @@ bool Device::create() { mapCache_->push_back(nullptr); // Create signal for HMM prefetch operation on device - if (HSA_STATUS_SUCCESS != hsa_signal_create(kInitSignalValueOne, 0, nullptr, &prefetch_signal_)) { + if (HSA_STATUS_SUCCESS != Hsa::signal_create(kInitSignalValueOne, 0, nullptr, &prefetch_signal_)) { return false; } if (AMD_LOG_LEVEL >= LOG_EXTRA_DEBUG) { uint8_t logMask[8] = {0}; hsa_flag_set64(logMask, HSA_AMD_LOG_FLAG_BLIT_KERNEL_PKTS); - hsa_amd_enable_logging(logMask, outFile); + Hsa::enable_logging(logMask, outFile); } return true; @@ -781,7 +770,7 @@ hsa_status_t Device::iterateGpuMemoryPoolCallback(hsa_amd_memory_pool_t pool, vo hsa_region_segment_t segment_type = (hsa_region_segment_t)0; hsa_status_t stat = - hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment_type); + Hsa::memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment_type); if (stat != HSA_STATUS_SUCCESS) { return stat; } @@ -793,7 +782,7 @@ hsa_status_t Device::iterateGpuMemoryPoolCallback(hsa_amd_memory_pool_t pool, vo if (dev->settings().enableLocalMemory_) { uint32_t global_flag = 0; hsa_status_t stat = - hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &global_flag); + Hsa::memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &global_flag); if (stat != HSA_STATUS_SUCCESS) { return stat; } @@ -811,8 +800,8 @@ hsa_status_t Device::iterateGpuMemoryPoolCallback(hsa_amd_memory_pool_t pool, vo // If cpu agent cannot access this pool, the device does not support large bar. hsa_amd_memory_pool_access_t tmp{}; - hsa_amd_agent_memory_pool_get_info(dev->cpu_agent_info_->agent, pool, - HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &tmp); + Hsa::agent_memory_pool_get_info(dev->cpu_agent_info_->agent, pool, + HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &tmp); if (tmp == HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED) { dev->info_.largeBar_ = false; @@ -821,8 +810,8 @@ hsa_status_t Device::iterateGpuMemoryPoolCallback(hsa_amd_memory_pool_t pool, vo } // Query the recommended granularity for this pool. - stat = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, - &(dev->info_.virtualMemAllocGranularity_)); + stat = Hsa::memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, + &(dev->info_.virtualMemAllocGranularity_)); if (stat != HSA_STATUS_SUCCESS) { LogPrintfError( "Cannot query HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE info" @@ -855,7 +844,7 @@ hsa_status_t Device::iterateCpuMemoryPoolCallback(hsa_amd_memory_pool_t pool, vo hsa_region_segment_t segment_type = (hsa_region_segment_t)0; hsa_status_t stat = - hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment_type); + Hsa::memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment_type); if (stat != HSA_STATUS_SUCCESS) { LogPrintfError("HSA_AMD_MEMORY_POOL_INFO_SEGMENT query failed with %x", stat); return stat; @@ -866,7 +855,7 @@ hsa_status_t Device::iterateCpuMemoryPoolCallback(hsa_amd_memory_pool_t pool, vo case HSA_REGION_SEGMENT_GLOBAL: { uint32_t global_flag = 0; stat = - hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &global_flag); + Hsa::memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &global_flag); if (stat != HSA_STATUS_SUCCESS) { LogPrintfError("HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS query failed with %x", stat); break; @@ -962,7 +951,7 @@ bool Sampler::create(const amd::Sampler& owner) { fillSampleDescriptor(samplerDescriptor, owner); hsa_status_t status = - hsa_ext_sampler_create_v2(dev_.getBackendDevice(), &samplerDescriptor, &hsa_sampler); + Hsa::sampler_create(dev_.getBackendDevice(), &samplerDescriptor, &hsa_sampler); if (HSA_STATUS_SUCCESS != status) { DevLogPrintfError("Sampler creation failed with status: %d \n", status); @@ -975,7 +964,7 @@ bool Sampler::create(const amd::Sampler& owner) { return true; } -Sampler::~Sampler() { hsa_ext_sampler_destroy(dev_.getBackendDevice(), hsa_sampler); } +Sampler::~Sampler() { Hsa::sampler_destroy(dev_.getBackendDevice(), hsa_sampler); } Memory* Device::getGpuMemory(amd::Memory* mem) const { return static_cast(mem->getDeviceMemory(*this)); @@ -996,15 +985,15 @@ bool Device::populateOCLDeviceConstants() { ::strncpy(info_.name_, isa().targetId(), sizeof(info_.name_) - 1); char device_name[64] = {0}; - if (HSA_STATUS_SUCCESS == hsa_agent_get_info(bkendDevice_, - (hsa_agent_info_t)HSA_AMD_AGENT_INFO_PRODUCT_NAME, - device_name)) { + if (HSA_STATUS_SUCCESS == Hsa::agent_get_info(bkendDevice_, + (hsa_agent_info_t)HSA_AMD_AGENT_INFO_PRODUCT_NAME, + device_name)) { ::strncpy(info_.boardName_, device_name, sizeof(info_.boardName_) - 1); } char unique_id[32] = {0}; if (HSA_STATUS_SUCCESS == - hsa_agent_get_info(bkendDevice_, static_cast(HSA_AMD_AGENT_INFO_UUID), + Hsa::agent_get_info(bkendDevice_, static_cast(HSA_AMD_AGENT_INFO_UUID), unique_id)) { // ROCr gives the UUID info in the format GPU-XXXX with length 20 bytes // Strip the first 4 bytes and store only the 16 bytes representing UUID @@ -1013,11 +1002,11 @@ bool Device::populateOCLDeviceConstants() { } } if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, - (amd::IS_HIP) - ? (hsa_agent_info_t)HSA_AMD_AGENT_INFO_COOPERATIVE_COMPUTE_UNIT_COUNT - : (hsa_agent_info_t)HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT, - &info_.maxComputeUnits_)) { + Hsa::agent_get_info(bkendDevice_, + (amd::IS_HIP) + ? (hsa_agent_info_t)HSA_AMD_AGENT_INFO_COOPERATIVE_COMPUTE_UNIT_COUNT + : (hsa_agent_info_t)HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT, + &info_.maxComputeUnits_)) { return false; } assert(info_.maxComputeUnits_ > 0); @@ -1026,8 +1015,8 @@ bool Device::populateOCLDeviceConstants() { settings().enableWgpMode_ ? info_.maxComputeUnits_ / 2 : info_.maxComputeUnits_; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT, - &info_.maxPhysicalComputeUnits_)) { + Hsa::agent_get_info(bkendDevice_, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT, + &info_.maxPhysicalComputeUnits_)) { return false; } assert(info_.maxPhysicalComputeUnits_ > 0); @@ -1035,9 +1024,9 @@ bool Device::populateOCLDeviceConstants() { info_.maxPhysicalComputeUnits_ = settings().enableWgpMode_ ? info_.maxPhysicalComputeUnits_ / 2 : info_.maxPhysicalComputeUnits_; - if (HSA_STATUS_SUCCESS != hsa_agent_get_info(bkendDevice_, - (hsa_agent_info_t)HSA_AMD_AGENT_INFO_CACHELINE_SIZE, - &info_.globalMemCacheLineSize_)) { + if (HSA_STATUS_SUCCESS != Hsa::agent_get_info(bkendDevice_, + (hsa_agent_info_t)HSA_AMD_AGENT_INFO_CACHELINE_SIZE, + &info_.globalMemCacheLineSize_)) { return false; } info_.globalMemCacheLineSize_ = @@ -1045,7 +1034,7 @@ bool Device::populateOCLDeviceConstants() { uint32_t cachesize[4] = {0}; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, HSA_AGENT_INFO_CACHE_SIZE, cachesize)) { + Hsa::agent_get_info(bkendDevice_, HSA_AGENT_INFO_CACHE_SIZE, cachesize)) { return false; } assert(cachesize[0] > 0); @@ -1060,8 +1049,8 @@ bool Device::populateOCLDeviceConstants() { (settings().doublePrecision_) ? 1 : 0; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_MAX_CLOCK_FREQUENCY, - &info_.maxEngineClockFrequency_)) { + Hsa::agent_get_info(bkendDevice_, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_MAX_CLOCK_FREQUENCY, + &info_.maxEngineClockFrequency_)) { return false; } @@ -1072,37 +1061,37 @@ bool Device::populateOCLDeviceConstants() { } if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_MEMORY_MAX_FREQUENCY, - &info_.maxMemoryClockFrequency_)) { + Hsa::agent_get_info(bkendDevice_, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_MEMORY_MAX_FREQUENCY, + &info_.maxMemoryClockFrequency_)) { return false; } uint64_t wallClockFrequency = 0; // in Hz if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_TIMESTAMP_FREQUENCY, - &wallClockFrequency)) { + Hsa::agent_get_info(bkendDevice_, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_TIMESTAMP_FREQUENCY, + &wallClockFrequency)) { LogWarning("HSA_AMD_AGENT_INFO_TIMESTAMP_FREQUENCY cannot be queried. Ignored!"); } info_.wallClockFrequency_ = static_cast(wallClockFrequency / 1000); // in KHz if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, - static_cast(HSA_AMD_AGENT_INFO_DRIVER_NODE_ID), - &info_.driverNodeId_)) { + Hsa::agent_get_info(bkendDevice_, + static_cast(HSA_AMD_AGENT_INFO_DRIVER_NODE_ID), + &info_.driverNodeId_)) { return false; } if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, - static_cast(HSA_AMD_AGENT_INFO_NUM_SDMA_ENG), - &info_.numSDMAengines_)) { + Hsa::agent_get_info(bkendDevice_, + static_cast(HSA_AMD_AGENT_INFO_NUM_SDMA_ENG), + &info_.numSDMAengines_)) { return false; } uint64_t scratchLimitMax = 0; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_SCRATCH_LIMIT_MAX, - &scratchLimitMax)) { + Hsa::agent_get_info(bkendDevice_, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_SCRATCH_LIMIT_MAX, + &scratchLimitMax)) { LogWarning("HSA_AMD_AGENT_INFO_SCRATCH_LIMIT_MAX cannot be queried!"); return false; } @@ -1112,7 +1101,7 @@ bool Device::populateOCLDeviceConstants() { checkAtomicSupport(); assert(cpu_agent_info_->fine_grain_pool.handle != 0); - if (HSA_STATUS_SUCCESS != hsa_amd_agent_iterate_memory_pools( + if (HSA_STATUS_SUCCESS != Hsa::agent_iterate_memory_pools( bkendDevice_, Device::iterateGpuMemoryPoolCallback, this)) { return false; } @@ -1124,8 +1113,8 @@ bool Device::populateOCLDeviceConstants() { hsa_status_t err; // Can another GPU (agent) have access to the current GPU memory pool (gpuvm_segment_)? hsa_amd_memory_pool_access_t access; - err = hsa_amd_agent_memory_pool_get_info(agent, gpuvm_segment_, - HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &access); + err = Hsa::agent_memory_pool_get_info(agent, gpuvm_segment_, + HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &access); if (err != HSA_STATUS_SUCCESS) { continue; } @@ -1147,23 +1136,23 @@ bool Device::populateOCLDeviceConstants() { } size_t group_segment_size = 0; - if (HSA_STATUS_SUCCESS != hsa_amd_memory_pool_get_info(group_segment_, - HSA_AMD_MEMORY_POOL_INFO_SIZE, - &group_segment_size)) { + if (HSA_STATUS_SUCCESS != Hsa::memory_pool_get_info(group_segment_, + HSA_AMD_MEMORY_POOL_INFO_SIZE, + &group_segment_size)) { return false; } assert(group_segment_size > 0); // Find SDMA read mask if (HSA_STATUS_SUCCESS != - hsa_amd_memory_copy_engine_status(getCpuAgent(), getBackendDevice(), &maxSdmaReadMask_)) { + Hsa::memory_copy_engine_status(getCpuAgent(), getBackendDevice(), &maxSdmaReadMask_)) { return false; } assert(maxSdmaReadMask_ > 0 && "No SDMA engines available for Read"); // Find SDMA write mask if (HSA_STATUS_SUCCESS != - hsa_amd_memory_copy_engine_status(getBackendDevice(), getCpuAgent(), &maxSdmaWriteMask_)) { + Hsa::memory_copy_engine_status(getBackendDevice(), getCpuAgent(), &maxSdmaWriteMask_)) { return false; } assert(maxSdmaWriteMask_ > 0 && "No SDMA engines available for Write"); @@ -1176,7 +1165,7 @@ bool Device::populateOCLDeviceConstants() { uint8_t memory_properties[8]; // Get the memory property from ROCr. if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_MEMORY_PROPERTIES, + Hsa::agent_get_info(bkendDevice_, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_MEMORY_PROPERTIES, memory_properties)) { LogError("HSA_AGENT_INFO_AMD_MEMORY_PROPERTIES query failed"); } @@ -1188,9 +1177,9 @@ bool Device::populateOCLDeviceConstants() { if (settings().enableLocalMemory_ && gpuvm_segment_.handle != 0) { size_t global_segment_size = 0; - if (HSA_STATUS_SUCCESS != hsa_amd_memory_pool_get_info(gpuvm_segment_, - HSA_AMD_MEMORY_POOL_INFO_SIZE, - &global_segment_size)) { + if (HSA_STATUS_SUCCESS != Hsa::memory_pool_get_info(gpuvm_segment_, + HSA_AMD_MEMORY_POOL_INFO_SIZE, + &global_segment_size)) { return false; } @@ -1213,8 +1202,8 @@ bool Device::populateOCLDeviceConstants() { info_.maxMemAllocSize_ = static_cast(gpuvm_segment_max_alloc_); if (HSA_STATUS_SUCCESS != - hsa_amd_memory_pool_get_info(gpuvm_segment_, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, - &alloc_granularity_)) { + Hsa::memory_pool_get_info(gpuvm_segment_, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, + &alloc_granularity_)) { return false; } @@ -1231,9 +1220,9 @@ bool Device::populateOCLDeviceConstants() { uint64_t(info_.globalMemSize_ * std::min(GPU_SINGLE_ALLOC_PERCENT, 100u) / 100u); if (HSA_STATUS_SUCCESS != - hsa_amd_memory_pool_get_info(cpu_agent_info_->fine_grain_pool, - HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, - &alloc_granularity_)) { + Hsa::memory_pool_get_info(cpu_agent_info_->fine_grain_pool, + HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, + &alloc_granularity_)) { return false; } } @@ -1252,7 +1241,7 @@ bool Device::populateOCLDeviceConstants() { uint32_t max_work_group_size = 0; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, HSA_AGENT_INFO_WORKGROUP_MAX_SIZE, &max_work_group_size)) { + Hsa::agent_get_info(bkendDevice_, HSA_AGENT_INFO_WORKGROUP_MAX_SIZE, &max_work_group_size)) { return false; } assert(max_work_group_size > 0); @@ -1262,7 +1251,7 @@ bool Device::populateOCLDeviceConstants() { uint16_t max_workgroup_size[3] = {0, 0, 0}; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, HSA_AGENT_INFO_WORKGROUP_MAX_DIM, &max_workgroup_size)) { + Hsa::agent_get_info(bkendDevice_, HSA_AGENT_INFO_WORKGROUP_MAX_DIM, &max_workgroup_size)) { return false; } assert(max_workgroup_size[0] != 0 && max_workgroup_size[1] != 0 && max_workgroup_size[2] != 0); @@ -1309,9 +1298,9 @@ bool Device::populateOCLDeviceConstants() { info_.spirVersions_ = ""; uint16_t major, minor; - if (hsa_agent_get_info(bkendDevice_, HSA_AGENT_INFO_VERSION_MAJOR, &major) != + if (Hsa::agent_get_info(bkendDevice_, HSA_AGENT_INFO_VERSION_MAJOR, &major) != HSA_STATUS_SUCCESS || - hsa_agent_get_info(bkendDevice_, HSA_AGENT_INFO_VERSION_MINOR, &minor) != + Hsa::agent_get_info(bkendDevice_, HSA_AGENT_INFO_VERSION_MINOR, &minor) != HSA_STATUS_SUCCESS) { return false; } @@ -1366,7 +1355,7 @@ bool Device::populateOCLDeviceConstants() { uint8_t hsa_extensions[128]; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, HSA_AGENT_INFO_EXTENSIONS, hsa_extensions)) { + Hsa::agent_get_info(bkendDevice_, HSA_AGENT_INFO_EXTENSIONS, hsa_extensions)) { return false; } @@ -1375,16 +1364,16 @@ bool Device::populateOCLDeviceConstants() { if (image_is_supported) { // Images if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, - static_cast(HSA_EXT_AGENT_INFO_MAX_SAMPLER_HANDLERS), - &info_.maxSamplers_)) { + Hsa::agent_get_info(bkendDevice_, + static_cast(HSA_EXT_AGENT_INFO_MAX_SAMPLER_HANDLERS), + &info_.maxSamplers_)) { return false; } if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, - static_cast(HSA_EXT_AGENT_INFO_MAX_IMAGE_RD_HANDLES), - &info_.maxReadImageArgs_)) { + Hsa::agent_get_info(bkendDevice_, + static_cast(HSA_EXT_AGENT_INFO_MAX_IMAGE_RD_HANDLES), + &info_.maxReadImageArgs_)) { return false; } @@ -1392,17 +1381,17 @@ bool Device::populateOCLDeviceConstants() { info_.maxWriteImageArgs_ = 8; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, - static_cast(HSA_EXT_AGENT_INFO_MAX_IMAGE_RORW_HANDLES), - &info_.maxReadWriteImageArgs_)) { + Hsa::agent_get_info(bkendDevice_, + static_cast(HSA_EXT_AGENT_INFO_MAX_IMAGE_RORW_HANDLES), + &info_.maxReadWriteImageArgs_)) { return false; } uint32_t image_max_dim[3]; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, - static_cast(HSA_EXT_AGENT_INFO_IMAGE_2D_MAX_ELEMENTS), - &image_max_dim)) { + Hsa::agent_get_info(bkendDevice_, + static_cast(HSA_EXT_AGENT_INFO_IMAGE_2D_MAX_ELEMENTS), + &image_max_dim)) { return false; } @@ -1410,9 +1399,9 @@ bool Device::populateOCLDeviceConstants() { info_.image2DMaxHeight_ = image_max_dim[1]; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, - static_cast(HSA_EXT_AGENT_INFO_IMAGE_3D_MAX_ELEMENTS), - &image_max_dim)) { + Hsa::agent_get_info(bkendDevice_, + static_cast(HSA_EXT_AGENT_INFO_IMAGE_3D_MAX_ELEMENTS), + &image_max_dim)) { return false; } @@ -1422,9 +1411,9 @@ bool Device::populateOCLDeviceConstants() { uint32_t max_array_size = 0; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, - static_cast(HSA_EXT_AGENT_INFO_IMAGE_ARRAY_MAX_LAYERS), - &max_array_size)) { + Hsa::agent_get_info(bkendDevice_, + static_cast(HSA_EXT_AGENT_INFO_IMAGE_ARRAY_MAX_LAYERS), + &max_array_size)) { return false; } @@ -1432,9 +1421,9 @@ bool Device::populateOCLDeviceConstants() { uint32_t max_image1da_width = 0; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, - static_cast(HSA_EXT_AGENT_INFO_IMAGE_1DA_MAX_ELEMENTS), - &max_image1da_width)) { + Hsa::agent_get_info(bkendDevice_, + static_cast(HSA_EXT_AGENT_INFO_IMAGE_1DA_MAX_ELEMENTS), + &max_image1da_width)) { return false; } @@ -1442,9 +1431,9 @@ bool Device::populateOCLDeviceConstants() { uint32_t max_image2da_width[2] = {0, 0}; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, - static_cast(HSA_EXT_AGENT_INFO_IMAGE_2DA_MAX_ELEMENTS), - &max_image2da_width)) { + Hsa::agent_get_info(bkendDevice_, + static_cast(HSA_EXT_AGENT_INFO_IMAGE_2DA_MAX_ELEMENTS), + &max_image2da_width)) { return false; } @@ -1453,17 +1442,17 @@ bool Device::populateOCLDeviceConstants() { uint32_t max_image1d_width = 0; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, - static_cast(HSA_EXT_AGENT_INFO_IMAGE_1D_MAX_ELEMENTS), - &max_image1d_width)) { + Hsa::agent_get_info(bkendDevice_, + static_cast(HSA_EXT_AGENT_INFO_IMAGE_1D_MAX_ELEMENTS), + &max_image1d_width)) { return false; } info_.image1DMaxWidth_ = max_image1d_width; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, - static_cast(HSA_EXT_AGENT_INFO_IMAGE_1DB_MAX_ELEMENTS), - &image_max_dim)) { + Hsa::agent_get_info(bkendDevice_, + static_cast(HSA_EXT_AGENT_INFO_IMAGE_1DB_MAX_ELEMENTS), + &image_max_dim)) { return false; } info_.imageMaxBufferSize_ = (amd::IS_HIP) ? image_max_dim[0] : (1 << 27); @@ -1502,31 +1491,31 @@ bool Device::populateOCLDeviceConstants() { info_.simdWidth_ = isa().simdWidth(); info_.simdInstructionWidth_ = isa().simdInstructionWidth(); if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, HSA_AGENT_INFO_WAVEFRONT_SIZE, &info_.wavefrontWidth_)) { + Hsa::agent_get_info(bkendDevice_, HSA_AGENT_INFO_WAVEFRONT_SIZE, &info_.wavefrontWidth_)) { return false; } if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, - static_cast(HSA_AMD_AGENT_INFO_MEMORY_WIDTH), - &info_.vramBusBitWidth_)) { + Hsa::agent_get_info(bkendDevice_, + static_cast(HSA_AMD_AGENT_INFO_MEMORY_WIDTH), + &info_.vramBusBitWidth_)) { return false; } info_.globalMemChannels_ = info_.vramBusBitWidth_ / 32; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, - static_cast(HSA_AMD_AGENT_INFO_NUM_SIMDS_PER_CU), - &info_.simdPerCU_)) { + Hsa::agent_get_info(bkendDevice_, + static_cast(HSA_AMD_AGENT_INFO_NUM_SIMDS_PER_CU), + &info_.simdPerCU_)) { return false; } uint32_t max_waves_per_cu = 0; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, - static_cast(HSA_AMD_AGENT_INFO_MAX_WAVES_PER_CU), - &max_waves_per_cu)) { + Hsa::agent_get_info(bkendDevice_, + static_cast(HSA_AMD_AGENT_INFO_MAX_WAVES_PER_CU), + &max_waves_per_cu)) { return false; } @@ -1539,16 +1528,16 @@ bool Device::populateOCLDeviceConstants() { uint32_t cache_sizes[4]; /* FIXIT [skudchad] - Seems like hardcoded in HSA backend so 0*/ if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, static_cast(HSA_AGENT_INFO_CACHE_SIZE), - cache_sizes)) { + Hsa::agent_get_info(bkendDevice_, static_cast(HSA_AGENT_INFO_CACHE_SIZE), + cache_sizes)) { return false; } uint32_t asic_revision = 0; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, - static_cast(HSA_AMD_AGENT_INFO_ASIC_REVISION), - &asic_revision)) { + Hsa::agent_get_info(bkendDevice_, + static_cast(HSA_AMD_AGENT_INFO_ASIC_REVISION), + &asic_revision)) { return false; } info_.asicRevision_ = asic_revision; @@ -1619,27 +1608,27 @@ bool Device::populateOCLDeviceConstants() { // Generic support for HMM interfaces if (HSA_STATUS_SUCCESS != - hsa_system_get_info(HSA_AMD_SYSTEM_INFO_SVM_SUPPORTED, &info_.hmmSupported_)) { + Hsa::system_get_info(HSA_AMD_SYSTEM_INFO_SVM_SUPPORTED, &info_.hmmSupported_)) { LogError("HSA_AMD_SYSTEM_INFO_SVM_SUPPORTED query failed. HMM will be disabled"); } // This capability should be available with xnack enabled - if (HSA_STATUS_SUCCESS != hsa_system_get_info(HSA_AMD_SYSTEM_INFO_SVM_ACCESSIBLE_BY_DEFAULT, - &info_.hmmCpuMemoryAccessible_)) { + if (HSA_STATUS_SUCCESS != Hsa::system_get_info(HSA_AMD_SYSTEM_INFO_SVM_ACCESSIBLE_BY_DEFAULT, + &info_.hmmCpuMemoryAccessible_)) { LogError("HSA_AMD_SYSTEM_INFO_SVM_ACCESSIBLE_BY_DEFAULT query failed."); } // HMM specific capability for CPU direct access to device memory if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, - static_cast(HSA_AMD_AGENT_INFO_SVM_DIRECT_HOST_ACCESS), - &info_.hmmDirectHostAccess_)) { + Hsa::agent_get_info(bkendDevice_, + static_cast(HSA_AMD_AGENT_INFO_SVM_DIRECT_HOST_ACCESS), + &info_.hmmDirectHostAccess_)) { LogError("HSA_AMD_AGENT_INFO_SVM_DIRECT_HOST_ACCESS query failed."); } if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, static_cast(HSA_AMD_AGENT_INFO_NUM_XCC), - &info_.numberOfXccs_)) { + Hsa::agent_get_info(bkendDevice_, static_cast(HSA_AMD_AGENT_INFO_NUM_XCC), + &info_.numberOfXccs_)) { LogError("HSA_AMD_AGENT_INFO_NUM_XCC query failed."); } @@ -1656,7 +1645,7 @@ bool Device::populateOCLDeviceConstants() { info_.virtualMemoryManagement_ = false; if (HIP_VMEM_MANAGE_SUPPORT) { if (HSA_STATUS_SUCCESS != - hsa_system_get_info( + Hsa::system_get_info( static_cast(HSA_AMD_SYSTEM_INFO_VIRTUAL_MEM_API_SUPPORTED), &info_.virtualMemoryManagement_)) { LogError("HSA_AMD_SYSTEM_INFO_VIRTUAL_MEM_API_SUPPORTED query failed "); @@ -1711,9 +1700,9 @@ bool Device::globalFreeMemory(size_t* freeMemory) const { uint64_t globalAvailMemory; // Queries memory available in bytes across all global pools owned by the agent if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, - static_cast(HSA_AMD_AGENT_INFO_MEMORY_AVAIL), - &globalAvailMemory)) { + Hsa::agent_get_info(bkendDevice_, + static_cast(HSA_AMD_AGENT_INFO_MEMORY_AVAIL), + &globalAvailMemory)) { LogError("HSA_AMD_AGENT_INFO_MEMORY_AVAIL query failed."); return false; } @@ -2015,7 +2004,8 @@ void* Device::hostAlloc(size_t size, size_t alignment, MemorySegment mem_seg, } hsa_amd_memory_pool_t pool = getHostMemoryPool(mem_seg, static_cast(agentInfo)); - hsa_status_t stat = hsa_amd_memory_pool_allocate(pool, size, memFlags, &ptr); + hsa_status_t stat = Hsa::memory_pool_allocate(pool, size, memFlags, &ptr); + ClPrint(amd::LOG_DEBUG, amd::LOG_MEM, "Allocate hsa host memory %p, size 0x%zx," " numa_node = %d, mem_seg = %d", @@ -2025,7 +2015,7 @@ void* Device::hostAlloc(size_t size, size_t alignment, MemorySegment mem_seg, return nullptr; } - stat = hsa_amd_agents_allow_access(gpu_agents_.size(), &gpu_agents_[0], nullptr, ptr); + stat = Hsa::agents_allow_access(gpu_agents_.size(), &gpu_agents_[0], nullptr, ptr); if (stat != HSA_STATUS_SUCCESS) { LogPrintfError("Fail hsa_amd_agents_allow_access with err %d", stat); hostFree(ptr, size); @@ -2079,7 +2069,7 @@ void* Device::hostNumaAlloc(size_t size, size_t alignment, MemorySegment mem_seg void* Device::hostLock(void* hostMem, size_t size, const MemorySegment memSegment) const { hsa_amd_memory_pool_t pool = getHostMemoryPool(memSegment); void* deviceMemory = nullptr; - hsa_status_t status = hsa_amd_memory_lock_to_pool( + hsa_status_t status = Hsa::memory_lock_to_pool( hostMem, size, const_cast(&bkendDevice_), 1, pool, 0, &deviceMemory); ClPrint(amd::LOG_DEBUG, amd::LOG_MEM, "Locking to pool %p, size 0x%zx, hostMem = %p," @@ -2098,7 +2088,7 @@ bool Device::deviceAllowAccess(void* ptr) const { std::lock_guard lock(lock_allow_access_); if (!p2pAgents().empty()) { hsa_status_t stat = - hsa_amd_agents_allow_access(p2pAgents().size(), p2pAgents().data(), nullptr, ptr); + Hsa::agents_allow_access(p2pAgents().size(), p2pAgents().data(), nullptr, ptr); if (stat != HSA_STATUS_SUCCESS) { LogPrintfError("Allow p2p access failed - hsa_amd_agents_allow_access with err %d", stat); return false; @@ -2114,7 +2104,7 @@ bool Device::allowPeerAccess(device::Memory* memory) const { if (!p2pAgents().empty()) { void* ptr = reinterpret_cast(memory->virtualAddress()); hsa_agent_t agent = getBackendDevice(); - hsa_status_t stat = hsa_amd_agents_allow_access(1, &agent, nullptr, ptr); + hsa_status_t stat = Hsa::agents_allow_access(1, &agent, nullptr, ptr); if (stat != HSA_STATUS_SUCCESS) { LogPrintfError("Allow p2p access failed - hsa_amd_agents_allow_access with err: %d", stat); return false; @@ -2128,7 +2118,7 @@ uint64_t Device::deviceVmemAlloc(size_t size, uint64_t flags) const { // We only allow pinned memory at this time. hsa_status_t hsa_status = - hsa_amd_vmem_handle_create(gpuvm_segment_, size, MEMORY_TYPE_PINNED, flags, &hsa_vmem_handle); + Hsa::vmem_handle_create(gpuvm_segment_, size, MEMORY_TYPE_PINNED, flags, &hsa_vmem_handle); if (hsa_status != HSA_STATUS_SUCCESS) { LogPrintfError("Failed hsa_amd_vmem_handle_create! Failed with hsa status: %d \n", hsa_status); } @@ -2140,7 +2130,7 @@ void Device::deviceVmemRelease(uint64_t mem_handle) const { hsa_amd_vmem_alloc_handle_t hsa_vmem_handle{}; hsa_vmem_handle.handle = mem_handle; - hsa_status_t hsa_status = hsa_amd_vmem_handle_release(hsa_vmem_handle); + hsa_status_t hsa_status = Hsa::vmem_handle_release(hsa_vmem_handle); if (hsa_status != HSA_STATUS_SUCCESS) { LogPrintfError("Failed hsa_amd_vmem_handle_release! Failed with hsa status: %d \n", hsa_status); } @@ -2149,8 +2139,8 @@ void Device::deviceVmemRelease(uint64_t mem_handle) const { void* Device::reserveMemory(size_t size, size_t alignment) const { void* ptr = nullptr; // Reserves non registered VA memory using HSA APIs. - hsa_status_t status = hsa_amd_vmem_address_reserve_align(&ptr, size, 0, alignment, - HSA_AMD_VMEM_ADDRESS_NO_REGISTER); + hsa_status_t status = Hsa::vmem_address_reserve_align(&ptr, size, 0, alignment, + HSA_AMD_VMEM_ADDRESS_NO_REGISTER); ClPrint(amd::LOG_DEBUG, amd::LOG_MEM, "Reserve hsa device memory %p, size 0x%zx", ptr, size); if (status != HSA_STATUS_SUCCESS) { LogError("Fail to reserve memory"); @@ -2160,7 +2150,7 @@ void* Device::reserveMemory(size_t size, size_t alignment) const { } void Device::releaseMemory(void* ptr, size_t size) const { - hsa_status_t hsa_status = hsa_amd_vmem_address_free(ptr, size); + hsa_status_t hsa_status = Hsa::vmem_address_free(ptr, size); ClPrint(amd::LOG_DEBUG, amd::LOG_MEM, "Free hsa reserved memory %p", ptr); if (hsa_status != HSA_STATUS_SUCCESS) { LogError("hsa_amd_vmem_address_free failed \n"); @@ -2189,7 +2179,7 @@ void* Device::deviceLocalAlloc(size_t size, const AllocationFlags& flags) const } void* ptr = nullptr; - hsa_status_t stat = hsa_amd_memory_pool_allocate(pool, size, hsa_mem_flags, &ptr); + hsa_status_t stat = Hsa::memory_pool_allocate(pool, size, hsa_mem_flags, &ptr); ClPrint(amd::LOG_DEBUG, amd::LOG_MEM, "Allocate hsa device memory %p, size 0x%zx, hsa_mem_flags 0x%xh", ptr, size, hsa_mem_flags); @@ -2207,7 +2197,7 @@ void* Device::deviceLocalAlloc(size_t size, const AllocationFlags& flags) const } void Device::memFree(void* ptr, size_t size) const { - hsa_status_t stat = hsa_amd_memory_pool_free(ptr); + hsa_status_t stat = Hsa::memory_pool_free(ptr); ClPrint(amd::LOG_DEBUG, amd::LOG_MEM, "Free hsa memory %p", ptr); if (stat != HSA_STATUS_SUCCESS) { LogError("Fail freeing local memory"); @@ -2284,7 +2274,7 @@ void* Device::virtualAlloc(void* req_addr, size_t size, size_t alignment) { // Reserves the address using HSA APIs, with requested address. // There is no guarantee that we will get the requested address. hsa_status_t hsa_status = - hsa_amd_vmem_address_reserve(&vptr, size, reinterpret_cast(req_addr), 0); + Hsa::vmem_address_reserve(&vptr, size, reinterpret_cast(req_addr), 0); if (hsa_status != HSA_STATUS_SUCCESS) { LogPrintfError("Failed hsa_amd_vmem_address_reserve. Failed with status: %d \n", hsa_status); return nullptr; @@ -2309,7 +2299,7 @@ bool Device::virtualFree(void* addr) { return false; } - hsa_status_t hsa_status = hsa_amd_vmem_address_free(memObj->getSvmPtr(), memObj->getSize()); + hsa_status_t hsa_status = Hsa::vmem_address_free(memObj->getSvmPtr(), memObj->getSize()); if (hsa_status != HSA_STATUS_SUCCESS) { LogPrintfError("Failed hsa_amd_vmem_address_free. Failed with status:%d \n", hsa_status); return false; @@ -2325,7 +2315,7 @@ bool Device::SetMemAccess(void* va_addr, size_t va_size, VmmAccess access_flags, desc.agent_handle = access_location == VmmLocationType::kDevice ? getBackendDevice() : getCpuAgent(); - if ((hsa_status = hsa_amd_vmem_set_access(va_addr, va_size, &desc, 1)) != HSA_STATUS_SUCCESS) { + if ((hsa_status = Hsa::vmem_set_access(va_addr, va_size, &desc, 1)) != HSA_STATUS_SUCCESS) { LogPrintfError("Failed hsa_amd_vmem_set_access. Failed with status:%d \n", hsa_status); return false; } @@ -2344,7 +2334,7 @@ bool Device::GetMemAccess(void* va_addr, VmmAccess* access_flags_ptr) const { return false; } - if ((hsa_status = hsa_amd_vmem_get_access(va_mem_obj->getSvmPtr(), &perms, getBackendDevice())) != + if ((hsa_status = Hsa::vmem_get_access(va_mem_obj->getSvmPtr(), &perms, getBackendDevice())) != HSA_STATUS_SUCCESS) { LogPrintfError("Failed hsa_amd_vmem_get_access. Failed with status:%d \n", hsa_status); return false; @@ -2367,7 +2357,7 @@ bool Device::ExportShareableVMMHandle(amd::Memory& amd_mem_obj, int flags, void* return false; } - if ((hsa_status = hsa_amd_vmem_export_shareable_handle(&dmabuf_fd, hsa_vmem_handle, flags)) != + if ((hsa_status = Hsa::vmem_export_shareable_handle(&dmabuf_fd, hsa_vmem_handle, flags)) != HSA_STATUS_SUCCESS) { LogPrintfError("Failed hsa_vmem_export_shareable_handle with status: %d \n", hsa_status); return false; @@ -2389,7 +2379,7 @@ bool Device::ImportShareableHSAHandle(void* osHandle, uint64_t* hsa_handle_ptr) } int dmabuf_fd = static_cast(reinterpret_cast(osHandle)); - if ((hsa_status = hsa_amd_vmem_import_shareable_handle(dmabuf_fd, &hsa_vmem_handle)) != + if ((hsa_status = Hsa::vmem_import_shareable_handle(dmabuf_fd, &hsa_vmem_handle)) != HSA_STATUS_SUCCESS) { LogPrintfError("Failed hsa_amd_vmem_import_shareable_handle with status: %d \n", hsa_status); return false; @@ -2489,7 +2479,7 @@ bool Device::SetSvmAttributesInt(const void* dev_ptr, size_t count, amd::MemoryA } hsa_status_t status = - hsa_amd_svm_attributes_set(const_cast(dev_ptr), count, attr.data(), attr.size()); + Hsa::svm_attributes_set(const_cast(dev_ptr), count, attr.data(), attr.size()); if (status != HSA_STATUS_SUCCESS) { LogPrintfError("hsa_amd_svm_attributes_set() failed. Advice: %d, status: %d", advice, status); return false; @@ -2527,7 +2517,7 @@ bool Device::GetSvmAttributes(void** data, size_t* data_sizes, int* attributes, ptr_info.size = sizeof(hsa_amd_pointer_info_t); // Query ptr type to see if it's a HMM allocation hsa_status_t status = - hsa_amd_pointer_info(const_cast(dev_ptr), &ptr_info, nullptr, nullptr, nullptr); + Hsa::pointer_info(const_cast(dev_ptr), &ptr_info, nullptr, nullptr, nullptr); // The call should never fail in ROCR, but just check for an error and continue if (status != HSA_STATUS_SUCCESS) { LogError("hsa_amd_pointer_info() failed"); @@ -2586,7 +2576,7 @@ bool Device::GetSvmAttributes(void** data, size_t* data_sizes, int* attributes, } hsa_status_t status = - hsa_amd_svm_attributes_get(const_cast(dev_ptr), count, attr.data(), attr.size()); + Hsa::svm_attributes_get(const_cast(dev_ptr), count, attr.data(), attr.size()); if (status != HSA_STATUS_SUCCESS) { LogError("hsa_amd_svm_attributes_get() failed"); return false; @@ -2703,7 +2693,7 @@ bool Device::GetSvmAttributes(void** data, size_t* data_sizes, int* attributes, size_t Device::ScratchLimitCurrent() const { uint64_t scratchLimitCurrent = 0; hsa_status_t ret = - hsa_agent_get_info(bkendDevice_, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_SCRATCH_LIMIT_CURRENT, + Hsa::agent_get_info(bkendDevice_, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_SCRATCH_LIMIT_CURRENT, &scratchLimitCurrent); if (HSA_STATUS_SUCCESS != ret) { LogPrintfError("HSA_AMD_AGENT_INFO_SCRATCH_LIMIT_CURRENT cannot be queried! Err: 0x%xh", ret); @@ -2713,7 +2703,7 @@ size_t Device::ScratchLimitCurrent() const { }; bool Device::UpdateScratchLimitCurrent(size_t limit) const { - hsa_status_t ret = hsa_amd_agent_set_async_scratch_limit(bkendDevice_, limit); + hsa_status_t ret = Hsa::agent_set_async_scratch_limit(bkendDevice_, limit); if (HSA_STATUS_SUCCESS != ret) { LogPrintfError("hsa_amd_agent_set_async_scratch_limit(%zu) failed with err 0x%xh", limit, ret); return false; @@ -2735,11 +2725,11 @@ bool Device::SvmAllocInit(void* memory, size_t size) const { if (info().hmmSupported_) { // Initialize signal for the barrier - hsa_signal_store_relaxed(prefetch_signal_, kInitSignalValueOne); + Hsa::signal_store_relaxed(prefetch_signal_, kInitSignalValueOne); // Initiate a prefetch command which should force memory update in HMM hsa_status_t status = - hsa_amd_svm_prefetch_async(memory, size, getBackendDevice(), 0, nullptr, prefetch_signal_); + Hsa::svm_prefetch_async(memory, size, getBackendDevice(), 0, nullptr, prefetch_signal_); if (status != HSA_STATUS_SUCCESS) { LogError("hsa_amd_svm_prefetch_async() failed"); return false; @@ -2812,7 +2802,7 @@ bool Device::IsHwEventReady(const amd::Event& event, bool wait, amd::SyncPolicy auto signal = reinterpret_cast(hw_event)->signal_; ClPrint(amd::LOG_INFO, amd::LOG_SIG, "Check HW event = 0x%lx", signal.handle); - return (hsa_signal_load_relaxed(signal) == 0); + return (Hsa::signal_load_relaxed(signal) == 0); } // ================================================================================================ @@ -2918,7 +2908,7 @@ hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue, // is no queue. uint32_t queue_max_packets = 0; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(bkendDevice_, HSA_AGENT_INFO_QUEUE_MAX_SIZE, &queue_max_packets)) { + Hsa::agent_get_info(bkendDevice_, HSA_AGENT_INFO_QUEUE_MAX_SIZE, &queue_max_packets)) { DevLogError("Cannot get hsa agent info \n"); return nullptr; } @@ -2932,9 +2922,9 @@ hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue, queue_type = HSA_QUEUE_TYPE_COOPERATIVE; } - while (hsa_queue_create(bkendDevice_, queue_size, queue_type, callbackQueue, this, - std::numeric_limits::max(), std::numeric_limits::max(), - &queue) != HSA_STATUS_SUCCESS) { + while (Hsa::queue_create(bkendDevice_, queue_size, queue_type, callbackQueue, this, + std::numeric_limits::max(), std::numeric_limits::max(), + &queue) != HSA_STATUS_SUCCESS) { queue_size >>= 1; if (queue_size < 64) { // if a queue with the same requested priority available from the pool, returns it here @@ -2948,10 +2938,10 @@ hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue, // default priority is normal so no need to set it again if (queue_priority != HSA_AMD_QUEUE_PRIORITY_NORMAL) { - hsa_status_t st = hsa_amd_queue_set_priority(queue, queue_priority); + hsa_status_t st = Hsa::queue_set_priority(queue, queue_priority); if (st != HSA_STATUS_SUCCESS) { DevLogError("Device::acquireQueue: hsa_amd_queue_set_priority failed!"); - hsa_queue_destroy(queue); + Hsa::queue_destroy(queue); return nullptr; } } @@ -2961,7 +2951,7 @@ hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue, "size %d with priority %d, cooperative: %i", queue, queue->base_address, queue_size, queue_priority, coop_queue); - hsa_amd_profiling_set_profiler_enabled(queue, 1); + Hsa::profiling_set_profiler_enabled(queue, 1); if (cuMask.size() != 0 || info_.globalCUMask_.size() != 0) { std::stringstream ss; ss << std::hex; @@ -3023,10 +3013,10 @@ hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue, } hsa_status_t status = - hsa_amd_queue_cu_set_mask(queue, final_mask.size() * 32, final_mask.data()); + Hsa::queue_cu_set_mask(queue, final_mask.size() * 32, final_mask.data()); if (status != HSA_STATUS_SUCCESS) { DevLogError("Device::acquireQueue: hsa_amd_queue_cu_set_mask failed!"); - hsa_queue_destroy(queue); + Hsa::queue_destroy(queue); return nullptr; } if (cuMask.size() != 0) { @@ -3098,14 +3088,14 @@ void Device::releaseQueue(hsa_queue_t* queue, const std::vector& cuMas ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Deleting hardware queue %p with refCount 0", queue->base_address); qIter = it.erase(qIter); - hsa_queue_destroy(queue); + Hsa::queue_destroy(queue); } } } if (coop_queue) { // cooperative queue ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Deleting CG enabled hardware queue %p ", queue->base_address); - hsa_queue_destroy(queue); + Hsa::queue_destroy(queue); } } @@ -3173,7 +3163,7 @@ bool Device::findLinkInfo(const hsa_amd_memory_pool_t& pool, // Retrieve the hops between 2 devices. int32_t hops = 0; - hsa_status_t hsa_status = hsa_amd_agent_memory_pool_get_info( + hsa_status_t hsa_status = Hsa::agent_memory_pool_get_info( bkendDevice_, pool, HSA_AMD_AGENT_MEMORY_POOL_INFO_NUM_LINK_HOPS, &hops); if (hsa_status != HSA_STATUS_SUCCESS) { @@ -3221,7 +3211,7 @@ bool Device::findLinkInfo(const hsa_amd_memory_pool_t& pool, // Retrieve link info on the pool. std::vector link_info(hops); - hsa_status = hsa_amd_agent_memory_pool_get_info( + hsa_status = Hsa::agent_memory_pool_get_info( bkendDevice_, pool, HSA_AMD_AGENT_MEMORY_POOL_INFO_LINK_INFO, link_info.data()); if (hsa_status != HSA_STATUS_SUCCESS) { @@ -3361,7 +3351,7 @@ hsa_status_t Device::BackendErrorCallBackHandler(const hsa_amd_event_t* event, v void Device::RegisterBackendErrorCb() { // Register ROCclr Error Callback hsa_status_t hsa_error = HSA_STATUS_SUCCESS; - hsa_error = hsa_amd_register_system_event_handler(BackendErrorCallBackHandler, nullptr); + hsa_error = Hsa::register_system_event_handler(BackendErrorCallBackHandler, nullptr); if (hsa_error != HSA_STATUS_SUCCESS) { LogError("Cannot Register Call back event handler"); } @@ -3409,10 +3399,10 @@ void Device::ReleaseGlobalSignal(void* signal) const { bool Device::CreateUserEvent(amd::UserEvent* event) const { std::unique_ptr signal(new ProfilingSignal()); if ((signal == nullptr) || - (HSA_STATUS_SUCCESS != hsa_signal_create(0, 0, nullptr, &signal->signal_))) { + (HSA_STATUS_SUCCESS != Hsa::signal_create(0, 0, nullptr, &signal->signal_))) { return false; } - hsa_signal_silent_store_relaxed(signal->signal_, kInitSignalValueOne); + Hsa::signal_silent_store_relaxed(signal->signal_, kInitSignalValueOne); event->SetHwEvent(signal.release()); return true; } @@ -3421,14 +3411,14 @@ bool Device::CreateUserEvent(amd::UserEvent* event) const { void Device::SetUserEvent(amd::UserEvent* event) const { auto signal = reinterpret_cast(event->HwEvent()); assert(signal != nullptr && "Can't have user event without hw event!"); - hsa_signal_silent_store_relaxed(signal->signal_, 0); + Hsa::signal_silent_store_relaxed(signal->signal_, 0); } // ================================================================================================ bool Device::IsValidAllocation(const void* dev_ptr, size_t size, hsa_amd_pointer_info_t* ptr_info) { // Query ptr type to see if it's a HMM allocation hsa_status_t status = - hsa_amd_pointer_info(const_cast(dev_ptr), ptr_info, nullptr, nullptr, nullptr); + Hsa::pointer_info(const_cast(dev_ptr), ptr_info, nullptr, nullptr, nullptr); // The call should never fail in ROCR, but just check for an error and continue if (status != HSA_STATUS_SUCCESS) { LogError("hsa_amd_pointer_info() failed"); @@ -3513,14 +3503,14 @@ void Device::RemoveKernel(Kernel& gpuKernel) const { // ================================================================================================ ProfilingSignal::~ProfilingSignal() { if (signal_.handle != 0) { - if (hsa_signal_load_relaxed(signal_) > 0 + if (Hsa::signal_load_relaxed(signal_) > 0 && !(HIP_SKIP_ABORT_ON_GPU_ERROR && amd::Device::IsGPUInError())) { LogError("Runtime shouldn't destroy a signal that is still busy!"); - if (hsa_signal_wait_scacquire(signal_, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne, + if (Hsa::signal_wait_scacquire(signal_, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne, kUnlimitedWait, HSA_WAIT_STATE_BLOCKED) != 0) { } } - hsa_signal_destroy(signal_); + Hsa::signal_destroy(signal_); } } @@ -3578,11 +3568,11 @@ void callbackQueue(hsa_status_t status, hsa_queue_t* queue, void* data) { } // Abort on device exceptions. const char* errorMsg = 0; - hsa_status_string(status, &errorMsg); + Hsa::status_string(status, &errorMsg); if (status == HSA_STATUS_ERROR_OUT_OF_RESOURCES) { size_t global_available_mem = 0; if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(dev->getBackendDevice(), + Hsa::agent_get_info(dev->getBackendDevice(), static_cast(HSA_AMD_AGENT_INFO_MEMORY_AVAIL), &global_available_mem)) { LogError("HSA_AMD_AGENT_INFO_MEMORY_AVAIL query failed."); diff --git a/projects/clr/rocclr/device/rocm/rocdevice.hpp b/projects/clr/rocclr/device/rocm/rocdevice.hpp index 414272da1a..944505f658 100644 --- a/projects/clr/rocclr/device/rocm/rocdevice.hpp +++ b/projects/clr/rocclr/device/rocm/rocdevice.hpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2009 - 2023 Advanced Micro Devices, Inc. +/* Copyright (c) 2009 - 2025 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -32,17 +32,13 @@ #include "thread/monitor.hpp" #include "utils/versions.hpp" +#include "device/rocm/rocrctx.hpp" #include "device/rocm/rocsettings.hpp" #include "device/rocm/rocvirtual.hpp" #include "device/rocm/rocdefs.hpp" #include "device/rocm/rocprintf.hpp" #include "device/rocm/rocglinterop.hpp" -#include "hsa/hsa.h" -#include "hsa/hsa_ext_image.h" -#include "hsa/hsa_ext_amd.h" -#include "hsa/hsa_ven_amd_loader.h" - #include #include #include @@ -305,12 +301,6 @@ class NullDevice : public amd::Device { #endif #endif - protected: - //! Initialize compiler instance and handle - static bool initCompiler(bool isOffline); - //! destroy compiler instance and handle - static bool destroyCompiler(); - private: static constexpr bool offlineDevice_ = true; }; diff --git a/projects/clr/rocclr/device/rocm/rocglinterop.hpp b/projects/clr/rocclr/device/rocm/rocglinterop.hpp index 478e45cd10..c19dd1be8b 100644 --- a/projects/clr/rocclr/device/rocm/rocglinterop.hpp +++ b/projects/clr/rocclr/device/rocm/rocglinterop.hpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2016 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2016 - 2025 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -36,9 +36,9 @@ typedef _XDisplay Display; typedef __GLXcontextRec* GLXContext; #endif +#include "device/rocm/rocdevice.hpp" #include "device/rocm/mesa_glinterop.h" #include "device/rocm/rocregisters.hpp" -#include "hsa/hsa_ext_amd.h" namespace amd::roc { diff --git a/projects/clr/rocclr/device/rocm/rockernel.cpp b/projects/clr/rocclr/device/rocm/rockernel.cpp index 40c0490b53..91f46b59b9 100644 --- a/projects/clr/rocclr/device/rocm/rockernel.cpp +++ b/projects/clr/rocclr/device/rocm/rockernel.cpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2009 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2009 - 2025 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -19,7 +19,6 @@ THE SOFTWARE. */ #include "rockernel.hpp" -#include "hsa/amd_hsa_kernel_code.h" #include @@ -46,7 +45,7 @@ bool Kernel::postLoad() { hsa_status_t hsaStatus; hsa_executable_symbol_t symbol; hsa_agent_t agent = program()->rocDevice().getBackendDevice(); - hsaStatus = hsa_executable_get_symbol_by_name(program()->hsaExecutable(), symbolName().c_str(), + hsaStatus = Hsa::executable_get_symbol_by_name(program()->hsaExecutable(), symbolName().c_str(), &agent, &symbol); if (hsaStatus != HSA_STATUS_SUCCESS) { DevLogPrintfError("Cannot Get Symbol : %s, failed with hsa_status: %d \n", symbolName().c_str(), @@ -54,7 +53,7 @@ bool Kernel::postLoad() { return false; } - hsaStatus = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, + hsaStatus = Hsa::executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kernelCodeHandle_); if (hsaStatus != HSA_STATUS_SUCCESS) { DevLogPrintfError(" Cannot Get Symbol Info: %s, failed with hsa_status: %d \n ", @@ -62,7 +61,7 @@ bool Kernel::postLoad() { return false; } - hsaStatus = hsa_executable_symbol_get_info( + hsaStatus = Hsa::executable_symbol_get_info( symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_DYNAMIC_CALLSTACK, &kernelHasDynamicCallStack_); if (hsaStatus != HSA_STATUS_SUCCESS) { DevLogPrintfError(" Cannot Get Dynamic callstack info, failed with hsa_status: %d \n ", @@ -80,7 +79,7 @@ bool Kernel::postLoad() { // retrieve the kernel code object handle of such a kernel. The address of the variable and the // kernel code object handle are known only after the hsa executable is loaded. The below code // copies the kernel code object handle to the address of the variable. - hsaStatus = hsa_executable_get_symbol_by_name(program()->hsaExecutable(), + hsaStatus = Hsa::executable_get_symbol_by_name(program()->hsaExecutable(), RuntimeHandle().c_str(), &agent, &kernelSymbol); if (hsaStatus != HSA_STATUS_SUCCESS) { DevLogPrintfError("Cannot get Kernel Symbol by name: %s, failed with hsa_status: %d \n", @@ -88,7 +87,7 @@ bool Kernel::postLoad() { return false; } - hsaStatus = hsa_executable_symbol_get_info( + hsaStatus = Hsa::executable_symbol_get_info( kernelSymbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE, &variable_size); if (hsaStatus != HSA_STATUS_SUCCESS) { DevLogPrintfError( @@ -96,7 +95,7 @@ bool Kernel::postLoad() { return false; } - hsaStatus = hsa_executable_symbol_get_info( + hsaStatus = Hsa::executable_symbol_get_info( kernelSymbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS, &variable_address); if (hsaStatus != HSA_STATUS_SUCCESS) { DevLogPrintfError("[ROC][Kernel] Cannot get Kernel Address, failed with hsa_status: %d \n", @@ -107,7 +106,7 @@ bool Kernel::postLoad() { const struct RuntimeHandle runtime_handle = { kernelCodeHandle_, WorkitemPrivateSegmentByteSize(), WorkgroupGroupSegmentByteSize()}; hsaStatus = - hsa_memory_copy(reinterpret_cast(variable_address), &runtime_handle, variable_size); + Hsa::memory_copy(reinterpret_cast(variable_address), &runtime_handle, variable_size); if (hsaStatus != HSA_STATUS_SUCCESS) { DevLogPrintfError("[ROC][Kernel] HSA Memory copy failed, failed with hsa_status: %d \n", @@ -121,8 +120,8 @@ bool Kernel::postLoad() { // We set the value to HSA if the value is uninitialized uint32_t wavefront_size = workGroupInfo_.wavefrontPerSIMD_; if (wavefront_size == 0 && - hsa_agent_get_info(program()->rocDevice().getBackendDevice(), HSA_AGENT_INFO_WAVEFRONT_SIZE, - &wavefront_size) != HSA_STATUS_SUCCESS) { + Hsa::agent_get_info(program()->rocDevice().getBackendDevice(), HSA_AGENT_INFO_WAVEFRONT_SIZE, + &wavefront_size) != HSA_STATUS_SUCCESS) { DevLogPrintfError("[ROC][Kernel] Cannot get Wavefront Size, failed with hsa_status: %d \n", hsaStatus); return false; diff --git a/projects/clr/rocclr/device/rocm/rocmemory.cpp b/projects/clr/rocclr/device/rocm/rocmemory.cpp index 136b8f69b0..743f33173f 100644 --- a/projects/clr/rocclr/device/rocm/rocmemory.cpp +++ b/projects/clr/rocclr/device/rocm/rocmemory.cpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2008 - 2023 Advanced Micro Devices, Inc. +/* Copyright (c) 2008 - 2025 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -212,8 +212,8 @@ hsa_status_t Memory::interopMapBuffer(amd::Os::FileDesc fdn) { #else auto fd = fdn; #endif - hsa_status_t status = hsa_amd_interop_map_buffer(1, &agent, fd, 0, &size, &interop_deviceMemory_, - &metadata_size, (const void**)&metadata); + hsa_status_t status = Hsa::interop_map_buffer(1, &agent, fd, 0, &size, &interop_deviceMemory_, + &metadata_size, (const void**)&metadata); ClPrint(amd::LOG_DEBUG, amd::LOG_MEM, "Map Interop memory %p, size 0x%zx", interop_deviceMemory_, size); deviceMemory_ = static_cast(interop_deviceMemory_); // + out.buf_offset; @@ -254,7 +254,7 @@ bool Memory::createInteropBuffer(GLenum targetType, int miplevel) { hsa_agent_t agent = dev().getBackendDevice(); uint32_t id; - hsa_agent_get_info(agent, static_cast(HSA_AMD_AGENT_INFO_CHIP_ID), &id); + Hsa::agent_get_info(agent, static_cast(HSA_AMD_AGENT_INFO_CHIP_ID), &id); static constexpr int MaxMetadataSizeDwords = 64; static constexpr int MaxMetadataSizeBytes = MaxMetadataSizeDwords * sizeof(int); @@ -293,7 +293,7 @@ bool Memory::createInteropBuffer(GLenum targetType, int miplevel) { void Memory::destroyInteropBuffer() { assert(kind_ == MEMORY_KIND_INTEROP && "Memory must be interop type."); - hsa_amd_interop_unmap_buffer(interop_deviceMemory_); + Hsa::interop_unmap_buffer(interop_deviceMemory_); ClPrint(amd::LOG_DEBUG, amd::LOG_MEM, "Unmap GL memory %p", deviceMemory_); deviceMemory_ = nullptr; } @@ -622,7 +622,7 @@ Buffer::~Buffer() { if (owner()->ipcShared()) { // Detach the memory from HSA - auto hsa_status = hsa_amd_ipc_memory_detach(owner()->getSvmPtr()); + auto hsa_status = Hsa::ipc_memory_detach(owner()->getSvmPtr()); if (hsa_status != HSA_STATUS_SUCCESS) { LogPrintfError("HSA failed to detach memory with status: %d", hsa_status); } @@ -665,9 +665,8 @@ void Buffer::destroy() { dev().memFree(deviceMemory_, size()); } } else if (memFlags & ROCCLR_MEM_HSA_SIGNAL_MEMORY) { - if (HSA_STATUS_SUCCESS != hsa_signal_destroy(signal_)) { - ClPrint(amd::LOG_ERROR, amd::LOG_MEM, - "hsa_signal_destroy failed"); + if (HSA_STATUS_SUCCESS != Hsa::signal_destroy(signal_)) { + ClPrint(amd::LOG_ERROR, amd::LOG_MEM, "hsa_signal_destroy failed"); } deviceMemory_ = nullptr; } else { @@ -680,7 +679,7 @@ void Buffer::destroy() { if (memFlags & CL_MEM_USE_HOST_PTR) { // unlock svm host pointer from memory pool if (!dev().info().hmmSupported_) { - hsa_amd_memory_unlock(owner()->getSvmPtr()); + Hsa::memory_unlock(owner()->getSvmPtr()); } // destroy system memory if (!(amd::Os::releaseMemory(deviceMemory_, size()))) { @@ -723,14 +722,14 @@ void Buffer::destroy() { if (needUnlockHostMem) { if (memFlags & (CL_MEM_USE_HOST_PTR | CL_MEM_ALLOC_HOST_PTR)) { - if (dev().agent_profile() != HSA_PROFILE_FULL) hsa_amd_memory_unlock(owner()->getHostMem()); + if (dev().agent_profile() != HSA_PROFILE_FULL) Hsa::memory_unlock(owner()->getHostMem()); } } } if (memFlags & CL_MEM_USE_HOST_PTR) { if (dev().agent_profile() == HSA_PROFILE_FULL) { - hsa_memory_deregister(owner()->getHostMem(), size()); + Hsa::memory_deregister(owner()->getHostMem(), size()); } } } @@ -759,7 +758,7 @@ bool Buffer::create(bool alloc_local) { // Extra 1 for the current device const uint32_t ipc_agents_num = dev().p2pAgents().size() + 1; // Retrieve the devPtr from the handle - auto hsa_status = hsa_amd_ipc_memory_attach( + auto hsa_status = Hsa::ipc_memory_attach( reinterpret_cast( reinterpret_cast(owner())->Handle()), owner()->getSize(), ipc_agents_num, dev().IpcAgents(), &orig_dev_ptr); @@ -832,14 +831,14 @@ bool Buffer::create(bool alloc_local) { } else if (memFlags & ROCCLR_MEM_HSA_SIGNAL_MEMORY) { // TODO: ROCr will introduce a new attribute enum that implies a non-blocking signal, // replace "HSA_AMD_SIGNAL_AMD_GPU_ONLY" with this new enum when it is ready. - if (HSA_STATUS_SUCCESS != hsa_amd_signal_create(kInitSignalValueOne, 0, nullptr, - HSA_AMD_SIGNAL_AMD_GPU_ONLY, &signal_)) { + if (HSA_STATUS_SUCCESS != Hsa::signal_create(kInitSignalValueOne, 0, nullptr, + HSA_AMD_SIGNAL_AMD_GPU_ONLY, &signal_)) { ClPrint(amd::LOG_ERROR, amd::LOG_MEM, "hsa_amd_signal_create signal failed"); return false; } volatile hsa_signal_value_t* signalValuePtr = nullptr; - if (HSA_STATUS_SUCCESS != hsa_amd_signal_value_pointer(signal_, &signalValuePtr)) { + if (HSA_STATUS_SUCCESS != Hsa::signal_value_pointer(signal_, &signalValuePtr)) { ClPrint(amd::LOG_ERROR, amd::LOG_MEM, "hsa_amd_signal_value_pointer failed"); return false; @@ -1005,7 +1004,7 @@ bool Buffer::create(bool alloc_local) { deviceMemory_ = owner()->getHostMem(); if (memFlags & CL_MEM_USE_HOST_PTR) { - hsa_memory_register(deviceMemory_, size()); + Hsa::memory_register(deviceMemory_, size()); } return deviceMemory_ != nullptr; @@ -1042,7 +1041,7 @@ bool Buffer::ExportHandle(void* handle) const { orig_dev_ptr = owner()->getHostMem(); } - auto hsa_status = hsa_amd_ipc_memory_create(orig_dev_ptr, owner()->getSize(), + auto hsa_status = Hsa::ipc_memory_create(orig_dev_ptr, owner()->getSize(), reinterpret_cast(handle)); if (hsa_status != HSA_STATUS_SUCCESS) { LogPrintfError("Failed to create memory for IPC, failed with hsa_status: %d", hsa_status); @@ -1061,7 +1060,7 @@ bool Buffer::GetFDHandleForMem(void* dev_ptr, size_t size, bool vmm, void* handl hsa_amd_vmem_alloc_handle_t mem_handle; // Retrieve the corresponding phys_mem handle for the mapped dev_ptr. - hsa_status_t hsa_status = hsa_amd_vmem_retain_alloc_handle(&mem_handle, dev_ptr); + hsa_status_t hsa_status = Hsa::vmem_retain_alloc_handle(&mem_handle, dev_ptr); if (hsa_status != HSA_STATUS_SUCCESS) { LogPrintfError("Cannot retain alloc handle for dev_ptr: 0x%x hsa returned status: %d", dev_ptr, hsa_status); @@ -1069,7 +1068,7 @@ bool Buffer::GetFDHandleForMem(void* dev_ptr, size_t size, bool vmm, void* handl } // Now, retrieve the shareable handle (fd in linux) for the phys_mem handle. - hsa_status = hsa_amd_vmem_export_shareable_handle(&dmabuffd, mem_handle, 0); + hsa_status = Hsa::vmem_export_shareable_handle(&dmabuffd, mem_handle, 0); if (hsa_status != HSA_STATUS_SUCCESS) { LogPrintfError("Cannot get shareable handle for mem_handle: %lu, hsa returned status: %d", mem_handle, hsa_status); @@ -1077,7 +1076,7 @@ bool Buffer::GetFDHandleForMem(void* dev_ptr, size_t size, bool vmm, void* handl } } else { // Retrieve a shareable handle for the device ptr. - hsa_status_t hsa_status = hsa_amd_portable_export_dmabuf(dev_ptr, size, &dmabuffd, &offset); + hsa_status_t hsa_status = Hsa::portable_export_dmabuf(dev_ptr, size, &dmabuffd, &offset); if (hsa_status != HSA_STATUS_SUCCESS) { LogPrintfError( "Cannot export a portable fd for dev_ptr: 0x%x with size: %lu," @@ -1244,8 +1243,8 @@ bool Image::createInteropImage() { originalDeviceMemory_ = deviceMemory_; if (obj->getGLTarget() == GL_TEXTURE_BUFFER) { - hsa_status_t err = hsa_ext_image_create(dev().getBackendDevice(), &imageDescriptor_, - originalDeviceMemory_, permission_, &hsaImageObject_); + hsa_status_t err = Hsa::image_create(dev().getBackendDevice(), &imageDescriptor_, + originalDeviceMemory_, permission_, &hsaImageObject_); return (err == HSA_STATUS_SUCCESS); } @@ -1263,8 +1262,8 @@ bool Image::createInteropImage() { } hsa_status_t err = - hsa_amd_image_create(dev().getBackendDevice(), &imageDescriptor_, amdImageDesc_, - originalDeviceMemory_, permission_, &hsaImageObject_); + Hsa::image_create(dev().getBackendDevice(), &imageDescriptor_, amdImageDesc_, + originalDeviceMemory_, permission_, &hsaImageObject_); if (err != HSA_STATUS_SUCCESS) return false; return true; @@ -1302,8 +1301,8 @@ bool Image::create(bool alloc_local) { } // Get memory size requirement for device specific image. - hsa_status_t status = hsa_ext_image_data_get_info(dev().getBackendDevice(), &imageDescriptor_, - permission_, &deviceImageInfo_); + hsa_status_t status = Hsa::image_data_get_info(dev().getBackendDevice(), &imageDescriptor_, + permission_, &deviceImageInfo_); if (status != HSA_STATUS_SUCCESS) { LogPrintfError("Fail to allocate image memory, failed with hsa_status: %d", status); @@ -1340,8 +1339,8 @@ bool Image::create(bool alloc_local) { assert(amd::isMultipleOf(deviceMemory_, static_cast(deviceImageInfo_.alignment))); - status = hsa_ext_image_create(dev().getBackendDevice(), &imageDescriptor_, deviceMemory_, - permission_, &hsaImageObject_); + status = Hsa::image_create(dev().getBackendDevice(), &imageDescriptor_, deviceMemory_, + permission_, &hsaImageObject_); if (status != HSA_STATUS_SUCCESS) { LogPrintfError("[OCL] Fail to allocate image memory, failed with hsa_status: %d \n", status); @@ -1388,7 +1387,7 @@ bool Image::createView(const Memory& parent) { rowPitch = elementSize * amd::alignUp(rowPitch, (dev().info().imagePitchAlignment_ / elementSize)); - status = hsa_ext_image_create_with_layout( + status = Hsa::image_create_with_layout( dev().getBackendDevice(), &imageDescriptor_, deviceMemory_, permission_, HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, rowPitch, 0, &hsaImageObject_); @@ -1409,7 +1408,7 @@ bool Image::createView(const Memory& parent) { break; } hsa_ext_image_t hsaImage; - if (HSA_STATUS_SUCCESS == hsa_ext_image_create_with_layout( + if (HSA_STATUS_SUCCESS == Hsa::image_create_with_layout( dev().getBackendDevice(), &imageDescriptor_, deviceMemory_, permission_, HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, tryPitch, 0, &hsaImage)) { @@ -1417,8 +1416,8 @@ bool Image::createView(const Memory& parent) { LogWarning("[OCL] will use copy image"); workaround = true; // Free the image. - hsa_ext_image_destroy(dev().getBackendDevice(), hsaImage); - hsa_ext_image_destroy(dev().getBackendDevice(), hsaImageObject_); + Hsa::image_destroy(dev().getBackendDevice(), hsaImage); + Hsa::image_destroy(dev().getBackendDevice(), hsaImageObject_); hsaImageObject_.handle = 0; break; } @@ -1436,11 +1435,11 @@ bool Image::createView(const Memory& parent) { } } else if (kind_ == MEMORY_KIND_INTEROP) { amdImageDesc_ = static_cast(parent.owner()->getDeviceMemory(dev()))->amdImageDesc_; - status = hsa_amd_image_create(dev().getBackendDevice(), &imageDescriptor_, amdImageDesc_, - deviceMemory_, permission_, &hsaImageObject_); + status = Hsa::image_create(dev().getBackendDevice(), &imageDescriptor_, amdImageDesc_, + deviceMemory_, permission_, &hsaImageObject_); } else { - status = hsa_ext_image_create(dev().getBackendDevice(), &imageDescriptor_, deviceMemory_, - permission_, &hsaImageObject_); + status = Hsa::image_create(dev().getBackendDevice(), &imageDescriptor_, deviceMemory_, + permission_, &hsaImageObject_); } if (status != HSA_STATUS_SUCCESS) { @@ -1537,7 +1536,7 @@ void Image::destroy() { delete copyImageBuffer_; if (hsaImageObject_.handle != 0 && ownsHsaImageObject_) { - hsa_status_t status = hsa_ext_image_destroy(dev().getBackendDevice(), hsaImageObject_); + hsa_status_t status = Hsa::image_destroy(dev().getBackendDevice(), hsaImageObject_); assert(status == HSA_STATUS_SUCCESS); } // Don't destroy memory if it's a view. Parent will destroy the original allocation. diff --git a/projects/clr/rocclr/device/rocm/rocprintf.cpp b/projects/clr/rocclr/device/rocm/rocprintf.cpp index 4ed427f3a0..cef7ab5b56 100644 --- a/projects/clr/rocclr/device/rocm/rocprintf.cpp +++ b/projects/clr/rocclr/device/rocm/rocprintf.cpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2010 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2010 - 2025 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -405,7 +405,7 @@ bool PrintfDbg::init(bool printfEnabled) { // Copy offset and number of bytes available for printf data // into the corresponding location in the debug buffer - hsa_status_t err = hsa_memory_copy(dbgBuffer_, sysMem, 2 * sizeof(uint32_t)); + hsa_status_t err = Hsa::memory_copy(dbgBuffer_, sysMem, 2 * sizeof(uint32_t)); if (err != HSA_STATUS_SUCCESS) { LogPrintfError( "\n Can't copy offset and bytes available data to dgbBuffer_," diff --git a/projects/clr/rocclr/device/rocm/rocprogram.cpp b/projects/clr/rocclr/device/rocm/rocprogram.cpp index c0e0a952ce..63502560fc 100644 --- a/projects/clr/rocclr/device/rocm/rocprogram.cpp +++ b/projects/clr/rocclr/device/rocm/rocprogram.cpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2008 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2008 - 2025 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -23,8 +23,6 @@ #include "utils/options.hpp" #include "rockernel.hpp" -#include "hsa/amd_hsa_kernel_code.h" - #include #include #include @@ -37,7 +35,7 @@ namespace amd::roc { static inline const char* hsa_strerror(hsa_status_t status) { const char* str = nullptr; - if (hsa_status_string(status, &str) == HSA_STATUS_SUCCESS) { + if (Hsa::status_string(status, &str) == HSA_STATUS_SUCCESS) { return str; } return "Unknown error"; @@ -46,10 +44,10 @@ static inline const char* hsa_strerror(hsa_status_t status) { Program::~Program() { // Destroy the executable. if (hsaExecutable_.handle != 0) { - hsa_executable_destroy(hsaExecutable_); + Hsa::executable_destroy(hsaExecutable_); } if (hsaCodeObjectReader_.handle != 0) { - hsa_code_object_reader_destroy(hsaCodeObjectReader_); + Hsa::code_object_reader_destroy(hsaCodeObjectReader_); } releaseClBinary(); } @@ -102,7 +100,7 @@ bool Program::defineGlobalVar(const char* name, void* dptr) { hsa_agent_t hsa_device = rocDevice().getBackendDevice(); hsa_status_t status = - hsa_executable_agent_global_variable_define(hsaExecutable_, hsa_device, name, dptr); + Hsa::executable_agent_global_variable_define(hsaExecutable_, hsa_device, name, dptr); if (status != HSA_STATUS_SUCCESS) { buildLog_ += "Error: Could not define global variable : "; buildLog_ += hsa_strerror(status); @@ -134,7 +132,7 @@ bool Program::createGlobalVarObj(amd::Memory** amd_mem_obj, void** device_pptr, /* Find HSA Symbol by name */ status = - hsa_executable_get_symbol_by_name(hsaExecutable_, global_name, &hsa_device, &global_symbol); + Hsa::executable_get_symbol_by_name(hsaExecutable_, global_name, &hsa_device, &global_symbol); if (status != HSA_STATUS_SUCCESS) { buildLog_ += "Error: Failed to find the Symbol by Name: "; buildLog_ += hsa_strerror(status); @@ -144,7 +142,7 @@ bool Program::createGlobalVarObj(amd::Memory** amd_mem_obj, void** device_pptr, /* Find HSA Symbol Type */ status = - hsa_executable_symbol_get_info(global_symbol, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &sym_type); + Hsa::executable_symbol_get_info(global_symbol, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &sym_type); if (status != HSA_STATUS_SUCCESS) { buildLog_ += "Error: Failed to find the Symbol Type : "; buildLog_ += hsa_strerror(status); @@ -161,7 +159,7 @@ bool Program::createGlobalVarObj(amd::Memory** amd_mem_obj, void** device_pptr, } /* Retrieve the size of the variable */ - status = hsa_executable_symbol_get_info(global_symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE, + status = Hsa::executable_symbol_get_info(global_symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE, bytes); if (status != HSA_STATUS_SUCCESS) { @@ -174,7 +172,7 @@ bool Program::createGlobalVarObj(amd::Memory** amd_mem_obj, void** device_pptr, // Handle size 0 symbols if (*bytes != 0) { // Find HSA Symbol Address - status = hsa_executable_symbol_get_info( + status = Hsa::executable_symbol_get_info( global_symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS, device_pptr); if (status != HSA_STATUS_SUCCESS) { buildLog_ += "Error: Failed to find the Symbol Address : "; @@ -288,7 +286,7 @@ bool LightningProgram::setKernels(void* binary, size_t binSize, amd::Os::FileDes hsa_agent_t agent = rocDevice().getBackendDevice(); hsa_status_t status; - status = hsa_executable_create_alt(HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, + status = Hsa::executable_create_alt(HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, nullptr, &hsaExecutable_); if (status != HSA_STATUS_SUCCESS) { buildLog_ += "Error: Executable for AMD HSA Code Object isn't created: "; @@ -300,7 +298,7 @@ bool LightningProgram::setKernels(void* binary, size_t binSize, amd::Os::FileDes // Load the code object, either with file descriptor and offset // or binary image and binary size with URI // or binary image and binary size - status = hsa_code_object_reader_create_from_memory(binary, binSize, &hsaCodeObjectReader_); + status = Hsa::code_object_reader_create_from_memory(binary, binSize, &hsaCodeObjectReader_); if (status != HSA_STATUS_SUCCESS) { buildLog_ += "Error: AMD HSA Code Object Reader create failed: "; buildLog_ += hsa_strerror(status); @@ -308,7 +306,7 @@ bool LightningProgram::setKernels(void* binary, size_t binSize, amd::Os::FileDes return false; } - status = hsa_executable_load_agent_code_object(hsaExecutable_, agent, hsaCodeObjectReader_, + status = Hsa::executable_load_agent_code_object(hsaExecutable_, agent, hsaCodeObjectReader_, nullptr, nullptr); if (status != HSA_STATUS_SUCCESS) { buildLog_ += "Error: AMD HSA Code Object loading failed: "; @@ -318,7 +316,7 @@ bool LightningProgram::setKernels(void* binary, size_t binSize, amd::Os::FileDes } // Freeze the executable. - status = hsa_executable_freeze(hsaExecutable_, nullptr); + status = Hsa::executable_freeze(hsaExecutable_, nullptr); if (status != HSA_STATUS_SUCCESS) { buildLog_ += "Error: Freezing the executable failed: "; buildLog_ += hsa_strerror(status); diff --git a/projects/clr/rocclr/device/rocm/rocrctx.cpp b/projects/clr/rocclr/device/rocm/rocrctx.cpp new file mode 100644 index 0000000000..63308eed8a --- /dev/null +++ b/projects/clr/rocclr/device/rocm/rocrctx.cpp @@ -0,0 +1,143 @@ +/* Copyright (c) 2025 Advanced Micro Devices, Inc. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in 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: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + 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 + AUTHORS 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 IN + THE SOFTWARE. */ + +#include "os/os.hpp" +#include "utils/flags.hpp" +#include "rocrctx.hpp" + +namespace amd { +namespace roc { + +std::once_flag Hsa::initialized; +RocrEntryPoints Hsa::cep_; +bool Hsa::is_ready_ = false; + +bool Hsa::LoadLib() { +#if defined(ROCR_DYN_DLL) + static const char* rocr_lib_name = WINDOWS_SWITCH("hsa-runtime64.dll", "hsa-runtime64.so.1"); + cep_.handle = Os::loadLibrary(rocr_lib_name); + if (nullptr == cep_.handle) { + ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "Failed to load COMGR library."); + return false; + } +#endif + GET_ROCR_SYMBOL(hsa_init) + GET_ROCR_SYMBOL(hsa_shut_down) + GET_ROCR_SYMBOL(hsa_system_get_info) + GET_ROCR_SYMBOL(hsa_iterate_agents) + GET_ROCR_SYMBOL(hsa_agent_get_info) + GET_ROCR_SYMBOL(hsa_queue_create) + GET_ROCR_SYMBOL(hsa_queue_destroy) + GET_ROCR_SYMBOL(hsa_queue_load_read_index_scacquire) + GET_ROCR_SYMBOL(hsa_queue_load_read_index_relaxed) + GET_ROCR_SYMBOL(hsa_queue_load_write_index_relaxed) + GET_ROCR_SYMBOL(hsa_queue_add_write_index_screlease) + GET_ROCR_SYMBOL(hsa_memory_register) + GET_ROCR_SYMBOL(hsa_memory_deregister) + GET_ROCR_SYMBOL(hsa_memory_copy) + GET_ROCR_SYMBOL(hsa_signal_create) + GET_ROCR_SYMBOL(hsa_signal_destroy) + GET_ROCR_SYMBOL(hsa_signal_load_relaxed) + GET_ROCR_SYMBOL(hsa_signal_store_relaxed) + GET_ROCR_SYMBOL(hsa_signal_silent_store_relaxed) + GET_ROCR_SYMBOL(hsa_signal_store_screlease) + GET_ROCR_SYMBOL(hsa_signal_wait_scacquire) + GET_ROCR_SYMBOL(hsa_signal_add_relaxed) + GET_ROCR_SYMBOL(hsa_signal_subtract_relaxed) + GET_ROCR_SYMBOL(hsa_isa_get_info_alt) + GET_ROCR_SYMBOL(hsa_agent_iterate_isas) + GET_ROCR_SYMBOL(hsa_system_get_major_extension_table) + GET_ROCR_SYMBOL(hsa_status_string) + GET_ROCR_SYMBOL(hsa_executable_create_alt) + GET_ROCR_SYMBOL(hsa_executable_destroy) + GET_ROCR_SYMBOL(hsa_executable_get_info) + GET_ROCR_SYMBOL(hsa_code_object_reader_destroy) + GET_ROCR_SYMBOL(hsa_code_object_reader_create_from_memory) + GET_ROCR_SYMBOL(hsa_executable_load_agent_code_object) + GET_ROCR_SYMBOL(hsa_executable_agent_global_variable_define) + GET_ROCR_SYMBOL(hsa_executable_get_symbol_by_name) + GET_ROCR_SYMBOL(hsa_executable_symbol_get_info) + GET_ROCR_SYMBOL(hsa_executable_freeze) + // AMD extensions + GET_ROCR_SYMBOL(hsa_amd_coherency_set_type) + GET_ROCR_SYMBOL(hsa_amd_profiling_set_profiler_enabled) + GET_ROCR_SYMBOL(hsa_amd_profiling_async_copy_enable) + GET_ROCR_SYMBOL(hsa_amd_profiling_get_dispatch_time) + GET_ROCR_SYMBOL(hsa_amd_profiling_get_async_copy_time) + GET_ROCR_SYMBOL(hsa_amd_signal_async_handler) + GET_ROCR_SYMBOL(hsa_amd_queue_cu_set_mask) + GET_ROCR_SYMBOL(hsa_amd_memory_pool_get_info) + GET_ROCR_SYMBOL(hsa_amd_agent_iterate_memory_pools) + GET_ROCR_SYMBOL(hsa_amd_memory_pool_allocate) + GET_ROCR_SYMBOL(hsa_amd_memory_pool_free) + GET_ROCR_SYMBOL(hsa_amd_memory_async_copy) + GET_ROCR_SYMBOL(hsa_amd_memory_async_copy_on_engine) + GET_ROCR_SYMBOL(hsa_amd_memory_copy_engine_status) + GET_ROCR_SYMBOL(hsa_amd_agent_memory_pool_get_info) + GET_ROCR_SYMBOL(hsa_amd_agents_allow_access) + GET_ROCR_SYMBOL(hsa_amd_memory_unlock) + GET_ROCR_SYMBOL(hsa_amd_interop_map_buffer) + GET_ROCR_SYMBOL(hsa_amd_interop_unmap_buffer) + GET_ROCR_SYMBOL(hsa_amd_image_create) + GET_ROCR_SYMBOL(hsa_amd_pointer_info) + GET_ROCR_SYMBOL(hsa_amd_ipc_memory_create) + GET_ROCR_SYMBOL(hsa_amd_ipc_memory_attach) + GET_ROCR_SYMBOL(hsa_amd_ipc_memory_detach) + GET_ROCR_SYMBOL(hsa_amd_signal_create) + GET_ROCR_SYMBOL(hsa_amd_register_system_event_handler) + GET_ROCR_SYMBOL(hsa_amd_queue_set_priority) + GET_ROCR_SYMBOL(hsa_amd_memory_async_copy_rect) + GET_ROCR_SYMBOL(hsa_amd_memory_lock_to_pool) + GET_ROCR_SYMBOL(hsa_amd_signal_value_pointer) + GET_ROCR_SYMBOL(hsa_amd_svm_attributes_set) + GET_ROCR_SYMBOL(hsa_amd_svm_attributes_get) + GET_ROCR_SYMBOL(hsa_amd_svm_prefetch_async) + GET_ROCR_SYMBOL(hsa_amd_portable_export_dmabuf) + GET_ROCR_SYMBOL(hsa_amd_portable_close_dmabuf) + GET_ROCR_SYMBOL(hsa_amd_vmem_address_reserve) + GET_ROCR_SYMBOL(hsa_amd_vmem_address_free) + GET_ROCR_SYMBOL(hsa_amd_vmem_handle_create) + GET_ROCR_SYMBOL(hsa_amd_vmem_handle_release) + GET_ROCR_SYMBOL(hsa_amd_vmem_map) + GET_ROCR_SYMBOL(hsa_amd_vmem_unmap) + GET_ROCR_SYMBOL(hsa_amd_vmem_set_access) + GET_ROCR_SYMBOL(hsa_amd_vmem_get_access) + GET_ROCR_SYMBOL(hsa_amd_vmem_export_shareable_handle) + GET_ROCR_SYMBOL(hsa_amd_vmem_import_shareable_handle) + GET_ROCR_SYMBOL(hsa_amd_vmem_retain_alloc_handle) + GET_ROCR_SYMBOL(hsa_amd_agent_set_async_scratch_limit) + GET_ROCR_SYMBOL(hsa_amd_vmem_address_reserve_align) + GET_ROCR_SYMBOL(hsa_amd_enable_logging) + GET_ROCR_SYMBOL(hsa_amd_memory_get_preferred_copy_engine) + + // Image extensions + GET_ROCR_SYMBOL(hsa_ext_image_data_get_info) + GET_ROCR_SYMBOL(hsa_ext_image_create) + GET_ROCR_SYMBOL(hsa_ext_image_import) + GET_ROCR_SYMBOL(hsa_ext_image_export) + GET_ROCR_SYMBOL(hsa_ext_image_destroy) + GET_ROCR_SYMBOL(hsa_ext_sampler_create_v2) + GET_ROCR_SYMBOL(hsa_ext_sampler_destroy) + GET_ROCR_SYMBOL(hsa_ext_image_create_with_layout) + is_ready_ = true; + return true; +} +} // namespace roc +} // namespace amd diff --git a/projects/clr/rocclr/device/rocm/rocrctx.hpp b/projects/clr/rocclr/device/rocm/rocrctx.hpp new file mode 100644 index 0000000000..73a1cded43 --- /dev/null +++ b/projects/clr/rocclr/device/rocm/rocrctx.hpp @@ -0,0 +1,571 @@ +/* Copyright (c) 2025 Advanced Micro Devices, Inc. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in 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: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + 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 + AUTHORS 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 IN + THE SOFTWARE. */ + +#pragma once + +#include +#include "top.hpp" + +#ifdef ROCR_DYN_DLL +#include "hsa.h" +#include "hsa_ext_image.h" +#include "hsa_ext_amd.h" +#include "amd_hsa_signal.h" +#include "hsa_ven_amd_loader.h" +#include "hsa_ven_amd_aqlprofile.h" +#else +#include "hsa/hsa.h" +#include "hsa/hsa_ext_image.h" +#include "hsa/hsa_ext_amd.h" +#include "hsa/amd_hsa_signal.h" +#include "hsa/hsa_ven_amd_loader.h" +#include "hsa/hsa_ven_amd_aqlprofile.h" +#endif + +namespace amd { +namespace roc { + +struct RocrEntryPoints { + void* handle; + + // Core functionality + decltype(hsa_init)* hsa_init_; + decltype(hsa_shut_down)* hsa_shut_down_; + decltype(hsa_system_get_info)* hsa_system_get_info_; + decltype(hsa_iterate_agents)* hsa_iterate_agents_; + decltype(hsa_agent_get_info)* hsa_agent_get_info_; + decltype(hsa_queue_create)* hsa_queue_create_; + decltype(hsa_queue_destroy)* hsa_queue_destroy_; + decltype(hsa_queue_load_read_index_scacquire)* hsa_queue_load_read_index_scacquire_; + decltype(hsa_queue_load_read_index_relaxed)* hsa_queue_load_read_index_relaxed_; + decltype(hsa_queue_load_write_index_relaxed)* hsa_queue_load_write_index_relaxed_; + decltype(hsa_queue_add_write_index_screlease)* hsa_queue_add_write_index_screlease_; + decltype(hsa_memory_register)* hsa_memory_register_; + decltype(hsa_memory_deregister)* hsa_memory_deregister_; + decltype(hsa_memory_copy)* hsa_memory_copy_; + decltype(hsa_signal_create)* hsa_signal_create_; + decltype(hsa_signal_destroy)* hsa_signal_destroy_; + decltype(hsa_signal_load_relaxed)* hsa_signal_load_relaxed_; + decltype(hsa_signal_store_relaxed)* hsa_signal_store_relaxed_; + decltype(hsa_signal_silent_store_relaxed)* hsa_signal_silent_store_relaxed_; + decltype(hsa_signal_store_screlease)* hsa_signal_store_screlease_; + decltype(hsa_signal_wait_scacquire)* hsa_signal_wait_scacquire_; + decltype(hsa_signal_add_relaxed)* hsa_signal_add_relaxed_; + decltype(hsa_signal_subtract_relaxed)* hsa_signal_subtract_relaxed_; + decltype(hsa_isa_get_info_alt)* hsa_isa_get_info_alt_; + decltype(hsa_agent_iterate_isas)* hsa_agent_iterate_isas_; + decltype(hsa_system_get_major_extension_table)* hsa_system_get_major_extension_table_; + decltype(hsa_status_string)* hsa_status_string_; + decltype(hsa_executable_create_alt)* hsa_executable_create_alt_; + decltype(hsa_executable_destroy)* hsa_executable_destroy_; + decltype(hsa_executable_get_info)* hsa_executable_get_info_; + decltype(hsa_code_object_reader_destroy)* hsa_code_object_reader_destroy_; + decltype(hsa_code_object_reader_create_from_memory)* hsa_code_object_reader_create_from_memory_; + decltype(hsa_executable_load_agent_code_object)* hsa_executable_load_agent_code_object_; + decltype(hsa_executable_agent_global_variable_define)* + hsa_executable_agent_global_variable_define_; + decltype(hsa_executable_get_symbol_by_name)* hsa_executable_get_symbol_by_name_; + decltype(hsa_executable_symbol_get_info)* hsa_executable_symbol_get_info_; + decltype(hsa_executable_freeze)* hsa_executable_freeze_; + // AMD extensions + decltype(hsa_amd_coherency_set_type)* hsa_amd_coherency_set_type_; + decltype(hsa_amd_profiling_set_profiler_enabled)* hsa_amd_profiling_set_profiler_enabled_; + decltype(hsa_amd_profiling_async_copy_enable)* hsa_amd_profiling_async_copy_enable_; + decltype(hsa_amd_profiling_get_dispatch_time)* hsa_amd_profiling_get_dispatch_time_; + decltype(hsa_amd_profiling_get_async_copy_time)* hsa_amd_profiling_get_async_copy_time_; + decltype(hsa_amd_signal_async_handler)* hsa_amd_signal_async_handler_; + decltype(hsa_amd_queue_cu_set_mask)* hsa_amd_queue_cu_set_mask_; + decltype(hsa_amd_memory_pool_get_info)* hsa_amd_memory_pool_get_info_; + decltype(hsa_amd_agent_iterate_memory_pools)* hsa_amd_agent_iterate_memory_pools_; + decltype(hsa_amd_memory_pool_allocate)* hsa_amd_memory_pool_allocate_; + decltype(hsa_amd_memory_pool_free)* hsa_amd_memory_pool_free_; + decltype(hsa_amd_memory_async_copy)* hsa_amd_memory_async_copy_; + decltype(hsa_amd_memory_async_copy_on_engine)* hsa_amd_memory_async_copy_on_engine_; + decltype(hsa_amd_memory_copy_engine_status)* hsa_amd_memory_copy_engine_status_; + decltype(hsa_amd_agent_memory_pool_get_info)* hsa_amd_agent_memory_pool_get_info_; + decltype(hsa_amd_agents_allow_access)* hsa_amd_agents_allow_access_; + decltype(hsa_amd_memory_unlock)* hsa_amd_memory_unlock_; + decltype(hsa_amd_interop_map_buffer)* hsa_amd_interop_map_buffer_; + decltype(hsa_amd_interop_unmap_buffer)* hsa_amd_interop_unmap_buffer_; + decltype(hsa_amd_image_create)* hsa_amd_image_create_; + decltype(hsa_amd_pointer_info)* hsa_amd_pointer_info_; + decltype(hsa_amd_ipc_memory_create)* hsa_amd_ipc_memory_create_; + decltype(hsa_amd_ipc_memory_attach)* hsa_amd_ipc_memory_attach_; + decltype(hsa_amd_ipc_memory_detach)* hsa_amd_ipc_memory_detach_; + decltype(hsa_amd_signal_create)* hsa_amd_signal_create_; + decltype(hsa_amd_register_system_event_handler)* hsa_amd_register_system_event_handler_; + decltype(hsa_amd_queue_set_priority)* hsa_amd_queue_set_priority_; + decltype(hsa_amd_memory_async_copy_rect)* hsa_amd_memory_async_copy_rect_; + decltype(hsa_amd_memory_lock_to_pool)* hsa_amd_memory_lock_to_pool_; + decltype(hsa_amd_signal_value_pointer)* hsa_amd_signal_value_pointer_; + decltype(hsa_amd_svm_attributes_set)* hsa_amd_svm_attributes_set_; + decltype(hsa_amd_svm_attributes_get)* hsa_amd_svm_attributes_get_; + decltype(hsa_amd_svm_prefetch_async)* hsa_amd_svm_prefetch_async_; + decltype(hsa_amd_portable_export_dmabuf)* hsa_amd_portable_export_dmabuf_; + decltype(hsa_amd_portable_close_dmabuf)* hsa_amd_portable_close_dmabuf_; // CLR doesn't use it? + decltype(hsa_amd_vmem_address_reserve)* hsa_amd_vmem_address_reserve_; + decltype(hsa_amd_vmem_address_free)* hsa_amd_vmem_address_free_; + decltype(hsa_amd_vmem_handle_create)* hsa_amd_vmem_handle_create_; + decltype(hsa_amd_vmem_handle_release)* hsa_amd_vmem_handle_release_; + decltype(hsa_amd_vmem_map)* hsa_amd_vmem_map_; + decltype(hsa_amd_vmem_unmap)* hsa_amd_vmem_unmap_; + decltype(hsa_amd_vmem_set_access)* hsa_amd_vmem_set_access_; + decltype(hsa_amd_vmem_get_access)* hsa_amd_vmem_get_access_; + decltype(hsa_amd_vmem_export_shareable_handle)* hsa_amd_vmem_export_shareable_handle_; + decltype(hsa_amd_vmem_import_shareable_handle)* hsa_amd_vmem_import_shareable_handle_; + decltype(hsa_amd_vmem_retain_alloc_handle)* hsa_amd_vmem_retain_alloc_handle_; + decltype(hsa_amd_agent_set_async_scratch_limit)* hsa_amd_agent_set_async_scratch_limit_; + decltype(hsa_amd_vmem_address_reserve_align)* hsa_amd_vmem_address_reserve_align_; + decltype(hsa_amd_enable_logging)* hsa_amd_enable_logging_; + decltype(hsa_amd_memory_get_preferred_copy_engine)* hsa_amd_memory_get_preferred_copy_engine_; + // Image extensions + decltype(hsa_ext_image_data_get_info)* hsa_ext_image_data_get_info_; + decltype(hsa_ext_image_create)* hsa_ext_image_create_; + decltype(hsa_ext_image_import)* hsa_ext_image_import_; + decltype(hsa_ext_image_export)* hsa_ext_image_export_; + decltype(hsa_ext_image_destroy)* hsa_ext_image_destroy_; + decltype(hsa_ext_sampler_create_v2)* hsa_ext_sampler_create_v2_; + decltype(hsa_ext_sampler_destroy)* hsa_ext_sampler_destroy_; + decltype(hsa_ext_image_create_with_layout)* hsa_ext_image_create_with_layout_; +}; + +#ifdef ROCR_DYN_DLL +#define ROCR_DYN(NAME) cep_.NAME##_ +#define GET_ROCR_SYMBOL(NAME) \ + cep_.NAME##_ = reinterpret_cast(Os::getSymbol(cep_.handle, #NAME)); \ + if (nullptr == cep_.NAME##_) { \ + ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "Failed to load ROCR function %s", #NAME); \ + return false; \ + } +#define GET_ROCR_OPTIONAL_SYMBOL(NAME) \ + cep_.NAME = reinterpret_cast(Os::getSymbol(cep_.handle, #NAME)); +#else +#define ROCR_DYN(NAME) NAME +#define GET_ROCR_SYMBOL(NAME) +#define GET_ROCR_OPTIONAL_SYMBOL(NAME) +#endif + +class Hsa : public amd::AllStatic { + public: + static std::once_flag initialized; + + static bool LoadLib(); + + static bool IsReady() { return is_ready_; } + + static hsa_status_t init() { return ROCR_DYN(hsa_init)(); } + static hsa_status_t shut_down() { return ROCR_DYN(hsa_shut_down)(); } + static hsa_status_t system_get_info(hsa_system_info_t attribute, void* value) { + return ROCR_DYN(hsa_system_get_info)(attribute, value); + } + static hsa_status_t iterate_agents(hsa_status_t (*callback)(hsa_agent_t agent, void* data), + void* data) { + return ROCR_DYN(hsa_iterate_agents)(callback, data); + } + static hsa_status_t agent_get_info(hsa_agent_t agent, hsa_agent_info_t attribute, void* value) { + return ROCR_DYN(hsa_agent_get_info)(agent, attribute, value); + } + static hsa_status_t queue_create(hsa_agent_t agent, 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) { + return ROCR_DYN(hsa_queue_create)(agent, size, type, callback, data, private_segment_size, + group_segment_size, queue); + } + static hsa_status_t queue_destroy(hsa_queue_t* queue) { + return ROCR_DYN(hsa_queue_destroy)(queue); + } + static uint64_t queue_load_read_index_scacquire(const hsa_queue_t* queue) { + return ROCR_DYN(hsa_queue_load_read_index_scacquire)(queue); + } + static uint64_t queue_load_read_index_relaxed(const hsa_queue_t* queue) { + return ROCR_DYN(hsa_queue_load_read_index_relaxed)(queue); + } + static uint64_t queue_load_write_index_relaxed(const hsa_queue_t* queue) { + return ROCR_DYN(hsa_queue_load_write_index_relaxed)(queue); + } + static uint64_t queue_add_write_index_screlease(const hsa_queue_t* queue, uint64_t value) { + return ROCR_DYN(hsa_queue_add_write_index_screlease)(queue, value); + } + static hsa_status_t memory_register(void* ptr, size_t size) { + return ROCR_DYN(hsa_memory_register)(ptr, size); + } + static hsa_status_t memory_deregister(void* ptr, size_t size) { + return ROCR_DYN(hsa_memory_deregister)(ptr, size); + } + static hsa_status_t memory_copy(void* dst, const void* src, size_t size) { + return ROCR_DYN(hsa_memory_copy)(dst, src, size); + } + static hsa_status_t signal_create(hsa_signal_value_t initial_value, uint32_t num_consumers, + const hsa_agent_t* consumers, hsa_signal_t* signal) { + return ROCR_DYN(hsa_signal_create)(initial_value, num_consumers, consumers, signal); + } + static hsa_status_t signal_destroy(hsa_signal_t signal) { + return ROCR_DYN(hsa_signal_destroy)(signal); + } + static hsa_signal_value_t signal_load_relaxed(hsa_signal_t signal) { + return ROCR_DYN(hsa_signal_load_relaxed)(signal); + } + static void signal_silent_store_relaxed(hsa_signal_t signal, hsa_signal_value_t value) { + ROCR_DYN(hsa_signal_silent_store_relaxed)(signal, value); + } + static void signal_store_relaxed(hsa_signal_t signal, hsa_signal_value_t value) { + ROCR_DYN(hsa_signal_store_relaxed)(signal, value); + } + static void signal_store_screlease(hsa_signal_t signal, hsa_signal_value_t value) { + ROCR_DYN(hsa_signal_store_screlease)(signal, value); + } + static hsa_signal_value_t signal_wait_scacquire(hsa_signal_t signal, + hsa_signal_condition_t condition, + hsa_signal_value_t compare_value, + uint64_t timeout_hint, + hsa_wait_state_t wait_state_hint) { + return ROCR_DYN(hsa_signal_wait_scacquire)(signal, condition, compare_value, timeout_hint, + wait_state_hint); + } + static void signal_add_relaxed(hsa_signal_t signal, hsa_signal_value_t value) { + ROCR_DYN(hsa_signal_add_relaxed)(signal, value); + } + static void signal_subtract_relaxed(hsa_signal_t signal, hsa_signal_value_t value) { + ROCR_DYN(hsa_signal_subtract_relaxed)(signal, value); + } + static hsa_status_t isa_get_info_alt(hsa_isa_t isa, hsa_isa_info_t attribute, void* value) { + return ROCR_DYN(hsa_isa_get_info_alt)(isa, attribute, value); + } + static hsa_status_t agent_iterate_isas(hsa_agent_t agent, + hsa_status_t (*callback)(hsa_isa_t isa, void* data), void* data) { + return ROCR_DYN(hsa_agent_iterate_isas)(agent, callback, data); + } + static hsa_status_t system_get_major_extension_table(uint16_t extension, uint16_t version_major, + size_t table_length, void* table) { + return ROCR_DYN(hsa_system_get_major_extension_table)(extension, version_major, + table_length, table); + } + static hsa_status_t status_string(hsa_status_t status, const char** status_string) { + return ROCR_DYN(hsa_status_string)(status, status_string); + } + static hsa_status_t executable_create_alt( + hsa_profile_t profile, hsa_default_float_rounding_mode_t default_float_rounding_mode, + const char* options, hsa_executable_t* executable) { + return ROCR_DYN(hsa_executable_create_alt)(profile, default_float_rounding_mode, options, + executable); + } + static hsa_status_t executable_destroy(hsa_executable_t executable) { + return ROCR_DYN(hsa_executable_destroy)(executable); + } + static hsa_status_t executable_get_info(hsa_executable_t executable, + hsa_executable_info_t attribute, void* value) { + return ROCR_DYN(hsa_executable_get_info)(executable, attribute, value); + } + static hsa_status_t code_object_reader_destroy(hsa_code_object_reader_t code_object_reader) { + return ROCR_DYN(hsa_code_object_reader_destroy)(code_object_reader); + } + static hsa_status_t code_object_reader_create_from_memory( + const void* code_object, size_t size, hsa_code_object_reader_t* code_object_reader) { + return ROCR_DYN(hsa_code_object_reader_create_from_memory)(code_object, size, + code_object_reader); + } + static hsa_status_t 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) { + return ROCR_DYN(hsa_executable_load_agent_code_object)(executable, agent, code_object_reader, + options, loaded_code_object); + } + static hsa_status_t executable_agent_global_variable_define(hsa_executable_t executable, + hsa_agent_t agent, const char* variable_name, void* address) { + return ROCR_DYN(hsa_executable_agent_global_variable_define)(executable, agent, + variable_name, address); + } + static hsa_status_t executable_get_symbol_by_name(hsa_executable_t executable, + const char* symbol_name, const hsa_agent_t* agent, hsa_executable_symbol_t* symbol) { + return ROCR_DYN(hsa_executable_get_symbol_by_name)(executable, symbol_name, agent, symbol); + } + static hsa_status_t executable_symbol_get_info(hsa_executable_symbol_t executable_symbol, + hsa_executable_symbol_info_t attribute, void* value) { + return ROCR_DYN(hsa_executable_symbol_get_info)(executable_symbol, attribute, value); + } + static hsa_status_t executable_freeze(hsa_executable_t executable, const char* options) { + return ROCR_DYN(hsa_executable_freeze)(executable, options); + } + // AMD extensions + static hsa_status_t coherency_set_type(hsa_agent_t agent, hsa_amd_coherency_type_t type) { + return ROCR_DYN(hsa_amd_coherency_set_type)(agent, type); + } + static hsa_status_t profiling_set_profiler_enabled(hsa_queue_t* queue, int enable) { + return ROCR_DYN(hsa_amd_profiling_set_profiler_enabled)(queue, enable); + } + static hsa_status_t profiling_async_copy_enable(bool enable) { + return ROCR_DYN(hsa_amd_profiling_async_copy_enable)(enable); + } + static hsa_status_t profiling_get_dispatch_time(hsa_agent_t agent, hsa_signal_t signal, + hsa_amd_profiling_dispatch_time_t* time) { + return ROCR_DYN(hsa_amd_profiling_get_dispatch_time)(agent, signal, time); + } + static hsa_status_t profiling_get_async_copy_time(hsa_signal_t signal, + hsa_amd_profiling_async_copy_time_t* time) { + return ROCR_DYN(hsa_amd_profiling_get_async_copy_time)(signal, time); + } + static hsa_status_t signal_async_handler(hsa_signal_t signal, hsa_signal_condition_t cond, + hsa_signal_value_t value, hsa_amd_signal_handler handler, + void* arg) { + return ROCR_DYN(hsa_amd_signal_async_handler)(signal, cond, value, handler, arg); + } + static hsa_status_t queue_cu_set_mask(const hsa_queue_t* queue, uint32_t num_cu_mask_count, + const uint32_t* cu_mask) { + return ROCR_DYN(hsa_amd_queue_cu_set_mask)(queue, num_cu_mask_count, cu_mask); + } + static hsa_status_t memory_pool_get_info(hsa_amd_memory_pool_t memory_pool, + hsa_amd_memory_pool_info_t attribute, void* value) { + return ROCR_DYN(hsa_amd_memory_pool_get_info)(memory_pool, attribute, value); + } + static hsa_status_t agent_iterate_memory_pools( + hsa_agent_t agent, hsa_status_t (*callback)(hsa_amd_memory_pool_t memory_pool, void* data), + void* data) { + return ROCR_DYN(hsa_amd_agent_iterate_memory_pools)(agent, callback, data); + } + static hsa_status_t memory_pool_allocate(hsa_amd_memory_pool_t memory_pool, size_t size, + uint32_t flags, void** ptr) { + return ROCR_DYN(hsa_amd_memory_pool_allocate)(memory_pool, size, flags, ptr); + } + static hsa_status_t memory_pool_free(void* ptr) { + return ROCR_DYN(hsa_amd_memory_pool_free)(ptr); + } + static hsa_status_t memory_async_copy(void* dst, hsa_agent_t dst_agent, const void* src, + hsa_agent_t src_agent, size_t size, + uint32_t num_dep_signals, const hsa_signal_t* dep_signals, + hsa_signal_t completion_signal) { + return ROCR_DYN(hsa_amd_memory_async_copy)(dst, dst_agent, src, src_agent, size, + num_dep_signals, dep_signals, completion_signal); + } + static hsa_status_t memory_async_copy_on_engine( + void* dst, hsa_agent_t dst_agent, const void* src, hsa_agent_t src_agent, size_t size, + uint32_t num_dep_signals, const hsa_signal_t* dep_signals, hsa_signal_t completion_signal, + hsa_amd_sdma_engine_id_t engine_id, bool force_copy_on_sdma) { + return ROCR_DYN(hsa_amd_memory_async_copy_on_engine)( + dst, dst_agent, src, src_agent, size, num_dep_signals, dep_signals, completion_signal, + engine_id, force_copy_on_sdma); + } + static hsa_status_t memory_copy_engine_status(hsa_agent_t dst_agent, + hsa_agent_t src_agent, uint32_t* engine_ids_mask) { + return ROCR_DYN(hsa_amd_memory_copy_engine_status)(dst_agent, src_agent, engine_ids_mask); + } + static hsa_status_t agent_memory_pool_get_info(hsa_agent_t agent, + hsa_amd_memory_pool_t memory_pool, hsa_amd_agent_memory_pool_info_t attribute, void* value) { + return ROCR_DYN(hsa_amd_agent_memory_pool_get_info)(agent, memory_pool, attribute, value); + } + static hsa_status_t agents_allow_access(uint32_t num_agents, const hsa_agent_t* agents, + const uint32_t* flags, const void* ptr) { + return ROCR_DYN(hsa_amd_agents_allow_access)(num_agents, agents, flags, ptr); + } + static hsa_status_t memory_lock_to_pool(void* host_ptr, size_t size, hsa_agent_t* agents, + int num_agent, hsa_amd_memory_pool_t pool, uint32_t flags, void** agent_ptr) { + return ROCR_DYN(hsa_amd_memory_lock_to_pool)(host_ptr, size, agents, num_agent, pool, flags, + agent_ptr); + } + static hsa_status_t memory_unlock(void* host_ptr) { + return ROCR_DYN(hsa_amd_memory_unlock)(host_ptr); + } + static hsa_status_t interop_map_buffer(uint32_t num_agents, hsa_agent_t* agents, + int interop_handle, uint32_t flags, size_t* size, + void** ptr, size_t* metadata_size, const void** metadata) { + return ROCR_DYN(hsa_amd_interop_map_buffer)(num_agents, agents, interop_handle, flags, size, + ptr, metadata_size, metadata); + } + static hsa_status_t interop_unmap_buffer(void* ptr) { + return ROCR_DYN(hsa_amd_interop_unmap_buffer)(ptr); + } + static hsa_status_t pointer_info(const void* ptr, hsa_amd_pointer_info_t* info, + void* (*alloc)(size_t), uint32_t* num_agents_accessible, hsa_agent_t** accessible) { + return ROCR_DYN(hsa_amd_pointer_info)(ptr, info, alloc, num_agents_accessible, accessible); + } + static hsa_status_t ipc_memory_create(void* ptr, size_t len, hsa_amd_ipc_memory_t* handle) { + return ROCR_DYN(hsa_amd_ipc_memory_create)(ptr, len, handle); + } + static hsa_status_t ipc_memory_attach(const hsa_amd_ipc_memory_t* handle, size_t len, + uint32_t num_agents, const hsa_agent_t* mapping_agents, void** mapped_ptr) { + return ROCR_DYN(hsa_amd_ipc_memory_attach)( + handle, len, num_agents, mapping_agents, mapped_ptr); + } + static hsa_status_t ipc_memory_detach(void* mapped_ptr) { + return ROCR_DYN(hsa_amd_ipc_memory_detach)(mapped_ptr); + } + static hsa_status_t signal_create(hsa_signal_value_t initial_value, uint32_t num_consumers, + const hsa_agent_t* consumers, uint64_t attributes, hsa_signal_t* signal) { + return ROCR_DYN(hsa_amd_signal_create)(initial_value, num_consumers, consumers, attributes, + signal); + } + static hsa_status_t register_system_event_handler(hsa_amd_system_event_callback_t callback, + void* data) { + return ROCR_DYN(hsa_amd_register_system_event_handler)(callback, data); + } + static hsa_status_t queue_set_priority(hsa_queue_t* queue, hsa_amd_queue_priority_t priority) { + return ROCR_DYN(hsa_amd_queue_set_priority)(queue, priority); + } + static hsa_status_t memory_async_copy_rect( + const hsa_pitched_ptr_t* dst, const hsa_dim3_t* dst_offset, const hsa_pitched_ptr_t* src, + const hsa_dim3_t* src_offset, const hsa_dim3_t* range, hsa_agent_t copy_agent, + hsa_amd_copy_direction_t dir, uint32_t num_dep_signals, const hsa_signal_t* dep_signals, + hsa_signal_t completion_signal) { + return ROCR_DYN(hsa_amd_memory_async_copy_rect)(dst, dst_offset, src, src_offset, range, + copy_agent, dir, num_dep_signals, dep_signals, + completion_signal); + } + static hsa_status_t signal_value_pointer(hsa_signal_t signal, + volatile hsa_signal_value_t** value_ptr) { + return ROCR_DYN(hsa_amd_signal_value_pointer)(signal, value_ptr); + } + static hsa_status_t svm_attributes_set(void* ptr, size_t size, + hsa_amd_svm_attribute_pair_t* attribute_list, size_t attribute_count) { + return ROCR_DYN(hsa_amd_svm_attributes_set)(ptr, size, attribute_list, attribute_count); + } + static hsa_status_t svm_attributes_get(void* ptr, size_t size, + hsa_amd_svm_attribute_pair_t* attribute_list, size_t attribute_count) { + return ROCR_DYN(hsa_amd_svm_attributes_get)(ptr, size, attribute_list, attribute_count); + } + static hsa_status_t svm_prefetch_async(void* ptr, size_t size, hsa_agent_t agent, + uint32_t num_dep_signals, const hsa_signal_t* dep_signals, hsa_signal_t completion_signal) { + return ROCR_DYN(hsa_amd_svm_prefetch_async)(ptr, size, agent, num_dep_signals, + dep_signals, completion_signal); + } + static hsa_status_t portable_export_dmabuf(const void* ptr, size_t size, int* dmabuf, + uint64_t* offset) { + return ROCR_DYN(hsa_amd_portable_export_dmabuf)(ptr, size, dmabuf, offset); + } + static hsa_status_t vmem_address_reserve(void** ptr, size_t size, uint64_t address, + uint64_t flags) { + return ROCR_DYN(hsa_amd_vmem_address_reserve)(ptr, size, address, flags); + } + static hsa_status_t vmem_address_free(void* ptr, size_t size) { + return ROCR_DYN(hsa_amd_vmem_address_free)(ptr, size); + } + static hsa_status_t vmem_handle_create(hsa_amd_memory_pool_t pool, size_t size, + hsa_amd_memory_type_t type, uint64_t flags, hsa_amd_vmem_alloc_handle_t* memory_handle) { + return ROCR_DYN(hsa_amd_vmem_handle_create)(pool, size, type, flags, memory_handle); + } + static hsa_status_t vmem_handle_release(hsa_amd_vmem_alloc_handle_t memory_handle) { + return ROCR_DYN(hsa_amd_vmem_handle_release)(memory_handle); + } + static hsa_status_t vmem_map(void* va, size_t size, size_t in_offset, + hsa_amd_vmem_alloc_handle_t memory_handle, uint64_t flags) { + return ROCR_DYN(hsa_amd_vmem_map)(va, size, in_offset, memory_handle, flags); + } + static hsa_status_t vmem_unmap(void* va, size_t size) { + return ROCR_DYN(hsa_amd_vmem_unmap)(va, size); + } + static hsa_status_t vmem_set_access(void* va, size_t size, + const hsa_amd_memory_access_desc_t* desc, const size_t desc_cnt) { + return ROCR_DYN(hsa_amd_vmem_set_access)(va, size, desc, desc_cnt); + } + static hsa_status_t vmem_get_access(void* va, hsa_access_permission_t* flags, + const hsa_agent_t agent_handle) { + return ROCR_DYN(hsa_amd_vmem_get_access)(va, flags, agent_handle); + } + static hsa_status_t vmem_export_shareable_handle(int* dmabuf_fd, + hsa_amd_vmem_alloc_handle_t handle, uint64_t flags) { + return ROCR_DYN(hsa_amd_vmem_export_shareable_handle)(dmabuf_fd, handle, flags); + } + static hsa_status_t vmem_import_shareable_handle(int dmabuf_fd, + hsa_amd_vmem_alloc_handle_t* handle) { + return ROCR_DYN(hsa_amd_vmem_import_shareable_handle)(dmabuf_fd, handle); + } + static hsa_status_t vmem_retain_alloc_handle(hsa_amd_vmem_alloc_handle_t* allocHandle, + void* addr) { + return ROCR_DYN(hsa_amd_vmem_retain_alloc_handle)(allocHandle, addr); + } + static hsa_status_t agent_set_async_scratch_limit(hsa_agent_t agent, size_t threshold) { + return ROCR_DYN(hsa_amd_agent_set_async_scratch_limit)(agent, threshold); + } + static hsa_status_t vmem_address_reserve_align(void** ptr, size_t size, uint64_t address, + uint64_t alignment, uint64_t flags) { + return ROCR_DYN(hsa_amd_vmem_address_reserve_align)(ptr, size, address, alignment, flags); + } + static hsa_status_t enable_logging(uint8_t* flags, void* file) { + return ROCR_DYN(hsa_amd_enable_logging)(flags, file); + } + static hsa_status_t memory_get_preferred_copy_engine(hsa_agent_t dst_agent, + hsa_agent_t src_agent, uint32_t* recommended_ids_mask) { + return ROCR_DYN(hsa_amd_memory_get_preferred_copy_engine)( + dst_agent, src_agent, recommended_ids_mask); + } + + // Image extensions + static hsa_status_t image_create(hsa_agent_t agent, + const hsa_ext_image_descriptor_t* image_descriptor, + const hsa_amd_image_descriptor_t* image_layout, + const void* image_data, hsa_access_permission_t access_permission, + hsa_ext_image_t* image) { + return ROCR_DYN(hsa_amd_image_create)(agent, image_descriptor, image_layout, image_data, + access_permission, image); + } + static hsa_status_t image_data_get_info( + hsa_agent_t agent, const hsa_ext_image_descriptor_t* image_descriptor, + hsa_access_permission_t access_permission, hsa_ext_image_data_info_t* image_data_info) { + return ROCR_DYN(hsa_ext_image_data_get_info)(agent, image_descriptor, access_permission, + image_data_info); + } + static hsa_status_t image_create(hsa_agent_t agent, + const hsa_ext_image_descriptor_t* image_descriptor, + const void* image_data, + hsa_access_permission_t access_permission, + hsa_ext_image_t* image) { + return ROCR_DYN(hsa_ext_image_create)(agent, image_descriptor, image_data, + access_permission, image); + } + static hsa_status_t image_import(hsa_agent_t agent, const void* src_memory, + size_t src_row_pitch, size_t src_slice_pitch, + hsa_ext_image_t dst_image, + const hsa_ext_image_region_t* image_region) { + return ROCR_DYN(hsa_ext_image_import)(agent, src_memory, src_row_pitch, src_slice_pitch, + dst_image, image_region); + } + static hsa_status_t image_export(hsa_agent_t agent, hsa_ext_image_t src_image, + void* dst_memory, size_t dst_row_pitch, + size_t dst_slice_pitch, + const hsa_ext_image_region_t* image_region) { + return ROCR_DYN(hsa_ext_image_export)(agent, src_image, dst_memory, dst_row_pitch, + dst_slice_pitch, image_region); + } + static hsa_status_t image_destroy(hsa_agent_t agent, hsa_ext_image_t image) { + return ROCR_DYN(hsa_ext_image_destroy)(agent, image); + } + static hsa_status_t sampler_create(hsa_agent_t agent, + const hsa_ext_sampler_descriptor_v2_t* sampler_descriptor, hsa_ext_sampler_t* sampler) { + return ROCR_DYN(hsa_ext_sampler_create_v2)(agent, sampler_descriptor, sampler); + } + static hsa_status_t sampler_destroy(hsa_agent_t agent, hsa_ext_sampler_t sampler) { + return ROCR_DYN(hsa_ext_sampler_destroy)(agent, sampler); + } + static hsa_status_t image_create_with_layout( + hsa_agent_t agent, const hsa_ext_image_descriptor_t* image_descriptor, const void* image_data, + hsa_access_permission_t access_permission, hsa_ext_image_data_layout_t image_data_layout, + size_t image_data_row_pitch, size_t image_data_slice_pitch, hsa_ext_image_t* image) { + return ROCR_DYN(hsa_ext_image_create_with_layout)( + agent, image_descriptor, image_data, access_permission, image_data_layout, + image_data_row_pitch, image_data_slice_pitch, image); + } + + private: + static RocrEntryPoints cep_; + static bool is_ready_; +}; + +} // namespace roc +} // namespace amd diff --git a/projects/clr/rocclr/device/rocm/rocsignal.cpp b/projects/clr/rocclr/device/rocm/rocsignal.cpp index 05d9bcfd29..22b5668bf1 100644 --- a/projects/clr/rocclr/device/rocm/rocsignal.cpp +++ b/projects/clr/rocclr/device/rocm/rocsignal.cpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2021 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2021 - 2025 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -22,12 +22,14 @@ #include "utils/flags.hpp" #include "utils/debug.hpp" #include "rocsignal.hpp" +#include "device/rocm/rocrctx.hpp" + namespace amd::roc { -Signal::~Signal() { hsa_signal_destroy(signal_); } +Signal::~Signal() { Hsa::signal_destroy(signal_); } bool Signal::Init(const amd::Device& dev, uint64_t init, device::Signal::WaitState ws) { - hsa_status_t status = hsa_signal_create(init, 0, nullptr, &signal_); + hsa_status_t status = Hsa::signal_create(init, 0, nullptr, &signal_); if (status != HSA_STATUS_SUCCESS) { return false; } @@ -38,10 +40,10 @@ bool Signal::Init(const amd::Device& dev, uint64_t init, device::Signal::WaitSta } uint64_t Signal::Wait(uint64_t value, device::Signal::Condition c, uint64_t timeout) { - return hsa_signal_wait_scacquire(signal_, static_cast(c), value, timeout, - static_cast(ws_)); + return Hsa::signal_wait_scacquire(signal_, static_cast(c), value, timeout, + static_cast(ws_)); } -void Signal::Reset(uint64_t value) { hsa_signal_store_screlease(signal_, value); } +void Signal::Reset(uint64_t value) { Hsa::signal_store_screlease(signal_, value); } }; // namespace amd::roc \ No newline at end of file diff --git a/projects/clr/rocclr/device/rocm/rocsignal.hpp b/projects/clr/rocclr/device/rocm/rocsignal.hpp index 0f7e4093c0..c120c1204d 100644 --- a/projects/clr/rocclr/device/rocm/rocsignal.hpp +++ b/projects/clr/rocclr/device/rocm/rocsignal.hpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2021 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2021 - 2025 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -20,10 +20,9 @@ #pragma once +#include "device/rocm/rocdevice.hpp" #include "device/devsignal.hpp" -#include "hsa/hsa.h" - namespace amd::roc { class Signal : public device::Signal { diff --git a/projects/clr/rocclr/device/rocm/rocurilocator.cpp b/projects/clr/rocclr/device/rocm/rocurilocator.cpp index f9221c459e..9e81db9abe 100644 --- a/projects/clr/rocclr/device/rocm/rocurilocator.cpp +++ b/projects/clr/rocclr/device/rocm/rocurilocator.cpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2021 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2021 - 2025 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -28,7 +28,7 @@ hsa_status_t UriLocator::createUriRangeTable() { auto execCb = [](hsa_executable_t exec, void* data) -> hsa_status_t { int execState = 0; hsa_status_t status; - status = hsa_executable_get_info(exec, HSA_EXECUTABLE_INFO_STATE, &execState); + status = Hsa::executable_get_info(exec, HSA_EXECUTABLE_INFO_STATE, &execState); if (status != HSA_STATUS_SUCCESS) return status; if (execState != HSA_EXECUTABLE_STATE_FROZEN) return status; @@ -146,8 +146,8 @@ UriLocator::UriInfo UriLocator::lookUpUri(uint64_t device_pc) { if (!init_) { hsa_status_t result; - result = hsa_system_get_major_extension_table(HSA_EXTENSION_AMD_LOADER, 1, sizeof(fn_table_), - &fn_table_); + result = Hsa::system_get_major_extension_table(HSA_EXTENSION_AMD_LOADER, 1, sizeof(fn_table_), + &fn_table_); if (result != HSA_STATUS_SUCCESS) return errorstate; result = createUriRangeTable(); if (result != HSA_STATUS_SUCCESS) { diff --git a/projects/clr/rocclr/device/rocm/rocurilocator.hpp b/projects/clr/rocclr/device/rocm/rocurilocator.hpp index ff41cddfe5..94a14f93d0 100644 --- a/projects/clr/rocclr/device/rocm/rocurilocator.hpp +++ b/projects/clr/rocclr/device/rocm/rocurilocator.hpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2019 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2019 - 2025 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -22,7 +22,6 @@ #if defined(__clang__) #if __has_feature(address_sanitizer) #include "device/devurilocator.hpp" -#include "hsa/hsa_ven_amd_loader.h" #include namespace amd::roc { diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/device/rocm/rocvirtual.cpp index 89153e20fc..339fec50c6 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.cpp +++ b/projects/clr/rocclr/device/rocm/rocvirtual.cpp @@ -34,9 +34,6 @@ #include "platform/sampler.hpp" #include "utils/debug.hpp" #include "os/os.hpp" -#include "hsa/amd_hsa_kernel_code.h" -#include "hsa/amd_hsa_queue.h" -#include "hsa/amd_hsa_signal.h" #include #include @@ -146,10 +143,10 @@ void Timestamp::checkGpuTime() { amd_signal_t* amdSignal = reinterpret_cast(it->signal_.handle); if (it->engine_ == HwQueueEngine::Compute) { - hsa_amd_profiling_get_dispatch_time(gpu()->gpu_device(), it->signal_, &time); + Hsa::profiling_get_dispatch_time(gpu()->gpu_device(), it->signal_, &time); } else { hsa_amd_profiling_async_copy_time_t time_sdma = {}; - hsa_amd_profiling_get_async_copy_time(it->signal_, &time_sdma); + Hsa::profiling_get_async_copy_time(it->signal_, &time_sdma); time.start = time_sdma.start; time.end = time_sdma.end; } @@ -198,8 +195,8 @@ bool HsaAmdSignalHandler(hsa_signal_value_t value, void* arg) { ts->setParsedCommand(head); for (auto it : headTs->Signals()) { hsa_signal_value_t complete_val = (headTs->GetCallbackSignal().handle != 0) ? 1 : 0; - if (int64_t val = hsa_signal_load_relaxed(it->signal_) > complete_val) { - hsa_status_t result = hsa_amd_signal_async_handler( + if (int64_t val = Hsa::signal_load_relaxed(it->signal_) > complete_val) { + hsa_status_t result = Hsa::signal_async_handler( headTs->Signals()[0]->signal_, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne, &HsaAmdSignalHandler, ts); if (HSA_STATUS_SUCCESS != result) { @@ -239,7 +236,7 @@ bool HsaAmdSignalHandler(hsa_signal_value_t value, void* arg) { // Reset API callback signal. It will release AQL queue and start commands processing if (callback_signal.handle != 0 && isBlocking) { - hsa_signal_subtract_relaxed(callback_signal, 1); + Hsa::signal_subtract_relaxed(callback_signal, 1); } // Return false, so the callback will not be called again for this signal @@ -379,12 +376,12 @@ bool VirtualGPU::HwQueueTracker::CreateSignal(ProfilingSignal* signal, bool inte interrupt |= !AMD_DIRECT_DISPATCH || !gpu_.dev().ActiveWait(); // Check if the interrupt was requested for the signal if (interrupt && settings.system_scope_signal_) { - if (HSA_STATUS_SUCCESS != hsa_signal_create(0, 0, nullptr, &signal->signal_)) { + if (HSA_STATUS_SUCCESS != Hsa::signal_create(0, 0, nullptr, &signal->signal_)) { return false; } } else { if (HSA_STATUS_SUCCESS != - hsa_amd_signal_create(0, 0, nullptr, HSA_AMD_SIGNAL_AMD_GPU_ONLY, &signal->signal_)) { + Hsa::signal_create(0, 0, nullptr, HSA_AMD_SIGNAL_AMD_GPU_ONLY, &signal->signal_)) { return false; } } @@ -436,9 +433,10 @@ hsa_signal_t VirtualGPU::HwQueueTracker::ActiveSignal(hsa_signal_value_t init_va // Peep signal +2 ahead to see if its done auto temp_id = (current_id_ + 2) % signal_list_.size(); + // If GPU is still busy with processing or if timestamps havent been saved out, // then add more signals to avoid more frequent stalls - if (hsa_signal_load_relaxed(signal_list_[temp_id]->signal_) > 0 || + if (Hsa::signal_load_relaxed(signal_list_[temp_id]->signal_) > 0 || !signal_list_[temp_id]->flags_.done_) { std::unique_ptr signal(new ProfilingSignal()); if ((signal != nullptr) && CreateSignal(signal.get())) { @@ -509,7 +507,7 @@ hsa_signal_t VirtualGPU::HwQueueTracker::ActiveSignal(hsa_signal_value_t init_va } ProfilingSignal* prof_signal = signal_list_[current_id_]; // Reset the signal and return - hsa_signal_silent_store_relaxed(prof_signal->signal_, init_val); + Hsa::signal_silent_store_relaxed(prof_signal->signal_, init_val); prof_signal->flags_.done_ = false; prof_signal->engine_ = engine_; prof_signal->flags_.isPacketDispatch_ = false; @@ -542,12 +540,12 @@ hsa_signal_t VirtualGPU::HwQueueTracker::ActiveSignal(hsa_signal_value_t init_va ts->SetCallbackSignal(prof_signal->signal_, blocking); // Blocks AQL queue from further processing if (blocking) { - hsa_signal_add_relaxed(prof_signal->signal_, 1); + Hsa::signal_add_relaxed(prof_signal->signal_, 1); init_value += 1; } } gpu_.QueuedAsyncHandlers()++; - hsa_status_t result = hsa_amd_signal_async_handler( + hsa_status_t result = Hsa::signal_async_handler( prof_signal->signal_, HSA_SIGNAL_CONDITION_LT, init_value, &HsaAmdSignalHandler, ts); if (HSA_STATUS_SUCCESS != result) { LogError("hsa_amd_signal_async_handler() failed to set the handler!"); @@ -605,7 +603,7 @@ std::vector& VirtualGPU::HwQueueTracker::WaitingSignal(HwQueueEngi // Validate all signals for the wait and skip already completed for (uint32_t i = 0; i < external_signals_.size(); ++i) { // Early signal status check - if (hsa_signal_load_relaxed(external_signals_[i]->signal_) > 0) { + if (Hsa::signal_load_relaxed(external_signals_[i]->signal_) > 0) { const Settings& settings = gpu_.dev().settings(); if (settings.cpu_wait_for_signal_) { // Wait on CPU for completion if requested @@ -631,7 +629,7 @@ bool VirtualGPU::HwQueueTracker::CpuWaitForSignal(ProfilingSignal* signal) { ts->checkGpuTime(); ts->release(); signal->ts_ = nullptr; - } else if (hsa_signal_load_relaxed(signal->signal_) > 0) { + } else if (Hsa::signal_load_relaxed(signal->signal_) > 0) { amd::ScopedLock lock(signal->LockSignalOps()); ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "Host wait on completion_signal=0x%zx", signal->signal_.handle); @@ -647,7 +645,7 @@ bool VirtualGPU::HwQueueTracker::CpuWaitForSignal(ProfilingSignal* signal) { // ================================================================================================ void VirtualGPU::HwQueueTracker::ResetCurrentSignal() { // Reset the signal and return - hsa_signal_silent_store_relaxed(signal_list_[current_id_]->signal_, 0); + Hsa::signal_silent_store_relaxed(signal_list_[current_id_]->signal_, 0); // Fallback to the previous signal current_id_ = (current_id_ == 0) ? (signal_list_.size() - 1) : (current_id_ - 1); } @@ -936,8 +934,8 @@ void VirtualGPU::AnalyzeAqlQueue() const { const uint32_t queueSize = gpu_queue_->size; const uint32_t queueMask = queueSize - 1; const uint32_t sw_queue_size = queueMask; - uint64_t index = hsa_queue_load_write_index_relaxed(gpu_queue_); - uint64_t read = hsa_queue_load_read_index_relaxed(gpu_queue_); + uint64_t index = Hsa::queue_load_write_index_relaxed(gpu_queue_); + uint64_t read = Hsa::queue_load_read_index_relaxed(gpu_queue_); if (index > read) { int valid_packet_idx = 0; constexpr int kAqlSearchWindow = 32; @@ -1010,7 +1008,7 @@ bool VirtualGPU::dispatchGenericAqlPacket(AqlPacket* packet, uint16_t header, ui const uint32_t sw_queue_size = queueMask; // Check for queue full and wait if needed. - uint64_t index = hsa_queue_add_write_index_screlease(gpu_queue_, 1); + uint64_t index = Hsa::queue_add_write_index_screlease(gpu_queue_, 1); fence_dirty_ = true; if (addSystemScope_) { @@ -1057,7 +1055,7 @@ bool VirtualGPU::dispatchGenericAqlPacket(AqlPacket* packet, uint16_t header, ui } // Make sure the slot is free for usage - while ((index - hsa_queue_load_read_index_scacquire(gpu_queue_)) >= sw_queue_size) { + while ((index - Hsa::queue_load_read_index_scacquire(gpu_queue_)) >= sw_queue_size) { // Active spin - no yield } @@ -1104,9 +1102,9 @@ bool VirtualGPU::dispatchGenericAqlPacket(AqlPacket* packet, uint16_t header, ui reinterpret_cast(packet)->kernarg_address, reinterpret_cast(packet)->completion_signal, reinterpret_cast(packet)->reserved2, - hsa_queue_load_read_index_scacquire(gpu_queue_), index); + Hsa::queue_load_read_index_scacquire(gpu_queue_), index); - hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index); + Hsa::signal_store_screlease(gpu_queue_->doorbell_signal, index); // Mark the flag indicating if a dispatch is outstanding. // We are not waiting after every dispatch. @@ -1180,8 +1178,8 @@ bool VirtualGPU::dispatchGenericAqlPacketBatch(const std::vector& pa size_t batchSize = 1; while (processedPackets < numPackets) { - uint64_t currentReadIndex = hsa_queue_load_read_index_scacquire(gpu_queue_); - uint64_t currentWriteIndex = hsa_queue_load_write_index_relaxed(gpu_queue_); + uint64_t currentReadIndex = Hsa::queue_load_read_index_scacquire(gpu_queue_); + uint64_t currentWriteIndex = Hsa::queue_load_write_index_relaxed(gpu_queue_); if (currentWriteIndex - currentReadIndex >= kGpuLagPackets) { //GPU is busy, so we can copy more packets @@ -1200,10 +1198,10 @@ bool VirtualGPU::dispatchGenericAqlPacketBatch(const std::vector& pa } // Now reserve space for the batch - uint64_t startIndex = hsa_queue_add_write_index_screlease(gpu_queue_, batchSize); + uint64_t startIndex = Hsa::queue_add_write_index_screlease(gpu_queue_, batchSize); // Make sure the slot is free for usage - while ((startIndex - hsa_queue_load_read_index_scacquire(gpu_queue_)) >= sw_queue_size) { + while ((startIndex - Hsa::queue_load_read_index_scacquire(gpu_queue_)) >= sw_queue_size) { // Active spin - no yield } @@ -1321,7 +1319,7 @@ bool VirtualGPU::dispatchGenericAqlPacketBatch(const std::vector& pa reinterpret_cast(packet)->kernarg_address, reinterpret_cast(packet)->completion_signal, reinterpret_cast(packet)->reserved2, - hsa_queue_load_read_index_scacquire(gpu_queue_), index); + Hsa::queue_load_read_index_scacquire(gpu_queue_), index); } } } @@ -1331,7 +1329,7 @@ bool VirtualGPU::dispatchGenericAqlPacketBatch(const std::vector& pa packet_store_release(reinterpret_cast(aql_loc), firstPacketHeader, firstPacketRest); // Ring doorbell for this batch - hsa_signal_store_screlease(gpu_queue_->doorbell_signal, startIndex); + Hsa::signal_store_screlease(gpu_queue_->doorbell_signal, startIndex); processedPackets += batchSize; @@ -1429,8 +1427,8 @@ void VirtualGPU::dispatchBarrierPacket(uint16_t packetHeader, bool skipSignal, } } - uint64_t index = hsa_queue_add_write_index_screlease(gpu_queue_, 1); - uint64_t read = hsa_queue_load_read_index_relaxed(gpu_queue_); + uint64_t index = Hsa::queue_add_write_index_screlease(gpu_queue_, 1); + uint64_t read = Hsa::queue_load_read_index_relaxed(gpu_queue_); fence_dirty_ = true; auto cache_state = extractAqlBits(packetHeader, HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE, @@ -1450,13 +1448,13 @@ void VirtualGPU::dispatchBarrierPacket(uint16_t packetHeader, bool skipSignal, fence_dirty_ = false; } - while ((index - hsa_queue_load_read_index_scacquire(gpu_queue_)) >= queueMask); + while ((index - Hsa::queue_load_read_index_scacquire(gpu_queue_)) >= queueMask); hsa_barrier_and_packet_t* aql_loc = &(reinterpret_cast(gpu_queue_->base_address))[index & queueMask]; *aql_loc = barrier_packet_; packet_store_release(reinterpret_cast(aql_loc), packetHeader, 0); - hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index); + Hsa::signal_store_screlease(gpu_queue_->doorbell_signal, index); ClPrint(amd::LOG_DEBUG, amd::LOG_AQL, "SWq=0x%zx, HWq=0x%zx, id=%d, BarrierAND Header = 0x%x (type=%d, barrier=%d, acquire=%d," " release=%d), " @@ -1534,18 +1532,18 @@ void VirtualGPU::dispatchBarrierValuePacket(uint16_t packetHeader, bool resolveD fence_dirty_ = false; } - uint64_t index = hsa_queue_add_write_index_screlease(gpu_queue_, 1); - uint64_t read = hsa_queue_load_read_index_relaxed(gpu_queue_); + uint64_t index = Hsa::queue_add_write_index_screlease(gpu_queue_, 1); + uint64_t read = Hsa::queue_load_read_index_relaxed(gpu_queue_); TrackQueueProgress(barrier_value_packet_, index); - while ((index - hsa_queue_load_read_index_scacquire(gpu_queue_)) >= queueMask); + while ((index - Hsa::queue_load_read_index_scacquire(gpu_queue_)) >= queueMask); hsa_amd_barrier_value_packet_t* aql_loc = &(reinterpret_cast( gpu_queue_->base_address))[index & queueMask]; *aql_loc = barrier_value_packet_; packet_store_release(reinterpret_cast(aql_loc), packetHeader, rest); - hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index); + Hsa::signal_store_screlease(gpu_queue_->doorbell_signal, index); ClPrint(amd::LOG_DEBUG, amd::LOG_AQL, "SWq=0x%zx, HWq=0x%zx, id=%d, BarrierValue Header = 0x%x AmdFormat = 0x%x " @@ -1687,7 +1685,7 @@ VirtualGPU::~VirtualGPU() { delete printfdbg_; if (nullptr != schedulerQueue_) { - hsa_queue_destroy(schedulerQueue_); + Hsa::queue_destroy(schedulerQueue_); } if (nullptr != virtualQueue_) { @@ -1742,7 +1740,7 @@ bool VirtualGPU::create() { // Initialize timestamp conversion factor if (Timestamp::getGpuTicksToTime() == 0) { uint64_t frequency; - hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &frequency); + Hsa::system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &frequency); Timestamp::setGpuTicksToTime(1e9 / double(frequency)); } @@ -1771,7 +1769,7 @@ bool VirtualGPU::create() { VirtualGPU::ManagedBuffer::~ManagedBuffer() { for (auto& it : pool_signal_) { if (it.handle != 0) { - hsa_signal_destroy(it); + Hsa::signal_destroy(it); } } if (pool_base_ != nullptr) { @@ -1803,7 +1801,7 @@ bool VirtualGPU::ManagedBuffer::Create(Device::MemorySegment mem_segment) { } hsa_agent_t agent = gpu_.dev().getBackendDevice(); for (auto& it : pool_signal_) { - if (HSA_STATUS_SUCCESS != hsa_signal_create(0, 1, &agent, &it)) { + if (HSA_STATUS_SUCCESS != Hsa::signal_create(0, 1, &agent, &it)) { return false; } } @@ -1827,7 +1825,7 @@ address VirtualGPU::ManagedBuffer::Acquire(uint32_t size, uint32_t alignment) { return result; } else { // Reset the signal for the barrier packet - hsa_signal_silent_store_relaxed(pool_signal_[active_chunk_], kInitSignalValueOne); + Hsa::signal_silent_store_relaxed(pool_signal_[active_chunk_], kInitSignalValueOne); ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, "Issue barrier to flush chunk %d", active_chunk_); // Currently don't skip wait signal check, because SDMA engine cna be used in staging copy @@ -2306,8 +2304,8 @@ void VirtualGPU::submitSvmPrefetchAsync(amd::SvmPrefetchAsyncCommand& cmd) { // Initiate a prefetch command hsa_status_t status = - hsa_amd_svm_prefetch_async(const_cast(cmd.dev_ptr()), cmd.count(), agent, - wait_events.size(), wait_events.data(), active); + Hsa::svm_prefetch_async(const_cast(cmd.dev_ptr()), cmd.count(), agent, + wait_events.size(), wait_events.data(), active); ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "HSA prefetch async dev_ptr=0x%zx, count=%d, wait_event=0x%zx, " "completion_signal=0x%zx", @@ -3157,7 +3155,7 @@ void VirtualGPU::submitVirtualMap(amd::VirtualMapCommand& vcmd) { // Map the physical to virtual address the hsa api hsa_amd_vmem_alloc_handle_t opaque_hsa_handle; opaque_hsa_handle.handle = phys_mem_obj->getUserData().hsa_handle; - if ((hsa_status = hsa_amd_vmem_map(vaddr_sub_obj->getSvmPtr(), vcmd.size(), + if ((hsa_status = Hsa::vmem_map(vaddr_sub_obj->getSvmPtr(), vcmd.size(), vaddr_sub_obj->getOffset(), opaque_hsa_handle, 0)) == HSA_STATUS_SUCCESS) { assert(amd::MemObjMap::FindMemObj(vcmd.ptr()) == nullptr); @@ -3175,7 +3173,7 @@ void VirtualGPU::submitVirtualMap(amd::VirtualMapCommand& vcmd) { assert(vaddr_sub_obj != nullptr); // Unmap the object, since the physical addr is set. - if ((hsa_status = hsa_amd_vmem_unmap(vaddr_sub_obj->getSvmPtr(), vcmd.size())) == + if ((hsa_status = Hsa::vmem_unmap(vaddr_sub_obj->getSvmPtr(), vcmd.size())) == HSA_STATUS_SUCCESS) { // assert the va is mapped and needs to be removed vaddr_sub_obj->getContext().devices()[0]->DestroyVirtualBuffer(vaddr_sub_obj); @@ -3273,9 +3271,9 @@ bool VirtualGPU::createSchedulerParam() { while (true) { // The queue is written by multiple threads of the scheduler kernel if (HSA_STATUS_SUCCESS != - hsa_queue_create(gpu_device(), 2048, HSA_QUEUE_TYPE_MULTI, callbackQueue, &roc_device_, - std::numeric_limits::max(), std::numeric_limits::max(), - &schedulerQueue_)) { + Hsa::queue_create(gpu_device(), 2048, HSA_QUEUE_TYPE_MULTI, callbackQueue, &roc_device_, + std::numeric_limits::max(), std::numeric_limits::max(), + &schedulerQueue_)) { break; } @@ -3283,7 +3281,7 @@ bool VirtualGPU::createSchedulerParam() { } if (nullptr != schedulerQueue_) { - hsa_queue_destroy(schedulerQueue_); + Hsa::queue_destroy(schedulerQueue_); schedulerQueue_ = nullptr; } diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.hpp b/projects/clr/rocclr/device/rocm/rocvirtual.hpp index 3ade98d3a7..987a1fc826 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.hpp +++ b/projects/clr/rocclr/device/rocm/rocvirtual.hpp @@ -25,11 +25,7 @@ #include "rocdevice.hpp" #include "utils/flags.hpp" #include "utils/util.hpp" -#include "hsa/hsa.h" -#include "hsa/hsa_ext_image.h" -#include "hsa/hsa_ext_amd.h" #include "rocprintf.hpp" -#include "hsa/hsa_ven_amd_aqlprofile.h" #include "rocsched.hpp" #include "device/device.hpp" #include "os/os.hpp" @@ -56,13 +52,13 @@ inline bool WaitForSignal(hsa_signal_t signal, bool active_wait = false, bool yi wait_state = HSA_WAIT_STATE_ACTIVE; } - if (hsa_signal_load_relaxed(signal) > 0) { + if (Hsa::signal_load_relaxed(signal) > 0) { // When it is blocked wait, we wait in active state for 100 us before proceeding to wait in // blocked state indefinitely. if (!active_wait) { ClPrint(amd::LOG_INFO, amd::LOG_SIG, "Host active wait for Signal = (0x%lx) for %d ns", signal.handle, kTimeout100us); - if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne, + if (Hsa::signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne, kTimeout100us, HSA_WAIT_STATE_ACTIVE) != 0) { if (HIP_SKIP_ABORT_ON_GPU_ERROR && amd::Device::IsGPUInError()) { ClPrint(amd::LOG_ERROR, amd::LOG_SIG, @@ -76,7 +72,7 @@ inline bool WaitForSignal(hsa_signal_t signal, bool active_wait = false, bool yi // This is unlimited wait, but we wait for 4 secs and check if the device is // unstable, if so we return, otherwise we continue to wait in the while loop. - while (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne, + while (Hsa::signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne, kTimeout4Secs, wait_state) != 0) { if (HIP_SKIP_ABORT_ON_GPU_ERROR && amd::Device::IsGPUInError()) { ClPrint(amd::LOG_ERROR, amd::LOG_SIG, @@ -98,7 +94,7 @@ inline void fetchSignalTime(hsa_signal_t signal, hsa_agent_t gpu_device, uint64_ uint64_t* end) { if (start != nullptr && end != nullptr) { hsa_amd_profiling_dispatch_time_t time = {}; - hsa_amd_profiling_get_dispatch_time(gpu_device, signal, &time); + Hsa::profiling_get_dispatch_time(gpu_device, signal, &time); *start = time.start; *end = time.end; } @@ -307,7 +303,7 @@ class VirtualGPU : public device::VirtualDevice { bool GetSDMAProfiling() { return sdma_profiling_; } void SetSDMAProfiling(bool profile) { sdma_profiling_ = profile; - hsa_amd_profiling_async_copy_enable(profile); + Hsa::profiling_async_copy_enable(profile); } private: @@ -558,7 +554,7 @@ class VirtualGPU : public device::VirtualDevice { if ((last_write_index_ == 0) && (last_completion_signal_.handle == 0)) { result = true; } else { - result = (hsa_signal_load_relaxed(last_completion_signal_) == 0); + result = (Hsa::signal_load_relaxed(last_completion_signal_) == 0); } } return result;