[clang] 55916de - [clang][HIP][OpenMP] Add warning if mixed HIP / OpenMP offloading
Michael Halkenhaeuser via cfe-commits
cfe-commits at lists.llvm.org
Wed Mar 29 02:46:05 PDT 2023
Author: Michael Halkenhaeuser
Date: 2023-03-29T11:31:17+02:00
New Revision: 55916de2d37742fe334c0726ccf9e584bdaed09f
URL: https://github.com/llvm/llvm-project/commit/55916de2d37742fe334c0726ccf9e584bdaed09f
DIFF: https://github.com/llvm/llvm-project/commit/55916de2d37742fe334c0726ccf9e584bdaed09f.diff
LOG: [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
Added:
clang/test/SemaOpenMP/hip-omp-mix.cpp
Modified:
clang/include/clang/Basic/DiagnosticGroups.td
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/lib/Sema/SemaOpenMP.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td
index 2b6525ca5f866..0d2829d64501f 100644
--- 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 79225036f7f2b..eedc8c18f93dc 100644
--- 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 1cd263b8a5b1c..08ca536e34631 100644
--- 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
index 0000000000000..a05bcf49c0e06
--- /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
+ {}
+}
More information about the cfe-commits
mailing list