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

rocr: Flags to alloc queue buf/struct in dev mem

Browse Source 
This builds on a prior change that allowed for allocating
a user-mode queue's packet buffer in device memory to also
allocate the queue struct in device memory. This provides
additional latency benefits particularly for cases where
dispatches are performed from the GPU itself. Flags are
added to support the various use cases.


[ROCm/ROCR-Runtime commit: 6e3c375bf1]
This commit is contained in:
Tony Gutierrez
2025-01-15 12:13:55 -08:00
committed by Yat Sin, David
parent 18404ba8a8
commit 6f37386eb2
23 changed files with 220 additions and 131 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocr-runtime/libhsakmt/src/memory.c
-1
View File
@@ -150,7 +150,6 @@ HSAKMT_STATUS HSAKMTAPI hsaKmtAllocMemoryAlign(HSAuint32 PreferredNode,
*MemoryAddress = NULL;
if ((MemFlags.ui32.CoarseGrain && MemFlags.ui32.ExtendedCoherent) ||
(MemFlags.ui32.CoarseGrain && MemFlags.ui32.Uncached) ||
(MemFlags.ui32.ExtendedCoherent && MemFlags.ui32.Uncached))
return HSAKMT_STATUS_INVALID_PARAMETER;
projects/rocr-runtime/runtime/hsa-runtime/core/inc/agent.h
+22 -22
View File
@@ -255,29 +255,29 @@ class Agent : public Checked<0xF6BC25EB17E6F917> {
virtual hsa_status_t IterateCache(hsa_status_t (*callback)(hsa_cache_t cache, void* data),
void* data) const = 0;
// @brief Create queue.
//
// @param [in] size Number of packets the queue is expected to hold. Must be a
// power of 2 greater than 0.
// @param [in] queue_type Queue type.
// @param [in] event_callback Callback invoked for every
// asynchronous event related to the newly created queue. May be NULL.The HSA
// runtime passes three arguments to the callback : a code identifying the
// event that triggered the invocation, a pointer to the queue where the event
// originated, and the application data.
// @param [in] data Application data that is passed to @p callback.
// @param [in] private_segment_size A hint to indicate the maximum expected
// private segment usage per work-item, in bytes.
// @param [in] group_segment_size A hint to indicate the maximum expected
// group segment usage per work-group, in bytes.
// @param[out] queue Memory location where the HSA runtime stores a pointer
// to the newly created queue.
//
// @retval HSA_STATUS_SUCCESS The queue has been created successfully.
virtual hsa_status_t QueueCreate(size_t size, hsa_queue_type32_t queue_type,
/// @brief Create queue.
///
/// @param [in] size Number of packets the queue is expected to hold. Must be a
/// power of 2 greater than 0.
/// @param [in] queue_type Queue type.
/// @param [in] flags Flags to specify queue attributes on creation.
/// @param [in] event_callback Callback invoked for every
/// asynchronous event related to the newly created queue. May be NULL.The HSA
/// runtime passes three arguments to the callback : a code identifying the
/// event that triggered the invocation, a pointer to the queue where the event
/// originated, and the application data.
/// @param [in] data Application data that is passed to @p callback.
/// @param [in] private_segment_size A hint to indicate the maximum expected
/// private segment usage per work-item, in bytes.
/// @param [in] group_segment_size A hint to indicate the maximum expected
/// group segment usage per work-group, in bytes.
/// @param[out] queue Memory location where the HSA runtime stores a pointer
/// to the newly created queue.
///
/// @retval HSA_STATUS_SUCCESS The queue has been created successfully.
virtual hsa_status_t QueueCreate(size_t size, hsa_queue_type32_t queue_type, uint64_t flags,
HsaEventCallback event_callback, void* data,
uint32_t private_segment_size,
uint32_t group_segment_size,
uint32_t private_segment_size, uint32_t group_segment_size,
Queue** queue) = 0;
// @brief Query the value of an attribute.
projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_aie_agent.h
+4 -5
View File
@@ -78,11 +78,10 @@ public:
hsa_status_t GetInfo(hsa_agent_info_t attribute, void *value) const override;
hsa_status_t QueueCreate(size_t size, hsa_queue_type32_t queue_type,
core::HsaEventCallback event_callback, void *data,
uint32_t private_segment_size,
uint32_t group_segment_size,
core::Queue **queue) override;
hsa_status_t QueueCreate(size_t size, hsa_queue_type32_t queue_type, uint64_t flags,
core::HsaEventCallback event_callback, void* data,
uint32_t private_segment_size, uint32_t group_segment_size,
core::Queue** queue) override;
// @brief Override from core::Agent.
const std::vector<const core::Isa*>& supported_isas() const override {
projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_aie_aql_queue.h
+2 -1
View File
@@ -68,7 +68,8 @@ class AieAqlQueue : public core::Queue,
return queue->IsType(&rtti_id());
}
AieAqlQueue(AieAgent *agent, size_t req_size_pkts, uint32_t node_id);
AieAqlQueue(core::SharedQueue* shared_queue, AieAgent* agent, size_t req_size_pkts,
uint32_t node_id, uint64_t flags);
~AieAqlQueue();
hsa_status_t Inactivate() override;
projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_aql_queue.h
+6 -4
View File
@@ -63,9 +63,9 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Doo
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,
ScratchInfo& scratch, core::HsaEventCallback callback,
void* err_data, bool is_kv = false);
AqlQueue(core::SharedQueue* shared_queue, GpuAgent* agent, size_t req_size_pkts,
HSAuint32 node_id, ScratchInfo& scratch, core::HsaEventCallback callback, void* err_data,
uint64_t flags, bool is_kv = false);
~AqlQueue();
@@ -236,7 +236,9 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Doo
// (De)allocates and (de)registers ring_buf_.
void AllocRegisteredRingBuffer(uint32_t queue_size_pkts);
void FreeRegisteredRingBuffer();
/// @brief Frees the queue's packet ring buffer and its queue struct.
void FreeQueueMemory();
/// @brief Abstracts the file handle use for double mapping queues.
void CloseRingBufferFD(const char* ring_buf_shm_path, int fd) const;
projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_cpu_agent.h
+2 -3
View File
@@ -100,10 +100,9 @@ class CpuAgent : public core::Agent {
hsa_status_t GetInfo(hsa_agent_info_t attribute, void* value) const override;
// @brief Override from core::Agent.
hsa_status_t QueueCreate(size_t size, hsa_queue_type32_t queue_type,
hsa_status_t QueueCreate(size_t size, hsa_queue_type32_t queue_type, uint64_t flags,
core::HsaEventCallback event_callback, void* data,
uint32_t private_segment_size,
uint32_t group_segment_size,
uint32_t private_segment_size, uint32_t group_segment_size,
core::Queue** queue) override;
// @brief Override from core::Agent.
projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_gpu_agent.h
+5 -6
View File
@@ -328,10 +328,9 @@ class GpuAgent : public GpuAgentInt {
hsa_status_t GetInfo(hsa_agent_info_t attribute, void* value) const override;
// @brief Override from core::Agent.
hsa_status_t QueueCreate(size_t size, hsa_queue_type32_t queue_type,
hsa_status_t QueueCreate(size_t size, hsa_queue_type32_t queue_type, uint64_t flags,
core::HsaEventCallback event_callback, void* data,
uint32_t private_segment_size,
uint32_t group_segment_size,
uint32_t private_segment_size, uint32_t group_segment_size,
core::Queue** queue) override;
// @brief Decrement GWS ref count.
@@ -436,7 +435,7 @@ class GpuAgent : public GpuAgentInt {
if (t0_.GPUClockCounter == t1_.GPUClockCounter) SyncClocks();
}
// @brief Override from AMD::GpuAgentInt.
/// @brief Override from AMD::GpuAgentInt.
__forceinline bool is_xgmi_cpu_gpu() const { return xgmi_cpu_gpu_; }
const size_t MAX_SCRATCH_APERTURE_PER_XCC = (1ULL << 32);
@@ -830,8 +829,8 @@ class GpuAgent : public GpuAgentInt {
// structure for stochastic sampling
pcs_data_t pcs_stochastic_data_;
// @bried XGMI CPU<->GPU
bool xgmi_cpu_gpu_;
/// @brief XGMI CPU<->GPU
bool xgmi_cpu_gpu_ = false;
};
} // namespace amd
projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_memory_region.h
+2 -5
View File
@@ -171,7 +171,7 @@ class MemoryRegion : public core::MemoryRegion {
return static_cast<uint32_t>(mem_props_.MemoryClockMax);
}
__forceinline size_t GetPageSize() const { return kPageSize(); }
__forceinline static size_t GetPageSize() { return kPageSize_; }
__forceinline const HsaMemFlags &mem_flags() const { return mem_flag_; }
__forceinline const HsaMemMapFlags &map_flags() const { return map_flag_; }
@@ -199,10 +199,7 @@ private:
// fragments of the block routing to the same MemoryRegion.
mutable KernelMutex access_lock_;
static __forceinline const size_t& kPageSize() {
static size_t kPageSize_ = sysconf(_SC_PAGESIZE);
return kPageSize_;
}
static const size_t kPageSize_;
// Determine access type allowed to requesting device
hsa_amd_memory_pool_access_t GetAccessInfo(const core::Agent& agent,
projects/rocr-runtime/runtime/hsa-runtime/core/inc/host_queue.h
+2 -2
View File
@@ -54,8 +54,8 @@ class HostQueue : public Queue {
public:
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);
HostQueue(core::SharedQueue* shared_queue, hsa_region_t region, uint32_t ring_size,
hsa_queue_type32_t type, uint32_t features, hsa_signal_t doorbell_signal);
~HostQueue();
projects/rocr-runtime/runtime/hsa-runtime/core/inc/intercept_queue.h
+8 -2
View File
@@ -63,12 +63,18 @@ class QueueWrapper : public Queue {
public:
std::unique_ptr<Queue> wrapped;
explicit QueueWrapper(std::unique_ptr<Queue> queue) : Queue(), wrapped(std::move(queue)) {
explicit QueueWrapper(std::unique_ptr<Queue> queue)
: Queue(static_cast<core::SharedQueue*>(core::Runtime::runtime_singleton_->system_allocator()(
sizeof(core::SharedQueue), 4096, 0, 0)),
0),
wrapped(std::move(queue)) {
memcpy(&amd_queue_, &wrapped->amd_queue_, sizeof(amd_queue_));
wrapped->set_public_handle(wrapped.get(), public_handle_);
}
~QueueWrapper() {}
~QueueWrapper() {
if (shared_queue_) core::Runtime::runtime_singleton_->system_deallocator()(shared_queue_);
}
hsa_status_t Inactivate() override { return wrapped->Inactivate(); }
hsa_status_t SetPriority(HSA_QUEUE_PRIORITY priority) override {
projects/rocr-runtime/runtime/hsa-runtime/core/inc/queue.h
+24 -22
View File
@@ -158,16 +158,6 @@ struct SharedQueue {
Queue* core_queue;
};
class LocalQueue {
public:
LocalQueue(int mem_flags) : local_queue_(mem_flags) {}
LocalQueue(int agent_node_id, int mem_flags) : local_queue_(agent_node_id, mem_flags) {}
SharedQueue* queue() const { return local_queue_.shared_object(); }
private:
Shared<SharedQueue> local_queue_;
};
/// @brief Class Queue which encapsulate user mode queues and
/// provides Api to access its Read, Write indices using Acquire,
/// Release and Relaxed semantics.
@@ -176,18 +166,18 @@ Queue is intended to be an pure interface class and may be wrapped or replaced
by tools.
All funtions other than Convert and public_handle must be virtual.
*/
class Queue : public Checked<0xFA3906A679F9DB49>, private LocalQueue {
class Queue : public Checked<0xFA3906A679F9DB49> {
public:
Queue(int mem_flags = 0) : LocalQueue(mem_flags), amd_queue_(queue()->amd_queue) {
queue()->core_queue = this;
public_handle_ = Convert(this);
pcie_write_ordering_ = false;
}
Queue(SharedQueue* shared_queue, uint64_t queue_flags)
: Queue(shared_queue, queue_flags, false) {}
Queue(int agent_node_id, int mem_flags) : LocalQueue(agent_node_id, mem_flags), amd_queue_(queue()->amd_queue) {
queue()->core_queue = this;
Queue(SharedQueue* shared_queue, uint64_t queue_flags, bool pcie_write_ordering)
: amd_queue_(shared_queue->amd_queue),
shared_queue_(shared_queue),
flags_(queue_flags),
pcie_write_ordering_(pcie_write_ordering) {
public_handle_ = Convert(this);
pcie_write_ordering_ = false;
shared_queue->core_queue = this;
}
virtual ~Queue() {}
@@ -386,9 +376,18 @@ class Queue : public Checked<0xFA3906A679F9DB49>, private LocalQueue {
bool IsType(rtti_t id) { return _IsA(id); }
bool needsPcieOrdering() const { return pcie_write_ordering_; }
/// @brief Used to determine if the queue's packet buffer was allocated
/// in the agent's local device memory.
bool IsDeviceMemRingBuf() const {
return (flags_ & HSA_AMD_QUEUE_CREATE_DEVICE_MEM_RING_BUF) != 0;
}
/// @brief Used to determine if the queue descriptor was allocated in
/// the agent's local device memory.
bool IsDeviceMemQueueDescriptor() const {
return (flags_ & HSA_AMD_QUEUE_CREATE_DEVICE_MEM_QUEUE_DESCRIPTOR) != 0;
}
void setPcieOrdering(bool val) { pcie_write_ordering_ = val; }
bool needsPcieOrdering() const { return pcie_write_ordering_; }
protected:
static void set_public_handle(Queue* ptr, hsa_queue_t* handle) {
@@ -400,6 +399,8 @@ class Queue : public Checked<0xFA3906A679F9DB49>, private LocalQueue {
virtual bool _IsA(rtti_t id) const = 0;
SharedQueue* shared_queue_;
hsa_queue_t* public_handle_;
/// Next available queue id.
@@ -410,7 +411,8 @@ class Queue : public Checked<0xFA3906A679F9DB49>, private LocalQueue {
// HSA Queue ID - used to bind a unique ID
static std::atomic<uint64_t> hsa_queue_counter_;
bool pcie_write_ordering_;
const uint64_t flags_;
bool pcie_write_ordering_ = false;
DISALLOW_COPY_AND_ASSIGN(Queue);
};
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aie_agent.cpp
+22 -6
View File
@@ -261,11 +261,16 @@ hsa_status_t AieAgent::GetInfo(hsa_agent_info_t attribute, void *value) const {
return HSA_STATUS_SUCCESS;
}
hsa_status_t AieAgent::QueueCreate(size_t size, hsa_queue_type32_t queue_type,
core::HsaEventCallback event_callback,
void *data, uint32_t private_segment_size,
uint32_t group_segment_size,
core::Queue **queue) {
hsa_status_t AieAgent::QueueCreate(size_t size, hsa_queue_type32_t queue_type, uint64_t flags,
core::HsaEventCallback event_callback, void* data,
uint32_t private_segment_size, uint32_t group_segment_size,
core::Queue** queue) {
if ((flags & HSA_AMD_QUEUE_CREATE_DEVICE_MEM_RING_BUF) != 0 ||
(flags & HSA_AMD_QUEUE_CREATE_DEVICE_MEM_QUEUE_DESCRIPTOR) != 0) {
// AIE agents do not currently support queue creation in device memory.
return HSA_STATUS_ERROR_INVALID_ARGUMENT;
}
if (!IsPowerOfTwo(size)) {
return HSA_STATUS_ERROR_INVALID_ARGUMENT;
}
@@ -274,7 +279,18 @@ hsa_status_t AieAgent::QueueCreate(size_t size, hsa_queue_type32_t queue_type,
return HSA_STATUS_ERROR_INVALID_ARGUMENT;
}
auto aql_queue(new AieAqlQueue(this, size, node_id()));
core::SharedQueue* shared_queue =
static_cast<core::SharedQueue*>(core::Runtime::runtime_singleton_->system_allocator()(
sizeof(core::SharedQueue), MemoryRegion::GetPageSize(), 0, node_id()));
if (!shared_queue) return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
auto aql_queue(new AieAqlQueue(shared_queue, this, size, node_id(), flags));
if (aql_queue == nullptr) {
core::Runtime::runtime_singleton_->system_deallocator()(shared_queue);
return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
}
*queue = aql_queue;
return HSA_STATUS_SUCCESS;
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aie_aql_queue.cpp
+11 -7
View File
@@ -65,10 +65,13 @@
namespace rocr {
namespace AMD {
AieAqlQueue::AieAqlQueue(AieAgent *agent, size_t req_size_pkts,
uint32_t node_id)
: Queue(0, 0), LocalSignal(0, false), DoorbellSignal(signal()),
agent_(*agent), active_(false) {
AieAqlQueue::AieAqlQueue(core::SharedQueue* shared_queue, AieAgent* agent, size_t req_size_pkts,
uint32_t node_id, uint64_t flags)
: Queue(shared_queue, flags),
LocalSignal(0, false),
DoorbellSignal(signal()),
agent_(*agent),
active_(false) {
if (agent_.device_type() != core::Agent::DeviceType::kAmdAieDevice) {
throw AMD::hsa_exception(
HSA_STATUS_ERROR_INVALID_AGENT,
@@ -105,9 +108,10 @@ AieAqlQueue::AieAqlQueue(AieAgent *agent, size_t req_size_pkts,
AieAqlQueue::~AieAqlQueue() {
AieAqlQueue::Inactivate();
if (ring_buf_) {
agent_.system_deallocator()(ring_buf_);
}
if (ring_buf_) agent_.system_deallocator()(ring_buf_);
if (shared_queue_) core::Runtime::runtime_singleton_->system_deallocator()(shared_queue_);
}
hsa_status_t AieAqlQueue::Inactivate() {
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp
+24 -14
View File
@@ -73,12 +73,12 @@
namespace rocr {
namespace AMD {
#define SCRATCH_ALT_RATIO 4
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),
AqlQueue::AqlQueue(core::SharedQueue* shared_queue, GpuAgent* agent, size_t req_size_pkts,
HSAuint32 node_id, ScratchInfo& scratch, core::HsaEventCallback callback,
void* err_data, uint64_t flags, bool is_kv)
: Queue(shared_queue, flags, !agent->is_xgmi_cpu_gpu()),
LocalSignal(0, false),
DoorbellSignal(signal()),
ring_buf_(nullptr),
@@ -129,7 +129,7 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr
// Allocate the AQL packet ring buffer.
AllocRegisteredRingBuffer(queue_size_pkts);
if (ring_buf_ == nullptr) throw std::bad_alloc();
MAKE_NAMED_SCOPE_GUARD(RingGuard, [&]() { FreeRegisteredRingBuffer(); });
MAKE_NAMED_SCOPE_GUARD(RingGuard, [&]() { FreeQueueMemory(); });
// Fill the ring buffer with invalid packet headers.
// Leave packet content uninitialized to help track errors.
@@ -350,7 +350,6 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr
if (!core::Runtime::runtime_singleton_->flag().cu_mask_skip_init()) SetCUMasking(0, nullptr);
active_ = true;
setPcieOrdering(agent->is_xgmi_cpu_gpu());
PM4IBGuard.Dismiss();
RingGuard.Dismiss();
@@ -386,10 +385,11 @@ AqlQueue::~AqlQueue() {
if (queue_scratch_.main_queue_base) agent_->ReleaseQueueMainScratch(queue_scratch_);
if (queue_scratch_.alt_queue_base) agent_->ReleaseQueueAltScratch(queue_scratch_);
FreeRegisteredRingBuffer();
exception_signal_->WaitingDec();
exception_signal_->DestroySignal();
HSA::hsa_signal_destroy(amd_queue_.queue_inactive_signal);
FreeQueueMemory();
if (core::g_use_interrupt_wait) {
ScopedAcquire<KernelMutex> lock(&queue_lock());
queue_count()--;
@@ -737,9 +737,10 @@ void AqlQueue::AllocRegisteredRingBuffer(uint32_t queue_size_pkts) {
ring_buf_alloc_bytes_ = queue_size_pkts * sizeof(core::AqlPacket);
assert(IsMultipleOf(ring_buf_alloc_bytes_, 4096) && "Ring buffer sizes must be 4KiB aligned.");
if (core::Runtime::runtime_singleton_->flag().dev_mem_queue()) {
ring_buf_ = agent_->finegrain_allocator()(ring_buf_alloc_bytes_,
core::MemoryRegion::AllocateUncached);
if (IsDeviceMemRingBuf()) {
ring_buf_ = agent_->coarsegrain_allocator()(
ring_buf_alloc_bytes_,
core::MemoryRegion::AllocateExecutable | core::MemoryRegion::AllocateUncached);
} else {
ring_buf_ = agent_->system_allocator()(
ring_buf_alloc_bytes_, 0x1000,
@@ -755,7 +756,16 @@ void AqlQueue::AllocRegisteredRingBuffer(uint32_t queue_size_pkts) {
}
}
void AqlQueue::FreeRegisteredRingBuffer() {
void AqlQueue::FreeQueueMemory() {
if (shared_queue_) {
if (IsDeviceMemQueueDescriptor())
agent_->coarsegrain_deallocator()(shared_queue_);
else
core::Runtime::runtime_singleton_->system_deallocator()(shared_queue_);
shared_queue_ = nullptr;
}
if ((agent_->profile() == HSA_PROFILE_FULL) && queue_full_workaround_) {
#ifdef __linux__
munmap(ring_buf_, ring_buf_alloc_bytes_);
@@ -767,8 +777,8 @@ void AqlQueue::FreeRegisteredRingBuffer() {
#endif
} else {
if (ring_buf_) {
if (core::Runtime::runtime_singleton_->flag().dev_mem_queue()) {
agent_->finegrain_deallocator()(ring_buf_);
if (IsDeviceMemRingBuf()) {
agent_->coarsegrain_deallocator()(ring_buf_);
} else {
agent_->system_deallocator()(ring_buf_);
}
@@ -1664,7 +1674,7 @@ void AqlQueue::ExecutePM4(uint32_t* cmd_data, size_t cmd_size_b, hsa_fence_scope
// Overwrite the AQL invalid header (first dword) last.
// This prevents the slot from being read until it's fully written.
memcpy(&queue_slot[1], &slot_data[1], slot_size_b - sizeof(uint32_t));
if (core::Runtime::runtime_singleton_->flag().dev_mem_queue() && !agent_->is_xgmi_cpu_gpu()) {
if (IsDeviceMemRingBuf() && needsPcieOrdering()) {
// Ensure the packet body is written as header may get reordered when writing over PCIE
_mm_sfence();
}
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp
+1 -1
View File
@@ -890,7 +890,7 @@ void BlitKernel::PopulateQueue(uint64_t index, uint64_t code_handle, void* args,
std::atomic_thread_fence(std::memory_order_acquire);
queue_buffer[index & queue_bitmask_] = packet;
std::atomic_thread_fence(std::memory_order_release);
if (core::Runtime::runtime_singleton_->flag().dev_mem_queue() && !queue_->needsPcieOrdering()) {
if (queue_->IsDeviceMemRingBuf() && queue_->needsPcieOrdering()) {
// Ensure the packet body is written as header may get reordered when writing over PCIE
_mm_sfence();
}
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_cpu_agent.cpp
+3 -4
View File
@@ -424,10 +424,9 @@ hsa_status_t CpuAgent::GetInfo(hsa_agent_info_t attribute, void* value) const {
return HSA_STATUS_SUCCESS;
}
hsa_status_t CpuAgent::QueueCreate(size_t size, hsa_queue_type32_t queue_type,
core::HsaEventCallback event_callback,
void* data, uint32_t private_segment_size,
uint32_t group_segment_size,
hsa_status_t CpuAgent::QueueCreate(size_t size, hsa_queue_type32_t queue_type, uint64_t flags,
core::HsaEventCallback event_callback, void* data,
uint32_t private_segment_size, uint32_t group_segment_size,
core::Queue** queue) {
// No HW AQL packet processor on CPU device.
return HSA_STATUS_ERROR;
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp
+24 -7
View File
@@ -706,7 +706,8 @@ core::Queue* GpuAgent::CreateInterceptibleQueue(void (*callback)(hsa_status_t st
uint32_t size = std::max(in_size, minAqlSize_);
size = std::min(size, maxAqlSize_);
QueueCreate(size, HSA_QUEUE_TYPE_MULTI, callback, data, 0, 0, &queue);
QueueCreate(size, HSA_QUEUE_TYPE_MULTI, HSA_AMD_QUEUE_CREATE_SYSTEM_MEM, callback, data, 0, 0,
&queue);
if (queue != nullptr)
core::Runtime::runtime_singleton_->InternalQueueCreateNotify(core::Queue::Convert(queue),
this->public_handle());
@@ -1666,10 +1667,9 @@ hsa_status_t GpuAgent::GetInfo(hsa_agent_info_t attribute, void* value) const {
return HSA_STATUS_SUCCESS;
}
hsa_status_t GpuAgent::QueueCreate(size_t size, hsa_queue_type32_t queue_type,
core::HsaEventCallback event_callback,
void* data, uint32_t private_segment_size,
uint32_t group_segment_size,
hsa_status_t GpuAgent::QueueCreate(size_t size, hsa_queue_type32_t queue_type, uint64_t flags,
core::HsaEventCallback event_callback, void* data,
uint32_t private_segment_size, uint32_t group_segment_size,
core::Queue** queue) {
// Handle GWS queues.
if (queue_type == HSA_QUEUE_TYPE_COOPERATIVE) {
@@ -1739,9 +1739,26 @@ hsa_status_t GpuAgent::QueueCreate(size_t size, hsa_queue_type32_t queue_type,
// ensured.
queues_[QueueUtility].touch();
bool dev_mem_queue_descriptor = (flags & HSA_AMD_QUEUE_CREATE_DEVICE_MEM_QUEUE_DESCRIPTOR) != 0;
// Create an HW AQL queue
auto aql_queue =
new AqlQueue(this, size, node_id(), scratch, event_callback, data, is_kv_device_);
core::SharedQueue* shared_queue = nullptr;
if (dev_mem_queue_descriptor) {
shared_queue = static_cast<core::SharedQueue*>(
finegrain_allocator()(sizeof(core::SharedQueue), core::MemoryRegion::AllocateUncached));
} else {
shared_queue =
static_cast<core::SharedQueue*>(core::Runtime::runtime_singleton_->system_allocator()(
sizeof(core::SharedQueue), MemoryRegion::GetPageSize(),
isMES() ? (MemoryRegion::AllocateGTTAccess | MemoryRegion::AllocateNonPaged) : 0,
node_id()));
}
if (!shared_queue) return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
auto aql_queue = new AqlQueue(shared_queue, this, size, node_id(), scratch, event_callback, data,
flags, is_kv_device_);
*queue = aql_queue;
aql_queues_.push_back(aql_queue);
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_memory_region.cpp
+9 -8
View File
@@ -56,6 +56,7 @@ namespace AMD {
// Tracks aggregate size of system memory available on platform
size_t MemoryRegion::max_sysmem_alloc_size_ = 0;
const size_t MemoryRegion::kPageSize_ = sysconf(_SC_PAGESIZE);
bool MemoryRegion::RegisterMemory(void* ptr, size_t size, const HsaMemFlags& MemFlags) {
assert(ptr != NULL);
@@ -123,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 = GetPageSize();
mem_flag_.ui32.NoSubstitute = 0;
mem_flag_.ui32.HostAccess = 1;
mem_flag_.ui32.CachePolicy = HSA_CACHING_CACHED;
@@ -136,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()), GetPageSize());
// Keep track of total system memory available
// @note: System memory is surfaced as both coarse
@@ -148,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_, GetPageSize()));
}
MemoryRegion::~MemoryRegion() {}
@@ -175,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, GetPageSize());
return owner()->driver().AllocateMemory(*this, alloc_flags, address, size,
agent_node_id);
@@ -276,7 +277,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) = GetPageSize();
break;
default:
*((size_t*)value) = 0;
@@ -289,7 +290,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) = GetPageSize();
break;
default:
*((size_t*)value) = 0;
@@ -358,12 +359,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) = GetPageSize();
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()
? GetPageSize()
: fragment_allocator_.default_block_size();
break;
default:
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/host_queue.cpp
+4 -3
View File
@@ -48,9 +48,9 @@
namespace rocr {
namespace core {
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) {
HostQueue::HostQueue(core::SharedQueue* shared_queue, hsa_region_t region, uint32_t ring_size,
hsa_queue_type32_t type, uint32_t features, hsa_signal_t doorbell_signal)
: Queue(shared_queue, 0), size_(ring_size) {
HSA::hsa_memory_register(this, sizeof(HostQueue));
MAKE_NAMED_SCOPE_GUARD(registerGuard,
[&]() { HSA::hsa_memory_deregister(this, sizeof(HostQueue)); });
@@ -93,6 +93,7 @@ HostQueue::HostQueue(hsa_region_t region, uint32_t ring_size, hsa_queue_type32_t
}
HostQueue::~HostQueue() {
HSA::hsa_memory_free(shared_queue_);
HSA::hsa_memory_free(ring_);
HSA::hsa_memory_deregister(this, sizeof(HostQueue));
}
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp
+13 -2
View File
@@ -742,8 +742,13 @@ hsa_status_t hsa_queue_create(
if (callback == nullptr) callback = core::Queue::DefaultErrorHandler;
uint64_t queue_create_flags = 0;
if (core::Runtime::runtime_singleton_->flag().dev_mem_queue_buf())
queue_create_flags = HSA_AMD_QUEUE_CREATE_DEVICE_MEM_RING_BUF;
core::Queue* cmd_queue = nullptr;
status = agent->QueueCreate(size, type, callback, data, private_segment_size,
status = agent->QueueCreate(size, type, queue_create_flags, callback, data, private_segment_size,
group_segment_size, &cmd_queue);
if (status != HSA_STATUS_SUCCESS) return status;
@@ -774,7 +779,13 @@ hsa_status_t hsa_soft_queue_create(hsa_region_t region, uint32_t size,
const core::Signal* signal = core::Signal::Convert(doorbell_signal);
IS_VALID(signal);
core::HostQueue* host_queue = new core::HostQueue(region, size, type, features, doorbell_signal);
void* shared_queue = nullptr;
hsa_status_t err = HSA::hsa_memory_allocate(region, sizeof(core::SharedQueue), &shared_queue);
if (err != HSA_STATUS_SUCCESS) return err;
assert(shared_queue && "Queue struct is NULL when creating host queue.");
core::HostQueue* host_queue = new core::HostQueue(static_cast<core::SharedQueue*>(shared_queue),
region, size, type, features, doorbell_signal);
*queue = core::Queue::Convert(host_queue);
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/intercept_queue.cpp
+3 -3
View File
@@ -256,7 +256,7 @@ uint64_t InterceptQueue::Submit(const AqlPacket* packets, uint64_t count) {
// Submit barrier which will wake async queue processing.
ring[barrier & mask].packet.body = {};
ring[barrier & mask].barrier_and.completion_signal = Signal::Convert(async_doorbell_);
if (Runtime::runtime_singleton_->flag().dev_mem_queue() && !needsPcieOrdering()) {
if (wrapped->IsDeviceMemRingBuf() && needsPcieOrdering()) {
// Ensure the packet body is written as header may get reordered when writing over PCIE
_mm_sfence();
}
@@ -303,7 +303,7 @@ uint64_t InterceptQueue::Submit(const AqlPacket* packets, uint64_t count) {
++packets_index;
}
if (write_index != 0) {
if (Runtime::runtime_singleton_->flag().dev_mem_queue() && !needsPcieOrdering()) {
if (wrapped->IsDeviceMemRingBuf() && needsPcieOrdering()) {
// Ensure the packet body is written as header may get reordered when writing over PCIE
_mm_sfence();
}
@@ -372,7 +372,7 @@ void InterceptQueue::StoreRelaxed(hsa_signal_value_t value) {
Cursor.pkt_index = i;
auto& handler = interceptors[Cursor.interceptor_index];
handler.first(&ring[i & mask], 1, i, handler.second, PacketWriter);
if (Runtime::runtime_singleton_->flag().dev_mem_queue() && !needsPcieOrdering()) {
if (IsDeviceMemRingBuf() && needsPcieOrdering()) {
// Ensure the packet body is written as header may get reordered when writing over PCIE
_mm_sfence();
}
projects/rocr-runtime/runtime/hsa-runtime/core/util/flag.h
+3 -3
View File
@@ -257,7 +257,7 @@ class Flag {
override_cpu_affinity_ = (var == "0") ? false : true;
var = os::GetEnvVar("HSA_ALLOCATE_QUEUE_DEV_MEM");
dev_mem_queue_ = (var == "1") ? true : false;
dev_mem_queue_buf_ = (var == "1") ? true : false;
var = os::GetEnvVar("HSA_WAIT_ANY_DEBUG");
wait_any_ = (var == "1") ? true : false;
@@ -396,7 +396,7 @@ class Flag {
size_t pc_sampling_max_device_buffer_size() const { return pc_sampling_max_device_buffer_size_; }
bool dev_mem_queue() const { return dev_mem_queue_; }
bool dev_mem_queue_buf() const { return dev_mem_queue_buf_; }
uint32_t signal_abort_timeout() const { return signal_abort_timeout_; }
@@ -433,7 +433,7 @@ class Flag {
bool enable_mwaitx_;
bool enable_ipc_mode_legacy_;
bool wait_any_;
bool dev_mem_queue_;
bool dev_mem_queue_buf_;
uint32_t signal_abort_timeout_;
int async_events_thread_priority_;
bool enable_3d_swizzle_ = false;
projects/rocr-runtime/runtime/hsa-runtime/inc/hsa_ext_amd.h
+26
View File
@@ -2805,6 +2805,32 @@ typedef enum hsa_amd_queue_priority_s {
hsa_status_t HSA_API hsa_amd_queue_set_priority(hsa_queue_t* queue,
hsa_amd_queue_priority_t priority);
/**
* @brief Queue creation attributes.
*/
typedef enum {
/**
* The queue's packet buffer and queue descriptor struct should be
* allocated in system memory (default). Mutually exclusive with
* HSA_AMD_QUEUE_CREATE_DEVICE_MEM_RING_BUF and
* HSA_AMD_QUEUE_CREATE_DEVICE_MEM_QUEUE_DESCRIPTOR.
*/
HSA_AMD_QUEUE_CREATE_SYSTEM_MEM = 0,
/**
* The queue's packet buffer should be allocated in the agent's
* fine-grain device memory region.
*/
HSA_AMD_QUEUE_CREATE_DEVICE_MEM_RING_BUF = (1 << 0),
/**
* The queue desciptor struct should be allocated in the agent's
* fine-grain device memory region. Not supported for devices
* connected via PCIe because the CPU's atomic read-modify-write
* operations cannot be promoted to PCIe atomic read-modify-write
* operations.
*/
HSA_AMD_QUEUE_CREATE_DEVICE_MEM_QUEUE_DESCRIPTOR = (1 << 1),
} hsa_amd_queue_create_flag_t;
/** @} */
/** \addtogroup memory Memory