[clang][HIP][OpenMP] Add warning if mixed HIP / OpenMP offloading

Adds a warning, issued by the clang semantic analysis, if HIP and OpenMP target offloading is requested concurrently.
That is, if HIP language mode is active but OpenMP target directives are encountered.
Previously, a user might not have been aware that target directives are ignored in such a case.

Generation of this warning is (lit-)tested via "make check-clang-semaopenmp".
The warning can be ignored via "-Wno-hip-omp-target-directives".

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

diff --git a/clang/include/clang/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td
index 2b6525c..0d2829d 100644 (file)
--- a/clang/include/clang/Basic/DiagnosticGroups.td
+++ b/clang/include/clang/Basic/DiagnosticGroups.td
@@ -1321,6 +1321,9 @@ def CudaUnknownVersion: DiagGroup<"unknown-cuda-version">;
 // ignored by CUDA.
 def HIPOnly : DiagGroup<"hip-only">;
 
+// Warning about mixed HIP and OpenMP compilation / target offloading.
+def HIPOpenMPOffloading: DiagGroup<"hip-omp-target-directives">;
+
 // Warnings which cause linking of the runtime libraries like
 // libc and the CRT to be skipped.
 def AVRRtlibLinkingQuirks : DiagGroup<"avr-rtlib-linking-quirks">;

diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 7922503..eedc8c1 100644 (file)
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -8634,6 +8634,10 @@ def note_cuda_device_builtin_surftex_should_be_template_class : Note<
 def err_hip_invalid_args_builtin_mangled_name : Error<
     "invalid argument: symbol must be a device-side function or global variable">;
 
+def warn_hip_omp_target_directives : Warning<
+  "HIP does not support OpenMP target directives; directive has been ignored">,
+  InGroup<HIPOpenMPOffloading>, DefaultError;
+
 def warn_non_pod_vararg_with_format_string : Warning<
   "cannot pass %select{non-POD|non-trivial}0 object of type %1 to variadic "
   "%select{function|block|method|constructor}2; expected type from format "

diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 1cd263b..08ca536 100644 (file)
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -6122,6 +6122,11 @@ StmtResult Sema::ActOnOpenMPExecutableDirective(
                             BindKind, StartLoc))
     return StmtError();
 
+  // Report affected OpenMP target offloading behavior when in HIP lang-mode.
+  if (getLangOpts().HIP && (isOpenMPTargetExecutionDirective(Kind) ||
+                            isOpenMPTargetDataManagementDirective(Kind)))
+    Diag(StartLoc, diag::warn_hip_omp_target_directives);
+
   llvm::SmallVector<OMPClause *, 8> ClausesWithImplicit;
   VarsWithInheritedDSAType VarsWithInheritedDSA;
   bool ErrorFound = false;
@@ -13286,6 +13291,10 @@ StmtResult Sema::ActOnOpenMPTeamsDirective(ArrayRef<OMPClause *> Clauses,
   if (!AStmt)
     return StmtError();
 
+  // Report affected OpenMP target offloading behavior when in HIP lang-mode.
+  if (getLangOpts().HIP && (DSAStack->getParentDirective() == OMPD_target))
+    Diag(StartLoc, diag::warn_hip_omp_target_directives);
+
   auto *CS = cast<CapturedStmt>(AStmt);
   // 1.2.2 OpenMP Language Terminology
   // Structured block - An executable statement with a single entry at the
@@ -22855,6 +22864,11 @@ bool Sema::ActOnStartOpenMPDeclareTargetContext(
     Diag(DTCI.Loc, diag::err_omp_region_not_file_context);
     return false;
   }
+
+  // Report affected OpenMP target offloading behavior when in HIP lang-mode.
+  if (getLangOpts().HIP)
+    Diag(DTCI.Loc, diag::warn_hip_omp_target_directives);
+
   DeclareTargetNesting.push_back(DTCI);
   return true;
 }
@@ -22927,6 +22941,10 @@ void Sema::ActOnOpenMPDeclareTargetName(NamedDecl *ND, SourceLocation Loc,
       (ND->isUsed(/*CheckUsedAttr=*/false) || ND->isReferenced()))
     Diag(Loc, diag::warn_omp_declare_target_after_first_use);
 
+  // Report affected OpenMP target offloading behavior when in HIP lang-mode.
+  if (getLangOpts().HIP)
+    Diag(Loc, diag::warn_hip_omp_target_directives);
+
   // Explicit declare target lists have precedence.
   const unsigned Level = -1;
 

diff --git a/clang/test/SemaOpenMP/hip-omp-mix.cpp b/clang/test/SemaOpenMP/hip-omp-mix.cpp
new file mode 100644 (file)
index 0000000..a05bcf4
--- /dev/null
+++ b/clang/test/SemaOpenMP/hip-omp-mix.cpp
@@ -0,0 +1,165 @@
+// REQUIRES: x86-registered-target
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 %s -x c++ -fopenmp -fsyntax-only -verify=host
+// host-no-diagnostics
+
+// RUN: %clang_cc1 %s -x hip -fopenmp -fsyntax-only -verify=device
+// device-error@#01 {{HIP does not support OpenMP target directives}}
+// device-error@#02 {{HIP does not support OpenMP target directives}}
+// device-error@#03 {{HIP does not support OpenMP target directives}}
+// device-error@#04 {{HIP does not support OpenMP target directives}}
+// device-error@#05 {{HIP does not support OpenMP target directives}}
+// device-error@#06 {{HIP does not support OpenMP target directives}}
+// device-error@#07 {{HIP does not support OpenMP target directives}}
+// device-error@#08 {{HIP does not support OpenMP target directives}}
+// device-error@#09 {{HIP does not support OpenMP target directives}}
+// device-error@#10 {{HIP does not support OpenMP target directives}}
+// device-error@#11 {{HIP does not support OpenMP target directives}}
+// device-error@#12 {{HIP does not support OpenMP target directives}}
+// device-error@#13 {{HIP does not support OpenMP target directives}}
+// device-error@#14 {{HIP does not support OpenMP target directives}}
+// device-error@#15 {{HIP does not support OpenMP target directives}}
+// device-error@#16 {{HIP does not support OpenMP target directives}}
+// device-error@#17 {{HIP does not support OpenMP target directives}}
+// device-error@#18 {{HIP does not support OpenMP target directives}}
+// device-error@#19 {{HIP does not support OpenMP target directives}}
+// device-error@#20 {{HIP does not support OpenMP target directives}}
+// device-error@#21 {{HIP does not support OpenMP target directives}}
+// device-error@#22 {{HIP does not support OpenMP target directives}}
+// device-error@#23 {{HIP does not support OpenMP target directives}}
+// device-error@#24 {{HIP does not support OpenMP target directives}}
+
+void test01() {
+#pragma omp target // #01
+  ;
+}
+
+
+void test02() {
+#pragma omp target parallel // #02
+  ;
+}
+
+void test03() {
+#pragma omp target parallel for // #03
+  for (int i = 0; i < 1; ++i);
+}
+
+void test04(int x) {
+#pragma omp target data map(x) // #04
+  ;
+}
+
+void test05(int * x, int n) {
+#pragma omp target enter data map(to:x[:n]) // #05
+}
+
+void test06(int * x, int n) {
+#pragma omp target exit data map(from:x[:n]) // #06
+}
+
+void test07(int * x, int n) {
+#pragma omp target update to(x[:n]) // #07
+}
+
+#pragma omp declare target (test07) // #08
+void test08() {
+
+}
+
+#pragma omp begin declare target // #09
+void test09_1() {
+
+}
+
+void test09_2() {
+
+}
+#pragma omp end declare target
+
+void test10(int n) {
+  #pragma omp target parallel // #10
+  for (int i = 0; i < n; ++i)
+    ;
+}
+
+void test11(int n) {
+  #pragma omp target parallel for // #11
+  for (int i = 0; i < n; ++i)
+    ;
+}
+
+void test12(int n) {
+  #pragma omp target parallel for simd // #12
+  for (int i = 0; i < n; ++i)
+    ;
+}
+
+void test13(int n) {
+  #pragma omp target parallel loop // #13
+  for (int i = 0; i < n; ++i)
+    ;
+}
+
+void test14(int n) {
+  #pragma omp target simd // #14
+  for (int i = 0; i < n; ++i)
+    ;
+}
+
+void test15(int n) {
+  #pragma omp target teams // #15
+  for (int i = 0; i < n; ++i)
+    ;
+}
+
+void test16(int n) {
+  #pragma omp target teams distribute // #16
+  for (int i = 0; i < n; ++i)
+    ;
+}
+
+void test17(int n) {
+  #pragma omp target teams distribute simd // #17
+  for (int i = 0; i < n; ++i)
+    ;
+}
+
+void test18(int n) {
+  #pragma omp target teams loop // #18
+  for (int i = 0; i < n; ++i)
+    ;
+}
+
+void test19(int n) {
+  #pragma omp target teams distribute parallel for // #19
+  for (int i = 0; i < n; ++i)
+    ;
+}
+
+void test20(int n) {
+  #pragma omp target teams distribute parallel for simd // #20
+  for (int i = 0; i < n; ++i)
+    ;
+}
+
+void test21() {
+#pragma omp target // #21
+  {
+#pragma omp teams // #22
+    {}
+  }
+}
+
+void test22() {
+#pragma omp target // #23
+#pragma omp teams // #24
+  {}
+}
+
+void test23() {
+// host code
+#pragma omp teams
+  {}
+}