[clang] c835630 - [OPENMP]Fix PR49098: respect firstprivate of declare target variable.
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Wed Apr 28 05:40:08 PDT 2021
Author: Alexey Bataev
Date: 2021-04-28T05:39:10-07:00
New Revision: c835630c25a4f9925517949579f66a43b113fbc9
URL: https://github.com/llvm/llvm-project/commit/c835630c25a4f9925517949579f66a43b113fbc9
DIFF: https://github.com/llvm/llvm-project/commit/c835630c25a4f9925517949579f66a43b113fbc9.diff
LOG: [OPENMP]Fix PR49098: respect firstprivate of declare target variable.
Need to respect mapping/privatization of declare target variables in the
target regions if explicitly specified by the user.
Differential Revision: https://reviews.llvm.org/D99530
Added:
Modified:
clang/lib/Sema/SemaOpenMP.cpp
clang/test/OpenMP/target_firstprivate_codegen.cpp
Removed:
################################################################################
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 2a45a095eb0fd..25ee46d95aa55 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -2187,15 +2187,11 @@ VarDecl *Sema::isOpenMPCapturedDecl(ValueDecl *D, bool CheckScopeInfo,
//
if (VD && !VD->hasLocalStorage() &&
(getCurCapturedRegion() || getCurBlock() || getCurLambda())) {
- if (isInOpenMPDeclareTargetContext()) {
- // Try to mark variable as declare target if it is used in capturing
- // regions.
- if (LangOpts.OpenMP <= 45 &&
- !OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD))
- checkDeclIsAllowedInOpenMPTarget(nullptr, VD);
- return nullptr;
- }
if (isInOpenMPTargetExecutionDirective()) {
+ DSAStackTy::DSAVarData DVarTop =
+ DSAStack->getTopDSA(D, DSAStack->isClauseParsingMode());
+ if (DVarTop.CKind != OMPC_unknown && DVarTop.RefExpr)
+ return VD;
// If the declaration is enclosed in a 'declare target' directive,
// then it should not be captured.
//
@@ -2220,6 +2216,14 @@ VarDecl *Sema::isOpenMPCapturedDecl(ValueDecl *D, bool CheckScopeInfo,
if (Regions[CSI->OpenMPCaptureLevel] != OMPD_task)
return VD;
}
+ if (isInOpenMPDeclareTargetContext()) {
+ // Try to mark variable as declare target if it is used in capturing
+ // regions.
+ if (LangOpts.OpenMP <= 45 &&
+ !OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD))
+ checkDeclIsAllowedInOpenMPTarget(nullptr, VD);
+ return nullptr;
+ }
}
if (CheckScopeInfo) {
diff --git a/clang/test/OpenMP/target_firstprivate_codegen.cpp b/clang/test/OpenMP/target_firstprivate_codegen.cpp
index 8522cac3afa3c..3535371b1f60b 100644
--- a/clang/test/OpenMP/target_firstprivate_codegen.cpp
+++ b/clang/test/OpenMP/target_firstprivate_codegen.cpp
@@ -43,6 +43,9 @@ struct TT {
tx X;
ty Y;
};
+#pragma omp declare target
+int ga = 5;
+#pragma omp end declare target
// CHECK-DAG: [[TT:%.+]] = type { i64, i8 }
// CHECK-DAG: [[TTII:%.+]] = type { i32, i32 }
@@ -52,9 +55,9 @@ struct TT {
// TCHECK-DAG: [[TTII:%.+]] = type { i32, i32 }
// TCHECK-DAG: [[S1:%.+]] = type { double }
-// CHECK-DAG: [[FP_E:@__omp_offloading_firstprivate_.+_e_l76]] = internal global [[TTII]] zeroinitializer
-// CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [2 x i{{32|64}}] [i[[SZ:32|64]] 4, i{{64|32}} {{8|4}}]
-// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 288, i64 49]
+// CHECK-DAG: [[FP_E:@__omp_offloading_firstprivate_.+_e_l79]] = internal global [[TTII]] zeroinitializer
+// CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [3 x i{{32|64}}] [i[[SZ:32|64]] 4, i{{64|32}} {{8|4}}, i[[SZ:32|64]] 4]
+// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 288, i64 49, i64 288]
// CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [9 x i64] [i64 288, i64 161, i64 800, i64 161, i64 161, i64 800, i64 800, i64 161, i64 161]
// CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [2 x i{{32|64}}] [i{{32|64}} 0, i{{32|64}} 8]
// CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 37]
@@ -76,7 +79,7 @@ int foo(int n, double *ptr) {
const TT<int, int> e = {n, n};
int *p __attribute__ ((aligned (64))) = &a;
-#pragma omp target firstprivate(a, p)
+#pragma omp target firstprivate(a, p, ga)
{
}
@@ -91,8 +94,8 @@ int foo(int n, double *ptr) {
// CHECK: [[D:%.+]] = alloca [[TT]],
// CHECK: [[P:%.+]] = alloca i32*, align 64
// CHECK: [[ACAST:%.+]] = alloca i{{[0-9]+}},
- // CHECK: [[BASE_PTR_ARR:%.+]] = alloca [2 x i8*],
- // CHECK: [[PTR_ARR:%.+]] = alloca [2 x i8*],
+ // CHECK: [[BASE_PTR_ARR:%.+]] = alloca [3 x i8*],
+ // CHECK: [[PTR_ARR:%.+]] = alloca [3 x i8*],
// CHECK: [[A2CAST:%.+]] = alloca i{{[0-9]+}},
// CHECK: [[BASE_PTR_ARR2:%.+]] = alloca [9 x i8*],
// CHECK: [[PTR_ARR2:%.+]] = alloca [9 x i8*],
@@ -116,29 +119,37 @@ int foo(int n, double *ptr) {
// CHECK-32: store i{{[0-9]+}} [[AVAL]], i{{[0-9]+}}* [[ACAST]],
// CHECK: [[ACAST_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[ACAST]],
// CHECK: [[P_PTR:%.+]] = load i32*, i32** [[P]], align 64
- // CHECK: [[BASE_PTR_GEP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BASE_PTR_ARR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: [[BASE_PTR_GEP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BASE_PTR_ARR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
// CHECK: [[ACAST_TOPTR:%.+]] = bitcast i8** [[BASE_PTR_GEP]] to i{{[0-9]+}}*
// CHECK: store i{{[0-9]+}} [[ACAST_VAL]], i{{[0-9]+}}* [[ACAST_TOPTR]],
- // CHECK: [[PTR_GEP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTR_ARR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: [[PTR_GEP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTR_ARR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
// CHECK: [[ACAST_TOPTR2:%.+]] = bitcast i8** [[PTR_GEP]] to i{{[0-9]+}}*
// CHECK: store i{{[0-9]+}} [[ACAST_VAL]], i{{[0-9]+}}* [[ACAST_TOPTR2]],
- // CHECK: [[BASE_PTR_GEP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BASE_PTR_ARR]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // CHECK: [[BASE_PTR_GEP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BASE_PTR_ARR]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
// CHECK: [[PCAST_TOPTR:%.+]] = bitcast i8** [[BASE_PTR_GEP]] to i32***
// CHECK: store i32** [[P]], i32*** [[PCAST_TOPTR]],
- // CHECK: [[PTR_GEP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTR_ARR]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // CHECK: [[PTR_GEP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTR_ARR]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
// CHECK: [[PCAST_TOPTR2:%.+]] = bitcast i8** [[PTR_GEP]] to i32**
// CHECK: store i32* [[P_PTR]], i32** [[PCAST_TOPTR2]],
- // CHECK: [[BASE_PTR_GEP_ARG:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BASE_PTR_ARR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
- // CHECK: [[PTR_GEP_ARG:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTR_ARR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
- // CHECK: {{.+}} = call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 -1, {{.+}}, i32 2, i8** [[BASE_PTR_GEP_ARG]], i8** [[PTR_GEP_ARG]], i[[SZ]]* getelementptr inbounds ([2 x i[[SZ]]], [2 x i[[SZ]]]* [[SIZET]], i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAPT]], i32 0, i32 0), i8** null, i8** null)
-
- // TCHECK: define weak void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], i32** nonnull align {{[0-9]+}} dereferenceable({{[0-9]+}}) [[P_IN:%.+]])
+ // CHECK: [[BASE_PTR_GEP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BASE_PTR_ARR]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+ // CHECK: [[PCAST_TOPTR:%.+]] = bitcast i8** [[BASE_PTR_GEP]] to i{{64|32}}*
+ // CHECK: store i{{64|32}} [[GA_VAL:%.*]], i{{64|32}}* [[PCAST_TOPTR]],
+ // CHECK: [[PTR_GEP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTR_ARR]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+ // CHECK: [[PCAST_TOPTR2:%.+]] = bitcast i8** [[PTR_GEP]] to i{{64|32}}*
+ // CHECK: store i{{64|32}} [[GA_VAL]], i{{64|32}}* [[PCAST_TOPTR2]],
+ // CHECK: [[BASE_PTR_GEP_ARG:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BASE_PTR_ARR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: [[PTR_GEP_ARG:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTR_ARR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: {{.+}} = call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 -1, {{.+}}, i32 3, i8** [[BASE_PTR_GEP_ARG]], i8** [[PTR_GEP_ARG]], i[[SZ]]* getelementptr inbounds ([3 x i[[SZ]]], [3 x i[[SZ]]]* [[SIZET]], i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPT]], i32 0, i32 0), i8** null, i8** null)
+
+ // TCHECK: define weak void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], i32** nonnull align {{[0-9]+}} dereferenceable({{[0-9]+}}) [[P_IN:%.+]], i{{[0-9]+}} [[GA_IN:%.+]])
// TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
// TCHECK: [[P_ADDR:%.+]] = alloca i32**,
+ // TCHECK: [[GA_ADDR:%.+]] = alloca i{{64|32}},
// TCHECK: [[P_PRIV:%.+]] = alloca i32*,
// TCHECK-NOT: alloca i{{[0-9]+}}
// TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
// TCHECK: store i32** [[P_IN]], i32*** [[P_ADDR]],
+ // TCHECK: store i{{[0-9]+}} [[GA_IN]], i{{[0-9]+}}* [[GA_ADDR]],
// TCHECK-NOT: store i{{[0-9]+}} %
// TCHECK: ret void
More information about the cfe-commits
mailing list