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 @@ -1321,6 +1321,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 @@ -8634,6 +8634,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, 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 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -6122,6 +6122,11 @@ 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 ClausesWithImplicit; VarsWithInheritedDSAType VarsWithInheritedDSA; bool ErrorFound = false; @@ -13286,6 +13291,10 @@ 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(AStmt); // 1.2.2 OpenMP Language Terminology // Structured block - An executable statement with a single entry at the @@ -22855,6 +22864,11 @@ 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 @@ (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 --- /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 + {} +}