[Clang][OpenMP] Fix mapping of arrays of structs with members with mappers #142511
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
…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.
abhinavgaba
commented
Jun 3, 2025
| // 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 |
There was a problem hiding this comment.
The only difference vs the existing test for array-sections is in this line, where we use &sa, instead of the GEP for &sa[0].
| sa[1].e = 111; | ||
| sa[1].f.a = 222; | ||
|
|
||
| #pragma omp target map(tofrom : sa) |
There was a problem hiding this comment.
The test is identical to the array-section version, with the only difference being this line:
| sa[1].e = 111; | ||
| sa[1].f.a = 222; | ||
|
|
||
| #pragma omp target map(tofrom : sa) |
There was a problem hiding this comment.
The only source difference vs the array-section version of the test is in this line:
| @@ -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); | |||
There was a problem hiding this comment.
Using x instead of y in these comparisons was a bug in the test. The fixed test passes with the compiler code change.
abhinavgaba
changed the title
llvmbot
added
clang
|
@llvm/pr-subscribers-offload Author: Abhinav Gaba (abhinavgaba) ChangesThis 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. Patch is 28.13 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/142511.diff 7 Files Affected:
diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst
index 91b89a0946555..61a6d07415181 100644
--- a/clang/docs/ReleaseNotes.rst
+++ b/clang/docs/ReleaseNotes.rst
@@ -1065,6 +1065,9 @@ OpenMP Support
open parenthesis. (#GH139665)
- An error is now emitted when OpenMP ``collapse`` and ``ordered`` clauses have
an argument larger than what can fit within a 64-bit integer.
+- 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
^^^^^^^^^^^^
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index dd185f2ff254b..a8ef521b7eefc 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -22060,20 +22060,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);
}
diff --git a/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_ast_dump.cpp b/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_ast_dump.cpp
new file mode 100644
index 0000000000000..3e55793dc2596
--- /dev/null
+++ b/clang/test/OpenMP/target_map_array_of_structs_with_nested_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)
+ {
+ 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'
diff --git a/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp b/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp
new file mode 100644
index 0000000000000..5df1e958ad55a
--- /dev/null
+++ b/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp
@@ -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 i...
[truncated]
|
|
@llvm/pr-subscribers-clang Author: Abhinav Gaba (abhinavgaba) ChangesThis 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. Patch is 28.13 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/142511.diff 7 Files Affected:
diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst
index 91b89a0946555..61a6d07415181 100644
--- a/clang/docs/ReleaseNotes.rst
+++ b/clang/docs/ReleaseNotes.rst
@@ -1065,6 +1065,9 @@ OpenMP Support
open parenthesis. (#GH139665)
- An error is now emitted when OpenMP ``collapse`` and ``ordered`` clauses have
an argument larger than what can fit within a 64-bit integer.
+- 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
^^^^^^^^^^^^
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index dd185f2ff254b..a8ef521b7eefc 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -22060,20 +22060,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);
}
diff --git a/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_ast_dump.cpp b/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_ast_dump.cpp
new file mode 100644
index 0000000000000..3e55793dc2596
--- /dev/null
+++ b/clang/test/OpenMP/target_map_array_of_structs_with_nested_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)
+ {
+ 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'
diff --git a/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp b/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp
new file mode 100644
index 0000000000000..5df1e958ad55a
--- /dev/null
+++ b/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp
@@ -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 i...
[truncated]
|
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 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.