Make OpenMP codegen work in the presence of scalar dependences

[grosser]

Summary

This change requires D11536 and D11554. It is still a proof-of-concept for the
following two reasons:

• We did not yet instrument the calls in the RegionGenerator that can possibly introduce stack allocated memory

• It would be nice to find a less intrusive solution, e.g. by just looking at the memory locations in the polyhedral model. With our recent change to distinguish PHI and normal scalar memory, this may be possible.

[grosser]

Update to the lastest changes in Polly

The patches that this patch relied on have been committed to Polly, hence it
can now be directly applied.

[jdoerfert]

Is this still relevant? If so I will take a closer look and give you feedback.

[grosser]

Another solution has already been comitted.

include/polly/CodeGen/BlockGenerators.h

public:
 
   /// @brief Simple vector of instructions to store escape users.
   using EscapeUserVectorTy = SmallVector<Instruction *, 4>;
 
   /// @brief Map type to resolve escaping users for scalar instructions.
   ///
   /// @see The EscapeMap member.
   using EscapeUsersAllocaMapTy =
-      DenseMap<Instruction *, std::pair<AllocaInst *, EscapeUserVectorTy>>;
+      DenseMap<Instruction *, std::pair<Value *, EscapeUserVectorTy>>;
 
   ///@}
 
   /// @brief Create a generator for basic blocks.
   ///
   /// @param Builder     The LLVM-IR Builder used to generate the statement. The
   ///                    code is generated at the location, the Builder points
   ///                    to.
public:
   ///
   /// @param Stmt      The block statement to code generate.
   /// @param GlobalMap A mapping from old values to their new values
   ///                  (for values recalculated in the new ScoP, but not
   ///                  within this basic block).
   /// @param LTS       A map from old loops to new induction variables as SCEVs.
   void copyStmt(ScopStmt &Stmt, ValueMapT &GlobalMap, LoopToScevMapT &LTS);
 
+  /// @brief Return (and create) the allocations needed for a given ScopStmt
+  ///
+  /// In some cases, e.g. the generation of OpenMP parallel code, it is
+  /// sometimes necessary to know ahead of time which memory locations will be
+  /// used when code-generating scalar dependences of PHI nodes.
+  /// This function returns a vector containing such memory locations and
+  /// introduces them if they have not yet been allocated.
+  ///
+  /// @param Stmt      The statement we look at
+  /// @param ScalarMap Map from scalars to their demoted location.
+  /// @param PHIOpMap  Map from PHIs to their demoted operand location.
+  std::vector<Value *> getScalarAllocationsForStmt(const ScopStmt &Stmt,
+                                                   ScalarAllocaMapTy &ScalarMap,
+                                                   ScalarAllocaMapTy &PHIOpMap);
+
   /// @brief Finalize the code generation for the SCoP @p S.
   ///
   /// This will initialize and finalize the scalar variables we demoted during
   /// the code generation.
   ///
   /// @see createScalarInitialization(Region &, ValueMapT &)
   /// @see createScalarFinalization(Region &)
   void finalizeSCoP(Scop &S, ValueMapT &VMap);
protected:
 
   /// @brief Return the alloca for @p ScalarBase in @p Map.
   ///
   /// If no alloca was mapped to @p ScalarBase in @p Map a new one is created
   /// and named after @p ScalarBase with the suffix @p NameExt.
   ///
   /// @param ScalarBase The demoted scalar instruction.
   /// @param Map        The map we should look for a mapped alloca instruction.
+  /// @param GlobalMap  A map that contains possible lobal rewrites for the
+  ///                   scalar or global map.
   /// @param NameExt    The suffix we add to the name of a new created alloca.
   /// @param IsNew      If set it will hold true iff the alloca was created.
   ///
   /// @returns The alloca for @p ScalarBase in @p Map.
-  AllocaInst *getOrCreateAlloca(Instruction *ScalarBase, ScalarAllocaMapTy &Map,
-                                const char *NameExt = ".s2a",
-                                bool *IsNew = nullptr);
+  Value *getOrCreateAlloca(Instruction *ScalarBase, ScalarAllocaMapTy &Map,
+                           ValueMapT &GlobalMap, const char *NameExt = ".s2a",
 
   /// @brief Generate reload of scalars demoted to memory and needed by @p Inst.
   ///
-  /// @param Stmt  The statement we generate code for.
-  /// @param Inst  The instruction that might need reloaded values.
-  /// @param BBMap A mapping from old values to their new values in this block.
-  virtual void generateScalarLoads(ScopStmt &Stmt, const Instruction *Inst,
-                                   ValueMapT &BBMap);
+  /// @param Stmt      The statement we generate code for.
+  /// @param Inst      The instruction that might need reloaded values.
+  /// @param BBMap     A mapping from old values to their new values in this
+  ///                  block.
+  /// @param GlobalMap A mapping from old values to their new values
+  ///                  (for values recalculated in the new ScoP, but not
+  ///                  within this basic block).
+  virtual void generateScalarLoads(const ScopStmt &Stmt,
+                                   const Instruction *Inst, ValueMapT &BBMap,
+                                   ValueMapT &GlobalMap,
+                                   std::vector<Value *> *Addresses = nullptr);
 
   /// @brief Generate the scalar stores for the given statement.
   ///
   /// After the statement @p Stmt was copied all inner-SCoP scalar dependences
   /// starting in @p Stmt (hence all scalar write accesses in @p Stmt) need to
   /// be demoted to memory.
   ///
-  /// @param Stmt  The statement we generate code for.
-  /// @param BB    The basic block we generate code for.
-  /// @param BBMap A mapping from old values to their new values in this block.
-  /// @param GlobalMap A mapping for globally replaced values.
-  virtual void generateScalarStores(ScopStmt &Stmt, BasicBlock *BB,
-                                    ValueMapT &BBMAp, ValueMapT &GlobalMap);
+  /// @param Stmt      The statement we generate code for.
+  /// @param BB        The basic block we generate code for.
+  /// @param BBMap     A mapping from old values to their new values in this
+  ///                  block.
+  /// @param GlobalMap A mapping from old values to their new values
+  ///                  (for values recalculated in the new ScoP, but not
+  ///                  within this basic block).
+  virtual void generateScalarStores(const ScopStmt &Stmt, BasicBlock *BB,
+                                    ValueMapT &BBMAp, ValueMapT &GlobalMap,
+                                    std::vector<Value *> *Addresses = nullptr);
 
   /// @brief Handle users of @p Inst outside the SCoP.
   ///
-  /// @param R        The current SCoP region.
-  /// @param Inst     The current instruction we check.
-  /// @param InstCopy The copy of the instruction @p Inst in the optimized SCoP.
-  void handleOutsideUsers(const Region &R, Instruction *Inst, Value *InstCopy);
+  /// @param R         The current SCoP region.
+  /// @param Inst      The current instruction we check.
+  /// @param InstCopy  The copy of the instruction @p Inst in the optimized
+  ///                  SCoP.
+  /// @param GlobalMap A map that contains possible lobal rewrites for the
+  ///                  scalar or global map.
+  void handleOutsideUsers(const Region &R, Instruction *Inst, Value *InstCopy,
+                          ValueMapT &GlobalMap,
+
+                          std::vector<Value *> *Addresses = nullptr);
 
   /// @brief Initialize the memory of demoted scalars.
   ///
   /// If a PHI node was demoted and one of its predecessor blocks was outside
   /// the SCoP we need to initialize the memory cell we demoted the PHI into
   /// with the value corresponding to that predecessor. As a SCoP is a
   /// __single__ entry region there is at most one such predecessor.
   void createScalarInitialization(Region &R, ValueMapT &VMap);
private:
   /// @param LTS        A map from old loops to new induction variables as
   /// SCEVs.
   void addOperandToPHI(ScopStmt &Stmt, const PHINode *PHI, PHINode *PHICopy,
                        BasicBlock *IncomingBB, ValueMapT &GlobalMap,
                        LoopToScevMapT &LTS);
 
   /// @brief Generate reload of scalars demoted to memory and needed by @p Inst.
   ///
-  /// @param Stmt  The statement we generate code for.
-  /// @param Inst  The instruction that might need reloaded values.
-  /// @param BBMap A mapping from old values to their new values in this block.
-  virtual void generateScalarLoads(ScopStmt &Stmt, const Instruction *Inst,
-                                   ValueMapT &BBMap) override;
+  /// @param Stmt      The statement we generate code for.
+  /// @param Inst      The instruction that might need reloaded values.
+  /// @param BBMap     A mapping from old values to their new values in this
+  ///                  block.
+  /// @param GlobalMap A mapping from old values to their new values
+  ///                  (for values recalculated in the new ScoP, but not
+  ///                  within this basic block).
+  virtual void
+  generateScalarLoads(const ScopStmt &Stmt, const Instruction *Inst,
+                      ValueMapT &BBMap, ValueMapT &GlobalMap,
+                      std::vector<Value *> *Addresses = nullptr) override;
 
   /// @brief Generate the scalar stores for the given statement.
   ///
   /// After the statement @p Stmt was copied all inner-SCoP scalar dependences
   /// starting in @p Stmt (hence all scalar write accesses in @p Stmt) need to
   /// be demoted to memory.
   ///
   /// @param Stmt  The statement we generate code for.
   /// @param BB    The basic block we generate code for.
   /// @param BBMap A mapping from old values to their new values in this block.
   /// @param GlobalMap  A mapping from old values to their new values
   ///                   (for values recalculated in the new ScoP, but not
   ///                   within this basic block).
-  virtual void generateScalarStores(ScopStmt &Stmt, BasicBlock *BB,
-                                    ValueMapT &BBMAp,
-                                    ValueMapT &GlobalMap) override;
+  virtual void
+  generateScalarStores(const ScopStmt &Stmt, BasicBlock *BB, ValueMapT &BBMAp,
+                       ValueMapT &GlobalMap,
+                       std::vector<Value *> *Addresses = nullptr) override;
 
   /// @brief Copy a single PHI instruction.
   ///
   /// This copies a single PHI instruction and updates references to old values
   /// with references to new values, as defined by GlobalMap and BBMap.
   ///
   /// @param Stmt      The statement to code generate.
   /// @param PHI       The PHI instruction to copy.

include/polly/CodeGen/IslNodeBuilder.h

public:
 
   /// @brief Finalize code generation for the SCoP @p S.
   ///
   /// @see BlockGenerator::finalizeSCoP(Scop &S)
   void finalizeSCoP(Scop &S) { BlockGen.finalizeSCoP(S, ValueMap); }
 
   IslExprBuilder &getExprBuilder() { return ExprBuilder; }
 
+  static int findValuesInBlock(struct FindValuesUser &User,
+                               const ScopStmt *Stmt, const BasicBlock *BB);
+
 private:
   Scop &S;
   PollyIRBuilder &Builder;
   ScopAnnotator &Annotator;
 
   /// @brief A SCEVExpander to create llvm values from SCEVs.
   SCEVExpander Rewriter;
 

lib/CodeGen/BlockGenerators.cpp

Value *BlockGenerator::generateScalarStore(ScopStmt &Stmt,
   return NewStore;
 }
 
 void BlockGenerator::copyInstruction(ScopStmt &Stmt, const Instruction *Inst,
                                      ValueMapT &BBMap, ValueMapT &GlobalMap,
                                      LoopToScevMapT &LTS) {
 
   // First check for possible scalar dependences for this instruction.
-  generateScalarLoads(Stmt, Inst, BBMap);
+  generateScalarLoads(Stmt, Inst, BBMap, GlobalMap);
 
   // Terminator instructions control the control flow. They are explicitly
   // expressed in the clast and do not need to be copied.
   if (Inst->isTerminator())
     return;
 
   Loop *L = getLoopForInst(Inst);
   if ((Stmt.isBlockStmt() || !Stmt.getRegion()->contains(L)) &&
void BlockGenerator::copyBB(ScopStmt &Stmt, BasicBlock *BB, BasicBlock *CopyBB,
 
   // After a basic block was copied store all scalars that escape this block
   // in their alloca. First the scalars that have dependences inside the SCoP,
   // then the ones that might escape the SCoP.
   generateScalarStores(Stmt, BB, BBMap, GlobalMap);
 
   const Region &R = Stmt.getParent()->getRegion();
   for (Instruction &Inst : *BB)
-    handleOutsideUsers(R, &Inst, BBMap[&Inst]);
+    handleOutsideUsers(R, &Inst, BBMap[&Inst], GlobalMap);
 }
 
-AllocaInst *BlockGenerator::getOrCreateAlloca(Instruction *ScalarBase,
-                                              ScalarAllocaMapTy &Map,
-                                              const char *NameExt,
-                                              bool *IsNew) {
+Value *BlockGenerator::getOrCreateAlloca(Instruction *ScalarBase,
+                                         ValueMapT &GlobalMap,
+                                         const char *NameExt, bool *IsNew) {
 
   // Check if an alloca was cached for the base instruction.
   AllocaInst *&Addr = Map[ScalarBase];
 
   // If needed indicate if it was found already or will be created.
   if (IsNew)
     *IsNew = (Addr == nullptr);
 
   // If no alloca was found create one and insert it in the entry block.
   if (!Addr) {
     auto *Ty = ScalarBase->getType();
     Addr = new AllocaInst(Ty, ScalarBase->getName() + NameExt);
     Addr->insertBefore(EntryBB->getFirstInsertionPt());
   }
 
-  return Addr;
+  Value *UpdatedAddr = Addr;
+
+  if (Value *New = GlobalMap.lookup(UpdatedAddr))
+    UpdatedAddr = New;
+
+  return UpdatedAddr;
+}
+
+std::vector<Value *>
+BlockGenerator::getScalarAllocationsForStmt(const ScopStmt &Stmt,
+                                            ScalarAllocaMapTy &ScalarMap,
+                                            ScalarAllocaMapTy &PHIOpMap) {
+  ValueMapT GlobalMap;
+  ValueMapT BBMap;
+  std::vector<Value *> Allocations;
+  BasicBlock *BB = Stmt.getBasicBlock();
+
+  for (auto &Inst : *BB)
+    generateScalarLoads(Stmt, &Inst, BBMap, GlobalMap, &Allocations);
+
+  generateScalarStores(Stmt, BB, BBMap, GlobalMap, &Allocations);
+
+  auto &R = Stmt.getParent()->getRegion();
+  for (auto &Inst : *BB)
+    handleOutsideUsers(R, &Inst, nullptr, GlobalMap, &Allocations);
+
+  return Allocations;
 }
 
 void BlockGenerator::handleOutsideUsers(const Region &R, Instruction *Inst,
-                                        Value *InstCopy) {
+                                        Value *InstCopy, ValueMapT &GlobalMap,
+                                        std::vector<Value *> *Addresses) {
   BasicBlock *ExitBB = R.getExit();
 
   EscapeUserVectorTy EscapeUsers;
   for (User *U : Inst->users()) {
 
     // Non-instruction user will never escape.
     Instruction *UI = dyn_cast<Instruction>(U);
     if (!UI)
void BlockGenerator::handleOutsideUsers(const Region &R, Instruction *Inst,
   // created by an already handled scalar dependence we have to initialize it
   // also. Lastly, if the instruction was copied multiple times we already did
   // this and can exit.
   if (EscapeMap.count(Inst))
     return;
 
   // Get or create an escape alloca for this instruction.
   bool IsNew;
-  AllocaInst *ScalarAddr =
-      getOrCreateAlloca(Inst, ScalarMap, ".escape", &IsNew);
+  Value *ScalarAddr =
+      getOrCreateAlloca(Inst, ScalarMap, GlobalMap, ".escape", &IsNew);
+
+  if (Addresses) {
+    Addresses->push_back(ScalarAddr);
+    return;
+  }
 
   // Remember that this instruction has escape uses and the escape alloca.
   EscapeMap[Inst] = std::make_pair(ScalarAddr, std::move(EscapeUsers));
 
   // If the escape alloca was just created store the instruction in there,
   // otherwise that happened already.
   if (IsNew) {
     assert(InstCopy && "Except PHIs every instruction should have a copy!");
     Builder.CreateStore(InstCopy, ScalarAddr);
   }
 }
 
-void BlockGenerator::generateScalarLoads(ScopStmt &Stmt,
+void BlockGenerator::generateScalarLoads(const ScopStmt &Stmt,
                                          const Instruction *Inst,
-                                         ValueMapT &BBMap) {
+                                         ValueMapT &BBMap, ValueMapT &GlobalMap,
+                                         std::vector<Value *> *Addresses) {
   auto *MAL = Stmt.lookupAccessesFor(Inst);
 
   if (!MAL)
     return;
 
   for (MemoryAccess &MA : *MAL) {
-    AllocaInst *Address;
+    Value *Address;
     if (!MA.isScalar() || !MA.isRead())
       continue;
 
     auto Base = cast<Instruction>(MA.getBaseAddr());
 
     if (MA.getScopArrayInfo()->isPHI())
-      Address = getOrCreateAlloca(Base, PHIOpMap, ".phiops");
+      Address = getOrCreateAlloca(Base, PHIOpMap, GlobalMap, ".phiops");
     else
-      Address = getOrCreateAlloca(Base, ScalarMap, ".s2a");
+      Address = getOrCreateAlloca(Base, ScalarMap, GlobalMap, ".s2a");
 
+    if (Addresses) {
+      Addresses->push_back(Address);
+    } else {
       BBMap[Base] = Builder.CreateLoad(Address, Address->getName() + ".reload");
     }
   }
+}
 
 Value *BlockGenerator::getNewScalarValue(Value *ScalarValue, const Region &R,
                                          ScalarAllocaMapTy &ReloadMap,
                                          ValueMapT &BBMap,
                                          ValueMapT &GlobalMap) {
   // If the value we want to store is an instruction we might have demoted it
   // in order to make it accessible here. In such a case a reload is
   // necessary. If it is no instruction it will always be a value that
assert(ReloadMap.count(ScalarValueInst) &&
          "ScalarInst not mapped in the block and not in the given reload map!");
   Value *ReloadAddr = ReloadMap[ScalarValueInst];
   ScalarValue =
       Builder.CreateLoad(ReloadAddr, ReloadAddr->getName() + ".reload");
 
   return ScalarValue;
 }
 
-void BlockGenerator::generateScalarStores(ScopStmt &Stmt, BasicBlock *BB,
+void BlockGenerator::generateScalarStores(const ScopStmt &Stmt, BasicBlock *BB,
                                           ValueMapT &BBMap,
-                                          ValueMapT &GlobalMap) {
+                                          ValueMapT &GlobalMap,
+                                          std::vector<Value *> *Addresses) {
   const Region &R = Stmt.getParent()->getRegion();
 
   assert(Stmt.isBlockStmt() && BB == Stmt.getBasicBlock() &&
          "Region statements need to use the generateScalarStores() "
          "function in the RegionGenerator");
 
   for (MemoryAccess *MA : Stmt) {
     if (!MA->isScalar() || MA->isRead())
       continue;
 
     Instruction *Base = cast<Instruction>(MA->getBaseAddr());
     Instruction *Inst = MA->getAccessInstruction();
 
     Value *Val = nullptr;
-    AllocaInst *Address = nullptr;
+    Value *Address = nullptr;
 
     if (MA->getScopArrayInfo()->isPHI()) {
       PHINode *BasePHI = dyn_cast<PHINode>(Base);
       int PHIIdx = BasePHI->getBasicBlockIndex(BB);
-      Address = getOrCreateAlloca(Base, PHIOpMap, ".phiops");
+      Address = getOrCreateAlloca(Base, PHIOpMap, GlobalMap, ".phiops");
       Val = BasePHI->getIncomingValue(PHIIdx);
     } else {
-      Address = getOrCreateAlloca(Base, ScalarMap, ".s2a");
+      Address = getOrCreateAlloca(Base, ScalarMap, GlobalMap, ".s2a");
       Val = Inst;
     }
+
+    if (Addresses) {
+      Addresses->push_back(Address);
+    } else {
       Val = getNewScalarValue(Val, R, ScalarMap, BBMap, GlobalMap);
       Builder.CreateStore(Val, Address);
     }
   }
+}
 
 void BlockGenerator::createScalarInitialization(Region &R,
                                                 ValueMapT &GlobalMap) {
   // The split block __just before__ the region and optimized region.
   BasicBlock *SplitBB = R.getEnteringBlock();
   BranchInst *SplitBBTerm = cast<BranchInst>(SplitBB->getTerminator());
   assert(SplitBBTerm->getNumSuccessors() == 2 && "Bad region entering block!");
 
void BlockGenerator::createScalarFinalization(Region &R) {
 
   Builder.SetInsertPoint(OptExitBB->getTerminator());
   for (const auto &EscapeMapping : EscapeMap) {
     // Extract the escaping instruction and the escaping users as well as the
     // alloca the instruction was demoted to.
     Instruction *EscapeInst = EscapeMapping.getFirst();
     const auto &EscapeMappingValue = EscapeMapping.getSecond();
     const EscapeUserVectorTy &EscapeUsers = EscapeMappingValue.second;
-    AllocaInst *ScalarAddr = EscapeMappingValue.first;
+    Value *ScalarAddr = EscapeMappingValue.first;
 
     // Reload the demoted instruction in the optimized version of the SCoP.
     Instruction *EscapeInstReload =
         Builder.CreateLoad(ScalarAddr, EscapeInst->getName() + ".final_reload");
 
     // Create the merge PHI that merges the optimized and unoptimized version.
     PHINode *MergePHI = PHINode::Create(EscapeInst->getType(), 2,
                                         EscapeInst->getName() + ".merge");
void RegionGenerator::copyStmt(ScopStmt &Stmt, ValueMapT &GlobalMap,
   // the copied instructions.
   for (auto &BBMap : RegionMaps)
     GlobalMap.insert(BBMap.second.begin(), BBMap.second.end());
 
   // Reset the old insert point for the build.
   Builder.SetInsertPoint(ExitBBCopy->begin());
 }
 
-void RegionGenerator::generateScalarLoads(ScopStmt &Stmt,
+void RegionGenerator::generateScalarLoads(const ScopStmt &Stmt,
                                           const Instruction *Inst,
-                                          ValueMapT &BBMap) {
+                                          ValueMapT &BBMap,
+                                          ValueMapT &GlobalMap,
+                                          std::vector<Value *> *Addresses) {
 
   // Inside a non-affine region PHI nodes are copied not demoted. Once the
   // phi is copied it will reload all inputs from outside the region, hence
   // we do not need to generate code for the read access of the operands of a
   // PHI.
   if (isa<PHINode>(Inst))
     return;
 
-  return BlockGenerator::generateScalarLoads(Stmt, Inst, BBMap);
+  return BlockGenerator::generateScalarLoads(Stmt, Inst, BBMap, GlobalMap,
+                                             Addresses);
 }
 
-void RegionGenerator::generateScalarStores(ScopStmt &Stmt, BasicBlock *BB,
+void RegionGenerator::generateScalarStores(const ScopStmt &Stmt, BasicBlock *BB,
                                            ValueMapT &BBMap,
-                                           ValueMapT &GlobalMap) {
+                                           ValueMapT &GlobalMap,
+                                           std::vector<Value *> *Addresses) {
   const Region &R = Stmt.getParent()->getRegion();
 
   Region *StmtR = Stmt.getRegion();
   assert(StmtR && "Block statements need to use the generateScalarStores() "
                   "function in the BlockGenerator");
 
   for (MemoryAccess *MA : Stmt) {
 
     if (!MA->isScalar() || MA->isRead())
       continue;
 
     Instruction *ScalarBase = cast<Instruction>(MA->getBaseAddr());
     Instruction *ScalarInst = MA->getAccessInstruction();
     PHINode *ScalarBasePHI = dyn_cast<PHINode>(ScalarBase);
 
     Value *Val = nullptr;
-    AllocaInst *ScalarAddr = nullptr;
+    Value *ScalarAddr = nullptr;
 
     if (MA->getScopArrayInfo()->isPHI()) {
       int PHIIdx = ScalarBasePHI->getBasicBlockIndex(BB);
-      ScalarAddr = getOrCreateAlloca(ScalarBase, PHIOpMap, ".phiops");
+      ScalarAddr =
+          getOrCreateAlloca(ScalarBase, PHIOpMap, GlobalMap, ".phiops");
       Val = ScalarBasePHI->getIncomingValue(PHIIdx);
     } else {
-      ScalarAddr = getOrCreateAlloca(ScalarBase, ScalarMap, ".s2a");
+      ScalarAddr = getOrCreateAlloca(ScalarBase, ScalarMap, GlobalMap, ".s2a");
       Val = ScalarInst;
     }
 
+    if (Addresses) {
+      Addresses->push_back(ScalarAddr);
+    } else {
       Val = getNewScalarValue(Val, R, ScalarMap, BBMap, GlobalMap);
       Builder.CreateStore(Val, ScalarAddr);
     }
   }
+}
 
 void RegionGenerator::addOperandToPHI(ScopStmt &Stmt, const PHINode *PHI,
                                       PHINode *PHICopy, BasicBlock *IncomingBB,
                                       ValueMapT &GlobalMap,
                                       LoopToScevMapT &LTS) {
   Region *StmtR = Stmt.getRegion();
 
   // If the incoming block was not yet copied mark this PHI as incomplete.
if (StmtR->contains(IncomingBB)) {
     Value *Op = PHI->getIncomingValueForBlock(IncomingBB);
     OpCopy =
         getNewValue(Stmt, Op, BBCopyMap, GlobalMap, LTS, getLoopForInst(PHI));
   } else {
 
     if (PHICopy->getBasicBlockIndex(BBCopy) >= 0)
       return;
 
-    AllocaInst *PHIOpAddr =
-        getOrCreateAlloca(const_cast<PHINode *>(PHI), PHIOpMap, ".phiops");
+    Value *PHIOpAddr = getOrCreateAlloca(const_cast<PHINode *>(PHI), PHIOpMap,
+                                         GlobalMap, ".phiops");
     OpCopy = new LoadInst(PHIOpAddr, PHIOpAddr->getName() + ".reload",
                           BlockMap[IncomingBB]->getTerminator());
   }
 
   assert(OpCopy && "Incoming PHI value was not copied properly");
   assert(BBCopy && "Incoming PHI block was not copied properly");
   PHICopy->addIncoming(OpCopy, BBCopy);
 }

lib/CodeGen/IslNodeBuilder.cpp

 }
 
 struct FindValuesUser {
   LoopInfo &LI;
   ScalarEvolution &SE;
   Region &R;
   SetVector<Value *> &Values;
   SetVector<const SCEV *> &SCEVs;
+  IslNodeBuilder *NodeBuilder;
 };
 
 /// @brief Extract the values and SCEVs needed to generate code for a block.
-static int findValuesInBlock(struct FindValuesUser &User, const ScopStmt *Stmt,
+int IslNodeBuilder::findValuesInBlock(struct FindValuesUser &User,
+                                      const ScopStmt *Stmt,
                                       const BasicBlock *BB) {
   // Check all the operands of instructions in the basic block.
   for (const Instruction &Inst : *BB) {
     for (Value *SrcVal : Inst.operands()) {
       if (Instruction *OpInst = dyn_cast<Instruction>(SrcVal))
         if (canSynthesize(OpInst, &User.LI, &User.SE, &User.R)) {
           User.SCEVs.insert(
               User.SE.getSCEVAtScope(OpInst, User.LI.getLoopFor(BB)));
           continue;
         }
       if (Instruction *OpInst = dyn_cast<Instruction>(SrcVal))
         if (Stmt->getParent()->getRegion().contains(OpInst))
           continue;
 
       if (isa<Instruction>(SrcVal) || isa<Argument>(SrcVal))
         User.Values.insert(SrcVal);
     }
   }
+
+  auto ScalarLocations = User.NodeBuilder->BlockGen.getScalarAllocationsForStmt(
+      *Stmt, User.NodeBuilder->ScalarMap, User.NodeBuilder->PHIOpMap);
+
+  User.Values.insert(ScalarLocations.begin(), ScalarLocations.end());
   return 0;
 }
 
 /// Extract the values and SCEVs needed to generate code for a ScopStmt.
 ///
 /// This function extracts a ScopStmt from a given isl_set and computes the
 /// Values this statement depends on as well as a set of SCEV expressions that
 /// need to be synthesized when generating code for this statment.
 static isl_stat findValuesInStmt(isl_set *Set, void *UserPtr) {
   isl_id *Id = isl_set_get_tuple_id(Set);
   struct FindValuesUser &User = *static_cast<struct FindValuesUser *>(UserPtr);
   const ScopStmt *Stmt = static_cast<const ScopStmt *>(isl_id_get_user(Id));
 
   if (Stmt->isBlockStmt())
-    findValuesInBlock(User, Stmt, Stmt->getBasicBlock());
+    IslNodeBuilder::findValuesInBlock(User, Stmt, Stmt->getBasicBlock());
   else {
     assert(Stmt->isRegionStmt() &&
            "Stmt was neither block nor region statement");
     for (const BasicBlock *BB : Stmt->getRegion()->blocks())
-      findValuesInBlock(User, Stmt, BB);
+      IslNodeBuilder::findValuesInBlock(User, Stmt, BB);
   }
 
   isl_id_free(Id);
   isl_set_free(Set);
   return isl_stat_ok;
 }
 
 void IslNodeBuilder::getReferencesInSubtree(__isl_keep isl_ast_node *For,
                                             SetVector<Value *> &Values,
                                             SetVector<const Loop *> &Loops) {
 
   SetVector<const SCEV *> SCEVs;
-  struct FindValuesUser FindValues = {LI, SE, S.getRegion(), Values, SCEVs};
+  struct FindValuesUser FindValues = {LI,     SE,    S.getRegion(),
+                                      Values, SCEVs, this};
 
   for (const auto &I : IDToValue)
     Values.insert(I.second);
 
   for (const auto &I : OutsideLoopIterations)
     Values.insert(cast<SCEVUnknown>(I.second)->getValue());
 
   isl_union_set *Schedule = isl_union_map_domain(IslAstInfo::getSchedule(For));

test/Isl/CodeGen/OpenMP/scalars-in-offloaded-section.ll

+; RUN: opt %loadPolly -polly-ignore-aliasing -polly-ast -analyze -polly-parallel -polly-parallel-force -polly-allow-nonaffine < %s | FileCheck %s -check-prefix=AST
+; RUN: opt %loadPolly -polly-ignore-aliasing -polly-codegen -polly-parallel -polly-parallel-force -polly-allow-nonaffine -S < %s | FileCheck %s -check-prefix=IR
+target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+
+; AST: {
+; AST: Stmt_bb3();
+; AST: #pragma simd
+; AST: #pragma omp parallel for
+; AST: for (int c0 = 0; c0 <= p_0; c0 += 1)
+; AST: Stmt_bb6(c0);
+; AST: }
+
+; Check that we correctly pass scalar dependences (here trough %tmp4) from the
+; host code to the device code.
+;
+; IR:   %7 = getelementptr inbounds { i64, double*, double*, i64* }, { i64, double*, double*, i64* }* %polly.par.userContext, i32 0, i32 3
+; IR: store i64* %tmp4.s2a, i64** %7
+
+
+define void @hoge(double* noalias %arg, double* noalias %arg1, double* noalias dereferenceable(8) %arg2) {
+bb:
+  br label %bb3
+
+bb3:                                              ; preds = %bb
+  %tmp = bitcast double* %arg2 to i64*
+  %tmp4 = load i64, i64* %tmp, align 8
+  br i1 false, label %bb12, label %bb5
+
+bb5:                                              ; preds = %bb3
+  br label %bb6
+
+bb6:                                              ; preds = %bb6, %bb5
+  %tmp7 = phi double* [ %arg, %bb5 ], [ %tmp9, %bb6 ]
+  %tmp8 = bitcast double* %tmp7 to i64*
+  store i64 %tmp4, i64* %tmp8, align 8
+  %tmp9 = getelementptr inbounds double, double* %tmp7, i64 1
+  %tmp10 = icmp eq double* %tmp9, %arg1
+  br i1 %tmp10, label %bb11, label %bb6
+
+bb11:                                             ; preds = %bb6
+  br label %bb12
+
+bb12:                                             ; preds = %bb11, %bb3
+  ret void
+}