diff --git a/hipamd/src/hip_graph_internal.cpp b/hipamd/src/hip_graph_internal.cpp index 9e7627974b..67fa238ce9 100644 --- a/hipamd/src/hip_graph_internal.cpp +++ b/hipamd/src/hip_graph_internal.cpp @@ -494,51 +494,65 @@ hipError_t GraphExec::Init() { 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; + if (parallelLists_.size() == 1) { + size_t kernArgSizeForGraph = 0; + // 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(); + } } } - } - 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)); - } + auto device = g_devices[ihipGetDevice()]->devices()[0]; + if (device->info().largeBar_) { + // Pad kernel argument buffer with sentinal size bytes to do a readback later + kernArgSizeForGraph += sizeof(int); + 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; + 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; + 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; + } + // Form GPU packet capture for the kernel node. + kernelNode->CaptureAndFormPacket(kernArgOffset); + } + } + + if (device_kernarg_pool_) { + // Write HDP_MEM_COHERENCY_FLUSH_CNTL reg to initiate flush read to HDP mem. Verify mem + // by readback of sentinal value at the tail end of the kernarg surface (allocated above) + // This needs to be done for PCIE connected devices only. HDP path is disabled for XGMI + // between CPU<->GPU + if (!device->isXgmi()) { + static int host_val = 1; + address dev_ptr = kernarg_pool_graph_ + kernarg_pool_size_graph_ - sizeof(int); + *dev_ptr = host_val; + *device->info().hdpMemFlushCntl = 1; + while (*dev_ptr != host_val); + host_val++; } - // Enable GPU packet capture for the kernel node. - kernelnode->EnableCapturing(kernArgOffset); } } return status; @@ -561,9 +575,11 @@ hipError_t FillCommands(std::vector>& parallelLists, } node->UpdateEventWaitLists(waitList); } + std::vector rootNodes = clonedGraph->GetRootNodes(); ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] RootCommand get launched on stream (stream:%p)\n", stream); + for (auto& root : rootNodes) { //If rootnode is launched on to the same stream dont add dependency if (root->GetQueue() != stream) { @@ -653,14 +669,16 @@ hipError_t GraphExec::Run(hipStream_t stream) { } else { repeatLaunch_ = true; } + if (parallelLists_.size() == 1) { if (device_kernarg_pool_) { - // If kernelArgs are in device memory flush the HDP. + // If kernelArgs are in device memory flush/invalidate L2 amd::Command* startCommand = nullptr; startCommand = new amd::Marker(*hip_stream, false); startCommand->enqueue(); startCommand->release(); } + for (int i = 0; i < topoOrder_.size(); i++) { if (DEBUG_CLR_GRAPH_PACKET_CAPTURE && topoOrder_[i]->GetType() == hipGraphNodeTypeKernel) { hip_stream->vdev()->dispatchAqlPacket(topoOrder_[i]->GetAqlPacket()); @@ -670,6 +688,7 @@ hipError_t GraphExec::Run(hipStream_t stream) { topoOrder_[i]->EnqueueCommands(stream); } } + if (DEBUG_CLR_GRAPH_PACKET_CAPTURE) { amd::Command* endCommand = nullptr; endCommand = new amd::Marker(*hip_stream, false); diff --git a/hipamd/src/hip_graph_internal.hpp b/hipamd/src/hip_graph_internal.hpp index eaef79bf5b..86f3544344 100644 --- a/hipamd/src/hip_graph_internal.hpp +++ b/hipamd/src/hip_graph_internal.hpp @@ -463,6 +463,7 @@ struct Graph { graphExeUserObjs.insert(userObj); } } + Graph* clone(std::unordered_map& clonedNodes) const; Graph* clone() const; void GenerateDOT(std::ostream& fout, hipGraphDebugDotFlags flag) { @@ -612,6 +613,7 @@ struct GraphExec { } return clonedNode; } + address allocKernArg(size_t size, size_t alignment) { assert(alignment != 0); address result = nullptr; @@ -622,12 +624,19 @@ struct GraphExec { } return result; } + // check executable graphs validity static bool isGraphExecValid(GraphExec* pGraphExec); std::vector& GetNodes() { return topoOrder_; } - hip::Stream* GetAvailableStreams() { return parallel_streams_[currentQueueIndex_++]; } + hip::Stream* GetAvailableStreams() { + if (currentQueueIndex_ < parallel_streams_.size()) { + return parallel_streams_[currentQueueIndex_++]; + } + return nullptr; + } + void ResetQueueIndex() { currentQueueIndex_ = 0; } hipError_t Init(); hipError_t CreateStreams(uint32_t num_streams); @@ -791,12 +800,12 @@ class GraphKernelNode : public GraphNode { out << "];"; } - void EnableCapturing(address kernArgOffset) { + void CaptureAndFormPacket(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. + // Enqueue command to capture GPU Packet. Packet is not sent to hardware queue. command->submit(*(command->queue())->vdev()); command->release(); } diff --git a/rocclr/device/device.hpp b/rocclr/device/device.hpp index 69de8521df..81bf3d0ed9 100644 --- a/rocclr/device/device.hpp +++ b/rocclr/device/device.hpp @@ -1758,6 +1758,11 @@ class Device : public RuntimeObject { return NULL; } + virtual bool isXgmi() const { + ShouldNotCallThis(); + return false; + } + virtual bool deviceAllowAccess(void* dst) const { ShouldNotCallThis(); return true; diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp index 5ab54a02d9..873dacf03b 100644 --- a/rocclr/device/rocm/rocvirtual.cpp +++ b/rocclr/device/rocm/rocvirtual.cpp @@ -521,7 +521,6 @@ std::vector& VirtualGPU::HwQueueTracker::WaitingSignal(HwQueueEngi // Validate all signals for the wait and skip already completed for (uint32_t i = 0; i < external_signals_.size(); ++i) { // Early signal status check - if (hsa_signal_load_relaxed(external_signals_[i]->signal_) > 0) { const Settings& settings = gpu_.dev().settings(); // Actively wait on CPU to avoid extra overheads of signal tracking on GPU. @@ -829,6 +828,13 @@ bool VirtualGPU::dispatchGenericAqlPacket( // 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_); + if (addSystemScope_) { + header &= ~(HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE | + HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE); + header |= (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE | + HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE); + addSystemScope_ = false; + } auto expected_fence_state = extractAqlBits(header, HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE, HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE); @@ -898,6 +904,10 @@ bool VirtualGPU::dispatchGenericAqlPacket( hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index - 1); + // Mark the flag indicating if a dispatch is outstanding. + // We are not waiting after every dispatch. + hasPendingDispatch_ = true; + // Wait on signal ? if (blocking) { LogInfo("Runtime reachead the AQL queue limit. SW is much ahead of HW. Blocking AQL queue!"); @@ -929,18 +939,14 @@ void VirtualGPU::dispatchBlockingWait() { bool VirtualGPU::dispatchAqlPacket(hsa_kernel_dispatch_packet_t* packet, uint16_t header, uint16_t rest, bool blocking, bool capturing, const uint8_t* aqlPacket) { - dispatchBlockingWait(); 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_); - } amd::Os::fastMemcpy(const_cast(aqlPacket), packet, sizeof(hsa_kernel_dispatch_packet_t)); return true; } else { + dispatchBlockingWait(); return dispatchGenericAqlPacket(packet, header, rest, blocking); } } @@ -1053,12 +1059,16 @@ void VirtualGPU::dispatchBarrierPacket(uint16_t packetHeader, bool skipSignal, } inline bool VirtualGPU::dispatchAqlPacket(uint8_t* aqlpacket) { + dispatchBlockingWait(); 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)) { + if (activity_prof::IsEnabled(OP_ID_DISPATCH) || + (roc_device_.info().queueProperties_ & CL_QUEUE_PROFILING_ENABLE)) { packet->reserved2 = activity_prof::correlation_id; + // Get active signal for current dispatch if profiling is necessary + packet->completion_signal = Barriers().ActiveSignal(kInitSignalValueOne, timestamp_); } dispatchGenericAqlPacket(packet, packet->header, packet->setup, false); return true; @@ -3227,13 +3237,8 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, aqlHeaderWithOrder &= kAqlHeaderMask; } - if (addSystemScope_ || (vcmd != nullptr && - vcmd->getEventScope() == amd::Device::kCacheStateSystem)) { - aqlHeaderWithOrder &= ~(HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE | - HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE); - aqlHeaderWithOrder |= (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE | - HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE); - addSystemScope_ = false; + if (vcmd != nullptr && vcmd->getEventScope() == amd::Device::kCacheStateSystem) { + addSystemScope_ = true; } // If profiling is enabled, store the correlation ID in the dispatch packet. The profiler can @@ -3268,10 +3273,6 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, } } - // Mark the flag indicating if a dispatch is outstanding. - // We are not waiting after every dispatch. - hasPendingDispatch_ = true; - // Output printf buffer if (!printfDbg()->output(*this, printfEnabled, gpuKernel.printfInfo())) { LogError("\nCould not print data from the printf buffer!");