[clang] [llvm] [OpenMP]Generate implicit default mapper for mapping array section. (PR #101101)
via cfe-commits
cfe-commits at lists.llvm.org
Thu Aug 1 23:49:02 PDT 2024
https://github.com/jyu2-git updated https://github.com/llvm/llvm-project/pull/101101
>From 1f6c326c6b93ebdbb9f4b35adfc903409ab67620 Mon Sep 17 00:00:00 2001
From: Jennifer Yu <jennifer.yu at intel.com>
Date: Mon, 29 Jul 2024 14:18:48 -0700
Subject: [PATCH 1/6] Generate implicit default mapper for mapping array
section.
This is only for struct containing nested structs with custom mappers.
Add three functions:
1>buildImplicitMap: build map for default mapper
2>buildImplicitMapper: build default mapper.
3 processImplicitMapperWithMaps: go throuth each map clause create
mapper as needed.
In processImplicitMapsWithDefaultMappers, when nested user defined
mapper is found, create a clause list (ClausesNeedImplicitMapper)
to generat mapper for the corespoing map clause.
---
clang/include/clang/AST/OpenMPClause.h | 2 +-
clang/lib/Sema/SemaOpenMP.cpp | 160 ++++++++
...et_map_pointer_defalut_mapper_ast_dump.cpp | 34 ++
...get_map_pointer_defalut_mapper_codegen.cpp | 356 ++++++++++++++++++
...eclare_mapper_nested_default_mappers_1.cpp | 34 ++
5 files changed, 585 insertions(+), 1 deletion(-)
create mode 100644 clang/test/OpenMP/target_map_pointer_defalut_mapper_ast_dump.cpp
create mode 100644 clang/test/OpenMP/target_map_pointer_defalut_mapper_codegen.cpp
create mode 100644 offload/test/mapping/declare_mapper_nested_default_mappers_1.cpp
diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h
index 325a1baa44614..ffd4e09d73468 100644
--- a/clang/include/clang/AST/OpenMPClause.h
+++ b/clang/include/clang/AST/OpenMPClause.h
@@ -5556,6 +5556,7 @@ class OMPMappableExprListClause : public OMPVarListClause<T>,
MapperIdInfo = MapperId;
}
+public:
/// Get the user-defined mapper references that are in the trailing objects of
/// the class.
MutableArrayRef<Expr *> getUDMapperRefs() {
@@ -5588,7 +5589,6 @@ class OMPMappableExprListClause : public OMPVarListClause<T>,
std::copy(DMDs.begin(), DMDs.end(), getUDMapperRefs().begin());
}
-public:
/// Return the number of unique base declarations in this clause.
unsigned getUniqueDeclarationsNum() const { return NumUniqueDeclarations; }
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 9c80b3eec914c..8a02f7eb71725 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -5753,6 +5753,155 @@ static ExprResult buildUserDefinedMapperRef(Sema &SemaRef, Scope *S,
QualType Type,
Expr *UnresolvedMapper);
+static std::pair<DeclRefExpr *, VarDecl *>
+buildImplicitMap(Sema &S, QualType BaseType, DSAStackTy *Stack,
+ SmallVectorImpl<OMPClause *> &Maps) {
+
+ const RecordDecl *RD = BaseType->getAsRecordDecl();
+ // AST context is RD's ParentASTContext().
+ ASTContext &Ctx = RD->getParentASTContext();
+ // DeclContext is RD's DeclContext.
+ DeclContext *DCT = const_cast<DeclContext *>(RD->getDeclContext());
+ SourceRange Range = RD->getSourceRange();
+ DeclarationNameInfo ImplicitName;
+ // Dummy variable _s for Mapper.
+ ImplicitName.setName(
+ Ctx.DeclarationNames.getIdentifier(&Ctx.Idents.get("_s")));
+ DeclarationName VN = ImplicitName.getName();
+ TypeSourceInfo *TInfo =
+ Ctx.getTrivialTypeSourceInfo(BaseType, Range.getEnd());
+ VarDecl *VD =
+ VarDecl::Create(Ctx, DCT, Range.getEnd(), Range.getEnd(),
+ VN.getAsIdentifierInfo(), BaseType, TInfo, SC_None);
+ DeclRefExpr *MapperVarRef =
+ buildDeclRefExpr(S, VD, BaseType, SourceLocation());
+
+ // Create implicit map clause for mapper.
+ SmallVector<Expr *, 4> SExprs;
+ for (auto *FD : RD->fields()) {
+ Expr *BE = S.BuildMemberExpr(
+ MapperVarRef, /*IsArrow=*/false, Range.getBegin(),
+ NestedNameSpecifierLoc(), Range.getBegin(), FD,
+ DeclAccessPair::make(FD, FD->getAccess()),
+ /*HadMultipleCandidates=*/false,
+ DeclarationNameInfo(FD->getDeclName(), FD->getSourceRange().getBegin()),
+ FD->getType(), VK_LValue, OK_Ordinary);
+ SExprs.push_back(BE);
+ }
+ CXXScopeSpec MapperIdScopeSpec;
+ DeclarationNameInfo MapperId;
+ OpenMPDirectiveKind DKind = Stack->getCurrentDirective();
+
+ OMPClause *MapClasue = S.OpenMP().ActOnOpenMPMapClause(
+ nullptr, OMPC_MAP_MODIFIER_unknown, SourceLocation(), MapperIdScopeSpec,
+ MapperId, DKind == OMPD_target_enter_data ? OMPC_MAP_to : OMPC_MAP_tofrom,
+ /*IsMapTypeImplicit=*/true, SourceLocation(), SourceLocation(), SExprs,
+ OMPVarListLocTy());
+ Maps.push_back(MapClasue);
+ return {MapperVarRef, VD};
+}
+
+static void buildImplicitMapper(Sema &S, QualType BaseType, DSAStackTy *Stack,
+ SmallVectorImpl<Expr *> &UDMapperRefs) {
+
+ // Build impilicit map for mapper
+ SmallVector<OMPClause *, 4> Maps;
+ VarDecl *VD;
+ DeclRefExpr *MapperVarRef;
+ std::tie(MapperVarRef, VD) = buildImplicitMap(S, BaseType, Stack, Maps);
+
+ const RecordDecl *RD = BaseType->getAsRecordDecl();
+ // AST context is RD's ParentASTContext().
+ ASTContext &Ctx = RD->getParentASTContext();
+ // DeclContext is RD's DeclContext.
+ DeclContext *DCT = const_cast<DeclContext *>(RD->getDeclContext());
+
+ // Create implicit default mapper for "RD".
+ DeclarationName MapperId;
+ auto &DeclNames = Ctx.DeclarationNames;
+ MapperId = DeclNames.getIdentifier(&Ctx.Idents.get("default"));
+ OMPDeclareMapperDecl *DMD = OMPDeclareMapperDecl::Create(
+ Ctx, DCT, SourceLocation(), MapperId, BaseType, MapperId, Maps, nullptr);
+ Scope *Scope = S.getScopeForContext(DCT);
+ if (Scope)
+ S.PushOnScopeChains(DMD, Scope, /*AddToContext*/ false);
+ DCT->addDecl(DMD);
+ DMD->setAccess(clang::AS_none);
+ VD->setDeclContext(DMD);
+ VD->setLexicalDeclContext(DMD);
+ DMD->addDecl(VD);
+ DMD->setMapperVarRef(MapperVarRef);
+ FieldDecl *FD = *RD->field_begin();
+ // create mapper refence.
+ DeclRefExpr *UDMapperRef =
+ DeclRefExpr::Create(Ctx, NestedNameSpecifierLoc{}, FD->getLocation(), DMD,
+ false, SourceLocation(), BaseType, VK_LValue);
+ UDMapperRefs.push_back(UDMapperRef);
+}
+
+static void
+processImplicitMapperWithMaps(Sema &S, DSAStackTy *Stack,
+ llvm::DenseMap<const Expr *, QualType> &MET,
+ SmallVectorImpl<OMPClause *> &Clauses) {
+
+ if (Stack->getCurrentDirective() == OMPD_unknown)
+ // declare mapper.
+ return;
+
+ for (int Cnt = 0, EndCnt = Clauses.size(); Cnt < EndCnt; ++Cnt) {
+ auto *C = dyn_cast<OMPMapClause>(Clauses[Cnt]);
+ if (!C || C->isImplicit())
+ continue;
+ SmallVector<Expr *, 4> UDMapperRefs;
+ auto *MI = C->mapperlist_begin();
+ auto *UDMapperRefI = C->getUDMapperRefs().begin();
+ for (auto I = C->varlist_begin(), End = C->varlist_end(); I != End;
+ ++I, ++MI, ++UDMapperRefI) {
+ // Expression is mapped using mapper - skip it.
+ if (*MI) {
+ UDMapperRefs.push_back(*UDMapperRefI);
+ continue;
+ }
+ Expr *E = *I;
+ if (MET.find(E) == MET.end()) {
+ UDMapperRefs.push_back(*UDMapperRefI);
+ continue;
+ }
+ // Array section - need to check for the mapping of the array section
+ // element.
+ QualType BaseType = E->getType().getCanonicalType();
+ if (BaseType->isSpecificBuiltinType(BuiltinType::ArraySection)) {
+ const auto *OASE = cast<ArraySectionExpr>(E->IgnoreParenImpCasts());
+ QualType BType = ArraySectionExpr::getBaseOriginalType(OASE->getBase());
+ QualType ElemType;
+ if (const auto *ATy = BType->getAsArrayTypeUnsafe())
+ ElemType = ATy->getElementType();
+ else
+ ElemType = BType->getPointeeType();
+ BaseType = ElemType.getCanonicalType();
+ }
+ CXXScopeSpec MapperIdScopeSpec;
+ DeclarationNameInfo DefaultMapperId;
+ DefaultMapperId.setName(S.Context.DeclarationNames.getIdentifier(
+ &S.Context.Idents.get("default")));
+ DefaultMapperId.setLoc(SourceLocation());
+ ExprResult ER = buildUserDefinedMapperRef(
+ S, Stack->getCurScope(), MapperIdScopeSpec, DefaultMapperId, BaseType,
+ /*UnresolvedMapper=*/nullptr);
+ if (ER.get()) {
+ UDMapperRefs.push_back(ER.get());
+ continue;
+ }
+ buildImplicitMapper(S, BaseType, Stack, UDMapperRefs);
+ }
+ if (!UDMapperRefs.empty()) {
+ assert(UDMapperRefs.size() == C->varlist_size());
+ // Update mapper in C->mapper_lists.
+ C->setUDMapperRefs(UDMapperRefs);
+ }
+ }
+}
+
/// Perform DFS through the structure/class data members trying to find
/// member(s) with user-defined 'default' mapper and generate implicit map
/// clauses for such members with the found 'default' mapper.
@@ -5763,6 +5912,8 @@ processImplicitMapsWithDefaultMappers(Sema &S, DSAStackTy *Stack,
if (S.getLangOpts().OpenMP < 50)
return;
SmallVector<OMPClause *, 4> ImplicitMaps;
+ SmallVector<OMPClause *, 4> ClausesNeedImplicitMapper;
+ llvm::DenseMap<const Expr *, QualType> ExprsNeedMapper;
for (int Cnt = 0, EndCnt = Clauses.size(); Cnt < EndCnt; ++Cnt) {
auto *C = dyn_cast<OMPMapClause>(Clauses[Cnt]);
if (!C)
@@ -5831,6 +5982,12 @@ processImplicitMapsWithDefaultMappers(Sema &S, DSAStackTy *Stack,
}
// Found default mapper.
if (It->second) {
+ if (isa<ArraySectionExpr>(E)) {
+ // For array section, mapper needs to be created.
+ ClausesNeedImplicitMapper.push_back(C);
+ ExprsNeedMapper.insert({E, BaseType});
+ continue;
+ }
auto *OE = new (S.Context) OpaqueValueExpr(E->getExprLoc(), CanonType,
VK_LValue, OK_Ordinary, E);
OE->setIsUnique(/*V=*/true);
@@ -5886,6 +6043,9 @@ processImplicitMapsWithDefaultMappers(Sema &S, DSAStackTy *Stack,
SubExprs, OMPVarListLocTy()))
Clauses.push_back(NewClause);
}
+ if (!ClausesNeedImplicitMapper.empty())
+ processImplicitMapperWithMaps(S, Stack, ExprsNeedMapper,
+ ClausesNeedImplicitMapper);
}
namespace {
diff --git a/clang/test/OpenMP/target_map_pointer_defalut_mapper_ast_dump.cpp b/clang/test/OpenMP/target_map_pointer_defalut_mapper_ast_dump.cpp
new file mode 100644
index 0000000000000..d7fcf96145722
--- /dev/null
+++ b/clang/test/OpenMP/target_map_pointer_defalut_mapper_ast_dump.cpp
@@ -0,0 +1,34 @@
+//RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -fopenmp -ast-dump %s | FileCheck %s --check-prefix=DUM
+
+typedef struct {
+ int a;
+} C;
+#pragma omp declare mapper(C s) map(to : s.a)
+
+typedef struct {
+ int e;
+ C f;
+ int h;
+} D;
+
+void foo() {
+ D sa[10];
+ sa[1].e = 111;
+ sa[1].f.a = 222;
+
+#pragma omp target map(tofrom : sa [0:2])
+ {
+ sa[1].e = 333;
+ sa[2].f.a = 444;
+ }
+}
+
+// DUM: -OMPDeclareMapperDecl{{.*}}<<invalid sloc>> <invalid sloc>
+// DUM-NEXT: |-OMPMapClause {{.*}}<<invalid sloc>> <implicit>
+// DUM-NEXT: | |-MemberExpr {{.*}}<line:9:3> 'int' lvalue .e
+// DUM-NEXT: | | `-DeclRefExpr {{.*}}<<invalid sloc>> 'D' lvalue Var {{.*}} '_s' 'D'
+// DUM-NEXT: | |-MemberExpr {{.*}}<line:10:3> 'C' lvalue .f {{.*}}
+// DUM-NEXT: | | `-DeclRefExpr {{.*}}<<invalid sloc>> 'D' lvalue Var {{.*}} '_s' 'D'
+// DUM-NEXT: | `-MemberExpr {{.*}}<line:11:3> 'int' lvalue .h {{.*}}
+// DUM-NEXT: | `-DeclRefExpr {{.*}}<<invalid sloc>> 'D' lvalue Var {{.*}} '_s' 'D'
+// DUM-NEXT: `-VarDecl {{.*}} <line:12:1> col:1 used _s 'D'
diff --git a/clang/test/OpenMP/target_map_pointer_defalut_mapper_codegen.cpp b/clang/test/OpenMP/target_map_pointer_defalut_mapper_codegen.cpp
new file mode 100644
index 0000000000000..98345ca39ace2
--- /dev/null
+++ b/clang/test/OpenMP/target_map_pointer_defalut_mapper_codegen.cpp
@@ -0,0 +1,356 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+typedef struct {
+ int a;
+} C;
+#pragma omp declare mapper(C s) map(to : s.a)
+
+typedef struct {
+ int e;
+ C f;
+ int h;
+} D;
+
+void foo() {
+ D sa[10];
+ sa[1].e = 111;
+ sa[1].f.a = 222;
+
+#pragma omp target map(tofrom : sa [0:2])
+ {
+ sa[1].e = 333;
+ sa[1].f.a = 444;
+ }
+}
+#endif
+// CHECK-LABEL: define {{[^@]+}}@_Z3foov
+// CHECK-SAME: () #[[ATTR0:[0-9]+]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[SA:%.*]] = alloca [10 x %struct.D], align 4
+// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8
+// CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8
+// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8
+// CHECK-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8
+// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x %struct.D], ptr [[SA]], i64 0, i64 1
+// CHECK-NEXT: [[E:%.*]] = getelementptr inbounds [[STRUCT_D:%.*]], ptr [[ARRAYIDX]], i32 0, i32 0
+// CHECK-NEXT: store i32 111, ptr [[E]], align 4
+// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [10 x %struct.D], ptr [[SA]], i64 0, i64 1
+// CHECK-NEXT: [[F:%.*]] = getelementptr inbounds [[STRUCT_D]], ptr [[ARRAYIDX1]], i32 0, i32 1
+// CHECK-NEXT: [[A:%.*]] = getelementptr inbounds [[STRUCT_C:%.*]], ptr [[F]], i32 0, i32 0
+// CHECK-NEXT: store i32 222, ptr [[A]], align 4
+// CHECK-NEXT: [[ARRAYIDX2:%.*]] = getelementptr inbounds [10 x %struct.D], ptr [[SA]], i64 0, i64 0
+// CHECK-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK-NEXT: store ptr [[SA]], ptr [[TMP0]], align 8
+// CHECK-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK-NEXT: store ptr [[ARRAYIDX2]], ptr [[TMP1]], align 8
+// CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
+// CHECK-NEXT: store ptr @.omp_mapper._ZTS1D.default, ptr [[TMP2]], align 8
+// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0
+// CHECK-NEXT: store i32 3, ptr [[TMP5]], align 4
+// CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1
+// CHECK-NEXT: store i32 1, ptr [[TMP6]], align 4
+// CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2
+// CHECK-NEXT: store ptr [[TMP3]], ptr [[TMP7]], align 8
+// CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3
+// CHECK-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 8
+// CHECK-NEXT: [[TMP9:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4
+// CHECK-NEXT: store ptr @.offload_sizes, ptr [[TMP9]], align 8
+// CHECK-NEXT: [[TMP10:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5
+// CHECK-NEXT: store ptr @.offload_maptypes, ptr [[TMP10]], align 8
+// CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6
+// CHECK-NEXT: store ptr null, ptr [[TMP11]], align 8
+// CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7
+// CHECK-NEXT: store ptr [[DOTOFFLOAD_MAPPERS]], ptr [[TMP12]], align 8
+// CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8
+// CHECK-NEXT: store i64 0, ptr [[TMP13]], align 8
+// CHECK-NEXT: [[TMP14:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9
+// CHECK-NEXT: store i64 0, ptr [[TMP14]], align 8
+// CHECK-NEXT: [[TMP15:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10
+// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP15]], align 4
+// CHECK-NEXT: [[TMP16:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11
+// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4
+// CHECK-NEXT: [[TMP17:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12
+// CHECK-NEXT: store i32 0, ptr [[TMP17]], align 4
+// CHECK-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l26.region_id, ptr [[KERNEL_ARGS]])
+// CHECK-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0
+// CHECK-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
+// CHECK: omp_offload.failed:
+// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l26(ptr [[SA]]) #[[ATTR3:[0-9]+]]
+// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT]]
+// CHECK: omp_offload.cont:
+// CHECK-NEXT: ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l26
+// CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(120) [[SA:%.*]]) #[[ATTR1:[0-9]+]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[SA_ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: store ptr [[SA]], ptr [[SA_ADDR]], align 8
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SA_ADDR]], align 8
+// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x %struct.D], ptr [[TMP0]], i64 0, i64 1
+// CHECK-NEXT: [[E:%.*]] = getelementptr inbounds [[STRUCT_D:%.*]], ptr [[ARRAYIDX]], i32 0, i32 0
+// CHECK-NEXT: store i32 333, ptr [[E]], align 4
+// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [10 x %struct.D], ptr [[TMP0]], i64 0, i64 1
+// CHECK-NEXT: [[F:%.*]] = getelementptr inbounds [[STRUCT_D]], ptr [[ARRAYIDX1]], i32 0, i32 1
+// CHECK-NEXT: [[A:%.*]] = getelementptr inbounds [[STRUCT_C:%.*]], ptr [[F]], i32 0, i32 0
+// CHECK-NEXT: store i32 444, ptr [[A]], align 4
+// CHECK-NEXT: ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@.omp_mapper._ZTS1D.default
+// CHECK-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]], ptr noundef [[TMP2:%.*]], i64 noundef [[TMP3:%.*]], i64 noundef [[TMP4:%.*]], ptr noundef [[TMP5:%.*]]) #[[ATTR2:[0-9]+]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: [[DOTADDR2:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: [[DOTADDR3:%.*]] = alloca i64, align 8
+// CHECK-NEXT: [[DOTADDR4:%.*]] = alloca i64, align 8
+// CHECK-NEXT: [[DOTADDR5:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 8
+// CHECK-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 8
+// CHECK-NEXT: store ptr [[TMP2]], ptr [[DOTADDR2]], align 8
+// CHECK-NEXT: store i64 [[TMP3]], ptr [[DOTADDR3]], align 8
+// CHECK-NEXT: store i64 [[TMP4]], ptr [[DOTADDR4]], align 8
+// CHECK-NEXT: store ptr [[TMP5]], ptr [[DOTADDR5]], align 8
+// CHECK-NEXT: [[TMP6:%.*]] = load i64, ptr [[DOTADDR3]], align 8
+// CHECK-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTADDR]], align 8
+// CHECK-NEXT: [[TMP8:%.*]] = load ptr, ptr [[DOTADDR1]], align 8
+// CHECK-NEXT: [[TMP9:%.*]] = load ptr, ptr [[DOTADDR2]], align 8
+// CHECK-NEXT: [[TMP10:%.*]] = udiv exact i64 [[TMP6]], 12
+// CHECK-NEXT: [[TMP11:%.*]] = getelementptr [[STRUCT_D:%.*]], ptr [[TMP9]], i64 [[TMP10]]
+// CHECK-NEXT: [[TMP12:%.*]] = load i64, ptr [[DOTADDR4]], align 8
+// CHECK-NEXT: [[TMP13:%.*]] = load ptr, ptr [[DOTADDR5]], align 8
+// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP10]], 1
+// CHECK-NEXT: [[TMP14:%.*]] = and i64 [[TMP12]], 8
+// CHECK-NEXT: [[TMP15:%.*]] = icmp ne ptr [[TMP8]], [[TMP9]]
+// CHECK-NEXT: [[TMP16:%.*]] = and i64 [[TMP12]], 16
+// CHECK-NEXT: [[TMP17:%.*]] = icmp ne i64 [[TMP16]], 0
+// CHECK-NEXT: [[TMP18:%.*]] = and i1 [[TMP15]], [[TMP17]]
+// CHECK-NEXT: [[TMP19:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP18]]
+// CHECK-NEXT: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP14]], 0
+// CHECK-NEXT: [[TMP20:%.*]] = and i1 [[TMP19]], [[DOTOMP_ARRAY__INIT__DELETE]]
+// CHECK-NEXT: br i1 [[TMP20]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]]
+// CHECK: .omp.array..init:
+// CHECK-NEXT: [[TMP21:%.*]] = mul nuw i64 [[TMP10]], 12
+// CHECK-NEXT: [[TMP22:%.*]] = and i64 [[TMP12]], -4
+// CHECK-NEXT: [[TMP23:%.*]] = or i64 [[TMP22]], 512
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[TMP8]], ptr [[TMP9]], i64 [[TMP21]], i64 [[TMP23]], ptr [[TMP13]])
+// CHECK-NEXT: br label [[OMP_ARRAYMAP_HEAD]]
+// CHECK: omp.arraymap.head:
+// CHECK-NEXT: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP9]], [[TMP11]]
+// CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISEMPTY]], label [[OMP_DONE:%.*]], label [[OMP_ARRAYMAP_BODY:%.*]]
+// CHECK: omp.arraymap.body:
+// CHECK-NEXT: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP9]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END25:%.*]] ]
+// CHECK-NEXT: [[E:%.*]] = getelementptr inbounds [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
+// CHECK-NEXT: [[F:%.*]] = getelementptr inbounds [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 1
+// CHECK-NEXT: [[H:%.*]] = getelementptr inbounds [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 2
+// CHECK-NEXT: [[TMP24:%.*]] = getelementptr i32, ptr [[H]], i32 1
+// CHECK-NEXT: [[TMP25:%.*]] = ptrtoint ptr [[TMP24]] to i64
+// CHECK-NEXT: [[TMP26:%.*]] = ptrtoint ptr [[E]] to i64
+// CHECK-NEXT: [[TMP27:%.*]] = sub i64 [[TMP25]], [[TMP26]]
+// CHECK-NEXT: [[TMP28:%.*]] = sdiv exact i64 [[TMP27]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
+// CHECK-NEXT: [[TMP29:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP7]])
+// CHECK-NEXT: [[TMP30:%.*]] = shl i64 [[TMP29]], 48
+// CHECK-NEXT: [[TMP31:%.*]] = add nuw i64 0, [[TMP30]]
+// CHECK-NEXT: [[TMP32:%.*]] = and i64 [[TMP12]], 3
+// CHECK-NEXT: [[TMP33:%.*]] = icmp eq i64 [[TMP32]], 0
+// CHECK-NEXT: br i1 [[TMP33]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]]
+// CHECK: omp.type.alloc:
+// CHECK-NEXT: [[TMP34:%.*]] = and i64 [[TMP31]], -4
+// CHECK-NEXT: br label [[OMP_TYPE_END:%.*]]
+// CHECK: omp.type.alloc.else:
+// CHECK-NEXT: [[TMP35:%.*]] = icmp eq i64 [[TMP32]], 1
+// CHECK-NEXT: br i1 [[TMP35]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]]
+// CHECK: omp.type.to:
+// CHECK-NEXT: [[TMP36:%.*]] = and i64 [[TMP31]], -3
+// CHECK-NEXT: br label [[OMP_TYPE_END]]
+// CHECK: omp.type.to.else:
+// CHECK-NEXT: [[TMP37:%.*]] = icmp eq i64 [[TMP32]], 2
+// CHECK-NEXT: br i1 [[TMP37]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]]
+// CHECK: omp.type.from:
+// CHECK-NEXT: [[TMP38:%.*]] = and i64 [[TMP31]], -2
+// CHECK-NEXT: br label [[OMP_TYPE_END]]
+// CHECK: omp.type.end:
+// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP34]], [[OMP_TYPE_ALLOC]] ], [ [[TMP36]], [[OMP_TYPE_TO]] ], [ [[TMP38]], [[OMP_TYPE_FROM]] ], [ [[TMP31]], [[OMP_TYPE_TO_ELSE]] ]
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP28]], i64 [[OMP_MAPTYPE]], ptr null)
+// CHECK-NEXT: [[TMP39:%.*]] = add nuw i64 281474976711171, [[TMP30]]
+// CHECK-NEXT: [[TMP40:%.*]] = and i64 [[TMP12]], 3
+// CHECK-NEXT: [[TMP41:%.*]] = icmp eq i64 [[TMP40]], 0
+// CHECK-NEXT: br i1 [[TMP41]], label [[OMP_TYPE_ALLOC6:%.*]], label [[OMP_TYPE_ALLOC_ELSE7:%.*]]
+// CHECK: omp.type.alloc6:
+// CHECK-NEXT: [[TMP42:%.*]] = and i64 [[TMP39]], -4
+// CHECK-NEXT: br label [[OMP_TYPE_END11:%.*]]
+// CHECK: omp.type.alloc.else7:
+// CHECK-NEXT: [[TMP43:%.*]] = icmp eq i64 [[TMP40]], 1
+// CHECK-NEXT: br i1 [[TMP43]], label [[OMP_TYPE_TO8:%.*]], label [[OMP_TYPE_TO_ELSE9:%.*]]
+// CHECK: omp.type.to8:
+// CHECK-NEXT: [[TMP44:%.*]] = and i64 [[TMP39]], -3
+// CHECK-NEXT: br label [[OMP_TYPE_END11]]
+// CHECK: omp.type.to.else9:
+// CHECK-NEXT: [[TMP45:%.*]] = icmp eq i64 [[TMP40]], 2
+// CHECK-NEXT: br i1 [[TMP45]], label [[OMP_TYPE_FROM10:%.*]], label [[OMP_TYPE_END11]]
+// CHECK: omp.type.from10:
+// CHECK-NEXT: [[TMP46:%.*]] = and i64 [[TMP39]], -2
+// CHECK-NEXT: br label [[OMP_TYPE_END11]]
+// CHECK: omp.type.end11:
+// CHECK-NEXT: [[OMP_MAPTYPE12:%.*]] = phi i64 [ [[TMP42]], [[OMP_TYPE_ALLOC6]] ], [ [[TMP44]], [[OMP_TYPE_TO8]] ], [ [[TMP46]], [[OMP_TYPE_FROM10]] ], [ [[TMP39]], [[OMP_TYPE_TO_ELSE9]] ]
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 4, i64 [[OMP_MAPTYPE12]], ptr null)
+// CHECK-NEXT: [[TMP47:%.*]] = add nuw i64 281474976711171, [[TMP30]]
+// CHECK-NEXT: [[TMP48:%.*]] = and i64 [[TMP12]], 3
+// CHECK-NEXT: [[TMP49:%.*]] = icmp eq i64 [[TMP48]], 0
+// CHECK-NEXT: br i1 [[TMP49]], label [[OMP_TYPE_ALLOC13:%.*]], label [[OMP_TYPE_ALLOC_ELSE14:%.*]]
+// CHECK: omp.type.alloc13:
+// CHECK-NEXT: [[TMP50:%.*]] = and i64 [[TMP47]], -4
+// CHECK-NEXT: br label [[OMP_TYPE_END18:%.*]]
+// CHECK: omp.type.alloc.else14:
+// CHECK-NEXT: [[TMP51:%.*]] = icmp eq i64 [[TMP48]], 1
+// CHECK-NEXT: br i1 [[TMP51]], label [[OMP_TYPE_TO15:%.*]], label [[OMP_TYPE_TO_ELSE16:%.*]]
+// CHECK: omp.type.to15:
+// CHECK-NEXT: [[TMP52:%.*]] = and i64 [[TMP47]], -3
+// CHECK-NEXT: br label [[OMP_TYPE_END18]]
+// CHECK: omp.type.to.else16:
+// CHECK-NEXT: [[TMP53:%.*]] = icmp eq i64 [[TMP48]], 2
+// CHECK-NEXT: br i1 [[TMP53]], label [[OMP_TYPE_FROM17:%.*]], label [[OMP_TYPE_END18]]
+// CHECK: omp.type.from17:
+// CHECK-NEXT: [[TMP54:%.*]] = and i64 [[TMP47]], -2
+// CHECK-NEXT: br label [[OMP_TYPE_END18]]
+// CHECK: omp.type.end18:
+// CHECK-NEXT: [[OMP_MAPTYPE19:%.*]] = phi i64 [ [[TMP50]], [[OMP_TYPE_ALLOC13]] ], [ [[TMP52]], [[OMP_TYPE_TO15]] ], [ [[TMP54]], [[OMP_TYPE_FROM17]] ], [ [[TMP47]], [[OMP_TYPE_TO_ELSE16]] ]
+// CHECK-NEXT: call void @.omp_mapper._ZTS1C.default(ptr [[TMP7]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[F]], i64 4, i64 [[OMP_MAPTYPE19]], ptr null) #[[ATTR3]]
+// CHECK-NEXT: [[TMP55:%.*]] = add nuw i64 281474976711171, [[TMP30]]
+// CHECK-NEXT: [[TMP56:%.*]] = and i64 [[TMP12]], 3
+// CHECK-NEXT: [[TMP57:%.*]] = icmp eq i64 [[TMP56]], 0
+// CHECK-NEXT: br i1 [[TMP57]], label [[OMP_TYPE_ALLOC20:%.*]], label [[OMP_TYPE_ALLOC_ELSE21:%.*]]
+// CHECK: omp.type.alloc20:
+// CHECK-NEXT: [[TMP58:%.*]] = and i64 [[TMP55]], -4
+// CHECK-NEXT: br label [[OMP_TYPE_END25]]
+// CHECK: omp.type.alloc.else21:
+// CHECK-NEXT: [[TMP59:%.*]] = icmp eq i64 [[TMP56]], 1
+// CHECK-NEXT: br i1 [[TMP59]], label [[OMP_TYPE_TO22:%.*]], label [[OMP_TYPE_TO_ELSE23:%.*]]
+// CHECK: omp.type.to22:
+// CHECK-NEXT: [[TMP60:%.*]] = and i64 [[TMP55]], -3
+// CHECK-NEXT: br label [[OMP_TYPE_END25]]
+// CHECK: omp.type.to.else23:
+// CHECK-NEXT: [[TMP61:%.*]] = icmp eq i64 [[TMP56]], 2
+// CHECK-NEXT: br i1 [[TMP61]], label [[OMP_TYPE_FROM24:%.*]], label [[OMP_TYPE_END25]]
+// CHECK: omp.type.from24:
+// CHECK-NEXT: [[TMP62:%.*]] = and i64 [[TMP55]], -2
+// CHECK-NEXT: br label [[OMP_TYPE_END25]]
+// CHECK: omp.type.end25:
+// CHECK-NEXT: [[OMP_MAPTYPE26:%.*]] = phi i64 [ [[TMP58]], [[OMP_TYPE_ALLOC20]] ], [ [[TMP60]], [[OMP_TYPE_TO22]] ], [ [[TMP62]], [[OMP_TYPE_FROM24]] ], [ [[TMP55]], [[OMP_TYPE_TO_ELSE23]] ]
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[H]], i64 4, i64 [[OMP_MAPTYPE26]], ptr null)
+// CHECK-NEXT: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1
+// CHECK-NEXT: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP11]]
+// CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]]
+// CHECK: omp.arraymap.exit:
+// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY27:%.*]] = icmp sgt i64 [[TMP10]], 1
+// CHECK-NEXT: [[TMP63:%.*]] = and i64 [[TMP12]], 8
+// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP63]], 0
+// CHECK-NEXT: [[TMP64:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY27]], [[DOTOMP_ARRAY__DEL__DELETE]]
+// CHECK-NEXT: br i1 [[TMP64]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
+// CHECK: .omp.array..del:
+// CHECK-NEXT: [[TMP65:%.*]] = mul nuw i64 [[TMP10]], 12
+// CHECK-NEXT: [[TMP66:%.*]] = and i64 [[TMP12]], -4
+// CHECK-NEXT: [[TMP67:%.*]] = or i64 [[TMP66]], 512
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[TMP8]], ptr [[TMP9]], i64 [[TMP65]], i64 [[TMP67]], ptr [[TMP13]])
+// CHECK-NEXT: br label [[OMP_DONE]]
+// CHECK: omp.done:
+// CHECK-NEXT: ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@.omp_mapper._ZTS1C.default
+// CHECK-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]], ptr noundef [[TMP2:%.*]], i64 noundef [[TMP3:%.*]], i64 noundef [[TMP4:%.*]], ptr noundef [[TMP5:%.*]]) #[[ATTR2]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: [[DOTADDR1:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: [[DOTADDR2:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: [[DOTADDR3:%.*]] = alloca i64, align 8
+// CHECK-NEXT: [[DOTADDR4:%.*]] = alloca i64, align 8
+// CHECK-NEXT: [[DOTADDR5:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 8
+// CHECK-NEXT: store ptr [[TMP1]], ptr [[DOTADDR1]], align 8
+// CHECK-NEXT: store ptr [[TMP2]], ptr [[DOTADDR2]], align 8
+// CHECK-NEXT: store i64 [[TMP3]], ptr [[DOTADDR3]], align 8
+// CHECK-NEXT: store i64 [[TMP4]], ptr [[DOTADDR4]], align 8
+// CHECK-NEXT: store ptr [[TMP5]], ptr [[DOTADDR5]], align 8
+// CHECK-NEXT: [[TMP6:%.*]] = load i64, ptr [[DOTADDR3]], align 8
+// CHECK-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTADDR]], align 8
+// CHECK-NEXT: [[TMP8:%.*]] = load ptr, ptr [[DOTADDR1]], align 8
+// CHECK-NEXT: [[TMP9:%.*]] = load ptr, ptr [[DOTADDR2]], align 8
+// CHECK-NEXT: [[TMP10:%.*]] = udiv exact i64 [[TMP6]], 4
+// CHECK-NEXT: [[TMP11:%.*]] = getelementptr [[STRUCT_C:%.*]], ptr [[TMP9]], i64 [[TMP10]]
+// CHECK-NEXT: [[TMP12:%.*]] = load i64, ptr [[DOTADDR4]], align 8
+// CHECK-NEXT: [[TMP13:%.*]] = load ptr, ptr [[DOTADDR5]], align 8
+// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP10]], 1
+// CHECK-NEXT: [[TMP14:%.*]] = and i64 [[TMP12]], 8
+// CHECK-NEXT: [[TMP15:%.*]] = icmp ne ptr [[TMP8]], [[TMP9]]
+// CHECK-NEXT: [[TMP16:%.*]] = and i64 [[TMP12]], 16
+// CHECK-NEXT: [[TMP17:%.*]] = icmp ne i64 [[TMP16]], 0
+// CHECK-NEXT: [[TMP18:%.*]] = and i1 [[TMP15]], [[TMP17]]
+// CHECK-NEXT: [[TMP19:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP18]]
+// CHECK-NEXT: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP14]], 0
+// CHECK-NEXT: [[TMP20:%.*]] = and i1 [[TMP19]], [[DOTOMP_ARRAY__INIT__DELETE]]
+// CHECK-NEXT: br i1 [[TMP20]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]]
+// CHECK: .omp.array..init:
+// CHECK-NEXT: [[TMP21:%.*]] = mul nuw i64 [[TMP10]], 4
+// CHECK-NEXT: [[TMP22:%.*]] = and i64 [[TMP12]], -4
+// CHECK-NEXT: [[TMP23:%.*]] = or i64 [[TMP22]], 512
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[TMP8]], ptr [[TMP9]], i64 [[TMP21]], i64 [[TMP23]], ptr [[TMP13]])
+// CHECK-NEXT: br label [[OMP_ARRAYMAP_HEAD]]
+// CHECK: omp.arraymap.head:
+// CHECK-NEXT: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP9]], [[TMP11]]
+// CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISEMPTY]], label [[OMP_DONE:%.*]], label [[OMP_ARRAYMAP_BODY:%.*]]
+// CHECK: omp.arraymap.body:
+// CHECK-NEXT: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP9]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END:%.*]] ]
+// CHECK-NEXT: [[A:%.*]] = getelementptr inbounds [[STRUCT_C]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
+// CHECK-NEXT: [[TMP24:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP7]])
+// CHECK-NEXT: [[TMP25:%.*]] = shl i64 [[TMP24]], 48
+// CHECK-NEXT: [[TMP26:%.*]] = add nuw i64 1, [[TMP25]]
+// CHECK-NEXT: [[TMP27:%.*]] = and i64 [[TMP12]], 3
+// CHECK-NEXT: [[TMP28:%.*]] = icmp eq i64 [[TMP27]], 0
+// CHECK-NEXT: br i1 [[TMP28]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]]
+// CHECK: omp.type.alloc:
+// CHECK-NEXT: [[TMP29:%.*]] = and i64 [[TMP26]], -4
+// CHECK-NEXT: br label [[OMP_TYPE_END]]
+// CHECK: omp.type.alloc.else:
+// CHECK-NEXT: [[TMP30:%.*]] = icmp eq i64 [[TMP27]], 1
+// CHECK-NEXT: br i1 [[TMP30]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]]
+// CHECK: omp.type.to:
+// CHECK-NEXT: [[TMP31:%.*]] = and i64 [[TMP26]], -3
+// CHECK-NEXT: br label [[OMP_TYPE_END]]
+// CHECK: omp.type.to.else:
+// CHECK-NEXT: [[TMP32:%.*]] = icmp eq i64 [[TMP27]], 2
+// CHECK-NEXT: br i1 [[TMP32]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]]
+// CHECK: omp.type.from:
+// CHECK-NEXT: [[TMP33:%.*]] = and i64 [[TMP26]], -2
+// CHECK-NEXT: br label [[OMP_TYPE_END]]
+// CHECK: omp.type.end:
+// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP29]], [[OMP_TYPE_ALLOC]] ], [ [[TMP31]], [[OMP_TYPE_TO]] ], [ [[TMP33]], [[OMP_TYPE_FROM]] ], [ [[TMP26]], [[OMP_TYPE_TO_ELSE]] ]
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[A]], i64 4, i64 [[OMP_MAPTYPE]], ptr null)
+// CHECK-NEXT: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_C]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1
+// CHECK-NEXT: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP11]]
+// CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]]
+// CHECK: omp.arraymap.exit:
+// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY6:%.*]] = icmp sgt i64 [[TMP10]], 1
+// CHECK-NEXT: [[TMP34:%.*]] = and i64 [[TMP12]], 8
+// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP34]], 0
+// CHECK-NEXT: [[TMP35:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY6]], [[DOTOMP_ARRAY__DEL__DELETE]]
+// CHECK-NEXT: br i1 [[TMP35]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
+// CHECK: .omp.array..del:
+// CHECK-NEXT: [[TMP36:%.*]] = mul nuw i64 [[TMP10]], 4
+// CHECK-NEXT: [[TMP37:%.*]] = and i64 [[TMP12]], -4
+// CHECK-NEXT: [[TMP38:%.*]] = or i64 [[TMP37]], 512
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[TMP8]], ptr [[TMP9]], i64 [[TMP36]], i64 [[TMP38]], ptr [[TMP13]])
+// CHECK-NEXT: br label [[OMP_DONE]]
+// CHECK: omp.done:
+// CHECK-NEXT: ret void
+//
diff --git a/offload/test/mapping/declare_mapper_nested_default_mappers_1.cpp b/offload/test/mapping/declare_mapper_nested_default_mappers_1.cpp
new file mode 100644
index 0000000000000..1658ce5f6070e
--- /dev/null
+++ b/offload/test/mapping/declare_mapper_nested_default_mappers_1.cpp
@@ -0,0 +1,34 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+extern "C" int printf(const char *, ...);
+
+typedef struct {
+ int a;
+} C;
+#pragma omp declare mapper(C s) map(to : s.a)
+
+typedef struct {
+ int e;
+ C f;
+ int h;
+} D;
+
+int main() {
+ D sa[10];
+ sa[1].e = 111;
+ sa[1].f.a = 222;
+
+ // CHECK: 111 222
+ printf("%d %d \n", sa[1].e, sa[1].f.a);
+#pragma omp target map(tofrom : sa[0 : 2])
+ {
+ // CHECK: 111
+ printf("%d \n", sa[1].e);
+ sa[0].e = 333;
+ sa[1].f.a = 444;
+ // CHECK: 333 444
+ printf("%d %d \n", sa[0].e, sa[1].f.a);
+ }
+ // CHECK: 333 222
+ printf("%d %d \n", sa[0].e, sa[1].f.a);
+}
>From 28fb0bd68f1feb2284c9000f2f36c0d91552f8ff Mon Sep 17 00:00:00 2001
From: Jennifer Yu <jennifer.yu at intel.com>
Date: Tue, 30 Jul 2024 10:11:20 -0700
Subject: [PATCH 2/6] Thanks Alexey for the review. This is address his
comments.
Rename the tests.
---
clang/lib/Sema/SemaOpenMP.cpp | 7 ++++---
...ump.cpp => target_map_nest_defalut_mapper_ast_dump.cpp} | 4 ++--
...egen.cpp => target_map_nest_defalut_mapper_codegen.cpp} | 0
3 files changed, 6 insertions(+), 5 deletions(-)
rename clang/test/OpenMP/{target_map_pointer_defalut_mapper_ast_dump.cpp => target_map_nest_defalut_mapper_ast_dump.cpp} (96%)
rename clang/test/OpenMP/{target_map_pointer_defalut_mapper_codegen.cpp => target_map_nest_defalut_mapper_codegen.cpp} (100%)
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 8a02f7eb71725..f57fe3f9fcd5a 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -5820,7 +5820,7 @@ static void buildImplicitMapper(Sema &S, QualType BaseType, DSAStackTy *Stack,
DeclarationName MapperId;
auto &DeclNames = Ctx.DeclarationNames;
MapperId = DeclNames.getIdentifier(&Ctx.Idents.get("default"));
- OMPDeclareMapperDecl *DMD = OMPDeclareMapperDecl::Create(
+ auto *DMD = OMPDeclareMapperDecl::Create(
Ctx, DCT, SourceLocation(), MapperId, BaseType, MapperId, Maps, nullptr);
Scope *Scope = S.getScopeForContext(DCT);
if (Scope)
@@ -5848,7 +5848,7 @@ processImplicitMapperWithMaps(Sema &S, DSAStackTy *Stack,
// declare mapper.
return;
- for (int Cnt = 0, EndCnt = Clauses.size(); Cnt < EndCnt; ++Cnt) {
+ for (int Cnt : llvm::seq<int>(0, Clauses.size())) {
auto *C = dyn_cast<OMPMapClause>(Clauses[Cnt]);
if (!C || C->isImplicit())
continue;
@@ -5895,7 +5895,8 @@ processImplicitMapperWithMaps(Sema &S, DSAStackTy *Stack,
buildImplicitMapper(S, BaseType, Stack, UDMapperRefs);
}
if (!UDMapperRefs.empty()) {
- assert(UDMapperRefs.size() == C->varlist_size());
+ assert(UDMapperRefs.size() == C->varlist_size() &&
+ "Unexpected number of default mappers.");
// Update mapper in C->mapper_lists.
C->setUDMapperRefs(UDMapperRefs);
}
diff --git a/clang/test/OpenMP/target_map_pointer_defalut_mapper_ast_dump.cpp b/clang/test/OpenMP/target_map_nest_defalut_mapper_ast_dump.cpp
similarity index 96%
rename from clang/test/OpenMP/target_map_pointer_defalut_mapper_ast_dump.cpp
rename to clang/test/OpenMP/target_map_nest_defalut_mapper_ast_dump.cpp
index d7fcf96145722..c088fc1d3c86f 100644
--- a/clang/test/OpenMP/target_map_pointer_defalut_mapper_ast_dump.cpp
+++ b/clang/test/OpenMP/target_map_nest_defalut_mapper_ast_dump.cpp
@@ -18,8 +18,8 @@ void foo() {
#pragma omp target map(tofrom : sa [0:2])
{
- sa[1].e = 333;
- sa[2].f.a = 444;
+ sa[0].e = 333;
+ sa[1].f.a = 444;
}
}
diff --git a/clang/test/OpenMP/target_map_pointer_defalut_mapper_codegen.cpp b/clang/test/OpenMP/target_map_nest_defalut_mapper_codegen.cpp
similarity index 100%
rename from clang/test/OpenMP/target_map_pointer_defalut_mapper_codegen.cpp
rename to clang/test/OpenMP/target_map_nest_defalut_mapper_codegen.cpp
>From d64a63c0ecd4e33cdb9719e4f34b922a0433d8f9 Mon Sep 17 00:00:00 2001
From: Jennifer Yu <jennifer.yu at intel.com>
Date: Tue, 30 Jul 2024 10:51:08 -0700
Subject: [PATCH 3/6] Fix format.
---
clang/lib/Sema/SemaOpenMP.cpp | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index f57fe3f9fcd5a..1203ff70b5981 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -5820,8 +5820,8 @@ static void buildImplicitMapper(Sema &S, QualType BaseType, DSAStackTy *Stack,
DeclarationName MapperId;
auto &DeclNames = Ctx.DeclarationNames;
MapperId = DeclNames.getIdentifier(&Ctx.Idents.get("default"));
- auto *DMD = OMPDeclareMapperDecl::Create(
- Ctx, DCT, SourceLocation(), MapperId, BaseType, MapperId, Maps, nullptr);
+ auto *DMD = OMPDeclareMapperDecl::Create(Ctx, DCT, SourceLocation(), MapperId,
+ BaseType, MapperId, Maps, nullptr);
Scope *Scope = S.getScopeForContext(DCT);
if (Scope)
S.PushOnScopeChains(DMD, Scope, /*AddToContext*/ false);
>From f03b49a4dcce8cfc9d8e1d93a8ca4a771c67d9ae Mon Sep 17 00:00:00 2001
From: Jennifer Yu <jennifer.yu at intel.com>
Date: Wed, 31 Jul 2024 12:51:33 -0700
Subject: [PATCH 4/6] Thanks Alexey for the review.
This is generate implicit mapper when map is created instead after map
created.
---
clang/include/clang/AST/OpenMPClause.h | 2 +-
clang/lib/Sema/SemaOpenMP.cpp | 321 +++++++++++++------------
2 files changed, 163 insertions(+), 160 deletions(-)
diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h
index ffd4e09d73468..325a1baa44614 100644
--- a/clang/include/clang/AST/OpenMPClause.h
+++ b/clang/include/clang/AST/OpenMPClause.h
@@ -5556,7 +5556,6 @@ class OMPMappableExprListClause : public OMPVarListClause<T>,
MapperIdInfo = MapperId;
}
-public:
/// Get the user-defined mapper references that are in the trailing objects of
/// the class.
MutableArrayRef<Expr *> getUDMapperRefs() {
@@ -5589,6 +5588,7 @@ class OMPMappableExprListClause : public OMPVarListClause<T>,
std::copy(DMDs.begin(), DMDs.end(), getUDMapperRefs().begin());
}
+public:
/// Return the number of unique base declarations in this clause.
unsigned getUniqueDeclarationsNum() const { return NumUniqueDeclarations; }
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 1203ff70b5981..74e6f730b775d 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -5753,156 +5753,6 @@ static ExprResult buildUserDefinedMapperRef(Sema &SemaRef, Scope *S,
QualType Type,
Expr *UnresolvedMapper);
-static std::pair<DeclRefExpr *, VarDecl *>
-buildImplicitMap(Sema &S, QualType BaseType, DSAStackTy *Stack,
- SmallVectorImpl<OMPClause *> &Maps) {
-
- const RecordDecl *RD = BaseType->getAsRecordDecl();
- // AST context is RD's ParentASTContext().
- ASTContext &Ctx = RD->getParentASTContext();
- // DeclContext is RD's DeclContext.
- DeclContext *DCT = const_cast<DeclContext *>(RD->getDeclContext());
- SourceRange Range = RD->getSourceRange();
- DeclarationNameInfo ImplicitName;
- // Dummy variable _s for Mapper.
- ImplicitName.setName(
- Ctx.DeclarationNames.getIdentifier(&Ctx.Idents.get("_s")));
- DeclarationName VN = ImplicitName.getName();
- TypeSourceInfo *TInfo =
- Ctx.getTrivialTypeSourceInfo(BaseType, Range.getEnd());
- VarDecl *VD =
- VarDecl::Create(Ctx, DCT, Range.getEnd(), Range.getEnd(),
- VN.getAsIdentifierInfo(), BaseType, TInfo, SC_None);
- DeclRefExpr *MapperVarRef =
- buildDeclRefExpr(S, VD, BaseType, SourceLocation());
-
- // Create implicit map clause for mapper.
- SmallVector<Expr *, 4> SExprs;
- for (auto *FD : RD->fields()) {
- Expr *BE = S.BuildMemberExpr(
- MapperVarRef, /*IsArrow=*/false, Range.getBegin(),
- NestedNameSpecifierLoc(), Range.getBegin(), FD,
- DeclAccessPair::make(FD, FD->getAccess()),
- /*HadMultipleCandidates=*/false,
- DeclarationNameInfo(FD->getDeclName(), FD->getSourceRange().getBegin()),
- FD->getType(), VK_LValue, OK_Ordinary);
- SExprs.push_back(BE);
- }
- CXXScopeSpec MapperIdScopeSpec;
- DeclarationNameInfo MapperId;
- OpenMPDirectiveKind DKind = Stack->getCurrentDirective();
-
- OMPClause *MapClasue = S.OpenMP().ActOnOpenMPMapClause(
- nullptr, OMPC_MAP_MODIFIER_unknown, SourceLocation(), MapperIdScopeSpec,
- MapperId, DKind == OMPD_target_enter_data ? OMPC_MAP_to : OMPC_MAP_tofrom,
- /*IsMapTypeImplicit=*/true, SourceLocation(), SourceLocation(), SExprs,
- OMPVarListLocTy());
- Maps.push_back(MapClasue);
- return {MapperVarRef, VD};
-}
-
-static void buildImplicitMapper(Sema &S, QualType BaseType, DSAStackTy *Stack,
- SmallVectorImpl<Expr *> &UDMapperRefs) {
-
- // Build impilicit map for mapper
- SmallVector<OMPClause *, 4> Maps;
- VarDecl *VD;
- DeclRefExpr *MapperVarRef;
- std::tie(MapperVarRef, VD) = buildImplicitMap(S, BaseType, Stack, Maps);
-
- const RecordDecl *RD = BaseType->getAsRecordDecl();
- // AST context is RD's ParentASTContext().
- ASTContext &Ctx = RD->getParentASTContext();
- // DeclContext is RD's DeclContext.
- DeclContext *DCT = const_cast<DeclContext *>(RD->getDeclContext());
-
- // Create implicit default mapper for "RD".
- DeclarationName MapperId;
- auto &DeclNames = Ctx.DeclarationNames;
- MapperId = DeclNames.getIdentifier(&Ctx.Idents.get("default"));
- auto *DMD = OMPDeclareMapperDecl::Create(Ctx, DCT, SourceLocation(), MapperId,
- BaseType, MapperId, Maps, nullptr);
- Scope *Scope = S.getScopeForContext(DCT);
- if (Scope)
- S.PushOnScopeChains(DMD, Scope, /*AddToContext*/ false);
- DCT->addDecl(DMD);
- DMD->setAccess(clang::AS_none);
- VD->setDeclContext(DMD);
- VD->setLexicalDeclContext(DMD);
- DMD->addDecl(VD);
- DMD->setMapperVarRef(MapperVarRef);
- FieldDecl *FD = *RD->field_begin();
- // create mapper refence.
- DeclRefExpr *UDMapperRef =
- DeclRefExpr::Create(Ctx, NestedNameSpecifierLoc{}, FD->getLocation(), DMD,
- false, SourceLocation(), BaseType, VK_LValue);
- UDMapperRefs.push_back(UDMapperRef);
-}
-
-static void
-processImplicitMapperWithMaps(Sema &S, DSAStackTy *Stack,
- llvm::DenseMap<const Expr *, QualType> &MET,
- SmallVectorImpl<OMPClause *> &Clauses) {
-
- if (Stack->getCurrentDirective() == OMPD_unknown)
- // declare mapper.
- return;
-
- for (int Cnt : llvm::seq<int>(0, Clauses.size())) {
- auto *C = dyn_cast<OMPMapClause>(Clauses[Cnt]);
- if (!C || C->isImplicit())
- continue;
- SmallVector<Expr *, 4> UDMapperRefs;
- auto *MI = C->mapperlist_begin();
- auto *UDMapperRefI = C->getUDMapperRefs().begin();
- for (auto I = C->varlist_begin(), End = C->varlist_end(); I != End;
- ++I, ++MI, ++UDMapperRefI) {
- // Expression is mapped using mapper - skip it.
- if (*MI) {
- UDMapperRefs.push_back(*UDMapperRefI);
- continue;
- }
- Expr *E = *I;
- if (MET.find(E) == MET.end()) {
- UDMapperRefs.push_back(*UDMapperRefI);
- continue;
- }
- // Array section - need to check for the mapping of the array section
- // element.
- QualType BaseType = E->getType().getCanonicalType();
- if (BaseType->isSpecificBuiltinType(BuiltinType::ArraySection)) {
- const auto *OASE = cast<ArraySectionExpr>(E->IgnoreParenImpCasts());
- QualType BType = ArraySectionExpr::getBaseOriginalType(OASE->getBase());
- QualType ElemType;
- if (const auto *ATy = BType->getAsArrayTypeUnsafe())
- ElemType = ATy->getElementType();
- else
- ElemType = BType->getPointeeType();
- BaseType = ElemType.getCanonicalType();
- }
- CXXScopeSpec MapperIdScopeSpec;
- DeclarationNameInfo DefaultMapperId;
- DefaultMapperId.setName(S.Context.DeclarationNames.getIdentifier(
- &S.Context.Idents.get("default")));
- DefaultMapperId.setLoc(SourceLocation());
- ExprResult ER = buildUserDefinedMapperRef(
- S, Stack->getCurScope(), MapperIdScopeSpec, DefaultMapperId, BaseType,
- /*UnresolvedMapper=*/nullptr);
- if (ER.get()) {
- UDMapperRefs.push_back(ER.get());
- continue;
- }
- buildImplicitMapper(S, BaseType, Stack, UDMapperRefs);
- }
- if (!UDMapperRefs.empty()) {
- assert(UDMapperRefs.size() == C->varlist_size() &&
- "Unexpected number of default mappers.");
- // Update mapper in C->mapper_lists.
- C->setUDMapperRefs(UDMapperRefs);
- }
- }
-}
-
/// Perform DFS through the structure/class data members trying to find
/// member(s) with user-defined 'default' mapper and generate implicit map
/// clauses for such members with the found 'default' mapper.
@@ -5983,12 +5833,6 @@ processImplicitMapsWithDefaultMappers(Sema &S, DSAStackTy *Stack,
}
// Found default mapper.
if (It->second) {
- if (isa<ArraySectionExpr>(E)) {
- // For array section, mapper needs to be created.
- ClausesNeedImplicitMapper.push_back(C);
- ExprsNeedMapper.insert({E, BaseType});
- continue;
- }
auto *OE = new (S.Context) OpaqueValueExpr(E->getExprLoc(), CanonType,
VK_LValue, OK_Ordinary, E);
OE->setIsUnique(/*V=*/true);
@@ -6044,9 +5888,6 @@ processImplicitMapsWithDefaultMappers(Sema &S, DSAStackTy *Stack,
SubExprs, OMPVarListLocTy()))
Clauses.push_back(NewClause);
}
- if (!ClausesNeedImplicitMapper.empty())
- processImplicitMapperWithMaps(S, Stack, ExprsNeedMapper,
- ClausesNeedImplicitMapper);
}
namespace {
@@ -20957,6 +20798,150 @@ struct MappableVarListInfo {
};
} // namespace
+static std::pair<DeclRefExpr *, VarDecl *>
+buildImplicitMap(Sema &S, QualType BaseType, DSAStackTy *Stack,
+ SmallVectorImpl<OMPClause *> &Maps) {
+
+ const RecordDecl *RD = BaseType->getAsRecordDecl();
+ // AST context is RD's ParentASTContext().
+ ASTContext &Ctx = RD->getParentASTContext();
+ // DeclContext is RD's DeclContext.
+ DeclContext *DCT = const_cast<DeclContext *>(RD->getDeclContext());
+ SourceRange Range = RD->getSourceRange();
+ DeclarationNameInfo ImplicitName;
+ // Dummy variable _s for Mapper.
+ ImplicitName.setName(
+ Ctx.DeclarationNames.getIdentifier(&Ctx.Idents.get("_s")));
+ DeclarationName VN = ImplicitName.getName();
+ TypeSourceInfo *TInfo =
+ Ctx.getTrivialTypeSourceInfo(BaseType, Range.getEnd());
+ VarDecl *VD =
+ VarDecl::Create(Ctx, DCT, Range.getEnd(), Range.getEnd(),
+ VN.getAsIdentifierInfo(), BaseType, TInfo, SC_None);
+ DeclRefExpr *MapperVarRef =
+ buildDeclRefExpr(S, VD, BaseType, SourceLocation());
+
+ // Create implicit map clause for mapper.
+ SmallVector<Expr *, 4> SExprs;
+ for (auto *FD : RD->fields()) {
+ Expr *BE = S.BuildMemberExpr(
+ MapperVarRef, /*IsArrow=*/false, Range.getBegin(),
+ NestedNameSpecifierLoc(), Range.getBegin(), FD,
+ DeclAccessPair::make(FD, FD->getAccess()),
+ /*HadMultipleCandidates=*/false,
+ DeclarationNameInfo(FD->getDeclName(), FD->getSourceRange().getBegin()),
+ FD->getType(), VK_LValue, OK_Ordinary);
+ SExprs.push_back(BE);
+ }
+ CXXScopeSpec MapperIdScopeSpec;
+ DeclarationNameInfo MapperId;
+ OpenMPDirectiveKind DKind = Stack->getCurrentDirective();
+
+ OMPClause *MapClasue = S.OpenMP().ActOnOpenMPMapClause(
+ nullptr, OMPC_MAP_MODIFIER_unknown, SourceLocation(), MapperIdScopeSpec,
+ MapperId, DKind == OMPD_target_enter_data ? OMPC_MAP_to : OMPC_MAP_tofrom,
+ /*IsMapTypeImplicit=*/true, SourceLocation(), SourceLocation(), SExprs,
+ OMPVarListLocTy());
+ Maps.push_back(MapClasue);
+ return {MapperVarRef, VD};
+}
+
+static ExprResult buildImplicitMapper(Sema &S, QualType BaseType,
+ DSAStackTy *Stack) {
+
+ // Build impilicit map for mapper
+ SmallVector<OMPClause *, 4> Maps;
+ VarDecl *VD;
+ DeclRefExpr *MapperVarRef;
+ std::tie(MapperVarRef, VD) = buildImplicitMap(S, BaseType, Stack, Maps);
+
+ const RecordDecl *RD = BaseType->getAsRecordDecl();
+ // AST context is RD's ParentASTContext().
+ ASTContext &Ctx = RD->getParentASTContext();
+ // DeclContext is RD's DeclContext.
+ DeclContext *DCT = const_cast<DeclContext *>(RD->getDeclContext());
+
+ // Create implicit default mapper for "RD".
+ DeclarationName MapperId;
+ auto &DeclNames = Ctx.DeclarationNames;
+ MapperId = DeclNames.getIdentifier(&Ctx.Idents.get("default"));
+ auto *DMD = OMPDeclareMapperDecl::Create(Ctx, DCT, SourceLocation(), MapperId,
+ BaseType, MapperId, Maps, nullptr);
+ Scope *Scope = S.getScopeForContext(DCT);
+ if (Scope)
+ S.PushOnScopeChains(DMD, Scope, /*AddToContext*/ false);
+ DCT->addDecl(DMD);
+ DMD->setAccess(clang::AS_none);
+ VD->setDeclContext(DMD);
+ VD->setLexicalDeclContext(DMD);
+ DMD->addDecl(VD);
+ DMD->setMapperVarRef(MapperVarRef);
+ FieldDecl *FD = *RD->field_begin();
+ // create mapper refence.
+ return DeclRefExpr::Create(Ctx, NestedNameSpecifierLoc{}, FD->getLocation(),
+ DMD, false, SourceLocation(), BaseType, VK_LValue);
+}
+
+static bool IsImplicitMapperNeeded(Sema &S, DSAStackTy *Stack,
+ QualType CanonType, const Expr *E) {
+
+ // DFS over data members in structures/classes.
+ SmallVector<std::pair<QualType, FieldDecl *>, 4> Types(1,
+ {CanonType, nullptr});
+ llvm::DenseMap<const Type *, Expr *> Visited;
+ SmallVector<std::pair<FieldDecl *, unsigned>, 4> ParentChain(1, {nullptr, 1});
+ while (!Types.empty()) {
+ QualType BaseType;
+ FieldDecl *CurFD;
+ std::tie(BaseType, CurFD) = Types.pop_back_val();
+ while (ParentChain.back().second == 0)
+ ParentChain.pop_back();
+ --ParentChain.back().second;
+ if (BaseType.isNull())
+ continue;
+ // Only structs/classes are allowed to have mappers.
+ const RecordDecl *RD = BaseType.getCanonicalType()->getAsRecordDecl();
+ if (!RD)
+ continue;
+ auto It = Visited.find(BaseType.getTypePtr());
+ if (It == Visited.end()) {
+ // Try to find the associated user-defined mapper.
+ CXXScopeSpec MapperIdScopeSpec;
+ DeclarationNameInfo DefaultMapperId;
+ DefaultMapperId.setName(S.Context.DeclarationNames.getIdentifier(
+ &S.Context.Idents.get("default")));
+ DefaultMapperId.setLoc(E->getExprLoc());
+ ExprResult ER = buildUserDefinedMapperRef(
+ S, Stack->getCurScope(), MapperIdScopeSpec, DefaultMapperId, BaseType,
+ /*UnresolvedMapper=*/nullptr);
+ if (ER.isInvalid())
+ continue;
+ It = Visited.try_emplace(BaseType.getTypePtr(), ER.get()).first;
+ }
+ // Found default mapper.
+ if (It->second)
+ return true;
+ // Check for the "default" mapper for data members.
+ bool FirstIter = true;
+ for (FieldDecl *FD : RD->fields()) {
+ if (!FD)
+ continue;
+ QualType FieldTy = FD->getType();
+ if (FieldTy.isNull() ||
+ !(FieldTy->isStructureOrClassType() || FieldTy->isUnionType()))
+ continue;
+ if (FirstIter) {
+ FirstIter = false;
+ ParentChain.emplace_back(CurFD, 1);
+ } else {
+ ++ParentChain.back().second;
+ }
+ Types.emplace_back(FieldTy, FD);
+ }
+ }
+ return false;
+}
+
// Check the validity of the provided variable list for the provided clause kind
// \a CKind. In the check process the valid expressions, mappable expression
// components, variables, and user-defined mappers are extracted and used to
@@ -21256,6 +21241,24 @@ static void checkMappableExpressionList(
Type.getCanonicalType(), UnresolvedMapper);
if (ER.isInvalid())
continue;
+ if (!ER.get() && isa<ArraySectionExpr>(VE)) {
+ // Create implicit mapper as needed.
+ QualType BaseType = VE->getType().getCanonicalType();
+ if (BaseType->isSpecificBuiltinType(BuiltinType::ArraySection)) {
+ const auto *OASE = cast<ArraySectionExpr>(VE->IgnoreParenImpCasts());
+ QualType BType = ArraySectionExpr::getBaseOriginalType(OASE->getBase());
+ QualType ElemType;
+ if (const auto *ATy = BType->getAsArrayTypeUnsafe())
+ ElemType = ATy->getElementType();
+ else
+ ElemType = BType->getPointeeType();
+ BaseType = ElemType.getCanonicalType();
+ }
+ if (BaseType->getAsRecordDecl() &&
+ IsImplicitMapperNeeded(SemaRef, DSAS, BaseType, VE)) {
+ ER = buildImplicitMapper(SemaRef, BaseType, DSAS);
+ }
+ }
MVLI.UDMapperList.push_back(ER.get());
// Save the current expression.
>From 40d801c37af51a51f3849d9ff65f3ee476c15805 Mon Sep 17 00:00:00 2001
From: Jennifer Yu <jennifer.yu at intel.com>
Date: Wed, 31 Jul 2024 13:20:00 -0700
Subject: [PATCH 5/6] Remove unused code.
---
clang/lib/Sema/SemaOpenMP.cpp | 2 --
1 file changed, 2 deletions(-)
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 74e6f730b775d..a302a5ef17caf 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -5763,8 +5763,6 @@ processImplicitMapsWithDefaultMappers(Sema &S, DSAStackTy *Stack,
if (S.getLangOpts().OpenMP < 50)
return;
SmallVector<OMPClause *, 4> ImplicitMaps;
- SmallVector<OMPClause *, 4> ClausesNeedImplicitMapper;
- llvm::DenseMap<const Expr *, QualType> ExprsNeedMapper;
for (int Cnt = 0, EndCnt = Clauses.size(); Cnt < EndCnt; ++Cnt) {
auto *C = dyn_cast<OMPMapClause>(Clauses[Cnt]);
if (!C)
>From ec9097ec9a06a77cae763d989c38cc7a54593e9b Mon Sep 17 00:00:00 2001
From: Jennifer Yu <jennifer.yu at intel.com>
Date: Thu, 1 Aug 2024 23:25:51 -0700
Subject: [PATCH 6/6] Thanks Alexey for the review. This is to add new
function hasUserDefinedMapper.
---
clang/lib/Sema/SemaOpenMP.cpp | 120 +++++++++++++-----
...arget_map_nest_defalut_mapper_ast_dump.cpp | 2 +-
2 files changed, 88 insertions(+), 34 deletions(-)
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index a302a5ef17caf..3f1ad3c89b661 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -20796,26 +20796,15 @@ struct MappableVarListInfo {
};
} // namespace
-static std::pair<DeclRefExpr *, VarDecl *>
-buildImplicitMap(Sema &S, QualType BaseType, DSAStackTy *Stack,
- SmallVectorImpl<OMPClause *> &Maps) {
+static DeclRefExpr *buildImplicitMap(Sema &S, QualType BaseType,
+ DSAStackTy *Stack,
+ SmallVectorImpl<OMPClause *> &Maps) {
const RecordDecl *RD = BaseType->getAsRecordDecl();
- // AST context is RD's ParentASTContext().
- ASTContext &Ctx = RD->getParentASTContext();
- // DeclContext is RD's DeclContext.
- DeclContext *DCT = const_cast<DeclContext *>(RD->getDeclContext());
SourceRange Range = RD->getSourceRange();
DeclarationNameInfo ImplicitName;
// Dummy variable _s for Mapper.
- ImplicitName.setName(
- Ctx.DeclarationNames.getIdentifier(&Ctx.Idents.get("_s")));
- DeclarationName VN = ImplicitName.getName();
- TypeSourceInfo *TInfo =
- Ctx.getTrivialTypeSourceInfo(BaseType, Range.getEnd());
- VarDecl *VD =
- VarDecl::Create(Ctx, DCT, Range.getEnd(), Range.getEnd(),
- VN.getAsIdentifierInfo(), BaseType, TInfo, SC_None);
+ VarDecl *VD = buildVarDecl(S, Range.getEnd(), BaseType, "_s");
DeclRefExpr *MapperVarRef =
buildDeclRefExpr(S, VD, BaseType, SourceLocation());
@@ -20835,13 +20824,13 @@ buildImplicitMap(Sema &S, QualType BaseType, DSAStackTy *Stack,
DeclarationNameInfo MapperId;
OpenMPDirectiveKind DKind = Stack->getCurrentDirective();
- OMPClause *MapClasue = S.OpenMP().ActOnOpenMPMapClause(
+ OMPClause *MapClause = S.OpenMP().ActOnOpenMPMapClause(
nullptr, OMPC_MAP_MODIFIER_unknown, SourceLocation(), MapperIdScopeSpec,
MapperId, DKind == OMPD_target_enter_data ? OMPC_MAP_to : OMPC_MAP_tofrom,
/*IsMapTypeImplicit=*/true, SourceLocation(), SourceLocation(), SExprs,
OMPVarListLocTy());
- Maps.push_back(MapClasue);
- return {MapperVarRef, VD};
+ Maps.push_back(MapClause);
+ return MapperVarRef;
}
static ExprResult buildImplicitMapper(Sema &S, QualType BaseType,
@@ -20849,9 +20838,7 @@ static ExprResult buildImplicitMapper(Sema &S, QualType BaseType,
// Build impilicit map for mapper
SmallVector<OMPClause *, 4> Maps;
- VarDecl *VD;
- DeclRefExpr *MapperVarRef;
- std::tie(MapperVarRef, VD) = buildImplicitMap(S, BaseType, Stack, Maps);
+ DeclRefExpr *MapperVarRef = buildImplicitMap(S, BaseType, Stack, Maps);
const RecordDecl *RD = BaseType->getAsRecordDecl();
// AST context is RD's ParentASTContext().
@@ -20870,6 +20857,7 @@ static ExprResult buildImplicitMapper(Sema &S, QualType BaseType,
S.PushOnScopeChains(DMD, Scope, /*AddToContext*/ false);
DCT->addDecl(DMD);
DMD->setAccess(clang::AS_none);
+ auto *VD = cast<DeclRefExpr>(MapperVarRef)->getDecl();
VD->setDeclContext(DMD);
VD->setLexicalDeclContext(DMD);
DMD->addDecl(VD);
@@ -20880,18 +20868,86 @@ static ExprResult buildImplicitMapper(Sema &S, QualType BaseType,
DMD, false, SourceLocation(), BaseType, VK_LValue);
}
-static bool IsImplicitMapperNeeded(Sema &S, DSAStackTy *Stack,
+// Look up the user-defined mapper given the mapper name and mapper type,
+// return true if found one.
+static bool hasUserDefinedMapper(Sema &SemaRef, Scope *S,
+ CXXScopeSpec &MapperIdScopeSpec,
+ const DeclarationNameInfo &MapperId,
+ QualType Type) {
+ // Find all user-defined mappers with the given MapperId.
+ SmallVector<UnresolvedSet<8>, 4> Lookups;
+ LookupResult Lookup(SemaRef, MapperId, Sema::LookupOMPMapperName);
+ Lookup.suppressDiagnostics();
+ if (S)
+ while (S && SemaRef.LookupParsedName(Lookup, S, &MapperIdScopeSpec,
+ /*ObjectType=*/QualType())) {
+ NamedDecl *D = Lookup.getRepresentativeDecl();
+ while (S && !S->isDeclScope(D))
+ S = S->getParent();
+ if (S)
+ S = S->getParent();
+ Lookups.emplace_back();
+ Lookups.back().append(Lookup.begin(), Lookup.end());
+ Lookup.clear();
+ }
+ if (SemaRef.CurContext->isDependentContext() || Type->isDependentType() ||
+ Type->isInstantiationDependentType() ||
+ Type->containsUnexpandedParameterPack() ||
+ filterLookupForUDReductionAndMapper<bool>(Lookups, [](ValueDecl *D) {
+ return !D->isInvalidDecl() &&
+ (D->getType()->isDependentType() ||
+ D->getType()->isInstantiationDependentType() ||
+ D->getType()->containsUnexpandedParameterPack());
+ }))
+ return false;
+ // Perform argument dependent lookup.
+ SourceLocation Loc = MapperId.getLoc();
+ if (SemaRef.getLangOpts().CPlusPlus && !MapperIdScopeSpec.isSet())
+ argumentDependentLookup(SemaRef, MapperId, Loc, Type, Lookups);
+ if (filterLookupForUDReductionAndMapper<ValueDecl *>(
+ Lookups, [&SemaRef, Type](ValueDecl *D) -> ValueDecl * {
+ if (!D->isInvalidDecl() &&
+ SemaRef.Context.hasSameType(D->getType(), Type))
+ return D;
+ return nullptr;
+ }))
+ return true;
+ // Find the first user-defined mapper with a type derived from the desired
+ // type.
+ if (auto *VD = filterLookupForUDReductionAndMapper<ValueDecl *>(
+ Lookups, [&SemaRef, Type, Loc](ValueDecl *D) -> ValueDecl * {
+ if (!D->isInvalidDecl() &&
+ SemaRef.IsDerivedFrom(Loc, Type, D->getType()) &&
+ !Type.isMoreQualifiedThan(D->getType()))
+ return D;
+ return nullptr;
+ })) {
+ CXXBasePaths Paths(/*FindAmbiguities=*/true, /*RecordPaths=*/true,
+ /*DetectVirtual=*/false);
+ if (SemaRef.IsDerivedFrom(Loc, Type, VD->getType(), Paths)) {
+ if (!Paths.isAmbiguous(SemaRef.Context.getCanonicalType(
+ VD->getType().getUnqualifiedType()))) {
+ if (SemaRef.CheckBaseClassAccess(
+ Loc, VD->getType(), Type, Paths.front(),
+ /*DiagID=*/0) != Sema::AR_inaccessible) {
+ return true;
+ }
+ }
+ }
+ }
+ return false;
+}
+
+static bool isImplicitMapperNeeded(Sema &S, DSAStackTy *Stack,
QualType CanonType, const Expr *E) {
// DFS over data members in structures/classes.
SmallVector<std::pair<QualType, FieldDecl *>, 4> Types(1,
{CanonType, nullptr});
- llvm::DenseMap<const Type *, Expr *> Visited;
+ llvm::DenseMap<const Type *, bool> Visited;
SmallVector<std::pair<FieldDecl *, unsigned>, 4> ParentChain(1, {nullptr, 1});
while (!Types.empty()) {
- QualType BaseType;
- FieldDecl *CurFD;
- std::tie(BaseType, CurFD) = Types.pop_back_val();
+ auto [BaseType, CurFD] = Types.pop_back_val();
while (ParentChain.back().second == 0)
ParentChain.pop_back();
--ParentChain.back().second;
@@ -20909,12 +20965,10 @@ static bool IsImplicitMapperNeeded(Sema &S, DSAStackTy *Stack,
DefaultMapperId.setName(S.Context.DeclarationNames.getIdentifier(
&S.Context.Idents.get("default")));
DefaultMapperId.setLoc(E->getExprLoc());
- ExprResult ER = buildUserDefinedMapperRef(
- S, Stack->getCurScope(), MapperIdScopeSpec, DefaultMapperId, BaseType,
- /*UnresolvedMapper=*/nullptr);
- if (ER.isInvalid())
- continue;
- It = Visited.try_emplace(BaseType.getTypePtr(), ER.get()).first;
+ bool HasUDMapper =
+ hasUserDefinedMapper(S, Stack->getCurScope(), MapperIdScopeSpec,
+ DefaultMapperId, BaseType);
+ It = Visited.try_emplace(BaseType.getTypePtr(), HasUDMapper).first;
}
// Found default mapper.
if (It->second)
@@ -21253,7 +21307,7 @@ static void checkMappableExpressionList(
BaseType = ElemType.getCanonicalType();
}
if (BaseType->getAsRecordDecl() &&
- IsImplicitMapperNeeded(SemaRef, DSAS, BaseType, VE)) {
+ isImplicitMapperNeeded(SemaRef, DSAS, BaseType, VE)) {
ER = buildImplicitMapper(SemaRef, BaseType, DSAS);
}
}
diff --git a/clang/test/OpenMP/target_map_nest_defalut_mapper_ast_dump.cpp b/clang/test/OpenMP/target_map_nest_defalut_mapper_ast_dump.cpp
index c088fc1d3c86f..b2fb8fbaa50b4 100644
--- a/clang/test/OpenMP/target_map_nest_defalut_mapper_ast_dump.cpp
+++ b/clang/test/OpenMP/target_map_nest_defalut_mapper_ast_dump.cpp
@@ -31,4 +31,4 @@ void foo() {
// DUM-NEXT: | | `-DeclRefExpr {{.*}}<<invalid sloc>> 'D' lvalue Var {{.*}} '_s' 'D'
// DUM-NEXT: | `-MemberExpr {{.*}}<line:11:3> 'int' lvalue .h {{.*}}
// DUM-NEXT: | `-DeclRefExpr {{.*}}<<invalid sloc>> 'D' lvalue Var {{.*}} '_s' 'D'
-// DUM-NEXT: `-VarDecl {{.*}} <line:12:1> col:1 used _s 'D'
+// DUM-NEXT: `-VarDecl {{.*}} <line:12:1> col:1 implicit used _s 'D'
More information about the cfe-commits
mailing list