creates array allocas for values that are identified as spanning multiple subcfgs
| 710 | |
| 711 | // creates array allocas for values that are identified as spanning multiple subcfgs |
| 712 | void SubCFG::arrayifyMultiSubCfgValues( |
| 713 | llvm::DenseMap<llvm::Instruction *, llvm::AllocaInst *> &InstAllocaMap, |
| 714 | llvm::DenseMap<llvm::Instruction *, llvm::AllocaInst *> &BaseInstAllocaMap, |
| 715 | llvm::DenseMap<llvm::Instruction *, llvm::SmallVector<llvm::Instruction *, 8>> |
| 716 | &ContInstReplicaMap, |
| 717 | llvm::ArrayRef<SubCFG> SubCFGs, llvm::Instruction *AllocaIP, llvm::Value *ReqdArrayElements, |
| 718 | VectorizationInfo &VecInfo) { |
| 719 | llvm::SmallPtrSet<llvm::BasicBlock *, 16> OtherCFGBlocks; |
| 720 | for (auto &Cfg : SubCFGs) { |
| 721 | if (&Cfg != this) |
| 722 | OtherCFGBlocks.insert(Cfg.Blocks_.begin(), Cfg.Blocks_.end()); |
| 723 | } |
| 724 | |
| 725 | for (auto *BB : Blocks_) { |
| 726 | for (auto &I : *BB) { |
| 727 | if (&I == ContIdx_) |
| 728 | continue; |
| 729 | if (InstAllocaMap.lookup(&I)) |
| 730 | continue; |
| 731 | // if any use is in another subcfg |
| 732 | if (utils::anyOfUsers<llvm::Instruction>(&I, [&OtherCFGBlocks, &I](auto *UI) { |
| 733 | return (UI->getParent() != I.getParent() || |
| 734 | UI->getParent() == I.getParent() && UI->comesBefore(&I)) && |
| 735 | OtherCFGBlocks.contains(UI->getParent()); |
| 736 | })) { |
| 737 | // load from an alloca, just widen alloca |
| 738 | if (auto *LInst = llvm::dyn_cast<llvm::LoadInst>(&I)) |
| 739 | if (auto *Alloca = utils::getLoopStateAllocaForLoad(*LInst)) { |
| 740 | InstAllocaMap.insert({&I, Alloca}); |
| 741 | continue; |
| 742 | } |
| 743 | // GEP from already widened alloca: reuse alloca |
| 744 | if (auto *GEP = llvm::dyn_cast<llvm::GetElementPtrInst>(&I)) |
| 745 | if (GEP->hasMetadata(hipsycl::compiler::MDKind::Arrayified)) { |
| 746 | InstAllocaMap.insert({&I, llvm::cast<llvm::AllocaInst>(GEP->getPointerOperand())}); |
| 747 | continue; |
| 748 | } |
| 749 | |
| 750 | auto Shape = VecInfo.getVectorShape(I); |
| 751 | #ifndef HIPSYCL_NO_PHIS_IN_SPLIT |
| 752 | // if value is uniform, just store to 1-wide alloca |
| 753 | if (Shape.isUniform()) { |
| 754 | HIPSYCL_DEBUG_INFO << "[SubCFG] Value uniform, store to single element alloca " << I |
| 755 | << "\n"; |
| 756 | auto *Alloca = utils::arrayifyInstruction(AllocaIP, &I, ContIdx_, nullptr); |
| 757 | InstAllocaMap.insert({&I, Alloca}); |
| 758 | VecInfo.setVectorShape(*Alloca, VectorShape::uni()); |
| 759 | continue; |
| 760 | } |
| 761 | #endif |
| 762 | // if contiguous, and can be recalculated, don't arrayify but store |
| 763 | // uniform values and insts required for recalculation |
| 764 | if (Shape.isContiguousOrStrided()) { |
| 765 | if (dontArrayifyContiguousValues(I, BaseInstAllocaMap, ContInstReplicaMap, AllocaIP, |
| 766 | ReqdArrayElements, ContIdx_, VecInfo)) { |
| 767 | HIPSYCL_DEBUG_INFO << "[SubCFG] Not arrayifying " << I << "\n"; |
| 768 | continue; |
| 769 | } |
no test coverage detected