alok created this revision.
alok added reviewers: aprantl, djtodoro, jmorse, jini.susan.
alok added a project: debug-info.
Herald added subscribers: guansong, hiraditya, yaxunl.
Herald added a project: All.
alok requested review of this revision.
Herald added a reviewer: jdoerfert.
Herald added subscribers: llvm-commits, cfe-commits, sstefan1.
Herald added projects: clang, LLVM.
In case of OpenMP, compilers generate encapsulates code present in
OpenMP construct to artificial functions. This is done to apply parallelism
to block of code. In context of these blocks, currently containing scope
variables are not accessible. This is due to new artificial function DIE being
in global scope.
As from user point of view, containing scope is same lexical scope, there
must be correct DIE hierarchy for artificial functions, which should be child
of containing scope.
Please consider below testcase.
1 #include <stdio.h>
2 #include <stdlib.h>
3
4 int global_var1;
5 int global_var2 = 99;
6 int foo(int n) {
7 int same_var = 5;
8 int other_var = 21;
9 int share = 9, priv, i;
10 global_var1 = 99;
11
12 if (n < 2)
13 return n;
14 else {
15 int same_var = rand() % 5;
16 int local_var = 31;
17 #pragma omp parallel for
18 for (i = 0; i < n; i++) {
19 share += i; // <-------------- (A)
20 }
21 return share;
22 }
23 }
24
25 int main() {
26 int n = 10;
27 printf("foo(%d) = %d\n", n, foo(n));
28 return 0;
29 }
Please consider the line# 19, user expects variables "same_var", "local_var",
"other_var" to be
accessible inside debugger but which is not possible.
(gdb) p same_var
No symbol "same_var" in current context.
(gdb) p other_var
No symbol "other_var" in current context.
(gdb) p local_var
No symbol "local_var" in current context.
After current patch.
(gdb) thr 1
[Switching to thread 1 (Thread 0x7ffff7c1c400 (LWP 17992))]
#0 .omp_outlined._debug__ (.global_tid.=0x7fffffffdad0,
.bound_tid.=0x7fffffffdac8, n=@0x7fffffffdf18: 10, share=@0x7fffffffdf0c: 9) at
1.c:19
19 share += i;
(gdb) p same_var
$1 = 3
(gdb) p local_var
$2 = 31
(gdb) p other_var
$3 = 21
Repository:
rG LLVM Github Monorepo
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,67 @@
+// 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-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: ".omp_outlined..1", 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
