AI/rocm-systems
2
0
Çatalla 0
Kod Konular Değişiklik İstekleri İşlemler Paketler Projeler Sürüm Viki Aktivite

SWDEV-547108 - Add dll loader for Windows build (#1004)

Kaynağa Gözat 
The build of ROCR backend will be enabled by default in Windows.
It requires the dll loader until ROCR dll will be always available in Windows for any configuration.
Bu işleme şunda yer alıyor:
German Andryeyev
2025-09-19 11:25:30 -04:00
işlemeyi yapan: GitHub
ebeveyn 18b4b84a9f
işleme ea89ddd589
19 değiştirilmiş dosya ile 1130 ekleme ve 419 silme
Yama Dosyasını İndir Diff Dosyasını İndir Tüm dosyaları genişlet Tüm dosyaları daralt
projects/clr/rocclr/cmake/ROCclrHSA.cmake
+50 -10
Dosyayı Görüntüle
@@ -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
projects/clr/rocclr/device/rocm/rocblit.cpp
+18 -18
Dosyayı Görüntüle
@@ -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<address>(dst) + dstOffset), dstAgent,
(reinterpret_cast<const_address>(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<const amd::Memory*>(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<address>(src), &info, nullptr, nullptr, nullptr)) {
Hsa::pointer_info(const_cast<address>(src), &info, nullptr, nullptr, nullptr)) {
srcAgent = info.agentOwner;
}
}
if (static_cast<const amd::Memory*>(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<uint64_t>(-1ULL);
}
projects/clr/rocclr/device/rocm/roccounters.cpp
+3 -3
Dosyayı Görüntüle
@@ -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) {
projects/clr/rocclr/device/rocm/roccounters.hpp
+2 -15
Dosyayı Görüntüle
@@ -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");
projects/clr/rocclr/device/rocm/rocdevice.cpp
+208 -218
Dosyayı Görüntüle
Dosya farkı çok büyük olduğundan ihmal edildi Fark Yükle
projects/clr/rocclr/device/rocm/rocdevice.hpp
+2 -12
Dosyayı Görüntüle
@@ -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 <atomic>
#include <iostream>
#include <vector>
@@ -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;
};
projects/clr/rocclr/device/rocm/rocglinterop.hpp
+2 -2
Dosyayı Görüntüle
@@ -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 {
projects/clr/rocclr/device/rocm/rockernel.cpp
+10 -11
Dosyayı Görüntüle
@@ -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 <algorithm>
@@ -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<void*>(variable_address), &runtime_handle, variable_size);
Hsa::memory_copy(reinterpret_cast<void*>(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;
projects/clr/rocclr/device/rocm/rocmemory.cpp
+37 -38
Dosyayı Görüntüle
@@ -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<char*>(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_agent_info_t>(HSA_AMD_AGENT_INFO_CHIP_ID), &id);
Hsa::agent_get_info(agent, static_cast<hsa_agent_info_t>(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<const hsa_amd_ipc_memory_t*>(
reinterpret_cast<const amd::IpcBuffer*>(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<hsa_amd_ipc_memory_t*>(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<size_t>(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<Image*>(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.
projects/clr/rocclr/device/rocm/rocprintf.cpp
+2 -2
Dosyayı Görüntüle
@@ -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_,"
projects/clr/rocclr/device/rocm/rocprogram.cpp
+13 -15
Dosyayı Görüntüle
@@ -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 <string>
#include <vector>
#include <cstring>
@@ -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);
projects/clr/rocclr/device/rocm/rocrctx.cpp
+143
Dosyayı Görüntüle
@@ -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
projects/clr/rocclr/device/rocm/rocrctx.hpp
+571
Dosyayı Görüntüle
@@ -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 <mutex>
#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<decltype(NAME)*>(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<t_##NAME>(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
projects/clr/rocclr/device/rocm/rocsignal.cpp
+8 -6
Dosyayı Görüntüle
@@ -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<hsa_signal_condition_t>(c), value, timeout,
static_cast<hsa_wait_state_t>(ws_));
return Hsa::signal_wait_scacquire(signal_, static_cast<hsa_signal_condition_t>(c), value, timeout,
static_cast<hsa_wait_state_t>(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
projects/clr/rocclr/device/rocm/rocsignal.hpp
+2 -3
Dosyayı Görüntüle
@@ -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 {
projects/clr/rocclr/device/rocm/rocurilocator.cpp
+4 -4
Dosyayı Görüntüle
@@ -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) {
projects/clr/rocclr/device/rocm/rocurilocator.hpp
+1 -2
Dosyayı Görüntüle
@@ -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 <vector>
namespace amd::roc {
projects/clr/rocclr/device/rocm/rocvirtual.cpp
+48 -50
Dosyayı Görüntüle
@@ -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 <fstream>
#include <limits>
@@ -146,10 +143,10 @@ void Timestamp::checkGpuTime() {
amd_signal_t* amdSignal = reinterpret_cast<amd_signal_t*>(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<ProfilingSignal> 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<hsa_signal_t>& 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<hsa_kernel_dispatch_packet_t*>(packet)->kernarg_address,
reinterpret_cast<hsa_kernel_dispatch_packet_t*>(packet)->completion_signal,
reinterpret_cast<hsa_kernel_dispatch_packet_t*>(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<AqlPacket*>& 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<AqlPacket*>& 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<AqlPacket*>& pa
reinterpret_cast<hsa_kernel_dispatch_packet_t*>(packet)->kernarg_address,
reinterpret_cast<hsa_kernel_dispatch_packet_t*>(packet)->completion_signal,
reinterpret_cast<hsa_kernel_dispatch_packet_t*>(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<AqlPacket*>& pa
packet_store_release(reinterpret_cast<uint32_t*>(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<hsa_barrier_and_packet_t*>(gpu_queue_->base_address))[index & queueMask];
*aql_loc = barrier_packet_;
packet_store_release(reinterpret_cast<uint32_t*>(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<hsa_amd_barrier_value_packet_t*>(
gpu_queue_->base_address))[index & queueMask];
*aql_loc = barrier_value_packet_;
packet_store_release(reinterpret_cast<uint32_t*>(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<void*>(cmd.dev_ptr()), cmd.count(), agent,
wait_events.size(), wait_events.data(), active);
Hsa::svm_prefetch_async(const_cast<void*>(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<uint>::max(), std::numeric_limits<uint>::max(),
&schedulerQueue_)) {
Hsa::queue_create(gpu_device(), 2048, HSA_QUEUE_TYPE_MULTI, callbackQueue, &roc_device_,
std::numeric_limits<uint>::max(), std::numeric_limits<uint>::max(),
&schedulerQueue_)) {
break;
}
@@ -3283,7 +3281,7 @@ bool VirtualGPU::createSchedulerParam() {
}
if (nullptr != schedulerQueue_) {
hsa_queue_destroy(schedulerQueue_);
Hsa::queue_destroy(schedulerQueue_);
schedulerQueue_ = nullptr;
}
projects/clr/rocclr/device/rocm/rocvirtual.hpp
+6 -10
Dosyayı Görüntüle
@@ -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;