[OpenMP][clang] Set num_threads 'strict' to unsupported on GPUs

clang/include/clang/Basic/DiagnosticCommonKinds.td

@@ -433,6 +433,8 @@ 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 err_omp_gpu_unsupported_clause: Error<"clause '%0' is currently not supported on a GPU">;
+def err_omp_gpu_unsupported_modifier_for_clause: Error<"modifier '%0' is currently not supported on a GPU for the '%1' clause">;
 
 // 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);
@@ -12263,7 +12268,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::err_omp_gpu_unsupported_clause)
+      << getOpenMPClauseName(OMPC_message);
+  return nullptr;
+}
+
+llvm::Value *
+CGOpenMPRuntimeGPU::emitSeverityClause(OpenMPSeverityClauseKind Severity,
+                                       SourceLocation Loc) {
+  CGM.getDiags().Report(Loc, diag::err_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::err_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-error@+2 {{modifier 'strict' is currently not supported on a GPU for the 'num_threads' clause}}
+#endif
+  #pragma omp parallel num_threads(strict: tx(20)) severity(fatal) message("msg")
+  {
+  }
+
+  short b = 1;
+#ifdef TARGET
+  // expected-error@+2 {{modifier 'strict' is currently not supported on a GPU for the 'num_threads' clause}}
+#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-error@+2 {{modifier 'strict' is currently not supported on a GPU for the 'num_threads' clause}}
+#endif
+  #pragma omp target parallel num_threads(strict: n) message("msg")
+  {
+  }
+
+#ifdef TARGET
+  // expected-error@+2 {{modifier 'strict' is currently not supported on a GPU for the 'num_threads' clause}}
+#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-error@+2 {{modifier 'strict' is currently not supported on a GPU for the 'num_threads' clause}}
+#endif
+    #pragma omp parallel num_threads(strict: n-b) severity(warning) message("msg")
+    {
+      this->a = (double)b + 1.5;
+    }
+
+#ifdef TARGET
+    // expected-error@+2 {{modifier 'strict' is currently not supported on a GPU for the 'num_threads' clause}}
+#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;
+}