caio/clang-p2996
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[Clang][OpenMP] Fix mapping of arrays of structs with members with mappers (#142511)

Browse Source 
This builds upon #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.
This commit is contained in:
Abhinav Gaba
2025-06-11 12:03:55 -07:00
committed by GitHub
parent 0c62571d9f
commit 02b6849cf1
7 changed files with 388 additions and 16 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
clang/docs/ReleaseNotes.rst
View File

@@ -1102,6 +1102,9 @@ OpenMP Support
- An error is now emitted when OpenMP ``collapse`` and ``ordered`` clauses have
an argument larger than what can fit within a 64-bit integer.
- Added support for private variable reduction.
- Fixed mapping of arrays of structs containing nested structs with user defined
mappers, by using compiler-generated default mappers for the outer structs for
such maps.
Improvements
^^^^^^^^^^^^

38
clang/lib/Sema/SemaOpenMP.cpp
View File

@@ -22057,20 +22057,34 @@ 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();
// If no user-defined mapper is found, we need to create an implicit one for
// arrays/array-sections on structs that have members that have
// user-defined mappers. This is needed to ensure that the mapper for the
// member is invoked when mapping each element of the array/array-section.
if (!ER.get()) {
QualType BaseType;
if (isa<ArraySectionExpr>(VE)) {
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();
}
} else if (VE->getType()->isArrayType()) {
const ArrayType *AT = VE->getType()->getAsArrayTypeUnsafe();
const QualType ElemType = AT->getElementType();
BaseType = ElemType.getCanonicalType();
}
if (BaseType->getAsRecordDecl() &&
if (!BaseType.isNull() && BaseType->getAsRecordDecl() &&
isImplicitMapperNeeded(SemaRef, DSAS, BaseType, VE)) {
ER = buildImplicitMapper(SemaRef, BaseType, DSAS);
}

34
clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_ast_dump.cpp Normal file
View File

@@ -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)
{
sa[0].e = 333;
sa[1].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 implicit used _s 'D'

323
clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp Normal file
View File

@@ -0,0 +1,323 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --check-globals --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _ --global-value-regex "\.offload_.*"
// 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)
{
sa[1].e = 333;
sa[1].f.a = 444;
}
}
#endif
//.
// CHECK: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 120]
// CHECK: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 35]
//.
// 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 nuw [[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 nuw [[STRUCT_D]], ptr [[ARRAYIDX1]], i32 0, i32 1
// CHECK-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_C:%.*]], ptr [[F]], i32 0, i32 0
// CHECK-NEXT: store i32 222, ptr [[A]], align 4
// 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 [[SA]], 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 nuw [[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 nuw [[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 nuw [[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 nuw [[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 nuw [[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 nuw [[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 nuw [[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 nuw [[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 nuw [[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 nuw [[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 nuw [[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 nuw [[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 nuw [[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, !nonnull [[META5:![0-9]+]], !align [[META6:![0-9]+]]
// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x %struct.D], ptr [[TMP0]], i64 0, i64 1
// CHECK-NEXT: [[E:%.*]] = getelementptr inbounds nuw [[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 nuw [[STRUCT_D]], ptr [[ARRAYIDX1]], i32 0, i32 1
// CHECK-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[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: [[TMP6:%.*]] = udiv exact i64 [[TMP3]], 12
// CHECK-NEXT: [[TMP7:%.*]] = getelementptr [[STRUCT_D:%.*]], ptr [[TMP2]], i64 [[TMP6]]
// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP6]], 1
// CHECK-NEXT: [[TMP8:%.*]] = and i64 [[TMP4]], 8
// CHECK-NEXT: [[TMP9:%.*]] = icmp ne ptr [[TMP1]], [[TMP2]]
// CHECK-NEXT: [[TMP10:%.*]] = and i64 [[TMP4]], 16
// CHECK-NEXT: [[TMP11:%.*]] = icmp ne i64 [[TMP10]], 0
// CHECK-NEXT: [[TMP12:%.*]] = and i1 [[TMP9]], [[TMP11]]
// CHECK-NEXT: [[TMP13:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP12]]
// CHECK-NEXT: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP8]], 0
// CHECK-NEXT: [[TMP14:%.*]] = and i1 [[TMP13]], [[DOTOMP_ARRAY__INIT__DELETE]]
// CHECK-NEXT: br i1 [[TMP14]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]]
// CHECK: .omp.array..init:
// CHECK-NEXT: [[TMP15:%.*]] = mul nuw i64 [[TMP6]], 12
// CHECK-NEXT: [[TMP16:%.*]] = and i64 [[TMP4]], -4
// CHECK-NEXT: [[TMP17:%.*]] = or i64 [[TMP16]], 512
// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP15]], i64 [[TMP17]], ptr [[TMP5]])
// CHECK-NEXT: br label [[OMP_ARRAYMAP_HEAD]]
// CHECK: omp.arraymap.head:
// CHECK-NEXT: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP7]]
// 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 [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END20:%.*]] ]
// CHECK-NEXT: [[E:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
// CHECK-NEXT: [[F:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 1
// CHECK-NEXT: [[H:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 2
// CHECK-NEXT: [[TMP18:%.*]] = getelementptr i32, ptr [[H]], i32 1
// CHECK-NEXT: [[TMP19:%.*]] = ptrtoint ptr [[TMP18]] to i64
// CHECK-NEXT: [[TMP20:%.*]] = ptrtoint ptr [[E]] to i64
// CHECK-NEXT: [[TMP21:%.*]] = sub i64 [[TMP19]], [[TMP20]]
// CHECK-NEXT: [[TMP22:%.*]] = sdiv exact i64 [[TMP21]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
// CHECK-NEXT: [[TMP23:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]])
// CHECK-NEXT: [[TMP24:%.*]] = shl i64 [[TMP23]], 48
// CHECK-NEXT: [[TMP25:%.*]] = add nuw i64 0, [[TMP24]]
// CHECK-NEXT: [[TMP26:%.*]] = and i64 [[TMP4]], 3
// CHECK-NEXT: [[TMP27:%.*]] = icmp eq i64 [[TMP26]], 0
// CHECK-NEXT: br i1 [[TMP27]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]]
// CHECK: omp.type.alloc:
// CHECK-NEXT: [[TMP28:%.*]] = and i64 [[TMP25]], -4
// CHECK-NEXT: br label [[OMP_TYPE_END:%.*]]
// CHECK: omp.type.alloc.else:
// CHECK-NEXT: [[TMP29:%.*]] = icmp eq i64 [[TMP26]], 1
// CHECK-NEXT: br i1 [[TMP29]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]]
// CHECK: omp.type.to:
// CHECK-NEXT: [[TMP30:%.*]] = and i64 [[TMP25]], -3
// CHECK-NEXT: br label [[OMP_TYPE_END]]
// CHECK: omp.type.to.else:
// CHECK-NEXT: [[TMP31:%.*]] = icmp eq i64 [[TMP26]], 2
// CHECK-NEXT: br i1 [[TMP31]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]]
// CHECK: omp.type.from:
// CHECK-NEXT: [[TMP32:%.*]] = and i64 [[TMP25]], -2
// CHECK-NEXT: br label [[OMP_TYPE_END]]
// CHECK: omp.type.end:
// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP28]], [[OMP_TYPE_ALLOC]] ], [ [[TMP30]], [[OMP_TYPE_TO]] ], [ [[TMP32]], [[OMP_TYPE_FROM]] ], [ [[TMP25]], [[OMP_TYPE_TO_ELSE]] ]
// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP22]], i64 [[OMP_MAPTYPE]], ptr null)
// CHECK-NEXT: [[TMP33:%.*]] = add nuw i64 281474976711171, [[TMP24]]
// CHECK-NEXT: [[TMP34:%.*]] = and i64 [[TMP4]], 3
// CHECK-NEXT: [[TMP35:%.*]] = icmp eq i64 [[TMP34]], 0
// CHECK-NEXT: br i1 [[TMP35]], label [[OMP_TYPE_ALLOC1:%.*]], label [[OMP_TYPE_ALLOC_ELSE2:%.*]]
// CHECK: omp.type.alloc1:
// CHECK-NEXT: [[TMP36:%.*]] = and i64 [[TMP33]], -4
// CHECK-NEXT: br label [[OMP_TYPE_END6:%.*]]
// CHECK: omp.type.alloc.else2:
// CHECK-NEXT: [[TMP37:%.*]] = icmp eq i64 [[TMP34]], 1
// CHECK-NEXT: br i1 [[TMP37]], label [[OMP_TYPE_TO3:%.*]], label [[OMP_TYPE_TO_ELSE4:%.*]]
// CHECK: omp.type.to3:
// CHECK-NEXT: [[TMP38:%.*]] = and i64 [[TMP33]], -3
// CHECK-NEXT: br label [[OMP_TYPE_END6]]
// CHECK: omp.type.to.else4:
// CHECK-NEXT: [[TMP39:%.*]] = icmp eq i64 [[TMP34]], 2
// CHECK-NEXT: br i1 [[TMP39]], label [[OMP_TYPE_FROM5:%.*]], label [[OMP_TYPE_END6]]
// CHECK: omp.type.from5:
// CHECK-NEXT: [[TMP40:%.*]] = and i64 [[TMP33]], -2
// CHECK-NEXT: br label [[OMP_TYPE_END6]]
// CHECK: omp.type.end6:
// CHECK-NEXT: [[OMP_MAPTYPE7:%.*]] = phi i64 [ [[TMP36]], [[OMP_TYPE_ALLOC1]] ], [ [[TMP38]], [[OMP_TYPE_TO3]] ], [ [[TMP40]], [[OMP_TYPE_FROM5]] ], [ [[TMP33]], [[OMP_TYPE_TO_ELSE4]] ]
// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 4, i64 [[OMP_MAPTYPE7]], ptr null)
// CHECK-NEXT: [[TMP41:%.*]] = add nuw i64 281474976711171, [[TMP24]]
// CHECK-NEXT: [[TMP42:%.*]] = and i64 [[TMP4]], 3
// CHECK-NEXT: [[TMP43:%.*]] = icmp eq i64 [[TMP42]], 0
// CHECK-NEXT: br i1 [[TMP43]], label [[OMP_TYPE_ALLOC8:%.*]], label [[OMP_TYPE_ALLOC_ELSE9:%.*]]
// CHECK: omp.type.alloc8:
// CHECK-NEXT: [[TMP44:%.*]] = and i64 [[TMP41]], -4
// CHECK-NEXT: br label [[OMP_TYPE_END13:%.*]]
// CHECK: omp.type.alloc.else9:
// CHECK-NEXT: [[TMP45:%.*]] = icmp eq i64 [[TMP42]], 1
// CHECK-NEXT: br i1 [[TMP45]], label [[OMP_TYPE_TO10:%.*]], label [[OMP_TYPE_TO_ELSE11:%.*]]
// CHECK: omp.type.to10:
// CHECK-NEXT: [[TMP46:%.*]] = and i64 [[TMP41]], -3
// CHECK-NEXT: br label [[OMP_TYPE_END13]]
// CHECK: omp.type.to.else11:
// CHECK-NEXT: [[TMP47:%.*]] = icmp eq i64 [[TMP42]], 2
// CHECK-NEXT: br i1 [[TMP47]], label [[OMP_TYPE_FROM12:%.*]], label [[OMP_TYPE_END13]]
// CHECK: omp.type.from12:
// CHECK-NEXT: [[TMP48:%.*]] = and i64 [[TMP41]], -2
// CHECK-NEXT: br label [[OMP_TYPE_END13]]
// CHECK: omp.type.end13:
// CHECK-NEXT: [[OMP_MAPTYPE14:%.*]] = phi i64 [ [[TMP44]], [[OMP_TYPE_ALLOC8]] ], [ [[TMP46]], [[OMP_TYPE_TO10]] ], [ [[TMP48]], [[OMP_TYPE_FROM12]] ], [ [[TMP41]], [[OMP_TYPE_TO_ELSE11]] ]
// CHECK-NEXT: call void @.omp_mapper._ZTS1C.default(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[F]], i64 4, i64 [[OMP_MAPTYPE14]], ptr null) #[[ATTR3]]
// CHECK-NEXT: [[TMP49:%.*]] = add nuw i64 281474976711171, [[TMP24]]
// CHECK-NEXT: [[TMP50:%.*]] = and i64 [[TMP4]], 3
// CHECK-NEXT: [[TMP51:%.*]] = icmp eq i64 [[TMP50]], 0
// CHECK-NEXT: br i1 [[TMP51]], label [[OMP_TYPE_ALLOC15:%.*]], label [[OMP_TYPE_ALLOC_ELSE16:%.*]]
// CHECK: omp.type.alloc15:
// CHECK-NEXT: [[TMP52:%.*]] = and i64 [[TMP49]], -4
// CHECK-NEXT: br label [[OMP_TYPE_END20]]
// CHECK: omp.type.alloc.else16:
// CHECK-NEXT: [[TMP53:%.*]] = icmp eq i64 [[TMP50]], 1
// CHECK-NEXT: br i1 [[TMP53]], label [[OMP_TYPE_TO17:%.*]], label [[OMP_TYPE_TO_ELSE18:%.*]]
// CHECK: omp.type.to17:
// CHECK-NEXT: [[TMP54:%.*]] = and i64 [[TMP49]], -3
// CHECK-NEXT: br label [[OMP_TYPE_END20]]
// CHECK: omp.type.to.else18:
// CHECK-NEXT: [[TMP55:%.*]] = icmp eq i64 [[TMP50]], 2
// CHECK-NEXT: br i1 [[TMP55]], label [[OMP_TYPE_FROM19:%.*]], label [[OMP_TYPE_END20]]
// CHECK: omp.type.from19:
// CHECK-NEXT: [[TMP56:%.*]] = and i64 [[TMP49]], -2
// CHECK-NEXT: br label [[OMP_TYPE_END20]]
// CHECK: omp.type.end20:
// CHECK-NEXT: [[OMP_MAPTYPE21:%.*]] = phi i64 [ [[TMP52]], [[OMP_TYPE_ALLOC15]] ], [ [[TMP54]], [[OMP_TYPE_TO17]] ], [ [[TMP56]], [[OMP_TYPE_FROM19]] ], [ [[TMP49]], [[OMP_TYPE_TO_ELSE18]] ]
// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[H]], i64 4, i64 [[OMP_MAPTYPE21]], 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]], [[TMP7]]
// CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]]
// CHECK: omp.arraymap.exit:
// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY22:%.*]] = icmp sgt i64 [[TMP6]], 1
// CHECK-NEXT: [[TMP57:%.*]] = and i64 [[TMP4]], 8
// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP57]], 0
// CHECK-NEXT: [[TMP58:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY22]], [[DOTOMP_ARRAY__DEL__DELETE]]
// CHECK-NEXT: br i1 [[TMP58]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
// CHECK: .omp.array..del:
// CHECK-NEXT: [[TMP59:%.*]] = mul nuw i64 [[TMP6]], 12
// CHECK-NEXT: [[TMP60:%.*]] = and i64 [[TMP4]], -4
// CHECK-NEXT: [[TMP61:%.*]] = or i64 [[TMP60]], 512
// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP59]], i64 [[TMP61]], ptr [[TMP5]])
// 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: [[TMP6:%.*]] = udiv exact i64 [[TMP3]], 4
// CHECK-NEXT: [[TMP7:%.*]] = getelementptr [[STRUCT_C:%.*]], ptr [[TMP2]], i64 [[TMP6]]
// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP6]], 1
// CHECK-NEXT: [[TMP8:%.*]] = and i64 [[TMP4]], 8
// CHECK-NEXT: [[TMP9:%.*]] = icmp ne ptr [[TMP1]], [[TMP2]]
// CHECK-NEXT: [[TMP10:%.*]] = and i64 [[TMP4]], 16
// CHECK-NEXT: [[TMP11:%.*]] = icmp ne i64 [[TMP10]], 0
// CHECK-NEXT: [[TMP12:%.*]] = and i1 [[TMP9]], [[TMP11]]
// CHECK-NEXT: [[TMP13:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP12]]
// CHECK-NEXT: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP8]], 0
// CHECK-NEXT: [[TMP14:%.*]] = and i1 [[TMP13]], [[DOTOMP_ARRAY__INIT__DELETE]]
// CHECK-NEXT: br i1 [[TMP14]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]]
// CHECK: .omp.array..init:
// CHECK-NEXT: [[TMP15:%.*]] = mul nuw i64 [[TMP6]], 4
// CHECK-NEXT: [[TMP16:%.*]] = and i64 [[TMP4]], -4
// CHECK-NEXT: [[TMP17:%.*]] = or i64 [[TMP16]], 512
// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP15]], i64 [[TMP17]], ptr [[TMP5]])
// CHECK-NEXT: br label [[OMP_ARRAYMAP_HEAD]]
// CHECK: omp.arraymap.head:
// CHECK-NEXT: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP7]]
// 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 [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END:%.*]] ]
// CHECK-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_C]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
// CHECK-NEXT: [[TMP18:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]])
// CHECK-NEXT: [[TMP19:%.*]] = shl i64 [[TMP18]], 48
// CHECK-NEXT: [[TMP20:%.*]] = add nuw i64 1, [[TMP19]]
// CHECK-NEXT: [[TMP21:%.*]] = and i64 [[TMP4]], 3
// CHECK-NEXT: [[TMP22:%.*]] = icmp eq i64 [[TMP21]], 0
// CHECK-NEXT: br i1 [[TMP22]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]]
// CHECK: omp.type.alloc:
// CHECK-NEXT: [[TMP23:%.*]] = and i64 [[TMP20]], -4
// CHECK-NEXT: br label [[OMP_TYPE_END]]
// CHECK: omp.type.alloc.else:
// CHECK-NEXT: [[TMP24:%.*]] = icmp eq i64 [[TMP21]], 1
// CHECK-NEXT: br i1 [[TMP24]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]]
// CHECK: omp.type.to:
// CHECK-NEXT: [[TMP25:%.*]] = and i64 [[TMP20]], -3
// CHECK-NEXT: br label [[OMP_TYPE_END]]
// CHECK: omp.type.to.else:
// CHECK-NEXT: [[TMP26:%.*]] = icmp eq i64 [[TMP21]], 2
// CHECK-NEXT: br i1 [[TMP26]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]]
// CHECK: omp.type.from:
// CHECK-NEXT: [[TMP27:%.*]] = and i64 [[TMP20]], -2
// CHECK-NEXT: br label [[OMP_TYPE_END]]
// CHECK: omp.type.end:
// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP23]], [[OMP_TYPE_ALLOC]] ], [ [[TMP25]], [[OMP_TYPE_TO]] ], [ [[TMP27]], [[OMP_TYPE_FROM]] ], [ [[TMP20]], [[OMP_TYPE_TO_ELSE]] ]
// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], 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]], [[TMP7]]
// CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]]
// CHECK: omp.arraymap.exit:
// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY1:%.*]] = icmp sgt i64 [[TMP6]], 1
// CHECK-NEXT: [[TMP28:%.*]] = and i64 [[TMP4]], 8
// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP28]], 0
// CHECK-NEXT: [[TMP29:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY1]], [[DOTOMP_ARRAY__DEL__DELETE]]
// CHECK-NEXT: br i1 [[TMP29]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
// CHECK: .omp.array..del:
// CHECK-NEXT: [[TMP30:%.*]] = mul nuw i64 [[TMP6]], 4
// CHECK-NEXT: [[TMP31:%.*]] = and i64 [[TMP4]], -4
// CHECK-NEXT: [[TMP32:%.*]] = or i64 [[TMP31]], 512
// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP30]], i64 [[TMP32]], ptr [[TMP5]])
// CHECK-NEXT: br label [[OMP_DONE]]
// CHECK: omp.done:
// CHECK-NEXT: ret void
//

0
clang/test/OpenMP/target_map_nest_defalut_mapper_ast_dump.cpp → clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_ast_dump.cpp
View File

0
clang/test/OpenMP/target_map_nest_defalut_mapper_codegen.cpp → clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_codegen.cpp
View File

6
offload/test/mapping/declare_mapper_nested_default_mappers_array.cpp
View File

@@ -4,8 +4,6 @@
// RUN: %libomptarget-compilexx-run-and-check-x86_64-unknown-linux-gnu
// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
// UNSUPPORTED: clang
#include <cstdio>
#include <cstdlib>
@@ -50,7 +48,7 @@ int main() {
sa[1].h = N;
printf("%d %d %d %4.5f %d\n", sa[1].e, sa[1].f.a, sa[1].f.c.a, sa[1].f.b[1],
sa[1].f.b == &x[0] ? 1 : 0);
sa[1].f.b == &y[0] ? 1 : 0);
// CHECK: 111 222 777 20.00000 1
__intptr_t p = reinterpret_cast<__intptr_t>(&y[0]);
@@ -65,6 +63,6 @@ int main() {
sa[1].f.b[1] = 40;
}
printf("%d %d %d %4.5f %d\n", sa[1].e, sa[1].f.a, sa[1].f.c.a, sa[1].f.b[1],
sa[1].f.b == &x[0] ? 1 : 0);
sa[1].f.b == &y[0] ? 1 : 0);
// CHECK: 333 222 777 40.00000 1
}