From b0e6f99ad7620af5eede5ea472c9543407cc3a73 Mon Sep 17 00:00:00 2001 From: Anusha GodavarthySurya Date: Sat, 17 Jun 2023 13:48:05 +0000 Subject: [PATCH] SWDEV-392732 - Initial commit for graph doorbell optimization(AQL Buffering) Change-Id: I451725006c54c249dc530c55d2af2a31594bf49b --- hipamd/src/hip_graph_internal.cpp | 16 ++++- hipamd/src/hip_graph_internal.hpp | 5 ++ rocclr/device/rocm/rocvirtual.cpp | 114 +++++++++++++++++++++++++++--- rocclr/device/rocm/rocvirtual.hpp | 5 +- rocclr/platform/command.cpp | 11 +-- rocclr/platform/command.hpp | 10 ++- rocclr/utils/flags.hpp | 7 +- 7 files changed, 149 insertions(+), 19 deletions(-) diff --git a/hipamd/src/hip_graph_internal.cpp b/hipamd/src/hip_graph_internal.cpp index f98eca1c05..cd8bf3c657 100644 --- a/hipamd/src/hip_graph_internal.cpp +++ b/hipamd/src/hip_graph_internal.cpp @@ -611,8 +611,20 @@ hipError_t hipGraphExec::Run(hipStream_t stream) { cmd->enqueue(); cmd->release(); } - for (auto& node : topoOrder_) { - node->EnqueueCommands(stream); + for (int i = 0; i < topoOrder_.size(); i++) { + if (DEBUG_CLR_GRAPH_ENABLE_BUFFERING) { + // Enable buffering for graph with single branch + if (parallelLists_.size() == 1) { + // Peep through the next node. If current and next node are kernel then enable AQL + // buffering + if (((i + 1) != topoOrder_.size()) && + topoOrder_[i]->GetType() == hipGraphNodeTypeKernel && + topoOrder_[i + 1]->GetType() == hipGraphNodeTypeKernel) { + topoOrder_[i]->EnableBuffering(); + } + } + } + topoOrder_[i]->EnqueueCommands(stream); } if (endCommand != nullptr) { endCommand->enqueue(); diff --git a/hipamd/src/hip_graph_internal.hpp b/hipamd/src/hip_graph_internal.hpp index 1c2b20cac0..709a1d8457 100644 --- a/hipamd/src/hip_graph_internal.hpp +++ b/hipamd/src/hip_graph_internal.hpp @@ -326,6 +326,11 @@ struct hipGraphNode : public hipGraphNodeDOTAttribute { command->release(); } } + virtual void EnableBuffering() { + for (auto& command : commands_) { + command->setBufferingState(true); + } + } ihipGraph* GetParentGraph() { return parentGraph_; } virtual ihipGraph* GetChildGraph() { return nullptr; } void SetParentGraph(ihipGraph* graph) { parentGraph_ = graph; } diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp index ed46499988..dbac1d3df3 100644 --- a/rocclr/device/rocm/rocvirtual.cpp +++ b/rocclr/device/rocm/rocvirtual.cpp @@ -810,6 +810,74 @@ static inline void packet_store_release(uint32_t* packet, uint16_t header, uint1 __atomic_store_n(packet, header | (rest << 16), __ATOMIC_RELEASE); } +bool VirtualGPU::dispatchAqlBuffer() { + size_t size = aqlBuffer_.size(); + if (size > 0) { + const uint32_t queueSize = gpu_queue_->size; + const uint32_t queueMask = queueSize - 1; + 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_, size); + uint64_t read = hsa_queue_load_read_index_relaxed(gpu_queue_); + + while (index - hsa_queue_load_read_index_scacquire(gpu_queue_) >= sw_queue_size - size) { + amd::Os::yield(); + } + + if (timestamp_ != nullptr) { + for (uint i = 0; i < size; i++) { + // Get active signal for current dispatch if profiling is necessary + aqlBuffer_[i].completion_signal = Barriers().ActiveSignal(kInitSignalValueOne, timestamp_); + } + } + for (uint i = 0; i < size; i++) { + ClPrint( + amd::LOG_DEBUG, amd::LOG_AQL, + "HWq=0x%zx, Dispatch AQL Buffer Header = " + "0x%x (type=%d, barrier=%d, acquire=%d, release=%d), " + "setup=%d, grid=[%zu, %zu, %zu], workgroup=[%zu, %zu, %zu], private_seg_size=%zu, " + "group_seg_size=%zu, kernel_obj=0x%zx, kernarg_address=0x%zx, completion_signal=0x%zx", + gpu_queue_->base_address, aqlBuffer_[i].header, + extractAqlBits(aqlBuffer_[i].header, HSA_PACKET_HEADER_TYPE, + HSA_PACKET_HEADER_WIDTH_TYPE), + extractAqlBits(aqlBuffer_[i].header, HSA_PACKET_HEADER_BARRIER, + HSA_PACKET_HEADER_WIDTH_BARRIER), + extractAqlBits(aqlBuffer_[i].header, HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE, + HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE), + extractAqlBits(aqlBuffer_[i].header, HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE, + HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE), + aqlBuffer_[i].setup, (aqlBuffer_[i]).grid_size_x, (aqlBuffer_[i]).grid_size_y, + (aqlBuffer_[i]).grid_size_z, (aqlBuffer_[i]).workgroup_size_x, + (aqlBuffer_[i]).workgroup_size_y, (aqlBuffer_[i]).workgroup_size_z, + (aqlBuffer_[i]).private_segment_size, (aqlBuffer_[i]).group_segment_size, + (aqlBuffer_[i]).kernel_object, (aqlBuffer_[i]).kernarg_address, + (aqlBuffer_[i]).completion_signal); + } + uint16_t firstPacketHeader = aqlBuffer_.front().header; + aqlBuffer_.front().header = kInvalidAql; + hsa_kernel_dispatch_packet_t* aql_loc = + &((hsa_kernel_dispatch_packet_t*)(gpu_queue_->base_address))[index & queueMask]; + size_t size_before_wrap = queueSize - (index % queueSize); + if (size_before_wrap < size) { + amd::Os::fastMemcpy(aql_loc, &aqlBuffer_[0], sizeof(hsa_kernel_dispatch_packet_t) * size_before_wrap); + hsa_kernel_dispatch_packet_t* aql_loc_0 = + &((hsa_kernel_dispatch_packet_t*)(gpu_queue_->base_address))[0]; + amd::Os::fastMemcpy(aql_loc_0, &aqlBuffer_[size_before_wrap], + sizeof(hsa_kernel_dispatch_packet_t) * (size - size_before_wrap)); + } else { + amd::Os::fastMemcpy(aql_loc, &aqlBuffer_[0], sizeof(hsa_kernel_dispatch_packet_t) * size); + } + + packet_store_release(reinterpret_cast(aql_loc), firstPacketHeader, + aqlBuffer_.front().setup); + hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index); + + aqlBuffer_.clear(); + } + return true; +} + // ================================================================================================ template bool VirtualGPU::dispatchGenericAqlPacket( @@ -919,13 +987,36 @@ void VirtualGPU::dispatchBlockingWait() { } // ================================================================================================ -bool VirtualGPU::dispatchAqlPacket( - hsa_kernel_dispatch_packet_t* packet, uint16_t header, uint16_t rest, bool blocking) { +bool VirtualGPU::dispatchAqlPacket(hsa_kernel_dispatch_packet_t* packet, uint16_t header, + uint16_t rest, bool blocking, bool buffering) { + static size_t initialAQLBufferSize = 1; dispatchBlockingWait(); - - return dispatchGenericAqlPacket(packet, header, rest, blocking); + if (buffering == true) { + aqlBuffer_.push_back(*packet); + aqlBuffer_.back().header = header; + aqlBuffer_.back().setup = rest; + if (!(aqlBuffer_.size() >= + initialAQLBufferSize)) { // Buffer maximum of AQL Buffer size packets once it + // exceeds send them for dispatch + return true; + } + } else if (!aqlBuffer_.empty()) { // If buffering is disabled and AQLBuffer is not empty then + // make sure current packet is added to the buffer for dispatch + aqlBuffer_.push_back(*packet); + aqlBuffer_.back().header = header; + aqlBuffer_.back().setup = rest; + } + if (aqlBuffer_.empty()) { + return dispatchGenericAqlPacket(packet, header, rest, blocking); + } else { + ClPrint(amd::LOG_DEBUG, amd::LOG_CODE, "Dispath AQL Buffer size:%ld", aqlBuffer_.size()); + // Increment buffer size ^2 until DEBUG_CLR_GRAPH_MAX_AQL_BUFFER_SIZE + if (initialAQLBufferSize < DEBUG_CLR_GRAPH_MAX_AQL_BUFFER_SIZE) { + initialAQLBufferSize = initialAQLBufferSize << 1; + } + return dispatchAqlBuffer(); + } } - // ================================================================================================ bool VirtualGPU::dispatchAqlPacket( hsa_barrier_and_packet_t* packet, uint16_t header, uint16_t rest, bool blocking) { @@ -1355,6 +1446,11 @@ void* VirtualGPU::allocKernArg(size_t size, size_t alignment) { //! That means the app didn't call clFlush/clFinish for very long time. // Reset the signal for the barrier packet hsa_signal_silent_store_relaxed(kernarg_pool_signal_[active_chunk_], kInitSignalValueOne); + // dispatch any buffered AQL packets + bool status = dispatchAqlBuffer(); + if (!status) { + LogError("dispatch Aql Buffer failed!"); + } // Dispatch a barrier packet into the queue dispatchBarrierPacket(kBarrierPacketHeader, true, kernarg_pool_signal_[active_chunk_]); // Get the next chunk @@ -3171,10 +3267,10 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, } // Dispatch the packet - if (!dispatchAqlPacket( - &dispatchPacket, aqlHeaderWithOrder, - (sizes.dimensions() << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS), - GPU_FLUSH_ON_EXECUTION)) { + if (!dispatchAqlPacket(&dispatchPacket, aqlHeaderWithOrder, + (sizes.dimensions() << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS), + GPU_FLUSH_ON_EXECUTION, + (vcmd != nullptr) ? vcmd->getBufferingState() : false)) { return false; } } diff --git a/rocclr/device/rocm/rocvirtual.hpp b/rocclr/device/rocm/rocvirtual.hpp index 2197f8996f..e34c075625 100644 --- a/rocclr/device/rocm/rocvirtual.hpp +++ b/rocclr/device/rocm/rocvirtual.hpp @@ -424,13 +424,15 @@ class VirtualGPU : public device::VirtualDevice { //! Dispatches a barrier with blocking HSA signals void dispatchBlockingWait(); + bool dispatchAqlBuffer(); bool dispatchAqlPacket(hsa_kernel_dispatch_packet_t* packet, uint16_t header, - uint16_t rest, bool blocking = true); + uint16_t rest, bool blocking = true, bool buffering = false); bool dispatchAqlPacket(hsa_barrier_and_packet_t* packet, uint16_t header, uint16_t rest, bool blocking = true); template bool dispatchGenericAqlPacket(AqlPacket* packet, uint16_t header, uint16_t rest, bool blocking, size_t size = 1); + void dispatchBarrierPacket(uint16_t packetHeader, bool skipSignal = false, hsa_signal_t signal = hsa_signal_t{0}); bool dispatchCounterAqlPacket(hsa_ext_amd_aql_pm4_packet_t* packet, const uint32_t gfxVersion, @@ -564,5 +566,6 @@ class VirtualGPU : public device::VirtualDevice { int fence_state_; //!< Fence scope //!< kUnknown/kFlushedToDevice/kFlushedToSystem bool fence_dirty_; //!< Fence modified flag + std::vector aqlBuffer_; //!< AQL packet buffer for graphs }; } diff --git a/rocclr/platform/command.cpp b/rocclr/platform/command.cpp index 35ec511329..0371d8dc57 100644 --- a/rocclr/platform/command.cpp +++ b/rocclr/platform/command.cpp @@ -310,15 +310,18 @@ bool Event::notifyCmdQueue(bool cpu_wait) { const Event::EventWaitList Event::nullWaitList(0); // ================================================================================================ -Command::Command(HostQueue& queue, cl_command_type type, - const EventWaitList& eventWaitList, uint32_t commandWaitBits, const Event* waitingEvent) - : Event(queue, activity_prof::IsEnabled(activity_prof::OperationId(type)) || - queue.properties().test(CL_QUEUE_PROFILING_ENABLE) || Agent::shouldPostEventEvents()), +Command::Command(HostQueue& queue, cl_command_type type, const EventWaitList& eventWaitList, + uint32_t commandWaitBits, const Event* waitingEvent) + : Event(queue, + activity_prof::IsEnabled(activity_prof::OperationId(type)) || + queue.properties().test(CL_QUEUE_PROFILING_ENABLE) || + Agent::shouldPostEventEvents()), queue_(&queue), next_(nullptr), type_(type), data_(nullptr), waitingEvent_(waitingEvent), + buffering_(false), eventWaitList_(eventWaitList), commandWaitBits_(commandWaitBits) { // Retain the commands from the event wait list. diff --git a/rocclr/platform/command.hpp b/rocclr/platform/command.hpp index 97c501c901..b45c9067ae 100644 --- a/rocclr/platform/command.hpp +++ b/rocclr/platform/command.hpp @@ -270,6 +270,7 @@ class Command : public Event { cl_command_type type_; //!< This command's OpenCL type. void* data_; const Event* waitingEvent_; //!< Waiting event associated with the marker + bool buffering_; //!< Flag to enable/disable AQL buffering protected: bool cpu_wait_ = false; //!< If true, then the command was issued for CPU/GPU sync @@ -278,7 +279,7 @@ class Command : public Event { EventWaitList eventWaitList_; //! Force await completion of previous command - //! 0x1 - wait before enqueue, 0x2 - wait after, 0x3 - wait both. + //! 0x1 - wait before enqueue, 0x2 - wait after, 0x3 - wait both uint32_t commandWaitBits_; //! Construct a new command of the given OpenCL type. @@ -293,6 +294,7 @@ class Command : public Event { type_(type), data_(nullptr), waitingEvent_(nullptr), + buffering_(false), eventWaitList_(nullWaitList), commandWaitBits_(0) {} @@ -307,6 +309,12 @@ class Command : public Event { } public: + //! Returns AQL buffer state + bool getBufferingState() const { return buffering_; } + + //! Sets AQL buffer state + void setBufferingState(bool state) { buffering_ = state; } + //! Return the queue this command is enqueued into. HostQueue* queue() const { return queue_; } diff --git a/rocclr/utils/flags.hpp b/rocclr/utils/flags.hpp index 54fd585460..926d63bf0e 100644 --- a/rocclr/utils/flags.hpp +++ b/rocclr/utils/flags.hpp @@ -287,8 +287,11 @@ release(size_t, HIP_INITIAL_DM_SIZE, 8 * Mi, \ release(bool, HIP_FORCE_DEV_KERNARG, 0, \ "Force device mem for kernel args.") \ release(bool, DEBUG_CLR_USE_SDMA_QUERY, 0, \ - "Use SDMA query API to make copy decisions.") - + "Use SDMA query API to make copy decisions.") \ +release(uint, DEBUG_CLR_GRAPH_MAX_AQL_BUFFER_SIZE, 32, \ + "Size of AQL buffering queue") \ +release(bool, DEBUG_CLR_GRAPH_ENABLE_BUFFERING, false, \ + "Enable/Disable graph AQL buffering") namespace amd { extern bool IS_HIP;