Author: Alexey Bataev Date: 2020-08-27T17:07:53-04:00 New Revision: ba1de5f2f7b078f69d5f6b0fe3af4911f76bb8fd
URL: https://github.com/llvm/llvm-project/commit/ba1de5f2f7b078f69d5f6b0fe3af4911f76bb8fd DIFF: https://github.com/llvm/llvm-project/commit/ba1de5f2f7b078f69d5f6b0fe3af4911f76bb8fd.diff LOG: [OPENMP]Do not crash for globals in inner regions with outer target region. If the global variable is used in the target region,it is always captured, if not marked as declare target. Added: Modified: clang/lib/Sema/SemaOpenMP.cpp clang/test/OpenMP/target_codegen.cpp Removed: ################################################################################ diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index cd20b6bfdb98..7b62c841b48a 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -2417,17 +2417,20 @@ bool Sema::isOpenMPGlobalCapturedDecl(ValueDecl *D, unsigned Level, if (const auto *VD = dyn_cast<VarDecl>(D)) { if (!VD->hasLocalStorage()) { + if (isInOpenMPTargetExecutionDirective()) + return true; DSAStackTy::DSAVarData TopDVar = DSAStack->getTopDSA(D, /*FromParent=*/false); unsigned NumLevels = getOpenMPCaptureLevels(DSAStack->getDirective(Level)); if (Level == 0) return (NumLevels == CaptureLevel + 1) && TopDVar.CKind != OMPC_shared; - DSAStackTy::DSAVarData DVar = DSAStack->getImplicitDSA(D, Level - 1); - return DVar.CKind != OMPC_shared || - isOpenMPGlobalCapturedDecl( - D, Level - 1, - getOpenMPCaptureLevels(DSAStack->getDirective(Level - 1)) - 1); + do { + --Level; + DSAStackTy::DSAVarData DVar = DSAStack->getImplicitDSA(D, Level); + if (DVar.CKind != OMPC_shared) + return true; + } while (Level >= 0); } } return true; diff --git a/clang/test/OpenMP/target_codegen.cpp b/clang/test/OpenMP/target_codegen.cpp index 55b03aae00a5..c8570bdf6b65 100644 --- a/clang/test/OpenMP/target_codegen.cpp +++ b/clang/test/OpenMP/target_codegen.cpp @@ -706,6 +706,8 @@ int bar(int n){ // CHECK: [[IFEND]] +// OMP45: define internal void @__omp_offloading_{{.+}}_{{.+}}bar{{.+}}_l838(i[[SZ]] %{{.+}}) + // OMP45: define {{.*}}@{{.*}}zee{{.*}} // OMP45: [[LOCAL_THIS:%.+]] = alloca [[S2]]* @@ -803,6 +805,7 @@ int bar(int n){ // CHECK-DAG: load i16, i16* [[REF_AA]] // CHECK-DAG: getelementptr inbounds [10 x i32], [10 x i32]* [[REF_B]], i[[SZ]] 0, i[[SZ]] 2 +// OMP50: define internal void @__omp_offloading_{{.+}}_{{.+}}bar{{.+}}_l838(i[[SZ]] %{{.+}}) // OMP50: define {{.*}}@{{.*}}zee{{.*}} @@ -833,7 +836,11 @@ int bar(int n){ void bar () { #define pragma_target _Pragma("omp target") pragma_target -{} +{ + global = 0; +#pragma omp parallel shared(global) + global = 1; +} } class S2 { _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
