SWDEV-373173 - Add kernel metadata for WGP mode

Change-Id: Ic270e90c46938b975513ff2bc19b7bb6b9485f9d


[ROCm/clr commit: 0bd0c29b7d]

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;
   }

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; }