From 5b79a926e3e64ed8f6da0de74a82a129428b60f9 Mon Sep 17 00:00:00 2001
From: Anusha GodavarthySurya <anusha.godavarthysurya@amd.com>
Date: Mon, 2 Jan 2023 12:28:22 +0000
Subject: [PATCH] SWDEV-373173 - Add kernel metadata for WGP mode

Change-Id: Ic270e90c46938b975513ff2bc19b7bb6b9485f9d


[ROCm/clr commit: 0bd0c29b7d3a1749ae1020a45d20b082796241e4]
---
 projects/clr/rocclr/device/devkernel.cpp |  3 +++
 projects/clr/rocclr/device/devkernel.hpp | 13 ++++++++++---
 2 files changed, 13 insertions(+), 3 deletions(-)

diff --git a/projects/clr/rocclr/device/devkernel.cpp b/projects/clr/rocclr/device/devkernel.cpp
index 2313084bef..213f39c2d0 100644
--- a/projects/clr/rocclr/device/devkernel.cpp
+++ b/projects/clr/rocclr/device/devkernel.cpp
@@ -566,6 +566,9 @@ static amd_comgr_status_t populateKernelMetaV3(const amd_comgr_metadata_node_t k
     case KernelField::Kind:
       kernel->SetKernelKind(buf);
       break;
+    case KernelField::WgpMode:
+      kernel->SetWGPMode(buf.compare("true") == 0);
+      break;
     default:
       return AMD_COMGR_STATUS_ERROR;
   }
diff --git a/projects/clr/rocclr/device/devkernel.hpp b/projects/clr/rocclr/device/devkernel.hpp
index f2ad0df404..2e4066c3c8 100644
--- a/projects/clr/rocclr/device/devkernel.hpp
+++ b/projects/clr/rocclr/device/devkernel.hpp
@@ -241,7 +241,8 @@ enum class KernelField : uint8_t {
   MaxFlatWorkGroupSize    = 12,
   NumSpilledSGPRs         = 13,
   NumSpilledVGPRs         = 14,
-  Kind                    = 15
+  Kind                    = 15,
+  WgpMode                = 16
 };
 
 static const std::map<std::string,ArgField> ArgFieldMapV3 =
@@ -327,7 +328,8 @@ static const std::map<std::string,KernelField> KernelFieldMapV3 =
   {".max_flat_workgroup_size",    KernelField::MaxFlatWorkGroupSize},
   {".sgpr_spill_count",           KernelField::NumSpilledSGPRs},
   {".vgpr_spill_count",           KernelField::NumSpilledVGPRs},
-  {".kind",                       KernelField::Kind}
+  {".kind",                       KernelField::Kind},
+  {".workgroup_processor_mode",   KernelField::WgpMode}
 };
 
 #endif  // defined(USE_COMGR_LIBRARY)
@@ -385,7 +387,8 @@ class Kernel : public amd::HeapObject {
     bool uniformWorkGroupSize_;       //!< uniform work group size option
     size_t wavesPerSimdHint_;         //!< waves per simd hit
     int maxOccupancyPerCu_;           //!< Max occupancy per compute unit in threads
-    size_t constMemSize_;           //!< size of user-allocated constant memory 
+    size_t constMemSize_;           //!< size of user-allocated constant memory
+    bool isWGPMode_;                  //!< kernel compiled in WGP/cumode
   };
 
   //! Default constructor
@@ -513,6 +516,10 @@ class Kernel : public amd::HeapObject {
     kind_ = (kind == "init") ? Init : ((kind == "fini") ? Fini : Normal);
   }
 
+  void SetWGPMode(bool wgpMode) {
+    workGroupInfo_.isWGPMode_ = wgpMode;
+  }
+
   bool isInitKernel() const { return kind_ == Init; }
 
   bool isFiniKernel() const { return kind_ == Fini; }