AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

rocr: Ensure globals are initialized at first use

Browse Source 
When ROCr is built as a static library, global variables
were often not initialized to valid values at their first
use. This change addresses that problem.

Change-Id: I550fa41feb3bc04b9cc686bcfb4acf2a7b651a88
This commit is contained in:
Chris Freehill
2024-09-24 11:23:55 -05:00
committed by Chris Freehill
parent 80da7d5ee4
commit 9b13bcd0ac
36 changed files with 323 additions and 233 deletions
Download Patch File Download Diff File Expand all files Collapse all files
runtime/hsa-runtime/core/common/shared.cpp
+1 -2
View File
@@ -44,7 +44,6 @@
namespace rocr {
namespace core {
std::function<void*(size_t, size_t, uint32_t, int)> BaseShared::allocate_ = nullptr;
std::function<void(void*)> BaseShared::free_ = nullptr;
} // namespace core
} // namespace rocr
runtime/hsa-runtime/core/common/shared.h
+38 -22
View File
@@ -58,25 +58,36 @@ namespace core {
class BaseShared {
public:
static void SetAllocateAndFree(
const std::function<void*(size_t, size_t, uint32_t, int)>& allocate,
const std::function<void(void*)>& free) {
allocate_ = allocate;
free_ = free;
const std::function<void*(size_t, size_t, uint32_t, int)>& alloc,
const std::function<void(void*)>& fr) {
allocate_() = alloc;
free_() = fr;
}
protected:
static std::function<void*(size_t, size_t, uint32_t, int)> allocate_;
static std::function<void(void*)> free_;
static __forceinline std::function<void*(size_t, size_t, uint32_t, int)>&
allocate_() {
static std::function<void*(size_t, size_t, uint32_t, int)> alloc =
nullptr;
return alloc;
}
static __forceinline std::function<void(void*)>&
free_() {
static std::function<void(void*)> fr = nullptr;
return fr;
}
};
/// @brief Default Allocator for Shared. Ensures allocations are whole pages.
template <typename T> class PageAllocator : private BaseShared {
public:
__forceinline static T* alloc(int flags = 0) {
T* ret = reinterpret_cast<T*>(allocate_(AlignUp(sizeof(T), 4096), 4096, flags, 0));
T* ret = reinterpret_cast<T*>(
allocate_()(AlignUp(sizeof(T), 4096), 4096, flags, 0));
if (ret == nullptr) throw std::bad_alloc();
MAKE_NAMED_SCOPE_GUARD(throwGuard, [&]() { free_(ret); });
MAKE_NAMED_SCOPE_GUARD(throwGuard, [&]() { free_()(ret); });
new (ret) T;
@@ -85,10 +96,11 @@ template <typename T> class PageAllocator : private BaseShared {
}
__forceinline static T* alloc(int agent_node_id, int flags) {
T* ret = reinterpret_cast<T*>(allocate_(AlignUp(sizeof(T), 4096), 4096, flags, agent_node_id));
T* ret = reinterpret_cast<T*>(
allocate_()(AlignUp(sizeof(T), 4096), 4096, flags, agent_node_id));
if (ret == nullptr) throw std::bad_alloc();
MAKE_NAMED_SCOPE_GUARD(throwGuard, [&]() { free_(ret); });
MAKE_NAMED_SCOPE_GUARD(throwGuard, [&]() { free_()(ret); });
new (ret) T;
@@ -99,7 +111,7 @@ template <typename T> class PageAllocator : private BaseShared {
__forceinline static void free(T* ptr) {
if (ptr != nullptr) {
ptr->~T();
free_(ptr);
free_()(ptr);
}
}
};
@@ -110,7 +122,7 @@ template <typename T, typename Allocator = PageAllocator<T>>
class Shared final : private BaseShared {
public:
explicit Shared(Allocator* pool = nullptr, int flags = 0) : pool_(pool) {
assert(allocate_ != nullptr && free_ != nullptr &&
assert(allocate_() != nullptr && free_() != nullptr &&
"Shared object allocator is not set");
if (pool_)
@@ -120,7 +132,7 @@ class Shared final : private BaseShared {
}
explicit Shared(int agent_node_id, Allocator* pool = nullptr, int flags = 0) : pool_(pool) {
assert(allocate_ != nullptr && free_ != nullptr &&
assert(allocate_() != nullptr && free_() != nullptr &&
"Shared object allocator is not set");
if (pool_)
@@ -130,7 +142,8 @@ class Shared final : private BaseShared {
}
~Shared() {
assert(allocate_ != nullptr && free_ != nullptr && "Shared object allocator is not set");
assert(allocate_() != nullptr && free_() != nullptr &&
"Shared object allocator is not set");
if (pool_)
pool_->free(shared_object_);
@@ -164,19 +177,20 @@ class Shared final : private BaseShared {
template <typename T> class Shared<T, PageAllocator<T>> final : private BaseShared {
public:
Shared(int flags = 0) {
assert(allocate_ != nullptr && free_ != nullptr && "Shared object allocator is not set");
assert(allocate_() != nullptr && free_() != nullptr &&
"Shared object allocator is not set");
shared_object_ = PageAllocator<T>::alloc(flags);
}
Shared(int agent_node_id, int flags) {
assert(allocate_ != nullptr && free_ != nullptr && "Shared object allocator is not set");
assert(allocate_() != nullptr && free_() != nullptr && "Shared object allocator is not set");
shared_object_ = PageAllocator<T>::alloc(agent_node_id, flags);
}
~Shared() {
assert(allocate_ != nullptr && free_ != nullptr &&
assert(allocate_() != nullptr && free_() != nullptr &&
"Shared object allocator is not set");
PageAllocator<T>::free(shared_object_);
@@ -207,18 +221,19 @@ template <typename T, size_t Align> class SharedArray final : private BaseShared
SharedArray() : shared_object_(nullptr) {}
explicit SharedArray(size_t length) : shared_object_(nullptr), len(length) {
assert(allocate_ != nullptr && free_ != nullptr && "Shared object allocator is not set");
assert(allocate_() != nullptr && free_() != nullptr &&
"Shared object allocator is not set");
static_assert((__alignof(T) <= Align) || (Align == 0), "Align is less than alignof(T)");
shared_object_ =
reinterpret_cast<T*>(allocate_(sizeof(T) * length, Max(__alignof(T), Align), 0, 0));
reinterpret_cast<T*>(allocate_()(sizeof(T) * length, Max(__alignof(T), Align), 0, 0));
if (shared_object_ == nullptr) throw std::bad_alloc();
size_t i = 0;
MAKE_NAMED_SCOPE_GUARD(loopGuard, [&]() {
for (size_t t = 0; t < i - 1; t++) shared_object_[t].~T();
free_(shared_object_);
free_()(shared_object_);
});
for (; i < length; i++) new (&shared_object_[i]) T;
@@ -227,11 +242,12 @@ template <typename T, size_t Align> class SharedArray final : private BaseShared
}
~SharedArray() {
assert(allocate_ != nullptr && free_ != nullptr && "Shared object allocator is not set");
assert(allocate_() != nullptr && free_() != nullptr &&
"Shared object allocator is not set");
if (shared_object_ != nullptr) {
for (size_t i = 0; i < len; i++) shared_object_[i].~T();
free_(shared_object_);
free_()(shared_object_);
}
}
runtime/hsa-runtime/core/inc/amd_aie_aql_queue.h
+8 -4
View File
@@ -62,11 +62,11 @@ class AieAqlQueue : public core::Queue,
core::DoorbellSignal {
public:
static __forceinline bool IsType(core::Signal *signal) {
return signal->IsType(&rtti_id_);
return signal->IsType(&rtti_id());
}
static __forceinline bool IsType(core::Queue *queue) {
return queue->IsType(&rtti_id_);
return queue->IsType(&rtti_id());
}
AieAqlQueue() = delete;
@@ -123,7 +123,7 @@ public:
uint32_t queue_size_bytes_ = std::numeric_limits<uint32_t>::max();
protected:
bool _IsA(Queue::rtti_t id) const override { return id == &rtti_id_; }
bool _IsA(Queue::rtti_t id) const override { return id == &rtti_id(); }
private:
AieAgent &agent_;
@@ -144,7 +144,11 @@ private:
/// @brief Indicates if queue is active.
std::atomic<bool> active_;
static int rtti_id_;
static __forceinline int& rtti_id() {
static int rtti_id_ = 0;
return rtti_id_;
}
};
} // namespace AMD
runtime/hsa-runtime/core/inc/amd_aql_queue.h
+19 -9
View File
@@ -57,10 +57,10 @@ namespace AMD {
class AqlQueue : public core::Queue, private core::LocalSignal, public core::DoorbellSignal {
public:
static __forceinline bool IsType(core::Signal* signal) {
return signal->IsType(&rtti_id_);
return signal->IsType(&rtti_id());
}
static __forceinline bool IsType(core::Queue* queue) { return queue->IsType(&rtti_id_); }
static __forceinline bool IsType(core::Queue* queue) { return queue->IsType(&rtti_id()); }
// Acquires/releases queue resources and requests HW schedule/deschedule.
AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id,
@@ -228,7 +228,7 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Doo
void AsyncReclaimAltScratch();
protected:
bool _IsA(Queue::rtti_t id) const override { return id == &rtti_id_; }
bool _IsA(Queue::rtti_t id) const override { return id == &rtti_id(); }
private:
uint32_t ComputeRingBufferMinPkts();
@@ -331,18 +331,28 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Doo
std::vector<uint32_t> cu_mask_;
// Shared event used for queue errors
static HsaEvent* queue_event_;
static __forceinline HsaEvent*& queue_event() {
static HsaEvent* queue_event_ = nullptr;
return queue_event_;
}
// Queue count - used to ref count queue_event_
static std::atomic<uint32_t> queue_count_;
static __forceinline std::atomic<uint32_t>& queue_count() {
static std::atomic<uint32_t> queue_count_(0);
return queue_count_;
}
// Mutex for queue_event_ manipulation
static KernelMutex queue_lock_;
static __forceinline KernelMutex& queue_lock() {
static KernelMutex queue_lock_;
return queue_lock_;
}
// Async scratch single limit - may be modified after init
size_t async_scratch_single_limit_;
static int rtti_id_;
static __forceinline int& rtti_id() {
static int rtti_id_ = 0;
return rtti_id_;
}
// Forbid copying and moving of this object
DISALLOW_COPY_AND_ASSIGN(AqlQueue);
runtime/hsa-runtime/core/inc/amd_blit_kernel.h
+1
View File
@@ -117,6 +117,7 @@ class BlitKernel : public core::Blit {
virtual void GangLeader(bool gang_leader) override {}
virtual bool GangLeader() const override { return false; }
const uint16_t kInvalidPacketHeader = HSA_PACKET_TYPE_INVALID;
private:
union KernelArgs {
struct __ALIGNED__(16) {
runtime/hsa-runtime/core/inc/amd_memory_region.h
+5 -2
View File
@@ -167,7 +167,7 @@ class MemoryRegion : public core::MemoryRegion {
return static_cast<uint32_t>(mem_props_.MemoryClockMax);
}
__forceinline size_t GetPageSize() const { return kPageSize_; }
__forceinline size_t GetPageSize() const { return kPageSize(); }
__forceinline const HsaMemFlags &mem_flags() const { return mem_flag_; }
__forceinline const HsaMemMapFlags &map_flags() const { return map_flag_; }
@@ -195,7 +195,10 @@ private:
// fragments of the block routing to the same MemoryRegion.
mutable KernelMutex access_lock_;
static size_t kPageSize_;
static __forceinline const size_t& kPageSize() {
static size_t kPageSize_ = sysconf(_SC_PAGESIZE);
return kPageSize_;
}
// Determine access type allowed to requesting device
hsa_amd_memory_pool_access_t GetAccessInfo(const core::Agent& agent,
runtime/hsa-runtime/core/inc/default_signal.h
+12 -6
View File
@@ -59,7 +59,7 @@ class BusyWaitSignal : public Signal {
/// @brief Determines if a Signal* can be safely converted to BusyWaitSignal*
/// via static_cast.
static __forceinline bool IsType(Signal* ptr) {
return ptr->IsType(&rtti_id_);
return ptr->IsType(&rtti_id());
}
/// @brief See base class Signal.
@@ -154,10 +154,13 @@ class BusyWaitSignal : public Signal {
__forceinline HsaEvent* EopEvent() { return NULL; }
protected:
bool _IsA(rtti_t id) const { return id == &rtti_id_; }
bool _IsA(rtti_t id) const { return id == &rtti_id(); }
private:
static int rtti_id_;
static __forceinline int& rtti_id() {
static int rtti_id_ = 0;
return rtti_id_;
}
DISALLOW_COPY_AND_ASSIGN(BusyWaitSignal);
};
@@ -167,7 +170,7 @@ class DefaultSignal : private LocalSignal, public BusyWaitSignal {
public:
/// @brief Determines if a Signal* can be safely converted to BusyWaitSignal*
/// via static_cast.
static __forceinline bool IsType(Signal* ptr) { return ptr->IsType(&rtti_id_); }
static __forceinline bool IsType(Signal* ptr) { return ptr->IsType(&rtti_id()); }
/// @brief See base class Signal.
explicit DefaultSignal(hsa_signal_value_t initial_value, bool enableIPC = false)
@@ -175,12 +178,15 @@ class DefaultSignal : private LocalSignal, public BusyWaitSignal {
protected:
bool _IsA(rtti_t id) const {
if (id == &rtti_id_) return true;
if (id == &rtti_id()) return true;
return BusyWaitSignal::_IsA(id);
}
private:
static int rtti_id_;
static __forceinline int& rtti_id() {
static int rtti_id_ = 0;
return rtti_id_;
}
DISALLOW_COPY_AND_ASSIGN(DefaultSignal);
};
runtime/hsa-runtime/core/inc/host_queue.h
+10 -4
View File
@@ -52,7 +52,7 @@ namespace rocr {
namespace core {
class HostQueue : public Queue {
public:
static __forceinline bool IsType(core::Queue* queue) { return queue->IsType(&rtti_id_); }
static __forceinline bool IsType(core::Queue* queue) { return queue->IsType(&rtti_id()); }
HostQueue(hsa_region_t region, uint32_t ring_size, hsa_queue_type32_t type,
uint32_t features, hsa_signal_t doorbell_signal);
@@ -175,17 +175,23 @@ class HostQueue : public Queue {
void operator delete(void*, void*) {}
protected:
bool _IsA(Queue::rtti_t id) const override { return id == &rtti_id_; }
bool _IsA(Queue::rtti_t id) const override { return id == &rtti_id(); }
private:
static int rtti_id_;
static __forceinline int& rtti_id() {
static int rtti_id_ = 0;
return rtti_id_;
}
static const size_t kRingAlignment = 256;
const uint32_t size_;
void* ring_;
// Host queue id counter, starting from 0x80000000 to avoid overlaping
// with aql queue id.
static std::atomic<uint32_t> queue_count_;
static __forceinline std::atomic<uint32_t>& queue_count() {
static std::atomic<uint32_t> queue_count_;
return queue_count_;
}
DISALLOW_COPY_AND_ASSIGN(HostQueue);
};
runtime/hsa-runtime/core/inc/hsa_amd_tool_int.hpp
+6 -6
View File
@@ -32,7 +32,7 @@ __forceinline void notify_event_scratch_async_reclaim_end(const hsa_queue_t* que
__forceinline void notify_event_scratch_alloc_start(const hsa_queue_t* queue,
scratch_alloc_flag flags,
uint64_t dispatch_id) {
const auto& tool_table = core::hsa_api_table_.tools_api;
const auto& tool_table = core::hsa_api_table().tools_api;
if (!tool_table.hsa_amd_tool_scratch_event_alloc_start_fn) {
return;
}
@@ -49,7 +49,7 @@ __forceinline void notify_event_scratch_alloc_start(const hsa_queue_t* queue,
__forceinline void notify_event_scratch_alloc_end(const hsa_queue_t* queue,
scratch_alloc_flag flags, uint64_t dispatch_id,
size_t size, size_t num_slots) {
const auto& tool_table = core::hsa_api_table_.tools_api;
const auto& tool_table = core::hsa_api_table().tools_api;
if (!tool_table.hsa_amd_tool_scratch_event_alloc_end_fn) {
return;
}
@@ -69,7 +69,7 @@ __forceinline void notify_event_scratch_alloc_end(const hsa_queue_t* queue,
__forceinline void notify_event_scratch_free_start(const hsa_queue_t* queue,
scratch_alloc_flag flags) {
const auto& tool_table = core::hsa_api_table_.tools_api;
const auto& tool_table = core::hsa_api_table().tools_api;
if (!tool_table.hsa_amd_tool_scratch_event_free_start_fn) {
return;
}
@@ -86,7 +86,7 @@ __forceinline void notify_event_scratch_free_start(const hsa_queue_t* queue,
__forceinline void notify_event_scratch_free_end(const hsa_queue_t* queue,
scratch_alloc_flag flags) {
const auto& tool_table = core::hsa_api_table_.tools_api;
const auto& tool_table = core::hsa_api_table().tools_api;
if (!tool_table.hsa_amd_tool_scratch_event_free_end_fn) {
return;
}
@@ -103,7 +103,7 @@ __forceinline void notify_event_scratch_free_end(const hsa_queue_t* queue,
__forceinline void notify_event_scratch_async_reclaim_start(const hsa_queue_t* queue,
scratch_alloc_flag flags) {
const auto& tool_table = core::hsa_api_table_.tools_api;
const auto& tool_table = core::hsa_api_table().tools_api;
if (!tool_table.hsa_amd_tool_scratch_event_async_reclaim_start_fn) {
return;
}
@@ -120,7 +120,7 @@ __forceinline void notify_event_scratch_async_reclaim_start(const hsa_queue_t* q
__forceinline void notify_event_scratch_async_reclaim_end(const hsa_queue_t* queue,
scratch_alloc_flag flags) {
const auto& tool_table = core::hsa_api_table_.tools_api;
const auto& tool_table = core::hsa_api_table().tools_api;
if (!tool_table.hsa_amd_tool_scratch_event_async_reclaim_end_fn) {
return;
}
runtime/hsa-runtime/core/inc/hsa_api_trace_int.h
+2 -2
View File
@@ -73,8 +73,8 @@ namespace core {
void Reset();
};
extern HsaApiTable hsa_api_table_;
extern HsaApiTable hsa_internal_api_table_;
extern HsaApiTable& hsa_api_table();
extern HsaApiTable& hsa_internal_api_table();
void LoadInitialHsaApiTable();
} // namespace core
runtime/hsa-runtime/core/inc/hsa_table_interface.h
+5
View File
@@ -40,8 +40,13 @@
//
////////////////////////////////////////////////////////////////////////////////
#ifndef RUNTIME_HSA_RUNTIME_CORE_INC_HSA_TABLE_INTERFACE_H_
#define RUNTIME_HSA_RUNTIME_CORE_INC_HSA_TABLE_INTERFACE_H_
#include "inc/hsa_api_trace.h"
void hsa_table_interface_init(const HsaApiTable* apiTable);
const HsaApiTable* hsa_table_interface_get_table();
#endif // RUNTIME_HSA_RUNTIME_CORE_INC_HSA_TABLE_INTERFACE_H_
runtime/hsa-runtime/core/inc/intercept_queue.h
+7 -4
View File
@@ -272,14 +272,17 @@ class InterceptQueue : public QueueProxy, private LocalSignal, public DoorbellSi
/// @brief Provide information about the queue
hsa_status_t GetInfo(hsa_queue_info_attribute_t attribute, void* value) override;
static __forceinline bool IsType(core::Signal* signal) { return signal->IsType(&rtti_id_); }
static __forceinline bool IsType(core::Queue* queue) { return queue->IsType(&rtti_id_); }
static __forceinline bool IsType(core::Signal* signal) { return signal->IsType(&rtti_id()); }
static __forceinline bool IsType(core::Queue* queue) { return queue->IsType(&rtti_id()); }
protected:
bool _IsA(Queue::rtti_t id) const override { return id == &rtti_id_; }
bool _IsA(Queue::rtti_t id) const override { return id == &rtti_id(); }
private:
static int rtti_id_;
static __forceinline int& rtti_id() {
static int rtti_id_ = 0;
return rtti_id_;
}
};
} // namespace core
runtime/hsa-runtime/core/inc/interrupt_signal.h
+6 -3
View File
@@ -94,7 +94,7 @@ class InterruptSignal : private LocalSignal, public Signal {
/// @brief Determines if a Signal* can be safely converted to an
/// InterruptSignal* via static_cast.
static __forceinline bool IsType(Signal* ptr) {
return ptr->IsType(&rtti_id_);
return ptr->IsType(&rtti_id());
}
explicit InterruptSignal(hsa_signal_value_t initial_value,
@@ -191,7 +191,7 @@ class InterruptSignal : private LocalSignal, public Signal {
__forceinline HsaEvent* EopEvent() { return event_; }
protected:
bool _IsA(rtti_t id) const { return id == &rtti_id_; }
bool _IsA(rtti_t id) const { return id == &rtti_id(); }
private:
/// @variable KFD event on which the interrupt signal is based on.
@@ -202,7 +202,10 @@ class InterruptSignal : private LocalSignal, public Signal {
bool free_event_;
/// Used to obtain a globally unique value (address) for rtti.
static int rtti_id_;
static __forceinline int& rtti_id() {
static int rtti_id_ = 0;
return rtti_id_;
}
/// @brief Notify driver of signal value change if necessary.
__forceinline void SetEvent() {
runtime/hsa-runtime/core/inc/ipc_signal.h
+6 -3
View File
@@ -90,16 +90,19 @@ class IPCSignal : private SharedMemorySignal, public BusyWaitSignal {
/// @brief Determines if a Signal* can be safely converted to BusyWaitSignal*
/// via static_cast.
static __forceinline bool IsType(Signal* ptr) { return ptr->IsType(&rtti_id_); }
static __forceinline bool IsType(Signal* ptr) { return ptr->IsType(&rtti_id()); }
protected:
bool _IsA(rtti_t id) const {
if (id == &rtti_id_) return true;
if (id == &rtti_id()) return true;
return BusyWaitSignal::_IsA(id);
}
private:
static int rtti_id_;
static __forceinline int& rtti_id() {
static int rtti_id_ = 0;
return rtti_id_;
}
static KernelMutex lock_;
explicit IPCSignal(SharedMemorySignal&& abi_block)
runtime/hsa-runtime/core/inc/isa.h
+6 -9
View File
@@ -101,7 +101,7 @@ class Isa final: public amd::hsa::common::Signed<0xB13594F2BD8F212D> {
typedef std::tuple<int32_t, int32_t, int32_t> Version;
/// @brief Default destructor.
~Isa() {}
~Isa() = default;
/// @returns Handle equivalent of @p isa_object.
static hsa_isa_t Handle(const Isa *isa_object) {
@@ -228,17 +228,14 @@ class IsaRegistry final {
/// @brief IsaRegistry's map type.
typedef std::unordered_map<std::string, Isa> IsaMap;
/// @brief Supported instruction set architectures.
static const IsaMap supported_isas_;
/// @brief Default constructor
IsaRegistry() = delete;
/// @brief Default constructor - not available.
IsaRegistry();
/// @brief Default destructor - not available.
~IsaRegistry();
/// @brief Default destructor
~IsaRegistry() = default;
/// @returns Supported instruction set architectures.
static const IsaMap GetSupportedIsas();
static const IsaMap& GetSupportedIsas();
}; // class IsaRegistry
} // namespace core
runtime/hsa-runtime/core/inc/runtime.h
+4 -2
View File
@@ -583,8 +583,10 @@ class Runtime {
// Will be created before any user could call hsa_init but also could be
// destroyed before incorrectly written programs call hsa_shutdown.
static KernelMutex bootstrap_lock_;
static __forceinline KernelMutex& bootstrap_lock() {
static KernelMutex bootstrap_lock_;
return bootstrap_lock_;
}
Runtime();
Runtime(const Runtime&);
runtime/hsa-runtime/core/runtime/amd_aie_aql_queue.cpp
-2
View File
@@ -66,8 +66,6 @@
namespace rocr {
namespace AMD {
int AieAqlQueue::rtti_id_ = 0;
AieAqlQueue::AieAqlQueue(AieAgent *agent, size_t req_size_pkts,
uint32_t node_id)
: Queue(0, 0), LocalSignal(0, false), DoorbellSignal(signal()),
runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp
+19 -24
View File
@@ -73,11 +73,6 @@
namespace rocr {
namespace AMD {
HsaEvent* AqlQueue::queue_event_ = nullptr;
std::atomic<uint32_t> AqlQueue::queue_count_(0);
KernelMutex AqlQueue::queue_lock_;
int AqlQueue::rtti_id_ = 0;
AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, ScratchInfo& scratch,
core::HsaEventCallback callback, void* err_data, bool is_kv)
: Queue(agent->node_id(), agent->isMES() ? (MemoryRegion::AllocateGTTAccess | MemoryRegion::AllocateNonPaged) : 0),
@@ -236,11 +231,11 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr
}
MAKE_NAMED_SCOPE_GUARD(EventGuard, [&]() {
ScopedAcquire<KernelMutex> _lock(&queue_lock_);
queue_count_--;
if (queue_count_ == 0) {
core::InterruptSignal::DestroyEvent(queue_event_);
queue_event_ = nullptr;
ScopedAcquire<KernelMutex> _lock(&queue_lock());
queue_count()--;
if (queue_count() == 0) {
core::InterruptSignal::DestroyEvent(queue_event());
queue_event() = nullptr;
}
});
@@ -251,20 +246,20 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr
});
if (core::g_use_interrupt_wait) {
ScopedAcquire<KernelMutex> _lock(&queue_lock_);
queue_count_++;
if (queue_event_ == nullptr) {
assert(queue_count_ == 1 && "Inconsistency in queue event reference counting found.\n");
ScopedAcquire<KernelMutex> _lock(&queue_lock());
queue_count()++;
if (queue_event() == nullptr) {
assert(queue_count() == 1 && "Inconsistency in queue event reference counting found.\n");
queue_event_ = core::InterruptSignal::CreateEvent(HSA_EVENTTYPE_SIGNAL, false);
if (queue_event_ == nullptr)
queue_event() = core::InterruptSignal::CreateEvent(HSA_EVENTTYPE_SIGNAL, false);
if (queue_event() == nullptr)
throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES,
"Queue event creation failed.\n");
}
auto Signal = new core::InterruptSignal(0, queue_event_);
auto Signal = new core::InterruptSignal(0, queue_event());
assert(Signal != nullptr && "Should have thrown!\n");
amd_queue_.queue_inactive_signal = core::InterruptSignal::Convert(Signal);
exception_signal_ = new core::InterruptSignal(0, queue_event_);
exception_signal_ = new core::InterruptSignal(0, queue_event());
assert(exception_signal_ != nullptr && "Should have thrown!\n");
} else {
EventGuard.Dismiss();
@@ -284,7 +279,7 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr
if (core::Runtime::runtime_singleton_->KfdVersion().supports_exception_debugging) {
queue_rsrc.ErrorReason = &exception_signal_->signal_.value;
kmt_status = hsaKmtCreateQueueExt(node_id, HSA_QUEUE_COMPUTE_AQL, 100, priority_, 0, ring_buf_,
ring_buf_alloc_bytes_, queue_event_, &queue_rsrc);
ring_buf_alloc_bytes_, queue_event(), &queue_rsrc);
} else {
kmt_status = hsaKmtCreateQueueExt(node_id, HSA_QUEUE_COMPUTE_AQL, 100, priority_, 0, ring_buf_,
ring_buf_alloc_bytes_, NULL, &queue_rsrc);
@@ -382,11 +377,11 @@ AqlQueue::~AqlQueue() {
exception_signal_->DestroySignal();
HSA::hsa_signal_destroy(amd_queue_.queue_inactive_signal);
if (core::g_use_interrupt_wait) {
ScopedAcquire<KernelMutex> lock(&queue_lock_);
queue_count_--;
if (queue_count_ == 0) {
core::InterruptSignal::DestroyEvent(queue_event_);
queue_event_ = nullptr;
ScopedAcquire<KernelMutex> lock(&queue_lock());
queue_count()--;
if (queue_count() == 0) {
core::InterruptSignal::DestroyEvent(queue_event());
queue_event() = nullptr;
}
}
agent_->system_deallocator()(pm4_ib_buf_);
runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp
+29 -13
View File
@@ -54,7 +54,8 @@ namespace rocr {
namespace AMD {
static const uint16_t kInvalidPacketHeader = HSA_PACKET_TYPE_INVALID;
static std::string kBlitKernelSource(R"(
static std::string& kBlitKernelSource() {
static std::string kBlitKernelSource_(R"(
// Compatibility function for GFXIP 7.
function s_load_dword_offset(byte_offset)
@@ -874,29 +875,44 @@ L_FILL_PHASE_2_DONE:
s_endpgm
end
)");
return kBlitKernelSource_;
}
// Search kernel source for variable definition and return value.
int GetKernelSourceParam(const char* paramName) {
std::stringstream paramDef;
paramDef << "var " << paramName << " = ";
std::string::size_type paramDefLoc = kBlitKernelSource.find(paramDef.str());
std::string::size_type paramDefLoc =
kBlitKernelSource().find(paramDef.str());
assert(paramDefLoc != std::string::npos);
std::string::size_type paramValLoc = paramDefLoc + paramDef.str().size();
std::string::size_type paramEndLoc =
kBlitKernelSource.find('\n', paramDefLoc);
kBlitKernelSource().find('\n', paramDefLoc);
assert(paramDefLoc != std::string::npos);
std::string paramVal(&kBlitKernelSource[paramValLoc],
&kBlitKernelSource[paramEndLoc]);
std::string paramVal(&kBlitKernelSource()[paramValLoc],
&kBlitKernelSource()[paramEndLoc]);
return std::stoi(paramVal);
}
static int kCopyAlignedVecWidth = GetKernelSourceParam("kCopyAlignedVecWidth");
static int kCopyAlignedUnroll = GetKernelSourceParam("kCopyAlignedUnroll");
static int kCopyMisalignedUnroll = GetKernelSourceParam("kCopyMisalignedUnroll");
static int kFillVecWidth = GetKernelSourceParam("kFillVecWidth");
static int kFillUnroll = GetKernelSourceParam("kFillUnroll");
#define DEFINE_KERNEL_PARAM_FUNC(name) \
static int& name() { \
static std::once_flag initFlag; \
static int val; \
std::call_once(initFlag, [&]() { \
val = GetKernelSourceParam(#name); \
}); \
return val; \
}
// Use the macro to define the functions
DEFINE_KERNEL_PARAM_FUNC(kCopyAlignedVecWidth)
DEFINE_KERNEL_PARAM_FUNC(kCopyAlignedUnroll)
DEFINE_KERNEL_PARAM_FUNC(kCopyMisalignedUnroll)
DEFINE_KERNEL_PARAM_FUNC(kFillVecWidth)
DEFINE_KERNEL_PARAM_FUNC(kFillUnroll)
static unsigned extractAqlBits(unsigned v, unsigned pos, unsigned width) {
return (v >> pos) & ((1 << width) - 1);
@@ -1093,7 +1109,7 @@ hsa_status_t BlitKernel::SubmitLinearCopyCommand(
// Phase 2 (unrolled dwordx4 copy) ends when last whole block fits.
uint64_t phase2_block = num_workitems * sizeof(uint32_t) *
kCopyAlignedUnroll * kCopyAlignedVecWidth;
kCopyAlignedUnroll() * kCopyAlignedVecWidth();
uint64_t phase2_size = ((size - phase1_size) / phase2_block) * phase2_block;
// Phase 3 (dword copy) ends when last whole dword fits.
@@ -1125,7 +1141,7 @@ hsa_status_t BlitKernel::SubmitLinearCopyCommand(
uintptr_t src_start = uintptr_t(src);
uintptr_t dst_start = uintptr_t(dst);
uint64_t phase1_block =
num_workitems * sizeof(uint8_t) * kCopyMisalignedUnroll;
num_workitems * sizeof(uint8_t) * kCopyMisalignedUnroll();
uint64_t phase1_size = (size / phase1_block) * phase1_block;
args->copy_misaligned.phase1_src_start = src_start;
@@ -1164,7 +1180,7 @@ hsa_status_t BlitKernel::SubmitLinearFillCommand(void* ptr, uint32_t value,
uint64_t fill_size = count * sizeof(uint32_t);
uint64_t phase1_block =
num_workitems * sizeof(uint32_t) * kFillUnroll * kFillVecWidth;
num_workitems * sizeof(uint32_t) * kFillUnroll() * kFillVecWidth();
uint64_t phase1_size = (fill_size / phase1_block) * phase1_block;
KernelArgs* args = ObtainAsyncKernelCopyArg();
runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp
+4 -6
View File
@@ -70,6 +70,7 @@
#include "core/inc/amd_trap_handler_v1.h"
#include "core/inc/amd_blit_shaders.h"
#include "core/inc/hsa_api_trace_int.h"
// Generated header
#include "amd_trap_handler_v2.h"
#include "amd_blit_shaders_v2.h"
@@ -88,9 +89,6 @@
#define DEFAULT_SCRATCH_SINGLE_LIMIT_ASYNC_PER_XCC (1 << 30) // 1 GB
namespace rocr {
namespace core {
extern HsaApiTable hsa_internal_api_table_;
} // namespace core
namespace AMD {
const uint64_t CP_DMA_DATA_TRANSFER_CNT_MAX = (1 << 26);
@@ -1397,15 +1395,15 @@ hsa_status_t GpuAgent::GetInfo(hsa_agent_info_t attribute, void* value) const {
((uint8_t*)value)[index] |= 1 << subBit;
};
if (core::hsa_internal_api_table_.finalizer_api.hsa_ext_program_finalize_fn != NULL) {
if (core::hsa_internal_api_table().finalizer_api.hsa_ext_program_finalize_fn != NULL) {
setFlag(HSA_EXTENSION_FINALIZER);
}
if (core::hsa_internal_api_table_.image_api.hsa_ext_image_create_fn != NULL) {
if (core::hsa_internal_api_table().image_api.hsa_ext_image_create_fn != NULL) {
setFlag(HSA_EXTENSION_IMAGES);
}
if (core::hsa_internal_api_table_.pcs_api.hsa_ven_amd_pcs_iterate_configuration_fn != NULL) {
if (core::hsa_internal_api_table().pcs_api.hsa_ven_amd_pcs_iterate_configuration_fn != NULL) {
setFlag(HSA_EXTENSION_AMD_PC_SAMPLING);
}
runtime/hsa-runtime/core/runtime/amd_memory_region.cpp
+8 -9
View File
@@ -57,7 +57,6 @@ namespace AMD {
// Tracks aggregate size of system memory available on platform
size_t MemoryRegion::max_sysmem_alloc_size_ = 0;
size_t MemoryRegion::kPageSize_ = sysconf(_SC_PAGESIZE);
bool MemoryRegion::RegisterMemory(void* ptr, size_t size, const HsaMemFlags& MemFlags) {
assert(ptr != NULL);
@@ -125,7 +124,7 @@ MemoryRegion::MemoryRegion(bool fine_grain, bool kernarg, bool full_profile,
virtual_size_ = kGpuVmSize;
} else if (IsSystem()) {
mem_flag_.ui32.PageSize = MemoryRegion::kPageSize_;
mem_flag_.ui32.PageSize = MemoryRegion::kPageSize();
mem_flag_.ui32.NoSubstitute = 0;
mem_flag_.ui32.HostAccess = 1;
mem_flag_.ui32.CachePolicy = HSA_CACHING_CACHED;
@@ -138,7 +137,7 @@ MemoryRegion::MemoryRegion(bool fine_grain, bool kernarg, bool full_profile,
// Adjust allocatable size per page align
max_single_alloc_size_ = AlignDown(static_cast<size_t>(GetPhysicalSize()), kPageSize_);
max_single_alloc_size_ = AlignDown(static_cast<size_t>(GetPhysicalSize()), kPageSize());
// Keep track of total system memory available
// @note: System memory is surfaced as both coarse
@@ -150,7 +149,7 @@ MemoryRegion::MemoryRegion(bool fine_grain, bool kernarg, bool full_profile,
}
assert(GetVirtualSize() != 0);
assert(IsMultipleOf(max_single_alloc_size_, kPageSize_));
assert(IsMultipleOf(max_single_alloc_size_, kPageSize()));
}
MemoryRegion::~MemoryRegion() {}
@@ -177,7 +176,7 @@ hsa_status_t MemoryRegion::AllocateImpl(size_t& size, AllocateFlags alloc_flags,
return HSA_STATUS_ERROR_INVALID_ALLOCATION;
}
size = AlignUp(size, kPageSize_);
size = AlignUp(size, kPageSize());
return core::Runtime::runtime_singleton_->AgentDriver(owner()->driver_type)
.AllocateMemory(*this, alloc_flags, address, size, agent_node_id);
@@ -279,7 +278,7 @@ hsa_status_t MemoryRegion::GetInfo(hsa_region_info_t attribute,
case HSA_HEAPTYPE_DEVICE_SVM:
case HSA_HEAPTYPE_FRAME_BUFFER_PRIVATE:
case HSA_HEAPTYPE_FRAME_BUFFER_PUBLIC:
*((size_t*)value) = kPageSize_;
*((size_t*)value) = kPageSize();
break;
default:
*((size_t*)value) = 0;
@@ -292,7 +291,7 @@ hsa_status_t MemoryRegion::GetInfo(hsa_region_info_t attribute,
case HSA_HEAPTYPE_DEVICE_SVM:
case HSA_HEAPTYPE_FRAME_BUFFER_PRIVATE:
case HSA_HEAPTYPE_FRAME_BUFFER_PUBLIC:
*((size_t*)value) = kPageSize_;
*((size_t*)value) = kPageSize();
break;
default:
*((size_t*)value) = 0;
@@ -361,12 +360,12 @@ hsa_status_t MemoryRegion::GetPoolInfo(hsa_amd_memory_pool_info_t attribute,
case HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_REC_GRANULE:
switch (mem_props_.HeapType) {
case HSA_HEAPTYPE_SYSTEM:
*((size_t*)value) = kPageSize_;
*((size_t*)value) = kPageSize();
break;
case HSA_HEAPTYPE_FRAME_BUFFER_PRIVATE:
case HSA_HEAPTYPE_FRAME_BUFFER_PUBLIC:
*((size_t*)value) = core::Runtime::runtime_singleton_->flag().disable_fragment_alloc()
? kPageSize_
? kPageSize()
: fragment_allocator_.default_block_size();
break;
default:
runtime/hsa-runtime/core/runtime/default_signal.cpp
-3
View File
@@ -51,9 +51,6 @@
namespace rocr {
namespace core {
int DefaultSignal::rtti_id_ = 0;
int BusyWaitSignal::rtti_id_ = 0;
BusyWaitSignal::BusyWaitSignal(SharedSignal* abi_block, bool enableIPC)
: Signal(abi_block, enableIPC) {
signal_.kind = AMD_SIGNAL_KIND_USER;
runtime/hsa-runtime/core/runtime/host_queue.cpp
-3
View File
@@ -48,9 +48,6 @@
namespace rocr {
namespace core {
int HostQueue::rtti_id_ = 0;
std::atomic<uint32_t> HostQueue::queue_count_(0x80000000);
HostQueue::HostQueue(hsa_region_t region, uint32_t ring_size, hsa_queue_type32_t type,
uint32_t features, hsa_signal_t doorbell_signal)
: Queue(), size_(ring_size) {
runtime/hsa-runtime/core/runtime/hsa_api_trace.cpp
+10 -3
View File
@@ -64,8 +64,15 @@ hsa_status_t hsa_amd_runtime_queue_create_register(hsa_amd_runtime_queue_notifie
namespace core {
HsaApiTable hsa_api_table_;
HsaApiTable hsa_internal_api_table_;
HsaApiTable& hsa_api_table() {
static HsaApiTable table;
return table;
}
HsaApiTable& hsa_internal_api_table() {
static HsaApiTable table;
return table;
}
HsaApiTable::HsaApiTable() {
Init();
@@ -483,7 +490,7 @@ void HsaApiTable::UpdateTools() {
}
void LoadInitialHsaApiTable() {
hsa_table_interface_init(&hsa_api_table_.hsa_api);
hsa_table_interface_init(&hsa_api_table().hsa_api);
}
} // namespace core
runtime/hsa-runtime/core/runtime/hsa_ext_interface.cpp
+13 -12
View File
@@ -119,28 +119,28 @@ void ExtensionEntryPoints::InitPcSamplingExtTable() {
// Initialize Amd Ext table for Api related to Images
void ExtensionEntryPoints::InitAmdExtTable() {
hsa_api_table_.amd_ext_api.hsa_amd_image_create_fn = hsa_ext_null;
hsa_internal_api_table_.amd_ext_api.hsa_amd_image_create_fn = hsa_ext_null;
hsa_api_table().amd_ext_api.hsa_amd_image_create_fn = hsa_ext_null;
hsa_internal_api_table().amd_ext_api.hsa_amd_image_create_fn = hsa_ext_null;
}
// Update Amd Ext table for Api related to Images.
// @note: Interface should be updated when Amd Ext table
// begins hosting Api's from other extension libraries
void ExtensionEntryPoints::UpdateAmdExtTable(decltype(::hsa_amd_image_create)* func_ptr) {
assert(hsa_api_table_.amd_ext_api.hsa_amd_image_create_fn ==
assert(hsa_api_table().amd_ext_api.hsa_amd_image_create_fn ==
(decltype(hsa_amd_image_create)*)hsa_ext_null &&
"Duplicate load of extension import.");
assert(hsa_internal_api_table_.amd_ext_api.hsa_amd_image_create_fn ==
assert(hsa_internal_api_table().amd_ext_api.hsa_amd_image_create_fn ==
(decltype(hsa_amd_image_create)*)hsa_ext_null &&
"Duplicate load of extension import.");
hsa_api_table_.amd_ext_api.hsa_amd_image_create_fn = func_ptr;
hsa_internal_api_table_.amd_ext_api.hsa_amd_image_create_fn = func_ptr;
hsa_api_table().amd_ext_api.hsa_amd_image_create_fn = func_ptr;
hsa_internal_api_table().amd_ext_api.hsa_amd_image_create_fn = func_ptr;
}
void ExtensionEntryPoints::UnloadImage() {
InitAmdExtTable();
InitImageExtTable();
core::hsa_internal_api_table_.Reset();
core::hsa_internal_api_table().Reset();
#ifdef HSA_IMAGE_SUPPORT
rocr::image::ReleaseImageRsrcs();
#endif
@@ -172,7 +172,7 @@ void ExtensionEntryPoints::Unload() {
InitPcSamplingExtTable();
InitImageExtTable();
InitAmdExtTable();
core::hsa_internal_api_table_.Reset();
core::hsa_internal_api_table().Reset();
}
bool ExtensionEntryPoints::LoadImage() {
@@ -193,7 +193,7 @@ bool ExtensionEntryPoints::LoadImage() {
image_api.version.step_id = HSA_IMAGE_API_TABLE_STEP_VERSION;
// Update private copy of Api table with handle for Image extensions
hsa_internal_api_table_.CloneExts(&image_api,
hsa_internal_api_table().CloneExts(&image_api,
core::HsaApiTable::HSA_EXT_IMAGE_API_TABLE_ID);
// Update Amd Ext Api table Api that deals with Images
@@ -215,7 +215,8 @@ void ExtensionEntryPoints::LoadPcSampling() {
pcs_api.version.step_id = HSA_PC_SAMPLING_API_TABLE_STEP_VERSION;
// Update private copy of Api table with handle for Image extensions
hsa_internal_api_table_.CloneExts(&pcs_api, core::HsaApiTable::HSA_EXT_PC_SAMPLING_API_TABLE_ID);
hsa_internal_api_table().CloneExts(&pcs_api,
core::HsaApiTable::HSA_EXT_PC_SAMPLING_API_TABLE_ID);
#endif
}
@@ -287,12 +288,12 @@ bool ExtensionEntryPoints::LoadFinalizer(std::string library_name) {
finalizer_api.version.step_id = HSA_FINALIZER_API_TABLE_STEP_VERSION;
// Update handle of table of HSA extensions
hsa_internal_api_table_.CloneExts(&finalizer_api,
hsa_internal_api_table().CloneExts(&finalizer_api,
core::HsaApiTable::HSA_EXT_FINALIZER_API_TABLE_ID);
ptr = os::GetExportAddress(lib, "Load");
if (ptr != NULL) {
((Load_t)ptr)(&core::hsa_internal_api_table_.hsa_api);
((Load_t)ptr)(&core::hsa_internal_api_table().hsa_api);
}
return true;
runtime/hsa-runtime/core/runtime/intercept_queue.cpp
-2
View File
@@ -80,8 +80,6 @@ static const uint16_t kBarrierHeader = (HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKE
(HSA_FENCE_SCOPE_NONE << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
(HSA_FENCE_SCOPE_NONE << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
int InterceptQueue::rtti_id_ = 0;
bool InterceptQueue::IsPendingRetryPoint(uint64_t wrapped_current_read_index) const {
// This function is intended to determine if the last retry barrier packet
// has definitely not been processed in order to avoid putting multiple retry
runtime/hsa-runtime/core/runtime/interrupt_signal.cpp
-2
View File
@@ -74,8 +74,6 @@ void InterruptSignal::EventPool::free(HsaEvent* evt) {
events_.push_back(unique_event_ptr(evt));
}
int InterruptSignal::rtti_id_ = 0;
HsaEvent* InterruptSignal::CreateEvent(HSA_EVENTTYPE type, bool manual_reset) {
HsaEventDescriptor event_descriptor;
event_descriptor.EventType = type;
runtime/hsa-runtime/core/runtime/ipc_signal.cpp
-1
View File
@@ -50,7 +50,6 @@
namespace rocr {
namespace core {
int IPCSignal::rtti_id_ = 0;
KernelMutex IPCSignal::lock_;
SharedMemory::SharedMemory(const hsa_amd_ipc_memory_t* handle, size_t len) {
runtime/hsa-runtime/core/runtime/isa.cpp
+18 -12
View File
@@ -199,12 +199,14 @@ hsa_round_method_t Isa::GetRoundMethod(
}
const Isa *IsaRegistry::GetIsa(const std::string &full_name) {
auto isareg_iter = supported_isas_.find(full_name);
return isareg_iter == supported_isas_.end() ? nullptr : &isareg_iter->second;
auto isareg_iter = GetSupportedIsas().find(full_name);
return isareg_iter == GetSupportedIsas().end() ?
nullptr : &isareg_iter->second;
}
const Isa *IsaRegistry::GetIsa(const Isa::Version &version, IsaFeature sramecc, IsaFeature xnack) {
auto isareg_iter = std::find_if(supported_isas_.begin(), supported_isas_.end(),
auto isareg_iter = std::find_if(GetSupportedIsas().begin(),
GetSupportedIsas().end(),
[&](const IsaMap::value_type& isareg) {
return isareg.second.GetVersion() == version &&
(isareg.second.GetSramecc() == IsaFeature::Unsupported ||
@@ -212,13 +214,12 @@ const Isa *IsaRegistry::GetIsa(const Isa::Version &version, IsaFeature sramecc,
(isareg.second.GetXnack() == IsaFeature::Unsupported ||
isareg.second.GetXnack() == xnack);
});
return isareg_iter == supported_isas_.end() ? nullptr : &isareg_iter->second;
return isareg_iter == GetSupportedIsas().end() ?
nullptr : &isareg_iter->second;
}
const IsaRegistry::IsaMap IsaRegistry::supported_isas_ =
IsaRegistry::GetSupportedIsas();
const IsaRegistry::IsaMap IsaRegistry::GetSupportedIsas() {
const IsaRegistry::IsaMap& IsaRegistry::GetSupportedIsas() {
// agent, and vendor name length limit excluding terminating nul character.
constexpr size_t hsa_name_size = 63;
@@ -236,11 +237,16 @@ constexpr size_t hsa_name_size = 63;
amd_amdgpu_##maj##min##stp##_SRAMECC_##sramecc##_XNACK_##xnack##_WAVEFRONTSIZE_##wavefrontsize.GetIsaName(), \
amd_amdgpu_##maj##min##stp##_SRAMECC_##sramecc##_XNACK_##xnack##_WAVEFRONTSIZE_##wavefrontsize)); \
IsaMap supported_isas;
IsaFeature unsupported = IsaFeature::Unsupported;
IsaFeature any = IsaFeature::Any;
IsaFeature disabled = IsaFeature::Disabled;
IsaFeature enabled = IsaFeature::Enabled;
static IsaMap supported_isas;
if (supported_isas.size() > 0) {
return supported_isas;
}
const IsaFeature unsupported = IsaFeature::Unsupported;
const IsaFeature any = IsaFeature::Any;
const IsaFeature disabled = IsaFeature::Disabled;
const IsaFeature enabled = IsaFeature::Enabled;
// Target ID Version SRAMECC XNACK
ISAREG_ENTRY_GEN("gfx700", 7, 0, 0, unsupported, unsupported, 64)
runtime/hsa-runtime/core/runtime/runtime.cpp
+40 -23
View File
@@ -40,8 +40,6 @@
//
////////////////////////////////////////////////////////////////////////////////
#include "core/inc/runtime.h"
#include <algorithm>
#include <atomic>
#include <climits>
@@ -50,6 +48,7 @@
#include <string>
#include <vector>
#include <list>
#include <link.h>
#include <dlfcn.h>
#include <amdgpu_drm.h>
#include <sys/mman.h>
@@ -59,6 +58,9 @@
#include <thread>
#include <chrono>
#include "core/inc/runtime.h"
#include "core/inc/hsa_table_interface.h"
#if defined(HSA_ROCPROFILER_REGISTER) && HSA_ROCPROFILER_REGISTER > 0
#include <rocprofiler-register/rocprofiler-register.h>
#endif
@@ -77,6 +79,7 @@
#include "core/inc/exceptions.h"
#include "inc/hsa_ven_amd_aqlprofile.h"
#include "core/inc/amd_core_dump.hpp"
#include "core/inc/host_queue.h"
#ifndef HSA_VERSION_MAJOR
#define HSA_VERSION_MAJOR 1
@@ -97,16 +100,20 @@ ROCPROFILER_REGISTER_DEFINE_IMPORT(hsa, ROCP_REG_VERSION)
const char rocrbuildid[] __attribute__((used)) = "ROCR BUILD ID: " STRING(ROCR_BUILD_ID);
extern r_debug _amdgpu_r_debug;
extern void _loader_debug_state();
namespace rocr {
namespace core {
bool g_use_interrupt_wait = true;
bool g_use_mwaitx = true;
bool g_use_interrupt_wait;
bool g_use_mwaitx;
Runtime* Runtime::runtime_singleton_ = NULL;
KernelMutex Runtime::bootstrap_lock_;
static bool loaded = true;
__forceinline static bool& loaded() {
static bool loaded_ = true;
return loaded_;
}
class RuntimeCleanup {
public:
@@ -115,7 +122,7 @@ class RuntimeCleanup {
delete Runtime::runtime_singleton_;
}
loaded = false;
loaded() = false;
}
};
@@ -123,9 +130,9 @@ static RuntimeCleanup cleanup_at_unload_;
hsa_status_t Runtime::Acquire() {
// Check to see if HSA has been cleaned up (process exit)
if (!loaded) return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
if (!loaded()) return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
ScopedAcquire<KernelMutex> boot(&bootstrap_lock_);
ScopedAcquire<KernelMutex> boot(&bootstrap_lock());
if (runtime_singleton_ == NULL) {
memset(log_flags, 0, sizeof(log_flags));
@@ -153,9 +160,9 @@ hsa_status_t Runtime::Acquire() {
hsa_status_t Runtime::Release() {
// Check to see if HSA has been cleaned up (process exit)
if (!loaded) return HSA_STATUS_SUCCESS;
if (!loaded()) return HSA_STATUS_SUCCESS;
ScopedAcquire<KernelMutex> boot(&bootstrap_lock_);
ScopedAcquire<KernelMutex> boot(&bootstrap_lock());
if (runtime_singleton_ == nullptr) return HSA_STATUS_ERROR_NOT_INITIALIZED;
@@ -747,11 +754,11 @@ hsa_status_t Runtime::GetSystemInfo(hsa_system_info_t attribute, void* value) {
((uint8_t*)value)[index] |= 1 << subBit;
};
if (hsa_internal_api_table_.finalizer_api.hsa_ext_program_finalize_fn != NULL) {
if (hsa_internal_api_table().finalizer_api.hsa_ext_program_finalize_fn != NULL) {
setFlag(HSA_EXTENSION_FINALIZER);
}
if (hsa_internal_api_table_.image_api.hsa_ext_image_create_fn != NULL) {
if (hsa_internal_api_table().image_api.hsa_ext_image_create_fn != NULL) {
setFlag(HSA_EXTENSION_IMAGES);
}
@@ -1944,6 +1951,16 @@ Runtime::Runtime()
asyncSignals_.monitor_exceptions = false;
asyncExceptions_.monitor_exceptions = true;
g_use_interrupt_wait = true;
g_use_mwaitx = true;
::_amdgpu_r_debug = {10,
nullptr,
reinterpret_cast<uintptr_t>(
&_loader_debug_state),
r_debug::RT_CONSISTENT,
0};
log_file = stderr;
}
hsa_status_t Runtime::Load() {
@@ -2068,17 +2085,17 @@ void Runtime::LoadExtensions() {
// Skipping finalizer loading since finalizer is no longer distributed.
// LinkExts will expose the finalizer-not-present implementation.
// extensions_.LoadFinalizer(kFinalizerLib[os_index(os::current_os)]);
hsa_api_table_.LinkExts(&extensions_.finalizer_api,
hsa_api_table().LinkExts(&extensions_.finalizer_api,
core::HsaApiTable::HSA_EXT_FINALIZER_API_TABLE_ID);
// Update Hsa Api Table with handle of Image extension Apis
extensions_.LoadImage();
hsa_api_table_.LinkExts(&extensions_.image_api,
hsa_api_table().LinkExts(&extensions_.image_api,
core::HsaApiTable::HSA_EXT_IMAGE_API_TABLE_ID);
// Update Hsa Api Table with handle of PCS extension Apis
extensions_.LoadPcSampling();
hsa_api_table_.LinkExts(&extensions_.pcs_api,
hsa_api_table().LinkExts(&extensions_.pcs_api,
core::HsaApiTable::HSA_EXT_PC_SAMPLING_API_TABLE_ID);
}
@@ -2216,7 +2233,7 @@ void Runtime::LoadTools() {
#if defined(HSA_ROCPROFILER_REGISTER) && HSA_ROCPROFILER_REGISTER > 0
if (!flag().disable_tool_register()) {
auto* profiler_api_table_ = static_cast<void*>(&hsa_api_table_);
auto* profiler_api_table_ = static_cast<void*>(&hsa_api_table());
auto lib_id = rocprofiler_register_library_indentifier_t{};
auto rocp_reg_status =
rocprofiler_register_library_api_table("hsa", &ROCPROFILER_REGISTER_IMPORT_FUNC(hsa),
@@ -2272,8 +2289,8 @@ void Runtime::LoadTools() {
}
// Discover loaded tools.
std::vector<os::LibHandle> loaded = os::GetLoadedToolsLib();
for(auto& handle : loaded) {
std::vector<os::LibHandle> loaded_hds = os::GetLoadedToolsLib();
for(auto& handle : loaded_hds) {
const uint32_t* order = (const uint32_t*)os::GetExportAddress(handle, "HSA_AMD_TOOL_PRIORITY");
if(order) {
sorted.push_back(lib_t(handle, *order+env_count, os::GetLibraryName(handle)));
@@ -2333,8 +2350,8 @@ void Runtime::LoadTools() {
os::CloseLib(tool);
continue;
}
if (!ld(&hsa_api_table_.hsa_api,
hsa_api_table_.hsa_api.version.major_id,
if (!ld(&hsa_api_table().hsa_api,
hsa_api_table().hsa_api.version.major_id,
failed.size(), failed.data())) {
failed.push_back(lib.name_.c_str());
os::CloseLib(tool);
@@ -2375,7 +2392,7 @@ void Runtime::UnloadTools() {
}
// Reset API table in case some tool doesn't cleanup properly
hsa_api_table_.Reset();
hsa_api_table().Reset();
}
void Runtime::CloseTools() {
runtime/hsa-runtime/core/runtime/signal.cpp
+4 -4
View File
@@ -64,7 +64,7 @@ void SharedSignalPool_t::clear() {
capacity - free_list_.size());
}
for (auto& block : block_list_) free_(block.first);
for (auto& block : block_list_) free_()(block.first);
block_list_.clear();
free_list_.clear();
}
@@ -73,15 +73,15 @@ SharedSignal* SharedSignalPool_t::alloc() {
ScopedAcquire<HybridMutex> lock(&lock_);
if (free_list_.empty()) {
SharedSignal* block = reinterpret_cast<SharedSignal*>(
allocate_(block_size_ * sizeof(SharedSignal), __alignof(SharedSignal), 0, 0));
allocate_()(block_size_ * sizeof(SharedSignal), __alignof(SharedSignal), 0, 0));
if (block == nullptr) {
block_size_ = minblock_;
block = reinterpret_cast<SharedSignal*>(
allocate_(block_size_ * sizeof(SharedSignal), __alignof(SharedSignal), 0, 0));
allocate_()(block_size_ * sizeof(SharedSignal), __alignof(SharedSignal), 0, 0));
if (block == nullptr) throw std::bad_alloc();
}
MAKE_NAMED_SCOPE_GUARD(throwGuard, [&]() { free_(block); });
MAKE_NAMED_SCOPE_GUARD(throwGuard, [&]() { free_()(block); });
block_list_.push_back(std::make_pair(block, block_size_));
throwGuard.Dismiss();
runtime/hsa-runtime/image/image_runtime.cpp
+6 -9
View File
@@ -61,9 +61,6 @@
namespace rocr {
namespace image {
std::atomic<ImageRuntime*> ImageRuntime::instance_(NULL);
std::mutex ImageRuntime::instance_mutex_;
hsa_status_t FindKernelArgPool(hsa_amd_memory_pool_t pool, void* data) {
assert(data != nullptr);
@@ -154,13 +151,13 @@ hsa_status_t ImageRuntime::CreateImageManager(hsa_agent_t agent, void* data) {
}
ImageRuntime* ImageRuntime::instance() {
ImageRuntime* instance = instance_.load(std::memory_order_acquire);
ImageRuntime* instance = get_instance().load(std::memory_order_acquire);
if (instance == NULL) {
// Protect the initialization from multi threaded access.
std::lock_guard<std::mutex> lock(instance_mutex_);
std::lock_guard<std::mutex> lock(instance_mutex());
// Make sure we are not initializing it twice.
instance = instance_.load(std::memory_order_relaxed);
instance = get_instance().load(std::memory_order_relaxed);
if (instance != NULL) {
return instance;
}
@@ -194,19 +191,19 @@ ImageRuntime* ImageRuntime::CreateSingleton() {
assert(instance->kernarg_pool_.handle != 0);
assert(instance->image_managers_.size() != 0);
instance_.store(instance, std::memory_order_release);
get_instance().store(instance, std::memory_order_release);
return instance;
}
void ImageRuntime::DestroySingleton() {
ImageRuntime* instance = instance_.load(std::memory_order_acquire);
ImageRuntime* instance = get_instance().load(std::memory_order_acquire);
if (instance == NULL) {
return;
}
instance->Cleanup();
instance_.store(NULL, std::memory_order_release);
get_instance().store(NULL, std::memory_order_release);
delete instance;
}
runtime/hsa-runtime/image/image_runtime.h
+8 -2
View File
@@ -163,9 +163,15 @@ class ImageRuntime {
void Cleanup();
/// Pointer to singleton object.
static std::atomic<ImageRuntime*> instance_;
static __forceinline std::atomic<ImageRuntime*>& get_instance() {
static std::atomic<ImageRuntime*> instance_(NULL);
return instance_;
}
static std::mutex instance_mutex_;
static __forceinline std::mutex& instance_mutex() {
static std::mutex instance_mutex_;
return instance_mutex_;
}
/// @brief Contains mapping of agent and its corresponding ::ImageManager
/// object.
runtime/hsa-runtime/loader/executable.cpp
+14 -14
View File
@@ -68,7 +68,7 @@ using namespace rocr::amd::hsa::common;
// Having a side effect prevents call site optimization that allows removal of a noinline function call
// with no side effect.
__attribute__((noinline)) static void _loader_debug_state() {
__attribute__((noinline)) void _loader_debug_state() {
static volatile int function_needs_a_side_effect = 0;
function_needs_a_side_effect ^= 1;
}
@@ -84,12 +84,12 @@ __attribute__((noinline)) static void _loader_debug_state() {
// 9: New trap handler ABI. For gfx11: Save PC in ttmp11[22:7] ttmp6[31:0], and park the wave if stopped.
// 10: New trap handler ABI. Set status.skip_export when halting the wave.
// For gfx940, set ttmp6[31] = 0 if ttmp11[31] == 0.
HSA_API r_debug _amdgpu_r_debug = {10,
nullptr,
reinterpret_cast<uintptr_t>(&_loader_debug_state),
r_debug::RT_CONSISTENT,
0};
static link_map* r_debug_tail = nullptr;
HSA_API r_debug _amdgpu_r_debug;
static __forceinline link_map*& r_debug_tail() {
static link_map* r_debug_tail_ = nullptr;
return r_debug_tail_;
}
namespace rocr {
namespace amd {
@@ -175,7 +175,7 @@ void Loader::Destroy(Loader *loader)
// Loader resets the link_map, but the executables and loaded code objects are not deleted.
_amdgpu_r_debug.r_map = nullptr;
_amdgpu_r_debug.r_state = r_debug::RT_CONSISTENT;
r_debug_tail = nullptr;
r_debug_tail() = nullptr;
delete loader;
}
@@ -201,21 +201,21 @@ Executable* AmdHsaCodeLoader::CreateExecutable(
}
static void AddCodeObjectInfoIntoDebugMap(link_map* map) {
if (r_debug_tail) {
r_debug_tail->l_next = map;
map->l_prev = r_debug_tail;
if (r_debug_tail()) {
r_debug_tail()->l_next = map;
map->l_prev = r_debug_tail();
map->l_next = nullptr;
} else {
_amdgpu_r_debug.r_map = map;
map->l_prev = nullptr;
map->l_next = nullptr;
}
r_debug_tail = map;
r_debug_tail() = map;
}
static void RemoveCodeObjectInfoFromDebugMap(link_map* map) {
if (r_debug_tail == map) {
r_debug_tail = map->l_prev;
if (r_debug_tail() == map) {
r_debug_tail() = map->l_prev;
}
if (_amdgpu_r_debug.r_map == map) {
_amdgpu_r_debug.r_map = map->l_next;
runtime/hsa-runtime/pcs/pcs_runtime.cpp
+6 -8
View File
@@ -57,17 +57,15 @@ do { \
if ((ptr) == NULL) return HSA_STATUS_ERROR_INVALID_ARGUMENT; \
} while (false)
std::atomic<PcsRuntime*> PcsRuntime::instance_(NULL);
std::mutex PcsRuntime::instance_mutex_;
PcsRuntime* PcsRuntime::instance() {
PcsRuntime* instance = instance_.load(std::memory_order_acquire);
PcsRuntime* instance = get_instance().load(std::memory_order_acquire);
if (instance == NULL) {
// Protect the initialization from multi threaded access.
std::lock_guard<std::mutex> lock(instance_mutex_);
std::lock_guard<std::mutex> lock(instance_mutex());
// Make sure we are not initializing it twice.
instance = instance_.load(std::memory_order_relaxed);
instance = get_instance().load(std::memory_order_relaxed);
if (instance != NULL) {
return instance;
}
@@ -84,17 +82,17 @@ PcsRuntime* PcsRuntime::instance() {
PcsRuntime* PcsRuntime::CreateSingleton() {
PcsRuntime* instance = new PcsRuntime();
instance_.store(instance, std::memory_order_release);
get_instance().store(instance, std::memory_order_release);
return instance;
}
void PcsRuntime::DestroySingleton() {
PcsRuntime* instance = instance_.load(std::memory_order_acquire);
PcsRuntime* instance = get_instance().load(std::memory_order_acquire);
if (instance == NULL) {
return;
}
instance_.store(NULL, std::memory_order_release);
get_instance().store(NULL, std::memory_order_release);
delete instance;
}
runtime/hsa-runtime/pcs/pcs_runtime.h
+8 -3
View File
@@ -152,9 +152,14 @@ class PcsRuntime {
static PcsRuntime* CreateSingleton();
/// Pointer to singleton object.
static std::atomic<PcsRuntime*> instance_;
static std::mutex instance_mutex_;
static __forceinline std::atomic<PcsRuntime*>& get_instance() {
static std::atomic<PcsRuntime*> instance_(nullptr);
return instance_;
}
static __forceinline std::mutex& instance_mutex() {
static std::mutex instance_mutex_;
return instance_mutex_;
}
// Map of pc sampling sessions indexed by hsa_ven_amd_pcs_t handle
std::map<uint64_t, PcSamplingSession> pc_sampling_;
KernelMutex pc_sampling_lock_;