jyu2 created this revision.
jyu2 added reviewers: ABataev, mikerice, jdoerfert.
jyu2 requested review of this revision.
Herald added a subscriber: sstefan1.
Herald added a project: clang.
A new rule is added in 5.0:
If a list item appears in a reduction, lastprivate or linear clause
on a combined target construct then it is treated as if it also appears
in a map clause with a map-type of tofrom.
Currently map clauses for all capture variables are added implicitly.
But missing for list item of expression for array elements or array
sections.
The change is to add implicit map clause for array of elements used in
reduction clause. Skip adding map clause if the expression is not
mappable.
Noted: For linear and lastprivate, since only variable name is
accepted, the map has been added though capture variables.
To do so:
During the mappable checking, if error, ignore diagnose and skip
adding implicit map clause.
The changes:
1> Add code to generate implicit map in ActOnOpenMPExecutableDirective,
for omp 5.0 and up.
2> Add extra default parameter NoDiagnose in ActOnOpenMPMapClause:
Use that to skip error as well as skip adding implicit map during the
mappable checking.
Note: there are only three places need to be check for NoDiagnose. Rest
of them either the check is for < omp 5.0 or the error already generated for
reduction clause.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D108132
Files:
clang/include/clang/Sema/Sema.h
clang/lib/Sema/SemaOpenMP.cpp
clang/lib/Sema/TreeTransform.h
clang/test/OpenMP/reduction_implicit_map.cpp
Index: clang/test/OpenMP/reduction_implicit_map.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/reduction_implicit_map.cpp
@@ -0,0 +1,115 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ \
+// RUN: -triple powerpc64le-unknown-unknown -DCUDA \
+// RUN: -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o \
+// RUN: %t-ppc-host.bc
+
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ \
+// RUN: -triple nvptx64-unknown-unknown -DCUA \
+// RUN: -fopenmp-targets=nvptx64-nvidia-cuda -DCUDA -emit-llvm %s \
+// RUN: -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc \
+// RUN: -o - | FileCheck %s --check-prefix CHECK
+
+// RUN: %clang_cc1 -verify -fopenmp -x c++ \
+// RUN: -triple powerpc64le-unknown-unknown -DDIAG\
+// RUN: -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm \
+// RUN: %s -o - | FileCheck %s \
+// RUN: --check-prefix=CHECK1
+
+// RUN: %clang_cc1 -verify -fopenmp -x c++ \
+// RUN: -triple i386-unknown-unknown \
+// RUN: -fopenmp-targets=i386-pc-linux-gnu -emit-llvm \
+// RUN: %s -o - | FileCheck %s \
+// RUN: --check-prefix=CHECK2
+
+// expected-no-diagnostics
+
+#if defined(CUDA)
+int foo(int n) {
+ double *e;
+ //no error and no implicit map generated for e[:1]
+ #pragma omp target parallel reduction(+: e[:1])
+ *e=10;
+ ;
+ return 0;
+}
+// CHECK-NOT @.offload_maptypes
+// CHECK: call void @__kmpc_nvptx_end_reduce_nowait(
+#elif defined(DIAG)
+class S2 {
+ mutable int a;
+public:
+ S2():a(0) { }
+ S2(S2 &s2):a(s2.a) { }
+ S2 &operator +(S2 &s);
+};
+int bar() {
+ S2 o[5];
+ //no warnig "copyable and not guaranteed to be mapped correctly" and
+ //implicit map generated.
+#pragma omp target reduction(+:o[0])
+ for (int i = 0; i < 10; i++);
+ double b[10][10][10];
+ //no error no implicit map generated, the map for b is generated but not
+ //for b[0:2][2:4][1].
+#pragma omp target parallel for reduction(task, +: b[0:2][2:4][1])
+ for (long long i = 0; i < 10; ++i);
+ return 0;
+}
+// map for variable o
+// CHECK1: offload_sizes = private unnamed_addr constant [1 x i64] [i64 20]
+// CHECK1: offload_maptypes = private unnamed_addr constant [1 x i64] [i64 547]
+// map for b:
+// CHECK1: @.offload_sizes.1 = private unnamed_addr constant [1 x i64] [i64 8000]
+// CHECK1: @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 547]
+#else
+// generate implicit map for array elements or array sections in reduction
+// clause. In following case: the implicit map is generate for output[0]
+// with map size 4 and output[:3] with map size 12.
+void sum(int* input, int size, int* output)
+{
+ #pragma omp target teams distribute parallel for reduction(+:output[0]) map(to:input[0:size])
+ for(int i=0; i<size; i++)
+ output[0] += input[i];
+ #pragma omp target teams distribute parallel for reduction(+:output[:3]) map(to:input[0:size])
+ for(int i=0; i<size; i++)
+ output[0] += input[i];
+ int a[10];
+ #pragma omp target parallel reduction(+:a[:2])
+ for(int i=0; i<size; i++);
+ #pragma omp target parallel reduction(+:a[3])
+ for(int i=0; i<size; i++);
+}
+//CHECK2: @.offload_sizes = private unnamed_addr constant [2 x i64] [i64 4, i64 8]
+//CHECK2: @.offload_maptypes.10 = private unnamed_addr constant [2 x i64] [i64 800, i64 547]
+//CHECK2: @.offload_sizes.13 = private unnamed_addr constant [2 x i64] [i64 4, i64 4]
+//CHECK2: @.offload_maptypes.14 = private unnamed_addr constant [2 x i64] [i64 800, i64 547]
+//CHECK2: define dso_local void @_Z3sumPiiS_
+//CHECK2-NEXT: entry
+//CHECK2-NEXT: [[INP:%.*]] = alloca i32*
+//CHECK2-NEXT: [[SIZE:%.*]] = alloca i32
+//CHECK2-NEXT: [[OUTP:%.*]] = alloca i32*
+//CHECK2: [[OFFSIZE:%.*]] = alloca [3 x i64]
+//CHECK2: [[OFFSIZE10:%.*]] = alloca [3 x i64]
+//CHECK2: [[T15:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[OFFSIZE]], i32 0, i32 0
+//CHECK2-NEXT: store i64 4, i64* [[T15]]
+//CHECK2: [[T21:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[OFFSIZE]], i32 0, i32 1
+//CHECK2-NEXT: store i64 4, i64* [[T21]]
+//CHECK2: [[T53:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[OFFSIZE10]], i32 0, i32 0
+//CHECK2-NEXT: store i64 4, i64* [[T53]]
+//CHECK2: [[T59:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[OFFSIZE10]], i32 0, i32 1
+//CHECK2-NEXT: store i64 12, i64* [[T59]]
+#endif
+int main()
+{
+#if defined(CUDA)
+ int a = foo(10);
+#elif defined(DIAG)
+ int a = bar();
+#else
+ const int size = 100;
+ int *array = new int[size];
+ int result = 0;
+ sum(array, size, &result);
+#endif
+ return 0;
+}
Index: clang/lib/Sema/TreeTransform.h
===================================================================
--- clang/lib/Sema/TreeTransform.h
+++ clang/lib/Sema/TreeTransform.h
@@ -1929,10 +1929,10 @@
OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
const OMPVarListLocTy &Locs, ArrayRef<Expr *> UnresolvedMappers) {
- return getSema().ActOnOpenMPMapClause(MapTypeModifiers, MapTypeModifiersLoc,
- MapperIdScopeSpec, MapperId, MapType,
- IsMapTypeImplicit, MapLoc, ColonLoc,
- VarList, Locs, UnresolvedMappers);
+ return getSema().ActOnOpenMPMapClause(
+ MapTypeModifiers, MapTypeModifiersLoc, MapperIdScopeSpec, MapperId,
+ MapType, IsMapTypeImplicit, MapLoc, ColonLoc, VarList, Locs,
+ /*NoDiagnose=*/false, UnresolvedMappers);
}
/// Build a new OpenMP 'allocate' clause.
Index: clang/lib/Sema/SemaOpenMP.cpp
===================================================================
--- clang/lib/Sema/SemaOpenMP.cpp
+++ clang/lib/Sema/SemaOpenMP.cpp
@@ -5812,6 +5812,31 @@
ErrorFound = true;
}
}
+ // OpenMP 5.0 [2.19.7]
+ // If a list item appears in a reduction, lastprivate or linear
+ // clause on a combined target construct then it is treated as
+ // if it also appears in a map clause with a map-type of tofrom
+ if (getLangOpts().OpenMP >= 50 && Kind != OMPD_target &&
+ isOpenMPTargetExecutionDirective(Kind)) {
+ SmallVector<Expr *, 4> ImplicitExprs;
+ for (OMPClause *C : Clauses) {
+ if (auto *RC = dyn_cast<OMPReductionClause>(C))
+ for (Expr *E : RC->varlists())
+ if (!dyn_cast<DeclRefExpr>(E))
+ ImplicitExprs.emplace_back(E);
+ }
+ if (!ImplicitExprs.empty()) {
+ ArrayRef<Expr *> Exprs = ImplicitExprs;
+ CXXScopeSpec MapperIdScopeSpec;
+ DeclarationNameInfo MapperId;
+ if (OMPClause *Implicit = ActOnOpenMPMapClause(
+ OMPC_MAP_MODIFIER_unknown, SourceLocation(), MapperIdScopeSpec,
+ MapperId, OMPC_MAP_tofrom,
+ /*IsMapTypeImplicit=*/true, SourceLocation(), SourceLocation(),
+ Exprs, OMPVarListLocTy(), /*NoDiagnose=*/true))
+ ClausesWithImplicit.emplace_back(Implicit);
+ }
+ }
for (unsigned I = 0, E = DefaultmapKindNum; I < E; ++I) {
int ClauseKindCnt = -1;
for (ArrayRef<Expr *> ImplicitMap : ImplicitMaps[I]) {
@@ -18442,12 +18467,15 @@
static bool checkTypeMappable(SourceLocation SL, SourceRange SR, Sema &SemaRef,
DSAStackTy *Stack, QualType QTy,
- bool FullCheck = true) {
+ bool FullCheck = true, bool NoDiagnose = false) {
if (SemaRef.RequireCompleteType(SL, QTy, diag::err_incomplete_type))
return false;
if (FullCheck && !SemaRef.CurContext->isDependentContext() &&
- !QTy.isTriviallyCopyableType(SemaRef.Context))
+ !QTy.isTriviallyCopyableType(SemaRef.Context)) {
+ if (NoDiagnose)
+ return true;
SemaRef.Diag(SL, diag::warn_omp_non_trivial_type_mapped) << QTy << SR;
+ }
return true;
}
@@ -18732,7 +18760,10 @@
}
bool VisitOMPArraySectionExpr(OMPArraySectionExpr *OASE) {
- assert(!NoDiagnose && "Array sections cannot be implicitly mapped.");
+ // After OMP 5.0 Array section in reduction clause will be implicitly
+ // mapped
+ assert(!(SemaRef.getLangOpts().OpenMP < 50 && NoDiagnose) &&
+ "Array sections cannot be implicitly mapped.");
Expr *E = OASE->getBase()->IgnoreParenImpCasts();
QualType CurType =
OMPArraySectionExpr::getBaseOriginalType(E).getCanonicalType();
@@ -18775,6 +18806,8 @@
} else if (AllowUnitySizeArraySection && NotUnity) {
// A unity or whole array section is not allowed and that is not
// compatible with the properties of the current array section.
+ if (NoDiagnose)
+ return false;
SemaRef.Diag(
ELoc, diag::err_array_section_does_not_specify_contiguous_storage)
<< OASE->getSourceRange();
@@ -19318,7 +19351,7 @@
CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo MapperId,
ArrayRef<Expr *> UnresolvedMappers,
OpenMPMapClauseKind MapType = OMPC_MAP_unknown,
- bool IsMapTypeImplicit = false) {
+ bool IsMapTypeImplicit = false, bool NoDiagnose = false) {
// We only expect mappable expressions in 'to', 'from', and 'map' clauses.
assert((CKind == OMPC_map || CKind == OMPC_to || CKind == OMPC_from) &&
"Unexpected clause kind with mappable expressions!");
@@ -19399,7 +19432,7 @@
// components array with all the components identified in the process.
const Expr *BE = checkMapClauseExpressionBase(
SemaRef, SimpleExpr, CurComponents, CKind, DSAS->getCurrentDirective(),
- /*NoDiagnose=*/false);
+ /*NoDiagnose=*/NoDiagnose);
if (!BE)
continue;
@@ -19445,6 +19478,8 @@
// OpenMP 4.5 [2.10.5, target update Construct]
// threadprivate variables cannot appear in a from clause.
if (VD && DSAS->isThreadPrivate(VD)) {
+ if (NoDiagnose)
+ continue;
DSAStackTy::DSAVarData DVar = DSAS->getTopDSA(VD, /*FromParent=*/false);
SemaRef.Diag(ELoc, diag::err_omp_threadprivate_in_clause)
<< getOpenMPClauseName(CKind);
@@ -19505,7 +19540,7 @@
// OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.9]
// A list item must have a mappable type.
if (!checkTypeMappable(VE->getExprLoc(), VE->getSourceRange(), SemaRef,
- DSAS, Type))
+ DSAS, Type, /*FullCheck=*/true, NoDiagnose))
continue;
if (CKind == OMPC_map) {
@@ -19608,7 +19643,8 @@
CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo &MapperId,
OpenMPMapClauseKind MapType, bool IsMapTypeImplicit, SourceLocation MapLoc,
SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
- const OMPVarListLocTy &Locs, ArrayRef<Expr *> UnresolvedMappers) {
+ const OMPVarListLocTy &Locs, bool NoDiagnose,
+ ArrayRef<Expr *> UnresolvedMappers) {
OpenMPMapModifierKind Modifiers[] = {
OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown,
OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown};
@@ -19632,7 +19668,7 @@
MappableVarListInfo MVLI(VarList);
checkMappableExpressionList(*this, DSAStack, OMPC_map, MVLI, Locs.StartLoc,
MapperIdScopeSpec, MapperId, UnresolvedMappers,
- MapType, IsMapTypeImplicit);
+ MapType, IsMapTypeImplicit, NoDiagnose);
// We need to produce a map clause even if we don't have variables so that
// other diagnostics related with non-existing map clauses are accurate.
Index: clang/include/clang/Sema/Sema.h
===================================================================
--- clang/include/clang/Sema/Sema.h
+++ clang/include/clang/Sema/Sema.h
@@ -11254,6 +11254,7 @@
OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
SourceLocation MapLoc, SourceLocation ColonLoc,
ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs,
+ bool NoDiagnose = false,
ArrayRef<Expr *> UnresolvedMappers = llvm::None);
/// Called on well-formed 'num_teams' clause.
OMPClause *ActOnOpenMPNumTeamsClause(Expr *NumTeams, SourceLocation StartLoc,
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
