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

[mhalk]

Summary

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".

[scchan]

LGTM

[tra]

How is this different from compiling a C++ file with opemnp directives in it? AFAICT neither clang nor gcc issue anywarnings: https://godbolt.org/z/5Me3dnTdr

What makes the warnings necessary for HIP?

[mhalk]

In D145591#4179144, @tra wrote:

How is this different from compiling a C++ file with opemnp directives in it? AFAICT neither clang nor gcc issue anywarnings: https://godbolt.org/z/5Me3dnTdr

What makes the warnings necessary for HIP?

When using HIP target offloading it will "take precedence" over OpenMP target offloading.
That is, there is a use case which compiles "successfully" / without warning but will not execute code within OpenMP target directives.

For example: when using -x hip -fopenmp --offload-arch=... (+ other reasonable parameters) a HIP program with OpenMP target directives will compile without warning / error.
But only host-related OpenMP code will be executed properly.

Another hint that HIP and OpenMP target offloading do not work concurrently ATM is that when compiling with -x hip -fopenmp -fopenmp-targets=... clang will throw an error.

Hence, we want to signal this issue to users, whorun into problems and then have to debug their programs to find out about this.
AFAIK this issue can be alleviated by refactoring / splitting up your code, then linking it back togehter -- but you have to know about this.

[tra]

In D145591#4180908, @mhalk wrote:

For example: when using -x hip -fopenmp --offload-arch=... (+ other reasonable parameters) a HIP program with OpenMP target directives will compile without warning / error.
But only host-related OpenMP code will be executed properly.

Another hint that HIP and OpenMP target offloading do not work concurrently ATM is that when compiling with -x hip -fopenmp -fopenmp-targets=... clang will throw an error.

It sounds like what we want is to make -x hip and -fopenmp mutually exclusive, with a hard error when both are used. If you look at the problem as "-fopenmp completely breaks HIP compilation", a warning is a bit too weak of a measure, IMO.

We can restate it all as "mixing offloading modes is not supported" and generalize it to both CUDA and HIP.

@jhuber6 - WDYT?

[yaxunl]

In D145591#4182016, @tra wrote:

It sounds like what we want is to make -x hip and -fopenmp mutually exclusive, with a hard error when both are used. If you look at the problem as "-fopenmp completely breaks HIP compilation", a warning is a bit too weak of a measure, IMO.

We can restate it all as "mixing offloading modes is not supported" and generalize it to both CUDA and HIP.

-x hip and -fopenmp has been a valid combination. -fopenmp with -x hip allows non-offloading OpenMP directives in host code in HIP. It just ignores the offloading directives.

[tra]

In D145591#4182168, @yaxunl wrote:

-x hip and -fopenmp has been a valid combination. -fopenmp with -x hip allows non-offloading OpenMP directives in host code in HIP. It just ignores the offloading directives.

That brings me back to the earlier question -- what do we currently do when target directives are encountered when we compile a C++ source w/ OpenMP enabled and why HIP shold be handled differently.

If a warning makes sense for target directives with offloading disabled, that warning would be equally applicable to C/C++/CUDA & HIP. If that's not the case, what am I missing?

[jhuber6]

I'm not a fan of the same warning being copied in 24 places. Why do we set LangOpts.IsOpenMP on the GPU compilation side, couldn't we just filter out the -fopenmp or whatever it is for the HIP job?

[jdoerfert]

clang/include/clang/Basic/DiagnosticSemaKinds.td

8638

I doubt the ignored part very much.

[yaxunl]

In D145591#4182351, @tra wrote:

In D145591#4182168, @yaxunl wrote:

-x hip and -fopenmp has been a valid combination. -fopenmp with -x hip allows non-offloading OpenMP directives in host code in HIP. It just ignores the offloading directives.

That brings me back to the earlier question -- what do we currently do when target directives are encountered when we compile a C++ source w/ OpenMP enabled and why HIP shold be handled differently.

If a warning makes sense for target directives with offloading disabled, that warning would be equally applicable to C/C++/CUDA & HIP. If that's not the case, what am I missing?

OpenMP offloading directives (e.g. "omp target") create implicit GPU kernels which require OpenMP toolchain to create offloading actions to support them. For C/C++ programs, OpenMP toolchain can create offloading actions to support the OpenMP offloading directives. For HIP programs, HIP toolchain can only do offloading for HIP, not for OpenMP, therefore the OpenMP offloading directives are ignored. However, the other OpenMP directives e.g. "omp parallel for" works on CPU without offloading toolchain support. Currently, the clang driver and OpenMP/HIP runtime are not ready to support both HIP and OpenMP offloading at the same time.

[yaxunl]

In D145591#4182360, @jhuber6 wrote:

I'm not a fan of the same warning being copied in 24 places. Why do we set LangOpts.IsOpenMP on the GPU compilation side, couldn't we just filter out the -fopenmp or whatever it is for the HIP job?

We cannot filter out -fopenmp for HIP job because the host code in HIP program needs it to support "omp parallel for" etc. Filtering it will break existing HIP programs.

[jhuber6]

In D145591#4182748, @yaxunl wrote:

In D145591#4182360, @jhuber6 wrote:

I'm not a fan of the same warning being copied in 24 places. Why do we set LangOpts.IsOpenMP on the GPU compilation side, couldn't we just filter out the -fopenmp or whatever it is for the HIP job?

We cannot filter out -fopenmp for HIP job because the host code in HIP program needs it to support "omp parallel for" etc. Filtering it will break existing HIP programs.

I mean, couldn't we just prevent the -cc1 arguments for the HIP device compilation from using any OpenMP? Or is that breaking. I figured it was only supported for the CPU portion.

[yaxunl]

In D145591#4182788, @jhuber6 wrote:

In D145591#4182748, @yaxunl wrote:

In D145591#4182360, @jhuber6 wrote:

I'm not a fan of the same warning being copied in 24 places. Why do we set LangOpts.IsOpenMP on the GPU compilation side, couldn't we just filter out the -fopenmp or whatever it is for the HIP job?

We cannot filter out -fopenmp for HIP job because the host code in HIP program needs it to support "omp parallel for" etc. Filtering it will break existing HIP programs.

I mean, couldn't we just prevent the -cc1 arguments for the HIP device compilation from using any OpenMP? Or is that breaking. I figured it was only supported for the CPU portion.

We tried that before but it won't work since that will make users' code inconsistent for host/device compilation and cause all kinds of ODR violations.

[tra]

In D145591#4182737, @yaxunl wrote:

OpenMP offloading directives (e.g. "omp target") create implicit GPU kernels which require OpenMP toolchain to create offloading actions to support them. For C/C++ programs, OpenMP toolchain can create offloading actions to support the OpenMP offloading directives. For HIP programs, HIP toolchain can only do offloading for HIP, not for OpenMP, therefore the OpenMP offloading directives are ignored. However, the other OpenMP directives e.g. "omp parallel for" works on CPU without offloading toolchain support. Currently, the clang driver and OpenMP/HIP runtime are not ready to support both HIP and OpenMP offloading at the same time.

So, we essentially end up with GPU code intended to be used with different runtimes, and we only support one kind at a time. OK. I think I'm getting a rough idea of what's going on.

Ignoring target openmp directives is clearly wrong, given that we do compile with -fopenmp and would normally expect openmp parts to be honored.
Issuing a warning does not feel right. Producing output with a missing key functionality is a bit too severe for a warning, IMO.
Making it an error would make most sense to me, but I'll defer to OpenMP folks on this.

[jdoerfert]

I would emit an error. A warning only if we can ensure the code does something sensible. Right now, I doubt that is the case, similarly I doubt we actually ignore things.

[yaxunl]

HIP toolchain does not pass -fopenmp-targets=amdgcn-amd-amdhsa to clang -cc1 in host compilation. It does not pass -fopenmp-is-device to clang -cc1 in device compilation. Without these options clang will not generate OpenMP offloading code for amdgpu in device and host compilation. The host code should still function correctly.

I tend to emit warning instead of error. I am concerned that emitting error may cause regressions in existing HIP applications whereas warnings can be disabled.

[mhalk]

clang/include/clang/Basic/DiagnosticSemaKinds.td

8638

I just re-checked the LLVM IR of an example.
With -x c++ there are __kmpc_target_(de)init calls; with -x hip there are none.
https://godbolt.org/z/fhds46Pvc

From my point of view, "effectively", target directives seem to have been ignored.
Maybe this is a coincidence (there's definitely some uncertainty) or I am misinterpreting the IR.
I would be glad if you could shed some light on this, if possible / you find the time to do so.

[jdoerfert]

This comment has been deleted.

clang/include/clang/Basic/DiagnosticSemaKinds.td

8638

Line 512 in the HIP output: https://godbolt.org/z/n5s5jz7j9.

[yaxunl]

clang/include/clang/Basic/DiagnosticSemaKinds.td

8638

If I remove --offload-arch for the C++ program, I got the same result as for HIP:

https://godbolt.org/z/4TMzxdfj6

Does that mean we should emit the same error for C++ too?

[jdoerfert]

clang/include/clang/Basic/DiagnosticSemaKinds.td

8638

I honestly do not follow your logic.

As I said in the beginning: I doubt the directive is ignored.

My link showed you that it doesn't ignore the pragma, e.g., the code w/ and w/o the pragma have different results.
This is also true for C++, but that doesn't make the warning any more accurate.

[yaxunl]

clang/include/clang/Basic/DiagnosticSemaKinds.td

8638

I understand that the directive is not ignored and incorrect IR are emitted.

What I mean is that compiling "omp target" with HIP is equivalent to compiling C++ containing "omp target" without -fopenmp-targets specified in clang FE, therefore the error should be emitted in a more generic case. That is:

if "omp target" is compiled without -fopenmp-targets specified in FE, an error should be emitted.

[jdoerfert]

clang/include/clang/Basic/DiagnosticSemaKinds.td

8638

incorrect IR are emitted.

No, that IR is not incorrect. It is exactly what should be emitted.

if "omp target" is compiled without -fopenmp-targets specified in FE, an error should be emitted.

No it should not. OpenMP target w/o an offloading target is well defined OpenMP.

[mhalk]

In D145591#4182360, @jhuber6 wrote:

I'm not a fan of the same warning being copied in 24 places. Why do we set LangOpts.IsOpenMP on the GPU compilation side, couldn't we just filter out the -fopenmp or whatever it is for the HIP job?

If your concern was on the clang / SemA side, I found appropriate helper functions to reduce the number of checks (if (getLangOpts().HIP) { Diag(...) }) from 19 to 4.
But e.g. in the test we will still emit 24 messages, for affected directives.

I doubt reducing the number of diagnostic messages is worth the effort and esp. cost during compilation when tackled e.g. via clang-tooling?! (unsure)
The idea being: emission of a single diagnostic message and pointing out all (24) affected locations.

I may be completely wrong, but I suspect that users affected by "using HIP language but OpenMP target offloading behavior is altered" will (want to) disable this particular diagnostic once they are "aware".

clang/include/clang/Basic/DiagnosticSemaKinds.td

8639

In D145591#4182945, @jdoerfert wrote:

I would emit an error. A warning only if we can ensure the code does something sensible. Right now, I doubt that is the case, similarly I doubt we actually ignore things.

@yaxunl Just wanted to point out: If we added DefaultError here, the SemA would emit an error as preferred by Johannes.
But one could still disable it via -Wno-hip-omp-target-directives even with -Werror (or also demote it via -Wno-error=hip-omp-target-directives).

Would that be acceptable for everyone?

[yaxunl]

clang/include/clang/Basic/DiagnosticSemaKinds.td

8639

LGTM

[mhalk]

Patch rework, implementing the mentioned changes.

[jdoerfert]

clang/include/clang/Basic/DiagnosticSemaKinds.td

8639

That sounds sensible.

[yaxunl]

LGTM. Thanks

clang/include/clang/Basic/DiagnosticGroups.td

 
 // Warning about unknown CUDA SDK version.
 def CudaUnknownVersion: DiagGroup<"unknown-cuda-version">;
 
 // A warning group for warnings about features supported by HIP but
 // 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">;
 
 // A warning group for things that will change semantics in the future.
 def FutureCompat : DiagGroup<"future-compat">;
 
 def InvalidOrNonExistentDirectory : DiagGroup<"invalid-or-nonexistent-directory">;

clang/include/clang/Basic/DiagnosticSemaKinds.td

def err_cuda_device_builtin_surftex_ref_decl : Error<
     "type %1 declared here">;
 def note_cuda_device_builtin_surftex_should_be_template_class : Note<
     "%0 needs to be instantiated from a class template with proper "
     "template arguments">;
 
 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 "
   "string was %3">, InGroup<NonPODVarargs>, DefaultError;
 // The arguments to this diagnostic should match the warning above.
 def err_cannot_pass_objc_interface_to_vararg_format : Error<
   "cannot pass object with interface type %1 by value to variadic "
   "%select{function|block|method|constructor}2; expected type from format "

clang/lib/Sema/SemaOpenMP.cpp

if (const OMPBindClause *BC =
           OMPExecutableDirective::getSingleClause<OMPBindClause>(Clauses))
     BindKind = BC->getBindKind();
   // First check CancelRegion which is then used in checkNestingOfRegions.
   if (checkCancelRegion(*this, Kind, CancelRegion, StartLoc) ||
       checkNestingOfRegions(*this, DSAStack, Kind, DirName, CancelRegion,
                             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;
   ClausesWithImplicit.append(Clauses.begin(), Clauses.end());
   if (AStmt && !CurContext->isDependentContext() && Kind != OMPD_atomic &&
       Kind != OMPD_critical && Kind != OMPD_section && Kind != OMPD_master &&
       Kind != OMPD_masked && !isOpenMPLoopTransformationDirective(Kind)) {
     assert(isa<CapturedStmt>(AStmt) && "Captured statement expected");

 }
 
 StmtResult Sema::ActOnOpenMPTeamsDirective(ArrayRef<OMPClause *> Clauses,
                                            Stmt *AStmt, SourceLocation StartLoc,
                                            SourceLocation EndLoc) {
   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
   // top and a single exit at the bottom.
   // The point of exit cannot be a branch out of the structured block.
   // longjmp() and throw() must not violate the entry/exit criteria.
   CS->getCapturedDecl()->setNothrow();
 
if (!CurLexicalContext->isFileContext() &&
       !CurLexicalContext->isExternCXXContext() &&
       !isa<CXXRecordDecl>(CurLexicalContext) &&
       !isa<ClassTemplateDecl>(CurLexicalContext) &&
       !isa<ClassTemplatePartialSpecializationDecl>(CurLexicalContext) &&
       !isa<ClassTemplateSpecializationDecl>(CurLexicalContext)) {
     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;
 }
 
 const Sema::DeclareTargetContextInfo
 Sema::ActOnOpenMPEndDeclareTargetDirective() {
   assert(!DeclareTargetNesting.empty() &&
          "check isInOpenMPDeclareTargetContext() first!");
assert((isa<VarDecl>(ND) || isa<FunctionDecl>(ND) ||
          "Expected variable, function or function template.");
 
   // Diagnose marking after use as it may lead to incorrect diagnosis and
   // codegen.
   if (LangOpts.OpenMP >= 50 &&
       (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;
 
   auto *VD = cast<ValueDecl>(ND);
   std::optional<OMPDeclareTargetDeclAttr *> ActiveAttr =
       OMPDeclareTargetDeclAttr::getActiveAttr(VD);
   if (ActiveAttr && (*ActiveAttr)->getDevType() != DTCI.DT &&
       (*ActiveAttr)->getLevel() == Level) {

clang/test/SemaOpenMP/hip-omp-mix.cpp

+// 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
+  {}
+}