[OpenMP] [clang] Prioritize thread_limit clause over command-line option.

Change-Id: I15c68f3815148447e423a4a71c545ea36861f971

Author: dhruvachak

clang/lib/CodeGen/CodeGenModule.cpp

@@ -8295,17 +8295,14 @@ int CodeGenModule::getWorkGroupSizeSPMDHelper(const OMPExecutableDirective &D) {
     clang::Expr::EvalResult Result;
     if (ThreadLimitExpr->EvaluateAsInt(Result, getContext())) {
       int ThreadLimitEval = Result.Val.getInt().getExtValue();
-      if (ThreadLimitEval > 0 && ThreadLimitEval < ThreadLimit)
+      if (ThreadLimitEval > 0 && ThreadLimitEval <= ThreadLimit) {
         ThreadLimit = ThreadLimitEval;
+        // Prioritize value from clause over command-line option.
+        WorkGroupSz = ThreadLimit;
+      }
     }
   }
 
-  // If the command line work group size is less than any default or user
-  // specified thread limit then it is honored otherwise the thread limit
-  // determined above will be used.
-  if (WorkGroupSz > ThreadLimit)
-    WorkGroupSz = ThreadLimit;
-
   // Set the actual number of threads if the user requests a value different
   // then the default. If the value is greater than the currently computed
   // thread limit then cap the number of threads to the thread limit.

openmp/libomptarget/test/offloading/thread_limit.c

@@ -10,20 +10,78 @@
 // UNSUPPORTED: x86_64-pc-linux-gnu
 // UNSUPPORTED: x86_64-pc-linux-gnu-LTO
 
+// clang-format on
 int main() {
   int n = 1 << 20;
   int th = 12;
   int te = n / th;
+
 // DEFAULT: 12 (MaxFlatWorkGroupSize: 
 #pragma omp target
 #pragma omp teams loop num_teams(te), thread_limit(th)
   for (int i = 0; i < n; i++) {
   }
 
-// DEFAULT: 13 (MaxFlatWorkGroupSize: 
-  #pragma omp target
-  #pragma omp teams distribute parallel for simd num_teams(te), thread_limit(th+1) simdlen(64)
+// DEFAULT: 13 (MaxFlatWorkGroupSize:
+#pragma omp target
+#pragma omp teams distribute parallel for simd num_teams(te),                  \
+    thread_limit(th + 1) simdlen(64)
   for(int i = 0; i < n; i++) {
   }
+
+// DEFAULT: 128 (MaxFlatWorkGroupSize:
+#pragma omp target teams distribute parallel for thread_limit(128)
+  for (int i = 0; i < n; i++) {
+  }
+
+// DEFAULT: 512 (MaxFlatWorkGroupSize:
+#pragma omp target teams distribute parallel for thread_limit(512)
+  for (int i = 0; i < n; i++) {
+  }
+
+// DEFAULT: 1024 (MaxFlatWorkGroupSize:
+#pragma omp target teams distribute parallel for thread_limit(1024)
+  for (int i = 0; i < n; i++) {
+  }
+
+// DEFAULT: 128 (MaxFlatWorkGroupSize:
+#pragma omp target teams distribute parallel for num_threads(128)
+  for (int i = 0; i < n; i++) {
+  }
+
+// DEFAULT: 512 (MaxFlatWorkGroupSize:
+#pragma omp target teams distribute parallel for num_threads(512)
+  for (int i = 0; i < n; i++) {
+  }
+
+// DEFAULT: 1024 (MaxFlatWorkGroupSize:
+#pragma omp target teams distribute parallel for num_threads(1024)
+  for (int i = 0; i < n; i++) {
+  }
+
+// DEFAULT: 64 (MaxFlatWorkGroupSize:
+#pragma omp target teams distribute parallel for thread_limit(64)              \
+    num_threads(128)
+  for (int i = 0; i < n; i++) {
+  }
+
+// DEFAULT: 64 (MaxFlatWorkGroupSize:
+#pragma omp target teams distribute parallel for thread_limit(128)             \
+    num_threads(64)
+  for (int i = 0; i < n; i++) {
+  }
+
+// DEFAULT: 512 (MaxFlatWorkGroupSize:
+#pragma omp target teams distribute parallel for thread_limit(512)             \
+    num_threads(1024)
+  for (int i = 0; i < n; i++) {
+  }
+
+// DEFAULT: 512 (MaxFlatWorkGroupSize:
+#pragma omp target teams distribute parallel for thread_limit(1024)            \
+    num_threads(512)
+  for (int i = 0; i < n; i++) {
+  }
+
   return 0;
 }