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 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: 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: 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 ParallelRegionBlockSet; SmallVector Blocks; SmallVector 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/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/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 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 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 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(&BB->front()); @@ -2858,7 +2858,7 @@ /*IsDependSource=*/false)); Builder.CreateRetVoid(); - OMPBuilder.finalize(); + OMPBuilder.finalizeModule(); EXPECT_FALSE(verifyModule(*M, &errs())); AllocaInst *AllocInst = dyn_cast(&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: 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) {