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(LexicalBlockStack.back()); + else if (ParentCGF && ParentCGF->CurFn) + FDContext = ParentCGF->CurFn->getSubprogram(); + } else + FDContext = Unit; + if (!HasDecl) { // Use llvm function name. LinkageName = Fn->getName(); 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(&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 Privates) { + ArrayRef 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, 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/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(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/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()) { 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() ? 1 : 0); IntegerLiteral IfCond(getContext(), TrueOrFalse, getContext().getIntTypeForBitwidth(32, /*Signed=*/0), @@ -6430,8 +6432,8 @@ CGM.getMangledName(GlobalDecl(cast(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(); const auto *TL = S.getSingleClause(); 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/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: 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: 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(Context)) return getOrCreateModule(M); + if (auto *LS = dyn_cast(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;