diff --git a/clang/include/clang/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td --- a/clang/include/clang/Basic/DiagnosticGroups.td +++ b/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">; diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/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; + 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 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/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(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(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(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(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(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(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(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(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(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(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(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(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(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(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(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(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(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; diff --git a/clang/test/SemaOpenMP/hip-omp-mix.cpp b/clang/test/SemaOpenMP/hip-omp-mix.cpp new file mode 100644 --- /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-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 + {} +}