Author: abataev Date: Mon Oct 14 09:44:01 2019 New Revision: 374787 URL: http://llvm.org/viewvc/llvm-project?rev=374787&view=rev Log: [OPENMP]Fix codegen for private variably length vars in combined constructs.
If OpenMP construct includes several capturing regions and the variable is declared as private, the length of the inner variable length array is not captured in outer captured regions, only in the innermost region. Patch fixes this bug. Modified: cfe/trunk/include/clang/Sema/Sema.h cfe/trunk/lib/Sema/SemaExpr.cpp cfe/trunk/lib/Sema/SemaOpenMP.cpp cfe/trunk/test/OpenMP/target_teams_codegen.cpp Modified: cfe/trunk/include/clang/Sema/Sema.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=374787&r1=374786&r2=374787&view=diff ============================================================================== --- cfe/trunk/include/clang/Sema/Sema.h (original) +++ cfe/trunk/include/clang/Sema/Sema.h Mon Oct 14 09:44:01 2019 @@ -9067,6 +9067,10 @@ private: void adjustOpenMPTargetScopeIndex(unsigned &FunctionScopesIndex, unsigned Level) const; + /// Returns the number of scopes associated with the construct on the given + /// OpenMP level. + int getNumberOfConstructScopes(unsigned Level) const; + /// Push new OpenMP function region for non-capturing function. void pushOpenMPFunctionRegion(); Modified: cfe/trunk/lib/Sema/SemaExpr.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaExpr.cpp?rev=374787&r1=374786&r2=374787&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaExpr.cpp (original) +++ cfe/trunk/lib/Sema/SemaExpr.cpp Mon Oct 14 09:44:01 2019 @@ -16092,7 +16092,25 @@ bool Sema::tryCaptureVariable( // target region should not be captured outside the scope of the region. if (RSI->CapRegionKind == CR_OpenMP) { bool IsOpenMPPrivateDecl = isOpenMPPrivateDecl(Var, RSI->OpenMPLevel); - auto IsTargetCap = !IsOpenMPPrivateDecl && + // If the variable is private (i.e. not captured) and has variably + // modified type, we still need to capture the type for correct + // codegen in all regions, associated with the construct. Currently, + // it is captured in the innermost captured region only. + if (IsOpenMPPrivateDecl && Var->getType()->isVariablyModifiedType()) { + QualType QTy = Var->getType(); + if (ParmVarDecl *PVD = dyn_cast_or_null<ParmVarDecl>(Var)) + QTy = PVD->getOriginalType(); + for (int I = 1, E = getNumberOfConstructScopes(RSI->OpenMPLevel); + I < E; ++I) { + auto *OuterRSI = cast<CapturedRegionScopeInfo>( + FunctionScopes[FunctionScopesIndex - I]); + assert(RSI->OpenMPLevel == OuterRSI->OpenMPLevel && + "Wrong number of captured regions associated with the " + "OpenMP construct."); + captureVariablyModifiedType(Context, QTy, OuterRSI); + } + } + bool IsTargetCap = !IsOpenMPPrivateDecl && isOpenMPTargetCapturedDecl(Var, RSI->OpenMPLevel); // When we detect target captures we are looking from inside the // target region, therefore we need to propagate the capture from the Modified: cfe/trunk/lib/Sema/SemaOpenMP.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=374787&r1=374786&r2=374787&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaOpenMP.cpp (original) +++ cfe/trunk/lib/Sema/SemaOpenMP.cpp Mon Oct 14 09:44:01 2019 @@ -3435,6 +3435,10 @@ void Sema::ActOnOpenMPRegionStart(OpenMP } } +int Sema::getNumberOfConstructScopes(unsigned Level) const { + return getOpenMPCaptureLevels(DSAStack->getDirective(Level)); +} + int Sema::getOpenMPCaptureLevels(OpenMPDirectiveKind DKind) { SmallVector<OpenMPDirectiveKind, 4> CaptureRegions; getOpenMPCaptureRegions(CaptureRegions, DKind); Modified: cfe/trunk/test/OpenMP/target_teams_codegen.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_teams_codegen.cpp?rev=374787&r1=374786&r2=374787&view=diff ============================================================================== --- cfe/trunk/test/OpenMP/target_teams_codegen.cpp (original) +++ cfe/trunk/test/OpenMP/target_teams_codegen.cpp Mon Oct 14 09:44:01 2019 @@ -86,6 +86,7 @@ // TCHECK: @{{.+}} = weak constant [[ENTTY]] // TCHECK: @{{.+}} = weak constant [[ENTTY]] // TCHECK: @{{.+}} = weak constant [[ENTTY]] +// TCHECK: @{{.+}} = weak constant [[ENTTY]] // TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]] // Check if offloading descriptor is created. @@ -341,6 +342,13 @@ int foo(int n) { d.Y += 1; } + const int nn = 0; + #pragma omp target teams shared(nn) + #pragma omp parallel firstprivate(nn) + (void)nn; + #pragma omp target teams firstprivate(nn) + #pragma omp parallel shared(nn) + (void)nn; return a; } @@ -481,6 +489,19 @@ int foo(int n) { // CHECK: define internal {{.*}}void [[OMP_OUTLINED4]](i32* noalias %.global_tid., i32* noalias %.bound_tid., i[[SZ]] %{{.+}}, [10 x float]* {{.+}}, i[[SZ]] %{{.+}}, float* {{.+}}, [5 x [10 x double]]* {{.+}}, i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}}, double* {{.+}}, [[TT]]* {{.+}}) // To reduce complexity, we're only going as far as validating the signature of the outlined parallel function. +// CHECK: define {{.*}}void @__omp_offloading_{{.*}}foo{{.*}}_l346(i[[SZ]] %{{.+}}) +// CHECK: define internal void {{@.+}}(i32* {{.+}}, i32* {{.+}}, i[[SZ]] %{{.+}}) +// CHECK: define {{.*}}void @__omp_offloading_{{.*}}foo{{.*}}_l349(i[[SZ]] %{{.+}}) +// CHECK: define internal void {{@.+}}(i32* {{.+}}, i32* {{.+}}, i32* dereferenceable{{.+}}) + +void bazzzz(int n, int f[n]) { +// CHECK: define internal void @__omp_offloading_{{.+}}bazzzz{{.+}}_l501(i[[SZ]] %{{[^,]+}}) +// CHECK: [[VLA:%.+]] = load i[[SZ]], i[[SZ]]* % +// CHECK: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @{{.+}}, i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i[[SZ]])* @{{.+}} to void (i32*, i32*, ...)*), i[[SZ]] [[VLA]]) +#pragma omp target teams private(f) + ; +} + template<typename tx> tx ftemplate(int n) { tx a = 0; @@ -846,21 +867,4 @@ int bar(int n){ // CHECK: define internal {{.*}}void [[OMP_OUTLINED7]](i32* noalias %.global_tid., i32* noalias %.bound_tid., i[[SZ]] %{{.+}}, i[[SZ]] %{{.+}}, [10 x i32]* {{.+}}) // To reduce complexity, we're only going as far as validating the signature of the outlined parallel function. -void foo1() { - const int n = 0; - #pragma omp target teams shared(n) - #pragma omp parallel firstprivate(n) - (void)n; -} -void foo() { - const int n = 0; - #pragma omp target teams firstprivate(n) - #pragma omp parallel shared(n) - (void)n; -} - -// define {{.*}}void @__omp_offloading_{{.*}}foo1{{.*}}_l841(i[[SZ]] %{{.+}}) -// define internal void {{@.+}}(i32* {{.+}}, i32* {{.+}}, i[[SZ]] %{{.+}}) -// define {{.*}}void @__omp_offloading_{{.*}}foo1{{.*}}_l847(i[[SZ]] %{{.+}}) -// define internal void {{@.+}}(i32* {{.+}}, i32* {{.+}}, i32* dereferenceable{{.+}}) #endif _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits