[Openmp-commits] [openmp] 98c36f0 - Revert "[libomptarget] [amdgpu] Fix default setting of max flat workgroup size"

[Dhruva Chakrabarti via Openmp-commits]

Author: Dhruva Chakrabarti
Date: 2021-06-30T17:15:00-07:00
New Revision: 98c36f0079d46c83bf03a7bce1b9d7e22abe0cba

URL: https://github.com/llvm/llvm-project/commit/98c36f0079d46c83bf03a7bce1b9d7e22abe0cba
DIFF: https://github.com/llvm/llvm-project/commit/98c36f0079d46c83bf03a7bce1b9d7e22abe0cba.diff

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

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.

Reviewed By: JonChesterfield

Differential Revision: https://reviews.llvm.org/D105250

Added: 
    

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

Removed: 
    


################################################################################
diff  --git a/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp b/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
index 03fcc470984f..b04745529652 100644
--- a/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
@@ -1711,9 +1711,10 @@ __tgt_target_table *__tgt_rtl_load_binary_locked(int32_t device_id,
       // 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 @@ launchVals getLaunchVals(int ConstWGSize, int ExecutionMode, int EnvTeamLimit,
     }
   }
   // 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);