[OpenMP][clang] Set num_threads 'strict' to unsupported on GPUs (llvm#160659)

Setting the prescriptiveness of the num_threads clause to 'strict' and
having a corresponding check (with message and severity clauses) does
not align well with how OpenMP should be handled for GPUs.

The num_threads expression may be an arbitrary integer expression which
is evaluated on the target, in correspondance to the OpenMP spec. This
prevents the check from being done before launching the kernel,
especially considering that the num_threads clause is associated with
the parallel directive and that there may be multiple parallel
directives with different num_threads clauses in a single target region.
Acting on the result of the 'strict' check on the GPU would require
doing I/O on the GPU, which can introduce performance regressions.
Delaying any actions resulting from the 'strict' check and doing them on
the host after executing the target region involves additional data
copies and is not really semantically correct.

For now, the 'strict' modifier for the num_threads clause and its
associated message and severity clause are set to be unsupported on
GPUs. Targets other than GPUs still support the aforementioned features
in the context of an OpenMP target region.

Author: ro-i

clang/include/clang/Basic/DiagnosticCommonKinds.td

@@ -433,6 +433,12 @@ def err_omp_more_one_clause : Error<
   "directive '#pragma omp %0' cannot contain more than one '%1' clause%select{| with '%3' name modifier| with 'source' dependence}2">;
 def err_omp_required_clause : Error<
   "directive '#pragma omp %0' requires the '%1' clause">;
+def warn_omp_gpu_unsupported_clause: Warning<
+  "clause '%0' is currently not supported on a GPU; clause ignored">,
+  InGroup<OpenMPClauses>;
+def warn_omp_gpu_unsupported_modifier_for_clause: Warning<
+  "modifier '%0' is currently not supported on a GPU for the '%1' clause; modifier ignored">,
+  InGroup<OpenMPClauses>;
 
 // Static Analyzer Core
 def err_unknown_analyzer_checker_or_package : Error<

clang/lib/CodeGen/CGOpenMPRuntime.cpp

@@ -2703,7 +2703,8 @@ llvm::Value *CGOpenMPRuntime::emitForNext(CodeGenFunction &CGF,
 }
 
 llvm::Value *CGOpenMPRuntime::emitMessageClause(CodeGenFunction &CGF,
-                                                const Expr *Message) {
+                                                const Expr *Message,
+                                                SourceLocation Loc) {
   if (!Message)
     return llvm::ConstantPointerNull::get(CGF.VoidPtrTy);
   return CGF.EmitScalarExpr(Message);
@@ -2713,11 +2714,13 @@ llvm::Value *
 CGOpenMPRuntime::emitMessageClause(CodeGenFunction &CGF,
                                    const OMPMessageClause *MessageClause) {
   return emitMessageClause(
-      CGF, MessageClause ? MessageClause->getMessageString() : nullptr);
+      CGF, MessageClause ? MessageClause->getMessageString() : nullptr,
+      MessageClause->getBeginLoc());
 }
 
 llvm::Value *
-CGOpenMPRuntime::emitSeverityClause(OpenMPSeverityClauseKind Severity) {
+CGOpenMPRuntime::emitSeverityClause(OpenMPSeverityClauseKind Severity,
+                                    SourceLocation Loc) {
   // OpenMP 6.0, 10.4: "If no severity clause is specified then the effect is
   // as if sev-level is fatal."
   return llvm::ConstantInt::get(CGM.Int32Ty,
@@ -2727,13 +2730,15 @@ CGOpenMPRuntime::emitSeverityClause(OpenMPSeverityClauseKind Severity) {
 llvm::Value *
 CGOpenMPRuntime::emitSeverityClause(const OMPSeverityClause *SeverityClause) {
   return emitSeverityClause(SeverityClause ? SeverityClause->getSeverityKind()
-                                           : OMPC_SEVERITY_unknown);
+                                           : OMPC_SEVERITY_unknown,
+                            SeverityClause->getBeginLoc());
 }
 
 void CGOpenMPRuntime::emitNumThreadsClause(
     CodeGenFunction &CGF, llvm::Value *NumThreads, SourceLocation Loc,
     OpenMPNumThreadsClauseModifier Modifier, OpenMPSeverityClauseKind Severity,
-    const Expr *Message) {
+    SourceLocation SeverityLoc, const Expr *Message,
+    SourceLocation MessageLoc) {
   if (!CGF.HaveInsertPoint())
     return;
   llvm::SmallVector<llvm::Value *, 4> Args(
@@ -2745,8 +2750,8 @@ void CGOpenMPRuntime::emitNumThreadsClause(
   RuntimeFunction FnID = OMPRTL___kmpc_push_num_threads;
   if (Modifier == OMPC_NUMTHREADS_strict) {
     FnID = OMPRTL___kmpc_push_num_threads_strict;
-    Args.push_back(emitSeverityClause(Severity));
-    Args.push_back(emitMessageClause(CGF, Message));
+    Args.push_back(emitSeverityClause(Severity, SeverityLoc));
+    Args.push_back(emitMessageClause(CGF, Message, MessageLoc));
   }
   CGF.EmitRuntimeCall(
       OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(), FnID), Args);
@@ -12654,7 +12659,8 @@ llvm::Value *CGOpenMPSIMDRuntime::emitForNext(CodeGenFunction &CGF,
 void CGOpenMPSIMDRuntime::emitNumThreadsClause(
     CodeGenFunction &CGF, llvm::Value *NumThreads, SourceLocation Loc,
     OpenMPNumThreadsClauseModifier Modifier, OpenMPSeverityClauseKind Severity,
-    const Expr *Message) {
+    SourceLocation SeverityLoc, const Expr *Message,
+    SourceLocation MessageLoc) {
   llvm_unreachable("Not supported in SIMD-only mode");
 }
 

clang/lib/CodeGen/CGOpenMPRuntime.h

@@ -1049,11 +1049,13 @@ class CGOpenMPRuntime {
                                    Address UB, Address ST);
 
   virtual llvm::Value *emitMessageClause(CodeGenFunction &CGF,
-                                         const Expr *Message);
+                                         const Expr *Message,
+                                         SourceLocation Loc);
   virtual llvm::Value *emitMessageClause(CodeGenFunction &CGF,
                                          const OMPMessageClause *MessageClause);
 
-  virtual llvm::Value *emitSeverityClause(OpenMPSeverityClauseKind Severity);
+  virtual llvm::Value *emitSeverityClause(OpenMPSeverityClauseKind Severity,
+                                          SourceLocation Loc);
   virtual llvm::Value *
   emitSeverityClause(const OMPSeverityClause *SeverityClause);
 
@@ -1069,7 +1071,9 @@ class CGOpenMPRuntime {
       CodeGenFunction &CGF, llvm::Value *NumThreads, SourceLocation Loc,
       OpenMPNumThreadsClauseModifier Modifier = OMPC_NUMTHREADS_unknown,
       OpenMPSeverityClauseKind Severity = OMPC_SEVERITY_fatal,
-      const Expr *Message = nullptr);
+      SourceLocation SeverityLoc = SourceLocation(),
+      const Expr *Message = nullptr,
+      SourceLocation MessageLoc = SourceLocation());
 
   /// Emit call to void __kmpc_push_proc_bind(ident_t *loc, kmp_int32
   /// global_tid, int proc_bind) to generate code for 'proc_bind' clause.
@@ -1956,7 +1960,9 @@ class CGOpenMPSIMDRuntime final : public CGOpenMPRuntime {
       CodeGenFunction &CGF, llvm::Value *NumThreads, SourceLocation Loc,
       OpenMPNumThreadsClauseModifier Modifier = OMPC_NUMTHREADS_unknown,
       OpenMPSeverityClauseKind Severity = OMPC_SEVERITY_fatal,
-      const Expr *Message = nullptr) override;
+      SourceLocation SeverityLoc = SourceLocation(),
+      const Expr *Message = nullptr,
+      SourceLocation MessageLoc = SourceLocation()) override;
 
   /// Emit call to void __kmpc_push_proc_bind(ident_t *loc, kmp_int32
   /// global_tid, int proc_bind) to generate code for 'proc_bind' clause.

clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

@@ -899,10 +899,34 @@ void CGOpenMPRuntimeGPU::emitProcBindClause(CodeGenFunction &CGF,
   // Nothing to do.
 }
 
+llvm::Value *CGOpenMPRuntimeGPU::emitMessageClause(CodeGenFunction &CGF,
+                                                   const Expr *Message,
+                                                   SourceLocation Loc) {
+  CGM.getDiags().Report(Loc, diag::warn_omp_gpu_unsupported_clause)
+      << getOpenMPClauseName(OMPC_message);
+  return nullptr;
+}
+
+llvm::Value *
+CGOpenMPRuntimeGPU::emitSeverityClause(OpenMPSeverityClauseKind Severity,
+                                       SourceLocation Loc) {
+  CGM.getDiags().Report(Loc, diag::warn_omp_gpu_unsupported_clause)
+      << getOpenMPClauseName(OMPC_severity);
+  return nullptr;
+}
+
 void CGOpenMPRuntimeGPU::emitNumThreadsClause(
     CodeGenFunction &CGF, llvm::Value *NumThreads, SourceLocation Loc,
     OpenMPNumThreadsClauseModifier Modifier, OpenMPSeverityClauseKind Severity,
-    const Expr *Message) {
+    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.
 }
 

clang/lib/CodeGen/CGOpenMPRuntimeGPU.h

@@ -162,14 +162,24 @@ class CGOpenMPRuntimeGPU : public CGOpenMPRuntime {
                           llvm::omp::ProcBindKind ProcBind,
                           SourceLocation Loc) override;
 
+  // Currently unsupported on the device.
+  llvm::Value *emitMessageClause(CodeGenFunction &CGF, const Expr *Message,
+                                 SourceLocation Loc) override;
+
+  // Currently unsupported on the device.
+  virtual llvm::Value *emitSeverityClause(OpenMPSeverityClauseKind Severity,
+                                          SourceLocation Loc) override;
+
   /// Emits call to void __kmpc_push_num_threads(ident_t *loc, kmp_int32
   /// global_tid, kmp_int32 num_threads) to generate code for 'num_threads'
   /// clause.
   void emitNumThreadsClause(
       CodeGenFunction &CGF, llvm::Value *NumThreads, SourceLocation Loc,
       OpenMPNumThreadsClauseModifier Modifier = OMPC_NUMTHREADS_unknown,
       OpenMPSeverityClauseKind Severity = OMPC_SEVERITY_fatal,
-      const Expr *Message = nullptr) override;
+      SourceLocation SeverityLoc = SourceLocation(),
+      const Expr *Message = nullptr,
+      SourceLocation MessageLoc = SourceLocation()) override;
 
   /// This function ought to emit, in the general case, a call to
   // the openmp runtime kmpc_push_num_teams. In NVPTX backend it is not needed

clang/lib/CodeGen/CGStmtOpenMP.cpp

@@ -1622,22 +1622,30 @@ static void emitCommonOMPParallelDirective(
   // if sev-level is fatal."
   OpenMPSeverityClauseKind Severity = OMPC_SEVERITY_fatal;
   clang::Expr *Message = nullptr;
+  SourceLocation SeverityLoc = SourceLocation();
+  SourceLocation MessageLoc = SourceLocation();
+
   llvm::Function *OutlinedFn =
       CGF.CGM.getOpenMPRuntime().emitParallelOutlinedFunction(
           CGF, S, *CS->getCapturedDecl()->param_begin(), InnermostKind,
           CodeGen);
+
   if (const auto *NumThreadsClause = S.getSingleClause<OMPNumThreadsClause>()) {
     CodeGenFunction::RunCleanupsScope NumThreadsScope(CGF);
     NumThreads = CGF.EmitScalarExpr(NumThreadsClause->getNumThreads(),
                                     /*IgnoreResultAssign=*/true);
     Modifier = NumThreadsClause->getModifier();
-    if (const auto *MessageClause = S.getSingleClause<OMPMessageClause>())
+    if (const auto *MessageClause = S.getSingleClause<OMPMessageClause>()) {
       Message = MessageClause->getMessageString();
-    if (const auto *SeverityClause = S.getSingleClause<OMPSeverityClause>())
+      MessageLoc = MessageClause->getBeginLoc();
+    }
+    if (const auto *SeverityClause = S.getSingleClause<OMPSeverityClause>()) {
       Severity = SeverityClause->getSeverityKind();
+      SeverityLoc = SeverityClause->getBeginLoc();
+    }
     CGF.CGM.getOpenMPRuntime().emitNumThreadsClause(
         CGF, NumThreads, NumThreadsClause->getBeginLoc(), Modifier, Severity,
-        Message);
+        SeverityLoc, Message, MessageLoc);
   }
   if (const auto *ProcBindClause = S.getSingleClause<OMPProcBindClause>()) {
     CodeGenFunction::RunCleanupsScope ProcBindScope(CGF);

clang/test/OpenMP/amdgcn_parallel_num_threads_strict_messages.cpp

@@ -0,0 +1,108 @@
+// RUN: %clang_cc1 -DF1 -verify -fopenmp -fopenmp-version=60 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host-ppc.bc
+// RUN: %clang_cc1 -DF1 -DTARGET -verify -fopenmp -fopenmp-version=60 -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host-ppc.bc -o /dev/null
+// RUN: %clang_cc1 -DF2 -verify -fopenmp -fopenmp-version=60 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host-ppc.bc
+// RUN: %clang_cc1 -DF2 -DTARGET -verify -fopenmp -fopenmp-version=60 -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host-ppc.bc -o /dev/null
+// RUN: %clang_cc1 -DF3 -verify -fopenmp -fopenmp-version=60 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host-ppc.bc
+// RUN: %clang_cc1 -DF3 -DTARGET -verify -fopenmp -fopenmp-version=60 -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host-ppc.bc -o /dev/null
+
+#ifndef TARGET
+// expected-no-diagnostics
+#endif
+
+#ifdef F3
+template<typename tx>
+tx ftemplate(int n) {
+  tx a = 0;
+
+#ifdef TARGET
+  // expected-warning@+2 {{modifier 'strict' is currently not supported on a GPU for the 'num_threads' clause; modifier ignored}}
+#endif
+  #pragma omp parallel num_threads(strict: tx(20)) severity(fatal) message("msg")
+  {
+  }
+
+  short b = 1;
+#ifdef TARGET
+  // expected-warning@+2 {{modifier 'strict' is currently not supported on a GPU for the 'num_threads' clause; modifier ignored}}
+#endif
+  #pragma omp parallel num_threads(strict: b) severity(warning) message("msg")
+  {
+    a += b;
+  }
+
+  return a;
+}
+#endif
+
+#ifdef F2
+static
+int fstatic(int n) {
+
+#ifdef TARGET
+  // expected-warning@+2 {{modifier 'strict' is currently not supported on a GPU for the 'num_threads' clause; modifier ignored}}
+#endif
+  #pragma omp target parallel num_threads(strict: n) message("msg")
+  {
+  }
+
+#ifdef TARGET
+  // expected-warning@+2 {{modifier 'strict' is currently not supported on a GPU for the 'num_threads' clause; modifier ignored}}
+#endif
+  #pragma omp target parallel num_threads(strict: 32+n) severity(warning)
+  {
+  }
+
+  return n+1;
+}
+#endif
+
+#ifdef F1
+struct S1 {
+  double a;
+
+  int r1(int n){
+    int b = 1;
+
+#ifdef TARGET
+    // expected-warning@+2 {{modifier 'strict' is currently not supported on a GPU for the 'num_threads' clause; modifier ignored}}
+#endif
+    #pragma omp parallel num_threads(strict: n-b) severity(warning) message("msg")
+    {
+      this->a = (double)b + 1.5;
+    }
+
+#ifdef TARGET
+    // expected-warning@+2 {{modifier 'strict' is currently not supported on a GPU for the 'num_threads' clause; modifier ignored}}
+#endif
+    #pragma omp parallel num_threads(strict: 1024) severity(fatal)
+    {
+      this->a = 2.5;
+    }
+
+    return (int)a;
+  }
+};
+#endif
+
+int bar(int n){
+  int a = 0;
+
+#ifdef F1
+  #pragma omp target
+  {
+    S1 S;
+    a += S.r1(n);
+  }
+#endif
+
+#ifdef F2
+  a += fstatic(n);
+#endif
+
+#ifdef F3
+  #pragma omp target
+  a += ftemplate<int>(n);
+#endif
+
+  return a;
+}