diff --git a/projects/clr/hipamd/src/hip_context.cpp b/projects/clr/hipamd/src/hip_context.cpp index 3609f0312b..f9cd781639 100644 --- a/projects/clr/hipamd/src/hip_context.cpp +++ b/projects/clr/hipamd/src/hip_context.cpp @@ -128,17 +128,17 @@ hip::Stream* getStream(hipStream_t stream, bool wait) { } // ================================================================================================ -hip::Stream* getNullStream(amd::Context& ctx) { +hip::Stream* getNullStream(amd::Context& ctx, bool wait) { for (auto& it : g_devices) { if (it->asContext() == &ctx) { - return it->NullStream(); + return it->NullStream(wait); } } // If it's a pure SVM allocation with system memory access, then it shouldn't matter which device // runtime selects by default if (hip::host_context == &ctx) { // Return current... - return getNullStream(); + return getNullStream(wait); } return nullptr; } diff --git a/projects/clr/hipamd/src/hip_graph.cpp b/projects/clr/hipamd/src/hip_graph.cpp index c36e7c1d19..d31ef055e3 100644 --- a/projects/clr/hipamd/src/hip_graph.cpp +++ b/projects/clr/hipamd/src/hip_graph.cpp @@ -85,13 +85,7 @@ hipError_t ihipGraphAddKernelNode(hip::GraphNode** pGraphNode, hip::Graph* graph hip::GraphNode* const* pDependencies, size_t numDependencies, const hipKernelNodeParams* pNodeParams, const ihipExtKernelEvents* pNodeEvents = nullptr, - bool capture = true, int coopKernel = 0) { - if (pGraphNode == nullptr || graph == nullptr || - (numDependencies > 0 && pDependencies == nullptr) || pNodeParams == nullptr || - pNodeParams->func == nullptr) { - return hipErrorInvalidValue; - } - + bool capture = true, int coopKernel = 0, int devId = 0) { if (!hip::Graph::isGraphValid(graph)) { return hipErrorInvalidValue; } @@ -115,6 +109,9 @@ hipError_t ihipGraphAddKernelNode(hip::GraphNode** pGraphNode, hip::Graph* graph } *pGraphNode = new hip::GraphKernelNode(pNodeParams, pNodeEvents, coopKernel); + if (devId != 0) { + (*pGraphNode)->SetDeviceId(devId); + } status = ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies, capture); return status; } @@ -231,9 +228,9 @@ hipError_t capturehipLaunchKernel(hipStream_t& stream, const void*& hostFunction nodeParams.sharedMemBytes = sharedMemBytes; hip::GraphNode* pGraphNode; - hipError_t status = - ihipGraphAddKernelNode(&pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), - s->GetLastCapturedNodes().size(), &nodeParams); + hipError_t status = ihipGraphAddKernelNode( + &pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), + s->GetLastCapturedNodes().size(), &nodeParams, nullptr, true, 0, s->DeviceId()); if (status != hipSuccess) { return status; } @@ -338,9 +335,9 @@ hipError_t capturehipModuleLaunchKernel(hipStream_t& stream, hipFunction_t& f, u nodeParams.sharedMemBytes = sharedMemBytes; hip::GraphNode* pGraphNode; - hipError_t status = - ihipGraphAddKernelNode(&pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), - s->GetLastCapturedNodes().size(), &nodeParams); + hipError_t status = ihipGraphAddKernelNode( + &pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), + s->GetLastCapturedNodes().size(), &nodeParams, nullptr, true, 0, s->DeviceId()); if (status != hipSuccess) { return status; } @@ -373,7 +370,7 @@ hipError_t capturehipModuleLaunchCooperativeKernel(hipStream_t& stream, hipFunct hipError_t status = ihipGraphAddKernelNode(&pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), s->GetLastCapturedNodes().size(), &nodeParams, nullptr, true, - amd::NDRangeKernelCommand::CooperativeGroups); + amd::NDRangeKernelCommand::CooperativeGroups, s->DeviceId()); if (status != hipSuccess) { return status; } @@ -400,9 +397,9 @@ hipError_t capturehipLaunchByPtr(hipStream_t& stream, hipFunction_t func, dim3 b hip::GraphNode* pGraphNode; hip::Stream* s = reinterpret_cast(stream); - hipError_t status = - ihipGraphAddKernelNode(&pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), - s->GetLastCapturedNodes().size(), &nodeParams); + hipError_t status = ihipGraphAddKernelNode( + &pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), + s->GetLastCapturedNodes().size(), &nodeParams, nullptr, true, 0, s->DeviceId()); if (status != hipSuccess) { return status; } @@ -434,7 +431,7 @@ hipError_t capturehipLaunchCooperativeKernel(hipStream_t& stream, const void*& f hipError_t status = ihipGraphAddKernelNode(&pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), s->GetLastCapturedNodes().size(), &nodeParams, nullptr, true, - amd::NDRangeKernelCommand::CooperativeGroups); + amd::NDRangeKernelCommand::CooperativeGroups, s->DeviceId()); if (status != hipSuccess) { return status; } @@ -1293,14 +1290,14 @@ hipError_t hipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, HIP_INIT_API(hipGraphAddKernelNode, pGraphNode, graph, pDependencies, numDependencies, pNodeParams); if (pGraphNode == nullptr || graph == nullptr || pNodeParams == nullptr || - (numDependencies > 0 && pDependencies == nullptr)) { + (numDependencies > 0 && pDependencies == nullptr) || pNodeParams->func == nullptr) { HIP_RETURN(hipErrorInvalidValue); } hip::GraphNode* node; - hipError_t status = ihipGraphAddKernelNode( - &node, reinterpret_cast(graph), - reinterpret_cast(pDependencies), numDependencies, pNodeParams, - nullptr, false); + hipError_t status = + ihipGraphAddKernelNode(&node, reinterpret_cast(graph), + reinterpret_cast(pDependencies), + numDependencies, pNodeParams, nullptr, false); *pGraphNode = reinterpret_cast(node); HIP_RETURN(status); } diff --git a/projects/clr/hipamd/src/hip_graph_internal.cpp b/projects/clr/hipamd/src/hip_graph_internal.cpp index b30a77ecb7..4839b18af3 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.cpp +++ b/projects/clr/hipamd/src/hip_graph_internal.cpp @@ -359,9 +359,6 @@ hipError_t GraphExec::Init() { } if (DEBUG_CLR_GRAPH_PACKET_CAPTURE) { if (max_streams_ == 1) { - // Don't wait for other streams to finish. - // Capture stream is to capture AQL packet. - capture_stream_ = hip::getNullStream(false); // For graph nodes capture AQL packets to dispatch them directly during graph launch. status = CaptureAQLPackets(); } @@ -386,8 +383,6 @@ void GraphExec::GetKernelArgSizeForGraph(size_t& kernArgSizeForGraph) { GraphKernelArgManager* KernelArgManager = GetKernelArgManager(); KernelArgManager->retain(); childNode->SetKernelArgManager(KernelArgManager); - // Set capture stream for child graph - childNode->capture_stream_ = capture_stream_; if (childNode->GetChildGraph()->max_streams_ == 1) { childNode->GetKernelArgSizeForGraph(kernArgSizeForGraph); } @@ -408,7 +403,7 @@ hipError_t GraphExec::AllocKernelArgForGraphNode() { } } if (node->GraphCaptureEnabled()) { - status = node->CaptureAndFormPacket(capture_stream_, GetKernelArgManager()); + status = node->CaptureAndFormPacket(GetKernelArgManager()); } else if (node->GetType() == hipGraphNodeTypeGraph) { auto childNode = reinterpret_cast(node); if (childNode->GetChildGraph()->max_streams_ == 1) { @@ -428,9 +423,13 @@ hipError_t GraphExec::CaptureAQLPackets() { hipError_t status = hipSuccess; size_t kernArgSizeForGraph = 0; GetKernelArgSizeForGraph(kernArgSizeForGraph); - auto device = g_devices[ihipGetDevice()]->devices()[0]; + // When we support multi device graph lauch we need to allocate the kenel args on respective + // device for each kernel Assume graph has nodes of same device allocate kernel args on the device + // from the first node + auto device = g_devices[topoOrder_[0]->GetDeviceId()]->devices()[0]; // Add a larger initial pool to accomodate for any updates to kernel args - bool bStatus = kernArgManager_->AllocGraphKernargPool(kernArgSizeForGraph + kKernArgChunkSize); + bool bStatus = + kernArgManager_->AllocGraphKernargPool(kernArgSizeForGraph + kKernArgChunkSize, device); if (bStatus != true) { return hipErrorMemoryAllocation; } @@ -447,7 +446,7 @@ hipError_t GraphExec::CaptureAQLPackets() { hipError_t GraphExec::UpdateAQLPacket(hip::GraphNode* node) { hipError_t status = hipSuccess; if (max_streams_ == 1) { - status = node->CaptureAndFormPacket(capture_stream_, kernArgManager_); + status = node->CaptureAndFormPacket(kernArgManager_); } return status; } @@ -749,11 +748,10 @@ hipError_t GraphExec::Run(hip::Stream* launch_stream) { } // ================================================================================================ -bool GraphKernelArgManager::AllocGraphKernargPool(size_t pool_size) { +bool GraphKernelArgManager::AllocGraphKernargPool(size_t pool_size, amd::Device* device) { bool bStatus = true; assert(pool_size > 0); address graph_kernarg_base; - auto device = g_devices[ihipGetDevice()]->devices()[0]; // Current device is stored as part of tls. Save current device to destroy kernelArgs from the // callback thread. device_ = device; @@ -783,7 +781,7 @@ address GraphKernelArgManager::AllocKernArg(size_t size, size_t alignment) { kernarg_graph_.back().kernarg_pool_offset_ = pool_new_usage; } else { // If current chunck is full allocate new chunck with same size as current - bool bStatus = AllocGraphKernargPool(kernarg_graph_.back().kernarg_pool_size_); + bool bStatus = AllocGraphKernargPool(kernarg_graph_.back().kernarg_pool_size_, device_); if (bStatus == false) { return nullptr; } else { diff --git a/projects/clr/hipamd/src/hip_graph_internal.hpp b/projects/clr/hipamd/src/hip_graph_internal.hpp index 79f3522872..59e2e0c8a0 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.hpp +++ b/projects/clr/hipamd/src/hip_graph_internal.hpp @@ -159,8 +159,8 @@ class GraphKernelArgManager : public amd::ReferenceCountedObject, } } - // Allocate kernel arg pool for the given size. - bool AllocGraphKernargPool(size_t pool_size); + // Allocate kernel arg pool on device for the given size. + bool AllocGraphKernargPool(size_t pool_size, amd::Device* device); // Allocate kernel args from current chunck for given size and alignment. // If kernel arg pool is full allocate new chunck and alloc kern args from new pool. @@ -194,6 +194,7 @@ class GraphNode : public hipGraphNodeDOTAttribute { id_(nextID++), parentGraph_(nullptr), isEnabled_(1), + dev_id_(ihipGetDevice()), hipGraphNodeDOTAttribute(style, shape, label) { amd::ScopedLock lock(nodeSetLock_); nodeSet_.insert(this); @@ -209,6 +210,7 @@ class GraphNode : public hipGraphNodeDOTAttribute { amd::ScopedLock lock(nodeSetLock_); nodeSet_.insert(this); isEnabled_ = node.isEnabled_; + dev_id_ = node.dev_id_; } virtual ~GraphNode() { @@ -240,7 +242,8 @@ class GraphNode : public hipGraphNodeDOTAttribute { size_t GetKerArgSize() const { return alignedKernArgSize_; } size_t GetKernargSegmentByteSize() const { return kernargSegmentByteSize_; } size_t GetKernargSegmentAlignment() const { return kernargSegmentAlignment_; } - hipError_t CaptureAndFormPacket(hip::Stream* capture_stream, GraphKernelArgManager* kernArgMgr) { + hipError_t CaptureAndFormPacket(GraphKernelArgManager* kernArgMgr) { + auto capture_stream = hip::getNullStream(g_devices[dev_id_]->devices()[0]->context(), false); hipError_t status = CreateCommand(capture_stream); if (status != hipSuccess) { return status; @@ -433,10 +436,13 @@ class GraphNode : public hipGraphNodeDOTAttribute { if (DEBUG_HIP_GRAPH_DOT_PRINT) { out << "\nStreamId:" << stream_id_; out << "\nSignalIsRequired: " << ((signal_is_required_) ? "true" : "false"); + out << "\nDeviceId:" << dev_id_; } out << "\""; out << "];"; } + void SetDeviceId(int id) { dev_id_ = id; } + int GetDeviceId() const { return dev_id_; } protected: // Declare Graph and GraphExec as friends of node for simpler access to GraphNode fields @@ -465,6 +471,8 @@ class GraphNode : public hipGraphNodeDOTAttribute { size_t alignedKernArgSize_ = 256; //!< Aligned size required for kernel args size_t kernargSegmentByteSize_ = 512; //!< Kernel arg segment byte size size_t kernargSegmentAlignment_ = 256; //!< Kernel arg segment alignment + int dev_id_; //!< Device Id when node is created(dev id from capture stream/current device + //!< when explicitly added) }; class Graph { @@ -799,7 +807,6 @@ class GraphExec : public amd::ReferenceCountedObject, public Graph { //! Topological order of the graph doesn't include nodes embedded as part of the child graph std::vector topoOrder_; std::vector parallel_streams_; - hip::Stream* capture_stream_; uint64_t flags_ = 0; GraphKernelArgManager* kernArgManager_ = nullptr; //!< Kernel Arg manager for graph. int instantiateDeviceId_ = -1; @@ -906,7 +913,7 @@ class GraphKernelNode : public GraphNode { return; } for (auto& command : commands_) { - hipFunction_t func = getFunc(kernelParams_, ihipGetDevice()); + hipFunction_t func = getFunc(kernelParams_, dev_id_); hip::DeviceFunc* function = hip::DeviceFunc::asFunction(func); amd::Kernel* kernel = function->kernel(); amd::ScopedLock lock(function->dflock_); @@ -934,13 +941,14 @@ class GraphKernelNode : public GraphNode { if (DEBUG_HIP_GRAPH_DOT_PRINT) { out << "StreamId:" << stream_id_; out << "\nSignalIsRequired: " << ((signal_is_required_) ? "true" : "false"); + out << "\nDeviceId:" << dev_id_; } out << "\""; out << "];"; } virtual std::string GetLabel(hipGraphDebugDotFlags flag) override { - hipFunction_t func = getFunc(kernelParams_, ihipGetDevice()); + hipFunction_t func = getFunc(kernelParams_, dev_id_); hip::DeviceFunc* function = hip::DeviceFunc::asFunction(func); std::string label; char buffer[4096]; @@ -1010,14 +1018,14 @@ class GraphKernelNode : public GraphNode { hipError_t copyParams(const hipKernelNodeParams* pNodeParams) { hasHiddenHeap_ = false; - hipFunction_t func = getFunc(*pNodeParams, ihipGetDevice()); + hipFunction_t func = getFunc(*pNodeParams, dev_id_); if (!func) { return hipErrorInvalidDeviceFunction; } 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]; + auto device = g_devices[dev_id_]->devices()[0]; device::Kernel* devKernel = const_cast(kernel->getDeviceKernel(*device)); kernargSegmentByteSize_ = devKernel->KernargSegmentByteSize(); kernargSegmentAlignment_ = devKernel->KernargSegmentAlignment(); @@ -1140,15 +1148,14 @@ class GraphKernelNode : public GraphNode { GraphNode* clone() const override { return new GraphKernelNode(*this); } hipError_t CreateCommand(hip::Stream* stream) override { - int devID = hip::getDeviceID(stream->context()); - hipFunction_t func = getFunc(kernelParams_, devID); + hipFunction_t func = getFunc(kernelParams_, dev_id_); if (!func) { return hipErrorInvalidDeviceFunction; } hip::DeviceFunc* function = hip::DeviceFunc::asFunction(func); amd::Kernel* kernel = function->kernel(); amd::ScopedLock lock(function->dflock_); - hipError_t status = validateKernelParams(&kernelParams_, func, devID); + hipError_t status = validateKernelParams(&kernelParams_, func, dev_id_); if (hipSuccess != status) { return status; } @@ -1187,12 +1194,12 @@ class GraphKernelNode : public GraphNode { void GetParams(hipKernelNodeParams* params) { *params = kernelParams_; } hipError_t SetParams(const hipKernelNodeParams* params) { - hipFunction_t func = getFunc(kernelParams_, ihipGetDevice()); + hipFunction_t func = getFunc(kernelParams_, dev_id_); if (!func) { return hipErrorInvalidDeviceFunction; } // updates kernel params - hipError_t status = validateKernelParams(params, func, ihipGetDevice()); + hipError_t status = validateKernelParams(params, func, dev_id_); if (hipSuccess != status) { ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "[hipGraph] Failed to validateKernelParams"); return status; @@ -1213,7 +1220,7 @@ class GraphKernelNode : public GraphNode { hipError_t SetAttrParams(hipKernelNodeAttrID attr, const hipKernelNodeAttrValue* params) { hipDeviceProp_t prop = {0}; - hipError_t status = ihipGetDeviceProperties(&prop, ihipGetDevice()); + hipError_t status = ihipGetDeviceProperties(&prop, dev_id_); if (hipSuccess != status){ return status; } diff --git a/projects/clr/hipamd/src/hip_internal.hpp b/projects/clr/hipamd/src/hip_internal.hpp index e9ff22f984..45d7dbe225 100644 --- a/projects/clr/hipamd/src/hip_internal.hpp +++ b/projects/clr/hipamd/src/hip_internal.hpp @@ -654,7 +654,7 @@ public: /// and Blocking streams extern hip::Stream* getStream(hipStream_t stream, bool wait = true); /// Get default stream associated with the ROCclr context - extern hip::Stream* getNullStream(amd::Context&); + extern hip::Stream* getNullStream(amd::Context&, bool wait = true); /// Get default stream of the thread extern hip::Stream* getNullStream(bool wait = true); /// Get device ID associated with the ROCclr context