[clang] [llvm] [Clang][OpenMP] Fix mapping of arrays of structs with members with mappers (PR #142511)
Abhinav Gaba via llvm-commits
llvm-commits at lists.llvm.org
Wed Jun 11 11:00:03 PDT 2025
https://github.com/abhinavgaba updated https://github.com/llvm/llvm-project/pull/142511
>From 9d29860f69dec99aa435ea23c2f2559394ca771e Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Wed, 28 May 2025 14:30:18 -0700
Subject: [PATCH 1/2] Create implicit mappers when mapping arrays of structs
when a member has a mapper.
This builds upon #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.
---
clang/docs/ReleaseNotes.rst | 3 +
clang/lib/Sema/SemaOpenMP.cpp | 38 ++-
...of_structs_with_nested_mapper_ast_dump.cpp | 34 ++
..._of_structs_with_nested_mapper_codegen.cpp | 323 ++++++++++++++++++
...f_structs_with_nested_mapper_ast_dump.cpp} | 0
...of_structs_with_nested_mapper_codegen.cpp} | 0
...re_mapper_nested_default_mappers_array.cpp | 6 +-
7 files changed, 388 insertions(+), 16 deletions(-)
create mode 100644 clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_ast_dump.cpp
create mode 100644 clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp
rename clang/test/OpenMP/{target_map_nest_defalut_mapper_ast_dump.cpp => target_map_array_section_of_structs_with_nested_mapper_ast_dump.cpp} (100%)
rename clang/test/OpenMP/{target_map_nest_defalut_mapper_codegen.cpp => target_map_array_section_of_structs_with_nested_mapper_codegen.cpp} (100%)
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 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
+//
diff --git a/clang/test/OpenMP/target_map_nest_defalut_mapper_ast_dump.cpp b/clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_ast_dump.cpp
similarity index 100%
rename from clang/test/OpenMP/target_map_nest_defalut_mapper_ast_dump.cpp
rename to clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_ast_dump.cpp
diff --git a/clang/test/OpenMP/target_map_nest_defalut_mapper_codegen.cpp b/clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_codegen.cpp
similarity index 100%
rename from clang/test/OpenMP/target_map_nest_defalut_mapper_codegen.cpp
rename to clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_codegen.cpp
diff --git a/offload/test/mapping/declare_mapper_nested_default_mappers_array.cpp b/offload/test/mapping/declare_mapper_nested_default_mappers_array.cpp
index d545e98ef6c3e..93695d1b388ff 100644
--- a/offload/test/mapping/declare_mapper_nested_default_mappers_array.cpp
+++ b/offload/test/mapping/declare_mapper_nested_default_mappers_array.cpp
@@ -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
}
>From c8984b49eaffae61be9d786b501a7f9b1567e7a3 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Tue, 3 Jun 2025 01:29:36 -0700
Subject: [PATCH 2/2] Fix a minor typo in a LIT test.
---
.../target_map_array_of_structs_with_nested_mapper_ast_dump.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
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
index 3e55793dc2596..a5847709d3e76 100644
--- 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
@@ -31,4 +31,4 @@ void foo() {
// DUM-NEXT: | | `-DeclRefExpr {{.*}}<<invalid sloc>> 'D' lvalue Var {{.*}} '_s' 'D'
// DUM-NEXT: | `-MemberExpr {{.*}}<line:11:3> 'int' lvalue .h {{.*}}
// DUM-NEXT: | `-DeclRefExpr {{.*}}<<invalid sloc>> 'D' lvalue Var {{.*}} '_s' 'D'
-// ]DUM-NEXT: `-VarDecl {{.*}} <line:12:1> col:1 implicit used _s 'D'
+// DUM-NEXT: `-VarDecl {{.*}} <line:12:1> col:1 implicit used _s 'D'
More information about the llvm-commits
mailing list