[OpenMP] Migrate device code privatisation from Clang CodeGen to OMPIRBuilder

[TIFitis]

Summary

This patch migrates the UseDevicePtr and UseDeviceAddr clause related code for handling privatisation from Clang codegen to the OMPIRBuilder

Depends on D150860

[TIFitis]

clang/lib/CodeGen/CGStmtOpenMP.cpp

7170–7234

Currently clang maintains a map between Clang::ValueDecl and corresponding address in Info.CaptureDeviceAddrMap.

Please see the EmitOMPUseDeviceAddrClause function just below this to see how it does so.

I would however like to get rid of this map as it allows us to then get rid of another map in the MapInfosTy and a clang specific callback for maintaining these maps..

This map is only used to retain the mapping between a ValueDecl inside a map clause to its corresponding llvm address.

The highlighted code here tries to retrieve this llvm address back from the ValueDecl without using the map, but as you can it is a very hacky way of doing so. But my intuition says that there must be a much simpler method of retrieving this address as Clang must be storing it in some way in order for it to generate code.

It would be great if someone with more experience in Clang codegen could help me achieve this.

Also I am aware that I am probably doing a very poor job at explaining this. Please let me know if you have any questions or would like me to try and explain it with more clarity.

[TIFitis]

Rebased

[TIFitis]

Rebased

[TIFitis]

Fixed typo

[jdoerfert]

typo in the title.

clang/lib/CodeGen/CGOpenMPRuntime.h

1444

If it is an alloca (for sure) use llvm::AllocaInst.

clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp

418

Did we just change the order or is this non-deterministic? The latter is not acceptable, the former might be.

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

4133

What else could it be here?

4698

What if the PointerBCOrASCast and the GEP are both no-ops, and BPVal is an alloca. Would that cause a problem as we later only check if the type of the value is an alloca. That said, it seems dangerous to only rely on the type of the value.

[TIFitis]

Addressed reviewer comments.

clang/lib/CodeGen/CGOpenMPRuntime.h

1444

It's alloca for use_dev_ptr and usually GEPInst for use_dev_addr.
Updated comment to llvm address

clang/lib/CodeGen/CGStmtOpenMP.cpp

7170–7234

Please ignore the above comment. I've kept things simpler for this patch.

I'll try to have another patch in the future to focus on removing the clang specific maps and callbacks.

clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp

418

It's a change of order. I've used a llvm::SmallMapVector for determinism.

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

4133

For use_dev_ptr a new alloca is inserted.
For use_dev_addr it points to the GEP instruction from the mappers.

4698

I don't think the GEP can be a no-op here as it's from the mapper array that we are adding above.

That being said, at worst we would be emitting a spurious load and store instruction with no use.

[TIFitis]

Rebased

[TIFitis]

Ping for review :)

[jdoerfert]

LG

clang/lib/CodeGen/CGOpenMPRuntime.h

public:
   /// throughout a 'target data' region.
   class TargetDataInfo : public llvm::OpenMPIRBuilder::TargetDataInfo {
   public:
     explicit TargetDataInfo() : llvm::OpenMPIRBuilder::TargetDataInfo() {}
     explicit TargetDataInfo(bool RequiresDevicePointerInfo,
                             bool SeparateBeginEndCalls)
         : llvm::OpenMPIRBuilder::TargetDataInfo(RequiresDevicePointerInfo,
                                                 SeparateBeginEndCalls) {}
-    /// Map between the a declaration of a capture and the corresponding base
-    /// pointer address where the runtime returns the device pointers.
-    llvm::DenseMap<const ValueDecl *, Address> CaptureDeviceAddrMap;
+    /// Map between the a declaration of a capture and the corresponding new
+    /// alloca address where the runtime returns the device pointers.
+    llvm::DenseMap<const ValueDecl *, llvm::Value *> CaptureDeviceAddrMap;
   };
 
   /// Emit the target data mapping code associated with \a D.
   /// \param D Directive to emit.
   /// \param IfCond Expression evaluated in if clause associated with the
   /// target directive, or null if no device clause is used.
   /// \param Device Expression evaluated in device clause associated with the
   /// target directive, or null if no device clause is used.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

public:
   public:
     MappingExprInfo(const ValueDecl *MapDecl, const Expr *MapExpr = nullptr)
         : MapDecl(MapDecl), MapExpr(MapExpr) {}
 
     const ValueDecl *getMapDecl() const { return MapDecl; }
     const Expr *getMapExpr() const { return MapExpr; }
   };
 
+  using DeviceInfoTy = llvm::OpenMPIRBuilder::DeviceInfoTy;
   using MapBaseValuesArrayTy = llvm::OpenMPIRBuilder::MapValuesArrayTy;
   using MapValuesArrayTy = llvm::OpenMPIRBuilder::MapValuesArrayTy;
   using MapFlagsArrayTy = llvm::OpenMPIRBuilder::MapFlagsArrayTy;
   using MapDimArrayTy = llvm::OpenMPIRBuilder::MapDimArrayTy;
   using MapNonContiguousArrayTy =
       llvm::OpenMPIRBuilder::MapNonContiguousArrayTy;
   using MapExprsArrayTy = SmallVector<MappingExprInfo, 4>;
   using MapValueDeclsArrayTy = SmallVector<const ValueDecl *, 4>;
for (; I != CE; ++I) {
                     CGF.EmitCastToVoidPtr(LB.getPointer()));
                 break;
               }
             }
             assert(Size && "Failed to determine structure size");
             CombinedInfo.Exprs.emplace_back(MapDecl, MapExpr);
             CombinedInfo.BasePointers.push_back(BP.getPointer());
             CombinedInfo.DevicePtrDecls.push_back(nullptr);
+            CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None);
             CombinedInfo.Pointers.push_back(LB.getPointer());
             CombinedInfo.Sizes.push_back(CGF.Builder.CreateIntCast(
                 Size, CGF.Int64Ty, /*isSigned=*/true));
             CombinedInfo.Types.push_back(Flags);
             CombinedInfo.Mappers.push_back(nullptr);
             CombinedInfo.NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize
                                                                       : 1);
             LB = CGF.Builder.CreateConstGEP(ComponentLB, 1);
           }
           CombinedInfo.Exprs.emplace_back(MapDecl, MapExpr);
           CombinedInfo.BasePointers.push_back(BP.getPointer());
           CombinedInfo.DevicePtrDecls.push_back(nullptr);
+          CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None);
           CombinedInfo.Pointers.push_back(LB.getPointer());
           Size = CGF.Builder.CreatePtrDiff(
               CGF.Int8Ty, CGF.Builder.CreateConstGEP(HB, 1).getPointer(),
               CGF.EmitCastToVoidPtr(LB.getPointer()));
           CombinedInfo.Sizes.push_back(
               CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true));
           CombinedInfo.Types.push_back(Flags);
           CombinedInfo.Mappers.push_back(nullptr);
           CombinedInfo.NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize
                                                                     : 1);
           break;
         }
         llvm::Value *Size = getExprTypeSize(I->getAssociatedExpression());
         if (!IsMemberPointerOrAddr ||
             (Next == CE && MapType != OMPC_MAP_unknown)) {
           CombinedInfo.Exprs.emplace_back(MapDecl, MapExpr);
           CombinedInfo.BasePointers.push_back(BP.getPointer());
           CombinedInfo.DevicePtrDecls.push_back(nullptr);
+          CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None);
           CombinedInfo.Pointers.push_back(LB.getPointer());
           CombinedInfo.Sizes.push_back(
               CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true));
           CombinedInfo.NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize
                                                                     : 1);
 
           // If Mapper is valid, the last component inherits the mapper.
           bool HasMapper = Mapper && Next == CE;
void generateAllInfoForClauses(
     // struct has been processed.
     llvm::MapVector<CanonicalDeclPtr<const Decl>,
                     SmallVector<DeferredDevicePtrEntryTy, 4>>
         DeferredInfo;
     MapCombinedInfoTy UseDeviceDataCombinedInfo;
 
     auto &&UseDeviceDataCombinedInfoGen =
         [&UseDeviceDataCombinedInfo](const ValueDecl *VD, llvm::Value *Ptr,
-                                     CodeGenFunction &CGF) {
+                                     CodeGenFunction &CGF, bool IsDevAddr) {
           UseDeviceDataCombinedInfo.Exprs.push_back(VD);
           UseDeviceDataCombinedInfo.BasePointers.emplace_back(Ptr);
           UseDeviceDataCombinedInfo.DevicePtrDecls.emplace_back(VD);
+          UseDeviceDataCombinedInfo.DevicePointers.emplace_back(
+              IsDevAddr ? DeviceInfoTy::Address : DeviceInfoTy::Pointer);
           UseDeviceDataCombinedInfo.Pointers.push_back(Ptr);
           UseDeviceDataCombinedInfo.Sizes.push_back(
               llvm::Constant::getNullValue(CGF.Int64Ty));
           UseDeviceDataCombinedInfo.Types.push_back(
               OpenMPOffloadMappingFlags::OMP_MAP_RETURN_PARAM);
           UseDeviceDataCombinedInfo.Mappers.push_back(nullptr);
         };
 
auto &&MapInfoGen =
             if (IsDevAddr) {
               if (IE->isGLValue())
                 Ptr = CGF.EmitLValue(IE).getPointer(CGF);
               else
                 Ptr = CGF.EmitScalarExpr(IE);
             } else {
               Ptr = CGF.EmitLoadOfScalar(CGF.EmitLValue(IE), IE->getExprLoc());
             }
-            UseDeviceDataCombinedInfoGen(VD, Ptr, CGF);
+            UseDeviceDataCombinedInfoGen(VD, Ptr, CGF, IsDevAddr);
           }
         };
 
     auto &&IsMapInfoExist = [&Info](CodeGenFunction &CGF, const ValueDecl *VD,
                                     const Expr *IE, bool IsDevAddr) -> bool {
       // We potentially have map information for this declaration already.
       // Look for the first set of components that refer to it. If found,
       // return true.
auto &&IsMapInfoExist = [&Info](CodeGenFunction &CGF, const ValueDecl *VD,
           // If we found a map entry, signal that the pointer has to be
           // returned and move on to the next declaration. Exclude cases where
           // the base pointer is mapped as array subscript, array section or
           // array shaping. The base address is passed as a pointer to base in
           // this case and cannot be used as a base for use_device_ptr list
           // item.
           if (CI != Data.end()) {
             if (IsDevAddr) {
+              CI->ForDeviceAddr = IsDevAddr;
               CI->ReturnDevicePointer = true;
               Found = true;
               break;
             } else {
               auto PrevCI = std::next(CI->Components.rbegin());
               const auto *VarD = dyn_cast<VarDecl>(VD);
               if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() ||
                   isa<MemberExpr>(IE) ||
                   !VD->getType().getNonReferenceType()->isPointerType() ||
                   PrevCI == CI->Components.rend() ||
                   isa<MemberExpr>(PrevCI->getAssociatedExpression()) || !VarD ||
                   VarD->hasLocalStorage()) {
+                CI->ForDeviceAddr = IsDevAddr;
                 CI->ReturnDevicePointer = true;
                 Found = true;
                 break;
               }
             }
           }
         }
         return Found;
for (const auto &Data : Info) {
                    "Unexpected number of mapped base pointers.");
 
             const ValueDecl *RelevantVD =
                 L.Components.back().getAssociatedDeclaration();
             assert(RelevantVD &&
                    "No relevant declaration related with device pointer??");
 
             CurInfo.DevicePtrDecls[CurrentBasePointersIdx] = RelevantVD;
+            CurInfo.DevicePointers[CurrentBasePointersIdx] =
+                L.ForDeviceAddr ? DeviceInfoTy::Address : DeviceInfoTy::Pointer;
             CurInfo.Types[CurrentBasePointersIdx] |=
                 OpenMPOffloadMappingFlags::OMP_MAP_RETURN_PARAM;
           }
         }
       }
 
       // Append any pending zero-length pointers which are struct members and
       // used with use_device_ptr or use_device_addr.
for (const auto &Data : Info) {
             CurInfo.Types.push_back(
                 OpenMPOffloadMappingFlags::OMP_MAP_PTR_AND_OBJ |
                 OpenMPOffloadMappingFlags::OMP_MAP_RETURN_PARAM |
                 OpenMPOffloadMappingFlags::OMP_MAP_MEMBER_OF);
           }
           CurInfo.Exprs.push_back(L.VD);
           CurInfo.BasePointers.emplace_back(BasePtr);
           CurInfo.DevicePtrDecls.emplace_back(L.VD);
+          CurInfo.DevicePointers.emplace_back(
+              L.ForDeviceAddr ? DeviceInfoTy::Address : DeviceInfoTy::Pointer);
           CurInfo.Pointers.push_back(Ptr);
           CurInfo.Sizes.push_back(
               llvm::Constant::getNullValue(this->CGF.Int64Ty));
           CurInfo.Mappers.push_back(nullptr);
         }
       }
       // If there is an entry in PartialStruct it means we have a struct with
       // individual members mapped. Emit an extra combined entry.
void emitCombinedEntry(MapCombinedInfoTy &CombinedInfo,
     if (PartialStruct.HasCompleteRecord) {
       LBAddr = PartialStruct.LB;
       HBAddr = PartialStruct.LB;
     }
     CombinedInfo.Exprs.push_back(VD);
     // Base is the base of the struct
     CombinedInfo.BasePointers.push_back(PartialStruct.Base.getPointer());
     CombinedInfo.DevicePtrDecls.push_back(nullptr);
+    CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None);
     // Pointer is the address of the lowest element
     llvm::Value *LB = LBAddr.getPointer();
     const CXXMethodDecl *MD =
         CGF.CurFuncDecl ? dyn_cast<CXXMethodDecl>(CGF.CurFuncDecl) : nullptr;
     const CXXRecordDecl *RD = MD ? MD->getParent() : nullptr;
     bool HasBaseClass = RD && IsMapThis ? RD->getNumBases() > 0 : false;
     // There should not be a mapper for a combined entry.
     if (HasBaseClass) {
if (ThisCapture) {
       LValue ThisLVal =
           CGF.EmitLValueForFieldInitialization(VDLVal, ThisCapture);
       LValue ThisLValVal = CGF.EmitLValueForField(VDLVal, ThisCapture);
       LambdaPointers.try_emplace(ThisLVal.getPointer(CGF),
                                  VDLVal.getPointer(CGF));
       CombinedInfo.Exprs.push_back(VD);
       CombinedInfo.BasePointers.push_back(ThisLVal.getPointer(CGF));
       CombinedInfo.DevicePtrDecls.push_back(nullptr);
+      CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None);
       CombinedInfo.Pointers.push_back(ThisLValVal.getPointer(CGF));
       CombinedInfo.Sizes.push_back(
           CGF.Builder.CreateIntCast(CGF.getTypeSize(CGF.getContext().VoidPtrTy),
                                     CGF.Int64Ty, /*isSigned=*/true));
       CombinedInfo.Types.push_back(
           OpenMPOffloadMappingFlags::OMP_MAP_PTR_AND_OBJ |
           OpenMPOffloadMappingFlags::OMP_MAP_LITERAL |
           OpenMPOffloadMappingFlags::OMP_MAP_MEMBER_OF |
for (const LambdaCapture &LC : RD->captures()) {
       LValue VarLVal = CGF.EmitLValueForFieldInitialization(VDLVal, It->second);
       if (LC.getCaptureKind() == LCK_ByRef) {
         LValue VarLValVal = CGF.EmitLValueForField(VDLVal, It->second);
         LambdaPointers.try_emplace(VarLVal.getPointer(CGF),
                                    VDLVal.getPointer(CGF));
         CombinedInfo.Exprs.push_back(VD);
         CombinedInfo.BasePointers.push_back(VarLVal.getPointer(CGF));
         CombinedInfo.DevicePtrDecls.push_back(nullptr);
+        CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None);
         CombinedInfo.Pointers.push_back(VarLValVal.getPointer(CGF));
         CombinedInfo.Sizes.push_back(CGF.Builder.CreateIntCast(
             CGF.getTypeSize(
                 VD->getType().getCanonicalType().getNonReferenceType()),
             CGF.Int64Ty, /*isSigned=*/true));
       } else {
         RValue VarRVal = CGF.EmitLoadOfLValue(VarLVal, RD->getLocation());
         LambdaPointers.try_emplace(VarLVal.getPointer(CGF),
                                    VDLVal.getPointer(CGF));
         CombinedInfo.Exprs.push_back(VD);
         CombinedInfo.BasePointers.push_back(VarLVal.getPointer(CGF));
         CombinedInfo.DevicePtrDecls.push_back(nullptr);
+        CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None);
         CombinedInfo.Pointers.push_back(VarRVal.getScalarVal());
         CombinedInfo.Sizes.push_back(llvm::ConstantInt::get(CGF.Int64Ty, 0));
       }
       CombinedInfo.Types.push_back(
           OpenMPOffloadMappingFlags::OMP_MAP_PTR_AND_OBJ |
           OpenMPOffloadMappingFlags::OMP_MAP_LITERAL |
           OpenMPOffloadMappingFlags::OMP_MAP_MEMBER_OF |
           OpenMPOffloadMappingFlags::OMP_MAP_IMPLICIT);
void generateInfoForCapture(const CapturedStmt::Capture *Cap,
 
     // If this declaration appears in a is_device_ptr clause we just have to
     // pass the pointer by value. If it is a reference to a declaration, we just
     // pass its value.
     if (VD && (DevPointersMap.count(VD) || HasDevAddrsMap.count(VD))) {
       CombinedInfo.Exprs.push_back(VD);
       CombinedInfo.BasePointers.emplace_back(Arg);
       CombinedInfo.DevicePtrDecls.emplace_back(VD);
+      CombinedInfo.DevicePointers.emplace_back(DeviceInfoTy::Pointer);
       CombinedInfo.Pointers.push_back(Arg);
       CombinedInfo.Sizes.push_back(CGF.Builder.CreateIntCast(
           CGF.getTypeSize(CGF.getContext().VoidPtrTy), CGF.Int64Ty,
           /*isSigned=*/true));
       CombinedInfo.Types.push_back(
           OpenMPOffloadMappingFlags::OMP_MAP_LITERAL |
           OpenMPOffloadMappingFlags::OMP_MAP_TARGET_PARAM);
       CombinedInfo.Mappers.push_back(nullptr);
void generateDefaultMapInfo(const CapturedStmt::Capture &CI,
                               const FieldDecl &RI, llvm::Value *CV,
                               MapCombinedInfoTy &CombinedInfo) const {
     bool IsImplicit = true;
     // Do the default mapping.
     if (CI.capturesThis()) {
       CombinedInfo.Exprs.push_back(nullptr);
       CombinedInfo.BasePointers.push_back(CV);
       CombinedInfo.DevicePtrDecls.push_back(nullptr);
+      CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None);
       CombinedInfo.Pointers.push_back(CV);
       const auto *PtrTy = cast<PointerType>(RI.getType().getTypePtr());
       CombinedInfo.Sizes.push_back(
           CGF.Builder.CreateIntCast(CGF.getTypeSize(PtrTy->getPointeeType()),
                                     CGF.Int64Ty, /*isSigned=*/true));
       // Default map type.
       CombinedInfo.Types.push_back(OpenMPOffloadMappingFlags::OMP_MAP_TO |
                                    OpenMPOffloadMappingFlags::OMP_MAP_FROM);
     } else if (CI.capturesVariableByCopy()) {
       const VarDecl *VD = CI.getCapturedVar();
       CombinedInfo.Exprs.push_back(VD->getCanonicalDecl());
       CombinedInfo.BasePointers.push_back(CV);
       CombinedInfo.DevicePtrDecls.push_back(nullptr);
+      CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None);
       CombinedInfo.Pointers.push_back(CV);
       if (!RI.getType()->isAnyPointerType()) {
         // We have to signal to the runtime captures passed by value that are
         // not pointers.
         CombinedInfo.Types.push_back(
             OpenMPOffloadMappingFlags::OMP_MAP_LITERAL);
         CombinedInfo.Sizes.push_back(CGF.Builder.CreateIntCast(
             CGF.getTypeSize(RI.getType()), CGF.Int64Ty, /*isSigned=*/true));
if (CI.capturesThis()) {
       // default the value doesn't have to be retrieved. For an aggregate
       // type, the default is 'tofrom'.
       CombinedInfo.Types.push_back(getMapModifiersForPrivateClauses(CI));
       const VarDecl *VD = CI.getCapturedVar();
       auto I = FirstPrivateDecls.find(VD);
       CombinedInfo.Exprs.push_back(VD->getCanonicalDecl());
       CombinedInfo.BasePointers.push_back(CV);
       CombinedInfo.DevicePtrDecls.push_back(nullptr);
+      CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None);
       if (I != FirstPrivateDecls.end() && ElementType->isAnyPointerType()) {
         Address PtrAddr = CGF.EmitLoadOfReference(CGF.MakeAddrLValue(
             CV, ElementType, CGF.getContext().getDeclAlign(VD),
             AlignmentSource::Decl));
         CombinedInfo.Pointers.push_back(PtrAddr.getPointer());
       } else {
         CombinedInfo.Pointers.push_back(CV);
       }

 /// Emit the arrays used to pass the captures and map information to the
 /// offloading runtime library. If there is no map or capture information,
 /// return nullptr by reference.
 static void emitOffloadingArrays(
     CodeGenFunction &CGF, MappableExprsHandler::MapCombinedInfoTy &CombinedInfo,
     CGOpenMPRuntime::TargetDataInfo &Info, llvm::OpenMPIRBuilder &OMPBuilder,
     bool IsNonContiguous = false) {
   CodeGenModule &CGM = CGF.CGM;
-  ASTContext &Ctx = CGF.getContext();
 
   // Reset the array information.
   Info.clearArrayInfo();
   Info.NumberOfPtrs = CombinedInfo.BasePointers.size();
 
   using InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy;
   InsertPointTy AllocaIP(CGF.AllocaInsertPt->getParent(),
                          CGF.AllocaInsertPt->getIterator());
   InsertPointTy CodeGenIP(CGF.Builder.GetInsertBlock(),
                           CGF.Builder.GetInsertPoint());
 
   auto fillInfoMap = [&](MappableExprsHandler::MappingExprInfo &MapExpr) {
     return emitMappingInformation(CGF, OMPBuilder, MapExpr);
   };
   if (CGM.getCodeGenOpts().getDebugInfo() !=
       llvm::codegenoptions::NoDebugInfo) {
     CombinedInfo.Names.resize(CombinedInfo.Exprs.size());
     llvm::transform(CombinedInfo.Exprs, CombinedInfo.Names.begin(),
                     fillInfoMap);
   }
 
-  auto DeviceAddrCB = [&](unsigned int I, llvm::Value *BP, llvm::Value *BPVal) {
+  auto DeviceAddrCB = [&](unsigned int I, llvm::Value *NewDecl) {
     if (const ValueDecl *DevVD = CombinedInfo.DevicePtrDecls[I]) {
-      Address BPAddr(BP, BPVal->getType(),
-                     Ctx.getTypeAlignInChars(Ctx.VoidPtrTy));
-      Info.CaptureDeviceAddrMap.try_emplace(DevVD, BPAddr);
+      Info.CaptureDeviceAddrMap.try_emplace(DevVD, NewDecl);
     }
   };
 
   auto CustomMapperCB = [&](unsigned int I) {
     llvm::Value *MFunc = nullptr;
     if (CombinedInfo.Mappers[I]) {
       Info.HasMapper = true;
       MFunc = CGF.CGM.getOpenMPRuntime().getOrCreateUserDefinedMapperFunc(
for (CapturedStmt::const_capture_iterator CI = CS.capture_begin(),
       MappableExprsHandler::StructRangeInfoTy PartialStruct;
 
       // VLA sizes are passed to the outlined region by copy and do not have map
       // information associated.
       if (CI->capturesVariableArrayType()) {
         CurInfo.Exprs.push_back(nullptr);
         CurInfo.BasePointers.push_back(*CV);
         CurInfo.DevicePtrDecls.push_back(nullptr);
+        CurInfo.DevicePointers.push_back(
+            MappableExprsHandler::DeviceInfoTy::None);
         CurInfo.Pointers.push_back(*CV);
         CurInfo.Sizes.push_back(CGF.Builder.CreateIntCast(
             CGF.getTypeSize(RI->getType()), CGF.Int64Ty, /*isSigned=*/true));
         // Copy to the device as an argument. No need to retrieve it.
         CurInfo.Types.push_back(
             OpenMPOffloadMappingFlags::OMP_MAP_LITERAL |
             OpenMPOffloadMappingFlags::OMP_MAP_TARGET_PARAM |
             OpenMPOffloadMappingFlags::OMP_MAP_IMPLICIT);
case 3:
         CodeGen(CGF);
       }
       break;
     }
     return InsertPointTy(CGF.Builder.GetInsertBlock(),
                          CGF.Builder.GetInsertPoint());
   };
 
-  auto DeviceAddrCB = [&](unsigned int I, llvm::Value *BP, llvm::Value *BPVal) {
+  auto DeviceAddrCB = [&](unsigned int I, llvm::Value *NewDecl) {
     if (const ValueDecl *DevVD = CombinedInfo.DevicePtrDecls[I]) {
-      ASTContext &Ctx = CGF.getContext();
-      Address BPAddr(BP, BPVal->getType(),
-                     Ctx.getTypeAlignInChars(Ctx.VoidPtrTy));
-      Info.CaptureDeviceAddrMap.try_emplace(DevVD, BPAddr);
+      Info.CaptureDeviceAddrMap.try_emplace(DevVD, NewDecl);
     }
   };
 
   auto CustomMapperCB = [&](unsigned int I) {
     llvm::Value *MFunc = nullptr;
     if (CombinedInfo.Mappers[I]) {
       Info.HasMapper = true;
       MFunc = CGF.CGM.getOpenMPRuntime().getOrCreateUserDefinedMapperFunc(

clang/lib/CodeGen/CGStmtOpenMP.cpp

assert(Kind == OMPD_for || Kind == OMPD_section || Kind == OMPD_sections ||
          Kind == OMPD_target_parallel_for ||
          Kind == OMPD_teams_distribute_parallel_for ||
          Kind == OMPD_target_teams_distribute_parallel_for);
   return OMPCancelStack.getExitBlock();
 }
 
 void CodeGenFunction::EmitOMPUseDevicePtrClause(
     const OMPUseDevicePtrClause &C, OMPPrivateScope &PrivateScope,
-    const llvm::DenseMap<const ValueDecl *, Address> &CaptureDeviceAddrMap) {
-  auto OrigVarIt = C.varlist_begin();
-  auto InitIt = C.inits().begin();
-  for (const Expr *PvtVarIt : C.private_copies()) {
-    const auto *OrigVD =
-        cast<VarDecl>(cast<DeclRefExpr>(*OrigVarIt)->getDecl());
-    const auto *InitVD = cast<VarDecl>(cast<DeclRefExpr>(*InitIt)->getDecl());
-    const auto *PvtVD = cast<VarDecl>(cast<DeclRefExpr>(PvtVarIt)->getDecl());
+    const llvm::SmallMapVector<const llvm::Value *,
+                               std::pair<llvm::Value *, llvm::Value *>, 4>
+        &DevicePtrInfoMap) {
+  llvm::SmallDenseSet<CanonicalDeclPtr<const Decl>, 4> Processed;
+  for (const Expr *Ref : C.varlists()) {
+    const auto *OrigVD = cast<VarDecl>(cast<DeclRefExpr>(Ref)->getDecl());
+    if (!Processed.insert(OrigVD).second)
+      continue;
+
+    const llvm::Value *DevPtr = nullptr;
 
     // In order to identify the right initializer we need to match the
     // declaration used by the mapping logic. In some cases we may get
     // OMPCapturedExprDecl that refers to the original declaration.
-    const ValueDecl *MatchingVD = OrigVD;
-    if (const auto *OED = dyn_cast<OMPCapturedExprDecl>(MatchingVD)) {
+    if (const auto *OED = dyn_cast<OMPCapturedExprDecl>(OrigVD)) {
       // OMPCapturedExprDecl are used to privative fields of the current
       // structure.
       const auto *ME = cast<MemberExpr>(OED->getInit());
       assert(isa<CXXThisExpr>(ME->getBase()->IgnoreImpCasts()) &&
              "Base should be the current struct!");
-      MatchingVD = ME->getMemberDecl();
+      auto *MI = cast<llvm::Instruction>(EmitLValue(ME).getPointer(*this));
+      while (auto *LI = dyn_cast<llvm::LoadInst>(MI)) {
+        auto *PO = LI->getPointerOperand();
+        if (!isa<llvm::AllocaInst>(PO) && isa<llvm::Instruction>(PO))
+          MI = cast<llvm::Instruction>(PO);
+        else
+          break;
+      }
+      for (const auto &DevMap : DevicePtrInfoMap) {
+        auto *DI = cast<llvm::Instruction>(DevMap.first);
+        while (auto *LI = dyn_cast<llvm::LoadInst>(DI)) {
+          auto *PO = LI->getPointerOperand();
+          if (!isa<llvm::AllocaInst>(PO) && isa<llvm::Instruction>(PO))
+            DI = cast<llvm::Instruction>(PO);
+          else
+            break;
+        }
+        if (DI->isIdenticalTo(MI)) {
+          DevPtr = DevMap.first;
+          break;
+        }
+      }
+      if (MI->users().empty())
+        MI->eraseFromParent();
+    } else {
+      if (Ref->isGLValue()) {
+        auto *MI = cast<llvm::Instruction>(EmitScalarExpr(Ref));
+        while (auto *LI = dyn_cast<llvm::LoadInst>(MI)) {
+          auto *PO = LI->getPointerOperand();
+          if (!isa<llvm::AllocaInst>(PO) && isa<llvm::Instruction>(PO))
+            MI = cast<llvm::Instruction>(PO);
+          else
+            break;
+        }
+        for (const auto &DevMap : DevicePtrInfoMap) {
+          auto *DI = cast<llvm::Instruction>(DevMap.first);
+          while (auto *LI = dyn_cast<llvm::LoadInst>(DI)) {
+            auto *PO = LI->getPointerOperand();
+            if (!isa<llvm::AllocaInst>(PO) && isa<llvm::Instruction>(PO))
+              DI = cast<llvm::Instruction>(PO);
+            else
+              break;
+          }
+          if (DI->isIdenticalTo(MI)) {
+            DevPtr = DevMap.first;
+            break;
+          }
+        }
+        if (MI->users().empty())
+          MI->eraseFromParent();
+      } else {
+        DevPtr = EmitLValue(Ref).getPointer(*this);
+      }
     }
 
     // If we don't have information about the current list item, move on to
     // the next one.
-    auto InitAddrIt = CaptureDeviceAddrMap.find(MatchingVD);
-    if (InitAddrIt == CaptureDeviceAddrMap.end())
+    if (!DevPtr)
+      continue;
+    auto InitAddrIt = DevicePtrInfoMap.find(DevPtr);
+    if (InitAddrIt == DevicePtrInfoMap.end())
       continue;
 
-    // Initialize the temporary initialization variable with the address
-    // we get from the runtime library. We have to cast the source address
-    // because it is always a void *. References are materialized in the
-    // privatization scope, so the initialization here disregards the fact
-    // the original variable is a reference.
     llvm::Type *Ty = ConvertTypeForMem(OrigVD->getType().getNonReferenceType());
-    Address InitAddr = Builder.CreateElementBitCast(InitAddrIt->second, Ty);
-    setAddrOfLocalVar(InitVD, InitAddr);
-
-    // Emit private declaration, it will be initialized by the value we
-    // declaration we just added to the local declarations map.
-    EmitDecl(*PvtVD);
-
-    // The initialization variables reached its purpose in the emission
-    // of the previous declaration, so we don't need it anymore.
-    LocalDeclMap.erase(InitVD);
 
     // Return the address of the private variable.
-    bool IsRegistered =
-        PrivateScope.addPrivate(OrigVD, GetAddrOfLocalVar(PvtVD));
+    bool IsRegistered = PrivateScope.addPrivate(
+        OrigVD,
+        Address(InitAddrIt->second.second, Ty,
+                getContext().getTypeAlignInChars(getContext().VoidPtrTy)));
     assert(IsRegistered && "firstprivate var already registered as private");
     // Silence the warning about unused variable.
     (void)IsRegistered;
-
-    ++OrigVarIt;
-    ++InitIt;
   }
 }
 
 static const VarDecl *getBaseDecl(const Expr *Ref) {
   const Expr *Base = Ref->IgnoreParenImpCasts();
   while (const auto *OASE = dyn_cast<OMPArraySectionExpr>(Base))
     Base = OASE->getBase()->IgnoreParenImpCasts();
   while (const auto *ASE = dyn_cast<ArraySubscriptExpr>(Base))
     Base = ASE->getBase()->IgnoreParenImpCasts();
   return cast<VarDecl>(cast<DeclRefExpr>(Base)->getDecl());
 }
 
 void CodeGenFunction::EmitOMPUseDeviceAddrClause(
     const OMPUseDeviceAddrClause &C, OMPPrivateScope &PrivateScope,
-    const llvm::DenseMap<const ValueDecl *, Address> &CaptureDeviceAddrMap) {
+    const llvm::DenseMap<const ValueDecl *, llvm::Value *>
+        CaptureDeviceAddrMap) {
   llvm::SmallDenseSet<CanonicalDeclPtr<const Decl>, 4> Processed;
   for (const Expr *Ref : C.varlists()) {
     const VarDecl *OrigVD = getBaseDecl(Ref);
     if (!Processed.insert(OrigVD).second)
       continue;
     // In order to identify the right initializer we need to match the
     // declaration used by the mapping logic. In some cases we may get
     // OMPCapturedExprDecl that refers to the original declaration.
     const ValueDecl *MatchingVD = OrigVD;
     if (const auto *OED = dyn_cast<OMPCapturedExprDecl>(MatchingVD)) {
       // OMPCapturedExprDecl are used to privative fields of the current
       // structure.
       const auto *ME = cast<MemberExpr>(OED->getInit());
       assert(isa<CXXThisExpr>(ME->getBase()) &&
              "Base should be the current struct!");
       MatchingVD = ME->getMemberDecl();
     }
 
     // If we don't have information about the current list item, move on to
     // the next one.
     auto InitAddrIt = CaptureDeviceAddrMap.find(MatchingVD);
     if (InitAddrIt == CaptureDeviceAddrMap.end())
       continue;
 
-    Address PrivAddr = InitAddrIt->getSecond();
+    llvm::Type *Ty = ConvertTypeForMem(OrigVD->getType().getNonReferenceType());
+
+    Address PrivAddr =
+        Address(InitAddrIt->second, Ty,
+                getContext().getTypeAlignInChars(getContext().VoidPtrTy));
     // For declrefs and variable length array need to load the pointer for
     // correct mapping, since the pointer to the data was passed to the runtime.
     if (isa<DeclRefExpr>(Ref->IgnoreParenImpCasts()) ||
         MatchingVD->getType()->isArrayType()) {
       QualType PtrTy = getContext().getPointerType(
           OrigVD->getType().getNonReferenceType());
       PrivAddr = EmitLoadOfPointer(
           Builder.CreateElementBitCast(PrivAddr, ConvertTypeForMem(PtrTy)),
auto &&PrivCodeGen = [&](CodeGenFunction &CGF, PrePostActionTy &Action) {
       // needed.
       Action.Enter(CGF);
 
       if (PrivatizeDevicePointers) {
         OMPPrivateScope PrivateScope(CGF);
         // Emit all instances of the use_device_ptr clause.
         for (const auto *C : S.getClausesOfKind<OMPUseDevicePtrClause>())
           CGF.EmitOMPUseDevicePtrClause(*C, PrivateScope,
-                                        Info.CaptureDeviceAddrMap);
+                                        Info.DevicePtrInfoMap);
         for (const auto *C : S.getClausesOfKind<OMPUseDeviceAddrClause>())
           CGF.EmitOMPUseDeviceAddrClause(*C, PrivateScope,
                                          Info.CaptureDeviceAddrMap);
         (void)PrivateScope.Privatize();
         RCG(CGF);
       } else {
         // If we don't have target devices, don't bother emitting the data
         // mapping code.

clang/lib/CodeGen/CodeGenFunction.h

std::pair<bool, RValue> EmitOMPAtomicSimpleUpdateExpr(
       llvm::AtomicOrdering AO, SourceLocation Loc,
       const llvm::function_ref<RValue(RValue)> CommonGen);
   bool EmitOMPFirstprivateClause(const OMPExecutableDirective &D,
                                  OMPPrivateScope &PrivateScope);
   void EmitOMPPrivateClause(const OMPExecutableDirective &D,
                             OMPPrivateScope &PrivateScope);
   void EmitOMPUseDevicePtrClause(
       const OMPUseDevicePtrClause &C, OMPPrivateScope &PrivateScope,
-      const llvm::DenseMap<const ValueDecl *, Address> &CaptureDeviceAddrMap);
+      const llvm::SmallMapVector<const llvm::Value *,
+                                 std::pair<llvm::Value *, llvm::Value *>, 4>
+          &DevicePtrInfoMap);
   void EmitOMPUseDeviceAddrClause(
       const OMPUseDeviceAddrClause &C, OMPPrivateScope &PrivateScope,
-      const llvm::DenseMap<const ValueDecl *, Address> &CaptureDeviceAddrMap);
+      const llvm::DenseMap<const ValueDecl *, llvm::Value *>
+          CaptureDeviceAddrMap);
   /// Emit code for copyin clause in \a D directive. The next code is
   /// generated at the start of outlined functions for directives:
   /// \code
   /// threadprivate_var1 = master_threadprivate_var1;
   /// operator=(threadprivate_var2, master_threadprivate_var2);
   /// ...
   /// __kmpc_barrier(&loc, global_tid);
   /// \endcode

clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp

void foo(double *&arg) {
     a++;
     la++;
 
     // CK2:     [[BP1:%.+]] = getelementptr inbounds [3 x ptr], ptr %{{.+}}, i32 0, i32 1
     // CK2:     store ptr [[RVAL1:%.+]], ptr [[BP1]],
     // CK2:     [[BP2:%.+]] = getelementptr inbounds [3 x ptr], ptr %{{.+}}, i32 0, i32 2
     // CK2:     store ptr [[RVAL2:%.+]], ptr [[BP2]],
     // CK2:     call void @__tgt_target_data_begin{{.+}}[[MTYPE03]]
+    // CK2:     [[VAL1:%.+]] = load ptr, ptr [[BP1]],
+    // CK2:     store ptr [[VAL1]], ptr [[PVT1:%.+]],
     // CK2:     [[VAL2:%.+]] = load ptr, ptr [[BP2]],
     // CK2:     store ptr [[VAL2]], ptr [[PVT2:%.+]],
     // CK2:     store ptr [[PVT2]], ptr [[_PVT2:%.+]],
-    // CK2:     [[VAL1:%.+]] = load ptr, ptr [[BP1]],
-    // CK2:     store ptr [[VAL1]], ptr [[PVT1:%.+]],
     // CK2:     store ptr [[PVT1]], ptr [[_PVT1:%.+]],
     // CK2:     [[TT2:%.+]] = load ptr, ptr [[_PVT2]],
     // CK2:     [[_TT2:%.+]] = load ptr, ptr [[TT2]],
     // CK2:     getelementptr inbounds double, ptr [[_TT2]], i32 1
     // CK2:     [[TT1:%.+]] = load ptr, ptr [[_PVT1]],
     // CK2:     [[_TT1:%.+]] = load ptr, ptr [[TT1]],
     // CK2:     getelementptr inbounds double, ptr [[_TT1]], i32 1
     #pragma omp target data map(b[:10]) use_device_ptr(a, b)

llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h

class TargetDataInfo {
     bool RequiresDevicePointerInfo = false;
     /// Set to true if Clang emits separate runtime calls for the beginning and
     /// end of the region.  These calls might have separate map type arrays.
     bool SeparateBeginEndCalls = false;
 
   public:
     TargetDataRTArgs RTArgs;
 
+    SmallMapVector<const Value *, std::pair<Value *, Value *>, 4>
+        DevicePtrInfoMap;
+
     /// Indicate whether any user-defined mapper exists.
     bool HasMapper = false;
     /// The total number of pointers passed to the runtime library.
     unsigned NumberOfPtrs = 0u;
 
     explicit TargetDataInfo() {}
     explicit TargetDataInfo(bool RequiresDevicePointerInfo,
                             bool SeparateBeginEndCalls)
bool isValid() {
       return RTArgs.BasePointersArray && RTArgs.PointersArray &&
              RTArgs.SizesArray && RTArgs.MapTypesArray &&
              (!HasMapper || RTArgs.MappersArray) && NumberOfPtrs;
     }
     bool requiresDevicePointerInfo() { return RequiresDevicePointerInfo; }
     bool separateBeginEndCalls() { return SeparateBeginEndCalls; }
   };
 
+  enum class DeviceInfoTy { None, Pointer, Address };
   using MapValuesArrayTy = SmallVector<Value *, 4>;
+  using MapDeviceInfoArrayTy = SmallVector<DeviceInfoTy, 4>;
   using MapFlagsArrayTy = SmallVector<omp::OpenMPOffloadMappingFlags, 4>;
   using MapNamesArrayTy = SmallVector<Constant *, 4>;
   using MapDimArrayTy = SmallVector<uint64_t, 4>;
   using MapNonContiguousArrayTy = SmallVector<MapValuesArrayTy, 4>;
 
   /// This structure contains combined information generated for mappable
   /// clauses, including base pointers, pointers, sizes, map types, user-defined
   /// mappers, and non-contiguous information.
   struct MapInfosTy {
     struct StructNonContiguousInfo {
       bool IsNonContiguous = false;
       MapDimArrayTy Dims;
       MapNonContiguousArrayTy Offsets;
       MapNonContiguousArrayTy Counts;
       MapNonContiguousArrayTy Strides;
     };
     MapValuesArrayTy BasePointers;
     MapValuesArrayTy Pointers;
+    MapDeviceInfoArrayTy DevicePointers;
     MapValuesArrayTy Sizes;
     MapFlagsArrayTy Types;
     MapNamesArrayTy Names;
     StructNonContiguousInfo NonContigInfo;
 
     /// Append arrays in \a CurInfo.
     void append(MapInfosTy &CurInfo) {
       BasePointers.append(CurInfo.BasePointers.begin(),
                           CurInfo.BasePointers.end());
       Pointers.append(CurInfo.Pointers.begin(), CurInfo.Pointers.end());
+      DevicePointers.append(CurInfo.DevicePointers.begin(),
+                            CurInfo.DevicePointers.end());
       Sizes.append(CurInfo.Sizes.begin(), CurInfo.Sizes.end());
       Types.append(CurInfo.Types.begin(), CurInfo.Types.end());
       Names.append(CurInfo.Names.begin(), CurInfo.Names.end());
       NonContigInfo.Dims.append(CurInfo.NonContigInfo.Dims.begin(),
                                 CurInfo.NonContigInfo.Dims.end());
       NonContigInfo.Offsets.append(CurInfo.NonContigInfo.Offsets.begin(),
                                    CurInfo.NonContigInfo.Offsets.end());
       NonContigInfo.Counts.append(CurInfo.NonContigInfo.Counts.begin(),
void emitNonContiguousDescriptor(InsertPointTy AllocaIP,
                                    TargetDataInfo &Info);
 
   /// Emit the arrays used to pass the captures and map information to the
   /// offloading runtime library. If there is no map or capture information,
   /// return nullptr by reference.
   void emitOffloadingArrays(
       InsertPointTy AllocaIP, InsertPointTy CodeGenIP, MapInfosTy &CombinedInfo,
       TargetDataInfo &Info, bool IsNonContiguous = false,
-      function_ref<void(unsigned int, Value *, Value *)> DeviceAddrCB = nullptr,
+      function_ref<void(unsigned int, Value *)> DeviceAddrCB = nullptr,
       function_ref<Value *(unsigned int)> CustomMapperCB = nullptr);
 
   /// Creates offloading entry for the provided entry ID \a ID, address \a
   /// Addr, size \a Size, and flags \a Flags.
   void createOffloadEntry(Constant *ID, Constant *Addr, uint64_t Size,
                           int32_t Flags, GlobalValue::LinkageTypes);
 
   /// The kind of errors that can occur when emitting the offload entries and
public:
   OpenMPIRBuilder::InsertPointTy createTargetData(
       const LocationDescription &Loc, InsertPointTy AllocaIP,
       InsertPointTy CodeGenIP, Value *DeviceID, Value *IfCond,
       TargetDataInfo &Info,
       function_ref<MapInfosTy &(InsertPointTy CodeGenIP)> GenMapInfoCB,
       omp::RuntimeFunction *MapperFunc = nullptr,
       function_ref<InsertPointTy(InsertPointTy CodeGenIP, int BodyGenType)>
           BodyGenCB = nullptr,
-      function_ref<void(unsigned int, Value *, Value *)> DeviceAddrCB = nullptr,
+      function_ref<void(unsigned int, Value *)> DeviceAddrCB = nullptr,
       function_ref<Value *(unsigned int)> CustomMapperCB = nullptr);
 
   using TargetBodyGenCallbackTy = function_ref<InsertPointTy(
       InsertPointTy AllocaIP, InsertPointTy CodeGenIP)>;
 
   /// Generator for '#omp target'
   ///
   /// \param Loc where the target data construct was encountered.

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

 OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::createTargetData(
     const LocationDescription &Loc, InsertPointTy AllocaIP,
     InsertPointTy CodeGenIP, Value *DeviceID, Value *IfCond,
     TargetDataInfo &Info,
     function_ref<MapInfosTy &(InsertPointTy CodeGenIP)> GenMapInfoCB,
     omp::RuntimeFunction *MapperFunc,
     function_ref<InsertPointTy(InsertPointTy CodeGenIP, int BodyGenType)>
         BodyGenCB,
-    function_ref<void(unsigned int, Value *, Value *)> DeviceAddrCB,
+    function_ref<void(unsigned int, Value *)> DeviceAddrCB,
     function_ref<Value *(unsigned int)> CustomMapperCB) {
   if (!updateToLocation(Loc))
     return InsertPointTy();
 
   Builder.restoreIP(CodeGenIP);
   bool IsStandAlone = !BodyGenCB;
 
   // Generate the code for the opening of the data environment. Capture all the
if (IsStandAlone) {
       Builder.CreateCall(getOrCreateRuntimeFunctionPtr(*MapperFunc),
                          OffloadingArgs);
     } else {
       Function *beginMapperFunc = getOrCreateRuntimeFunctionPtr(
           omp::OMPRTL___tgt_target_data_begin_mapper);
 
       Builder.CreateCall(beginMapperFunc, OffloadingArgs);
 
+      for (auto DeviceMap : Info.DevicePtrInfoMap) {
+        if (isa<AllocaInst>(DeviceMap.second.second)) {
+          auto *LI =
+              Builder.CreateLoad(Builder.getPtrTy(), DeviceMap.second.first);
+          Builder.CreateStore(LI, DeviceMap.second.second);
+        }
+      }
+
       // If device pointer privatization is required, emit the body of the
       // region here. It will have to be duplicated: with and without
       // privatization.
       Builder.restoreIP(BodyGenCB(Builder.saveIP(), 1));
     }
   };
 
   // If we need device pointer privatization, we need to emit the body of the
Builder.CreateAlignedStore(
         DAddr, P, M.getDataLayout().getPrefTypeAlign(Builder.getInt8PtrTy()));
     ++L;
   }
 }
 
 void OpenMPIRBuilder::emitOffloadingArrays(
     InsertPointTy AllocaIP, InsertPointTy CodeGenIP, MapInfosTy &CombinedInfo,
     TargetDataInfo &Info, bool IsNonContiguous,
-    function_ref<void(unsigned int, Value *, Value *)> DeviceAddrCB,
+    function_ref<void(unsigned int, Value *)> DeviceAddrCB,
     function_ref<Value *(unsigned int)> CustomMapperCB) {
 
   // Reset the array information.
   Info.clearArrayInfo();
   Info.NumberOfPtrs = CombinedInfo.BasePointers.size();
 
   if (Info.NumberOfPtrs) {
     Builder.restoreIP(AllocaIP);
for (unsigned I = 0; I < Info.NumberOfPtrs; ++I) {
           Info.RTArgs.BasePointersArray, 0, I);
       BP = Builder.CreatePointerBitCastOrAddrSpaceCast(
           BP, BPVal->getType()->getPointerTo(/*AddrSpace=*/0));
       Builder.CreateAlignedStore(
           BPVal, BP,
           M.getDataLayout().getPrefTypeAlign(Builder.getInt8PtrTy()));
 
       if (Info.requiresDevicePointerInfo()) {
+        if (CombinedInfo.DevicePointers[I] == DeviceInfoTy::Pointer) {
+          CodeGenIP = Builder.saveIP();
+          Builder.restoreIP(AllocaIP);
+          Info.DevicePtrInfoMap[BPVal] = {
+              BP, Builder.CreateAlloca(Builder.getPtrTy())};
+          Builder.restoreIP(CodeGenIP);
+          assert(DeviceAddrCB);
+          DeviceAddrCB(I, Info.DevicePtrInfoMap[BPVal].second);
+        } else if (CombinedInfo.DevicePointers[I] == DeviceInfoTy::Address) {
+          Info.DevicePtrInfoMap[BPVal] = {BP, BP};
           assert(DeviceAddrCB);
-        DeviceAddrCB(I, BP, BPVal);
+          DeviceAddrCB(I, BP);
+        }
       }
 
       Value *PVal = CombinedInfo.Pointers[I];
       Value *P = Builder.CreateConstInBoundsGEP2_32(
           ArrayType::get(Builder.getInt8PtrTy(), Info.NumberOfPtrs),
           Info.RTArgs.PointersArray, 0, I);
       P = Builder.CreatePointerBitCastOrAddrSpaceCast(
           P, PVal->getType()->getPointerTo(/*AddrSpace=*/0));