alok updated this revision to Diff 428322.
alok added a comment.
Re-based.
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D124982/new/
https://reviews.llvm.org/D124982
Files:
clang/lib/CodeGen/CGDebugInfo.cpp
clang/lib/CodeGen/CGDebugInfo.h
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/lib/CodeGen/CGOpenMPRuntime.h
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
clang/lib/CodeGen/CGStmtOpenMP.cpp
clang/lib/CodeGen/CodeGenFunction.cpp
clang/test/OpenMP/debug_containing_scope.c
llvm/lib/CodeGen/AsmPrinter/DwarfCompileUnit.cpp
llvm/lib/CodeGen/AsmPrinter/DwarfUnit.cpp
llvm/lib/CodeGen/AsmPrinter/DwarfUnit.h
Index: llvm/lib/CodeGen/AsmPrinter/DwarfUnit.h
===================================================================
--- llvm/lib/CodeGen/AsmPrinter/DwarfUnit.h
+++ llvm/lib/CodeGen/AsmPrinter/DwarfUnit.h
@@ -237,6 +237,7 @@
DIE *getOrCreateNameSpace(const DINamespace *NS);
DIE *getOrCreateModule(const DIModule *M);
DIE *getOrCreateSubprogramDIE(const DISubprogram *SP, bool Minimal = false);
+ DIE *getOrCreateLexicalScopeDIE(const DILexicalBlock *LS);
void applySubprogramAttributes(const DISubprogram *SP, DIE &SPDie,
bool SkipSPAttributes = false);
Index: llvm/lib/CodeGen/AsmPrinter/DwarfUnit.cpp
===================================================================
--- llvm/lib/CodeGen/AsmPrinter/DwarfUnit.cpp
+++ llvm/lib/CodeGen/AsmPrinter/DwarfUnit.cpp
@@ -554,6 +554,8 @@
return getOrCreateSubprogramDIE(SP);
if (auto *M = dyn_cast<DIModule>(Context))
return getOrCreateModule(M);
+ if (auto *LS = dyn_cast<DILexicalBlock>(Context))
+ return getOrCreateLexicalScopeDIE(LS);
return getDIE(Context);
}
@@ -1181,6 +1183,17 @@
return &SPDie;
}
+DIE *DwarfUnit::getOrCreateLexicalScopeDIE(const DILexicalBlock *LS) {
+ DIE *ContextDIE = getOrCreateContextDIE(LS->getScope());
+
+ if (DIE *LSDie = getDIE(LS))
+ return LSDie;
+
+ DIE &LSDie = createAndAddDIE(dwarf::DW_TAG_lexical_block, *ContextDIE, LS);
+
+ return &LSDie;
+}
+
bool DwarfUnit::applySubprogramDefinitionAttributes(const DISubprogram *SP,
DIE &SPDie, bool Minimal) {
DIE *DeclDie = nullptr;
Index: llvm/lib/CodeGen/AsmPrinter/DwarfCompileUnit.cpp
===================================================================
--- llvm/lib/CodeGen/AsmPrinter/DwarfCompileUnit.cpp
+++ llvm/lib/CodeGen/AsmPrinter/DwarfCompileUnit.cpp
@@ -686,6 +686,7 @@
return nullptr;
auto ScopeDIE = DIE::get(DIEValueAllocator, dwarf::DW_TAG_lexical_block);
+ insertDIE(Scope->getScopeNode(), ScopeDIE);
if (Scope->isAbstractScope())
return ScopeDIE;
Index: clang/test/OpenMP/debug_containing_scope.c
===================================================================
--- /dev/null
+++ clang/test/OpenMP/debug_containing_scope.c
@@ -0,0 +1,68 @@
+// This testcase checks parent child relationship for OpenMP generated
+// functions.
+
+// REQUIRES: x86_64-linux
+
+// RUN: %clang_cc1 -debug-info-kind=constructor -DSHARED -x c -verify -triple x86_64-pc-linux-gnu -fopenmp -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK
+// expected-no-diagnostics
+
+// CHECK-DAG: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @8, i32 5, ptr @[[OUTLINED2:.omp_outlined..[0-9]]]
+// CHECK-LABEL: distinct !DICompileUnit
+// CHECK-DAG: [[FOO:![0-9]+]] = distinct !DISubprogram(name: "foo",
+// CHECK-DAG: [[LEX1:![0-9]+]] = distinct !DILexicalBlock(scope: [[FOO]]
+// CHECK-DAG: [[LEX2:![0-9]+]] = distinct !DILexicalBlock(scope: [[LEX1]]
+// CHECK-DAG: [[LEX3:![0-9]+]] = distinct !DILexicalBlock(scope: [[LEX2]]
+// CHECK-DAG: !DISubprogram(linkageName: ".omp_task_privates_map.", scope: [[LEX3]]
+// CHECK-DAG: !DISubprogram(linkageName: ".omp_task_entry.", scope: [[LEX3]]
+// CHECK-DAG: !DISubprogram(name: ".omp_outlined.", scope: [[LEX2]]
+// CHECK-DAG: !DISubprogram(name: ".omp_outlined._debug__", scope: [[LEX2]]
+// CHECK-DAG: !DISubprogram(name: "[[OUTLINED2]]", scope: [[LEX2]]
+
+extern int printf(const char *, ...);
+extern int rand(void);
+
+int global_var1;
+int global_var2 = 99;
+int foo(int n) {
+ int same_var = 5;
+ int other_var = 21;
+ int share = 9, priv, i;
+ global_var1 = 99;
+
+ if (n < 2)
+ return n;
+ else {
+ int same_var = rand() % 5;
+ int local_var = 31;
+#pragma omp task shared(share) private(priv)
+ {
+ priv = n;
+ printf("share = %d\n", share);
+ printf("global_var1 = %d\n", global_var1);
+ printf("global_var2 = %d\n", global_var2);
+ printf("same_var = %d\n", same_var);
+ printf("other_var = %d\n", other_var);
+ printf("local_var = %d\n", local_var);
+ share = priv + foo(n - 1);
+ }
+#pragma omp taskwait
+
+#pragma omp parallel for
+ for (i = 0; i < n; i++) {
+ share += i;
+ printf("share = %d\n", share);
+ printf("global_var1 = %d\n", global_var1);
+ printf("global_var2 = %d\n", global_var2);
+ printf("same_var = %d\n", same_var);
+ printf("other_var = %d\n", other_var);
+ printf("local_var = %d\n", local_var);
+ }
+ return share;
+ }
+}
+
+int main() {
+ int n = 10;
+ printf("foo(%d) = %d\n", n, foo(n));
+ return 0;
+}
Index: clang/lib/CodeGen/CodeGenFunction.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenFunction.cpp
+++ clang/lib/CodeGen/CodeGenFunction.cpp
@@ -1021,7 +1021,7 @@
// convention.
DI->emitFunctionStart(GD, Loc, StartLoc,
DI->getFunctionType(FD, RetTy, Args), CurFn,
- CurFuncIsThunk);
+ CurFuncIsThunk, ParentCGF);
}
if (ShouldInstrumentFunction()) {
Index: clang/lib/CodeGen/CGStmtOpenMP.cpp
===================================================================
--- clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -648,6 +648,7 @@
/*RegisterCastedArgsOnly=*/true,
CapturedStmtInfo->getHelperName(), Loc);
CodeGenFunction WrapperCGF(CGM, /*suppressNewContext=*/true);
+ WrapperCGF.ParentCGF = ParentCGF;
WrapperCGF.CapturedStmtInfo = CapturedStmtInfo;
Args.clear();
LocalAddrs.clear();
@@ -1544,7 +1545,8 @@
llvm::Value *NumThreads = nullptr;
llvm::Function *OutlinedFn =
CGF.CGM.getOpenMPRuntime().emitParallelOutlinedFunction(
- S, *CS->getCapturedDecl()->param_begin(), InnermostKind, CodeGen);
+ S, *CS->getCapturedDecl()->param_begin(), InnermostKind, CodeGen,
+ CGF);
if (const auto *NumThreadsClause = S.getSingleClause<OMPNumThreadsClause>()) {
CodeGenFunction::RunCleanupsScope NumThreadsScope(CGF);
NumThreads = CGF.EmitScalarExpr(NumThreadsClause->getNumThreads(),
@@ -4853,7 +4855,7 @@
};
llvm::Function *OutlinedFn = CGM.getOpenMPRuntime().emitTaskOutlinedFunction(
S, *I, *PartId, *TaskT, S.getDirectiveKind(), CodeGen, Data.Tied,
- Data.NumberOfParts);
+ Data.NumberOfParts, *this);
OMPLexicalScope Scope(*this, S, llvm::None,
!isOpenMPParallelDirective(S.getDirectiveKind()) &&
!isOpenMPSimdDirective(S.getDirectiveKind()));
@@ -5016,7 +5018,7 @@
};
llvm::Function *OutlinedFn = CGM.getOpenMPRuntime().emitTaskOutlinedFunction(
S, *I, *PartId, *TaskT, S.getDirectiveKind(), CodeGen, /*Tied=*/true,
- Data.NumberOfParts);
+ Data.NumberOfParts, *this);
llvm::APInt TrueOrFalse(32, S.hasClausesOfKind<OMPNowaitClause>() ? 1 : 0);
IntegerLiteral IfCond(getContext(), TrueOrFalse,
getContext().getIntTypeForBitwidth(32, /*Signed=*/0),
@@ -6430,8 +6432,8 @@
CGM.getMangledName(GlobalDecl(cast<FunctionDecl>(CGF.CurFuncDecl)));
// Emit target region as a standalone region.
- CGM.getOpenMPRuntime().emitTargetOutlinedFunction(S, ParentName, Fn, FnID,
- IsOffloadEntry, CodeGen);
+ CGM.getOpenMPRuntime().emitTargetOutlinedFunction(
+ S, ParentName, Fn, FnID, IsOffloadEntry, CodeGen, &CGF);
OMPLexicalScope Scope(CGF, S, OMPD_task);
auto &&SizeEmitter =
[IsOffloadEntry](CodeGenFunction &CGF,
@@ -6492,7 +6494,8 @@
const CapturedStmt *CS = S.getCapturedStmt(OMPD_teams);
llvm::Function *OutlinedFn =
CGF.CGM.getOpenMPRuntime().emitTeamsOutlinedFunction(
- S, *CS->getCapturedDecl()->param_begin(), InnermostKind, CodeGen);
+ S, *CS->getCapturedDecl()->param_begin(), InnermostKind, CodeGen,
+ CGF);
const auto *NT = S.getSingleClause<OMPNumTeamsClause>();
const auto *TL = S.getSingleClause<OMPThreadLimitClause>();
Index: clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
+++ clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
@@ -79,12 +79,15 @@
/// \param OutlinedFn Outlined function value to be defined by this call.
/// \param OutlinedFnID Outlined function ID value to be defined by this call.
/// \param IsOffloadEntry True if the outlined function is an offload entry.
+ /// \param CodeGen Object containing the target statements.
+ /// \param ParentCGF The CGF of parent/containing function.
/// An outlined function may not be an entry if, e.g. the if clause always
/// evaluates to false.
void emitNonSPMDKernel(const OMPExecutableDirective &D, StringRef ParentName,
llvm::Function *&OutlinedFn,
llvm::Constant *&OutlinedFnID, bool IsOffloadEntry,
- const RegionCodeGenTy &CodeGen);
+ const RegionCodeGenTy &CodeGen,
+ CodeGenFunction *ParentCGF = nullptr);
/// Emit outlined function specialized for the Single Program
/// Multiple Data programming model for applicable target directives on the
@@ -95,12 +98,14 @@
/// \param OutlinedFnID Outlined function ID value to be defined by this call.
/// \param IsOffloadEntry True if the outlined function is an offload entry.
/// \param CodeGen Object containing the target statements.
+ /// \param ParentCGF The CGF of parent/containing function.
/// An outlined function may not be an entry if, e.g. the if clause always
/// evaluates to false.
void emitSPMDKernel(const OMPExecutableDirective &D, StringRef ParentName,
llvm::Function *&OutlinedFn,
llvm::Constant *&OutlinedFnID, bool IsOffloadEntry,
- const RegionCodeGenTy &CodeGen);
+ const RegionCodeGenTy &CodeGen,
+ CodeGenFunction *ParentCGF = nullptr);
/// Emit outlined function for 'target' directive on the NVPTX
/// device.
@@ -109,14 +114,16 @@
/// \param OutlinedFn Outlined function value to be defined by this call.
/// \param OutlinedFnID Outlined function ID value to be defined by this call.
/// \param IsOffloadEntry True if the outlined function is an offload entry.
+ /// \param CodeGen Object containing the target statements.
+ /// \param ParentCGF The CGF of parent/containing function.
/// An outlined function may not be an entry if, e.g. the if clause always
/// evaluates to false.
- void emitTargetOutlinedFunction(const OMPExecutableDirective &D,
- StringRef ParentName,
- llvm::Function *&OutlinedFn,
- llvm::Constant *&OutlinedFnID,
- bool IsOffloadEntry,
- const RegionCodeGenTy &CodeGen) override;
+ void
+ emitTargetOutlinedFunction(const OMPExecutableDirective &D,
+ StringRef ParentName, llvm::Function *&OutlinedFn,
+ llvm::Constant *&OutlinedFnID, bool IsOffloadEntry,
+ const RegionCodeGenTy &CodeGen,
+ CodeGenFunction *ParentCGF = nullptr) override;
/// Emits code for parallel or serial call of the \a OutlinedFn with
/// variables captured in a record which address is stored in \a
@@ -215,11 +222,11 @@
/// \param InnermostKind Kind of innermost directive (for simple directives it
/// is a directive itself, for combined - its innermost directive).
/// \param CodeGen Code generation sequence for the \a D directive.
- llvm::Function *
- emitParallelOutlinedFunction(const OMPExecutableDirective &D,
- const VarDecl *ThreadIDVar,
- OpenMPDirectiveKind InnermostKind,
- const RegionCodeGenTy &CodeGen) override;
+ /// \param CodeGenFunction of outlining/containing function.
+ llvm::Function *emitParallelOutlinedFunction(
+ const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
+ OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen,
+ CodeGenFunction &ParentCGF) override;
/// Emits inlined function for the specified OpenMP teams
// directive.
@@ -230,11 +237,11 @@
/// \param InnermostKind Kind of innermost directive (for simple directives it
/// is a directive itself, for combined - its innermost directive).
/// \param CodeGen Code generation sequence for the \a D directive.
- llvm::Function *
- emitTeamsOutlinedFunction(const OMPExecutableDirective &D,
- const VarDecl *ThreadIDVar,
- OpenMPDirectiveKind InnermostKind,
- const RegionCodeGenTy &CodeGen) override;
+ /// \param CodeGenFunction of outlining/containing function.
+ llvm::Function *emitTeamsOutlinedFunction(
+ const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
+ OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen,
+ CodeGenFunction &ParentCGF) override;
/// Emits code for teams call of the \a OutlinedFn with
/// variables captured in a record which address is stored in \a
Index: clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
@@ -1008,11 +1008,12 @@
}
void CGOpenMPRuntimeGPU::emitNonSPMDKernel(const OMPExecutableDirective &D,
- StringRef ParentName,
- llvm::Function *&OutlinedFn,
- llvm::Constant *&OutlinedFnID,
- bool IsOffloadEntry,
- const RegionCodeGenTy &CodeGen) {
+ StringRef ParentName,
+ llvm::Function *&OutlinedFn,
+ llvm::Constant *&OutlinedFnID,
+ bool IsOffloadEntry,
+ const RegionCodeGenTy &CodeGen,
+ CodeGenFunction *ParentCGF) {
ExecutionRuntimeModesRAII ModeRAII(CurrentExecutionMode);
EntryFunctionState EST;
WrapperFunctionsMap.clear();
@@ -1041,7 +1042,7 @@
CodeGen.setAction(Action);
IsInTTDRegion = true;
emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
- IsOffloadEntry, CodeGen);
+ IsOffloadEntry, CodeGen, ParentCGF);
IsInTTDRegion = false;
}
@@ -1065,11 +1066,12 @@
}
void CGOpenMPRuntimeGPU::emitSPMDKernel(const OMPExecutableDirective &D,
- StringRef ParentName,
- llvm::Function *&OutlinedFn,
- llvm::Constant *&OutlinedFnID,
- bool IsOffloadEntry,
- const RegionCodeGenTy &CodeGen) {
+ StringRef ParentName,
+ llvm::Function *&OutlinedFn,
+ llvm::Constant *&OutlinedFnID,
+ bool IsOffloadEntry,
+ const RegionCodeGenTy &CodeGen,
+ CodeGenFunction *ParentCGF) {
ExecutionRuntimeModesRAII ModeRAII(
CurrentExecutionMode, RequiresFullRuntime,
CGM.getLangOpts().OpenMPCUDAForceFullRuntime ||
@@ -1098,7 +1100,7 @@
CodeGen.setAction(Action);
IsInTTDRegion = true;
emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
- IsOffloadEntry, CodeGen);
+ IsOffloadEntry, CodeGen, ParentCGF);
IsInTTDRegion = false;
}
@@ -1149,7 +1151,8 @@
void CGOpenMPRuntimeGPU::emitTargetOutlinedFunction(
const OMPExecutableDirective &D, StringRef ParentName,
llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,
- bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) {
+ bool IsOffloadEntry, const RegionCodeGenTy &CodeGen,
+ CodeGenFunction *ParentCGF) {
if (!IsOffloadEntry) // Nothing to do.
return;
@@ -1158,10 +1161,10 @@
bool Mode = supportsSPMDExecutionMode(CGM.getContext(), D);
if (Mode)
emitSPMDKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,
- CodeGen);
+ CodeGen, ParentCGF);
else
emitNonSPMDKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,
- CodeGen);
+ CodeGen, ParentCGF);
setPropertyExecutionMode(CGM, OutlinedFn->getName(), Mode);
}
@@ -1238,7 +1241,8 @@
llvm::Function *CGOpenMPRuntimeGPU::emitParallelOutlinedFunction(
const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
- OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen) {
+ OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen,
+ CodeGenFunction &ParentCGF) {
// Emit target region as a standalone region.
class NVPTXPrePostActionTy : public PrePostActionTy {
bool &IsInParallelRegion;
@@ -1262,7 +1266,7 @@
IsInTargetMasterThreadRegion = false;
auto *OutlinedFun =
cast<llvm::Function>(CGOpenMPRuntime::emitParallelOutlinedFunction(
- D, ThreadIDVar, InnermostKind, CodeGen));
+ D, ThreadIDVar, InnermostKind, CodeGen, ParentCGF));
IsInTargetMasterThreadRegion = PrevIsInTargetMasterThreadRegion;
IsInTTDRegion = PrevIsInTTDRegion;
if (getExecutionMode() != CGOpenMPRuntimeGPU::EM_SPMD &&
@@ -1315,7 +1319,8 @@
llvm::Function *CGOpenMPRuntimeGPU::emitTeamsOutlinedFunction(
const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
- OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen) {
+ OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen,
+ CodeGenFunction &ParentCGF) {
SourceLocation Loc = D.getBeginLoc();
const RecordDecl *GlobalizedRD = nullptr;
@@ -1376,7 +1381,7 @@
} Action(Loc, GlobalizedRD, MappedDeclsFields);
CodeGen.setAction(Action);
llvm::Function *OutlinedFun = CGOpenMPRuntime::emitTeamsOutlinedFunction(
- D, ThreadIDVar, InnermostKind, CodeGen);
+ D, ThreadIDVar, InnermostKind, CodeGen, ParentCGF);
return OutlinedFun;
}
Index: clang/lib/CodeGen/CGOpenMPRuntime.h
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.h
+++ clang/lib/CodeGen/CGOpenMPRuntime.h
@@ -328,14 +328,14 @@
/// \param OutlinedFnID Outlined function ID value to be defined by this call.
/// \param IsOffloadEntry True if the outlined function is an offload entry.
/// \param CodeGen Lambda codegen specific to an accelerator device.
+ /// \param ParentCGF The CGF of parent/containing function.
/// An outlined function may not be an entry if, e.g. the if clause always
/// evaluates to false.
- virtual void emitTargetOutlinedFunctionHelper(const OMPExecutableDirective &D,
- StringRef ParentName,
- llvm::Function *&OutlinedFn,
- llvm::Constant *&OutlinedFnID,
- bool IsOffloadEntry,
- const RegionCodeGenTy &CodeGen);
+ virtual void emitTargetOutlinedFunctionHelper(
+ const OMPExecutableDirective &D, StringRef ParentName,
+ llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,
+ bool IsOffloadEntry, const RegionCodeGenTy &CodeGen,
+ CodeGenFunction *ParentCGF = nullptr);
/// Emits object of ident_t type with info for source location.
/// \param Flags Flags for OpenMP location.
@@ -964,9 +964,11 @@
/// \param InnermostKind Kind of innermost directive (for simple directives it
/// is a directive itself, for combined - its innermost directive).
/// \param CodeGen Code generation sequence for the \a D directive.
+ /// \param CodeGenFunction of outlining/containing function.
virtual llvm::Function *emitParallelOutlinedFunction(
const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
- OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen);
+ OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen,
+ CodeGenFunction &ParentCGF);
/// Emits outlined function for the specified OpenMP teams directive
/// \a D. This outlined function has type void(*)(kmp_int32 *ThreadID,
@@ -976,9 +978,11 @@
/// \param InnermostKind Kind of innermost directive (for simple directives it
/// is a directive itself, for combined - its innermost directive).
/// \param CodeGen Code generation sequence for the \a D directive.
+ /// \param CodeGenFunction of outlining/containing function.
virtual llvm::Function *emitTeamsOutlinedFunction(
const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
- OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen);
+ OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen,
+ CodeGenFunction &ParentCGF);
/// Emits outlined function for the OpenMP task directive \a D. This
/// outlined function has type void(*)(kmp_int32 ThreadID, struct task_t*
@@ -994,12 +998,13 @@
/// \param Tied true if task is generated for tied task, false otherwise.
/// \param NumberOfParts Number of parts in untied task. Ignored for tied
/// tasks.
+ /// \param CodeGenFunction of outlining/containing function.
///
virtual llvm::Function *emitTaskOutlinedFunction(
const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
const VarDecl *PartIDVar, const VarDecl *TaskTVar,
OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen,
- bool Tied, unsigned &NumberOfParts);
+ bool Tied, unsigned &NumberOfParts, CodeGenFunction &ParentCGF);
/// Cleans up references to the objects in finished function.
///
@@ -1575,6 +1580,7 @@
/// \param OutlinedFnID Outlined function ID value to be defined by this call.
/// \param IsOffloadEntry True if the outlined function is an offload entry.
/// \param CodeGen Code generation sequence for the \a D directive.
+ /// \param ParentCGF The CGF of parent/containing function.
/// An outlined function may not be an entry if, e.g. the if clause always
/// evaluates to false.
virtual void emitTargetOutlinedFunction(const OMPExecutableDirective &D,
@@ -1582,7 +1588,8 @@
llvm::Function *&OutlinedFn,
llvm::Constant *&OutlinedFnID,
bool IsOffloadEntry,
- const RegionCodeGenTy &CodeGen);
+ const RegionCodeGenTy &CodeGen,
+ CodeGenFunction *ParentCGF = nullptr);
/// Emit the target offloading code associated with \a D. The emitted
/// code attempts offloading the execution to the device, an the event of
@@ -1942,11 +1949,11 @@
/// \param InnermostKind Kind of innermost directive (for simple directives it
/// is a directive itself, for combined - its innermost directive).
/// \param CodeGen Code generation sequence for the \a D directive.
- llvm::Function *
- emitParallelOutlinedFunction(const OMPExecutableDirective &D,
- const VarDecl *ThreadIDVar,
- OpenMPDirectiveKind InnermostKind,
- const RegionCodeGenTy &CodeGen) override;
+ /// \param CodeGenFunction of outlining/containing function.
+ llvm::Function *emitParallelOutlinedFunction(
+ const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
+ OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen,
+ CodeGenFunction &ParentCGF) override;
/// Emits outlined function for the specified OpenMP teams directive
/// \a D. This outlined function has type void(*)(kmp_int32 *ThreadID,
@@ -1956,11 +1963,11 @@
/// \param InnermostKind Kind of innermost directive (for simple directives it
/// is a directive itself, for combined - its innermost directive).
/// \param CodeGen Code generation sequence for the \a D directive.
- llvm::Function *
- emitTeamsOutlinedFunction(const OMPExecutableDirective &D,
- const VarDecl *ThreadIDVar,
- OpenMPDirectiveKind InnermostKind,
- const RegionCodeGenTy &CodeGen) override;
+ /// \param CodeGenFunction of outlining/containing function.
+ llvm::Function *emitTeamsOutlinedFunction(
+ const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
+ OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen,
+ CodeGenFunction &ParentCGF) override;
/// Emits outlined function for the OpenMP task directive \a D. This
/// outlined function has type void(*)(kmp_int32 ThreadID, struct task_t*
@@ -1976,12 +1983,13 @@
/// \param Tied true if task is generated for tied task, false otherwise.
/// \param NumberOfParts Number of parts in untied task. Ignored for tied
/// tasks.
+ /// \param CodeGenFunction of outlining/containing function.
///
llvm::Function *emitTaskOutlinedFunction(
const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
const VarDecl *PartIDVar, const VarDecl *TaskTVar,
OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen,
- bool Tied, unsigned &NumberOfParts) override;
+ bool Tied, unsigned &NumberOfParts, CodeGenFunction &ParentCGF) override;
/// Emits code for parallel or serial call of the \a OutlinedFn with
/// variables captured in a record which address is stored in \a
@@ -2415,14 +2423,15 @@
/// \param OutlinedFnID Outlined function ID value to be defined by this call.
/// \param IsOffloadEntry True if the outlined function is an offload entry.
/// \param CodeGen Code generation sequence for the \a D directive.
+ /// \param ParentCGF The CGF of parent/containing function.
/// An outlined function may not be an entry if, e.g. the if clause always
/// evaluates to false.
- void emitTargetOutlinedFunction(const OMPExecutableDirective &D,
- StringRef ParentName,
- llvm::Function *&OutlinedFn,
- llvm::Constant *&OutlinedFnID,
- bool IsOffloadEntry,
- const RegionCodeGenTy &CodeGen) override;
+ void
+ emitTargetOutlinedFunction(const OMPExecutableDirective &D,
+ StringRef ParentName, llvm::Function *&OutlinedFn,
+ llvm::Constant *&OutlinedFnID, bool IsOffloadEntry,
+ const RegionCodeGenTy &CodeGen,
+ CodeGenFunction *ParentCGF = nullptr) override;
/// Emit the target offloading code associated with \a D. The emitted
/// code attempts offloading the execution to the device, an the event of
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -1232,10 +1232,12 @@
static llvm::Function *emitParallelOrTeamsOutlinedFunction(
CodeGenModule &CGM, const OMPExecutableDirective &D, const CapturedStmt *CS,
const VarDecl *ThreadIDVar, OpenMPDirectiveKind InnermostKind,
- const StringRef OutlinedHelperName, const RegionCodeGenTy &CodeGen) {
+ const StringRef OutlinedHelperName, const RegionCodeGenTy &CodeGen,
+ CodeGenFunction &ParentCGF) {
assert(ThreadIDVar->getType()->isPointerType() &&
"thread id variable must be of type kmp_int32 *");
CodeGenFunction CGF(CGM, true);
+ CGF.ParentCGF = &ParentCGF;
bool HasCancel = false;
if (const auto *OPD = dyn_cast<OMPParallelDirective>(&D))
HasCancel = OPD->hasCancel();
@@ -1268,25 +1270,29 @@
llvm::Function *CGOpenMPRuntime::emitParallelOutlinedFunction(
const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
- OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen) {
+ OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen,
+ CodeGenFunction &ParentCGF) {
const CapturedStmt *CS = D.getCapturedStmt(OMPD_parallel);
return emitParallelOrTeamsOutlinedFunction(
- CGM, D, CS, ThreadIDVar, InnermostKind, getOutlinedHelperName(), CodeGen);
+ CGM, D, CS, ThreadIDVar, InnermostKind, getOutlinedHelperName(), CodeGen,
+ ParentCGF);
}
llvm::Function *CGOpenMPRuntime::emitTeamsOutlinedFunction(
const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
- OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen) {
+ OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen,
+ CodeGenFunction &ParentCGF) {
const CapturedStmt *CS = D.getCapturedStmt(OMPD_teams);
return emitParallelOrTeamsOutlinedFunction(
- CGM, D, CS, ThreadIDVar, InnermostKind, getOutlinedHelperName(), CodeGen);
+ CGM, D, CS, ThreadIDVar, InnermostKind, getOutlinedHelperName(), CodeGen,
+ ParentCGF);
}
llvm::Function *CGOpenMPRuntime::emitTaskOutlinedFunction(
const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
const VarDecl *PartIDVar, const VarDecl *TaskTVar,
OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen,
- bool Tied, unsigned &NumberOfParts) {
+ bool Tied, unsigned &NumberOfParts, CodeGenFunction &ParentCGF) {
auto &&UntiedCodeGen = [this, &D, TaskTVar](CodeGenFunction &CGF,
PrePostActionTy &) {
llvm::Value *ThreadID = getThreadID(CGF, D.getBeginLoc());
@@ -1320,6 +1326,7 @@
HasCancel = TD->hasCancel();
CodeGenFunction CGF(CGM, true);
+ CGF.ParentCGF = &ParentCGF;
CGOpenMPTaskOutlinedRegionInfo CGInfo(*CS, ThreadIDVar, CodeGen,
InnermostKind, HasCancel, Action);
CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo);
@@ -3492,13 +3499,12 @@
/// return 0;
/// }
/// \endcode
-static llvm::Function *
-emitProxyTaskFunction(CodeGenModule &CGM, SourceLocation Loc,
- OpenMPDirectiveKind Kind, QualType KmpInt32Ty,
- QualType KmpTaskTWithPrivatesPtrQTy,
- QualType KmpTaskTWithPrivatesQTy, QualType KmpTaskTQTy,
- QualType SharedsPtrTy, llvm::Function *TaskFunction,
- llvm::Value *TaskPrivatesMap) {
+static llvm::Function *emitProxyTaskFunction(
+ CodeGenModule &CGM, SourceLocation Loc, OpenMPDirectiveKind Kind,
+ QualType KmpInt32Ty, QualType KmpTaskTWithPrivatesPtrQTy,
+ QualType KmpTaskTWithPrivatesQTy, QualType KmpTaskTQTy,
+ QualType SharedsPtrTy, llvm::Function *TaskFunction,
+ llvm::Value *TaskPrivatesMap, CodeGenFunction &ParentCGF) {
ASTContext &C = CGM.getContext();
FunctionArgList Args;
ImplicitParamDecl GtidArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr, KmpInt32Ty,
@@ -3518,6 +3524,7 @@
CGM.SetInternalFunctionAttributes(GlobalDecl(), TaskEntry, TaskEntryFnInfo);
TaskEntry->setDoesNotRecurse();
CodeGenFunction CGF(CGM);
+ CGF.ParentCGF = &ParentCGF;
CGF.StartFunction(GlobalDecl(), KmpInt32Ty, TaskEntry, TaskEntryFnInfo, Args,
Loc, Loc);
@@ -3658,7 +3665,8 @@
static llvm::Value *
emitTaskPrivateMappingFunction(CodeGenModule &CGM, SourceLocation Loc,
const OMPTaskDataTy &Data, QualType PrivatesQTy,
- ArrayRef<PrivateDataTy> Privates) {
+ ArrayRef<PrivateDataTy> Privates,
+ CodeGenFunction &ParentCGF) {
ASTContext &C = CGM.getContext();
FunctionArgList Args;
ImplicitParamDecl TaskPrivatesArg(
@@ -3731,6 +3739,7 @@
TaskPrivatesMap->addFnAttr(llvm::Attribute::AlwaysInline);
}
CodeGenFunction CGF(CGM);
+ CGF.ParentCGF = &ParentCGF;
CGF.StartFunction(GlobalDecl(), C.VoidTy, TaskPrivatesMap,
TaskPrivatesMapFnInfo, Args, Loc, Loc);
@@ -4202,8 +4211,8 @@
std::next(TaskFunction->arg_begin(), 3)->getType();
if (!Privates.empty()) {
auto FI = std::next(KmpTaskTWithPrivatesQTyRD->field_begin());
- TaskPrivatesMap =
- emitTaskPrivateMappingFunction(CGM, Loc, Data, FI->getType(), Privates);
+ TaskPrivatesMap = emitTaskPrivateMappingFunction(
+ CGM, Loc, Data, FI->getType(), Privates, CGF);
TaskPrivatesMap = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
TaskPrivatesMap, TaskPrivatesMapTy);
} else {
@@ -4215,7 +4224,7 @@
llvm::Function *TaskEntry = emitProxyTaskFunction(
CGM, Loc, D.getDirectiveKind(), KmpInt32Ty, KmpTaskTWithPrivatesPtrQTy,
KmpTaskTWithPrivatesQTy, KmpTaskTQTy, SharedsPtrTy, TaskFunction,
- TaskPrivatesMap);
+ TaskPrivatesMap, CGF);
// Build call kmp_task_t * __kmpc_omp_task_alloc(ident_t *, kmp_int32 gtid,
// kmp_int32 flags, size_t sizeof_kmp_task_t, size_t sizeof_shareds,
@@ -6319,7 +6328,8 @@
void CGOpenMPRuntime::emitTargetOutlinedFunction(
const OMPExecutableDirective &D, StringRef ParentName,
llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,
- bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) {
+ bool IsOffloadEntry, const RegionCodeGenTy &CodeGen,
+ CodeGenFunction *ParentCGF) {
assert(!ParentName.empty() && "Invalid target region parent name!");
HasEmittedTargetRegion = true;
SmallVector<std::pair<const Expr *, const Expr *>, 4> Allocators;
@@ -6334,7 +6344,7 @@
OMPUsesAllocatorsActionTy UsesAllocatorAction(Allocators);
CodeGen.setAction(UsesAllocatorAction);
emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
- IsOffloadEntry, CodeGen);
+ IsOffloadEntry, CodeGen, ParentCGF);
}
void CGOpenMPRuntime::emitUsesAllocatorsInit(CodeGenFunction &CGF,
@@ -6391,7 +6401,8 @@
void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(
const OMPExecutableDirective &D, StringRef ParentName,
llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,
- bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) {
+ bool IsOffloadEntry, const RegionCodeGenTy &CodeGen,
+ CodeGenFunction *ParentCGF) {
// Create a unique name for the entry function using the source location
// information of the current target region. The name will be something like:
//
@@ -6418,6 +6429,7 @@
const CapturedStmt &CS = *D.getCapturedStmt(OMPD_target);
CodeGenFunction CGF(CGM, true);
+ CGF.ParentCGF = ParentCGF;
CGOpenMPTargetRegionInfo CGInfo(CS, CodeGen, EntryFnName);
CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo);
@@ -12801,13 +12813,15 @@
llvm::Function *CGOpenMPSIMDRuntime::emitParallelOutlinedFunction(
const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
- OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen) {
+ OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen,
+ CodeGenFunction &ParentCGF) {
llvm_unreachable("Not supported in SIMD-only mode");
}
llvm::Function *CGOpenMPSIMDRuntime::emitTeamsOutlinedFunction(
const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
- OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen) {
+ OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen,
+ CodeGenFunction &ParentCGF) {
llvm_unreachable("Not supported in SIMD-only mode");
}
@@ -12815,7 +12829,7 @@
const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
const VarDecl *PartIDVar, const VarDecl *TaskTVar,
OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen,
- bool Tied, unsigned &NumberOfParts) {
+ bool Tied, unsigned &NumberOfParts, CodeGenFunction &ParentCGF) {
llvm_unreachable("Not supported in SIMD-only mode");
}
@@ -13031,7 +13045,7 @@
void CGOpenMPSIMDRuntime::emitTargetOutlinedFunction(
const OMPExecutableDirective &D, StringRef ParentName,
llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,
- bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) {
+ bool IsOffloadEntry, const RegionCodeGenTy &CodeGen, CodeGenFunction *CGF) {
llvm_unreachable("Not supported in SIMD-only mode");
}
Index: clang/lib/CodeGen/CGDebugInfo.h
===================================================================
--- clang/lib/CodeGen/CGDebugInfo.h
+++ clang/lib/CodeGen/CGDebugInfo.h
@@ -437,9 +437,11 @@
/// start of a new function.
/// \param Loc The location of the function header.
/// \param ScopeLoc The location of the function body.
+ /// \param ParentCGF The CGF of parent/containing function.
void emitFunctionStart(GlobalDecl GD, SourceLocation Loc,
SourceLocation ScopeLoc, QualType FnType,
- llvm::Function *Fn, bool CurFnIsThunk);
+ llvm::Function *Fn, bool CurFnIsThunk,
+ CodeGenFunction *ParentCGF = nullptr);
/// Start a new scope for an inlined function.
void EmitInlineFunctionStart(CGBuilderTy &Builder, GlobalDecl GD);
Index: clang/lib/CodeGen/CGDebugInfo.cpp
===================================================================
--- clang/lib/CodeGen/CGDebugInfo.cpp
+++ clang/lib/CodeGen/CGDebugInfo.cpp
@@ -4044,7 +4044,8 @@
void CGDebugInfo::emitFunctionStart(GlobalDecl GD, SourceLocation Loc,
SourceLocation ScopeLoc, QualType FnType,
- llvm::Function *Fn, bool CurFuncIsThunk) {
+ llvm::Function *Fn, bool CurFuncIsThunk,
+ CodeGenFunction *ParentCGF) {
StringRef Name;
StringRef LinkageName;
@@ -4055,9 +4056,20 @@
llvm::DINode::DIFlags Flags = llvm::DINode::FlagZero;
llvm::DISubprogram::DISPFlags SPFlags = llvm::DISubprogram::SPFlagZero;
- llvm::DIFile *Unit = getOrCreateFile(Loc);
- llvm::DIScope *FDContext = Unit;
llvm::DINodeArray TParamsArray;
+ llvm::DIFile *Unit = getOrCreateFile(Loc);
+ llvm::DIScope *FDContext;
+
+ // Handle Parent Scope if ParentCGF is not NULL
+ if (ParentCGF) {
+ // Use LexicalBlock if present, otherwise use parent function
+ if (!LexicalBlockStack.empty())
+ FDContext = cast<llvm::DIScope>(LexicalBlockStack.back());
+ else if (ParentCGF && ParentCGF->CurFn)
+ FDContext = ParentCGF->CurFn->getSubprogram();
+ } else
+ FDContext = Unit;
+
if (!HasDecl) {
// Use llvm function name.
LinkageName = Fn->getName();
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
