[OpenMP] Generate implicit default mapper for mapping array section. #101101
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
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.
|
@llvm/pr-subscribers-offload Author: None (jyu2-git) ChangesThis is only for struct containing nested structs with custom mappers. Add three functions: In processImplicitMapsWithDefaultMappers, when nested user defined mapper is found, create a clause list (ClausesNeedImplicitMapper) to generat mapper for the corespoing map clause. Patch is 35.04 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/101101.diff 5 Files Affected:
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]], ...
[truncated]
|
|
@llvm/pr-subscribers-clang Author: None (jyu2-git) ChangesThis is only for struct containing nested structs with custom mappers. Add three functions: In processImplicitMapsWithDefaultMappers, when nested user defined mapper is found, create a clause list (ClausesNeedImplicitMapper) to generat mapper for the corespoing map clause. Patch is 35.04 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/101101.diff 5 Files Affected:
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]], ...
[truncated]
|
Rename the tests.
|
✅ With the latest revision this PR passed the C/C++ code formatter. |
This is generate implicit mapper when map is created instead after map created.
Thanks Alexey for the code review!! |
jyu2-git
added a commit
to jyu2-git/llvm-project
that referenced
this pull request
Aug 3, 2024
Add unspport. This is relate llvm#101101
Merged
jyu2-git
added a commit
that referenced
this pull request
Aug 3, 2024
Add unspport. This is relate #101101
abhinavgaba
added a commit
to abhinavgaba/llvm-project
that referenced
this pull request
Jun 3, 2025
…has a mapper. This builds upon llvm#101101, which used implicit compiler-generated mappers when mapping an array-section of structs with members that have user-defained default mappers. Now we do the same when mapping arrays of structs.
tomtor
pushed a commit
to tomtor/llvm-project
that referenced
this pull request
Jun 14, 2025
…ppers (llvm#142511) This builds upon llvm#101101 from @jyu2-git, which used compiler-generated mappers when mapping an array-section of structs with members that have user-defined default mappers. Now we do the same when mapping arrays of structs.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This is only for struct containing nested structs with user defined mappers.
Add four functions:
1>buildImplicitMap: build map for default mapper
2>buildImplicitMapper: build default mapper.
3>hasUserDefinedMapper for given mapper name and mapper type, lookup user defined map, if found one return true.
4>isImplicitMapperNeeded check if Mapper is needed
During create map, in checkMappableExpressionList, call isImplicitMapperNeeded when it return true, call buildImplicitMapper to generate implicit mapper and added to map clause.