Revert "[libomptarget] [amdgpu] Fix default setting of max flat workgroup size"

[dhruvachak]

Summary

This reverts commit 2240b41ee4f30fe938975677a0a5a2c5c26d271b.
A value of 0 for KernDescVal WG_Size implies it is unknown, so it should be
set to the default. The above change was made without this assumption.

[JonChesterfield]

git revert is also good. Thanks

openmp/libomptarget/plugins/amdgpu/src/rtl.cpp

 #include <assert.h>
 #include <cstdio>
 #include <cstdlib>
 #include <cstring>
 #include <elf.h>
 #include <fstream>
 #include <functional>
 #include <iostream>
 #include <libelf.h>
 #include <list>
 #include <memory>
 #include <mutex>
 #include <shared_mutex>
 #include <thread>
 #include <unordered_map>
 #include <vector>
 
if (err == HSA_STATUS_SUCCESS) {
       DP("KernDesc: Version: %d\n", KernDescVal.Version);
       DP("KernDesc: TSize: %d\n", KernDescVal.TSize);
       DP("KernDesc: WG_Size: %d\n", KernDescVal.WG_Size);
       DP("KernDesc: Mode: %d\n", KernDescVal.Mode);
 
       // Get ExecMode
       ExecModeVal = KernDescVal.Mode;
       DP("ExecModeVal %d\n", ExecModeVal);
-      // If KernDescVal.WG_Size is 0, it is equivalent to not
-      // specified. Hence, max_flat_workgroup_size is filtered out in
-      // getLaunchVals
+      if (KernDescVal.WG_Size == 0) {
+        KernDescVal.WG_Size = RTLDeviceInfoTy::Default_WG_Size;
+        DP("Setting KernDescVal.WG_Size to default %d\n", KernDescVal.WG_Size);
+      }
       WGSizeVal = KernDescVal.WG_Size;
       DP("WGSizeVal %d\n", WGSizeVal);
       check("Loading KernDesc computation property", err);
     } else {
       DP("Warning: Loading KernDesc '%s' - symbol not found, ", KernDescName);
 
       // Generic
       std::string ExecModeNameStr(e->name);
if (ExecutionMode == GENERIC) { // Add master warp for GENERIC
       DP("Adding master wavefront: +%d threads\n", RTLDeviceInfoTy::Warp_Size);
     }
     if (threadsPerGroup > RTLDeviceInfoTy::Max_WG_Size) { // limit to max
       threadsPerGroup = RTLDeviceInfoTy::Max_WG_Size;
       DP("Setting threads per block to maximum %d\n", threadsPerGroup);
     }
   }
   // check flat_max_work_group_size attr here
-  if (ConstWGSize > 0 && threadsPerGroup > ConstWGSize) {
+  if (threadsPerGroup > ConstWGSize) {
     threadsPerGroup = ConstWGSize;
     DP("Reduced threadsPerGroup to flat-attr-group-size limit %d\n",
        threadsPerGroup);
   }
   if (print_kernel_trace & STARTUP_DETAILS)
     fprintf(stderr, "threadsPerGroup: %d\n", threadsPerGroup);
   DP("Preparing %d threads\n", threadsPerGroup);