mhalk created this revision.
Herald added subscribers: guansong, yaxunl.
Herald added a project: All.
mhalk requested review of this revision.
Herald added a reviewer: jdoerfert.
Herald added subscribers: cfe-commits, sstefan1.
Herald added a project: clang.
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".
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D145591
Files:
clang/include/clang/Basic/DiagnosticGroups.td
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/lib/Sema/SemaOpenMP.cpp
clang/test/SemaOpenMP/hip-omp-mix.cpp
Index: clang/test/SemaOpenMP/hip-omp-mix.cpp
===================================================================
--- /dev/null
+++ 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-warning@#01 {{HIP does not support OpenMP target directives}}
+// device-warning@#02 {{HIP does not support OpenMP target directives}}
+// device-warning@#03 {{HIP does not support OpenMP target directives}}
+// device-warning@#04 {{HIP does not support OpenMP target directives}}
+// device-warning@#05 {{HIP does not support OpenMP target directives}}
+// device-warning@#06 {{HIP does not support OpenMP target directives}}
+// device-warning@#07 {{HIP does not support OpenMP target directives}}
+// device-warning@#08 {{HIP does not support OpenMP target directives}}
+// device-warning@#09 {{HIP does not support OpenMP target directives}}
+// device-warning@#10 {{HIP does not support OpenMP target directives}}
+// device-warning@#11 {{HIP does not support OpenMP target directives}}
+// device-warning@#12 {{HIP does not support OpenMP target directives}}
+// device-warning@#13 {{HIP does not support OpenMP target directives}}
+// device-warning@#14 {{HIP does not support OpenMP target directives}}
+// device-warning@#15 {{HIP does not support OpenMP target directives}}
+// device-warning@#16 {{HIP does not support OpenMP target directives}}
+// device-warning@#17 {{HIP does not support OpenMP target directives}}
+// device-warning@#18 {{HIP does not support OpenMP target directives}}
+// device-warning@#19 {{HIP does not support OpenMP target directives}}
+// device-warning@#20 {{HIP does not support OpenMP target directives}}
+// device-warning@#21 {{HIP does not support OpenMP target directives}}
+// device-warning@#22 {{HIP does not support OpenMP target directives}}
+// device-warning@#23 {{HIP does not support OpenMP target directives}}
+// device-warning@#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
+ {}
+}
Index: clang/lib/Sema/SemaOpenMP.cpp
===================================================================
--- clang/lib/Sema/SemaOpenMP.cpp
+++ clang/lib/Sema/SemaOpenMP.cpp
@@ -10635,6 +10635,9 @@
DSAStack))
return StmtError();
+ if (getLangOpts().HIP)
+ 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
@@ -10732,6 +10735,9 @@
DSAStack))
return StmtError();
+ if (getLangOpts().HIP)
+ 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
@@ -12959,6 +12965,9 @@
if (!AStmt)
return StmtError();
+ if (getLangOpts().HIP)
+ 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
@@ -13024,6 +13033,9 @@
if (!AStmt)
return StmtError();
+ if (getLangOpts().HIP)
+ 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
@@ -13055,6 +13067,9 @@
if (!AStmt)
return StmtError();
+ if (getLangOpts().HIP)
+ 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
@@ -13143,6 +13158,9 @@
if (!AStmt)
return StmtError();
+ if (getLangOpts().HIP)
+ Diag(StartLoc, diag::warn_hip_omp_target_directives);
+
assert(isa<CapturedStmt>(AStmt) && "Captured statement expected");
// OpenMP [2.12.2, target data Construct, Restrictions]
@@ -13173,6 +13191,9 @@
if (!AStmt)
return StmtError();
+ if (getLangOpts().HIP)
+ 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
@@ -13210,6 +13231,9 @@
if (!AStmt)
return StmtError();
+ if (getLangOpts().HIP)
+ 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
@@ -13247,6 +13271,9 @@
if (!AStmt)
return StmtError();
+ if (getLangOpts().HIP)
+ 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
@@ -13285,6 +13312,10 @@
if (!AStmt)
return StmtError();
+ // Report if the parent is a target directive while HIP offloading
+ 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
@@ -14064,6 +14095,9 @@
if (!AStmt)
return StmtError();
+ if (getLangOpts().HIP)
+ 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
@@ -14119,6 +14153,9 @@
if (!AStmt)
return StmtError();
+ if (getLangOpts().HIP)
+ 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
@@ -14393,6 +14430,9 @@
if (!AStmt)
return StmtError();
+ if (getLangOpts().HIP)
+ 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
@@ -14423,6 +14463,9 @@
if (!AStmt)
return StmtError();
+ if (getLangOpts().HIP)
+ 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
@@ -14466,6 +14509,9 @@
if (!AStmt)
return StmtError();
+ if (getLangOpts().HIP)
+ 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
@@ -14521,6 +14567,9 @@
if (!AStmt)
return StmtError();
+ if (getLangOpts().HIP)
+ 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
@@ -14580,6 +14629,9 @@
if (!AStmt)
return StmtError();
+ if (getLangOpts().HIP)
+ 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
@@ -22854,6 +22906,10 @@
Diag(DTCI.Loc, diag::err_omp_region_not_file_context);
return false;
}
+
+ if (getLangOpts().HIP)
+ Diag(DTCI.Loc, diag::warn_hip_omp_target_directives);
+
DeclareTargetNesting.push_back(DTCI);
return true;
}
@@ -22926,6 +22982,9 @@
(ND->isUsed(/*CheckUsedAttr=*/false) || ND->isReferenced()))
Diag(Loc, diag::warn_omp_declare_target_after_first_use);
+ if (getLangOpts().HIP)
+ Diag(Loc, diag::warn_hip_omp_target_directives);
+
// Explicit declare target lists have precedence.
const unsigned Level = -1;
Index: clang/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -8630,6 +8630,10 @@
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>;
+
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 "
Index: clang/include/clang/Basic/DiagnosticGroups.td
===================================================================
--- clang/include/clang/Basic/DiagnosticGroups.td
+++ clang/include/clang/Basic/DiagnosticGroups.td
@@ -1319,6 +1319,9 @@
// 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">;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
