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

SWDEV-421025 - Graph Kernel Node priority Attribute Set/Get

Browse Source 
Change-Id: I5c422728aa694c8dabb5cf9bade441101512a249
This commit is contained in:
Rahul Manocha
2024-01-11 22:21:44 +00:00
parent ea25e65834
commit f40c380cdb
2 changed files with 20 additions and 12 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/src/hip_graph.cpp
+6 -5
View File
@@ -1255,7 +1255,8 @@ hipError_t hipGraphInstantiateWithFlags(hipGraphExec_t* pGraphExec, hipGraph_t g
}
// invalid flag check
if (flags != 0 && flags != hipGraphInstantiateFlagAutoFreeOnLaunch) {
if (flags != 0 && flags != hipGraphInstantiateFlagAutoFreeOnLaunch &&
flags != hipGraphInstantiateFlagUseNodePriority) {
HIP_RETURN(hipErrorInvalidValue);
}
@@ -1395,7 +1396,8 @@ hipError_t hipGraphKernelNodeSetAttribute(hipGraphNode_t hNode, hipKernelNodeAtt
HIP_RETURN(hipErrorInvalidValue);
}
if (attr != hipKernelNodeAttributeAccessPolicyWindow &&
attr != hipKernelNodeAttributeCooperative) {
attr != hipKernelNodeAttributeCooperative &&
attr != hipLaunchAttributePriority) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(reinterpret_cast<hip::GraphKernelNode*>(hNode)->SetAttrParams(attr, value));
@@ -1408,7 +1410,8 @@ hipError_t hipGraphKernelNodeGetAttribute(hipGraphNode_t hNode, hipKernelNodeAtt
HIP_RETURN(hipErrorInvalidValue);
}
if (attr != hipKernelNodeAttributeAccessPolicyWindow &&
attr != hipKernelNodeAttributeCooperative) {
attr != hipKernelNodeAttributeCooperative &&
attr != hipLaunchAttributePriority) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(reinterpret_cast<hip::GraphKernelNode*>(hNode)->GetAttrParams(attr, value));
@@ -2870,5 +2873,3 @@ hipError_t hipGraphExecExternalSemaphoresWaitNodeSetParams(hipGraphExec_t hGraph
}
} // namespace hip
hipamd/src/hip_graph_internal.hpp
+14 -7
View File
@@ -817,14 +817,15 @@ class GraphKernelNode : public GraphNode {
"{\n%s\n| {ID | %d | %s\\<\\<\\<(%u,%u,%u),(%u,%u,%u),%u\\>\\>\\>}\n| {{node "
"handle | func handle} | {%p | %p}}\n| {accessPolicyWindow | {base_ptr | num_bytes | "
"hitRatio | hitProp | missProp} | {%p | %zu | %f | %d | %d}}\n| {cooperative | "
"%u}\n| {priority | 0}\n}",
"%u}\n| {priority | %d}\n}",
label_.c_str(), GetID(), function->name().c_str(), kernelParams_.gridDim.x,
kernelParams_.gridDim.y, kernelParams_.gridDim.z, kernelParams_.blockDim.x,
kernelParams_.blockDim.y, kernelParams_.blockDim.z,
kernelParams_.sharedMemBytes, this, kernelParams_.func,
kernelAttr_.accessPolicyWindow.base_ptr, kernelAttr_.accessPolicyWindow.num_bytes,
kernelAttr_.accessPolicyWindow.hitRatio, kernelAttr_.accessPolicyWindow.hitProp,
kernelAttr_.accessPolicyWindow.missProp, kernelAttr_.cooperative);
kernelAttr_.accessPolicyWindow.missProp, kernelAttr_.cooperative,
kernelAttr_.priority);
label = buffer;
}
else if (flag == hipGraphDebugDotFlagsKernelNodeAttributes) {
@@ -832,11 +833,12 @@ class GraphKernelNode : public GraphNode {
"{\n%s\n| {ID | %d | %s}\n"
"| {accessPolicyWindow | {base_ptr | num_bytes | "
"hitRatio | hitProp | missProp} |\n| {%p | %zu | %f | %d | %d}}\n| {cooperative | "
"%u}\n| {priority | 0}\n}",
"%u}\n| {priority | %d}\n}",
label_.c_str(), GetID(), function->name().c_str(),
kernelAttr_.accessPolicyWindow.base_ptr, kernelAttr_.accessPolicyWindow.num_bytes,
kernelAttr_.accessPolicyWindow.hitRatio, kernelAttr_.accessPolicyWindow.hitProp,
kernelAttr_.accessPolicyWindow.missProp, kernelAttr_.cooperative);
kernelAttr_.accessPolicyWindow.missProp, kernelAttr_.cooperative,
kernelAttr_.priority);
label = buffer;
}
else if (flag == hipGraphDebugDotFlagsKernelNodeParams) {
@@ -1072,7 +1074,10 @@ class GraphKernelNode : public GraphNode {
kernelAttr_.accessPolicyWindow.num_bytes = params->accessPolicyWindow.num_bytes;
} else if (attr == hipKernelNodeAttributeCooperative) {
kernelAttr_.cooperative = params->cooperative;
} else if (attr == hipLaunchAttributePriority) {
kernelAttr_.priority = params->priority;
}
kernelAttrInUse_ = attr;
return hipSuccess;
}
@@ -1087,6 +1092,8 @@ class GraphKernelNode : public GraphNode {
params->accessPolicyWindow.num_bytes = kernelAttr_.accessPolicyWindow.num_bytes;
} else if (attr == hipKernelNodeAttributeCooperative) {
params->cooperative = kernelAttr_.cooperative;
} else if (attr == hipLaunchAttributePriority) {
params->priority = kernelAttr_.priority;
}
return hipSuccess;
}
@@ -1110,6 +1117,9 @@ class GraphKernelNode : public GraphNode {
case hipKernelNodeAttributeCooperative:
kernelAttr_.cooperative = srcNode->kernelAttr_.cooperative;
break;
case hipLaunchAttributePriority:
kernelAttr_.priority = srcNode->kernelAttr_.priority;
break;
default:
return hipErrorInvalidValue;
}
@@ -2447,6 +2457,3 @@ class hipGraphExternalSemWaitNode : public GraphNode {
};
} // namespace hip