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...@intel.com> Date: Mon, 29 Jul 2024 14:18:48 -0700 Subject: [PATCH 1/7] 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...@intel.com> Date: Tue, 30 Jul 2024 10:11:20 -0700 Subject: [PATCH 2/7] 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...@intel.com> Date: Tue, 30 Jul 2024 10:51:08 -0700 Subject: [PATCH 3/7] 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...@intel.com> Date: Wed, 31 Jul 2024 12:51:33 -0700 Subject: [PATCH 4/7] 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...@intel.com> Date: Wed, 31 Jul 2024 13:20:00 -0700 Subject: [PATCH 5/7] 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...@intel.com> Date: Thu, 1 Aug 2024 23:25:51 -0700 Subject: [PATCH 6/7] 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' >From 5a855e3ef090663da2bafa23887e64486bb10d4d Mon Sep 17 00:00:00 2001 From: Jennifer Yu <jennifer...@intel.com> Date: Fri, 2 Aug 2024 11:49:59 -0700 Subject: [PATCH 7/7] Thanks Alexey for the review. This address his' comments. --- clang/docs/OpenMPSupport.rst | 2 ++ clang/docs/ReleaseNotes.rst | 4 +++ clang/lib/Sema/SemaOpenMP.cpp | 51 ++++++++++++++++++----------------- 3 files changed, 32 insertions(+), 25 deletions(-) diff --git a/clang/docs/OpenMPSupport.rst b/clang/docs/OpenMPSupport.rst index 5e63b2c0f0be6..375091e17af49 100644 --- a/clang/docs/OpenMPSupport.rst +++ b/clang/docs/OpenMPSupport.rst @@ -181,6 +181,8 @@ implementation. +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+ | device | user-defined mappers | :good:`done` | D56326,D58638,D58523,D58074,D60972,D59474 | +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+ +| device | map array-section with implicit mapper | :good:`done` | https://github.com/llvm/llvm-project/pull/101101 | ++------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+ | device | mapping lambda expression | :good:`done` | D51107 | +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+ | device | clause: use_device_addr for target data | :good:`done` | | diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst index 25a9f2836ba3d..7320b7384af87 100644 --- a/clang/docs/ReleaseNotes.rst +++ b/clang/docs/ReleaseNotes.rst @@ -295,6 +295,10 @@ Python Binding Changes OpenMP Support -------------- +Improvements +^^^^^^^^^^^^ +- Improve the handling of mapping array-section for struct containing nested structs with user defined mappers + Additional Information ====================== diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index 3f1ad3c89b661..fbf40f249e36a 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -20878,18 +20878,17 @@ static bool hasUserDefinedMapper(Sema &SemaRef, Scope *S, 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(); - } + 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() || @@ -20914,25 +20913,27 @@ static bool hasUserDefinedMapper(Sema &SemaRef, Scope *S, return true; // Find the first user-defined mapper with a type derived from the desired // type. - if (auto *VD = filterLookupForUDReductionAndMapper<ValueDecl *>( + 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; - } - } + }); + if (!VD) + return false; + CXXBasePaths Paths(/*FindAmbiguities=*/true, /*RecordPaths=*/true, + /*DetectVirtual=*/false); + if (SemaRef.IsDerivedFrom(Loc, Type, VD->getType(), Paths)) { + bool IsAmbiguous = !Paths.isAmbiguous(SemaRef.Context.getCanonicalType( + VD->getType().getUnqualifiedType())); + if (IsAmbiguous) + return false; + if (SemaRef.CheckBaseClassAccess( + Loc, VD->getType(), Type, Paths.front(), + /*DiagID=*/0) != Sema::AR_inaccessible) { + return true; } } return false; _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits