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

SWDEV-522841 - Graph nodes must be created/launched on device where they are captured/created (#108)

Browse Source 
[ROCm/clr commit: 2538d7f02b]
This commit is contained in:
Godavarthy Surya, Anusha
2025-04-29 22:20:39 +05:30
committed by GitHub
parent 6ab34e0924
commit e4a499f22e
5 changed files with 55 additions and 53 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/src/hip_context.cpp
+3 -3
View File
@@ -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;
}
projects/clr/hipamd/src/hip_graph.cpp
+20 -23
View File
@@ -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<hip::Stream*>(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<hip::Graph*>(graph),
reinterpret_cast<hip::GraphNode* const*>(pDependencies), numDependencies, pNodeParams,
nullptr, false);
hipError_t status =
ihipGraphAddKernelNode(&node, reinterpret_cast<hip::Graph*>(graph),
reinterpret_cast<hip::GraphNode* const*>(pDependencies),
numDependencies, pNodeParams, nullptr, false);
*pGraphNode = reinterpret_cast<hipGraphNode*>(node);
HIP_RETURN(status);
}
projects/clr/hipamd/src/hip_graph_internal.cpp
+10 -12
View File
@@ -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<hip::ChildGraphNode*>(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 {
projects/clr/hipamd/src/hip_graph_internal.hpp
+21 -14
View File
@@ -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<Node> topoOrder_;
std::vector<hip::Stream*> 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<device::Kernel*>(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;
}
projects/clr/hipamd/src/hip_internal.hpp
+1 -1
View File
@@ -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