| 21 | namespace { |
| 22 | using namespace hipsycl::compiler; |
| 23 | void markLoopParallel(llvm::Function &F, llvm::Loop *L) { |
| 24 | // LLVM < 12.0.1 might miscompile if conditionals in "parallel" loop (https://llvm.org/PR46666) |
| 25 | |
| 26 | // Mark memory accesses with access group |
| 27 | auto *MDAccessGroup = llvm::MDNode::getDistinct(F.getContext(), {}); |
| 28 | for (auto *BB : L->blocks()) { |
| 29 | for (auto &I : *BB) { |
| 30 | if (I.mayReadOrWriteMemory() && !I.hasMetadata(llvm::LLVMContext::MD_access_group)) { |
| 31 | utils::addAccessGroupMD(&I, MDAccessGroup); |
| 32 | } |
| 33 | } |
| 34 | } |
| 35 | |
| 36 | // make the access group parallel w.r.t the WI loop |
| 37 | utils::createParallelAccessesMdOrAddAccessGroup(&F, L, MDAccessGroup); |
| 38 | |
| 39 | // debug, check whether loop is really marked parallel. |
| 40 | if (HIPSYCL_DEBUG_LEVEL_INFO <= hipsycl::common::output_stream::get().get_debug_level()) { |
| 41 | if (utils::isAnnotatedParallel(L)) { |
| 42 | HIPSYCL_DEBUG_INFO << "[ParallelMarker] loop is parallel: " << L->getHeader()->getName() |
| 43 | << "\n"; |
| 44 | } else if (L->getLoopID()) { |
| 45 | assert(L->getLoopID()); |
| 46 | const llvm::Module *M = F.getParent(); |
| 47 | HIPSYCL_DEBUG_WARNING << "[ParallelMarker] failed to mark wi-loop as parallel, loop id for " |
| 48 | << L->getHeader()->getName(); |
| 49 | HIPSYCL_DEBUG_EXECUTE_WARNING( |
| 50 | L->getLoopID()->print(llvm::outs(), M); |
| 51 | for (auto &MDOp : llvm::drop_begin(L->getLoopID()->operands(), 1)) { |
| 52 | MDOp->print(llvm::outs(), M); |
| 53 | } |
| 54 | llvm::outs() << "\n"; |
| 55 | ) |
| 56 | } |
| 57 | } |
| 58 | } |
| 59 | |
| 60 | void addVectorizationHints(const llvm::Function &F, const llvm::TargetTransformInfo &TTI, |
| 61 | const llvm::Loop *L) { |
no test coverage detected