[OpenMP] Add support for handling declare target to clause when unified memory is required

[gtbercea]

Summary

This patch adds support for the handling of the variables under the declare target to clause.

The variables in this case are handled like link variables are. A pointer is created on the host and then mapped to the device. The runtime will then copy the address of the host variable in the device pointer.

[gtbercea]

ping

[ABataev]

Will review it next week, when I'm back to work, need to think about it a little bit.

[ABataev]

lib/CodeGen/CGExpr.cpp

2297–2319

I think it would be better to merge these 2 functions into 1 emitDeclTargetVarDeclLValue. It should return the correct address for link vars and to vars with unified memory.

lib/CodeGen/CGOpenMPRuntime.cpp

7458–7467

You can merge those pieces of code into one, no need to duplicate. Plus, I don't think we need a new flag for to with unified memory if we can express everything using the existing flags.

lib/CodeGen/CGOpenMPRuntime.h

1123–1127

Same here, better to merge these 2 functions into one getAddrOfDeclareTargetVar

test/OpenMP/nvptx_target_requires_unified_shared_memory.cpp

2

CHECK is the default prefix, no need to specify it.

2

We also need the checks for the device codegen since you're changing something not only in the host codegen, but in the device codegen too. Just extend this test to check for the codegen for the device.

[gtbercea]

• Merge MT_Link and MT_To with unified memory cases.
• Transform switch into if statements.
• Fix declare target attribute checks.

[gtbercea]

Still need to update the test but the rest of the code is updated.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

2540–2544

I think we can use the same suffix in all cases, instead of _link_ptr or _to_ptr just something like _ref_ptr would be good.

7427

Fix the comment too. Also, the name of the variable is not very good, better to rename it to something like MustBeReference or something similar.

9211–9212

Better to have if-else here

9246

|| (*Res == OMPDeclareTargetDeclAttr::MT_To && HasRequiresUnifiedSharedMemory) and fix the message too

lib/CodeGen/CGOpenMPRuntime.h

1128

Better to name it getAddrOfDeclareTargetVar

lib/CodeGen/CodeGenModule.cpp

2488–2489

Better to keep the assert but with a different condition.

[gtbercea]

• Address comments.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

9211–9212

Use assert here instead of else if and remove else with llvm_unreachable

lib/CodeGen/CGOpenMPRuntime.h

1128

Not renamed

[gtbercea]

• Address comments.

[ABataev]

LG

lib/CodeGen/CGDeclCXX.cpp

static void EmitDeclDestroy(CodeGenFunction &CGF, const VarDecl &D,
   // Honor __attribute__((no_destroy)) and bail instead of attempting
   // to emit a reference to a possibly nonexistent destructor, which
   // in turn can cause a crash. This will result in a global constructor
   // that isn't balanced out by a destructor call as intended by the
   // attribute. This also checks for -fno-c++-static-destructors and
   // bails even if the attribute is not present.
   if (D.isNoDestroy(CGF.getContext()))
     return;
-  
   CodeGenModule &CGM = CGF.CGM;
 
   // FIXME:  __attribute__((cleanup)) ?
 
   QualType Type = D.getType();
   QualType::DestructionKind DtorKind = Type.isDestructedType();
 
   switch (DtorKind) {

lib/CodeGen/CGExpr.cpp

 static LValue EmitThreadPrivateVarDeclLValue(
     CodeGenFunction &CGF, const VarDecl *VD, QualType T, Address Addr,
     llvm::Type *RealVarTy, SourceLocation Loc) {
   Addr = CGF.CGM.getOpenMPRuntime().getAddrOfThreadPrivate(CGF, VD, Addr, Loc);
   Addr = CGF.Builder.CreateElementBitCast(Addr, RealVarTy);
   return CGF.MakeAddrLValue(Addr, T, AlignmentSource::Decl);
 }
 
+static Address emitDeclTargetToVarDeclLValue(CodeGenFunction &CGF,
+                                             const VarDecl *VD, QualType T) {
+  llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
+      OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
+  if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_Link)
+    return Address::invalid();
+  assert(*Res == OMPDeclareTargetDeclAttr::MT_To && "Expected to clause");
+  QualType PtrTy = CGF.getContext().getPointerType(VD->getType());
+  Address Addr = CGF.CGM.getOpenMPRuntime().getAddrOfDeclareTargetToUnderUnifiedMem(VD);
+  return CGF.EmitLoadOfPointer(Addr, PtrTy->castAs<PointerType>());
+}
+
 static Address emitDeclTargetLinkVarDeclLValue(CodeGenFunction &CGF,
                                                const VarDecl *VD, QualType T) {
   llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
       OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
   if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_To)
     return Address::invalid();
   assert(*Res == OMPDeclareTargetDeclAttr::MT_Link && "Expected link clause");
   QualType PtrTy = CGF.getContext().getPointerType(VD->getType());
   Address Addr = CGF.CGM.getOpenMPRuntime().getAddrOfDeclareTargetLink(VD);
   return CGF.EmitLoadOfPointer(Addr, PtrTy->castAs<PointerType>());
 }
 
 Address
 CodeGenFunction::EmitLoadOfReference(LValue RefLVal,
                                      LValueBaseInfo *PointeeBaseInfo,
                                      TBAAAccessInfo *PointeeTBAAInfo) {
   llvm::LoadInst *Load = Builder.CreateLoad(RefLVal.getAddress(),
                                             RefLVal.isVolatile());
   CGM.DecorateInstructionWithTBAA(Load, RefLVal.getTBAAInfo());
static LValue EmitGlobalVarDeclLValue(CodeGenFunction &CGF,
   // If it's thread_local, emit a call to its wrapper function instead.
   if (VD->getTLSKind() == VarDecl::TLS_Dynamic &&
       CGF.CGM.getCXXABI().usesThreadWrapperFunction())
     return CGF.CGM.getCXXABI().EmitThreadLocalVarDeclLValue(CGF, VD, T);
   // Check if the variable is marked as declare target with link clause in
   // device codegen.
   if (CGF.getLangOpts().OpenMPIsDevice) {
     Address Addr = emitDeclTargetLinkVarDeclLValue(CGF, VD, T);
+    if (!Addr.isValid() &&
+        CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory())
+      Addr = emitDeclTargetToVarDeclLValue(CGF, VD, T);
     if (Addr.isValid())
       return CGF.MakeAddrLValue(Addr, T, AlignmentSource::Decl);
   }
 
   llvm::Value *V = CGF.CGM.GetAddrOfGlobalVar(VD);
   llvm::Type *RealVarTy = CGF.getTypes().ConvertTypeForMem(VD->getType());
   V = EmitBitCastOfLValueToProperType(CGF, V, RealVarTy);
   CharUnits Alignment = CGF.getContext().getDeclAlign(VD);

lib/CodeGen/CGOpenMPRuntime.h

public:
   /// \param VDAddr Address of the global variable \a VD.
   /// \param Loc Location of the reference to threadprivate var.
   /// \return Address of the threadprivate variable for the current thread.
   virtual Address getAddrOfThreadPrivate(CodeGenFunction &CGF,
                                          const VarDecl *VD,
                                          Address VDAddr,
                                          SourceLocation Loc);
 
+  virtual Address getAddrOfDeclareTargetToUnderUnifiedMem(const VarDecl *VD);
+
   /// Returns the address of the variable marked as declare target with link
   /// clause.
   virtual Address getAddrOfDeclareTargetLink(const VarDecl *VD);
 
   /// Emit a code for initialization of threadprivate variable. It emits
   /// a call to runtime library which adds initial value to the newly created
   /// threadprivate variable (if it is not constant) and registers destructor
   /// for the variable (if any).
   /// \param VD Threadprivate variable.
   /// \param VDAddr Address of the global variable \a VD.
   /// \param Loc Location of threadprivate declaration.
   /// \param PerformInit true if initialization expression is not constant.

lib/CodeGen/CGOpenMPRuntime.cpp

CGOpenMPRuntime::createDispatchNextFunction(unsigned IVSize, bool IVSigned) {
   return CGM.CreateRuntimeFunction(FnTy, Name);
 }
 
 Address CGOpenMPRuntime::getAddrOfDeclareTargetLink(const VarDecl *VD) {
   if (CGM.getLangOpts().OpenMPSimd)
     return Address::invalid();
   llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
       OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
-  if (Res && *Res == OMPDeclareTargetDeclAttr::MT_Link) {
+  if (Res && (*Res == OMPDeclareTargetDeclAttr::MT_Link ||
+              (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+               HasRequiresUnifiedSharedMemory))) {
     SmallString<64> PtrName;
     {
       llvm::raw_svector_ostream OS(PtrName);
+      if (*Res == OMPDeclareTargetDeclAttr::MT_Link)
         OS << CGM.getMangledName(GlobalDecl(VD)) << "_decl_tgt_link_ptr";
+      else
+        OS << CGM.getMangledName(GlobalDecl(VD)) << "_decl_tgt_to_ptr";
     }
     llvm::Value *Ptr = CGM.getModule().getNamedValue(PtrName);
     if (!Ptr) {
       QualType PtrTy = CGM.getContext().getPointerType(VD->getType());
       Ptr = getOrCreateInternalVariable(CGM.getTypes().ConvertTypeForMem(PtrTy),
                                         PtrName);
       if (!CGM.getLangOpts().OpenMPIsDevice) {
         auto *GV = cast<llvm::GlobalVariable>(Ptr);
         GV->setLinkage(llvm::GlobalValue::ExternalLinkage);
         GV->setInitializer(CGM.GetAddrOfGlobal(VD));
       }
       CGM.addUsedGlobal(cast<llvm::GlobalValue>(Ptr));
       registerTargetGlobalVariable(VD, cast<llvm::Constant>(Ptr));
     }
     return Address(Ptr, CGM.getContext().getDeclAlign(VD));
   }
   return Address::invalid();
 }
 
+Address
+CGOpenMPRuntime::getAddrOfDeclareTargetToUnderUnifiedMem(const VarDecl *VD) {
+  return getAddrOfDeclareTargetLink(VD);
+}
+
 llvm::Constant *
 CGOpenMPRuntime::getOrCreateThreadPrivateCache(const VarDecl *VD) {
   assert(!CGM.getLangOpts().OpenMPUseTLS ||
          !CGM.getContext().getTargetInfo().isTLSSupported());
   // Lookup the entry, lazily creating it if necessary.
   std::string Suffix = getName({"cache", ""});
   return getOrCreateInternalVariable(
       CGM.Int8PtrPtrTy, Twine(CGM.getMangledName(VD)).concat(Suffix));
static void getTargetEntryUniqueInfo(ASTContext &C, SourceLocation Loc,
   LineNum = PLoc.getLine();
 }
 
 bool CGOpenMPRuntime::emitDeclareTargetVarDefinition(const VarDecl *VD,
                                                      llvm::GlobalVariable *Addr,
                                                      bool PerformInit) {
   Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
       OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
-  if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_Link)
+  if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_Link ||
+      (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+       HasRequiresUnifiedSharedMemory))
     return CGM.getLangOpts().OpenMPIsDevice;
   VD = VD->getDefinition(CGM.getContext());
   if (VD && !DeclareTargetWithDefinition.insert(CGM.getMangledName(VD)).second)
     return CGM.getLangOpts().OpenMPIsDevice;
 
   QualType ASTTy = VD->getType();
 
   SourceLocation Loc = VD->getCanonicalDecl()->getBeginLoc();
for (const auto *E : OrderedEntries) {
     } else if (const auto *CE =
                    dyn_cast<OffloadEntriesInfoManagerTy::
                                 OffloadEntryInfoDeviceGlobalVar>(E)) {
       OffloadEntriesInfoManagerTy::OMPTargetGlobalVarEntryKind Flags =
           static_cast<OffloadEntriesInfoManagerTy::OMPTargetGlobalVarEntryKind>(
               CE->getFlags());
       switch (Flags) {
       case OffloadEntriesInfoManagerTy::OMPTargetGlobalVarEntryTo: {
+        if (CGM.getLangOpts().OpenMPIsDevice &&
+            CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory())
+          continue;
         if (!CE->getAddress()) {
           unsigned DiagID = CGM.getDiags().getCustomDiagID(
               DiagnosticsEngine::Error,
               "Offloading entry for declare target variable is incorrect: the "
               "address is invalid.");
           CGM.getDiags().Report(DiagID);
           continue;
         }
void generateInfoForComponentList(
     // ps, &(ps->p), sizeof(S2*), TARGET_PARAM
     // ps, &(ps->p), sizeof(double*), MEMBER_OF(4) (*)
     // &(ps->p), &(ps->p[0]), 33*sizeof(double), MEMBER_OF(4) | PTR_AND_OBJ | TO
     // (*) the struct this entry pertains to is the 4th element in the list
     //     of arguments, hence MEMBER_OF(4)
 
     // Track if the map information being generated is the first for a capture.
     bool IsCaptureFirstInfo = IsFirstComponentList;
     bool IsLink = false; // Is this variable a "declare target link"?
 
     // Scan the components from the base to the complete expression.
     auto CI = Components.rbegin();
     auto CE = Components.rend();
     auto I = CI;
 
     // Track if the map information being generated is the first for a list of
     // components.
if (isa<MemberExpr>(AssocExpr)) {
       BP = CGF.EmitOMPSharedLValue(AssocExpr).getAddress();
     } else {
       // The base is the reference to the variable.
       // BP = &Var.
       BP = CGF.EmitOMPSharedLValue(AssocExpr).getAddress();
       if (const auto *VD =
               dyn_cast_or_null<VarDecl>(I->getAssociatedDeclaration())) {
         if (llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
-                OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD))
+                OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD)) {
           if (*Res == OMPDeclareTargetDeclAttr::MT_Link) {
             IsLink = true;
             BP = CGF.CGM.getOpenMPRuntime().getAddrOfDeclareTargetLink(VD);
           }
+          if (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+              CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory()) {
+            // TODO: Make this into a flag for TO with unified memory.
+            IsLink = true;
+            BP = CGF.CGM.getOpenMPRuntime().getAddrOfDeclareTargetToUnderUnifiedMem(VD);
+          }
+        }
       }
 
       // If the variable is a pointer and is being dereferenced (i.e. is not
       // the last component), the base has to be the pointer itself, not its
       // reference. References are ignored for mapping purposes.
       QualType Ty =
           I->getAssociatedDeclaration()->getType().getNonReferenceType();
       if (Ty->isAnyPointerType() && std::next(I) != CE) {
if (const CXXDestructorDecl *Dtor = RD->getDestructor()) {
       scanForTargetRegionsFunctions(Dtor->getBody(), ParentName);
     }
   }
 
   // Do not to emit variable if it is not marked as declare target.
   llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
       OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(
           cast<VarDecl>(GD.getDecl()));
-  if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_Link) {
+  if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_Link ||
+      (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+       HasRequiresUnifiedSharedMemory)) {
     DeferredGlobalVariables.insert(cast<VarDecl>(GD.getDecl()));
     return true;
   }
   return false;
 }
 
 llvm::Constant *
 CGOpenMPRuntime::registerTargetFirstprivateCopy(CodeGenFunction &CGF,
void CGOpenMPRuntime::registerTargetGlobalVariable(const VarDecl *VD,
   // Register declare target variables.
   OffloadEntriesInfoManagerTy::OMPTargetGlobalVarEntryKind Flags;
   StringRef VarName;
   CharUnits VarSize;
   llvm::GlobalValue::LinkageTypes Linkage;
   switch (*Res) {
   case OMPDeclareTargetDeclAttr::MT_To:
     Flags = OffloadEntriesInfoManagerTy::OMPTargetGlobalVarEntryTo;
+    if (HasRequiresUnifiedSharedMemory) {
+      if (CGM.getLangOpts().OpenMPIsDevice) {
+        VarName = Addr->getName();
+        Addr = nullptr;
+      } else {
+        VarName = getAddrOfDeclareTargetToUnderUnifiedMem(VD).getName();
+        Addr = cast<llvm::Constant>(getAddrOfDeclareTargetToUnderUnifiedMem(VD).getPointer());
+      }
+      VarSize = CGM.getPointerSize();
+      Linkage = llvm::GlobalValue::WeakAnyLinkage;
+      break;
+    }
     VarName = CGM.getMangledName(VD);
     if (VD->hasDefinition(CGM.getContext()) != VarDecl::DeclarationOnly) {
       VarSize = CGM.getContext().getTypeSizeInChars(VD->getType());
       assert(!VarSize.isZero() && "Expected non-zero size of the variable");
     } else {
       VarSize = CharUnits::Zero();
     }
     Linkage = CGM.getLLVMLinkageVarDefinition(VD, /*IsConstant=*/false);
     // Temp solution to prevent optimizations of the internal variables.
     if (CGM.getLangOpts().OpenMPIsDevice && !VD->isExternallyVisible()) {
       std::string RefName = getName({VarName, "ref"});
       if (!CGM.GetGlobalValue(RefName)) {
         llvm::Constant *AddrRef =
             getOrCreateInternalVariable(Addr->getType(), RefName);
         auto *GVAddrRef = cast<llvm::GlobalVariable>(AddrRef);
         GVAddrRef->setConstant(/*Val=*/true);
         GVAddrRef->setLinkage(llvm::GlobalValue::InternalLinkage);
         GVAddrRef->setInitializer(Addr);
         CGM.addCompilerUsedGlobal(GVAddrRef);
       }
     }
     break;
   case OMPDeclareTargetDeclAttr::MT_Link:
     Flags = OffloadEntriesInfoManagerTy::OMPTargetGlobalVarEntryLink;
     if (CGM.getLangOpts().OpenMPIsDevice) {
       VarName = Addr->getName();
       Addr = nullptr;
     } else {
       VarName = getAddrOfDeclareTargetLink(VD).getName();
       Addr = cast<llvm::Constant>(getAddrOfDeclareTargetLink(VD).getPointer());
     }
     VarSize = CGM.getPointerSize();
     Linkage = llvm::GlobalValue::WeakAnyLinkage;

 }
 
 void CGOpenMPRuntime::emitDeferredTargetDecls() const {
   for (const VarDecl *VD : DeferredGlobalVariables) {
     llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
         OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
     if (!Res)
       continue;
-    if (*Res == OMPDeclareTargetDeclAttr::MT_To) {
+    if (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+        !HasRequiresUnifiedSharedMemory) {
       CGM.EmitGlobal(VD);
     } else {
-      assert(*Res == OMPDeclareTargetDeclAttr::MT_Link &&
+      assert((*Res == OMPDeclareTargetDeclAttr::MT_Link ||
+              *Res == OMPDeclareTargetDeclAttr::MT_To) &&
              "Expected to or link clauses.");
+      if (*Res == OMPDeclareTargetDeclAttr::MT_To)
+        (void)CGM.getOpenMPRuntime().getAddrOfDeclareTargetToUnderUnifiedMem(VD);
+      else
         (void)CGM.getOpenMPRuntime().getAddrOfDeclareTargetLink(VD);
     }
   }
 }
 
 void CGOpenMPRuntime::adjustTargetSpecificDataForLambdas(
     CodeGenFunction &CGF, const OMPExecutableDirective &D) const {
   assert(isOpenMPTargetExecutionDirective(D.getDirectiveKind()) &&
          " Expected target-based directive.");

lib/CodeGen/CodeGenModule.cpp

if (const auto *FD = dyn_cast<FunctionDecl>(Global)) {
     assert(VD->isFileVarDecl() && "Cannot emit local var decl as global.");
     if (VD->isThisDeclarationADefinition() != VarDecl::Definition &&
         !Context.isMSStaticDataMemberInlineDefinition(VD)) {
       if (LangOpts.OpenMP) {
         // Emit declaration of the must-be-emitted declare target variable.
         if (llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
                 OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD)) {
           if (*Res == OMPDeclareTargetDeclAttr::MT_To) {
+            if (getOpenMPRuntime().hasRequiresUnifiedSharedMemory())
+              (void)getOpenMPRuntime().getAddrOfDeclareTargetToUnderUnifiedMem(VD);
+            else
               (void)GetAddrOfGlobalVar(VD);
           } else {
             assert(*Res == OMPDeclareTargetDeclAttr::MT_Link &&
-                   "link claue expected.");
+                   "link clause expected.");
             (void)getOpenMPRuntime().getAddrOfDeclareTargetLink(VD);
           }
           return;
         }
       }
       // If this declaration may have caused an inline variable definition to
       // change linkage, make sure that it's emitted.
       if (Context.getInlineVariableDefinitionKind(VD) ==
           ASTContext::InlineVariableDefinitionKind::Strong)
         GetAddrOfGlobalVar(VD);
       return;

test/OpenMP/nvptx_target_requires_unified_shared_memory.cpp

 // Test declare target link under unified memory requirement.
 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK
 // expected-no-diagnostics
 
 #ifndef HEADER
 #define HEADER
 
 #define N 1000
 
 double var = 10.0;
+double to_var = 20.0;
 
 #pragma omp requires unified_shared_memory
 #pragma omp declare target link(var)
+#pragma omp declare target to(to_var)
 
 int bar(int n){
   double sum = 0;
 
 #pragma omp target
   for(int i = 0; i < n; i++) {
-    sum += var;
+    sum += var + to_var;
   }
 
   return sum;
 }
 
 // CHECK: [[VAR:@.+]] = global double 1.000000e+01
 // CHECK: [[VAR_DECL_TGT_LINK_PTR:@.+]] = global double* [[VAR]]
 
+// CHECK: [[TO_VAR:@.+]] = global double 2.000000e+01
+// CHECK: [[VAR_DECL_TGT_TO_PTR:@.+]] = global double* [[TO_VAR]]
+
 // CHECK: [[OFFLOAD_SIZES:@.+]] = private unnamed_addr constant [2 x i64] [i64 4, i64 8]
 // CHECK: [[OFFLOAD_MAPTYPES:@.+]] = private unnamed_addr constant [2 x i64] [i64 800, i64 800]
 
+// CHECK: [[OMP_OFFLOAD_ENTRY_LINK_VAR_PTR_NAME:@.+]] = internal unnamed_addr constant [22 x i8]
+// CHECK: [[OMP_OFFLOAD_ENTRY_LINK_VAR_PTR:@.+]] = weak constant %struct.__tgt_offload_entry { i8* bitcast (double** [[VAR_DECL_TGT_LINK_PTR]] to i8*), i8* getelementptr inbounds ([22 x i8], [22 x i8]* [[OMP_OFFLOAD_ENTRY_LINK_VAR_PTR_NAME]], i32 0, i32 0), i64 8, i32 1, i32 0 }, section ".omp_offloading.entries"
+
+// CHECK: [[OMP_OFFLOAD_ENTRY_TO_VAR_PTR_NAME:@.+]] = internal unnamed_addr constant [23 x i8]
+// CHECK: [[OMP_OFFLOAD_ENTRY_TO_VAR_PTR:@.+]] = weak constant %struct.__tgt_offload_entry { i8* bitcast (double** [[VAR_DECL_TGT_TO_PTR]] to i8*), i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[OMP_OFFLOAD_ENTRY_TO_VAR_PTR_NAME]], i32 0, i32 0), i64 8, i32 0, i32 0 }, section ".omp_offloading.entries"
+
+// CHECK: @llvm.used = appending global [2 x i8*] [i8* bitcast (double** [[VAR_DECL_TGT_LINK_PTR]] to i8*), i8* bitcast (double** [[VAR_DECL_TGT_TO_PTR]] to i8*)], section "llvm.metadata"
+
 // CHECK: [[N_CASTED:%.+]] = alloca i64
 // CHECK: [[SUM_CASTED:%.+]] = alloca i64
 
 // CHECK: [[OFFLOAD_BASEPTRS:%.+]] = alloca [2 x i8*]
 // CHECK: [[OFFLOAD_PTRS:%.+]] = alloca [2 x i8*]
 
 // CHECK: [[LOAD1:%.+]] = load i64, i64* [[N_CASTED]]
 // CHECK: [[LOAD2:%.+]] = load i64, i64* [[SUM_CASTED]]