From f40c380cdb58005e57931257b9c01e714fca6686 Mon Sep 17 00:00:00 2001
From: Rahul Manocha <rahul.manocha@amd.com>
Date: Thu, 11 Jan 2024 22:21:44 +0000
Subject: [PATCH] SWDEV-421025 - Graph Kernel Node priority Attribute Set/Get

Change-Id: I5c422728aa694c8dabb5cf9bade441101512a249
---
 hipamd/src/hip_graph.cpp          | 11 ++++++-----
 hipamd/src/hip_graph_internal.hpp | 21 ++++++++++++++-------
 2 files changed, 20 insertions(+), 12 deletions(-)

diff --git a/hipamd/src/hip_graph.cpp b/hipamd/src/hip_graph.cpp
index bbb572a030..7cfb1babde 100644
--- a/hipamd/src/hip_graph.cpp
+++ b/hipamd/src/hip_graph.cpp
@@ -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
-
-
diff --git a/hipamd/src/hip_graph_internal.hpp b/hipamd/src/hip_graph_internal.hpp
index 616a51f9c2..7f6543cb91 100644
--- a/hipamd/src/hip_graph_internal.hpp
+++ b/hipamd/src/hip_graph_internal.hpp
@@ -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
-
-
-