From e63c280d4d3ccfb4b27e50bd919b7e78c654a433 Mon Sep 17 00:00:00 2001 From: Anusha GodavarthySurya Date: Wed, 19 Jul 2023 10:10:37 +0000 Subject: [PATCH] SWDEV-422207 - Capture AQL Packets for graph Kernel nodes during graph Inst. And enqueue AQL packet during launch Change-Id: I1e5f7f9e2a70bd500d190193cb6ba0867f5a63e7 --- hipamd/src/hip_graph.cpp | 4 + hipamd/src/hip_graph_internal.cpp | 91 +++++++++++++++--- hipamd/src/hip_graph_internal.hpp | 102 ++++++++++++++------ rocclr/device/device.hpp | 7 ++ rocclr/device/pal/palvirtual.hpp | 2 + rocclr/device/rocm/rocvirtual.cpp | 150 +++++++++--------------------- rocclr/device/rocm/rocvirtual.hpp | 8 +- rocclr/platform/command.cpp | 3 +- rocclr/platform/command.hpp | 38 +++++--- rocclr/utils/flags.hpp | 6 +- 10 files changed, 237 insertions(+), 174 deletions(-) diff --git a/hipamd/src/hip_graph.cpp b/hipamd/src/hip_graph.cpp index 55f35053c6..95aa925f61 100644 --- a/hipamd/src/hip_graph.cpp +++ b/hipamd/src/hip_graph.cpp @@ -1177,6 +1177,10 @@ hipError_t hipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph, hip::GraphExec* ge; hipError_t status = ihipGraphInstantiate(&ge, reinterpret_cast(graph)); *pGraphExec = reinterpret_cast(ge); + if (DEBUG_CLR_GRAPH_PACKET_CAPTURE) { + // For graph nodes capture AQL packets to dispatch them directly during graph launch. + status = ge->CaptureAQLPackets(); + } HIP_RETURN(status); } diff --git a/hipamd/src/hip_graph_internal.cpp b/hipamd/src/hip_graph_internal.cpp index 70fe657064..5f5f40d4b6 100644 --- a/hipamd/src/hip_graph_internal.cpp +++ b/hipamd/src/hip_graph_internal.cpp @@ -480,7 +480,7 @@ hipError_t GraphExec::CreateStreams(uint32_t num_streams) { } hipError_t GraphExec::Init() { - hipError_t status; + hipError_t status = hipSuccess; size_t min_num_streams = 1; for (auto& node : topoOrder_) { @@ -493,11 +493,63 @@ hipError_t GraphExec::Init() { return status; } +hipError_t GraphExec::CaptureAQLPackets() { + hipError_t status = hipSuccess; + size_t KernArgSizeForGraph = 0; + bool GraphHasOnlyKerns = true; + // GPU packet capture is enabled for kernel nodes. Calculate the kernel arg size required for all + // graph kernel nodes to allocate + for (const auto& list : parallelLists_) { + hip::Stream* stream = GetAvailableStreams(); + for (auto& node : list) { + node->SetStream(stream, this); + if (node->GetType() == hipGraphNodeTypeKernel) { + KernArgSizeForGraph += reinterpret_cast(node)->GetKerArgSize(); + } else { + GraphHasOnlyKerns = false; + } + } + } + + auto device = g_devices[ihipGetDevice()]->devices()[0]; + const auto& info = device->info(); + // Enable allocating kerns on device memory if graph as only kernels. memcpy nodes require hdp + // flush. ToDo: Work on enabling device kern args later for all type of nodes for large bar + if (GraphHasOnlyKerns == true && info.largeBar_) { + kernarg_pool_graph_ = reinterpret_cast
(device->deviceLocalAlloc(KernArgSizeForGraph)); + device_kernarg_pool_ = true; + } else { + kernarg_pool_graph_ = reinterpret_cast
( + device->hostAlloc(KernArgSizeForGraph, 0, amd::Device::MemorySegment::kKernArg)); + } + + if (kernarg_pool_graph_ == nullptr) { + return hipErrorMemoryAllocation; + } + kernarg_pool_size_graph_ = KernArgSizeForGraph; + + for (auto& node : topoOrder_) { + if (node->GetType() == hipGraphNodeTypeKernel) { + auto kernelnode = reinterpret_cast(node); + status = node->CreateCommand(node->GetQueue()); + // From the kernel pool allocate the kern arg size required for the current kernel node. + address kernArgOffset = allocKernArg(kernelnode->GetKernargSegmentByteSize(), + kernelnode->GetKernargSegmentAlignment()); + if (kernArgOffset == nullptr) { + return hipErrorMemoryAllocation; + } + // Enable GPU packet capture for the kernel node. + kernelnode->EnableCapturing(kernArgOffset); + } + } + return status; +} + hipError_t FillCommands(std::vector>& parallelLists, std::unordered_map>& nodeWaitLists, std::vector& topoOrder, Graph* clonedGraph, amd::Command*& graphStart, amd::Command*& graphEnd, hip::Stream* stream) { - hipError_t status; + hipError_t status = hipSuccess; for (auto& node : topoOrder) { // TODO: clone commands from next launch status = node->CreateCommand(node->GetQueue()); @@ -578,7 +630,7 @@ void UpdateStream(std::vector>& parallelLists, hip::Stream* st } hipError_t GraphExec::Run(hipStream_t stream) { - hipError_t status; + hipError_t status = hipSuccess; if (hip::getStream(stream) == nullptr) { return hipErrorInvalidResourceHandle; @@ -603,19 +655,30 @@ hipError_t GraphExec::Run(hipStream_t stream) { repeatLaunch_ = true; } if (parallelLists_.size() == 1) { + if (device_kernarg_pool_) { + // If kernelArgs are in device memory flush the HDP. + amd::Command* startCommand = nullptr; + startCommand = new amd::Marker(*hip_stream, false); + startCommand->enqueue(); + startCommand->release(); + } for (int i = 0; i < topoOrder_.size(); i++) { - topoOrder_[i]->SetStream(hip_stream, this); - status = topoOrder_[i]->CreateCommand(topoOrder_[i]->GetQueue()); - if (DEBUG_CLR_GRAPH_ENABLE_BUFFERING) { - // Enable buffering for graph with single branch - // 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(); - } + if (DEBUG_CLR_GRAPH_PACKET_CAPTURE && topoOrder_[i]->GetType() == hipGraphNodeTypeKernel) { + hip_stream->vdev()->dispatchAqlPacket(topoOrder_[i]->GetAqlPacket()); + } else { + topoOrder_[i]->SetStream(hip_stream, this); + status = topoOrder_[i]->CreateCommand(topoOrder_[i]->GetQueue()); + topoOrder_[i]->EnqueueCommands(stream); } - topoOrder_[i]->EnqueueCommands(stream); + } + if (DEBUG_CLR_GRAPH_PACKET_CAPTURE) { + amd::Command* endCommand = nullptr; + endCommand = new amd::Marker(*hip_stream, false); + // Since the end command is for graph completion tracking, + // it may not need release scopes + endCommand->setEventScope(amd::Device::kCacheStateIgnore); + endCommand->enqueue(); + endCommand->release(); } } else { UpdateStream(parallelLists_, hip_stream, this); diff --git a/hipamd/src/hip_graph_internal.hpp b/hipamd/src/hip_graph_internal.hpp index 005e1194be..c47e82e8d2 100644 --- a/hipamd/src/hip_graph_internal.hpp +++ b/hipamd/src/hip_graph_internal.hpp @@ -182,6 +182,7 @@ struct GraphNode : public hipGraphNodeDOTAttribute { static std::unordered_set nodeSet_; static amd::Monitor nodeSetLock_; unsigned int isEnabled_; + uint8_t gpuPacket_[64]; //!< GPU Packet to enqueue during graph launch public: GraphNode(hipGraphNodeType type, std::string style = "", std::string shape = "", @@ -229,7 +230,8 @@ struct GraphNode : public hipGraphNodeDOTAttribute { } return true; } - + // Return gpu packet address to update with actual packet under capture. + uint8_t* GetAqlPacket() { return gpuPacket_; } hip::Stream* GetQueue() { return stream_; } virtual void SetStream(hip::Stream* stream, GraphExec* ptr = nullptr) { @@ -336,11 +338,6 @@ struct GraphNode : public hipGraphNodeDOTAttribute { command->release(); } } - virtual void EnableBuffering() { - for (auto& command : commands_) { - command->setBufferingState(true); - } - } Graph* GetParentGraph() { return parentGraph_; } virtual Graph* GetChildGraph() { return nullptr; } void SetParentGraph(Graph* graph) { parentGraph_ = graph; } @@ -567,7 +564,11 @@ struct GraphExec { static amd::Monitor graphExecSetLock_; uint64_t flags_ = 0; bool repeatLaunch_ = false; - + // Graph Kernel arg vars + bool device_kernarg_pool_ = false; + address kernarg_pool_graph_ = nullptr; + uint32_t kernarg_pool_size_graph_ = 0; + uint32_t kernarg_pool_cur_graph_offset_ = 0; public: GraphExec(std::vector& topoOrder, std::vector>& lists, std::unordered_map>& nodeWaitLists, struct Graph*& clonedGraph, @@ -592,6 +593,11 @@ struct GraphExec { hip::Stream::Destroy(stream); } } + // Release the kernel arg memory. + auto device = g_devices[ihipGetDevice()]->devices()[0]; + if (DEBUG_CLR_GRAPH_PACKET_CAPTURE) { + device->hostFree(kernarg_pool_graph_, kernarg_pool_size_graph_); + } amd::ScopedLock lock(graphExecSetLock_); graphExecSet_.erase(this); delete clonedGraph_; @@ -606,7 +612,16 @@ struct GraphExec { } return clonedNode; } - + address allocKernArg(size_t size, size_t alignment) { + assert(alignment != 0); + address result = nullptr; + result = amd::alignUp(kernarg_pool_graph_ + kernarg_pool_cur_graph_offset_, alignment); + const size_t pool_new_usage = (result + size) - kernarg_pool_graph_; + if (pool_new_usage <= kernarg_pool_size_graph_) { + kernarg_pool_cur_graph_offset_ = pool_new_usage; + } + return result; + } // check executable graphs validity static bool isGraphExecValid(GraphExec* pGraphExec); @@ -617,6 +632,8 @@ struct GraphExec { hipError_t Init(); hipError_t CreateStreams(uint32_t num_streams); hipError_t Run(hipStream_t stream); + // Capture GPU Packets from graph commands + hipError_t CaptureAQLPackets(); }; struct ChildGraphNode : public GraphNode { @@ -742,31 +759,48 @@ struct ChildGraphNode : public GraphNode { }; class GraphKernelNode : public GraphNode { - hipKernelNodeParams kernelParams_; - unsigned int numParams_; - hipKernelNodeAttrValue kernelAttr_; - unsigned int kernelAttrInUse_; - ihipExtKernelEvents kernelEvents_; + hipKernelNodeParams kernelParams_; //!< Kernel node parameters + unsigned int numParams_; //!< No. of kernel params as part of signature + hipKernelNodeAttrValue kernelAttr_; //!< Kernel node attributes + unsigned int kernelAttrInUse_; //!< Kernel attributes in use + ihipExtKernelEvents kernelEvents_; //!< Events for Ext launch kernel + size_t alignedKernArgSize_; //!< Aligned size required for kernel args + size_t kernargSegmentByteSize_; //!< Kernel arg segment byte size + size_t kernargSegmentAlignment_; //!< Kernel arg segment alignment public: - void PrintAttributes(std::ostream& out, hipGraphDebugDotFlags flag) { - out << "["; - out << "style"; - out << "=\""; - out << style_; - (flag == hipGraphDebugDotFlagsKernelNodeParams || - flag == hipGraphDebugDotFlagsKernelNodeAttributes) ? - out << "\n" : out << "\""; - out << "shape"; - out << "=\""; - out << GetShape(flag); - out << "\""; - out << "label"; - out << "=\""; - out << GetLabel(flag); - out << "\""; - out << "];"; + size_t GetKerArgSize() const { return alignedKernArgSize_; } + size_t GetKernargSegmentByteSize() const { return kernargSegmentByteSize_; } + size_t GetKernargSegmentAlignment() const { return kernargSegmentAlignment_; } + void PrintAttributes(std::ostream& out, hipGraphDebugDotFlags flag) { + out << "["; + out << "style"; + out << "=\""; + out << style_; + (flag == hipGraphDebugDotFlagsKernelNodeParams || + flag == hipGraphDebugDotFlagsKernelNodeAttributes) ? + out << "\n" : out << "\""; + out << "shape"; + out << "=\""; + out << GetShape(flag); + out << "\""; + out << "label"; + out << "=\""; + out << GetLabel(flag); + out << "\""; + out << "];"; + } + + void EnableCapturing(address kernArgOffset) { + for (auto& command : commands_) { + reinterpret_cast(command)->setCapturingState( + true, GetAqlPacket(), kernArgOffset); + // Enqueue command to capture GPU Packet. Packet is not sent to hardware queue. + + command->submit(*(command->queue())->vdev()); + command->release(); } + } std::string GetLabel(hipGraphDebugDotFlags flag) { hipFunction_t func = getFunc(kernelParams_, ihipGetDevice()); @@ -842,6 +876,14 @@ class GraphKernelNode : public GraphNode { } hip::DeviceFunc* function = hip::DeviceFunc::asFunction(func); amd::Kernel* kernel = function->kernel(); + if (DEBUG_CLR_GRAPH_PACKET_CAPTURE) { + auto device = g_devices[ihipGetDevice()]->devices()[0]; + device::Kernel* devKernel = const_cast(kernel->getDeviceKernel(*device)); + kernargSegmentByteSize_ = devKernel->KernargSegmentByteSize(); + kernargSegmentAlignment_ = devKernel->KernargSegmentAlignment(); + alignedKernArgSize_ = + amd::alignUp(devKernel->KernargSegmentByteSize(), devKernel->KernargSegmentAlignment()); + } const amd::KernelSignature& signature = kernel->signature(); numParams_ = signature.numParameters(); diff --git a/rocclr/device/device.hpp b/rocclr/device/device.hpp index 86e3caff1a..d8dc24a22e 100644 --- a/rocclr/device/device.hpp +++ b/rocclr/device/device.hpp @@ -1279,6 +1279,7 @@ class VirtualDevice : public amd::HeapObject { //! Returns fence state of the VirtualGPU virtual bool isFenceDirty() const = 0; + virtual bool dispatchAqlPacket(uint8_t* aqlpacket) = 0; //! Resets fence state of the VirtualGPU virtual void resetFenceDirty() = 0; @@ -1733,6 +1734,12 @@ class Device : public RuntimeObject { return NULL; } + virtual void* deviceLocalAlloc(size_t size, bool atomics = false, + bool pseudo_fine_grain = false) const { + ShouldNotCallThis(); + return NULL; + } + virtual bool deviceAllowAccess(void* dst) const { ShouldNotCallThis(); return true; diff --git a/rocclr/device/pal/palvirtual.hpp b/rocclr/device/pal/palvirtual.hpp index 08bfa0e69a..f1c078d640 100644 --- a/rocclr/device/pal/palvirtual.hpp +++ b/rocclr/device/pal/palvirtual.hpp @@ -344,6 +344,8 @@ class VirtualGPU : public device::VirtualDevice { bool isFenceDirty() const { return false; } + inline bool dispatchAqlPacket(uint8_t* aqlpacket) { return false; } + void resetFenceDirty() {} //! Returns GPU device object associated with this kernel diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp index 2a6548bc2f..7f0f433e76 100644 --- a/rocclr/device/rocm/rocvirtual.cpp +++ b/rocclr/device/rocm/rocvirtual.cpp @@ -815,74 +815,6 @@ 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 + size - 1); - - aqlBuffer_.clear(); - } - return true; -} - // ================================================================================================ template bool VirtualGPU::dispatchGenericAqlPacket( @@ -990,36 +922,23 @@ void VirtualGPU::dispatchBlockingWait() { } } } - // ================================================================================================ bool VirtualGPU::dispatchAqlPacket(hsa_kernel_dispatch_packet_t* packet, uint16_t header, - uint16_t rest, bool blocking, bool buffering) { - static size_t initialAQLBufferSize = 1; + uint16_t rest, bool blocking, bool capturing, + const uint8_t* aqlPacket) { dispatchBlockingWait(); - 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; + if (capturing == true) { + packet->header = header; + packet->setup = rest; + if (timestamp_ != nullptr) { + // Get active signal for current dispatch if profiling is necessary + packet->completion_signal = Barriers().ActiveSignal(kInitSignalValueOne, timestamp_); } - } 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); + amd::Os::fastMemcpy(const_cast(aqlPacket), packet, + sizeof(hsa_kernel_dispatch_packet_t)); + return true; } 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(); + return dispatchGenericAqlPacket(packet, header, rest, blocking); } } // ================================================================================================ @@ -1129,6 +1048,18 @@ void VirtualGPU::dispatchBarrierPacket(uint16_t packetHeader, bool skipSignal, barrier_packet_.dep_signal[4] = hsa_signal_t{}; } +inline bool VirtualGPU::dispatchAqlPacket(uint8_t* aqlpacket) { + auto packet = reinterpret_cast(aqlpacket); + // If rocprof tracing is enabled, store the correlation ID in the dispatch packet. + // The profiler can retrieve this correlation ID to attribute waves to specific dispatch + // locations. + if (activity_prof::IsEnabled(OP_ID_DISPATCH)) { + packet->reserved2 = activity_prof::correlation_id; + } + dispatchGenericAqlPacket(packet, packet->header, packet->setup, false); + return true; +} + // ================================================================================================ void VirtualGPU::dispatchBarrierValuePacket(uint16_t packetHeader, bool resolveDepSignal, hsa_signal_t signal, hsa_signal_value_t value, @@ -1449,15 +1380,9 @@ void* VirtualGPU::allocKernArg(size_t size, size_t alignment) { kernarg_pool_cur_offset_ = pool_new_usage; return result; } else { - //! We run out of the arguments space! //! 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 @@ -3196,8 +3121,12 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, // Find all parameters for the current kernel if (!kernel.parameters().deviceKernelArgs() || gpuKernel.isInternalKernel()) { // Allocate buffer to hold kernel arguments - argBuffer = reinterpret_cast
(allocKernArg(gpuKernel.KernargSegmentByteSize(), + if(vcmd != nullptr && vcmd->getCapturingState()) { + argBuffer = vcmd->getKernArgOffset(); + } else { + argBuffer = reinterpret_cast
(allocKernArg(gpuKernel.KernargSegmentByteSize(), gpuKernel.KernargSegmentAlignment())); + } // Load all kernel arguments nontemporalMemcpy(argBuffer, parameters, std::min(gpuKernel.KernargSegmentByteSize(), @@ -3272,13 +3201,20 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE); aql_packet->setup = sizes.dimensions() << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS; } - - // Dispatch the packet - if (!dispatchAqlPacket(&dispatchPacket, aqlHeaderWithOrder, - (sizes.dimensions() << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS), - GPU_FLUSH_ON_EXECUTION, - (vcmd != nullptr) ? vcmd->getBufferingState() : false)) { - return false; + if (vcmd == nullptr) { + // Dispatch the packet + if (!dispatchAqlPacket(&dispatchPacket, aqlHeaderWithOrder, + (sizes.dimensions() << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS), + GPU_FLUSH_ON_EXECUTION)) { + return false; + } + } else { + if (!dispatchAqlPacket(&dispatchPacket, aqlHeaderWithOrder, + (sizes.dimensions() << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS), + GPU_FLUSH_ON_EXECUTION, vcmd->getCapturingState(), + vcmd->getAqlPacket())) { + return false; + } } } diff --git a/rocclr/device/rocm/rocvirtual.hpp b/rocclr/device/rocm/rocvirtual.hpp index 107dd56976..cc90df0e06 100644 --- a/rocclr/device/rocm/rocvirtual.hpp +++ b/rocclr/device/rocm/rocvirtual.hpp @@ -424,9 +424,10 @@ 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, bool buffering = false); + inline bool dispatchAqlPacket(uint8_t* aqlpacket); + bool dispatchAqlPacket(hsa_kernel_dispatch_packet_t* packet, uint16_t header, uint16_t rest, + bool blocking = true, bool capturing = false, + const uint8_t* aqlPacket = nullptr); 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, @@ -566,6 +567,5 @@ 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 d67f951e29..06e6153fad 100644 --- a/rocclr/platform/command.cpp +++ b/rocclr/platform/command.cpp @@ -317,7 +317,6 @@ Command::Command(HostQueue& queue, cl_command_type type, const EventWaitList& ev type_(type), data_(nullptr), waitingEvent_(waitingEvent), - buffering_(false), eventWaitList_(eventWaitList), commandWaitBits_(commandWaitBits) { // Retain the commands from the event wait list. @@ -354,7 +353,7 @@ void Command::enqueue() { // Notify all commands about the waiter. Barrier will be sent in order to obtain // HSA signal for a wait on the current queue - for (const auto &event: eventWaitList()) { + for (const auto& event : eventWaitList()) { event->notifyCmdQueue(!kCpuWait); } diff --git a/rocclr/platform/command.hpp b/rocclr/platform/command.hpp index ddb9a9f170..6861546951 100644 --- a/rocclr/platform/command.hpp +++ b/rocclr/platform/command.hpp @@ -251,13 +251,12 @@ union CopyMetadata { */ class Command : public Event { private: - HostQueue* queue_; //!< The command queue this command is enqueue into - Command* next_; //!< Next GPU command in the queue list - Command* batch_head_ = nullptr; //!< The head of the batch commands - cl_command_type type_; //!< This command's OpenCL type. + HostQueue* queue_; //!< The command queue this command is enqueue into + Command* next_; //!< Next GPU command in the queue list + Command* batch_head_ = nullptr; //!< The head of the batch commands + 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 + const Event* waitingEvent_; //!< Waiting event associated with the marker protected: bool cpu_wait_ = false; //!< If true, then the command was issued for CPU/GPU sync @@ -281,7 +280,6 @@ class Command : public Event { type_(type), data_(nullptr), waitingEvent_(nullptr), - buffering_(false), eventWaitList_(nullWaitList), commandWaitBits_(0) {} @@ -296,12 +294,6 @@ 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_; } @@ -1083,6 +1075,10 @@ class NDRangeKernelCommand : public Command { uint32_t firstDevice_; //!< Device index of the first device in the gridc uint32_t numWorkgroups_; //!< Total number of workgroups in the current launch + bool capturing_ = false; //!< Flag to enable/disable graph gpu packet capture + uint8_t* gpuPacket_ = nullptr; //!< GPU packet to capture, when graph capturing is enabled + address kernArgOffset_ = nullptr; //!< KernelArg buffer to used when graph capturing is enabled + public: enum { CooperativeGroups = 0x01, @@ -1090,6 +1086,22 @@ class NDRangeKernelCommand : public Command { AnyOrderLaunch = 0x04, }; + //! Returns AQL buffer state + bool getCapturingState() const { return capturing_; } + + //! Sets AQL capture state, aql packet to capture and where to copy kernArgs + void setCapturingState(bool state, uint8_t* packet, address kernArgOffset) { + capturing_ = state; + gpuPacket_ = packet; + kernArgOffset_ = kernArgOffset; + } + + //! returns the graph executable object command belongs to. + const uint8_t* getAqlPacket() const { return gpuPacket_; } + + //! returns the graph executable object command belongs to. + const address getKernArgOffset() const { return kernArgOffset_; } + //! Construct an ExecuteKernel command NDRangeKernelCommand(HostQueue& queue, const EventWaitList& eventWaitList, Kernel& kernel, const NDRangeContainer& sizes, uint32_t sharedMemBytes = 0, diff --git a/rocclr/utils/flags.hpp b/rocclr/utils/flags.hpp index cf6917541d..2c7f698bf9 100644 --- a/rocclr/utils/flags.hpp +++ b/rocclr/utils/flags.hpp @@ -230,10 +230,8 @@ release(size_t, HIP_INITIAL_DM_SIZE, 8 * Mi, \ "Set initial heap size for device malloc.") \ release(bool, HIP_FORCE_DEV_KERNARG, 0, \ "Force device mem for kernel args.") \ -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") \ +release(bool, DEBUG_CLR_GRAPH_PACKET_CAPTURE, false, \ + "Enable/Disable graph packet capturing") \ release(cstring, HIPRTC_COMPILE_OPTIONS_APPEND, "", \ "Set compile options needed for hiprtc compilation") \ release(cstring, HIPRTC_LINK_OPTIONS_APPEND, "", \