domada created this revision.
domada added reviewers: jsjodin, skatrak, agozillon, kiranchandramohan,
kiranktp, NimishMishra, TIFitis, raghavendhra, dpalermo, jdoerfert.
domada added a project: OpenMP.
Herald added subscribers: gysit, Dinistro, bviyer, Moerafaat, zero9178,
bzcheeseman, awarzynski, sdasgup3, wenzhicui, wrengr, ormris, cota, teijeong,
rdzhabarov, tatianashp, msifontes, jurahul, Kayjukh, grosul1, Joonsoo,
liufengdb, aartbik, mgester, arpith-jacob, antiagainst, shauheen, rriddle,
mehdi_amini, hiraditya.
Herald added a reviewer: ftynse.
Herald added a reviewer: dcaballe.
Herald added a project: All.
domada requested review of this revision.
Herald added subscribers: llvm-commits, cfe-commits, jplehr, sstefan1,
stephenneuendorffer, nicolasvasilache.
Herald added projects: clang, MLIR, LLVM.
OMP offload module metadata should be created only when module generation is
finalized. If we finalize LLVM IR function we should not create omp offload
metadata.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D156423
Files:
clang/lib/CodeGen/CodeGenFunction.cpp
clang/test/OpenMP/irbuilder_omp_offload_metadata.c
llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
llvm/lib/Transforms/IPO/OpenMPOpt.cpp
llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
mlir/lib/Target/LLVMIR/ModuleTranslation.cpp
Index: mlir/lib/Target/LLVMIR/ModuleTranslation.cpp
===================================================================
--- mlir/lib/Target/LLVMIR/ModuleTranslation.cpp
+++ mlir/lib/Target/LLVMIR/ModuleTranslation.cpp
@@ -471,7 +471,7 @@
ModuleTranslation::~ModuleTranslation() {
if (ompBuilder)
- ompBuilder->finalize();
+ ompBuilder->finalizeModule();
}
void ModuleTranslation::forgetMapping(Region ®ion) {
Index: llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
===================================================================
--- llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
+++ llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
@@ -664,7 +664,7 @@
Builder.restoreIP(AfterIP);
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_NE(PrivAI, nullptr);
Function *OutlinedFn = PrivAI->getFunction();
@@ -760,7 +760,7 @@
Builder.restoreIP(AfterIP);
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_EQ(M->size(), 5U);
for (Function &OutlinedFn : *M) {
@@ -864,7 +864,7 @@
Builder.restoreIP(AfterIP);
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_EQ(M->size(), 6U);
for (Function &OutlinedFn : *M) {
@@ -976,7 +976,7 @@
Builder.restoreIP(AfterIP);
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_NE(PrivAI, nullptr);
Function *OutlinedFn = PrivAI->getFunction();
@@ -1093,7 +1093,7 @@
Builder.restoreIP(AfterIP);
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
@@ -1172,7 +1172,7 @@
Builder.restoreIP(AfterIP);
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
Function *OutlinedFn = Internal->getFunction();
@@ -1206,7 +1206,7 @@
Builder.restoreIP(Loop->getAfterIP());
ReturnInst *RetInst = Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
Loop->assertOK();
EXPECT_FALSE(verifyModule(*M, &errs()));
@@ -1307,7 +1307,7 @@
// Finalize the function and verify it.
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
}
@@ -1368,7 +1368,7 @@
CanonicalLoopInfo *Collapsed =
OMPBuilder.collapseLoops(DL, {OuterLoop, InnerLoop}, ComputeIP);
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
// Verify control flow and BB order.
@@ -1419,7 +1419,7 @@
std::vector<CanonicalLoopInfo *> GenLoops =
OMPBuilder.tileLoops(DL, {Loop}, {TileSize});
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
EXPECT_EQ(GenLoops.size(), 2u);
@@ -1486,7 +1486,7 @@
std::vector<CanonicalLoopInfo *> GenLoops = OMPBuilder.tileLoops(
DL, {OuterLoop, InnerLoop}, {OuterTileSize, InnerTileSize});
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
EXPECT_EQ(GenLoops.size(), 4u);
@@ -1588,7 +1588,7 @@
std::vector<CanonicalLoopInfo *> GenLoops =
OMPBuilder.tileLoops(DL, {OuterLoop, InnerLoop}, {TileSize0, TileSize1});
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
EXPECT_EQ(GenLoops.size(), 4u);
@@ -1741,7 +1741,7 @@
// Finalize the function.
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
}
@@ -1757,7 +1757,7 @@
/* Simdlen */ nullptr,
/* Safelen */ nullptr);
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
PassBuilder PB;
@@ -1798,7 +1798,7 @@
/* Simdlen */ nullptr,
/* Safelen */ nullptr);
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
PassBuilder PB;
@@ -1853,7 +1853,7 @@
ConstantInt::get(Type::getInt32Ty(Ctx), 3),
/* Safelen */ nullptr);
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
PassBuilder PB;
@@ -1888,7 +1888,7 @@
CLI, AlignedVars, /* IfCond */ nullptr, OrderKind::OMP_ORDER_concurrent,
/* Simdlen */ nullptr, ConstantInt::get(Type::getInt32Ty(Ctx), 3));
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
PassBuilder PB;
@@ -1924,7 +1924,7 @@
CLI, AlignedVars, /* IfCond */ nullptr, OrderKind::OMP_ORDER_unknown,
/* Simdlen */ nullptr, ConstantInt::get(Type::getInt32Ty(Ctx), 3));
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
PassBuilder PB;
@@ -1959,7 +1959,7 @@
ConstantInt::get(Type::getInt32Ty(Ctx), 2),
ConstantInt::get(Type::getInt32Ty(Ctx), 3));
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
PassBuilder PB;
@@ -2005,7 +2005,7 @@
ConstantInt::get(Type::getInt32Ty(Ctx), 3),
/* Safelen */ nullptr);
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
PassBuilder PB;
@@ -2042,7 +2042,7 @@
// Unroll the loop.
OMPBuilder.unrollLoopFull(DL, CLI);
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
PassBuilder PB;
@@ -2067,7 +2067,7 @@
OMPBuilder.unrollLoopPartial(DL, CLI, 5, &UnrolledLoop);
ASSERT_NE(UnrolledLoop, nullptr);
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
UnrolledLoop->assertOK();
@@ -2099,7 +2099,7 @@
// Unroll the loop.
OMPBuilder.unrollLoopHeuristic(DL, CLI);
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
PassBuilder PB;
@@ -2233,7 +2233,7 @@
OMPBuilder.applyWorkshareLoop(DL, CLI, AllocaIP, /*NeedsBarrier=*/true,
OMP_SCHEDULE_Static, ChunkSize);
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
BasicBlock *Entry = &F->getEntryBlock();
@@ -2422,7 +2422,7 @@
// Add a termination to our block and check that it is internally consistent.
Builder.restoreIP(EndIP);
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
}
@@ -2483,7 +2483,7 @@
// Add a termination to our block and check that it is internally consistent.
Builder.restoreIP(EndIP);
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
CallInst *InitCall = nullptr;
@@ -2773,7 +2773,7 @@
/*IsDependSource=*/true));
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
AllocaInst *AllocInst = dyn_cast<AllocaInst>(&BB->front());
@@ -2858,7 +2858,7 @@
/*IsDependSource=*/false));
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
AllocaInst *AllocInst = dyn_cast<AllocaInst>(&BB->front());
@@ -2953,7 +2953,7 @@
OMPBuilder.createOrderedThreadsSimd(Builder, BodyGenCB, FiniCB, true));
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
EXPECT_NE(EntryBB->getTerminator(), nullptr);
@@ -3024,7 +3024,7 @@
OMPBuilder.createOrderedThreadsSimd(Builder, BodyGenCB, FiniCB, false));
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
EXPECT_NE(EntryBB->getTerminator(), nullptr);
@@ -3307,7 +3307,7 @@
EXPECT_EQ(StoreofAtomic->getPointerOperand(), VVal);
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
}
@@ -3355,7 +3355,7 @@
EXPECT_EQ(StoreofAtomic->getValueOperand(), AtomicLoad);
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
}
@@ -3389,7 +3389,7 @@
EXPECT_TRUE(StoreofAtomic->isAtomic());
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
}
@@ -3429,7 +3429,7 @@
EXPECT_EQ(StoreofAtomic->getValueOperand(), ValToWrite);
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
}
@@ -3496,7 +3496,7 @@
EXPECT_EQ(UpdateTemp, Ld->getPointerOperand());
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
}
@@ -3562,7 +3562,7 @@
EXPECT_NE(Ld, nullptr);
EXPECT_EQ(UpdateTemp, Ld->getPointerOperand());
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
}
@@ -3629,7 +3629,7 @@
EXPECT_EQ(UpdateTemp, Ld->getPointerOperand());
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
}
@@ -3679,7 +3679,7 @@
EXPECT_EQ(St->getPointerOperand(), VVal);
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
}
@@ -3739,7 +3739,7 @@
EXPECT_EQ(AXCHG->getNewValOperand(), D);
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
}
@@ -3990,7 +3990,7 @@
EXPECT_EQ(Store8->getValueOperand(), Sel2);
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
EXPECT_FALSE(verifyModule(*M, &errs()));
}
@@ -4225,7 +4225,7 @@
Builder.restoreIP(AfterIP);
Builder.CreateRetVoid();
- OMPBuilder.finalize(F);
+ OMPBuilder.finalizeFunction(F);
// The IR must be valid.
EXPECT_FALSE(verifyModule(*M));
@@ -4476,7 +4476,7 @@
Builder.restoreIP(AfterIP);
Builder.CreateRetVoid();
- OMPBuilder.finalize(F);
+ OMPBuilder.finalizeFunction(F);
// The IR must be valid.
EXPECT_FALSE(verifyModule(*M));
@@ -5092,7 +5092,7 @@
OpenMPIRBuilder::LocationDescription OmpLoc({Builder.saveIP(), DL});
Builder.restoreIP(OMPBuilder.createTarget(OmpLoc, Builder.saveIP(), EntryInfo,
-1, -1, Inputs, BodyGenCB));
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
Builder.CreateRetVoid();
// Check the outlined call
@@ -5146,7 +5146,7 @@
OMPBuilder.createTarget(Loc, EntryIP, EntryInfo, /*NumTeams=*/-1,
/*NumThreads=*/-1, CapturedArgs, BodyGenCB));
Builder.CreateRetVoid();
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
// Check outlined function
EXPECT_FALSE(verifyModule(*M, &errs()));
@@ -5260,7 +5260,7 @@
Builder.restoreIP(OMPBuilder.createTask(
Loc, InsertPointTy(AllocaBB, AllocaBB->getFirstInsertionPt()),
BodyGenCB));
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
Builder.CreateRetVoid();
EXPECT_FALSE(verifyModule(*M, &errs()));
@@ -5352,7 +5352,7 @@
Builder.restoreIP(OMPBuilder.createTask(
Loc, InsertPointTy(AllocaBB, AllocaBB->getFirstInsertionPt()),
BodyGenCB));
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
Builder.CreateRetVoid();
EXPECT_FALSE(verifyModule(*M, &errs()));
@@ -5372,7 +5372,7 @@
Builder.restoreIP(OMPBuilder.createTask(
Loc, InsertPointTy(AllocaBB, AllocaBB->getFirstInsertionPt()), BodyGenCB,
/*Tied=*/false));
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
Builder.CreateRetVoid();
// Check for the `Tied` argument
@@ -5408,7 +5408,7 @@
Builder.restoreIP(OMPBuilder.createTask(
Loc, InsertPointTy(AllocaBB, AllocaBB->getFirstInsertionPt()), BodyGenCB,
/*Tied=*/false, /*Final*/ nullptr, /*IfCondition*/ nullptr, DDS));
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
Builder.CreateRetVoid();
// Check for the `NumDeps` argument
@@ -5476,7 +5476,7 @@
OpenMPIRBuilder::LocationDescription Loc(Builder.saveIP(), DL);
Builder.restoreIP(OMPBuilder.createTask(Loc, AllocaIP, BodyGenCB,
/*Tied=*/false, Final));
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
Builder.CreateRetVoid();
// Check for the `Tied` argument
@@ -5530,7 +5530,7 @@
Builder.restoreIP(OMPBuilder.createTask(Loc, AllocaIP, BodyGenCB,
/*Tied=*/false, /*Final=*/nullptr,
IfCondition));
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
Builder.CreateRetVoid();
EXPECT_FALSE(verifyModule(*M, &errs()));
@@ -5621,7 +5621,7 @@
Builder.restoreIP(OMPBuilder.createTaskgroup(
Loc, InsertPointTy(AllocaBB, AllocaBB->getFirstInsertionPt()),
BodyGenCB));
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
Builder.CreateRetVoid();
EXPECT_FALSE(verifyModule(*M, &errs()));
@@ -5714,7 +5714,7 @@
Builder.restoreIP(OMPBuilder.createTaskgroup(
Loc, InsertPointTy(AllocaBB, AllocaBB->getFirstInsertionPt()),
BodyGenCB));
- OMPBuilder.finalize();
+ OMPBuilder.finalizeModule();
Builder.CreateRetVoid();
EXPECT_FALSE(verifyModule(*M, &errs()));
Index: llvm/lib/Transforms/IPO/OpenMPOpt.cpp
===================================================================
--- llvm/lib/Transforms/IPO/OpenMPOpt.cpp
+++ llvm/lib/Transforms/IPO/OpenMPOpt.cpp
@@ -1222,7 +1222,7 @@
BranchInst::Create(AfterBB, AfterIP.getBlock());
// Perform the actual outlining.
- OMPInfoCache.OMPBuilder.finalize(OriginalFn);
+ OMPInfoCache.OMPBuilder.finalizeFunction(OriginalFn);
Function *OutlinedFn = MergableCIs.front()->getCaller();
Index: llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
===================================================================
--- llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -502,7 +502,7 @@
loadOffloadInfoMetadata(*M.get());
}
-void OpenMPIRBuilder::finalize(Function *Fn) {
+void OpenMPIRBuilder::finalizeFunction(Function *Fn) {
SmallPtrSet<BasicBlock *, 32> ParallelRegionBlockSet;
SmallVector<BasicBlock *, 32> Blocks;
SmallVector<OutlineInfo, 16> DeferredOutlines;
@@ -589,7 +589,10 @@
// Remove work items that have been completed.
OutlineInfos = std::move(DeferredOutlines);
+}
+void OpenMPIRBuilder::finalizeModule() {
+ finalizeFunction();
EmitMetadataErrorReportFunctionTy &&ErrorReportFn =
[](EmitMetadataErrorKind Kind,
const TargetRegionEntryInfo &EntryInfo) -> void {
@@ -5584,7 +5587,7 @@
auto &&GetMDString = [&C](StringRef V) { return MDString::get(C, V); };
// Create the offloading info metadata node.
- NamedMDNode *MD = M.getOrInsertNamedMetadata("omp_offload.info");
+ NamedMDNode *MD = M.getOrInsertNamedMetadata(ompOffloadInfoName);
auto &&TargetRegionMetadataEmitter =
[&C, MD, &OrderedEntries, &GetMDInt, &GetMDString](
const TargetRegionEntryInfo &EntryInfo,
Index: llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
===================================================================
--- llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -448,10 +448,14 @@
void setConfig(OpenMPIRBuilderConfig C) { Config = C; }
- /// Finalize the underlying module, e.g., by outlining regions.
+ /// Finalize the underlying function, e.g., by outlining regions.
/// \param Fn The function to be finalized. If not used,
/// all functions are finalized.
- void finalize(Function *Fn = nullptr);
+ void finalizeFunction(Function *Fn = nullptr);
+
+ /// Finalize the underlying module. Finalize all functions and create
+ /// offload metadata for the module
+ void finalizeModule();
/// Add attributes known for \p FnID to \p Fn.
void addAttributes(omp::RuntimeFunction FnID, Function &Fn);
Index: clang/test/OpenMP/irbuilder_omp_offload_metadata.c
===================================================================
--- /dev/null
+++ clang/test/OpenMP/irbuilder_omp_offload_metadata.c
@@ -0,0 +1,16 @@
+// This test checks if OpenMPIRBuilder generates the same number of omp offload
+// info nodes as Clang does. The wrong number of metadata nodes can provide
+// miscompilation of the device code for enabled OpenMPIRBuilder
+// RUN: %clang_cc1 -triple x86_64--unknown-linux-gnu -emit-llvm -fopenmp -fopenmp-enable-irbuilder -fopenmp-targets=amdgcn-amd-amdhsa -faddrsig %s -o - | FileCheck --check-prefix BUILDER %s
+// RUN: %clang_cc1 -triple x86_64--unknown-linux-gnu -emit-llvm -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -faddrsig %s -o - | FileCheck --check-prefix NOBUILDER %s
+
+void omp_offload_metadata_irbuilder_test() {
+int a[256];
+#pragma omp target parallel for
+ for (int i = 0; i < 256; i++) {
+ a[i] = i;
+ }
+}
+
+//BUILDER: !omp_offload.info = !{!{{[0-9]+}}}
+//NOBUILDER: !omp_offload.info = !{!{{[0-9]+}}}
Index: clang/lib/CodeGen/CodeGenFunction.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenFunction.cpp
+++ clang/lib/CodeGen/CodeGenFunction.cpp
@@ -96,7 +96,7 @@
// time of the CodeGenModule, because we have to ensure the IR has not yet
// been "emitted" to the outside, thus, modifications are still sensible.
if (CGM.getLangOpts().OpenMPIRBuilder && CurFn)
- CGM.getOpenMPRuntime().getOMPBuilder().finalize(CurFn);
+ CGM.getOpenMPRuntime().getOMPBuilder().finalizeFunction(CurFn);
}
// Map the LangOption for exception behavior into
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
