Skip to content

Commit d8b61dd

Browse files
jyu2-gitjyu2-git
authored
[OpenMP] Generate implicit default mapper for mapping array section. (#101101)
This is only for struct containing nested structs with user defined mappers. Add four functions: 1>buildImplicitMap: build map for default mapper 2>buildImplicitMapper: build default mapper. 3>hasUserDefinedMapper for given mapper name and mapper type, lookup user defined map, if found one return true. 4>isImplicitMapperNeeded check if Mapper is needed During create map, in checkMappableExpressionList, call isImplicitMapperNeeded when it return true, call buildImplicitMapper to generate implicit mapper and added to map clause. #101101
1 parent 1c7540c commit d8b61dd

File tree

6 files changed

+645
-0
lines changed

6 files changed

+645
-0
lines changed

‎clang/docs/OpenMPSupport.rst

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -181,6 +181,8 @@ implementation.
181181
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
182182
| device | user-defined mappers | :good:`done` | D56326,D58638,D58523,D58074,D60972,D59474 |
183183
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
184+
| device | map array-section with implicit mapper | :good:`done` | https://github.com/llvm/llvm-project/pull/101101 |
185+
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
184186
| device | mapping lambda expression | :good:`done` | D51107 |
185187
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
186188
| device | clause: use_device_addr for target data | :good:`done` | |

‎clang/docs/ReleaseNotes.rst

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -307,6 +307,10 @@ Python Binding Changes
307307
OpenMP Support
308308
--------------
309309

310+
Improvements
311+
^^^^^^^^^^^^
312+
- Improve the handling of mapping array-section for struct containing nested structs with user defined mappers
313+
310314
Additional Information
311315
======================
312316

‎clang/lib/Sema/SemaOpenMP.cpp

Lines changed: 215 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -20796,6 +20796,203 @@ struct MappableVarListInfo {
2079620796
};
2079720797
} // namespace
2079820798

20799+
static DeclRefExpr *buildImplicitMap(Sema &S, QualType BaseType,
20800+
DSAStackTy *Stack,
20801+
SmallVectorImpl<OMPClause *> &Maps) {
20802+
20803+
const RecordDecl *RD = BaseType->getAsRecordDecl();
20804+
SourceRange Range = RD->getSourceRange();
20805+
DeclarationNameInfo ImplicitName;
20806+
// Dummy variable _s for Mapper.
20807+
VarDecl *VD = buildVarDecl(S, Range.getEnd(), BaseType, "_s");
20808+
DeclRefExpr *MapperVarRef =
20809+
buildDeclRefExpr(S, VD, BaseType, SourceLocation());
20810+
20811+
// Create implicit map clause for mapper.
20812+
SmallVector<Expr *, 4> SExprs;
20813+
for (auto *FD : RD->fields()) {
20814+
Expr *BE = S.BuildMemberExpr(
20815+
MapperVarRef, /*IsArrow=*/false, Range.getBegin(),
20816+
NestedNameSpecifierLoc(), Range.getBegin(), FD,
20817+
DeclAccessPair::make(FD, FD->getAccess()),
20818+
/*HadMultipleCandidates=*/false,
20819+
DeclarationNameInfo(FD->getDeclName(), FD->getSourceRange().getBegin()),
20820+
FD->getType(), VK_LValue, OK_Ordinary);
20821+
SExprs.push_back(BE);
20822+
}
20823+
CXXScopeSpec MapperIdScopeSpec;
20824+
DeclarationNameInfo MapperId;
20825+
OpenMPDirectiveKind DKind = Stack->getCurrentDirective();
20826+
20827+
OMPClause *MapClause = S.OpenMP().ActOnOpenMPMapClause(
20828+
nullptr, OMPC_MAP_MODIFIER_unknown, SourceLocation(), MapperIdScopeSpec,
20829+
MapperId, DKind == OMPD_target_enter_data ? OMPC_MAP_to : OMPC_MAP_tofrom,
20830+
/*IsMapTypeImplicit=*/true, SourceLocation(), SourceLocation(), SExprs,
20831+
OMPVarListLocTy());
20832+
Maps.push_back(MapClause);
20833+
return MapperVarRef;
20834+
}
20835+
20836+
static ExprResult buildImplicitMapper(Sema &S, QualType BaseType,
20837+
DSAStackTy *Stack) {
20838+
20839+
// Build impilicit map for mapper
20840+
SmallVector<OMPClause *, 4> Maps;
20841+
DeclRefExpr *MapperVarRef = buildImplicitMap(S, BaseType, Stack, Maps);
20842+
20843+
const RecordDecl *RD = BaseType->getAsRecordDecl();
20844+
// AST context is RD's ParentASTContext().
20845+
ASTContext &Ctx = RD->getParentASTContext();
20846+
// DeclContext is RD's DeclContext.
20847+
DeclContext *DCT = const_cast<DeclContext *>(RD->getDeclContext());
20848+
20849+
// Create implicit default mapper for "RD".
20850+
DeclarationName MapperId;
20851+
auto &DeclNames = Ctx.DeclarationNames;
20852+
MapperId = DeclNames.getIdentifier(&Ctx.Idents.get("default"));
20853+
auto *DMD = OMPDeclareMapperDecl::Create(Ctx, DCT, SourceLocation(), MapperId,
20854+
BaseType, MapperId, Maps, nullptr);
20855+
Scope *Scope = S.getScopeForContext(DCT);
20856+
if (Scope)
20857+
S.PushOnScopeChains(DMD, Scope, /*AddToContext*/ false);
20858+
DCT->addDecl(DMD);
20859+
DMD->setAccess(clang::AS_none);
20860+
auto *VD = cast<DeclRefExpr>(MapperVarRef)->getDecl();
20861+
VD->setDeclContext(DMD);
20862+
VD->setLexicalDeclContext(DMD);
20863+
DMD->addDecl(VD);
20864+
DMD->setMapperVarRef(MapperVarRef);
20865+
FieldDecl *FD = *RD->field_begin();
20866+
// create mapper refence.
20867+
return DeclRefExpr::Create(Ctx, NestedNameSpecifierLoc{}, FD->getLocation(),
20868+
DMD, false, SourceLocation(), BaseType, VK_LValue);
20869+
}
20870+
20871+
// Look up the user-defined mapper given the mapper name and mapper type,
20872+
// return true if found one.
20873+
static bool hasUserDefinedMapper(Sema &SemaRef, Scope *S,
20874+
CXXScopeSpec &MapperIdScopeSpec,
20875+
const DeclarationNameInfo &MapperId,
20876+
QualType Type) {
20877+
// Find all user-defined mappers with the given MapperId.
20878+
SmallVector<UnresolvedSet<8>, 4> Lookups;
20879+
LookupResult Lookup(SemaRef, MapperId, Sema::LookupOMPMapperName);
20880+
Lookup.suppressDiagnostics();
20881+
while (S && SemaRef.LookupParsedName(Lookup, S, &MapperIdScopeSpec,
20882+
/*ObjectType=*/QualType())) {
20883+
NamedDecl *D = Lookup.getRepresentativeDecl();
20884+
while (S && !S->isDeclScope(D))
20885+
S = S->getParent();
20886+
if (S)
20887+
S = S->getParent();
20888+
Lookups.emplace_back();
20889+
Lookups.back().append(Lookup.begin(), Lookup.end());
20890+
Lookup.clear();
20891+
}
20892+
if (SemaRef.CurContext->isDependentContext() || Type->isDependentType() ||
20893+
Type->isInstantiationDependentType() ||
20894+
Type->containsUnexpandedParameterPack() ||
20895+
filterLookupForUDReductionAndMapper<bool>(Lookups, [](ValueDecl *D) {
20896+
return !D->isInvalidDecl() &&
20897+
(D->getType()->isDependentType() ||
20898+
D->getType()->isInstantiationDependentType() ||
20899+
D->getType()->containsUnexpandedParameterPack());
20900+
}))
20901+
return false;
20902+
// Perform argument dependent lookup.
20903+
SourceLocation Loc = MapperId.getLoc();
20904+
if (SemaRef.getLangOpts().CPlusPlus && !MapperIdScopeSpec.isSet())
20905+
argumentDependentLookup(SemaRef, MapperId, Loc, Type, Lookups);
20906+
if (filterLookupForUDReductionAndMapper<ValueDecl *>(
20907+
Lookups, [&SemaRef, Type](ValueDecl *D) -> ValueDecl * {
20908+
if (!D->isInvalidDecl() &&
20909+
SemaRef.Context.hasSameType(D->getType(), Type))
20910+
return D;
20911+
return nullptr;
20912+
}))
20913+
return true;
20914+
// Find the first user-defined mapper with a type derived from the desired
20915+
// type.
20916+
auto *VD = filterLookupForUDReductionAndMapper<ValueDecl *>(
20917+
Lookups, [&SemaRef, Type, Loc](ValueDecl *D) -> ValueDecl * {
20918+
if (!D->isInvalidDecl() &&
20919+
SemaRef.IsDerivedFrom(Loc, Type, D->getType()) &&
20920+
!Type.isMoreQualifiedThan(D->getType()))
20921+
return D;
20922+
return nullptr;
20923+
});
20924+
if (!VD)
20925+
return false;
20926+
CXXBasePaths Paths(/*FindAmbiguities=*/true, /*RecordPaths=*/true,
20927+
/*DetectVirtual=*/false);
20928+
if (SemaRef.IsDerivedFrom(Loc, Type, VD->getType(), Paths)) {
20929+
bool IsAmbiguous = !Paths.isAmbiguous(
20930+
SemaRef.Context.getCanonicalType(VD->getType().getUnqualifiedType()));
20931+
if (IsAmbiguous)
20932+
return false;
20933+
if (SemaRef.CheckBaseClassAccess(Loc, VD->getType(), Type, Paths.front(),
20934+
/*DiagID=*/0) != Sema::AR_inaccessible)
20935+
return true;
20936+
}
20937+
return false;
20938+
}
20939+
20940+
static bool isImplicitMapperNeeded(Sema &S, DSAStackTy *Stack,
20941+
QualType CanonType, const Expr *E) {
20942+
20943+
// DFS over data members in structures/classes.
20944+
SmallVector<std::pair<QualType, FieldDecl *>, 4> Types(1,
20945+
{CanonType, nullptr});
20946+
llvm::DenseMap<const Type *, bool> Visited;
20947+
SmallVector<std::pair<FieldDecl *, unsigned>, 4> ParentChain(1, {nullptr, 1});
20948+
while (!Types.empty()) {
20949+
auto [BaseType, CurFD] = Types.pop_back_val();
20950+
while (ParentChain.back().second == 0)
20951+
ParentChain.pop_back();
20952+
--ParentChain.back().second;
20953+
if (BaseType.isNull())
20954+
continue;
20955+
// Only structs/classes are allowed to have mappers.
20956+
const RecordDecl *RD = BaseType.getCanonicalType()->getAsRecordDecl();
20957+
if (!RD)
20958+
continue;
20959+
auto It = Visited.find(BaseType.getTypePtr());
20960+
if (It == Visited.end()) {
20961+
// Try to find the associated user-defined mapper.
20962+
CXXScopeSpec MapperIdScopeSpec;
20963+
DeclarationNameInfo DefaultMapperId;
20964+
DefaultMapperId.setName(S.Context.DeclarationNames.getIdentifier(
20965+
&S.Context.Idents.get("default")));
20966+
DefaultMapperId.setLoc(E->getExprLoc());
20967+
bool HasUDMapper =
20968+
hasUserDefinedMapper(S, Stack->getCurScope(), MapperIdScopeSpec,
20969+
DefaultMapperId, BaseType);
20970+
It = Visited.try_emplace(BaseType.getTypePtr(), HasUDMapper).first;
20971+
}
20972+
// Found default mapper.
20973+
if (It->second)
20974+
return true;
20975+
// Check for the "default" mapper for data members.
20976+
bool FirstIter = true;
20977+
for (FieldDecl *FD : RD->fields()) {
20978+
if (!FD)
20979+
continue;
20980+
QualType FieldTy = FD->getType();
20981+
if (FieldTy.isNull() ||
20982+
!(FieldTy->isStructureOrClassType() || FieldTy->isUnionType()))
20983+
continue;
20984+
if (FirstIter) {
20985+
FirstIter = false;
20986+
ParentChain.emplace_back(CurFD, 1);
20987+
} else {
20988+
++ParentChain.back().second;
20989+
}
20990+
Types.emplace_back(FieldTy, FD);
20991+
}
20992+
}
20993+
return false;
20994+
}
20995+
2079920996
// Check the validity of the provided variable list for the provided clause kind
2080020997
// \a CKind. In the check process the valid expressions, mappable expression
2080120998
// components, variables, and user-defined mappers are extracted and used to
@@ -21095,6 +21292,24 @@ static void checkMappableExpressionList(
2109521292
Type.getCanonicalType(), UnresolvedMapper);
2109621293
if (ER.isInvalid())
2109721294
continue;
21295+
if (!ER.get() && isa<ArraySectionExpr>(VE)) {
21296+
// Create implicit mapper as needed.
21297+
QualType BaseType = VE->getType().getCanonicalType();
21298+
if (BaseType->isSpecificBuiltinType(BuiltinType::ArraySection)) {
21299+
const auto *OASE = cast<ArraySectionExpr>(VE->IgnoreParenImpCasts());
21300+
QualType BType = ArraySectionExpr::getBaseOriginalType(OASE->getBase());
21301+
QualType ElemType;
21302+
if (const auto *ATy = BType->getAsArrayTypeUnsafe())
21303+
ElemType = ATy->getElementType();
21304+
else
21305+
ElemType = BType->getPointeeType();
21306+
BaseType = ElemType.getCanonicalType();
21307+
}
21308+
if (BaseType->getAsRecordDecl() &&
21309+
isImplicitMapperNeeded(SemaRef, DSAS, BaseType, VE)) {
21310+
ER = buildImplicitMapper(SemaRef, BaseType, DSAS);
21311+
}
21312+
}
2109821313
MVLI.UDMapperList.push_back(ER.get());
2109921314

2110021315
// Save the current expression.

‎clang/test/OpenMP/target_map_nest_defalut_mapper_ast_dump.cpp

Lines changed: 34 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,34 @@
1+
//RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -fopenmp -ast-dump %s | FileCheck %s --check-prefix=DUM
2+
3+
typedef struct {
4+
int a;
5+
} C;
6+
#pragma omp declare mapper(C s) map(to : s.a)
7+
8+
typedef struct {
9+
int e;
10+
C f;
11+
int h;
12+
} D;
13+
14+
void foo() {
15+
D sa[10];
16+
sa[1].e = 111;
17+
sa[1].f.a = 222;
18+
19+
#pragma omp target map(tofrom : sa [0:2])
20+
{
21+
sa[0].e = 333;
22+
sa[1].f.a = 444;
23+
}
24+
}
25+
26+
// DUM: -OMPDeclareMapperDecl{{.*}}<<invalid sloc>> <invalid sloc>
27+
// DUM-NEXT: |-OMPMapClause {{.*}}<<invalid sloc>> <implicit>
28+
// DUM-NEXT: | |-MemberExpr {{.*}}<line:9:3> 'int' lvalue .e
29+
// DUM-NEXT: | | `-DeclRefExpr {{.*}}<<invalid sloc>> 'D' lvalue Var {{.*}} '_s' 'D'
30+
// DUM-NEXT: | |-MemberExpr {{.*}}<line:10:3> 'C' lvalue .f {{.*}}
31+
// DUM-NEXT: | | `-DeclRefExpr {{.*}}<<invalid sloc>> 'D' lvalue Var {{.*}} '_s' 'D'
32+
// DUM-NEXT: | `-MemberExpr {{.*}}<line:11:3> 'int' lvalue .h {{.*}}
33+
// DUM-NEXT: | `-DeclRefExpr {{.*}}<<invalid sloc>> 'D' lvalue Var {{.*}} '_s' 'D'
34+
// DUM-NEXT: `-VarDecl {{.*}} <line:12:1> col:1 implicit used _s 'D'

0 commit comments

Comments
 (0)