From 339bb6a7dea035add36847da4aa89937450b1e75 Mon Sep 17 00:00:00 2001
From: Tao Sang <tao.sang@amd.com>
Date: Mon, 16 Jan 2023 18:36:44 -0500
Subject: [PATCH] SWDEV-377423 - Fix kernel attribute missing

Fix missing issue of kernel attributes including vec_type_hint,
work_group_size_hint and reqd_work_group_size.
Make WorkGroupInfo's meta attributes initialized before other parameters
are initialized.
This way workGroupInfo_'s compileSizeHint_, compileSize_ and
compileVecTypeHint_ will be valid when they are used to create kernel
signature in Kernel::createSignature().
Fix a typo of ".workgorup_size_hint".

Change-Id: I4a1ede2210a25596ad7a935cd4debb896e0147f8


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

diff --git a/projects/clr/rocclr/device/devkernel.cpp b/projects/clr/rocclr/device/devkernel.cpp
index 213f39c2d0..e598dca74f 100644
--- a/projects/clr/rocclr/device/devkernel.cpp
+++ b/projects/clr/rocclr/device/devkernel.cpp
@@ -1095,8 +1095,6 @@ bool Kernel::GetAttrCodePropMetadata() {
     return false;
   }
 
-  InitParameters(kernelMetaNode);
-
   // Set the workgroup information for the kernel
   workGroupInfo_.availableLDSSize_ = device().info().localMemSizePerCU_;
   workGroupInfo_.availableSGPRs_ = 104;
@@ -1144,12 +1142,13 @@ bool Kernel::GetAttrCodePropMetadata() {
                                                 static_cast<void*>(this));
   }
 
-
   if (status != AMD_COMGR_STATUS_SUCCESS) {
     LogError("Comgr Api failed with Status: \n");
     return false;
   }
 
+  InitParameters(kernelMetaNode);
+
   return true;
 }
 
diff --git a/projects/clr/rocclr/device/devkernel.hpp b/projects/clr/rocclr/device/devkernel.hpp
index 2e4066c3c8..3c244ff25b 100644
--- a/projects/clr/rocclr/device/devkernel.hpp
+++ b/projects/clr/rocclr/device/devkernel.hpp
@@ -315,7 +315,7 @@ static const std::map<std::string,KernelField> KernelFieldMapV3 =
 {
   {".symbol",                     KernelField::SymbolName},
   {".reqd_workgroup_size",        KernelField::ReqdWorkGroupSize},
-  {".workgorup_size_hint",        KernelField::WorkGroupSizeHint},
+  {".workgroup_size_hint",        KernelField::WorkGroupSizeHint},
   {".vec_type_hint",              KernelField::VecTypeHint},
   {".device_enqueue_symbol",      KernelField::DeviceEnqueueSymbol},
   {".kernarg_segment_size",       KernelField::KernargSegmentSize},