[offload][OpenMP] Re-enable strict modifier for num_threads

clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

@@ -922,13 +922,6 @@ void CGOpenMPRuntimeGPU::emitNumThreadsClause(
     OpenMPNumThreadsClauseModifier Modifier, OpenMPSeverityClauseKind Severity,
     SourceLocation SeverityLoc, const Expr *Message,
     SourceLocation MessageLoc) {
-  if (Modifier == OMPC_NUMTHREADS_strict) {
-    CGM.getDiags().Report(Loc,
-                          diag::warn_omp_gpu_unsupported_modifier_for_clause)
-        << "strict" << getOpenMPClauseName(OMPC_num_threads);
-    return;
-  }
-
   // Nothing to do.
 }
 
@@ -1236,9 +1229,9 @@ void CGOpenMPRuntimeGPU::emitParallelCall(
   if (!CGF.HaveInsertPoint())
     return;
 
-  auto &&ParallelGen = [this, Loc, OutlinedFn, CapturedVars, IfCond,
-                        NumThreads](CodeGenFunction &CGF,
-                                    PrePostActionTy &Action) {
+  auto &&ParallelGen = [this, Loc, OutlinedFn, CapturedVars, IfCond, NumThreads,
+                        NumThreadsModifier](CodeGenFunction &CGF,
+                                            PrePostActionTy &Action) {
     CGBuilderTy &Bld = CGF.Builder;
     llvm::Value *NumThreadsVal = NumThreads;
     llvm::Function *WFn = WrapperFunctionsMap[OutlinedFn];
@@ -1289,21 +1282,20 @@ void CGOpenMPRuntimeGPU::emitParallelCall(
       NumThreadsVal = Bld.CreateZExtOrTrunc(NumThreadsVal, CGF.Int32Ty);
 
     assert(IfCondVal && "Expected a value");
+    RuntimeFunction FnID = OMPRTL___kmpc_parallel_51;
     llvm::Value *RTLoc = emitUpdateLocation(CGF, Loc);
-    llvm::Value *Args[] = {
-        RTLoc,
-        getThreadID(CGF, Loc),
-        IfCondVal,
-        NumThreadsVal,
-        llvm::ConstantInt::get(CGF.Int32Ty, -1),
-        FnPtr,
-        ID,
-        Bld.CreateBitOrPointerCast(CapturedVarsAddrs.emitRawPointer(CGF),
-                                   CGF.VoidPtrPtrTy),
-        llvm::ConstantInt::get(CGM.SizeTy, CapturedVars.size())};
-    CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
-                            CGM.getModule(), OMPRTL___kmpc_parallel_51),
-                        Args);
+    llvm::SmallVector<llvm::Value *, 10> Args(
+        {RTLoc, getThreadID(CGF, Loc), IfCondVal, NumThreadsVal,
+         llvm::ConstantInt::get(CGF.Int32Ty, -1), FnPtr, ID,
+         Bld.CreateBitOrPointerCast(CapturedVarsAddrs.emitRawPointer(CGF),
+                                    CGF.VoidPtrPtrTy),
+         llvm::ConstantInt::get(CGM.SizeTy, CapturedVars.size())});
+    if (NumThreadsModifier == OMPC_NUMTHREADS_strict) {
+      FnID = OMPRTL___kmpc_parallel_60;
+      Args.append({llvm::ConstantInt::get(CGM.Int32Ty, true)});
+    }
+    CGF.EmitRuntimeCall(
+        OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(), FnID), Args);
   };
 
   RegionCodeGenTy RCG(ParallelGen);

clang/lib/CodeGen/CGOpenMPRuntimeGPU.h

@@ -245,7 +245,7 @@ class CGOpenMPRuntimeGPU : public CGOpenMPRuntime {
   /// \param NumThreads The value corresponding to the num_threads clause, if
   /// any, or nullptr.
   /// \param NumThreadsModifier The modifier of the num_threads clause, if
-  /// any, ignored otherwise. Currently unused on the device.
+  /// any, ignored otherwise.
   /// \param Severity The severity corresponding to the num_threads clause, if
   /// any, ignored otherwise. Currently unused on the device.
   /// \param Message The message string corresponding to the num_threads clause,

clang/test/AST/ByteCode/openmp.cpp

@@ -17,13 +17,12 @@ extern int omp_get_thread_num(void);
 
 int test2() {
   int x = 0;
-  int result[N] = {0};
+  int device_result[N] = {0};
 
-  #pragma omp parallel loop num_threads(strict: N) severity(warning) message("msg")
+  #pragma omp target parallel loop num_threads(strict: N)
   for (int i = 0; i < N; i++) {
     x = omp_get_thread_num();
-    result[i] = i + x;
+    device_result[i] = i + x;
   }
 }
 
-