[clang][OpenMP][DebugInfo] Debug support for variables in containing scope of OMP constructs

[alok]

Summary

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

[alok]

re-based.

[alok]

Re-based.

[bwyma]

If the parent function is inlined into multiple callers, is the outlined subprogram scope treated like any other nested lexical scope and duplicated?
If the outlined subprogram scope is duplicated when inlined, how will retained nodes on the outlined subprogram be handled in this case?
Does CodeViewDebug handle the nested subprogram scoping?
Can you add a test which validates the nested scoping is correct in the DWARF emission (covering the non-clang part of this patch)?

[krisb]

llvm/lib/CodeGen/AsmPrinter/DwarfCompileUnit.cpp

689

In case of compiling optimized code with inlining enabled, the Scope may be either a scope of an inlined instance (which may appear multiple times), an abstract scope or a concrete out-of-line scope. All of them will match a single ScopeNode, so this will override the map's value multiple times and getDIE() will likely return something that one may not expect.
I've been working on patches that make possible for local types, imports and static variables to be scoped in a lexical block, see D125693 for backend-related changes. May be you could find something useful there (see DwarfCompileUnit::getOrCreateLexicalBlockDIE() and DwarfCompileUnit::getOrCreateContextDIE()) or could help to review the patches.

clang/lib/CodeGen/CGDebugInfo.h

public:
 
   QualType getFunctionType(const FunctionDecl *FD, QualType RetTy,
                            const SmallVectorImpl<const VarDecl *> &Args);
 
   /// Emit a call to llvm.dbg.function.start to indicate
   /// 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);
   /// End an inlined function scope.
   void EmitInlineFunctionEnd(CGBuilderTy &Builder);
 
   /// Emit debug info for a function declaration.
   /// \p Fn is set only when a declaration for a debug call site gets created.

clang/lib/CodeGen/CGDebugInfo.cpp

CGDebugInfo::getFunctionType(const FunctionDecl *FD, QualType RetTy,
   for (const VarDecl *VD : Args)
     ArgTypes.push_back(VD->getType());
   return CGM.getContext().getFunctionType(RetTy, ArgTypes,
                                           FunctionProtoType::ExtProtoInfo(CC));
 }
 
 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;
 
   FnBeginRegionCount.push_back(LexicalBlockStack.size());
 
   const Decl *D = GD.getDecl();
   bool HasDecl = (D != nullptr);
 
   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();
   } else if (const auto *FD = dyn_cast<FunctionDecl>(D)) {
     // If there is a subprogram for this function available then use it.
     auto FI = SPCache.find(FD->getCanonicalDecl());
     if (FI != SPCache.end()) {
       auto *SP = dyn_cast_or_null<llvm::DISubprogram>(FI->second);

clang/lib/CodeGen/CGOpenMPRuntime.h

protected:
 
   /// Helper to emit outlined function for 'target' directive.
   /// \param D Directive to emit.
   /// \param ParentName Name of the function that encloses the target region.
   /// \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 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.
   ///
   llvm::Value *emitUpdateLocation(CodeGenFunction &CGF, SourceLocation Loc,
                                   unsigned Flags = 0);
 
   /// Emit the number of teams for a target directive.  Inspect the num_teams
public:
   /// Emits outlined function for the specified OpenMP parallel directive
   /// \a D. This outlined function has type void(*)(kmp_int32 *ThreadID,
   /// kmp_int32 BoundID, struct context_vars*).
   /// \param D OpenMP directive.
   /// \param ThreadIDVar Variable for thread id in the current OpenMP region.
   /// \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,
   /// kmp_int32 BoundID, struct context_vars*).
   /// \param D OpenMP directive.
   /// \param ThreadIDVar Variable for thread id in the current OpenMP region.
   /// \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*
   /// TaskT).
   /// \param D OpenMP directive.
   /// \param ThreadIDVar Variable for thread id in the current OpenMP region.
   /// \param PartIDVar Variable for partition id in the current OpenMP untied
   /// task region.
   /// \param TaskTVar Variable for task_t argument.
   /// \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 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.
   ///
   virtual void functionFinished(CodeGenFunction &CGF);
 
   /// Emits code for parallel or serial call of the \a OutlinedFn with
   /// variables captured in a record which address is stored in \a
   /// CapturedStruct.
public:
 
   /// Emit outilined function for 'target' directive.
   /// \param D Directive to emit.
   /// \param ParentName Name of the function that encloses the target region.
   /// \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 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,
                                           StringRef ParentName,
                                           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
   /// a failure it executes the host version outlined in \a OutlinedFn.
   /// \param D Directive to emit.
   /// \param OutlinedFn Host version of the code to be offloaded.
   /// \param OutlinedFnID ID of host version of the code to be offloaded.
   /// \param IfCond Expression evaluated in if clause associated with the target
public:
   /// Emits outlined function for the specified OpenMP parallel directive
   /// \a D. This outlined function has type void(*)(kmp_int32 *ThreadID,
   /// kmp_int32 BoundID, struct context_vars*).
   /// \param D OpenMP directive.
   /// \param ThreadIDVar Variable for thread id in the current OpenMP region.
   /// \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,
   /// kmp_int32 BoundID, struct context_vars*).
   /// \param D OpenMP directive.
   /// \param ThreadIDVar Variable for thread id in the current OpenMP region.
   /// \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*
   /// TaskT).
   /// \param D OpenMP directive.
   /// \param ThreadIDVar Variable for thread id in the current OpenMP region.
   /// \param PartIDVar Variable for partition id in the current OpenMP untied
   /// task region.
   /// \param TaskTVar Variable for task_t argument.
   /// \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 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
   /// CapturedStruct.
   /// \param OutlinedFn Outlined function to be run in parallel threads. Type of
   /// this function is void(*)(kmp_int32 *, kmp_int32, struct context_vars*).
   /// \param CapturedVars A pointer to the record with the references to
   /// variables used in \a OutlinedFn function.
public:
 
   /// Emit outilined function for 'target' directive.
   /// \param D Directive to emit.
   /// \param ParentName Name of the function that encloses the target region.
   /// \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 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
   /// a failure it executes the host version outlined in \a OutlinedFn.
   /// \param D Directive to emit.
   /// \param OutlinedFn Host version of the code to be offloaded.
   /// \param OutlinedFnID ID of host version of the code to be offloaded.
   /// \param IfCond Expression evaluated in if clause associated with the target

clang/lib/CodeGen/CGOpenMPRuntime.cpp

struct PushAndPopStackRAII {
   }
   llvm::OpenMPIRBuilder *OMPBuilder;
 };
 } // namespace
 
 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();
   else if (const auto *OPD = dyn_cast<OMPTargetParallelDirective>(&D))
     HasCancel = OPD->hasCancel();
   else if (const auto *OPSD = dyn_cast<OMPParallelSectionsDirective>(&D))
     HasCancel = OPSD->hasCancel();
   else if (const auto *OPFD = dyn_cast<OMPParallelForDirective>(&D))
static llvm::Function *emitParallelOrTeamsOutlinedFunction(
   CGOpenMPOutlinedRegionInfo CGInfo(*CS, ThreadIDVar, CodeGen, InnermostKind,
                                     HasCancel, OutlinedHelperName);
   CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo);
   return CGF.GenerateOpenMPCapturedStmtFunction(*CS, D.getBeginLoc());
 }
 
 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());
     llvm::Value *UpLoc = emitUpdateLocation(CGF, D.getBeginLoc());
     llvm::Value *TaskArgs[] = {
         UpLoc, ThreadID,
         CGF.EmitLoadOfPointerLValue(CGF.GetAddrOfLocalVar(TaskTVar),
                                     TaskTVar->getType()->castAs<PointerType>())
llvm::Function *CGOpenMPRuntime::emitTaskOutlinedFunction(
   else if (const auto *TD = dyn_cast<OMPTaskLoopDirective>(&D))
     HasCancel = TD->hasCancel();
   else if (const auto *TD = dyn_cast<OMPMasterTaskLoopDirective>(&D))
     HasCancel = TD->hasCancel();
   else if (const auto *TD = dyn_cast<OMPParallelMasterTaskLoopDirective>(&D))
     HasCancel = TD->hasCancel();
 
   CodeGenFunction CGF(CGM, true);
+  CGF.ParentCGF = &ParentCGF;
   CGOpenMPTaskOutlinedRegionInfo CGInfo(*CS, ThreadIDVar, CodeGen,
                                         InnermostKind, HasCancel, Action);
   CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo);
   llvm::Function *Res = CGF.GenerateCapturedStmtFunction(*CS);
   if (!Tied)
     NumberOfParts = Action.getNumberOfParts();
   return Res;
 }

 /// kmp_int32 .omp_task_entry.(kmp_int32 gtid, kmp_task_t *tt) {
 ///   TaskFunction(gtid, tt->part_id, &tt->privates, task_privates_map, tt,
 ///   For taskloops:
 ///   tt->task_data.lb, tt->task_data.ub, tt->task_data.st, tt->task_data.liter,
 ///   tt->reductions, tt->shareds);
 ///   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,
+    llvm::Value *TaskPrivatesMap, CodeGenFunction &ParentCGF) {
   ASTContext &C = CGM.getContext();
   FunctionArgList Args;
   ImplicitParamDecl GtidArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr, KmpInt32Ty,
                             ImplicitParamDecl::Other);
   ImplicitParamDecl TaskTypeArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
                                 KmpTaskTWithPrivatesPtrQTy.withRestrict(),
                                 ImplicitParamDecl::Other);
   Args.push_back(&GtidArg);
   Args.push_back(&TaskTypeArg);
   const auto &TaskEntryFnInfo =
       CGM.getTypes().arrangeBuiltinFunctionDeclaration(KmpInt32Ty, Args);
   llvm::FunctionType *TaskEntryTy =
       CGM.getTypes().GetFunctionType(TaskEntryFnInfo);
   std::string Name = CGM.getOpenMPRuntime().getName({"omp_task_entry", ""});
   auto *TaskEntry = llvm::Function::Create(
       TaskEntryTy, llvm::GlobalValue::InternalLinkage, Name, &CGM.getModule());
   CGM.SetInternalFunctionAttributes(GlobalDecl(), TaskEntry, TaskEntryFnInfo);
   TaskEntry->setDoesNotRecurse();
   CodeGenFunction CGF(CGM);
+  CGF.ParentCGF = &ParentCGF;
   CGF.StartFunction(GlobalDecl(), KmpInt32Ty, TaskEntry, TaskEntryFnInfo, Args,
                     Loc, Loc);
 
   // TaskFunction(gtid, tt->task_data.part_id, &tt->privates, task_privates_map,
   // tt,
   // For taskloops:
   // tt->task_data.lb, tt->task_data.ub, tt->task_data.st, tt->task_data.liter,
   // tt->task_data.shareds);

 ///   *priv1 = &.privates.priv1;
 ///   ...;
 ///   *privn = &.privates.privn;
 /// }
 /// \endcode
 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(
       C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
       C.getPointerType(PrivatesQTy).withConst().withRestrict(),
       ImplicitParamDecl::Other);
   Args.push_back(&TaskPrivatesArg);
   llvm::DenseMap<CanonicalDeclPtr<const VarDecl>, unsigned> PrivateVarsPos;
emitTaskPrivateMappingFunction(CodeGenModule &CGM, SourceLocation Loc,
   CGM.SetInternalFunctionAttributes(GlobalDecl(), TaskPrivatesMap,
                                     TaskPrivatesMapFnInfo);
   if (CGM.getLangOpts().Optimize) {
     TaskPrivatesMap->removeFnAttr(llvm::Attribute::NoInline);
     TaskPrivatesMap->removeFnAttr(llvm::Attribute::OptimizeNone);
     TaskPrivatesMap->addFnAttr(llvm::Attribute::AlwaysInline);
   }
   CodeGenFunction CGF(CGM);
+  CGF.ParentCGF = &ParentCGF;
   CGF.StartFunction(GlobalDecl(), C.VoidTy, TaskPrivatesMap,
                     TaskPrivatesMapFnInfo, Args, Loc, Loc);
 
   // *privi = &.privates.privi;
   LValue Base = CGF.EmitLoadOfPointerLValue(
       CGF.GetAddrOfLocalVar(&TaskPrivatesArg),
       TaskPrivatesArg.getType()->castAs<PointerType>());
   const auto *PrivatesQTyRD = cast<RecordDecl>(PrivatesQTy->getAsTagDecl());
CGOpenMPRuntime::emitTaskInit(CodeGenFunction &CGF, SourceLocation Loc,
   QualType SharedsPtrTy = C.getPointerType(SharedsTy);
 
   // Emit initial values for private copies (if any).
   llvm::Value *TaskPrivatesMap = nullptr;
   llvm::Type *TaskPrivatesMapTy =
       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 {
     TaskPrivatesMap = llvm::ConstantPointerNull::get(
         cast<llvm::PointerType>(TaskPrivatesMapTy));
   }
   // Build a proxy function kmp_int32 .omp_task_entry.(kmp_int32 gtid,
   // kmp_task_t *tt);
   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,
   // kmp_routine_entry_t *task_entry);
   // Task flags. Format is taken from
   // https://github.com/llvm/llvm-project/blob/main/openmp/runtime/src/kmp.h,
   // description of kmp_tasking_flags struct.
   enum {
void Exit(CodeGenFunction &CGF) override {
     }
   }
 };
 } // namespace
 
 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;
   for (const auto *C : D.getClausesOfKind<OMPUsesAllocatorsClause>()) {
     for (unsigned I = 0, E = C->getNumberOfAllocators(); I < E; ++I) {
       const OMPUsesAllocatorsClause::Data D = C->getAllocatorData(I);
       if (!D.AllocatorTraits)
         continue;
       Allocators.emplace_back(D.Allocator, D.AllocatorTraits);
     }
   }
   OMPUsesAllocatorsActionTy UsesAllocatorAction(Allocators);
   CodeGen.setAction(UsesAllocatorAction);
   emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
-                                   IsOffloadEntry, CodeGen);
+                                   IsOffloadEntry, CodeGen, ParentCGF);
 }
 
 void CGOpenMPRuntime::emitUsesAllocatorsInit(CodeGenFunction &CGF,
                                              const Expr *Allocator,
                                              const Expr *AllocatorTraits) {
   llvm::Value *ThreadId = getThreadID(CGF, Allocator->getExprLoc());
   ThreadId = CGF.Builder.CreateIntCast(ThreadId, CGF.IntTy, /*isSigned=*/true);
   // Use default memspace handle.
(void)CGF.EmitRuntimeCall(
       OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(),
                                             OMPRTL___kmpc_destroy_allocator),
       {ThreadId, AllocatorVal});
 }
 
 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:
   //
   // __omp_offloading_DD_FFFF_PP_lBB
   //
   // where DD_FFFF is an ID unique to the file (device and file IDs), PP is the
   // mangled name of the function that encloses the target region and BB is the
   // line number of the target region.
SmallString<64> EntryFnName;
     llvm::raw_svector_ostream OS(EntryFnName);
     OS << "__omp_offloading" << llvm::format("_%x", DeviceID)
        << llvm::format("_%x_", FileID) << ParentName << "_l" << Line;
   }
 
   const CapturedStmt &CS = *D.getCapturedStmt(OMPD_target);
 
   CodeGenFunction CGF(CGM, true);
+  CGF.ParentCGF = ParentCGF;
   CGOpenMPTargetRegionInfo CGInfo(CS, CodeGen, EntryFnName);
   CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo);
 
   if (BuildOutlinedFn)
     OutlinedFn = CGF.GenerateOpenMPCapturedStmtFunction(CS, D.getBeginLoc());
 
   // If this target outline function is not an offload entry, we don't need to
   // register it.
LValue LPLVal = CGF.MakeAddrLValue(
       Address(GV, GV->getValueType(), PrivLVal.getAlignment()),
       PrivLVal.getType().getNonReferenceType());
   llvm::Value *Res = CGF.EmitLoadOfScalar(LPLVal, Loc);
   CGF.EmitStoreOfScalar(Res, PrivLVal);
 }
 
 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");
 }
 
 llvm::Function *CGOpenMPSIMDRuntime::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) {
   llvm_unreachable("Not supported in SIMD-only mode");
 }
 
 void CGOpenMPSIMDRuntime::emitParallelCall(CodeGenFunction &CGF,
                                            SourceLocation Loc,
                                            llvm::Function *OutlinedFn,
                                            ArrayRef<llvm::Value *> CapturedVars,
                                            const Expr *IfCond,
void CGOpenMPSIMDRuntime::emitCancelCall(CodeGenFunction &CGF,
                                          SourceLocation Loc, const Expr *IfCond,
                                          OpenMPDirectiveKind CancelRegion) {
   llvm_unreachable("Not supported in SIMD-only mode");
 }
 
 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");
 }
 
 void CGOpenMPSIMDRuntime::emitTargetCall(
     CodeGenFunction &CGF, const OMPExecutableDirective &D,
     llvm::Function *OutlinedFn, llvm::Value *OutlinedFnID, const Expr *IfCond,
     llvm::PointerIntPair<const Expr *, 2, OpenMPDeviceClauseModifier> Device,
     llvm::function_ref<llvm::Value *(CodeGenFunction &CGF,

clang/lib/CodeGen/CGOpenMPRuntimeGPU.h

private:
 
   /// Emit outlined function specialized for the Fork-Join
   /// programming model for applicable target directives on the NVPTX device.
   /// \param D Directive to emit.
   /// \param ParentName Name of the function that encloses the target region.
   /// \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
   /// NVPTX device.
   /// \param D Directive to emit.
   /// \param ParentName Name of the function that encloses the target region.
   /// \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 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.
   /// \param D Directive to emit.
   /// \param ParentName Name of the function that encloses the target region.
   /// \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
   /// CapturedStruct.
   /// This call is for the Non-SPMD Execution Mode.
   /// \param OutlinedFn Outlined function to be run in parallel threads. Type of
   /// this function is void(*)(kmp_int32 *, kmp_int32, struct context_vars*).
   /// \param CapturedVars A pointer to the record with the references to
public:
   //  directive.
   /// \a D. This outlined function has type void(*)(kmp_int32 *ThreadID,
   /// kmp_int32 BoundID, struct context_vars*).
   /// \param D OpenMP directive.
   /// \param ThreadIDVar Variable for thread id in the current OpenMP region.
   /// \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.
   /// \a D. This outlined function has type void(*)(kmp_int32 *ThreadID,
   /// kmp_int32 BoundID, struct context_vars*).
   /// \param D OpenMP directive.
   /// \param ThreadIDVar Variable for thread id in the current OpenMP region.
   /// \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
   /// CapturedStruct.
   /// \param OutlinedFn Outlined function to be run by team masters. Type of
   /// this function is void(*)(kmp_int32 *, kmp_int32, struct context_vars*).
   /// \param CapturedVars A pointer to the record with the references to
   /// variables used in \a OutlinedFn function.

clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

static bool supportsLightweightRuntime(ASTContext &Ctx,
   default:
     break;
   }
   llvm_unreachable(
       "Unknown programming model for OpenMP directive on NVPTX target.");
 }
 
 void CGOpenMPRuntimeGPU::emitNonSPMDKernel(const OMPExecutableDirective &D,
-                                             StringRef ParentName,
-                                             llvm::Function *&OutlinedFn,
-                                             llvm::Constant *&OutlinedFnID,
-                                             bool IsOffloadEntry,
-                                             const RegionCodeGenTy &CodeGen) {
+                                           const RegionCodeGenTy &CodeGen,
+                                           CodeGenFunction *ParentCGF) {
   ExecutionRuntimeModesRAII ModeRAII(CurrentExecutionMode);
   EntryFunctionState EST;
   WrapperFunctionsMap.clear();
 
   // Emit target region as a standalone region.
   class NVPTXPrePostActionTy : public PrePostActionTy {
     CGOpenMPRuntimeGPU::EntryFunctionState &EST;
 
void Exit(CodeGenFunction &CGF) override {
           static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
       RT.clearLocThreadIdInsertPt(CGF);
       RT.emitKernelDeinit(CGF, EST, /* IsSPMD */ false);
     }
   } Action(EST);
   CodeGen.setAction(Action);
   IsInTTDRegion = true;
   emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
-                                   IsOffloadEntry, CodeGen);
+                                   IsOffloadEntry, CodeGen, ParentCGF);
   IsInTTDRegion = false;
 }
 
 void CGOpenMPRuntimeGPU::emitKernelInit(CodeGenFunction &CGF,
                                         EntryFunctionState &EST, bool IsSPMD) {
   CGBuilderTy &Bld = CGF.Builder;
   Bld.restoreIP(OMPBuilder.createTargetInit(Bld, IsSPMD, requiresFullRuntime()));
   IsInTargetMasterThreadRegion = IsSPMD;
   if (!IsSPMD)
     emitGenericVarsProlog(CGF, EST.Loc);
 }
 
 void CGOpenMPRuntimeGPU::emitKernelDeinit(CodeGenFunction &CGF,
                                           EntryFunctionState &EST,
                                           bool IsSPMD) {
   if (!IsSPMD)
     emitGenericVarsEpilog(CGF);
 
   CGBuilderTy &Bld = CGF.Builder;
   OMPBuilder.createTargetDeinit(Bld, IsSPMD, requiresFullRuntime());
 }
 
 void CGOpenMPRuntimeGPU::emitSPMDKernel(const OMPExecutableDirective &D,
-                                          StringRef ParentName,
-                                          llvm::Function *&OutlinedFn,
-                                          llvm::Constant *&OutlinedFnID,
-                                          bool IsOffloadEntry,
-                                          const RegionCodeGenTy &CodeGen) {
+                                        const RegionCodeGenTy &CodeGen,
+                                        CodeGenFunction *ParentCGF) {
   ExecutionRuntimeModesRAII ModeRAII(
       CurrentExecutionMode, RequiresFullRuntime,
       CGM.getLangOpts().OpenMPCUDAForceFullRuntime ||
           !supportsLightweightRuntime(CGM.getContext(), D));
   EntryFunctionState EST;
 
   // Emit target region as a standalone region.
   class NVPTXPrePostActionTy : public PrePostActionTy {
public:
     void Exit(CodeGenFunction &CGF) override {
       RT.clearLocThreadIdInsertPt(CGF);
       RT.emitKernelDeinit(CGF, EST, /* IsSPMD */ true);
     }
   } Action(*this, EST);
   CodeGen.setAction(Action);
   IsInTTDRegion = true;
   emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
-                                   IsOffloadEntry, CodeGen);
+                                   IsOffloadEntry, CodeGen, ParentCGF);
   IsInTTDRegion = false;
 }
 
 // Create a unique global variable to indicate the execution mode of this target
 // region. The execution mode is either 'generic', or 'spmd' depending on the
 // target directive. This variable is picked up by the offload library to setup
 // the device appropriately before kernel launch. If the execution mode is
 // 'generic', the runtime reserves one warp for the master, otherwise, all
void CGOpenMPRuntimeGPU::createOffloadEntry(llvm::Constant *ID,
 
   // Add a function attribute for the kernel.
   Fn->addFnAttr(llvm::Attribute::get(Ctx, "kernel"));
 }
 
 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;
 
   assert(!ParentName.empty() && "Invalid target region parent name!");
 
   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);
 }
 
 namespace {
 LLVM_ENABLE_BITMASK_ENUMS_IN_NAMESPACE();
 /// Enum for accesseing the reserved_2 field of the ident_t struct.
 enum ModeFlagsTy : unsigned {

 
 void CGOpenMPRuntimeGPU::emitNumTeamsClause(CodeGenFunction &CGF,
                                               const Expr *NumTeams,
                                               const Expr *ThreadLimit,
                                               SourceLocation Loc) {}
 
 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;
     bool PrevIsInParallelRegion;
 
   public:
     NVPTXPrePostActionTy(bool &IsInParallelRegion)
         : IsInParallelRegion(IsInParallelRegion) {}
     void Enter(CodeGenFunction &CGF) override {
       PrevIsInParallelRegion = IsInParallelRegion;
       IsInParallelRegion = true;
     }
     void Exit(CodeGenFunction &CGF) override {
       IsInParallelRegion = PrevIsInParallelRegion;
     }
   } Action(IsInParallelRegion);
   CodeGen.setAction(Action);
   bool PrevIsInTTDRegion = IsInTTDRegion;
   IsInTTDRegion = false;
   bool PrevIsInTargetMasterThreadRegion = IsInTargetMasterThreadRegion;
   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 &&
       !IsInParallelRegion) {
     llvm::Function *WrapperFun =
         createParallelDataSharingWrapper(OutlinedFun, D);
     WrapperFunctionsMap[OutlinedFun] = WrapperFun;
   }
getTeamsReductionVars(ASTContext &Ctx, const OMPExecutableDirective &D,
   for (const auto *C : D.getClausesOfKind<OMPReductionClause>()) {
     for (const Expr *E : C->privates())
       Vars.push_back(getPrivateItem(E));
   }
 }
 
 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;
   llvm::SmallVector<const ValueDecl *, 4> LastPrivatesReductions;
   llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> MappedDeclsFields;
   unsigned WarpSize = CGM.getTarget().getGridValue().GV_Warp_Size;
   // Globalize team reductions variable unconditionally in all modes.
   if (getExecutionMode() != CGOpenMPRuntimeGPU::EM_SPMD)
public:
     }
     void Exit(CodeGenFunction &CGF) override {
       static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime())
           .emitGenericVarsEpilog(CGF);
     }
   } Action(Loc, GlobalizedRD, MappedDeclsFields);
   CodeGen.setAction(Action);
   llvm::Function *OutlinedFun = CGOpenMPRuntime::emitTeamsOutlinedFunction(
-      D, ThreadIDVar, InnermostKind, CodeGen);
+      D, ThreadIDVar, InnermostKind, CodeGen, ParentCGF);
 
   return OutlinedFun;
 }
 
 void CGOpenMPRuntimeGPU::emitGenericVarsProlog(CodeGenFunction &CGF,
                                                  SourceLocation Loc,
                                                  bool WithSPMDCheck) {
   if (getDataSharingMode(CGM) != CGOpenMPRuntimeGPU::Generic &&

clang/lib/CodeGen/CGStmtOpenMP.cpp

CodeGenFunction::GenerateOpenMPCapturedStmtFunction(const CapturedStmt &S,
   FinishFunction(CD->getBodyRBrace());
   if (!NeedWrapperFunction)
     return F;
 
   FunctionOptions WrapperFO(&S, /*UIntPtrCastRequired=*/true,
                             /*RegisterCastedArgsOnly=*/true,
                             CapturedStmtInfo->getHelperName(), Loc);
   CodeGenFunction WrapperCGF(CGM, /*suppressNewContext=*/true);
+  WrapperCGF.ParentCGF = ParentCGF;
   WrapperCGF.CapturedStmtInfo = CapturedStmtInfo;
   Args.clear();
   LocalAddrs.clear();
   VLASizes.clear();
   llvm::Function *WrapperF =
       emitOutlinedFunctionPrologue(WrapperCGF, Args, LocalAddrs, VLASizes,
                                    WrapperCGF.CXXThisValue, WrapperFO);
   llvm::SmallVector<llvm::Value *, 4> CallArgs;

 static void emitCommonOMPParallelDirective(
     CodeGenFunction &CGF, const OMPExecutableDirective &S,
     OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen,
     const CodeGenBoundParametersTy &CodeGenBoundParameters) {
   const CapturedStmt *CS = S.getCapturedStmt(OMPD_parallel);
   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(),
                                     /*IgnoreResultAssign=*/true);
     CGF.CGM.getOpenMPRuntime().emitNumThreadsClause(
         CGF, NumThreads, NumThreadsClause->getBeginLoc());
   }
   if (const auto *ProcBindClause = S.getSingleClause<OMPProcBindClause>()) {
auto &&CodeGen = [&Data, &S, CS, &BodyGen, &LastprivateDstsOrigs,
 
     CGOpenMPRuntime::UntiedTaskLocalDeclsRAII LocalVarsScope(CGF,
                                                              UntiedLocalVars);
     Action.Enter(CGF);
     BodyGen(CGF);
   };
   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()));
   TaskGen(*this, OutlinedFn, Data);
 }
 
 static ImplicitParamDecl *
 createImplicitFirstprivateForType(ASTContext &C, OMPTaskDataTy &Data,
auto &&CodeGen = [&Data, &S, CS, &BodyGen, BPVD, PVD, SVD, MVD,
     }
 
     Action.Enter(CGF);
     OMPLexicalScope LexScope(CGF, S, OMPD_task, /*EmitPreInitStmt=*/false);
     BodyGen(CGF);
   };
   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),
                         SourceLocation());
 
   CGM.getOpenMPRuntime().emitTaskCall(*this, S.getBeginLoc(), S, OutlinedFn,
                                       SharedsTy, CapturedStruct, &IfCond, Data);
 }
if (const auto *D = dyn_cast<CXXConstructorDecl>(CGF.CurFuncDecl))
     ParentName = CGM.getMangledName(GlobalDecl(D, Ctor_Complete));
   else if (const auto *D = dyn_cast<CXXDestructorDecl>(CGF.CurFuncDecl))
     ParentName = CGM.getMangledName(GlobalDecl(D, Dtor_Complete));
   else
     ParentName =
         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,
                        const OMPLoopDirective &D) -> llvm::Value * {
     if (IsOffloadEntry) {
       OMPLoopScope(CGF, D);
       // Emit calculation of the iterations count.
       llvm::Value *NumIterations = CGF.EmitScalarExpr(D.getNumIterations());

 
 static void emitCommonOMPTeamsDirective(CodeGenFunction &CGF,
                                         const OMPExecutableDirective &S,
                                         OpenMPDirectiveKind InnermostKind,
                                         const RegionCodeGenTy &CodeGen) {
   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>();
   if (NT || TL) {
     const Expr *NumTeams = NT ? NT->getNumTeams() : nullptr;
     const Expr *ThreadLimit = TL ? TL->getThreadLimit() : nullptr;
 
     CGF.CGM.getOpenMPRuntime().emitNumTeamsClause(CGF, NumTeams, ThreadLimit,

clang/lib/CodeGen/CodeGenFunction.cpp

#undef SANITIZER
 
   // Emit subprogram debug descriptor.
   if (CGDebugInfo *DI = getDebugInfo()) {
     // Reconstruct the type from the argument list so that implicit parameters,
     // such as 'this' and 'vtt', show up in the debug info. Preserve the calling
     // convention.
     DI->emitFunctionStart(GD, Loc, StartLoc,
                           DI->getFunctionType(FD, RetTy, Args), CurFn,
-                          CurFuncIsThunk);
+                          CurFuncIsThunk, ParentCGF);
   }
 
   if (ShouldInstrumentFunction()) {
     if (CGM.getCodeGenOpts().InstrumentFunctions)
       CurFn->addFnAttr("instrument-function-entry", "__cyg_profile_func_enter");
     if (CGM.getCodeGenOpts().InstrumentFunctionsAfterInlining)
       CurFn->addFnAttr("instrument-function-entry-inlined",
                        "__cyg_profile_func_enter");

clang/test/OpenMP/debug_containing_scope.c

+// 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;
+}

llvm/lib/CodeGen/AsmPrinter/DwarfCompileUnit.cpp

 
 // Construct new DW_TAG_lexical_block for this scope and attach
 // DW_AT_low_pc/DW_AT_high_pc labels.
 DIE *DwarfCompileUnit::constructLexicalScopeDIE(LexicalScope *Scope) {
   if (DD->isLexicalScopeDIENull(Scope))
     return nullptr;
 
   auto ScopeDIE = DIE::get(DIEValueAllocator, dwarf::DW_TAG_lexical_block);
+  insertDIE(Scope->getScopeNode(), ScopeDIE);
   if (Scope->isAbstractScope())
     return ScopeDIE;
 
   attachRangesOrLowHighPC(*ScopeDIE, Scope->getRanges());
 
   return ScopeDIE;
 }
 

llvm/lib/CodeGen/AsmPrinter/DwarfUnit.h

public:
   /// This takes and attribute parameter because DW_AT_friend attributes are
   /// also type references.
   void addType(DIE &Entity, const DIType *Ty,
                dwarf::Attribute Attribute = dwarf::DW_AT_type);
 
   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);
 
   /// Creates type DIE with specific context.
   DIE *createTypeDIE(const DIScope *Context, DIE &ContextDIE, const DIType *Ty);
 
   /// Find existing DIE or create new DIE for the given type.

llvm/lib/CodeGen/AsmPrinter/DwarfUnit.cpp

DIE *DwarfUnit::getOrCreateContextDIE(const DIScope *Context) {
   if (auto *T = dyn_cast<DIType>(Context))
     return getOrCreateTypeDIE(T);
   if (auto *NS = dyn_cast<DINamespace>(Context))
     return getOrCreateNameSpace(NS);
   if (auto *SP = dyn_cast<DISubprogram>(Context))
     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);
 }
 
 DIE *DwarfUnit::createTypeDIE(const DICompositeType *Ty) {
   auto *Context = Ty->getScope();
   DIE *ContextDIE = getOrCreateContextDIE(Context);
 
   if (DIE *TyDIE = getDIE(Ty))
DIE *DwarfUnit::getOrCreateSubprogramDIE(const DISubprogram *SP, bool Minimal) {
   if (SP->isDefinition())
     return &SPDie;
 
   static_cast<DwarfUnit *>(SPDie.getUnit())
       ->applySubprogramAttributes(SP, SPDie);
   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;
   StringRef DeclLinkageName;
   if (auto *SPDecl = SP->getDeclaration()) {
     if (!Minimal) {
       DITypeRefArray DeclArgs, DefinitionArgs;
       DeclArgs = SPDecl->getType()->getTypeArray();