[Openmp-commits] [PATCH] D105250: Revert "[libomptarget] [amdgpu] Fix default setting of max flat workgroup size"

[Dhruva Chakrabarti via Phabricator via Openmp-commits]

This revision was landed with ongoing or failed builds.
This revision was automatically updated to reflect the committed changes.
Closed by commit rG98c36f0079d4: Revert "[libomptarget] [amdgpu] Fix default setting of max flat workgroup size" (authored by dhruvachak).

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D105250/new/

https://reviews.llvm.org/D105250

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


Index: openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
===================================================================
--- openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
+++ openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
@@ -1711,9 +1711,10 @@
       // 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);
@@ -1930,7 +1931,7 @@
     }
   }
   // 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);


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D105250.355732.patch
Type: text/x-patch
Size: 1157 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20210701/5a0bbc43/attachment.bin>