[CUDA][HIP] Fix host/device check with -fopenmp

[yaxunl]

Summary

CUDA/HIP program may be compiled with -fopenmp. In this case, -fopenmp is only passed to host compilation
to take advantages of multi-threads computation.

CUDA/HIP and OpenMP both use Sema::DeviceCallGraph to store functions to be analyzed and remove them
once they decide the function is sure to be emitted. CUDA/HIP and OpenMP have different functions to determine
if a function is sure to be emitted.

To check host/device correctly for CUDA/HIP when -fopenmp is enabled, there needs a unified logic to determine
whether a function is to be emitted. The logic needs to be aware of both CUDA and OpenMP logic.

This patch only affects CUDA/HIP program which are compiled with -fopenmp. It should have no effect on C++ programs
compiled with -fopenmp.

[rjmccall]

lib/Sema/Sema.cpp

1511 ↗

(On Diff #221040)

There's an overload of DenseMap::erase that just takes a key value, so this whole thing can be S.DeviceCallGraph.erase(OrigCallee);.

Why do we need to erase the entry instead of re-using it? If the call graphs are different for the two use-cases, is that conflict a problem for other reasons?

[yaxunl]

reuse the call tree.

[yaxunl]

lib/Sema/Sema.cpp

1511 ↗

(On Diff #221040)

I think you are right. we should reuse the call graph.

[rjmccall]

lib/Sema/Sema.cpp

1511 ↗

(On Diff #221040)

Okay. I believe DenseMap has a count that tells you how many entries there are for a key.

[yaxunl]

revised by John's comments

[rjmccall]

Okay. And it's okay to fall down to the code below when functions are used in both ways this way?

[yaxunl]

In D67837#1679670, @rjmccall wrote:

Okay. And it's okay to fall down to the code below when functions are used in both ways this way?

This part of code is for delayed checking of hostness. If a host function calls device function, we do not want to diagnose it unless we are sure the host func is to be emitted.

whether a func is emitted is determined by IsKnownEmitted function. HIP and OpenMP each has its own copy.

In this case, HIP thinks a static host func is not sure to be emitted, unless it has an external linkage.

However, OpenMP thinks a static host func is always emitted.

If we do not consolidate IsKnownEmitted function of HIP and OpenMP, and use the current fix, for HIP, some additional diagnostics will be emitted compared to without -fopenmp in certain cases.

It seems I should figure out whether the static host func is really emitted for HIP then consolidate IsKnownEmitted func for HIP and OpenMP.

[rjmccall]

If that's tractable, then that does seem like the best solution. You can commit this if you need a shorter-term fix.

[yaxunl]

Unify CUDA/HIP/OpenMP host/device check.

[rjmccall]

That looks great. I'm probably not the best person to validate that the OMP / CUDA logic is still right, but it seems like the right technical design.

[yaxunl]

ping
@ABataev Any comments? Thanks.

[ABataev]

lib/Sema/SemaDecl.cpp

17618 ↗

(On Diff #221829)

Are you going to handle #pragma omp declare target device_type(nohost) in Cuda mode correctly? Such functions should not be emitted on the host in OpenMP 5.0.

lib/Sema/SemaOpenMP.cpp

1569–1574 ↗

(On Diff #221829)

You can remove the whole function and use Sema::getEmissionStatus() instead.

1593 ↗

(On Diff #221829)

I think, CUDADiscarded case must be unreachable and it must be a case for llvm_unreachable() in case of the OpenMP device code.

1602–1607 ↗

(On Diff #221829)

You can remove the whole function and use Sema::getEmissionStatus() instead.

1647–1648 ↗

(On Diff #221829)

I would add an assert thet this function does not return CUDADiscarded value.

1684–1685 ↗

(On Diff #221829)

Also, it would be good to check that it cannot return CUDADiscarded here.

[yaxunl]

lib/Sema/SemaDecl.cpp

17618 ↗

(On Diff #221829)

will do

lib/Sema/SemaOpenMP.cpp

1569–1574 ↗

(On Diff #221829)

done

1593 ↗

(On Diff #221829)

done

1602–1607 ↗

(On Diff #221829)

done

1647–1648 ↗

(On Diff #221829)

done

1684–1685 ↗

(On Diff #221829)

In host compilation, if the source code is CUDA/HIP language, a CUDA/HIP device function may be discarded as CUDADiscarded and show up here. For C++, CUDADiscarded should not be returned here.

[yaxunl]

Revised by Alexey's comments.

[ABataev]

lib/Sema/SemaCUDA.cpp

604 ↗

(On Diff #223284)

I believe this function can be removed

lib/Sema/SemaDecl.cpp

17619 ↗

(On Diff #223284)

Better to use real type here, not auto

17623–17626 ↗

(On Diff #223284)

Enclose in braces

17626 ↗

(On Diff #223284)

Hmm, it must be marked as know-emitted only if S.DeviceKnownEmittedFns.count(FD) > 0. Otherwise, it must be unknown.

17629–17631 ↗

(On Diff #223284)

Enclose in braces, not goo to have else branch enclosed in braces and then branch without.

17641 ↗

(On Diff #223284)

Same here, it must be marked as know-emitted only if S.DeviceKnownEmittedFns.count(FD) > 0. Otherwise, it must be unknown.

17684–17689 ↗

(On Diff #223284)

Reformat the comment here

lib/Sema/SemaOpenMP.cpp

1629–1630 ↗

(On Diff #223284)

The same condition is checked twice, one of them must be for CalleeS, I believe

1674 ↗

(On Diff #223284)

Again, the same condition checked twice.

[yaxunl]

revised by Alexey's comments.

[ABataev]

lib/Sema/SemaCUDA.cpp

616 ↗

(On Diff #223712)

I assume, better to use cast here, not dyn_cast

645 ↗

(On Diff #223712)

Same here, just cast

lib/Sema/SemaOpenMP.cpp

1627–1628 ↗

(On Diff #223712)

Better to use real type instead of auto here

1670–1671 ↗

(On Diff #223712)

Real types instead of auto

[yaxunl]

revised by Alexey's comments.

lib/Sema/SemaCUDA.cpp

616 ↗

(On Diff #223712)

done

645 ↗

(On Diff #223712)

done

604 ↗

(On Diff #223284)

done

lib/Sema/SemaDecl.cpp

17619 ↗

(On Diff #223284)

done

17623–17626 ↗

(On Diff #223284)

done

17626 ↗

(On Diff #223284)

done

17629–17631 ↗

(On Diff #223284)

done

17641 ↗

(On Diff #223284)

done

17684–17689 ↗

(On Diff #223284)

done

lib/Sema/SemaOpenMP.cpp

1627–1628 ↗

(On Diff #223712)

done

1670–1671 ↗

(On Diff #223712)

done

1629–1630 ↗

(On Diff #223284)

done

1674 ↗

(On Diff #223284)

done

[ABataev]

LG

clang/include/clang/Sema/Sema.h

public:
   LiteralOperatorLookupResult LookupLiteralOperator(Scope *S, LookupResult &R,
                                                     ArrayRef<QualType> ArgTys,
                                                     bool AllowRaw,
                                                     bool AllowTemplate,
                                                     bool AllowStringTemplate,
                                                     bool DiagnoseMissing);
   bool isKnownName(StringRef name);
 
+  /// Status of the function emission on the CUDA/HIP/OpenMP host/device attrs.
+  enum class FunctionEmissionStatus {
+    Emitted,
+    CUDADiscarded,     // Discarded due to CUDA/HIP hostness
+    OMPDiscarded,      // Discarded due to OpenMP hostness
+    TemplateDiscarded, // Discarded due to uninstantiated templates
+    Unknown,
+  };
+  FunctionEmissionStatus getEmissionStatus(FunctionDecl *Decl);
+
+  // Whether the callee should be ignored in CUDA/HIP/OpenMP host/device check.
+  bool shouldIgnoreInHostDeviceCheck(FunctionDecl *Callee);
+
   void ArgumentDependentLookup(DeclarationName Name, SourceLocation Loc,
                                ArrayRef<Expr *> Args, ADLResult &Functions);
 
   void LookupVisibleDecls(Scope *S, LookupNameKind Kind,
                           VisibleDeclConsumer &Consumer,
                           bool IncludeGlobalScope = true,
                           bool LoadExternal = true);
   void LookupVisibleDecls(DeclContext *Ctx, LookupNameKind Kind,

clang/lib/Sema/SemaCUDA.cpp

if (It != Previous.end()) {
     }
     return;
   }
 
   NewD->addAttr(CUDAHostAttr::CreateImplicit(Context));
   NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
 }
 
-// Do we know that we will eventually codegen the given function?
-static bool IsKnownEmitted(Sema &S, FunctionDecl *FD) {
-  // Templates are emitted when they're instantiated.
-  if (FD->isDependentContext())
-    return false;
-
-  // When compiling for device, host functions are never emitted.  Similarly,
-  // when compiling for host, device and global functions are never emitted.
-  // (Technically, we do emit a host-side stub for global functions, but this
-  // doesn't count for our purposes here.)
-  Sema::CUDAFunctionTarget T = S.IdentifyCUDATarget(FD);
-  if (S.getLangOpts().CUDAIsDevice && T == Sema::CFT_Host)
-    return false;
-  if (!S.getLangOpts().CUDAIsDevice &&
-      (T == Sema::CFT_Device || T == Sema::CFT_Global))
-    return false;
-
-  // Check whether this function is externally visible -- if so, it's
-  // known-emitted.
-  //
-  // We have to check the GVA linkage of the function's *definition* -- if we
-  // only have a declaration, we don't know whether or not the function will be
-  // emitted, because (say) the definition could include "inline".
-  FunctionDecl *Def = FD->getDefinition();
-
-  if (Def &&
-      !isDiscardableGVALinkage(S.getASTContext().GetGVALinkageForFunction(Def)))
-    return true;
-
-  // Otherwise, the function is known-emitted if it's in our set of
-  // known-emitted functions.
-  return S.DeviceKnownEmittedFns.count(FD) > 0;
-}
-
 Sema::DeviceDiagBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc,
                                                    unsigned DiagID) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
   DeviceDiagBuilder::Kind DiagKind = [this] {
     switch (CurrentCUDATarget()) {
     case CFT_Global:
     case CFT_Device:
       return DeviceDiagBuilder::K_Immediate;
     case CFT_HostDevice:
       // An HD function counts as host code if we're compiling for host, and
       // device code if we're compiling for device.  Defer any errors in device
       // mode until the function is known-emitted.
       if (getLangOpts().CUDAIsDevice) {
-        return IsKnownEmitted(*this, dyn_cast<FunctionDecl>(CurContext))
+        return (getEmissionStatus(cast<FunctionDecl>(CurContext)) ==
+                FunctionEmissionStatus::Emitted)
                    ? DeviceDiagBuilder::K_ImmediateWithCallStack
                    : DeviceDiagBuilder::K_Deferred;
       }
       return DeviceDiagBuilder::K_Nop;
 
     default:
       return DeviceDiagBuilder::K_Nop;
     }
case CFT_Host:
       return DeviceDiagBuilder::K_Immediate;
     case CFT_HostDevice:
       // An HD function counts as host code if we're compiling for host, and
       // device code if we're compiling for device.  Defer any errors in device
       // mode until the function is known-emitted.
       if (getLangOpts().CUDAIsDevice)
         return DeviceDiagBuilder::K_Nop;
 
-      return IsKnownEmitted(*this, dyn_cast<FunctionDecl>(CurContext))
+      return (getEmissionStatus(cast<FunctionDecl>(CurContext)) ==
+              FunctionEmissionStatus::Emitted)
                  ? DeviceDiagBuilder::K_ImmediateWithCallStack
                  : DeviceDiagBuilder::K_Deferred;
     default:
       return DeviceDiagBuilder::K_Nop;
     }
   }();
   return DeviceDiagBuilder(DiagKind, Loc, DiagID,
                            dyn_cast<FunctionDecl>(CurContext), *this);
bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
   // FIXME: Is bailing out early correct here?  Should we instead assume that
   // the caller is a global initializer?
   FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext);
   if (!Caller)
     return true;
 
   // If the caller is known-emitted, mark the callee as known-emitted.
   // Otherwise, mark the call in our call graph so we can traverse it later.
-  bool CallerKnownEmitted = IsKnownEmitted(*this, Caller);
+  bool CallerKnownEmitted =
+      getEmissionStatus(Caller) == FunctionEmissionStatus::Emitted;
   if (CallerKnownEmitted) {
     // Host-side references to a __global__ function refer to the stub, so the
     // function itself is never emitted and therefore should not be marked.
-    if (getLangOpts().CUDAIsDevice || IdentifyCUDATarget(Callee) != CFT_Global)
-      markKnownEmitted(*this, Caller, Callee, Loc, IsKnownEmitted);
+    if (!shouldIgnoreInHostDeviceCheck(Callee))
+      markKnownEmitted(
+          *this, Caller, Callee, Loc, [](Sema &S, FunctionDecl *FD) {
+            return S.getEmissionStatus(FD) == FunctionEmissionStatus::Emitted;
+          });
   } else {
     // If we have
     //   host fn calls kernel fn calls host+device,
     // the HD function does not get instantiated on the host.  We model this by
     // omitting at the call to the kernel from the callgraph.  This ensures
     // that, when compiling for host, only HD functions actually called from the
     // host get marked as known-emitted.
-    if (getLangOpts().CUDAIsDevice || IdentifyCUDATarget(Callee) != CFT_Global)
+    if (!shouldIgnoreInHostDeviceCheck(Callee))
       DeviceCallGraph[Caller].insert({Callee, Loc});
   }
 
   DeviceDiagBuilder::Kind DiagKind = [this, Caller, Callee,
                                       CallerKnownEmitted] {
     switch (IdentifyCUDAPreference(Caller, Callee)) {
     case CFP_Never:
       return DeviceDiagBuilder::K_Immediate;

clang/lib/Sema/SemaDecl.cpp

if (PrevDecl && (isa<FunctionDecl>(PrevDecl) || isa<VarDecl>(PrevDecl))) {
     (void)WeakUndeclaredIdentifiers.insert(
       std::pair<IdentifierInfo*,WeakInfo>(AliasName, W));
   }
 }
 
 Decl *Sema::getObjCDeclContext() const {
   return (dyn_cast_or_null<ObjCContainerDecl>(CurContext));
 }
+
+Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD) {
+  // Templates are emitted when they're instantiated.
+  if (FD->isDependentContext())
+    return FunctionEmissionStatus::TemplateDiscarded;
+
+  FunctionEmissionStatus OMPES = FunctionEmissionStatus::Unknown;
+  if (LangOpts.OpenMPIsDevice) {
+    Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
+        OMPDeclareTargetDeclAttr::getDeviceType(FD->getCanonicalDecl());
+    if (DevTy.hasValue()) {
+      if (*DevTy == OMPDeclareTargetDeclAttr::DT_Host)
+        OMPES = FunctionEmissionStatus::OMPDiscarded;
+      else if (DeviceKnownEmittedFns.count(FD) > 0)
+        OMPES = FunctionEmissionStatus::Emitted;
+    }
+  } else if (LangOpts.OpenMP) {
+    // In OpenMP 4.5 all the functions are host functions.
+    if (LangOpts.OpenMP <= 45) {
+      OMPES = FunctionEmissionStatus::Emitted;
+    } else {
+      Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
+          OMPDeclareTargetDeclAttr::getDeviceType(FD->getCanonicalDecl());
+      // In OpenMP 5.0 or above, DevTy may be changed later by
+      // #pragma omp declare target to(*) device_type(*). Therefore DevTy
+      // having no value does not imply host. The emission status will be
+      // checked again at the end of compilation unit.
+      if (DevTy.hasValue()) {
+        if (*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost) {
+          OMPES = FunctionEmissionStatus::OMPDiscarded;
+        } else if (DeviceKnownEmittedFns.count(FD) > 0) {
+          OMPES = FunctionEmissionStatus::Emitted;
+        }
+      }
+    }
+  }
+  if (OMPES == FunctionEmissionStatus::OMPDiscarded ||
+      (OMPES == FunctionEmissionStatus::Emitted && !LangOpts.CUDA))
+    return OMPES;
+
+  if (LangOpts.CUDA) {
+    // When compiling for device, host functions are never emitted.  Similarly,
+    // when compiling for host, device and global functions are never emitted.
+    // (Technically, we do emit a host-side stub for global functions, but this
+    // doesn't count for our purposes here.)
+    Sema::CUDAFunctionTarget T = IdentifyCUDATarget(FD);
+    if (LangOpts.CUDAIsDevice && T == Sema::CFT_Host)
+      return FunctionEmissionStatus::CUDADiscarded;
+    if (!LangOpts.CUDAIsDevice &&
+        (T == Sema::CFT_Device || T == Sema::CFT_Global))
+      return FunctionEmissionStatus::CUDADiscarded;
+
+    // Check whether this function is externally visible -- if so, it's
+    // known-emitted.
+    //
+    // We have to check the GVA linkage of the function's *definition* -- if we
+    // only have a declaration, we don't know whether or not the function will
+    // be emitted, because (say) the definition could include "inline".
+    FunctionDecl *Def = FD->getDefinition();
+
+    if (Def &&
+        !isDiscardableGVALinkage(getASTContext().GetGVALinkageForFunction(Def))
+        && (!LangOpts.OpenMP || OMPES == FunctionEmissionStatus::Emitted))
+      return FunctionEmissionStatus::Emitted;
+  }
+
+  // Otherwise, the function is known-emitted if it's in our set of
+  // known-emitted functions.
+  return (DeviceKnownEmittedFns.count(FD) > 0)
+             ? FunctionEmissionStatus::Emitted
+             : FunctionEmissionStatus::Unknown;
+}
+
+bool Sema::shouldIgnoreInHostDeviceCheck(FunctionDecl *Callee) {
+  // Host-side references to a __global__ function refer to the stub, so the
+  // function itself is never emitted and therefore should not be marked.
+  // If we have host fn calls kernel fn calls host+device, the HD function
+  // does not get instantiated on the host. We model this by omitting at the
+  // call to the kernel from the callgraph. This ensures that, when compiling
+  // for host, only HD functions actually called from the host get marked as
+  // known-emitted.
+  return LangOpts.CUDA && !LangOpts.CUDAIsDevice &&
+         IdentifyCUDATarget(Callee) == CFT_Global;
+}

clang/lib/Sema/SemaOpenMP.cpp

 /// Status of the function emission on the host/device.
 enum class FunctionEmissionStatus {
   Emitted,
   Discarded,
   Unknown,
 };
 } // anonymous namespace
 
-/// Do we know that we will eventually codegen the given function?
-static FunctionEmissionStatus isKnownDeviceEmitted(Sema &S, FunctionDecl *FD) {
-  assert(S.LangOpts.OpenMP && S.LangOpts.OpenMPIsDevice &&
-         "Expected OpenMP device compilation.");
-  // Templates are emitted when they're instantiated.
-  if (FD->isDependentContext())
-    return FunctionEmissionStatus::Discarded;
-
-  Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
-      OMPDeclareTargetDeclAttr::getDeviceType(FD->getCanonicalDecl());
-  if (DevTy.hasValue())
-    return (*DevTy == OMPDeclareTargetDeclAttr::DT_Host)
-               ? FunctionEmissionStatus::Discarded
-               : FunctionEmissionStatus::Emitted;
-
-  // Otherwise, the function is known-emitted if it's in our set of
-  // known-emitted functions.
-  return (S.DeviceKnownEmittedFns.count(FD) > 0)
-             ? FunctionEmissionStatus::Emitted
-             : FunctionEmissionStatus::Unknown;
-}
-
 Sema::DeviceDiagBuilder Sema::diagIfOpenMPDeviceCode(SourceLocation Loc,
                                                      unsigned DiagID) {
   assert(LangOpts.OpenMP && LangOpts.OpenMPIsDevice &&
          "Expected OpenMP device compilation.");
-  FunctionEmissionStatus FES =
-      isKnownDeviceEmitted(*this, getCurFunctionDecl());
+  FunctionEmissionStatus FES = getEmissionStatus(getCurFunctionDecl());
   DeviceDiagBuilder::Kind Kind = DeviceDiagBuilder::K_Nop;
   switch (FES) {
   case FunctionEmissionStatus::Emitted:
     Kind = DeviceDiagBuilder::K_Immediate;
     break;
   case FunctionEmissionStatus::Unknown:
     Kind = isOpenMPDeviceDelayedContext(*this) ? DeviceDiagBuilder::K_Deferred
                                                : DeviceDiagBuilder::K_Immediate;
     break;
-  case FunctionEmissionStatus::Discarded:
+  case FunctionEmissionStatus::TemplateDiscarded:
+  case FunctionEmissionStatus::OMPDiscarded:
     Kind = DeviceDiagBuilder::K_Nop;
     break;
+  case FunctionEmissionStatus::CUDADiscarded:
+    llvm_unreachable("CUDADiscarded unexpected in OpenMP device compilation");
+    break;
   }
 
   return DeviceDiagBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this);
 }
 
-/// Do we know that we will eventually codegen the given function?
-static FunctionEmissionStatus isKnownHostEmitted(Sema &S, FunctionDecl *FD) {
-  assert(S.LangOpts.OpenMP && !S.LangOpts.OpenMPIsDevice &&
-         "Expected OpenMP host compilation.");
-  // In OpenMP 4.5 all the functions are host functions.
-  if (S.LangOpts.OpenMP <= 45)
-    return FunctionEmissionStatus::Emitted;
-
-  Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
-      OMPDeclareTargetDeclAttr::getDeviceType(FD->getCanonicalDecl());
-  if (DevTy.hasValue())
-    return (*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost)
-               ? FunctionEmissionStatus::Discarded
-               : FunctionEmissionStatus::Emitted;
-
-  // Otherwise, the function is known-emitted if it's in our set of
-  // known-emitted functions.
-  return (S.DeviceKnownEmittedFns.count(FD) > 0)
-             ? FunctionEmissionStatus::Emitted
-             : FunctionEmissionStatus::Unknown;
-}
-
 Sema::DeviceDiagBuilder Sema::diagIfOpenMPHostCode(SourceLocation Loc,
                                                    unsigned DiagID) {
   assert(LangOpts.OpenMP && !LangOpts.OpenMPIsDevice &&
          "Expected OpenMP host compilation.");
-  FunctionEmissionStatus FES =
-      isKnownHostEmitted(*this, getCurFunctionDecl());
+  FunctionEmissionStatus FES = getEmissionStatus(getCurFunctionDecl());
   DeviceDiagBuilder::Kind Kind = DeviceDiagBuilder::K_Nop;
   switch (FES) {
   case FunctionEmissionStatus::Emitted:
     Kind = DeviceDiagBuilder::K_Immediate;
     break;
   case FunctionEmissionStatus::Unknown:
     Kind = DeviceDiagBuilder::K_Deferred;
     break;
-  case FunctionEmissionStatus::Discarded:
+  case FunctionEmissionStatus::TemplateDiscarded:
+  case FunctionEmissionStatus::OMPDiscarded:
+  case FunctionEmissionStatus::CUDADiscarded:
     Kind = DeviceDiagBuilder::K_Nop;
     break;
   }
 
   return DeviceDiagBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this);
 }
 
 void Sema::checkOpenMPDeviceFunction(SourceLocation Loc, FunctionDecl *Callee,
                                      bool CheckForDelayedContext) {
   assert(LangOpts.OpenMP && LangOpts.OpenMPIsDevice &&
          "Expected OpenMP device compilation.");
   assert(Callee && "Callee may not be null.");
   Callee = Callee->getMostRecentDecl();
   FunctionDecl *Caller = getCurFunctionDecl();
 
   // host only function are not available on the device.
-  if (Caller &&
-      (isKnownDeviceEmitted(*this, Caller) == FunctionEmissionStatus::Emitted ||
+  if (Caller) {
+    FunctionEmissionStatus CallerS = getEmissionStatus(Caller);
+    FunctionEmissionStatus CalleeS = getEmissionStatus(Callee);
+    assert(CallerS != FunctionEmissionStatus::CUDADiscarded &&
+           CalleeS != FunctionEmissionStatus::CUDADiscarded &&
+           "CUDADiscarded unexpected in OpenMP device function check");
+    if ((CallerS == FunctionEmissionStatus::Emitted ||
          (!isOpenMPDeviceDelayedContext(*this) &&
-        isKnownDeviceEmitted(*this, Caller) ==
-            FunctionEmissionStatus::Unknown)) &&
-      isKnownDeviceEmitted(*this, Callee) ==
-          FunctionEmissionStatus::Discarded) {
-    StringRef HostDevTy =
-        getOpenMPSimpleClauseTypeName(OMPC_device_type, OMPC_DEVICE_TYPE_host);
+          CallerS == FunctionEmissionStatus::Unknown)) &&
+        CalleeS == FunctionEmissionStatus::OMPDiscarded) {
+      StringRef HostDevTy = getOpenMPSimpleClauseTypeName(
+          OMPC_device_type, OMPC_DEVICE_TYPE_host);
       Diag(Loc, diag::err_omp_wrong_device_function_call) << HostDevTy << 0;
       Diag(Callee->getAttr<OMPDeclareTargetDeclAttr>()->getLocation(),
            diag::note_omp_marked_device_type_here)
           << HostDevTy;
       return;
     }
+  }
   // If the caller is known-emitted, mark the callee as known-emitted.
   // Otherwise, mark the call in our call graph so we can traverse it later.
   if ((CheckForDelayedContext && !isOpenMPDeviceDelayedContext(*this)) ||
       (!Caller && !CheckForDelayedContext) ||
-      (Caller &&
-       isKnownDeviceEmitted(*this, Caller) == FunctionEmissionStatus::Emitted))
+      (Caller && getEmissionStatus(Caller) == FunctionEmissionStatus::Emitted))
     markKnownEmitted(*this, Caller, Callee, Loc,
                      [CheckForDelayedContext](Sema &S, FunctionDecl *FD) {
                        return CheckForDelayedContext &&
-                              isKnownDeviceEmitted(S, FD) ==
+                              S.getEmissionStatus(FD) ==
                                   FunctionEmissionStatus::Emitted;
                      });
   else if (Caller)
     DeviceCallGraph[Caller].insert({Callee, Loc});
 }
 
 void Sema::checkOpenMPHostFunction(SourceLocation Loc, FunctionDecl *Callee,
                                    bool CheckCaller) {
   assert(LangOpts.OpenMP && !LangOpts.OpenMPIsDevice &&
          "Expected OpenMP host compilation.");
   assert(Callee && "Callee may not be null.");
   Callee = Callee->getMostRecentDecl();
   FunctionDecl *Caller = getCurFunctionDecl();
 
   // device only function are not available on the host.
-  if (Caller &&
-      isKnownHostEmitted(*this, Caller) == FunctionEmissionStatus::Emitted &&
-      isKnownHostEmitted(*this, Callee) == FunctionEmissionStatus::Discarded) {
+  if (Caller) {
+    FunctionEmissionStatus CallerS = getEmissionStatus(Caller);
+    FunctionEmissionStatus CalleeS = getEmissionStatus(Callee);
+    assert(
+        (LangOpts.CUDA || (CallerS != FunctionEmissionStatus::CUDADiscarded &&
+                           CalleeS != FunctionEmissionStatus::CUDADiscarded)) &&
+        "CUDADiscarded unexpected in OpenMP host function check");
+    if (CallerS == FunctionEmissionStatus::Emitted &&
+        CalleeS == FunctionEmissionStatus::OMPDiscarded) {
       StringRef NoHostDevTy = getOpenMPSimpleClauseTypeName(
           OMPC_device_type, OMPC_DEVICE_TYPE_nohost);
       Diag(Loc, diag::err_omp_wrong_device_function_call) << NoHostDevTy << 1;
       Diag(Callee->getAttr<OMPDeclareTargetDeclAttr>()->getLocation(),
            diag::note_omp_marked_device_type_here)
           << NoHostDevTy;
       return;
     }
+  }
   // If the caller is known-emitted, mark the callee as known-emitted.
   // Otherwise, mark the call in our call graph so we can traverse it later.
+  if (!shouldIgnoreInHostDeviceCheck(Callee)) {
     if ((!CheckCaller && !Caller) ||
         (Caller &&
-       isKnownHostEmitted(*this, Caller) == FunctionEmissionStatus::Emitted))
+         getEmissionStatus(Caller) == FunctionEmissionStatus::Emitted))
       markKnownEmitted(
           *this, Caller, Callee, Loc, [CheckCaller](Sema &S, FunctionDecl *FD) {
             return CheckCaller &&
-                 isKnownHostEmitted(S, FD) == FunctionEmissionStatus::Emitted;
+                   S.getEmissionStatus(FD) == FunctionEmissionStatus::Emitted;
           });
     else if (Caller)
       DeviceCallGraph[Caller].insert({Callee, Loc});
   }
+}
 
 void Sema::checkOpenMPDeviceExpr(const Expr *E) {
   assert(getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice &&
          "OpenMP device compilation mode is expected.");
   QualType Ty = E->getType();
   if ((Ty->isFloat16Type() && !Context.getTargetInfo().hasFloat16Type()) ||
       ((Ty->isFloat128Type() ||
         (Ty->isRealFloatingType() && Context.getTypeSize(Ty) == 128)) &&

clang/test/CodeGenCUDA/openmp-target.cu

+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm \
+// RUN:            -fopenmp -fopenmp-version=50 -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm \
+// RUN:            -fopenmp -fopenmp-version=50 -o - -x c++ %s | FileCheck %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device \
+// RUN:            -emit-llvm -o - %s | FileCheck -check-prefixes=DEV %s
+
+// CHECK: declare{{.*}}@_Z7nohost1v()
+// DEV-NOT: _Z7nohost1v
+void nohost1() {}
+#pragma omp declare target to(nohost1) device_type(nohost)
+
+// CHECK: declare{{.*}}@_Z7nohost2v()
+// DEV-NOT: _Z7nohost2v
+void nohost2() {nohost1();}
+#pragma omp declare target to(nohost2) device_type(nohost)
+

clang/test/OpenMP/declare_target_messages.cpp

 
 #pragma omp declare target link(S) // expected-error {{'S' used in declare target directive is not a variable or a function name}}
 
 #pragma omp declare target (x, x) // expected-error {{'x' appears multiple times in clauses on the same declare target directive}}
 #pragma omp declare target to(x) to(x) // expected-error {{'x' appears multiple times in clauses on the same declare target directive}}
 #pragma omp declare target link(x) // expected-error {{'x' must not appear in both clauses 'to' and 'link'}}
 
 void bazz() {}
-#pragma omp declare target to(bazz) device_type(nohost) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}}
+#pragma omp declare target to(bazz) device_type(nohost) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} host5-note {{marked as 'device_type(nohost)' here}}
 void bazzz() {bazz();}
 #pragma omp declare target to(bazzz) device_type(nohost) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}}
-void any() {bazz();}
+void any() {bazz();} // host5-error {{function with 'device_type(nohost)' is not available on host}}
 void host1() {bazz();}
 #pragma omp declare target to(host1) device_type(host) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} dev5-note 2 {{marked as 'device_type(host)' here}}
 void host2() {bazz();}
 #pragma omp declare target to(host2)
 void device() {host1();}
 #pragma omp declare target to(device) device_type(nohost) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} host5-note 2 {{marked as 'device_type(nohost)' here}}
 void host3() {host1();}
 #pragma omp declare target to(host3)

clang/test/SemaCUDA/call-device-fn-from-host.cu

 // RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - \
 // RUN:   -verify -verify-ignore-unexpected=note
+// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - \
+// RUN:   -verify -verify-ignore-unexpected=note -fopenmp
 
 // Note: This test won't work with -fsyntax-only, because some of these errors
 // are emitted during codegen.
 
 #include "Inputs/cuda.h"
 
 __device__ void device_fn() {}
 // expected-note@-1 5 {{'device_fn' declared here}}

clang/test/SemaCUDA/host-device-constexpr.cu

 // RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs %s
-// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs %s -fcuda-is-device
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs %s \
+// RUN:            -fcuda-is-device
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs \
+// RUN:            -fopenmp %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs \
+// RUN:            -fopenmp %s -fcuda-is-device
 
 #include "Inputs/cuda.h"
 
 // Declares one function and pulls it into namespace ns:
 //
 //   __device__ int OverloadMe();
 //   namespace ns { using ::OverloadMe; }
 //

clang/test/SemaCUDA/openmp-static-func.cu

+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only \
+// RUN:            -verify -fopenmp %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only \
+// RUN:            -verify -fopenmp -x hip %s
+// expected-no-diagnostics
+
+// Tests there is no assertion in Sema::markKnownEmitted when fopenmp is used
+// with CUDA/HIP host compilation.
+
+static void f() {}
+
+static void g() { f(); }
+
+static void h() { g(); }

clang/test/SemaCUDA/openmp-target.cu

+// RUN: %clang_cc1 -triple x86_64 -verify=expected,dev \
+// RUN:            -verify-ignore-unexpected=note \
+// RUN:            -fopenmp -fopenmp-version=50 -o - %s
+// RUN: %clang_cc1 -triple x86_64 -verify -verify-ignore-unexpected=note\
+// RUN:            -fopenmp -fopenmp-version=50 -o - -x c++ %s
+// RUN: %clang_cc1 -triple x86_64 -verify=dev -verify-ignore-unexpected=note\
+// RUN:            -fcuda-is-device -o - %s
+
+#if __CUDA__
+#include "Inputs/cuda.h"
+__device__ void cu_devf();
+#endif
+
+void bazz() {}
+#pragma omp declare target to(bazz) device_type(nohost)
+void bazzz() {bazz();}
+#pragma omp declare target to(bazzz) device_type(nohost)
+void any() {bazz();} // expected-error {{function with 'device_type(nohost)' is not available on host}}
+void host1() {bazz();}
+#pragma omp declare target to(host1) device_type(host)
+void host2() {bazz();}
+#pragma omp declare target to(host2)
+void device() {host1();}
+#pragma omp declare target to(device) device_type(nohost)
+void host3() {host1();}
+#pragma omp declare target to(host3)
+
+#pragma omp declare target
+void any1() {any();}
+void any2() {host1();}
+void any3() {device();} // expected-error {{function with 'device_type(nohost)' is not available on host}}
+void any4() {any2();}
+#pragma omp end declare target
+
+void any5() {any();}
+void any6() {host1();}
+void any7() {device();} // expected-error {{function with 'device_type(nohost)' is not available on host}}
+void any8() {any2();}
+
+#if __CUDA__
+void cu_hostf() { cu_devf(); } // dev-error {{no matching function for call to 'cu_devf'}}
+__device__ void cu_devf2() { cu_hostf(); } // dev-error{{no matching function for call to 'cu_hostf'}}
+#endif