Create a new alloca of size \a NumElements at \a IPAllocas. The type is taken from \a ToArrayify. At \a InsertionPoint, a store is added that stores the \a ToArrayify value to the alloca element at \a Idx.
| 421 | // At \a InsertionPoint, a store is added that stores the \a ToArrayify |
| 422 | // value to the alloca element at \a Idx. |
| 423 | llvm::AllocaInst *arrayifyValue(llvm::Instruction *IPAllocas, llvm::Value *ToArrayify, |
| 424 | llvm::Instruction *InsertionPoint, llvm::Value *Idx, |
| 425 | llvm::Value *NumElements, llvm::MDTuple *MDAlloca) { |
| 426 | assert(Idx && "Valid WI-Index required"); |
| 427 | |
| 428 | if (!MDAlloca) |
| 429 | MDAlloca = llvm::MDNode::get(IPAllocas->getContext(), |
| 430 | {llvm::MDString::get(IPAllocas->getContext(), MDKind::LoopState)}); |
| 431 | |
| 432 | auto *T = ToArrayify->getType(); |
| 433 | llvm::IRBuilder AllocaBuilder{IPAllocas}; |
| 434 | auto *Alloca = AllocaBuilder.CreateAlloca(T, NumElements, ToArrayify->getName() + "_alloca"); |
| 435 | if (NumElements) |
| 436 | Alloca->setAlignment(llvm::Align{hipsycl::compiler::DefaultAlignment}); |
| 437 | Alloca->setMetadata(hipsycl::compiler::MDKind::Arrayified, MDAlloca); |
| 438 | |
| 439 | llvm::IRBuilder WriteBuilder{InsertionPoint}; |
| 440 | llvm::Value *StoreTarget = Alloca; |
| 441 | if (NumElements) { |
| 442 | auto *GEP = llvm::cast<llvm::GetElementPtrInst>(WriteBuilder.CreateInBoundsGEP( |
| 443 | Alloca->getAllocatedType(), Alloca, Idx, ToArrayify->getName() + "_gep")); |
| 444 | GEP->setMetadata(hipsycl::compiler::MDKind::Arrayified, MDAlloca); |
| 445 | StoreTarget = GEP; |
| 446 | } |
| 447 | WriteBuilder.CreateStore(ToArrayify, StoreTarget); |
| 448 | return Alloca; |
| 449 | } |
| 450 | |
| 451 | // see arrayifyValue. The store is inserted after the \a ToArrayify instruction |
| 452 | llvm::AllocaInst *arrayifyInstruction(llvm::Instruction *IPAllocas, llvm::Instruction *ToArrayify, |
no test coverage detected