CbC/CbC_llvm: clang/lib/CodeGen/CGStmtOpenMP.cpp comparison

comparison clang/lib/CodeGen/CGStmtOpenMP.cpp @ 221:79ff65ed7e25

LLVM12 Original

author	Shinji KONO <kono@ie.u-ryukyu.ac.jp>
date	Tue, 15 Jun 2021 19:15:29 +0900
parents	0572611fdcc8
children	5f17cb93ff66

comparison

equal deleted inserted replaced

-:42394fc6a535
+:79ff65ed7e25
 #include "clang/AST/Attr.h"
 #include "clang/AST/DeclOpenMP.h"
 #include "clang/AST/OpenMPClause.h"
 #include "clang/AST/Stmt.h"
 #include "clang/AST/StmtOpenMP.h"
+#include "clang/AST/StmtVisitor.h"
 #include "clang/Basic/OpenMPKinds.h"
 #include "clang/Basic/PrettyStackTrace.h"
+#include "llvm/Frontend/OpenMP/OMPConstants.h"
 #include "llvm/Frontend/OpenMP/OMPIRBuilder.h"
 #include "llvm/IR/Constants.h"
 #include "llvm/IR/Instructions.h"
 #include "llvm/Support/AtomicOrdering.h"
 using namespace clang;
 using namespace CodeGen;
 using namespace llvm::omp;
+static const VarDecl *getBaseDecl(const Expr *Ref);
 namespace {
 /// Lexical scope for OpenMP executable constructs, that handles correct codegen
 /// for captured expressions.
 class OMPLexicalScope : public CodeGenFunction::LexicalScope {
 };
 /// Private scope for OpenMP loop-based directives, that supports capturing
 /// of used expression from loop statement.
 class OMPLoopScope : public CodeGenFunction::RunCleanupsScope {
-void emitPreInitStmt(CodeGenFunction &CGF, const OMPLoopDirective &S) {
+void emitPreInitStmt(CodeGenFunction &CGF, const OMPLoopBasedDirective &S) {
+const DeclStmt *PreInits;
 CodeGenFunction::OMPMapVars PreCondVars;
-llvm::DenseSet<const VarDecl *> EmittedAsPrivate;
+if (auto *LD = dyn_cast<OMPLoopDirective>(&S)) {
-for (const auto *E : S.counters()) {
+llvm::DenseSet<const VarDecl *> EmittedAsPrivate;
-const auto *VD = cast<VarDecl>(cast<DeclRefExpr>(E)->getDecl());
+for (const auto *E : LD->counters()) {
-EmittedAsPrivate.insert(VD->getCanonicalDecl());
+const auto *VD = cast<VarDecl>(cast<DeclRefExpr>(E)->getDecl());
-(void)PreCondVars.setVarAddr(
+EmittedAsPrivate.insert(VD->getCanonicalDecl());
-CGF, VD, CGF.CreateMemTemp(VD->getType().getNonReferenceType()));
+(void)PreCondVars.setVarAddr(
-}
+CGF, VD, CGF.CreateMemTemp(VD->getType().getNonReferenceType()));
-// Mark private vars as undefs.
+}
-for (const auto *C : S.getClausesOfKind<OMPPrivateClause>()) {
+// Mark private vars as undefs.
-for (const Expr *IRef : C->varlists()) {
+for (const auto *C : LD->getClausesOfKind<OMPPrivateClause>()) {
-const auto *OrigVD = cast<VarDecl>(cast<DeclRefExpr>(IRef)->getDecl());
+for (const Expr *IRef : C->varlists()) {
-if (EmittedAsPrivate.insert(OrigVD->getCanonicalDecl()).second) {
+const auto *OrigVD =
-(void)PreCondVars.setVarAddr(
+cast<VarDecl>(cast<DeclRefExpr>(IRef)->getDecl());
-CGF, OrigVD,
+if (EmittedAsPrivate.insert(OrigVD->getCanonicalDecl()).second) {
-Address(llvm::UndefValue::get(
+(void)PreCondVars.setVarAddr(
-CGF.ConvertTypeForMem(CGF.getContext().getPointerType(
+CGF, OrigVD,
-OrigVD->getType().getNonReferenceType()))),
+Address(llvm::UndefValue::get(CGF.ConvertTypeForMem(
-CGF.getContext().getDeclAlign(OrigVD)));
+CGF.getContext().getPointerType(
+OrigVD->getType().getNonReferenceType()))),
+CGF.getContext().getDeclAlign(OrigVD)));
+}
 }
 }
-}
+(void)PreCondVars.apply(CGF);
-(void)PreCondVars.apply(CGF);
+// Emit init, __range and __end variables for C++ range loops.
-// Emit init, __range and __end variables for C++ range loops.
+(void)OMPLoopBasedDirective::doForAllLoops(
-const Stmt *Body =
+LD->getInnermostCapturedStmt()->getCapturedStmt(),
-S.getInnermostCapturedStmt()->getCapturedStmt()->IgnoreContainers();
+/*TryImperfectlyNestedLoops=*/true, LD->getLoopsNumber(),
-for (unsigned Cnt = 0; Cnt < S.getCollapsedNumber(); ++Cnt) {
+[&CGF](unsigned Cnt, const Stmt *CurStmt) {
-Body = OMPLoopDirective::tryToFindNextInnerLoop(
+if (const auto *CXXFor = dyn_cast<CXXForRangeStmt>(CurStmt)) {
-Body, /*TryImperfectlyNestedLoops=*/true);
+if (const Stmt *Init = CXXFor->getInit())
-if (auto *For = dyn_cast<ForStmt>(Body)) {
+CGF.EmitStmt(Init);
-Body = For->getBody();
+CGF.EmitStmt(CXXFor->getRangeStmt());
-} else {
+CGF.EmitStmt(CXXFor->getEndStmt());
-assert(isa<CXXForRangeStmt>(Body) &&
+}
-"Expected canonical for loop or range-based for loop.");
+return false;
-auto *CXXFor = cast<CXXForRangeStmt>(Body);
+});
-if (const Stmt *Init = CXXFor->getInit())
+PreInits = cast_or_null<DeclStmt>(LD->getPreInits());
-CGF.EmitStmt(Init);
+} else if (const auto *Tile = dyn_cast<OMPTileDirective>(&S)) {
-CGF.EmitStmt(CXXFor->getRangeStmt());
+PreInits = cast_or_null<DeclStmt>(Tile->getPreInits());
-CGF.EmitStmt(CXXFor->getEndStmt());
+} else {
-Body = CXXFor->getBody();
+llvm_unreachable("Unknown loop-based directive kind.");
 }
-}
+if (PreInits) {
-if (const auto *PreInits = cast_or_null<DeclStmt>(S.getPreInits())) {
 for (const auto *I : PreInits->decls())
 CGF.EmitVarDecl(cast<VarDecl>(*I));
 }
 PreCondVars.restore(CGF);
 }
 public:
-OMPLoopScope(CodeGenFunction &CGF, const OMPLoopDirective &S)
+OMPLoopScope(CodeGenFunction &CGF, const OMPLoopBasedDirective &S)
 : CodeGenFunction::RunCleanupsScope(CGF) {
 emitPreInitStmt(CGF, S);
 }
 };
 for (const Expr *E : UDP->varlists()) {
 const Decl *D = cast<DeclRefExpr>(E)->getDecl();
 if (const auto *OED = dyn_cast<OMPCapturedExprDecl>(D))
 CGF.EmitVarDecl(*OED);
 }
+} else if (const auto *UDP = dyn_cast<OMPUseDeviceAddrClause>(C)) {
+for (const Expr *E : UDP->varlists()) {
+const Decl *D = getBaseDecl(E);
+if (const auto *OED = dyn_cast<OMPCapturedExprDecl>(D))
+CGF.EmitVarDecl(*OED);
+}
 }
 }
 if (!isOpenMPSimdDirective(S.getDirectiveKind()))
 CGF.EmitOMPPrivateClause(S, InlinedShareds);
 if (const auto *TG = dyn_cast<OMPTaskgroupDirective>(&S)) {
 if (const Expr *E = TG->getReductionRef())
 CGF.EmitVarDecl(*cast<VarDecl>(cast<DeclRefExpr>(E)->getDecl()));
 }
+// Temp copy arrays for inscan reductions should not be emitted as they are
+// not used in simd only mode.
+llvm::DenseSet<CanonicalDeclPtr<const Decl>> CopyArrayTemps;
+for (const auto *C : S.getClausesOfKind<OMPReductionClause>()) {
+if (C->getModifier() != OMPC_REDUCTION_inscan)
+continue;
+for (const Expr *E : C->copy_array_temps())
+CopyArrayTemps.insert(cast<DeclRefExpr>(E)->getDecl());
+}
 const auto *CS = cast_or_null<CapturedStmt>(S.getAssociatedStmt());
 while (CS) {
 for (auto &C : CS->captures()) {
 if (C.capturesVariable() || C.capturesVariableByCopy()) {
 auto *VD = C.getCapturedVar();
+if (CopyArrayTemps.contains(VD))
+continue;
 assert(VD == VD->getCanonicalDecl() &&
 "Canonical decl must be captured.");
 DeclRefExpr DRE(CGF.getContext(), const_cast<VarDecl *>(VD),
 isCapturedVar(CGF, VD) ||
 (CGF.CapturedStmtInfo &&
 VLASizes.clear();
 llvm::Function *WrapperF =
 emitOutlinedFunctionPrologue(WrapperCGF, Args, LocalAddrs, VLASizes,
 WrapperCGF.CXXThisValue, WrapperFO);
 llvm::SmallVector<llvm::Value *, 4> CallArgs;
+auto *PI = F->arg_begin();
 for (const auto *Arg : Args) {
 llvm::Value *CallArg;
 auto I = LocalAddrs.find(Arg);
 if (I != LocalAddrs.end()) {
 LValue LV = WrapperCGF.MakeAddrLValue(
 I->second.second,
 I->second.first ? I->second.first->getType() : Arg->getType(),
 AlignmentSource::Decl);
+if (LV.getType()->isAnyComplexType())
+LV.setAddress(WrapperCGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
+LV.getAddress(WrapperCGF),
+PI->getType()->getPointerTo(
+LV.getAddress(WrapperCGF).getAddressSpace())));
 CallArg = WrapperCGF.EmitLoadOfScalar(LV, S.getBeginLoc());
 } else {
 auto EI = VLASizes.find(Arg);
 if (EI != VLASizes.end()) {
 CallArg = EI->second.second;
 AlignmentSource::Decl);
 CallArg = WrapperCGF.EmitLoadOfScalar(LV, S.getBeginLoc());
 }
 }
 CallArgs.emplace_back(WrapperCGF.EmitFromMemory(CallArg, Arg->getType()));
+++PI;
 }
 CGM.getOpenMPRuntime().emitOutlinedFunctionCall(WrapperCGF, Loc, F, CallArgs);
 WrapperCGF.FinishFunction();
 return WrapperF;
 }
 if (CopiedVars.size() == 1) {
 // At first check if current thread is a master thread. If it is, no
 // need to copy data.
 CopyBegin = createBasicBlock("copyin.not.master");
 CopyEnd = createBasicBlock("copyin.not.master.end");
+// TODO: Avoid ptrtoint conversion.
+auto *MasterAddrInt =
+Builder.CreatePtrToInt(MasterAddr.getPointer(), CGM.IntPtrTy);
+auto *PrivateAddrInt =
+Builder.CreatePtrToInt(PrivateAddr.getPointer(), CGM.IntPtrTy);
 Builder.CreateCondBr(
-Builder.CreateICmpNE(
+Builder.CreateICmpNE(MasterAddrInt, PrivateAddrInt), CopyBegin,
-Builder.CreatePtrToInt(MasterAddr.getPointer(), CGM.IntPtrTy),
+CopyEnd);
-Builder.CreatePtrToInt(PrivateAddr.getPointer(),
-CGM.IntPtrTy)),
-CopyBegin, CopyEnd);
 EmitBlock(CopyBegin);
 }
 const auto *SrcVD =
 cast<VarDecl>(cast<DeclRefExpr>(*ISrcRef)->getDecl());
 const auto *DestVD =
 EmitBlock(DoneBB, /*IsFinished=*/true);
 }
 void CodeGenFunction::EmitOMPReductionClauseInit(
 const OMPExecutableDirective &D,
-CodeGenFunction::OMPPrivateScope &PrivateScope) {
+CodeGenFunction::OMPPrivateScope &PrivateScope, bool ForInscan) {
 if (!HaveInsertPoint())
 return;
 SmallVector<const Expr *, 4> Shareds;
 SmallVector<const Expr *, 4> Privates;
 SmallVector<const Expr *, 4> ReductionOps;
 SmallVector<const Expr *, 4> RHSs;
 OMPTaskDataTy Data;
 SmallVector<const Expr *, 4> TaskLHSs;
 SmallVector<const Expr *, 4> TaskRHSs;
 for (const auto *C : D.getClausesOfKind<OMPReductionClause>()) {
+if (ForInscan != (C->getModifier() == OMPC_REDUCTION_inscan))
+continue;
 Shareds.append(C->varlist_begin(), C->varlist_end());
 Privates.append(C->privates().begin(), C->privates().end());
 ReductionOps.append(C->reduction_ops().begin(), C->reduction_ops().end());
 LHSs.append(C->lhs_exprs().begin(), C->lhs_exprs().end());
 RHSs.append(C->rhs_exprs().begin(), C->rhs_exprs().end());
 case OMPD_requires:
 case OMPD_declare_variant:
 case OMPD_begin_declare_variant:
 case OMPD_end_declare_variant:
 case OMPD_unknown:
+default:
 llvm_unreachable("Enexpected directive with task reductions.");
 }
 const auto *VD = cast<VarDecl>(cast<DeclRefExpr>(TaskRedRef)->getDecl());
 EmitVarDecl(*VD);
 llvm::SmallVector<const Expr *, 8> RHSExprs;
 llvm::SmallVector<const Expr *, 8> ReductionOps;
 bool HasAtLeastOneReduction = false;
 bool IsReductionWithTaskMod = false;
 for (const auto *C : D.getClausesOfKind<OMPReductionClause>()) {
+// Do not emit for inscan reductions.
+if (C->getModifier() == OMPC_REDUCTION_inscan)
+continue;
 HasAtLeastOneReduction = true;
 Privates.append(C->privates().begin(), C->privates().end());
 LHSExprs.append(C->lhs_exprs().begin(), C->lhs_exprs().end());
 RHSExprs.append(C->rhs_exprs().begin(), C->rhs_exprs().end());
 ReductionOps.append(C->reduction_ops().begin(), C->reduction_ops().end());
 CGF.GenerateOpenMPCapturedVars(*CS, CapturedVars);
 CGF.CGM.getOpenMPRuntime().emitParallelCall(CGF, S.getBeginLoc(), OutlinedFn,
 CapturedVars, IfCond);
 }
+static bool isAllocatableDecl(const VarDecl *VD) {
+const VarDecl *CVD = VD->getCanonicalDecl();
+if (!CVD->hasAttr<OMPAllocateDeclAttr>())
+return false;
+const auto *AA = CVD->getAttr<OMPAllocateDeclAttr>();
+// Use the default allocation.
+return !((AA->getAllocatorType() == OMPAllocateDeclAttr::OMPDefaultMemAlloc ||
+AA->getAllocatorType() == OMPAllocateDeclAttr::OMPNullMemAlloc) &&
+!AA->getAllocator());
+}
 static void emitEmptyBoundParameters(CodeGenFunction &,
 const OMPExecutableDirective &,
 llvm::SmallVectorImpl<llvm::Value *> &) {}
+Address CodeGenFunction::OMPBuilderCBHelpers::getAddressOfLocalVariable(
+CodeGenFunction &CGF, const VarDecl *VD) {
+CodeGenModule &CGM = CGF.CGM;
+auto &OMPBuilder = CGM.getOpenMPRuntime().getOMPBuilder();
+if (!VD)
+return Address::invalid();
+const VarDecl *CVD = VD->getCanonicalDecl();
+if (!isAllocatableDecl(CVD))
+return Address::invalid();
+llvm::Value *Size;
+CharUnits Align = CGM.getContext().getDeclAlign(CVD);
+if (CVD->getType()->isVariablyModifiedType()) {
+Size = CGF.getTypeSize(CVD->getType());
+// Align the size: ((size + align - 1) / align) * align
+Size = CGF.Builder.CreateNUWAdd(
+Size, CGM.getSize(Align - CharUnits::fromQuantity(1)));
+Size = CGF.Builder.CreateUDiv(Size, CGM.getSize(Align));
+Size = CGF.Builder.CreateNUWMul(Size, CGM.getSize(Align));
+} else {
+CharUnits Sz = CGM.getContext().getTypeSizeInChars(CVD->getType());
+Size = CGM.getSize(Sz.alignTo(Align));
+}
+const auto *AA = CVD->getAttr<OMPAllocateDeclAttr>();
+assert(AA->getAllocator() &&
+"Expected allocator expression for non-default allocator.");
+llvm::Value *Allocator = CGF.EmitScalarExpr(AA->getAllocator());
+// According to the standard, the original allocator type is a enum (integer).
+// Convert to pointer type, if required.
+if (Allocator->getType()->isIntegerTy())
+Allocator = CGF.Builder.CreateIntToPtr(Allocator, CGM.VoidPtrTy);
+else if (Allocator->getType()->isPointerTy())
+Allocator = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(Allocator,
+CGM.VoidPtrTy);
+llvm::Value *Addr = OMPBuilder.createOMPAlloc(
+CGF.Builder, Size, Allocator,
+getNameWithSeparators({CVD->getName(), ".void.addr"}, ".", "."));
+llvm::CallInst *FreeCI =
+OMPBuilder.createOMPFree(CGF.Builder, Addr, Allocator);
+CGF.EHStack.pushCleanup<OMPAllocateCleanupTy>(NormalAndEHCleanup, FreeCI);
+Addr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
+Addr,
+CGF.ConvertTypeForMem(CGM.getContext().getPointerType(CVD->getType())),
+getNameWithSeparators({CVD->getName(), ".addr"}, ".", "."));
+return Address(Addr, Align);
+}
+Address CodeGenFunction::OMPBuilderCBHelpers::getAddrOfThreadPrivate(
+CodeGenFunction &CGF, const VarDecl *VD, Address VDAddr,
+SourceLocation Loc) {
+CodeGenModule &CGM = CGF.CGM;
+if (CGM.getLangOpts().OpenMPUseTLS &&
+CGM.getContext().getTargetInfo().isTLSSupported())
+return VDAddr;
+llvm::OpenMPIRBuilder &OMPBuilder = CGM.getOpenMPRuntime().getOMPBuilder();
+llvm::Type *VarTy = VDAddr.getElementType();
+llvm::Value *Data =
+CGF.Builder.CreatePointerCast(VDAddr.getPointer(), CGM.Int8PtrTy);
+llvm::ConstantInt *Size = CGM.getSize(CGM.GetTargetTypeStoreSize(VarTy));
+std::string Suffix = getNameWithSeparators({"cache", ""});
+llvm::Twine CacheName = Twine(CGM.getMangledName(VD)).concat(Suffix);
+llvm::CallInst *ThreadPrivateCacheCall =
+OMPBuilder.createCachedThreadPrivate(CGF.Builder, Data, Size, CacheName);
+return Address(ThreadPrivateCacheCall, VDAddr.getAlignment());
+}
+std::string CodeGenFunction::OMPBuilderCBHelpers::getNameWithSeparators(
+ArrayRef<StringRef> Parts, StringRef FirstSeparator, StringRef Separator) {
+SmallString<128> Buffer;
+llvm::raw_svector_ostream OS(Buffer);
+StringRef Sep = FirstSeparator;
+for (StringRef Part : Parts) {
+OS << Sep << Part;
+Sep = Separator;
+}
+return OS.str().str();
+}
 void CodeGenFunction::EmitOMPParallelDirective(const OMPParallelDirective &S) {
-if (llvm::OpenMPIRBuilder *OMPBuilder = CGM.getOpenMPIRBuilder()) {
+if (CGM.getLangOpts().OpenMPIRBuilder) {
+llvm::OpenMPIRBuilder &OMPBuilder = CGM.getOpenMPRuntime().getOMPBuilder();
 // Check if we have any if clause associated with the directive.
 llvm::Value *IfCond = nullptr;
 if (const auto *C = S.getSingleClause<OMPIfClause>())
 IfCond = EmitScalarExpr(C->getCondition(),
 /*IgnoreResultAssign=*/true);
 // Privatization callback that performs appropriate action for
 // shared/private/firstprivate/lastprivate/copyin/... variables.
 //
 // TODO: This defaults to shared right now.
 auto PrivCB = [](InsertPointTy AllocaIP, InsertPointTy CodeGenIP,
-llvm::Value &Val, llvm::Value *&ReplVal) {
+llvm::Value &, llvm::Value &Val, llvm::Value *&ReplVal) {
 // The next line is appropriate only for variables (Val) with the
 // data-sharing attribute "shared".
 ReplVal = &Val;
 return CodeGenIP;
 CodeGenIP, ContinuationBB);
 };
 CGCapturedStmtInfo CGSI(*CS, CR_OpenMP);
 CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(*this, &CGSI);
-Builder.restoreIP(OMPBuilder->CreateParallel(Builder, BodyGenCB, PrivCB,
+llvm::OpenMPIRBuilder::InsertPointTy AllocaIP(
-FiniCB, IfCond, NumThreads,
+AllocaInsertPt->getParent(), AllocaInsertPt->getIterator());
-ProcBind, S.hasCancel()));
+Builder.restoreIP(
+OMPBuilder.createParallel(Builder, AllocaIP, BodyGenCB, PrivCB, FiniCB,
+IfCond, NumThreads, ProcBind, S.hasCancel()));
 return;
 }
 // Emit parallel region as a standalone region.
 auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &Action) {
 }
 // Check for outer lastprivate conditional update.
 checkForLastprivateConditionalUpdate(*this, S);
 }
+namespace {
+/// RAII to handle scopes for loop transformation directives.
+class OMPTransformDirectiveScopeRAII {
+OMPLoopScope *Scope = nullptr;
+CodeGenFunction::CGCapturedStmtInfo *CGSI = nullptr;
+CodeGenFunction::CGCapturedStmtRAII *CapInfoRAII = nullptr;
+public:
+OMPTransformDirectiveScopeRAII(CodeGenFunction &CGF, const Stmt *S) {
+if (const auto *Dir = dyn_cast<OMPLoopBasedDirective>(S)) {
+Scope = new OMPLoopScope(CGF, *Dir);
+CGSI = new CodeGenFunction::CGCapturedStmtInfo(CR_OpenMP);
+CapInfoRAII = new CodeGenFunction::CGCapturedStmtRAII(CGF, CGSI);
+}
+}
+~OMPTransformDirectiveScopeRAII() {
+if (!Scope)
+return;
+delete CapInfoRAII;
+delete CGSI;
+delete Scope;
+}
+};
+} // namespace
 static void emitBody(CodeGenFunction &CGF, const Stmt *S, const Stmt *NextLoop,
 int MaxLevel, int Level = 0) {
 assert(Level < MaxLevel && "Too deep lookup during loop body codegen.");
 const Stmt *SimplifiedS = S->IgnoreContainers();
 if (const auto *CS = dyn_cast<CompoundStmt>(SimplifiedS)) {
 for (const Stmt *CurStmt : CS->body())
 emitBody(CGF, CurStmt, NextLoop, MaxLevel, Level);
 return;
 }
 if (SimplifiedS == NextLoop) {
+if (auto *Dir = dyn_cast<OMPTileDirective>(SimplifiedS))
+SimplifiedS = Dir->getTransformedStmt();
+if (const auto *CanonLoop = dyn_cast<OMPCanonicalLoop>(SimplifiedS))
+SimplifiedS = CanonLoop->getLoopStmt();
 if (const auto *For = dyn_cast<ForStmt>(SimplifiedS)) {
 S = For->getBody();
 } else {
 assert(isa<CXXForRangeStmt>(SimplifiedS) &&
 "Expected canonical for loop or range-based for loop.");
 llvm::BasicBlock *NextBB = createBasicBlock("omp.body.next");
 EmitBranchOnBoolExpr(E, NextBB, Continue.getBlock(),
 getProfileCount(D.getBody()));
 EmitBlock(NextBB);
 }
+OMPPrivateScope InscanScope(*this);
+EmitOMPReductionClauseInit(D, InscanScope, /*ForInscan=*/true);
+bool IsInscanRegion = InscanScope.Privatize();
+if (IsInscanRegion) {
+// Need to remember the block before and after scan directive
+// to dispatch them correctly depending on the clause used in
+// this directive, inclusive or exclusive. For inclusive scan the natural
+// order of the blocks is used, for exclusive clause the blocks must be
+// executed in reverse order.
+OMPBeforeScanBlock = createBasicBlock("omp.before.scan.bb");
+OMPAfterScanBlock = createBasicBlock("omp.after.scan.bb");
+// No need to allocate inscan exit block, in simd mode it is selected in the
+// codegen for the scan directive.
+if (D.getDirectiveKind() != OMPD_simd && !getLangOpts().OpenMPSimd)
+OMPScanExitBlock = createBasicBlock("omp.exit.inscan.bb");
+OMPScanDispatch = createBasicBlock("omp.inscan.dispatch");
+EmitBranch(OMPScanDispatch);
+EmitBlock(OMPBeforeScanBlock);
+}
 // Emit loop variables for C++ range loops.
 const Stmt *Body =
 D.getInnermostCapturedStmt()->getCapturedStmt()->IgnoreContainers();
 // Emit loop body.
 emitBody(*this, Body,
-OMPLoopDirective::tryToFindNextInnerLoop(
+OMPLoopBasedDirective::tryToFindNextInnerLoop(
 Body, /*TryImperfectlyNestedLoops=*/true),
-D.getCollapsedNumber());
+D.getLoopsNumber());
+// Jump to the dispatcher at the end of the loop body.
+if (IsInscanRegion)
+EmitBranch(OMPScanExitBlock);
 // The end (updates/cleanups).
 EmitBlock(Continue.getBlock());
 BreakContinueStack.pop_back();
 }
+using EmittedClosureTy = std::pair<llvm::Function *, llvm::Value *>;
+/// Emit a captured statement and return the function as well as its captured
+/// closure context.
+static EmittedClosureTy emitCapturedStmtFunc(CodeGenFunction &ParentCGF,
+const CapturedStmt *S) {
+LValue CapStruct = ParentCGF.InitCapturedStruct(*S);
+CodeGenFunction CGF(ParentCGF.CGM, /*suppressNewContext=*/true);
+std::unique_ptr<CodeGenFunction::CGCapturedStmtInfo> CSI =
+std::make_unique<CodeGenFunction::CGCapturedStmtInfo>(*S);
+CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, CSI.get());
+llvm::Function *F = CGF.GenerateCapturedStmtFunction(*S);
+return {F, CapStruct.getPointer(ParentCGF)};
+}
+/// Emit a call to a previously captured closure.
+static llvm::CallInst *
+emitCapturedStmtCall(CodeGenFunction &ParentCGF, EmittedClosureTy Cap,
+llvm::ArrayRef<llvm::Value *> Args) {
+// Append the closure context to the argument.
+SmallVector<llvm::Value *> EffectiveArgs;
+EffectiveArgs.reserve(Args.size() + 1);
+llvm::append_range(EffectiveArgs, Args);
+EffectiveArgs.push_back(Cap.second);
+return ParentCGF.Builder.CreateCall(Cap.first, EffectiveArgs);
+}
+llvm::CanonicalLoopInfo *
+CodeGenFunction::EmitOMPCollapsedCanonicalLoopNest(const Stmt *S, int Depth) {
+assert(Depth == 1 && "Nested loops with OpenMPIRBuilder not yet implemented");
+EmitStmt(S);
+assert(OMPLoopNestStack.size() >= (size_t)Depth && "Found too few loops");
+// The last added loop is the outermost one.
+return OMPLoopNestStack.back();
+}
+void CodeGenFunction::EmitOMPCanonicalLoop(const OMPCanonicalLoop *S) {
+const Stmt *SyntacticalLoop = S->getLoopStmt();
+if (!getLangOpts().OpenMPIRBuilder) {
+// Ignore if OpenMPIRBuilder is not enabled.
+EmitStmt(SyntacticalLoop);
+return;
+}
+LexicalScope ForScope(*this, S->getSourceRange());
+// Emit init statements. The Distance/LoopVar funcs may reference variable
+// declarations they contain.
+const Stmt *BodyStmt;
+if (const auto *For = dyn_cast<ForStmt>(SyntacticalLoop)) {
+if (const Stmt *InitStmt = For->getInit())
+EmitStmt(InitStmt);
+BodyStmt = For->getBody();
+} else if (const auto *RangeFor =
+dyn_cast<CXXForRangeStmt>(SyntacticalLoop)) {
+if (const DeclStmt *RangeStmt = RangeFor->getRangeStmt())
+EmitStmt(RangeStmt);
+if (const DeclStmt *BeginStmt = RangeFor->getBeginStmt())
+EmitStmt(BeginStmt);
+if (const DeclStmt *EndStmt = RangeFor->getEndStmt())
+EmitStmt(EndStmt);
+if (const DeclStmt *LoopVarStmt = RangeFor->getLoopVarStmt())
+EmitStmt(LoopVarStmt);
+BodyStmt = RangeFor->getBody();
+} else
+llvm_unreachable("Expected for-stmt or range-based for-stmt");
+// Emit closure for later use. By-value captures will be captured here.
+const CapturedStmt *DistanceFunc = S->getDistanceFunc();
+EmittedClosureTy DistanceClosure = emitCapturedStmtFunc(*this, DistanceFunc);
+const CapturedStmt *LoopVarFunc = S->getLoopVarFunc();
+EmittedClosureTy LoopVarClosure = emitCapturedStmtFunc(*this, LoopVarFunc);
+// Call the distance function to get the number of iterations of the loop to
+// come.
+QualType LogicalTy = DistanceFunc->getCapturedDecl()
+->getParam(0)
+->getType()
+.getNonReferenceType();
+Address CountAddr = CreateMemTemp(LogicalTy, ".count.addr");
+emitCapturedStmtCall(*this, DistanceClosure, {CountAddr.getPointer()});
+llvm::Value *DistVal = Builder.CreateLoad(CountAddr, ".count");
+// Emit the loop structure.
+llvm::OpenMPIRBuilder &OMPBuilder = CGM.getOpenMPRuntime().getOMPBuilder();
+auto BodyGen = [&, this](llvm::OpenMPIRBuilder::InsertPointTy CodeGenIP,
+llvm::Value *IndVar) {
+Builder.restoreIP(CodeGenIP);
+// Emit the loop body: Convert the logical iteration number to the loop
+// variable and emit the body.
+const DeclRefExpr *LoopVarRef = S->getLoopVarRef();
+LValue LCVal = EmitLValue(LoopVarRef);
+Address LoopVarAddress = LCVal.getAddress(*this);
+emitCapturedStmtCall(*this, LoopVarClosure,
+{LoopVarAddress.getPointer(), IndVar});
+RunCleanupsScope BodyScope(*this);
+EmitStmt(BodyStmt);
+};
+llvm::CanonicalLoopInfo *CL =
+OMPBuilder.createCanonicalLoop(Builder, BodyGen, DistVal);
+// Finish up the loop.
+Builder.restoreIP(CL->getAfterIP());
+ForScope.ForceCleanup();
+// Remember the CanonicalLoopInfo for parent AST nodes consuming it.
+OMPLoopNestStack.push_back(CL);
+}
 void CodeGenFunction::EmitOMPInnerLoop(
-const Stmt &S, bool RequiresCleanup, const Expr *LoopCond,
+const OMPExecutableDirective &S, bool RequiresCleanup, const Expr *LoopCond,
 const Expr *IncExpr,
 const llvm::function_ref<void(CodeGenFunction &)> BodyGen,
 const llvm::function_ref<void(CodeGenFunction &)> PostIncGen) {
 auto LoopExit = getJumpDestInCurrentScope("omp.inner.for.end");
 // If attributes are attached, push to the basic block with them.
 const auto &OMPED = cast<OMPExecutableDirective>(S);
 const CapturedStmt *ICS = OMPED.getInnermostCapturedStmt();
 const Stmt *SS = ICS->getCapturedStmt();
 const AttributedStmt *AS = dyn_cast_or_null<AttributedStmt>(SS);
+OMPLoopNestStack.clear();
 if (AS)
 LoopStack.push(CondBlock, CGM.getContext(), CGM.getCodeGenOpts(),
 AS->getAttrs(), SourceLocToDebugLoc(R.getBegin()),
 SourceLocToDebugLoc(R.getEnd()));
 else
 }
 // Privatize extra loop counters used in loops for ordered(n) clauses.
 for (const auto *C : S.getClausesOfKind<OMPOrderedClause>()) {
 if (!C->getNumForLoops())
 continue;
-for (unsigned I = S.getCollapsedNumber(),
+for (unsigned I = S.getLoopsNumber(), E = C->getLoopNumIterations().size();
-E = C->getLoopNumIterations().size();
 I < E; ++I) {
 const auto *DRE = cast<DeclRefExpr>(C->getLoopCounter(I));
 const auto *VD = cast<VarDecl>(DRE->getDecl());
 // Override only those variables that can be captured to avoid re-emission
 // of the variables declared within the loops.
 LoopStack.setVectorizeEnable();
 emitSimdlenSafelenClause(*this, D, IsMonotonic);
 if (const auto *C = D.getSingleClause<OMPOrderClause>())
 if (C->getKind() == OMPC_ORDER_concurrent)
 LoopStack.setParallel(/*Enable=*/true);
+if ((D.getDirectiveKind() == OMPD_simd ||
+(getLangOpts().OpenMPSimd &&
+isOpenMPSimdDirective(D.getDirectiveKind()))) &&
+llvm::any_of(D.getClausesOfKind<OMPReductionClause>(),
+[](const OMPReductionClause *C) {
+return C->getModifier() == OMPC_REDUCTION_inscan;
+}))
+// Disable parallel access in case of prefix sum.
+LoopStack.setParallel(/*Enable=*/false);
 }
 void CodeGenFunction::EmitOMPSimdFinal(
 const OMPLoopDirective &D,
 const llvm::function_ref<llvm::Value *(CodeGenFunction &)> CondGen) {
 },
 [&S, &LoopScope](CodeGenFunction &CGF, PrePostActionTy &) {
 CGF.EmitOMPInnerLoop(
 S, LoopScope.requiresCleanups(), S.getCond(), S.getInc(),
 [&S](CodeGenFunction &CGF) {
-CGF.EmitOMPLoopBody(S, CodeGenFunction::JumpDest());
+emitOMPLoopBodyWithStopPoint(CGF, S,
-CGF.EmitStopPoint(&S);
+CodeGenFunction::JumpDest());
 },
 [](CodeGenFunction &) {});
 });
 CGF.EmitOMPSimdFinal(S, [](CodeGenFunction &) { return nullptr; });
 // Emit final copy of the lastprivate variables at the end of loops.
 CGF.EmitBlock(ContBlock, true);
 }
 }
 void CodeGenFunction::EmitOMPSimdDirective(const OMPSimdDirective &S) {
+ParentLoopDirectiveForScanRegion ScanRegion(*this, S);
+OMPFirstScanLoop = true;
 auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &Action) {
 emitOMPSimdRegion(CGF, S, Action);
 };
 {
 auto LPCRegion =
 OMPLexicalScope Scope(*this, S, OMPD_unknown);
 CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_simd, CodeGen);
 }
 // Check for outer lastprivate conditional update.
 checkForLastprivateConditionalUpdate(*this, S);
+}
+void CodeGenFunction::EmitOMPTileDirective(const OMPTileDirective &S) {
+// Emit the de-sugared statement.
+OMPTransformDirectiveScopeRAII TileScope(*this, &S);
+EmitStmt(S.getTransformedStmt());
 }
 void CodeGenFunction::EmitOMPOuterLoop(
 bool DynamicOrOrdered, bool IsMonotonic, const OMPLoopDirective &S,
 CodeGenFunction::OMPPrivateScope &LoopScope,
 // Start the loop with a block that tests the condition.
 llvm::BasicBlock *CondBlock = createBasicBlock("omp.dispatch.cond");
 EmitBlock(CondBlock);
 const SourceRange R = S.getSourceRange();
+OMPLoopNestStack.clear();
 LoopStack.push(CondBlock, SourceLocToDebugLoc(R.getBegin()),
 SourceLocToDebugLoc(R.getEnd()));
 llvm::Value *BoolCondVal = nullptr;
 if (!DynamicOrOrdered) {
 EmitIgnoredExpr(LoopArgs.NextLB);
 EmitIgnoredExpr(LoopArgs.NextUB);
 }
 EmitBranch(CondBlock);
+OMPLoopNestStack.clear();
 LoopStack.pop();
 // Emit the fall-through block.
 EmitBlock(LoopExit.getBlock());
 // Tell the runtime we are done.
 // specified, and if no monotonic modifier is specified, the effect will
 // be as if the monotonic modifier was specified.
 bool StaticChunkedOne = RT.isStaticChunked(ScheduleKind.Schedule,
 /* Chunked */ Chunk != nullptr) && HasChunkSizeOne &&
 isOpenMPLoopBoundSharingDirective(S.getDirectiveKind());
+bool IsMonotonic =
+Ordered ||
+((ScheduleKind.Schedule == OMPC_SCHEDULE_static ||
+ScheduleKind.Schedule == OMPC_SCHEDULE_unknown) &&
+!(ScheduleKind.M1 == OMPC_SCHEDULE_MODIFIER_nonmonotonic ||
+ScheduleKind.M2 == OMPC_SCHEDULE_MODIFIER_nonmonotonic)) ||
+ScheduleKind.M1 == OMPC_SCHEDULE_MODIFIER_monotonic ||
+ScheduleKind.M2 == OMPC_SCHEDULE_MODIFIER_monotonic;
 if ((RT.isStaticNonchunked(ScheduleKind.Schedule,
 /* Chunked */ Chunk != nullptr) ||
 StaticChunkedOne) &&
 !Ordered) {
 JumpDest LoopExit =
 getJumpDestInCurrentScope(createBasicBlock("omp.loop.exit"));
 emitCommonSimdLoop(
 *this, S,
-[&S](CodeGenFunction &CGF, PrePostActionTy &) {
+[&S, IsMonotonic](CodeGenFunction &CGF, PrePostActionTy &) {
 if (isOpenMPSimdDirective(S.getDirectiveKind())) {
-CGF.EmitOMPSimdInit(S, /*IsMonotonic=*/true);
+CGF.EmitOMPSimdInit(S, IsMonotonic);
 } else if (const auto *C = S.getSingleClause<OMPOrderClause>()) {
 if (C->getKind() == OMPC_ORDER_concurrent)
 CGF.LoopStack.setParallel(/*Enable=*/true);
 }
 },
 S, LoopScope.requiresCleanups(),
 StaticChunkedOne ? S.getCombinedParForInDistCond()
 : S.getCond(),
 StaticChunkedOne ? S.getDistInc() : S.getInc(),
 [&S, LoopExit](CodeGenFunction &CGF) {
-CGF.EmitOMPLoopBody(S, LoopExit);
+emitOMPLoopBodyWithStopPoint(CGF, S, LoopExit);
-CGF.EmitStopPoint(&S);
 },
 [](CodeGenFunction &) {});
 });
 EmitBlock(LoopExit.getBlock());
 // Tell the runtime we are done.
 CGF.CGM.getOpenMPRuntime().emitForStaticFinish(CGF, S.getEndLoc(),
 S.getDirectiveKind());
 };
 OMPCancelStack.emitExit(*this, S.getDirectiveKind(), CodeGen);
 } else {
-const bool IsMonotonic =
-Ordered || ScheduleKind.Schedule == OMPC_SCHEDULE_static ||
-ScheduleKind.Schedule == OMPC_SCHEDULE_unknown ||
-ScheduleKind.M1 == OMPC_SCHEDULE_MODIFIER_monotonic ||
-ScheduleKind.M2 == OMPC_SCHEDULE_MODIFIER_monotonic;
 // Emit the outer loop, which requests its work chunk [LB..UB] from
 // runtime and runs the inner loop to process it.
 const OMPLoopArguments LoopArguments(
 LB.getAddress(*this), UB.getAddress(*this), ST.getAddress(*this),
 IL.getAddress(*this), Chunk, EUB);
 llvm::Value *LBVal = CGF.Builder.getIntN(IVSize, 0);
 llvm::Value *UBVal = CGF.EmitScalarExpr(LS.getLastIteration());
 return {LBVal, UBVal};
 }
-void CodeGenFunction::EmitOMPForDirective(const OMPForDirective &S) {
+/// Emits internal temp array declarations for the directive with inscan
-bool HasLastprivates = false;
+/// reductions.
-auto &&CodeGen = [&S, &HasLastprivates](CodeGenFunction &CGF,
+/// The code is the following:
-PrePostActionTy &) {
+/// \code
-OMPCancelStackRAII CancelRegion(CGF, OMPD_for, S.hasCancel());
+/// size num_iters = <num_iters>;
+/// <type> buffer[num_iters];
+/// \endcode
+static void emitScanBasedDirectiveDecls(
+CodeGenFunction &CGF, const OMPLoopDirective &S,
+llvm::function_ref<llvm::Value *(CodeGenFunction &)> NumIteratorsGen) {
+llvm::Value *OMPScanNumIterations = CGF.Builder.CreateIntCast(
+NumIteratorsGen(CGF), CGF.SizeTy, /*isSigned=*/false);
+SmallVector<const Expr *, 4> Shareds;
+SmallVector<const Expr *, 4> Privates;
+SmallVector<const Expr *, 4> ReductionOps;
+SmallVector<const Expr *, 4> CopyArrayTemps;
+for (const auto *C : S.getClausesOfKind<OMPReductionClause>()) {
+assert(C->getModifier() == OMPC_REDUCTION_inscan &&
+"Only inscan reductions are expected.");
+Shareds.append(C->varlist_begin(), C->varlist_end());
+Privates.append(C->privates().begin(), C->privates().end());
+ReductionOps.append(C->reduction_ops().begin(), C->reduction_ops().end());
+CopyArrayTemps.append(C->copy_array_temps().begin(),
+C->copy_array_temps().end());
+}
+{
+// Emit buffers for each reduction variables.
+// ReductionCodeGen is required to emit correctly the code for array
+// reductions.
+ReductionCodeGen RedCG(Shareds, Shareds, Privates, ReductionOps);
+unsigned Count = 0;
+auto *ITA = CopyArrayTemps.begin();
+for (const Expr *IRef : Privates) {
+const auto *PrivateVD = cast<VarDecl>(cast<DeclRefExpr>(IRef)->getDecl());
+// Emit variably modified arrays, used for arrays/array sections
+// reductions.
+if (PrivateVD->getType()->isVariablyModifiedType()) {
+RedCG.emitSharedOrigLValue(CGF, Count);
+RedCG.emitAggregateType(CGF, Count);
+}
+CodeGenFunction::OpaqueValueMapping DimMapping(
+CGF,
+cast<OpaqueValueExpr>(
+cast<VariableArrayType>((*ITA)->getType()->getAsArrayTypeUnsafe())
+->getSizeExpr()),
+RValue::get(OMPScanNumIterations));
+// Emit temp buffer.
+CGF.EmitVarDecl(*cast<VarDecl>(cast<DeclRefExpr>(*ITA)->getDecl()));
+++ITA;
+++Count;
+}
+}
+}
+/// Emits the code for the directive with inscan reductions.
+/// The code is the following:
+/// \code
+/// #pragma omp ...
+/// for (i: 0..<num_iters>) {
+///   <input phase>;
+///   buffer[i] = red;
+/// }
+/// #pragma omp master // in parallel region
+/// for (int k = 0; k != ceil(log2(num_iters)); ++k)
+/// for (size cnt = last_iter; cnt >= pow(2, k); --k)
+///   buffer[i] op= buffer[i-pow(2,k)];
+/// #pragma omp barrier // in parallel region
+/// #pragma omp ...
+/// for (0..<num_iters>) {
+///   red = InclusiveScan ? buffer[i] : buffer[i-1];
+///   <scan phase>;
+/// }
+/// \endcode
+static void emitScanBasedDirective(
+CodeGenFunction &CGF, const OMPLoopDirective &S,
+llvm::function_ref<llvm::Value *(CodeGenFunction &)> NumIteratorsGen,
+llvm::function_ref<void(CodeGenFunction &)> FirstGen,
+llvm::function_ref<void(CodeGenFunction &)> SecondGen) {
+llvm::Value *OMPScanNumIterations = CGF.Builder.CreateIntCast(
+NumIteratorsGen(CGF), CGF.SizeTy, /*isSigned=*/false);
+SmallVector<const Expr *, 4> Privates;
+SmallVector<const Expr *, 4> ReductionOps;
+SmallVector<const Expr *, 4> LHSs;
+SmallVector<const Expr *, 4> RHSs;
+SmallVector<const Expr *, 4> CopyArrayElems;
+for (const auto *C : S.getClausesOfKind<OMPReductionClause>()) {
+assert(C->getModifier() == OMPC_REDUCTION_inscan &&
+"Only inscan reductions are expected.");
+Privates.append(C->privates().begin(), C->privates().end());
+ReductionOps.append(C->reduction_ops().begin(), C->reduction_ops().end());
+LHSs.append(C->lhs_exprs().begin(), C->lhs_exprs().end());
+RHSs.append(C->rhs_exprs().begin(), C->rhs_exprs().end());
+CopyArrayElems.append(C->copy_array_elems().begin(),
+C->copy_array_elems().end());
+}
+CodeGenFunction::ParentLoopDirectiveForScanRegion ScanRegion(CGF, S);
+{
+// Emit loop with input phase:
+// #pragma omp ...
+// for (i: 0..<num_iters>) {
+//   <input phase>;
+//   buffer[i] = red;
+// }
+CGF.OMPFirstScanLoop = true;
+CodeGenFunction::OMPLocalDeclMapRAII Scope(CGF);
+FirstGen(CGF);
+}
+// #pragma omp barrier // in parallel region
+auto &&CodeGen = [&S, OMPScanNumIterations, &LHSs, &RHSs, &CopyArrayElems,
+&ReductionOps,
+&Privates](CodeGenFunction &CGF, PrePostActionTy &Action) {
+Action.Enter(CGF);
+// Emit prefix reduction:
+// #pragma omp master // in parallel region
+// for (int k = 0; k <= ceil(log2(n)); ++k)
+llvm::BasicBlock *InputBB = CGF.Builder.GetInsertBlock();
+llvm::BasicBlock *LoopBB = CGF.createBasicBlock("omp.outer.log.scan.body");
+llvm::BasicBlock *ExitBB = CGF.createBasicBlock("omp.outer.log.scan.exit");
+llvm::Function *F =
+CGF.CGM.getIntrinsic(llvm::Intrinsic::log2, CGF.DoubleTy);
+llvm::Value *Arg =
+CGF.Builder.CreateUIToFP(OMPScanNumIterations, CGF.DoubleTy);
+llvm::Value *LogVal = CGF.EmitNounwindRuntimeCall(F, Arg);
+F = CGF.CGM.getIntrinsic(llvm::Intrinsic::ceil, CGF.DoubleTy);
+LogVal = CGF.EmitNounwindRuntimeCall(F, LogVal);
+LogVal = CGF.Builder.CreateFPToUI(LogVal, CGF.IntTy);
+llvm::Value *NMin1 = CGF.Builder.CreateNUWSub(
+OMPScanNumIterations, llvm::ConstantInt::get(CGF.SizeTy, 1));
+auto DL = ApplyDebugLocation::CreateDefaultArtificial(CGF, S.getBeginLoc());
+CGF.EmitBlock(LoopBB);
+auto *Counter = CGF.Builder.CreatePHI(CGF.IntTy, 2);
+// size pow2k = 1;
+auto *Pow2K = CGF.Builder.CreatePHI(CGF.SizeTy, 2);
+Counter->addIncoming(llvm::ConstantInt::get(CGF.IntTy, 0), InputBB);
+Pow2K->addIncoming(llvm::ConstantInt::get(CGF.SizeTy, 1), InputBB);
+// for (size i = n - 1; i >= 2 ^ k; --i)
+//   tmp[i] op= tmp[i-pow2k];
+llvm::BasicBlock *InnerLoopBB =
+CGF.createBasicBlock("omp.inner.log.scan.body");
+llvm::BasicBlock *InnerExitBB =
+CGF.createBasicBlock("omp.inner.log.scan.exit");
+llvm::Value *CmpI = CGF.Builder.CreateICmpUGE(NMin1, Pow2K);
+CGF.Builder.CreateCondBr(CmpI, InnerLoopBB, InnerExitBB);
+CGF.EmitBlock(InnerLoopBB);
+auto *IVal = CGF.Builder.CreatePHI(CGF.SizeTy, 2);
+IVal->addIncoming(NMin1, LoopBB);
+{
+CodeGenFunction::OMPPrivateScope PrivScope(CGF);
+auto *ILHS = LHSs.begin();
+auto *IRHS = RHSs.begin();
+for (const Expr *CopyArrayElem : CopyArrayElems) {
+const auto *LHSVD = cast<VarDecl>(cast<DeclRefExpr>(*ILHS)->getDecl());
+const auto *RHSVD = cast<VarDecl>(cast<DeclRefExpr>(*IRHS)->getDecl());
+Address LHSAddr = Address::invalid();
+{
+CodeGenFunction::OpaqueValueMapping IdxMapping(
+CGF,
+cast<OpaqueValueExpr>(
+cast<ArraySubscriptExpr>(CopyArrayElem)->getIdx()),
+RValue::get(IVal));
+LHSAddr = CGF.EmitLValue(CopyArrayElem).getAddress(CGF);
+}
+PrivScope.addPrivate(LHSVD, [LHSAddr]() { return LHSAddr; });
+Address RHSAddr = Address::invalid();
+{
+llvm::Value *OffsetIVal = CGF.Builder.CreateNUWSub(IVal, Pow2K);
+CodeGenFunction::OpaqueValueMapping IdxMapping(
+CGF,
+cast<OpaqueValueExpr>(
+cast<ArraySubscriptExpr>(CopyArrayElem)->getIdx()),
+RValue::get(OffsetIVal));
+RHSAddr = CGF.EmitLValue(CopyArrayElem).getAddress(CGF);
+}
+PrivScope.addPrivate(RHSVD, [RHSAddr]() { return RHSAddr; });
+++ILHS;
+++IRHS;
+}
+PrivScope.Privatize();
+CGF.CGM.getOpenMPRuntime().emitReduction(
+CGF, S.getEndLoc(), Privates, LHSs, RHSs, ReductionOps,
+{/*WithNowait=*/true, /*SimpleReduction=*/true, OMPD_unknown});
+}
+llvm::Value *NextIVal =
+CGF.Builder.CreateNUWSub(IVal, llvm::ConstantInt::get(CGF.SizeTy, 1));
+IVal->addIncoming(NextIVal, CGF.Builder.GetInsertBlock());
+CmpI = CGF.Builder.CreateICmpUGE(NextIVal, Pow2K);
+CGF.Builder.CreateCondBr(CmpI, InnerLoopBB, InnerExitBB);
+CGF.EmitBlock(InnerExitBB);
+llvm::Value *Next =
+CGF.Builder.CreateNUWAdd(Counter, llvm::ConstantInt::get(CGF.IntTy, 1));
+Counter->addIncoming(Next, CGF.Builder.GetInsertBlock());
+// pow2k <<= 1;
+llvm::Value *NextPow2K =
+CGF.Builder.CreateShl(Pow2K, 1, "", /*HasNUW=*/true);
+Pow2K->addIncoming(NextPow2K, CGF.Builder.GetInsertBlock());
+llvm::Value *Cmp = CGF.Builder.CreateICmpNE(Next, LogVal);
+CGF.Builder.CreateCondBr(Cmp, LoopBB, ExitBB);
+auto DL1 = ApplyDebugLocation::CreateDefaultArtificial(CGF, S.getEndLoc());
+CGF.EmitBlock(ExitBB);
+};
+if (isOpenMPParallelDirective(S.getDirectiveKind())) {
+CGF.CGM.getOpenMPRuntime().emitMasterRegion(CGF, CodeGen, S.getBeginLoc());
+CGF.CGM.getOpenMPRuntime().emitBarrierCall(
+CGF, S.getBeginLoc(), OMPD_unknown, /*EmitChecks=*/false,
+/*ForceSimpleCall=*/true);
+} else {
+RegionCodeGenTy RCG(CodeGen);
+RCG(CGF);
+}
+CGF.OMPFirstScanLoop = false;
+SecondGen(CGF);
+}
+static bool emitWorksharingDirective(CodeGenFunction &CGF,
+const OMPLoopDirective &S,
+bool HasCancel) {
+bool HasLastprivates;
+if (llvm::any_of(S.getClausesOfKind<OMPReductionClause>(),
+[](const OMPReductionClause *C) {
+return C->getModifier() == OMPC_REDUCTION_inscan;
+})) {
+const auto &&NumIteratorsGen = [&S](CodeGenFunction &CGF) {
+CodeGenFunction::OMPLocalDeclMapRAII Scope(CGF);
+OMPLoopScope LoopScope(CGF, S);
+return CGF.EmitScalarExpr(S.getNumIterations());
+};
+const auto &&FirstGen = [&S, HasCancel](CodeGenFunction &CGF) {
+CodeGenFunction::OMPCancelStackRAII CancelRegion(
+CGF, S.getDirectiveKind(), HasCancel);
+(void)CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(),
+emitForLoopBounds,
+emitDispatchForLoopBounds);
+// Emit an implicit barrier at the end.
+CGF.CGM.getOpenMPRuntime().emitBarrierCall(CGF, S.getBeginLoc(),
+OMPD_for);
+};
+const auto &&SecondGen = [&S, HasCancel,
+&HasLastprivates](CodeGenFunction &CGF) {
+CodeGenFunction::OMPCancelStackRAII CancelRegion(
+CGF, S.getDirectiveKind(), HasCancel);
+HasLastprivates = CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(),
+emitForLoopBounds,
+emitDispatchForLoopBounds);
+};
+if (!isOpenMPParallelDirective(S.getDirectiveKind()))
+emitScanBasedDirectiveDecls(CGF, S, NumIteratorsGen);
+emitScanBasedDirective(CGF, S, NumIteratorsGen, FirstGen, SecondGen);
+} else {
+CodeGenFunction::OMPCancelStackRAII CancelRegion(CGF, S.getDirectiveKind(),
+HasCancel);
 HasLastprivates = CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(),
 emitForLoopBounds,
 emitDispatchForLoopBounds);
+}
+return HasLastprivates;
+}
+static bool isSupportedByOpenMPIRBuilder(const OMPForDirective &S) {
+if (S.hasCancel())
+return false;
+for (OMPClause *C : S.clauses())
+if (!isa<OMPNowaitClause>(C))
+return false;
+return true;
+}
+void CodeGenFunction::EmitOMPForDirective(const OMPForDirective &S) {
+bool HasLastprivates = false;
+bool UseOMPIRBuilder =
+CGM.getLangOpts().OpenMPIRBuilder && isSupportedByOpenMPIRBuilder(S);
+auto &&CodeGen = [this, &S, &HasLastprivates,
+UseOMPIRBuilder](CodeGenFunction &CGF, PrePostActionTy &) {
+// Use the OpenMPIRBuilder if enabled.
+if (UseOMPIRBuilder) {
+// Emit the associated statement and get its loop representation.
+const Stmt *Inner = S.getRawStmt();
+llvm::CanonicalLoopInfo *CLI =
+EmitOMPCollapsedCanonicalLoopNest(Inner, 1);
+bool NeedsBarrier = !S.getSingleClause<OMPNowaitClause>();
+llvm::OpenMPIRBuilder &OMPBuilder =
+CGM.getOpenMPRuntime().getOMPBuilder();
+llvm::OpenMPIRBuilder::InsertPointTy AllocaIP(
+AllocaInsertPt->getParent(), AllocaInsertPt->getIterator());
+OMPBuilder.createWorkshareLoop(Builder, CLI, AllocaIP, NeedsBarrier);
+return;
+}
+HasLastprivates = emitWorksharingDirective(CGF, S, S.hasCancel());
 };
 {
 auto LPCRegion =
 CGOpenMPRuntime::LastprivateConditionalRAII::disable(*this, S);
 OMPLexicalScope Scope(*this, S, OMPD_unknown);
 CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_for, CodeGen,
 S.hasCancel());
 }
-// Emit an implicit barrier at the end.
+if (!UseOMPIRBuilder) {
-if (!S.getSingleClause<OMPNowaitClause>() || HasLastprivates)
+// Emit an implicit barrier at the end.
-CGM.getOpenMPRuntime().emitBarrierCall(*this, S.getBeginLoc(), OMPD_for);
+if (!S.getSingleClause<OMPNowaitClause>() || HasLastprivates)
+CGM.getOpenMPRuntime().emitBarrierCall(*this, S.getBeginLoc(), OMPD_for);
+}
 // Check for outer lastprivate conditional update.
 checkForLastprivateConditionalUpdate(*this, S);
 }
 void CodeGenFunction::EmitOMPForSimdDirective(const OMPForSimdDirective &S) {
 bool HasLastprivates = false;
 auto &&CodeGen = [&S, &HasLastprivates](CodeGenFunction &CGF,
 PrePostActionTy &) {
-HasLastprivates = CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(),
+HasLastprivates = emitWorksharingDirective(CGF, S, /*HasCancel=*/false);
-emitForLoopBounds,
-emitDispatchForLoopBounds);
 };
 {
 auto LPCRegion =
 CGOpenMPRuntime::LastprivateConditionalRAII::disable(*this, S);
 OMPLexicalScope Scope(*this, S, OMPD_unknown);
 OpaqueValueExpr UBRefExpr(S.getBeginLoc(), KmpInt32Ty, VK_LValue);
 CodeGenFunction::OpaqueValueMapping OpaqueUB(CGF, &UBRefExpr, UB);
 // Generate condition for loop.
 BinaryOperator *Cond = BinaryOperator::Create(
 C, &IVRefExpr, &UBRefExpr, BO_LE, C.BoolTy, VK_RValue, OK_Ordinary,
-S.getBeginLoc(), FPOptions(C.getLangOpts()));
+S.getBeginLoc(), FPOptionsOverride());
 // Increment for loop counter.
 UnaryOperator *Inc = UnaryOperator::Create(
 C, &IVRefExpr, UO_PreInc, KmpInt32Ty, VK_RValue, OK_Ordinary,
-S.getBeginLoc(), true, FPOptions(C.getLangOpts()));
+S.getBeginLoc(), true, FPOptionsOverride());
 auto &&BodyGen = [CapturedStmt, CS, &S, &IV](CodeGenFunction &CGF) {
 // Iterate through all sections and emit a switch construct:
 // switch (IV) {
 //   case 0:
 //     <SectionStmt[0]>;
 OMPD_unknown);
 }
 }
 void CodeGenFunction::EmitOMPSectionsDirective(const OMPSectionsDirective &S) {
+if (CGM.getLangOpts().OpenMPIRBuilder) {
+llvm::OpenMPIRBuilder &OMPBuilder = CGM.getOpenMPRuntime().getOMPBuilder();
+using InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy;
+using BodyGenCallbackTy = llvm::OpenMPIRBuilder::StorableBodyGenCallbackTy;
+auto FiniCB = [this](InsertPointTy IP) {
+OMPBuilderCBHelpers::FinalizeOMPRegion(*this, IP);
+};
+const CapturedStmt *ICS = S.getInnermostCapturedStmt();
+const Stmt *CapturedStmt = S.getInnermostCapturedStmt()->getCapturedStmt();
+const auto *CS = dyn_cast<CompoundStmt>(CapturedStmt);
+llvm::SmallVector<BodyGenCallbackTy, 4> SectionCBVector;
+if (CS) {
+for (const Stmt *SubStmt : CS->children()) {
+auto SectionCB = [this, SubStmt](InsertPointTy AllocaIP,
+InsertPointTy CodeGenIP,
+llvm::BasicBlock &FiniBB) {
+OMPBuilderCBHelpers::InlinedRegionBodyRAII IRB(*this, AllocaIP,
+FiniBB);
+OMPBuilderCBHelpers::EmitOMPRegionBody(*this, SubStmt, CodeGenIP,
+FiniBB);
+};
+SectionCBVector.push_back(SectionCB);
+}
+} else {
+auto SectionCB = [this, CapturedStmt](InsertPointTy AllocaIP,
+InsertPointTy CodeGenIP,
+llvm::BasicBlock &FiniBB) {
+OMPBuilderCBHelpers::InlinedRegionBodyRAII IRB(*this, AllocaIP, FiniBB);
+OMPBuilderCBHelpers::EmitOMPRegionBody(*this, CapturedStmt, CodeGenIP,
+FiniBB);
+};
+SectionCBVector.push_back(SectionCB);
+}
+// Privatization callback that performs appropriate action for
+// shared/private/firstprivate/lastprivate/copyin/... variables.
+//
+// TODO: This defaults to shared right now.
+auto PrivCB = [](InsertPointTy AllocaIP, InsertPointTy CodeGenIP,
+llvm::Value &, llvm::Value &Val, llvm::Value *&ReplVal) {
+// The next line is appropriate only for variables (Val) with the
+// data-sharing attribute "shared".
+ReplVal = &Val;
+return CodeGenIP;
+};
+CGCapturedStmtInfo CGSI(*ICS, CR_OpenMP);
+CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(*this, &CGSI);
+llvm::OpenMPIRBuilder::InsertPointTy AllocaIP(
+AllocaInsertPt->getParent(), AllocaInsertPt->getIterator());
+Builder.restoreIP(OMPBuilder.createSections(
+Builder, AllocaIP, SectionCBVector, PrivCB, FiniCB, S.hasCancel(),
+S.getSingleClause<OMPNowaitClause>()));
+return;
+}
 {
 auto LPCRegion =
 CGOpenMPRuntime::LastprivateConditionalRAII::disable(*this, S);
 OMPLexicalScope Scope(*this, S, OMPD_unknown);
 EmitSections(S);
 // Check for outer lastprivate conditional update.
 checkForLastprivateConditionalUpdate(*this, S);
 }
 void CodeGenFunction::EmitOMPSectionDirective(const OMPSectionDirective &S) {
-auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &) {
+if (CGM.getLangOpts().OpenMPIRBuilder) {
-CGF.EmitStmt(S.getInnermostCapturedStmt()->getCapturedStmt());
+llvm::OpenMPIRBuilder &OMPBuilder = CGM.getOpenMPRuntime().getOMPBuilder();
-};
+using InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy;
-OMPLexicalScope Scope(*this, S, OMPD_unknown);
-CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_section, CodeGen,
+const Stmt *SectionRegionBodyStmt = S.getAssociatedStmt();
-S.hasCancel());
+auto FiniCB = [this](InsertPointTy IP) {
+OMPBuilderCBHelpers::FinalizeOMPRegion(*this, IP);
+};
+auto BodyGenCB = [SectionRegionBodyStmt, this](InsertPointTy AllocaIP,
+InsertPointTy CodeGenIP,
+llvm::BasicBlock &FiniBB) {
+OMPBuilderCBHelpers::InlinedRegionBodyRAII IRB(*this, AllocaIP, FiniBB);
+OMPBuilderCBHelpers::EmitOMPRegionBody(*this, SectionRegionBodyStmt,
+CodeGenIP, FiniBB);
+};
+LexicalScope Scope(*this, S.getSourceRange());
+EmitStopPoint(&S);
+Builder.restoreIP(OMPBuilder.createSection(Builder, BodyGenCB, FiniCB));
+return;
+}
+LexicalScope Scope(*this, S.getSourceRange());
+EmitStopPoint(&S);
+EmitStmt(S.getAssociatedStmt());
 }
 void CodeGenFunction::EmitOMPSingleDirective(const OMPSingleDirective &S) {
 llvm::SmallVector<const Expr *, 8> CopyprivateVars;
 llvm::SmallVector<const Expr *, 8> DestExprs;
 }
 static void emitMaster(CodeGenFunction &CGF, const OMPExecutableDirective &S) {
 auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &Action) {
 Action.Enter(CGF);
-CGF.EmitStmt(S.getInnermostCapturedStmt()->getCapturedStmt());
+CGF.EmitStmt(S.getRawStmt());
 };
 CGF.CGM.getOpenMPRuntime().emitMasterRegion(CGF, CodeGen, S.getBeginLoc());
 }
 void CodeGenFunction::EmitOMPMasterDirective(const OMPMasterDirective &S) {
-if (llvm::OpenMPIRBuilder *OMPBuilder = CGM.getOpenMPIRBuilder()) {
+if (CGM.getLangOpts().OpenMPIRBuilder) {
+llvm::OpenMPIRBuilder &OMPBuilder = CGM.getOpenMPRuntime().getOMPBuilder();
 using InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy;
-const CapturedStmt *CS = S.getInnermostCapturedStmt();
+const Stmt *MasterRegionBodyStmt = S.getAssociatedStmt();
-const Stmt *MasterRegionBodyStmt = CS->getCapturedStmt();
 auto FiniCB = [this](InsertPointTy IP) {
 OMPBuilderCBHelpers::FinalizeOMPRegion(*this, IP);
 };
 OMPBuilderCBHelpers::InlinedRegionBodyRAII IRB(*this, AllocaIP, FiniBB);
 OMPBuilderCBHelpers::EmitOMPRegionBody(*this, MasterRegionBodyStmt,
 CodeGenIP, FiniBB);
 };
-CGCapturedStmtInfo CGSI(*CS, CR_OpenMP);
+LexicalScope Scope(*this, S.getSourceRange());
-CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(*this, &CGSI);
+EmitStopPoint(&S);
-Builder.restoreIP(OMPBuilder->CreateMaster(Builder, BodyGenCB, FiniCB));
+Builder.restoreIP(OMPBuilder.createMaster(Builder, BodyGenCB, FiniCB));
 return;
 }
-OMPLexicalScope Scope(*this, S, OMPD_unknown);
+LexicalScope Scope(*this, S.getSourceRange());
+EmitStopPoint(&S);
 emitMaster(*this, S);
 }
+static void emitMasked(CodeGenFunction &CGF, const OMPExecutableDirective &S) {
+auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &Action) {
+Action.Enter(CGF);
+CGF.EmitStmt(S.getRawStmt());
+};
+Expr *Filter = nullptr;
+if (const auto *FilterClause = S.getSingleClause<OMPFilterClause>())
+Filter = FilterClause->getThreadID();
+CGF.CGM.getOpenMPRuntime().emitMaskedRegion(CGF, CodeGen, S.getBeginLoc(),
+Filter);
+}
+void CodeGenFunction::EmitOMPMaskedDirective(const OMPMaskedDirective &S) {
+if (CGM.getLangOpts().OpenMPIRBuilder) {
+llvm::OpenMPIRBuilder &OMPBuilder = CGM.getOpenMPRuntime().getOMPBuilder();
+using InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy;
+const Stmt *MaskedRegionBodyStmt = S.getAssociatedStmt();
+const Expr *Filter = nullptr;
+if (const auto *FilterClause = S.getSingleClause<OMPFilterClause>())
+Filter = FilterClause->getThreadID();
+llvm::Value *FilterVal = Filter
+? EmitScalarExpr(Filter, CGM.Int32Ty)
+: llvm::ConstantInt::get(CGM.Int32Ty, /*V=*/0);
+auto FiniCB = [this](InsertPointTy IP) {
+OMPBuilderCBHelpers::FinalizeOMPRegion(*this, IP);
+};
+auto BodyGenCB = [MaskedRegionBodyStmt, this](InsertPointTy AllocaIP,
+InsertPointTy CodeGenIP,
+llvm::BasicBlock &FiniBB) {
+OMPBuilderCBHelpers::InlinedRegionBodyRAII IRB(*this, AllocaIP, FiniBB);
+OMPBuilderCBHelpers::EmitOMPRegionBody(*this, MaskedRegionBodyStmt,
+CodeGenIP, FiniBB);
+};
+LexicalScope Scope(*this, S.getSourceRange());
+EmitStopPoint(&S);
+Builder.restoreIP(
+OMPBuilder.createMasked(Builder, BodyGenCB, FiniCB, FilterVal));
+return;
+}
+LexicalScope Scope(*this, S.getSourceRange());
+EmitStopPoint(&S);
+emitMasked(*this, S);
+}
 void CodeGenFunction::EmitOMPCriticalDirective(const OMPCriticalDirective &S) {
-if (llvm::OpenMPIRBuilder *OMPBuilder = CGM.getOpenMPIRBuilder()) {
+if (CGM.getLangOpts().OpenMPIRBuilder) {
+llvm::OpenMPIRBuilder &OMPBuilder = CGM.getOpenMPRuntime().getOMPBuilder();
 using InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy;
-const CapturedStmt *CS = S.getInnermostCapturedStmt();
+const Stmt *CriticalRegionBodyStmt = S.getAssociatedStmt();
-const Stmt *CriticalRegionBodyStmt = CS->getCapturedStmt();
 const Expr *Hint = nullptr;
 if (const auto *HintClause = S.getSingleClause<OMPHintClause>())
 Hint = HintClause->getHint();
 // TODO: This is slightly different from what's currently being done in
 OMPBuilderCBHelpers::InlinedRegionBodyRAII IRB(*this, AllocaIP, FiniBB);
 OMPBuilderCBHelpers::EmitOMPRegionBody(*this, CriticalRegionBodyStmt,
 CodeGenIP, FiniBB);
 };
-CGCapturedStmtInfo CGSI(*CS, CR_OpenMP);
+LexicalScope Scope(*this, S.getSourceRange());
-CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(*this, &CGSI);
+EmitStopPoint(&S);
-Builder.restoreIP(OMPBuilder->CreateCritical(
+Builder.restoreIP(OMPBuilder.createCritical(
 Builder, BodyGenCB, FiniCB, S.getDirectiveName().getAsString(),
 HintInst));
 return;
 }
 auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &Action) {
 Action.Enter(CGF);
-CGF.EmitStmt(S.getInnermostCapturedStmt()->getCapturedStmt());
+CGF.EmitStmt(S.getAssociatedStmt());
 };
 const Expr *Hint = nullptr;
 if (const auto *HintClause = S.getSingleClause<OMPHintClause>())
 Hint = HintClause->getHint();
-OMPLexicalScope Scope(*this, S, OMPD_unknown);
+LexicalScope Scope(*this, S.getSourceRange());
+EmitStopPoint(&S);
 CGM.getOpenMPRuntime().emitCriticalRegion(*this,
 S.getDirectiveName().getAsString(),
 CodeGen, S.getBeginLoc(), Hint);
 }
 const OMPParallelForDirective &S) {
 // Emit directive as a combined directive that consists of two implicit
 // directives: 'parallel' with 'for' directive.
 auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &Action) {
 Action.Enter(CGF);
-OMPCancelStackRAII CancelRegion(CGF, OMPD_parallel_for, S.hasCancel());
+(void)emitWorksharingDirective(CGF, S, S.hasCancel());
-CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(), emitForLoopBounds,
-emitDispatchForLoopBounds);
 };
 {
+if (llvm::any_of(S.getClausesOfKind<OMPReductionClause>(),
+[](const OMPReductionClause *C) {
+return C->getModifier() == OMPC_REDUCTION_inscan;
+})) {
+const auto &&NumIteratorsGen = [&S](CodeGenFunction &CGF) {
+CodeGenFunction::OMPLocalDeclMapRAII Scope(CGF);
+CGCapturedStmtInfo CGSI(CR_OpenMP);
+CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGSI);
+OMPLoopScope LoopScope(CGF, S);
+return CGF.EmitScalarExpr(S.getNumIterations());
+};
+emitScanBasedDirectiveDecls(*this, S, NumIteratorsGen);
+}
 auto LPCRegion =
 CGOpenMPRuntime::LastprivateConditionalRAII::disable(*this, S);
 emitCommonOMPParallelDirective(*this, S, OMPD_for, CodeGen,
 emitEmptyBoundParameters);
 }
 const OMPParallelForSimdDirective &S) {
 // Emit directive as a combined directive that consists of two implicit
 // directives: 'parallel' with 'for' directive.
 auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &Action) {
 Action.Enter(CGF);
-CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(), emitForLoopBounds,
+(void)emitWorksharingDirective(CGF, S, /*HasCancel=*/false);
-emitDispatchForLoopBounds);
 };
 {
+if (llvm::any_of(S.getClausesOfKind<OMPReductionClause>(),
+[](const OMPReductionClause *C) {
+return C->getModifier() == OMPC_REDUCTION_inscan;
+})) {
+const auto &&NumIteratorsGen = [&S](CodeGenFunction &CGF) {
+CodeGenFunction::OMPLocalDeclMapRAII Scope(CGF);
+CGCapturedStmtInfo CGSI(CR_OpenMP);
+CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGSI);
+OMPLoopScope LoopScope(CGF, S);
+return CGF.EmitScalarExpr(S.getNumIterations());
+};
+emitScanBasedDirectiveDecls(*this, S, NumIteratorsGen);
+}
 auto LPCRegion =
 CGOpenMPRuntime::LastprivateConditionalRAII::disable(*this, S);
-emitCommonOMPParallelDirective(*this, S, OMPD_simd, CodeGen,
+emitCommonOMPParallelDirective(*this, S, OMPD_for_simd, CodeGen,
 emitEmptyBoundParameters);
 }
 // Check for outer lastprivate conditional update.
 checkForLastprivateConditionalUpdate(*this, S);
 }
 emitEmptyBoundParameters);
 }
 // Check for outer lastprivate conditional update.
 checkForLastprivateConditionalUpdate(*this, S);
 }
+namespace {
+/// Get the list of variables declared in the context of the untied tasks.
+class CheckVarsEscapingUntiedTaskDeclContext final
+: public ConstStmtVisitor<CheckVarsEscapingUntiedTaskDeclContext> {
+llvm::SmallVector<const VarDecl *, 4> PrivateDecls;
+public:
+explicit CheckVarsEscapingUntiedTaskDeclContext() = default;
+virtual ~CheckVarsEscapingUntiedTaskDeclContext() = default;
+void VisitDeclStmt(const DeclStmt *S) {
+if (!S)
+return;
+// Need to privatize only local vars, static locals can be processed as is.
+for (const Decl *D : S->decls()) {
+if (const auto *VD = dyn_cast_or_null<VarDecl>(D))
+if (VD->hasLocalStorage())
+PrivateDecls.push_back(VD);
+}
+}
+void VisitOMPExecutableDirective(const OMPExecutableDirective *) { return; }
+void VisitCapturedStmt(const CapturedStmt *) { return; }
+void VisitLambdaExpr(const LambdaExpr *) { return; }
+void VisitBlockExpr(const BlockExpr *) { return; }
+void VisitStmt(const Stmt *S) {
+if (!S)
+return;
+for (const Stmt *Child : S->children())
+if (Child)
+Visit(Child);
+}
+/// Swaps list of vars with the provided one.
+ArrayRef<const VarDecl *> getPrivateDecls() const { return PrivateDecls; }
+};
+} // anonymous namespace
 void CodeGenFunction::EmitOMPTaskBasedDirective(
 const OMPExecutableDirective &S, const OpenMPDirectiveKind CapturedRegion,
 const RegionCodeGenTy &BodyGen, const TaskGenTy &TaskGen,
 OMPTaskDataTy &Data) {
 ++IRef;
 ++IElemInitRef;
 }
 }
 // Get list of lastprivate variables (for taskloops).
-llvm::DenseMap<const VarDecl *, const DeclRefExpr *> LastprivateDstsOrigs;
+llvm::MapVector<const VarDecl *, const DeclRefExpr *> LastprivateDstsOrigs;
 for (const auto *C : S.getClausesOfKind<OMPLastprivateClause>()) {
 auto IRef = C->varlist_begin();
 auto ID = C->destination_exprs().begin();
 for (const Expr *IInit : C->private_copies()) {
 const auto *OrigVD = cast<VarDecl>(cast<DeclRefExpr>(*IRef)->getDecl());
 if (EmittedAsPrivate.insert(OrigVD->getCanonicalDecl()).second) {
 Data.LastprivateVars.push_back(*IRef);
 Data.LastprivateCopies.push_back(IInit);
 }
 LastprivateDstsOrigs.insert(
-{cast<VarDecl>(cast<DeclRefExpr>(*ID)->getDecl()),
+std::make_pair(cast<VarDecl>(cast<DeclRefExpr>(*ID)->getDecl()),
-cast<DeclRefExpr>(*IRef)});
+cast<DeclRefExpr>(*IRef)));
 ++IRef;
 ++ID;
 }
 }
 SmallVector<const Expr *, 4> LHSs;
 for (const auto *C : S.getClausesOfKind<OMPDependClause>()) {
 OMPTaskDataTy::DependData &DD =
 Data.Dependences.emplace_back(C->getDependencyKind(), C->getModifier());
 DD.DepExprs.append(C->varlist_begin(), C->varlist_end());
 }
+// Get list of local vars for untied tasks.
+if (!Data.Tied) {
+CheckVarsEscapingUntiedTaskDeclContext Checker;
+Checker.Visit(S.getInnermostCapturedStmt()->getCapturedStmt());
+Data.PrivateLocals.append(Checker.getPrivateDecls().begin(),
+Checker.getPrivateDecls().end());
+}
 auto &&CodeGen = [&Data, &S, CS, &BodyGen, &LastprivateDstsOrigs,
 CapturedRegion](CodeGenFunction &CGF,
 PrePostActionTy &Action) {
+llvm::MapVector<CanonicalDeclPtr<const VarDecl>,
+std::pair<Address, Address>>
+UntiedLocalVars;
 // Set proper addresses for generated private copies.
 OMPPrivateScope Scope(CGF);
 llvm::SmallVector<std::pair<const VarDecl *, Address>, 16> FirstprivatePtrs;
 if (!Data.PrivateVars.empty() || !Data.FirstprivateVars.empty() ||
-!Data.LastprivateVars.empty()) {
+!Data.LastprivateVars.empty() || !Data.PrivateLocals.empty()) {
-llvm::FunctionType *CopyFnTy = llvm::FunctionType::get(
-CGF.Builder.getVoidTy(), {CGF.Builder.getInt8PtrTy()}, true);
 enum { PrivatesParam = 2, CopyFnParam = 3 };
 llvm::Value *CopyFn = CGF.Builder.CreateLoad(
 CGF.GetAddrOfLocalVar(CS->getCapturedDecl()->getParam(CopyFnParam)));
 llvm::Value *PrivatesPtr = CGF.Builder.CreateLoad(CGF.GetAddrOfLocalVar(
 CS->getCapturedDecl()->getParam(PrivatesParam)));
 // Map privates.
 llvm::SmallVector<std::pair<const VarDecl *, Address>, 16> PrivatePtrs;
 llvm::SmallVector<llvm::Value *, 16> CallArgs;
+llvm::SmallVector<llvm::Type *, 4> ParamTypes;
 CallArgs.push_back(PrivatesPtr);
+ParamTypes.push_back(PrivatesPtr->getType());
 for (const Expr *E : Data.PrivateVars) {
 const auto *VD = cast<VarDecl>(cast<DeclRefExpr>(E)->getDecl());
 Address PrivatePtr = CGF.CreateMemTemp(
 CGF.getContext().getPointerType(E->getType()), ".priv.ptr.addr");
 PrivatePtrs.emplace_back(VD, PrivatePtr);
 CallArgs.push_back(PrivatePtr.getPointer());
+ParamTypes.push_back(PrivatePtr.getType());
 }
 for (const Expr *E : Data.FirstprivateVars) {
 const auto *VD = cast<VarDecl>(cast<DeclRefExpr>(E)->getDecl());
 Address PrivatePtr =
 CGF.CreateMemTemp(CGF.getContext().getPointerType(E->getType()),
 ".firstpriv.ptr.addr");
 PrivatePtrs.emplace_back(VD, PrivatePtr);
 FirstprivatePtrs.emplace_back(VD, PrivatePtr);
 CallArgs.push_back(PrivatePtr.getPointer());
+ParamTypes.push_back(PrivatePtr.getType());
 }
 for (const Expr *E : Data.LastprivateVars) {
 const auto *VD = cast<VarDecl>(cast<DeclRefExpr>(E)->getDecl());
 Address PrivatePtr =
 CGF.CreateMemTemp(CGF.getContext().getPointerType(E->getType()),
 ".lastpriv.ptr.addr");
 PrivatePtrs.emplace_back(VD, PrivatePtr);
 CallArgs.push_back(PrivatePtr.getPointer());
+ParamTypes.push_back(PrivatePtr.getType());
 }
+for (const VarDecl *VD : Data.PrivateLocals) {
+QualType Ty = VD->getType().getNonReferenceType();
+if (VD->getType()->isLValueReferenceType())
+Ty = CGF.getContext().getPointerType(Ty);
+if (isAllocatableDecl(VD))
+Ty = CGF.getContext().getPointerType(Ty);
+Address PrivatePtr = CGF.CreateMemTemp(
+CGF.getContext().getPointerType(Ty), ".local.ptr.addr");
+auto Result = UntiedLocalVars.insert(
+std::make_pair(VD, std::make_pair(PrivatePtr, Address::invalid())));
+// If key exists update in place.
+if (Result.second == false)
+*Result.first = std::make_pair(
+VD, std::make_pair(PrivatePtr, Address::invalid()));
+CallArgs.push_back(PrivatePtr.getPointer());
+ParamTypes.push_back(PrivatePtr.getType());
+}
+auto *CopyFnTy = llvm::FunctionType::get(CGF.Builder.getVoidTy(),
+ParamTypes, /*isVarArg=*/false);
+CopyFn = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
+CopyFn, CopyFnTy->getPointerTo());
 CGF.CGM.getOpenMPRuntime().emitOutlinedFunctionCall(
 CGF, S.getBeginLoc(), {CopyFnTy, CopyFn}, CallArgs);
 for (const auto &Pair : LastprivateDstsOrigs) {
 const auto *OrigVD = cast<VarDecl>(Pair.second->getDecl());
 DeclRefExpr DRE(CGF.getContext(), const_cast<VarDecl *>(OrigVD),
 }
 for (const auto &Pair : PrivatePtrs) {
 Address Replacement(CGF.Builder.CreateLoad(Pair.second),
 CGF.getContext().getDeclAlign(Pair.first));
 Scope.addPrivate(Pair.first, [Replacement]() { return Replacement; });
+}
+// Adjust mapping for internal locals by mapping actual memory instead of
+// a pointer to this memory.
+for (auto &Pair : UntiedLocalVars) {
+if (isAllocatableDecl(Pair.first)) {
+llvm::Value *Ptr = CGF.Builder.CreateLoad(Pair.second.first);
+Address Replacement(Ptr, CGF.getPointerAlign());
+Pair.second.first = Replacement;
+Ptr = CGF.Builder.CreateLoad(Replacement);
+Replacement = Address(Ptr, CGF.getContext().getDeclAlign(Pair.first));
+Pair.second.second = Replacement;
+} else {
+llvm::Value *Ptr = CGF.Builder.CreateLoad(Pair.second.first);
+Address Replacement(Ptr, CGF.getContext().getDeclAlign(Pair.first));
+Pair.second.first = Replacement;
+}
 }
 }
 if (Data.Reductions) {
 OMPPrivateScope FirstprivateScope(CGF);
 for (const auto &Pair : FirstprivatePtrs) {
 [Replacement]() { return Replacement; });
 }
 }
 (void)InRedScope.Privatize();
+CGOpenMPRuntime::UntiedTaskLocalDeclsRAII LocalVarsScope(CGF,
+UntiedLocalVars);
 Action.Enter(CGF);
 BodyGen(CGF);
 };
 llvm::Function *OutlinedFn = CGM.getOpenMPRuntime().emitTaskOutlinedFunction(
 S, *I, *PartId, *TaskT, S.getDirectiveKind(), CodeGen, Data.Tied,
 C, NestedNameSpecifierLoc(), SourceLocation(), InitVD,
 /*RefersToEnclosingVariableOrCapture=*/false, Loc, ElemType, VK_LValue);
 PrivateVD->setInitStyle(VarDecl::CInit);
 PrivateVD->setInit(ImplicitCastExpr::Create(C, ElemType, CK_LValueToRValue,
 InitRef, /*BasePath=*/nullptr,
-VK_RValue));
+VK_RValue, FPOptionsOverride()));
 Data.FirstprivateVars.emplace_back(OrigRef);
 Data.FirstprivateCopies.emplace_back(PrivateRef);
 Data.FirstprivateInits.emplace_back(InitRef);
 return OrigVD;
 }
 }
 OMPPrivateScope TargetScope(*this);
 VarDecl *BPVD = nullptr;
 VarDecl *PVD = nullptr;
 VarDecl *SVD = nullptr;
+VarDecl *MVD = nullptr;
 if (InputInfo.NumberOfTargetItems > 0) {
 auto *CD = CapturedDecl::Create(
 getContext(), getContext().getTranslationUnitDecl(), /*NumParams=*/0);
 llvm::APInt ArrSize(/*numBits=*/32, InputInfo.NumberOfTargetItems);
-QualType BaseAndPointersType = getContext().getConstantArrayType(
+QualType BaseAndPointerAndMapperType = getContext().getConstantArrayType(
 getContext().VoidPtrTy, ArrSize, nullptr, ArrayType::Normal,
 /*IndexTypeQuals=*/0);
 BPVD = createImplicitFirstprivateForType(
-getContext(), Data, BaseAndPointersType, CD, S.getBeginLoc());
+getContext(), Data, BaseAndPointerAndMapperType, CD, S.getBeginLoc());
 PVD = createImplicitFirstprivateForType(
-getContext(), Data, BaseAndPointersType, CD, S.getBeginLoc());
+getContext(), Data, BaseAndPointerAndMapperType, CD, S.getBeginLoc());
 QualType SizesType = getContext().getConstantArrayType(
 getContext().getIntTypeForBitwidth(/*DestWidth=*/64, /*Signed=*/1),
 ArrSize, nullptr, ArrayType::Normal,
 /*IndexTypeQuals=*/0);
 SVD = createImplicitFirstprivateForType(getContext(), Data, SizesType, CD,
 BPVD, [&InputInfo]() { return InputInfo.BasePointersArray; });
 TargetScope.addPrivate(PVD,
 [&InputInfo]() { return InputInfo.PointersArray; });
 TargetScope.addPrivate(SVD,
 [&InputInfo]() { return InputInfo.SizesArray; });
+// If there is no user-defined mapper, the mapper array will be nullptr. In
+// this case, we don't need to privatize it.
+if (!dyn_cast_or_null<llvm::ConstantPointerNull>(
+InputInfo.MappersArray.getPointer())) {
+MVD = createImplicitFirstprivateForType(
+getContext(), Data, BaseAndPointerAndMapperType, CD, S.getBeginLoc());
+TargetScope.addPrivate(MVD,
+[&InputInfo]() { return InputInfo.MappersArray; });
+}
 }
 (void)TargetScope.Privatize();
 // Build list of dependences.
 for (const auto *C : S.getClausesOfKind<OMPDependClause>()) {
 OMPTaskDataTy::DependData &DD =
 Data.Dependences.emplace_back(C->getDependencyKind(), C->getModifier());
 DD.DepExprs.append(C->varlist_begin(), C->varlist_end());
 }
-auto &&CodeGen = [&Data, &S, CS, &BodyGen, BPVD, PVD, SVD,
+auto &&CodeGen = [&Data, &S, CS, &BodyGen, BPVD, PVD, SVD, MVD,
 &InputInfo](CodeGenFunction &CGF, PrePostActionTy &Action) {
 // Set proper addresses for generated private copies.
 OMPPrivateScope Scope(CGF);
 if (!Data.FirstprivateVars.empty()) {
-llvm::FunctionType *CopyFnTy = llvm::FunctionType::get(
-CGF.Builder.getVoidTy(), {CGF.Builder.getInt8PtrTy()}, true);
 enum { PrivatesParam = 2, CopyFnParam = 3 };
 llvm::Value *CopyFn = CGF.Builder.CreateLoad(
 CGF.GetAddrOfLocalVar(CS->getCapturedDecl()->getParam(CopyFnParam)));
 llvm::Value *PrivatesPtr = CGF.Builder.CreateLoad(CGF.GetAddrOfLocalVar(
 CS->getCapturedDecl()->getParam(PrivatesParam)));
 // Map privates.
 llvm::SmallVector<std::pair<const VarDecl *, Address>, 16> PrivatePtrs;
 llvm::SmallVector<llvm::Value *, 16> CallArgs;
+llvm::SmallVector<llvm::Type *, 4> ParamTypes;
 CallArgs.push_back(PrivatesPtr);
+ParamTypes.push_back(PrivatesPtr->getType());
 for (const Expr *E : Data.FirstprivateVars) {
 const auto *VD = cast<VarDecl>(cast<DeclRefExpr>(E)->getDecl());
 Address PrivatePtr =
 CGF.CreateMemTemp(CGF.getContext().getPointerType(E->getType()),
 ".firstpriv.ptr.addr");
 PrivatePtrs.emplace_back(VD, PrivatePtr);
 CallArgs.push_back(PrivatePtr.getPointer());
+ParamTypes.push_back(PrivatePtr.getType());
 }
+auto *CopyFnTy = llvm::FunctionType::get(CGF.Builder.getVoidTy(),
+ParamTypes, /*isVarArg=*/false);
+CopyFn = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
+CopyFn, CopyFnTy->getPointerTo());
 CGF.CGM.getOpenMPRuntime().emitOutlinedFunctionCall(
 CGF, S.getBeginLoc(), {CopyFnTy, CopyFn}, CallArgs);
 for (const auto &Pair : PrivatePtrs) {
 Address Replacement(CGF.Builder.CreateLoad(Pair.second),
 CGF.getContext().getDeclAlign(Pair.first));
 CGF.GetAddrOfLocalVar(BPVD), /*Index=*/0);
 InputInfo.PointersArray = CGF.Builder.CreateConstArrayGEP(
 CGF.GetAddrOfLocalVar(PVD), /*Index=*/0);
 InputInfo.SizesArray = CGF.Builder.CreateConstArrayGEP(
 CGF.GetAddrOfLocalVar(SVD), /*Index=*/0);
+// If MVD is nullptr, the mapper array is not privatized
+if (MVD)
+InputInfo.MappersArray = CGF.Builder.CreateConstArrayGEP(
+CGF.GetAddrOfLocalVar(MVD), /*Index=*/0);
 }
 Action.Enter(CGF);
 OMPLexicalScope LexScope(CGF, S, OMPD_task, /*EmitPreInitStmt=*/false);
 BodyGen(CGF);
 if (const auto *UC = S.getSingleClause<OMPUpdateClause>()) {
 CGM.getOpenMPRuntime().emitUpdateClause(
 *this, DOLVal, UC->getDependencyKind(), UC->getBeginLoc());
 return;
 }
+}
+void CodeGenFunction::EmitOMPScanDirective(const OMPScanDirective &S) {
+if (!OMPParentLoopDirectiveForScan)
+return;
+const OMPExecutableDirective &ParentDir = *OMPParentLoopDirectiveForScan;
+bool IsInclusive = S.hasClausesOfKind<OMPInclusiveClause>();
+SmallVector<const Expr *, 4> Shareds;
+SmallVector<const Expr *, 4> Privates;
+SmallVector<const Expr *, 4> LHSs;
+SmallVector<const Expr *, 4> RHSs;
+SmallVector<const Expr *, 4> ReductionOps;
+SmallVector<const Expr *, 4> CopyOps;
+SmallVector<const Expr *, 4> CopyArrayTemps;
+SmallVector<const Expr *, 4> CopyArrayElems;
+for (const auto *C : ParentDir.getClausesOfKind<OMPReductionClause>()) {
+if (C->getModifier() != OMPC_REDUCTION_inscan)
+continue;
+Shareds.append(C->varlist_begin(), C->varlist_end());
+Privates.append(C->privates().begin(), C->privates().end());
+LHSs.append(C->lhs_exprs().begin(), C->lhs_exprs().end());
+RHSs.append(C->rhs_exprs().begin(), C->rhs_exprs().end());
+ReductionOps.append(C->reduction_ops().begin(), C->reduction_ops().end());
+CopyOps.append(C->copy_ops().begin(), C->copy_ops().end());
+CopyArrayTemps.append(C->copy_array_temps().begin(),
+C->copy_array_temps().end());
+CopyArrayElems.append(C->copy_array_elems().begin(),
+C->copy_array_elems().end());
+}
+if (ParentDir.getDirectiveKind() == OMPD_simd ||
+(getLangOpts().OpenMPSimd &&
+isOpenMPSimdDirective(ParentDir.getDirectiveKind()))) {
+// For simd directive and simd-based directives in simd only mode, use the
+// following codegen:
+// int x = 0;
+// #pragma omp simd reduction(inscan, +: x)
+// for (..) {
+//   <first part>
+//   #pragma omp scan inclusive(x)
+//   <second part>
+//  }
+// is transformed to:
+// int x = 0;
+// for (..) {
+//   int x_priv = 0;
+//   <first part>
+//   x = x_priv + x;
+//   x_priv = x;
+//   <second part>
+// }
+// and
+// int x = 0;
+// #pragma omp simd reduction(inscan, +: x)
+// for (..) {
+//   <first part>
+//   #pragma omp scan exclusive(x)
+//   <second part>
+// }
+// to
+// int x = 0;
+// for (..) {
+//   int x_priv = 0;
+//   <second part>
+//   int temp = x;
+//   x = x_priv + x;
+//   x_priv = temp;
+//   <first part>
+// }
+llvm::BasicBlock *OMPScanReduce = createBasicBlock("omp.inscan.reduce");
+EmitBranch(IsInclusive
+? OMPScanReduce
+: BreakContinueStack.back().ContinueBlock.getBlock());
+EmitBlock(OMPScanDispatch);
+{
+// New scope for correct construction/destruction of temp variables for
+// exclusive scan.
+LexicalScope Scope(*this, S.getSourceRange());
+EmitBranch(IsInclusive ? OMPBeforeScanBlock : OMPAfterScanBlock);
+EmitBlock(OMPScanReduce);
+if (!IsInclusive) {
+// Create temp var and copy LHS value to this temp value.
+// TMP = LHS;
+for (unsigned I = 0, E = CopyArrayElems.size(); I < E; ++I) {
+const Expr *PrivateExpr = Privates[I];
+const Expr *TempExpr = CopyArrayTemps[I];
+EmitAutoVarDecl(
+*cast<VarDecl>(cast<DeclRefExpr>(TempExpr)->getDecl()));
+LValue DestLVal = EmitLValue(TempExpr);
+LValue SrcLVal = EmitLValue(LHSs[I]);
+EmitOMPCopy(PrivateExpr->getType(), DestLVal.getAddress(*this),
+SrcLVal.getAddress(*this),
+cast<VarDecl>(cast<DeclRefExpr>(LHSs[I])->getDecl()),
+cast<VarDecl>(cast<DeclRefExpr>(RHSs[I])->getDecl()),
+CopyOps[I]);
+}
+}
+CGM.getOpenMPRuntime().emitReduction(
+*this, ParentDir.getEndLoc(), Privates, LHSs, RHSs, ReductionOps,
+{/*WithNowait=*/true, /*SimpleReduction=*/true, OMPD_simd});
+for (unsigned I = 0, E = CopyArrayElems.size(); I < E; ++I) {
+const Expr *PrivateExpr = Privates[I];
+LValue DestLVal;
+LValue SrcLVal;
+if (IsInclusive) {
+DestLVal = EmitLValue(RHSs[I]);
+SrcLVal = EmitLValue(LHSs[I]);
+} else {
+const Expr *TempExpr = CopyArrayTemps[I];
+DestLVal = EmitLValue(RHSs[I]);
+SrcLVal = EmitLValue(TempExpr);
+}
+EmitOMPCopy(PrivateExpr->getType(), DestLVal.getAddress(*this),
+SrcLVal.getAddress(*this),
+cast<VarDecl>(cast<DeclRefExpr>(LHSs[I])->getDecl()),
+cast<VarDecl>(cast<DeclRefExpr>(RHSs[I])->getDecl()),
+CopyOps[I]);
+}
+}
+EmitBranch(IsInclusive ? OMPAfterScanBlock : OMPBeforeScanBlock);
+OMPScanExitBlock = IsInclusive
+? BreakContinueStack.back().ContinueBlock.getBlock()
+: OMPScanReduce;
+EmitBlock(OMPAfterScanBlock);
+return;
+}
+if (!IsInclusive) {
+EmitBranch(BreakContinueStack.back().ContinueBlock.getBlock());
+EmitBlock(OMPScanExitBlock);
+}
+if (OMPFirstScanLoop) {
+// Emit buffer[i] = red; at the end of the input phase.
+const auto *IVExpr = cast<OMPLoopDirective>(ParentDir)
+.getIterationVariable()
+->IgnoreParenImpCasts();
+LValue IdxLVal = EmitLValue(IVExpr);
+llvm::Value *IdxVal = EmitLoadOfScalar(IdxLVal, IVExpr->getExprLoc());
+IdxVal = Builder.CreateIntCast(IdxVal, SizeTy, /*isSigned=*/false);
+for (unsigned I = 0, E = CopyArrayElems.size(); I < E; ++I) {
+const Expr *PrivateExpr = Privates[I];
+const Expr *OrigExpr = Shareds[I];
+const Expr *CopyArrayElem = CopyArrayElems[I];
+OpaqueValueMapping IdxMapping(
+*this,
+cast<OpaqueValueExpr>(
+cast<ArraySubscriptExpr>(CopyArrayElem)->getIdx()),
+RValue::get(IdxVal));
+LValue DestLVal = EmitLValue(CopyArrayElem);
+LValue SrcLVal = EmitLValue(OrigExpr);
+EmitOMPCopy(PrivateExpr->getType(), DestLVal.getAddress(*this),
+SrcLVal.getAddress(*this),
+cast<VarDecl>(cast<DeclRefExpr>(LHSs[I])->getDecl()),
+cast<VarDecl>(cast<DeclRefExpr>(RHSs[I])->getDecl()),
+CopyOps[I]);
+}
+}
+EmitBranch(BreakContinueStack.back().ContinueBlock.getBlock());
+if (IsInclusive) {
+EmitBlock(OMPScanExitBlock);
+EmitBranch(BreakContinueStack.back().ContinueBlock.getBlock());
+}
+EmitBlock(OMPScanDispatch);
+if (!OMPFirstScanLoop) {
+// Emit red = buffer[i]; at the entrance to the scan phase.
+const auto *IVExpr = cast<OMPLoopDirective>(ParentDir)
+.getIterationVariable()
+->IgnoreParenImpCasts();
+LValue IdxLVal = EmitLValue(IVExpr);
+llvm::Value *IdxVal = EmitLoadOfScalar(IdxLVal, IVExpr->getExprLoc());
+IdxVal = Builder.CreateIntCast(IdxVal, SizeTy, /*isSigned=*/false);
+llvm::BasicBlock *ExclusiveExitBB = nullptr;
+if (!IsInclusive) {
+llvm::BasicBlock *ContBB = createBasicBlock("omp.exclusive.dec");
+ExclusiveExitBB = createBasicBlock("omp.exclusive.copy.exit");
+llvm::Value *Cmp = Builder.CreateIsNull(IdxVal);
+Builder.CreateCondBr(Cmp, ExclusiveExitBB, ContBB);
+EmitBlock(ContBB);
+// Use idx - 1 iteration for exclusive scan.
+IdxVal = Builder.CreateNUWSub(IdxVal, llvm::ConstantInt::get(SizeTy, 1));
+}
+for (unsigned I = 0, E = CopyArrayElems.size(); I < E; ++I) {
+const Expr *PrivateExpr = Privates[I];
+const Expr *OrigExpr = Shareds[I];
+const Expr *CopyArrayElem = CopyArrayElems[I];
+OpaqueValueMapping IdxMapping(
+*this,
+cast<OpaqueValueExpr>(
+cast<ArraySubscriptExpr>(CopyArrayElem)->getIdx()),
+RValue::get(IdxVal));
+LValue SrcLVal = EmitLValue(CopyArrayElem);
+LValue DestLVal = EmitLValue(OrigExpr);
+EmitOMPCopy(PrivateExpr->getType(), DestLVal.getAddress(*this),
+SrcLVal.getAddress(*this),
+cast<VarDecl>(cast<DeclRefExpr>(LHSs[I])->getDecl()),
+cast<VarDecl>(cast<DeclRefExpr>(RHSs[I])->getDecl()),
+CopyOps[I]);
+}
+if (!IsInclusive) {
+EmitBlock(ExclusiveExitBB);
+}
+}
+EmitBranch((OMPFirstScanLoop == IsInclusive) ? OMPBeforeScanBlock
+: OMPAfterScanBlock);
+EmitBlock(OMPAfterScanBlock);
 }
 void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S,
 const CodeGenLoopTy &CodeGenLoop,
 Expr *IncExpr) {
 return Fn;
 }
 void CodeGenFunction::EmitOMPOrderedDirective(const OMPOrderedDirective &S) {
 if (S.hasClausesOfKind<OMPDependClause>()) {
-assert(!S.getAssociatedStmt() &&
+assert(!S.hasAssociatedStmt() &&
 "No associated statement must be in ordered depend construct.");
 for (const auto *DC : S.getClausesOfKind<OMPDependClause>())
 CGM.getOpenMPRuntime().emitDoacrossOrdered(*this, DC);
 return;
 }
 case OMPC_reduction:
 case OMPC_task_reduction:
 case OMPC_in_reduction:
 case OMPC_safelen:
 case OMPC_simdlen:
+case OMPC_sizes:
 case OMPC_allocator:
 case OMPC_allocate:
 case OMPC_collapse:
 case OMPC_default:
 case OMPC_seq_cst:
 case OMPC_defaultmap:
 case OMPC_uniform:
 case OMPC_to:
 case OMPC_from:
 case OMPC_use_device_ptr:
+case OMPC_use_device_addr:
 case OMPC_is_device_ptr:
 case OMPC_unified_address:
 case OMPC_unified_shared_memory:
 case OMPC_reverse_offload:
 case OMPC_dynamic_allocators:
 case OMPC_detach:
 case OMPC_inclusive:
 case OMPC_exclusive:
 case OMPC_uses_allocators:
 case OMPC_affinity:
+case OMPC_init:
+case OMPC_inbranch:
+case OMPC_notinbranch:
+case OMPC_link:
+case OMPC_use:
+case OMPC_novariants:
+case OMPC_nocontext:
+case OMPC_filter:
 llvm_unreachable("Clause is not allowed in 'omp atomic'.");
 }
 }
 void CodeGenFunction::EmitOMPAtomicDirective(const OMPAtomicDirective &S) {
 // if it is first).
 if (C->getClauseKind() != OMPC_seq_cst &&
 C->getClauseKind() != OMPC_acq_rel &&
 C->getClauseKind() != OMPC_acquire &&
 C->getClauseKind() != OMPC_release &&
-C->getClauseKind() != OMPC_relaxed) {
+C->getClauseKind() != OMPC_relaxed && C->getClauseKind() != OMPC_hint) {
 Kind = C->getClauseKind();
 break;
 }
 }
 if (!MemOrderingSpecified) {
 AO = llvm::AtomicOrdering::Acquire;
 }
 }
 }
-const Stmt *CS = S.getInnermostCapturedStmt()->IgnoreContainers();
+LexicalScope Scope(*this, S.getSourceRange());
-if (const auto *FE = dyn_cast<FullExpr>(CS))
+EmitStopPoint(S.getAssociatedStmt());
-enterFullExpression(FE);
+emitOMPAtomicExpr(*this, Kind, AO, S.isPostfixUpdate(), S.getX(), S.getV(),
-// Processing for statements under 'atomic capture'.
+S.getExpr(), S.getUpdateExpr(), S.isXLHSInRHSPart(),
-if (const auto *Compound = dyn_cast<CompoundStmt>(CS)) {
+S.getBeginLoc());
-for (const Stmt *C : Compound->body()) {
-if (const auto *FE = dyn_cast<FullExpr>(C))
-enterFullExpression(FE);
-}
-}
-auto &&CodeGen = [&S, Kind, AO, CS](CodeGenFunction &CGF,
-PrePostActionTy &) {
-CGF.EmitStopPoint(CS);
-emitOMPAtomicExpr(CGF, Kind, AO, S.isPostfixUpdate(), S.getX(), S.getV(),
-S.getExpr(), S.getUpdateExpr(), S.isXLHSInRHSPart(),
-S.getBeginLoc());
-};
-OMPLexicalScope Scope(*this, S, OMPD_unknown);
-CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_atomic, CodeGen);
 }
 static void emitCommonOMPTargetDirective(CodeGenFunction &CGF,
 const OMPExecutableDirective &S,
 const RegionCodeGenTy &CodeGen) {
 (void)PrivateScope.Privatize();
 if (isOpenMPTargetExecutionDirective(S.getDirectiveKind()))
 CGF.CGM.getOpenMPRuntime().adjustTargetSpecificDataForLambdas(CGF, S);
 CGF.EmitStmt(S.getCapturedStmt(OMPD_target)->getCapturedStmt());
+CGF.EnsureInsertPoint();
 }
 void CodeGenFunction::EmitOMPTargetDeviceFunction(CodeGenModule &CGM,
 StringRef ParentName,
 const OMPTargetDirective &S) {
 C->getNameModifier() == OMPD_cancel) {
 IfCond = C->getCondition();
 break;
 }
 }
-if (llvm::OpenMPIRBuilder *OMPBuilder = CGM.getOpenMPIRBuilder()) {
+if (CGM.getLangOpts().OpenMPIRBuilder) {
+llvm::OpenMPIRBuilder &OMPBuilder = CGM.getOpenMPRuntime().getOMPBuilder();
 // TODO: This check is necessary as we only generate `omp parallel` through
 // the OpenMPIRBuilder for now.
-if (S.getCancelRegion() == OMPD_parallel) {
+if (S.getCancelRegion() == OMPD_parallel ||
+S.getCancelRegion() == OMPD_sections ||
+S.getCancelRegion() == OMPD_section) {
 llvm::Value *IfCondition = nullptr;
 if (IfCond)
 IfCondition = EmitScalarExpr(IfCond,
 /*IgnoreResultAssign=*/true);
 return Builder.restoreIP(
-OMPBuilder->CreateCancel(Builder, IfCondition, S.getCancelRegion()));
+OMPBuilder.createCancel(Builder, IfCondition, S.getCancelRegion()));
 }
 }
 CGM.getOpenMPRuntime().emitCancelCall(*this, S.getBeginLoc(), IfCond,
 S.getCancelRegion());
 Kind == OMPD_target_teams_distribute_parallel_for);
 return OMPCancelStack.getExitBlock();
 }
 void CodeGenFunction::EmitOMPUseDevicePtrClause(
-const OMPClause &NC, OMPPrivateScope &PrivateScope,
+const OMPUseDevicePtrClause &C, OMPPrivateScope &PrivateScope,
 const llvm::DenseMap<const ValueDecl *, Address> &CaptureDeviceAddrMap) {
-const auto &C = cast<OMPUseDevicePtrClause>(NC);
 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());
 ++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) {
+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();
+// 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())
+PrivAddr =
+EmitLoadOfPointer(PrivAddr, getContext()
+.getPointerType(OrigVD->getType())
+->castAs<PointerType>());
+llvm::Type *RealTy =
+ConvertTypeForMem(OrigVD->getType().getNonReferenceType())
+->getPointerTo();
+PrivAddr = Builder.CreatePointerBitCastOrAddrSpaceCast(PrivAddr, RealTy);
+(void)PrivateScope.addPrivate(OrigVD, [PrivAddr]() { return PrivAddr; });
+}
+}
 // Generate the instructions for '#pragma omp target data' directive.
 void CodeGenFunction::EmitOMPTargetDataDirective(
 const OMPTargetDataDirective &S) {
-CGOpenMPRuntime::TargetDataInfo Info(/*RequiresDevicePointerInfo=*/true);
+CGOpenMPRuntime::TargetDataInfo Info(/*RequiresDevicePointerInfo=*/true,
+/*SeparateBeginEndCalls=*/true);
 // Create a pre/post action to signal the privatization of the device pointer.
 // This action can be replaced by the OpenMP runtime code generation to
 // deactivate privatization.
 bool PrivatizeDevicePointers = false;
 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);
+for (const auto *C : S.getClausesOfKind<OMPUseDeviceAddrClause>())
+CGF.EmitOMPUseDeviceAddrClause(*C, PrivateScope,
+Info.CaptureDeviceAddrMap);
 (void)PrivateScope.Privatize();
 RCG(CGF);
 } else {
+OMPLexicalScope Scope(CGF, S, OMPD_unknown);
 RCG(CGF);
 }
 };
 // Forward the provided action to the privatization codegen.
 },
 [&S, &LoopScope](CodeGenFunction &CGF, PrePostActionTy &) {
 CGF.EmitOMPInnerLoop(
 S, LoopScope.requiresCleanups(), S.getCond(), S.getInc(),
 [&S](CodeGenFunction &CGF) {
-CGF.EmitOMPLoopBody(S, CodeGenFunction::JumpDest());
+emitOMPLoopBodyWithStopPoint(CGF, S,
-CGF.EmitStopPoint(&S);
+CodeGenFunction::JumpDest());
 },
 [](CodeGenFunction &) {});
 });
 }
 // Emit: if (PreCond) - end.
 CGM.getOpenMPRuntime().emitTargetDataStandAloneCall(*this, S, IfCond, Device);
 }
 void CodeGenFunction::EmitSimpleOMPExecutableDirective(
 const OMPExecutableDirective &D) {
+if (const auto *SD = dyn_cast<OMPScanDirective>(&D)) {
+EmitOMPScanDirective(*SD);
+return;
+}
 if (!D.hasAssociatedStmt() || !D.getAssociatedStmt())
 return;
 auto &&CodeGen = [&D](CodeGenFunction &CGF, PrePostActionTy &Action) {
 OMPPrivateScope GlobalsScope(CGF);
 if (isOpenMPTaskingDirective(D.getDirectiveKind())) {
 }
 }
 }
 if (isOpenMPSimdDirective(D.getDirectiveKind())) {
 (void)GlobalsScope.Privatize();
+ParentLoopDirectiveForScanRegion ScanRegion(CGF, D);
 emitOMPSimdRegion(CGF, cast<OMPLoopDirective>(D), Action);
 } else {
 if (const auto *LD = dyn_cast<OMPLoopDirective>(&D)) {
 for (const Expr *E : LD->counters()) {
 const auto *VD = cast<VarDecl>(cast<DeclRefExpr>(E)->getDecl());
 }
 }
 for (const auto *C : D.getClausesOfKind<OMPOrderedClause>()) {
 if (!C->getNumForLoops())
 continue;
-for (unsigned I = LD->getCollapsedNumber(),
+for (unsigned I = LD->getLoopsNumber(),
 E = C->getLoopNumIterations().size();
 I < E; ++I) {
 if (const auto *VD = dyn_cast<OMPCapturedExprDecl>(
 cast<DeclRefExpr>(C->getLoopCounter(I))->getDecl())) {
 // Emit only those that were not explicitly referenced in clauses.
 }
 (void)GlobalsScope.Privatize();
 CGF.EmitStmt(D.getInnermostCapturedStmt()->getCapturedStmt());
 }
 };
-{
+if (D.getDirectiveKind() == OMPD_atomic ||
+D.getDirectiveKind() == OMPD_critical ||
+D.getDirectiveKind() == OMPD_section ||
+D.getDirectiveKind() == OMPD_master ||
+D.getDirectiveKind() == OMPD_masked) {
+EmitStmt(D.getAssociatedStmt());
+} else {
 auto LPCRegion =
 CGOpenMPRuntime::LastprivateConditionalRAII::disable(*this, D);
 OMPSimdLexicalScope Scope(*this, D);
 CGM.getOpenMPRuntime().emitInlinedDirective(
 *this,

Mercurial > hg > CbC > CbC_llvm

comparison clang/lib/CodeGen/CGStmtOpenMP.cpp @ 221:79ff65ed7e25