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 &region) {
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
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

Reply via email to