https://github.com/abhinavgaba updated https://github.com/llvm/llvm-project/pull/177059
>From c59e3ccc4b3c3cedd8ba21b67117bcd2e5c9038d Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Tue, 20 Jan 2026 15:28:09 -0800 Subject: [PATCH 1/4] [OpenMP][Mappers] Fix ref-count tracking for maps inserted by mappers. This is a fix for https://github.com/llvm/llvm-project/issues/61636. Ravi had this implemented downstream before he retired. The test is taken from @jdoerfert's WIP change in https://github.com/jdoerfert/llvm-project/commit/527bf4b1293f452aac1f872c4f771c7950b911f9. The change partially undoes the changes done in 0caf736d7e1d16d1059553fc28dbac31f0b9f788, so @alexey-bataev might need to take a look. Co-authored-by: Ravi Narayanaswamy <[email protected]> Co-authored-by: Johannes Doerfert <[email protected]> --- clang/lib/CodeGen/CGOpenMPRuntime.cpp | 6 +- llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 8 - offload/libomptarget/omptarget.cpp | 8 +- .../mapping/declare_mapper_target_checks.cpp | 164 ++++++++++++++++++ 4 files changed, 169 insertions(+), 17 deletions(-) create mode 100644 offload/test/mapping/declare_mapper_target_checks.cpp diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 01661ad54ee2f..1912374d59d78 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -10461,10 +10461,8 @@ getNestedDistributeDirective(ASTContext &Ctx, const OMPExecutableDirective &D) { /// void *base, void *begin, /// int64_t size, int64_t type, /// void *name = nullptr) { -/// // Allocate space for an array section first or add a base/begin for -/// // pointer dereference. -/// if ((size > 1 || (base != begin && maptype.IsPtrAndObj)) && -/// !maptype.IsDelete) +/// // Allocate space for an array section first. +/// if ((size > 1 || (base != begin)) && !maptype.IsDelete) /// __tgt_push_mapper_component(rt_mapper_handle, base, begin, /// size*sizeof(Ty), clearToFromMember(type)); /// // Map members. diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index 8d7a207a91f5a..0d9bc5c763d81 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -9466,14 +9466,6 @@ void OpenMPIRBuilder::emitUDMapperArrayInitOrDel( if (IsInit) { // base != begin? Value *BaseIsBegin = Builder.CreateICmpNE(Base, Begin); - // IsPtrAndObj? - Value *PtrAndObjBit = Builder.CreateAnd( - MapType, - Builder.getInt64( - static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>( - OpenMPOffloadMappingFlags::OMP_MAP_PTR_AND_OBJ))); - PtrAndObjBit = Builder.CreateIsNotNull(PtrAndObjBit); - BaseIsBegin = Builder.CreateAnd(BaseIsBegin, PtrAndObjBit); Cond = Builder.CreateOr(IsArray, BaseIsBegin); DeleteCond = Builder.CreateIsNull( DeleteBit, diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp index 960c5bc17df96..1821ec9f6954d 100644 --- a/offload/libomptarget/omptarget.cpp +++ b/offload/libomptarget/omptarget.cpp @@ -600,8 +600,7 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum, // then no argument is marked as TARGET_PARAM ("omp target data map" is not // associated with a target region, so there are no target parameters). This // may be considered a hack, we could revise the scheme in the future. - bool UpdateRef = - !(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) && !(FromMapper && I == 0); + bool UpdateRef = !(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF); MappingInfoTy::HDTTMapAccessorTy HDTTMap = Device.getMappingInfo().HostDataToTargetMap.getExclusiveAccessor(); @@ -1110,9 +1109,8 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum, void *HstPtrBegin = Args[I]; int64_t DataSize = ArgSizes[I]; bool IsImplicit = ArgTypes[I] & OMP_TGT_MAPTYPE_IMPLICIT; - bool UpdateRef = (!(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) || - (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) && - !(FromMapper && I == 0); + bool UpdateRef = !(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) || + (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ); bool ForceDelete = ArgTypes[I] & OMP_TGT_MAPTYPE_DELETE; bool HasPresentModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_PRESENT; bool HasHoldModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_OMPX_HOLD; diff --git a/offload/test/mapping/declare_mapper_target_checks.cpp b/offload/test/mapping/declare_mapper_target_checks.cpp new file mode 100644 index 0000000000000..3136138f9b030 --- /dev/null +++ b/offload/test/mapping/declare_mapper_target_checks.cpp @@ -0,0 +1,164 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic +#include <omp.h> +#include <stdio.h> + +#define TRUE 1 +#define FALSE 0 + +struct TY1 { + int i1, i2, i3; + static constexpr auto name = "TY1"; +}; +struct TY2 { + int i1, i2, i3; + static constexpr auto name = "TY2"; +}; + +// TY1 is not mapped, TY2 is +#pragma omp declare mapper(TY2 t) map(to : t.i1) map(from : t.i3) + +struct TY3 { + TY2 n; + static constexpr auto name = "TY3"; +}; +struct TY4 { + int a; + TY2 n; + int b; + static constexpr auto name = "TY4"; +}; + +static TY2 q2; +struct TY5 { + TY2 &n = q2; + static constexpr auto name = "TY5"; +}; + +template <typename T> int testType() { + T t1[2], t2[3], t3[4]; + for (int i = 0; i < 2; i++) + t1[i].i1 = t3[i].i1 = 1; + +#pragma omp target map(tofrom : t1, t2, t3) + for (int i = 0; i < 2; i++) { + t1[i].i3 = t3[i].i3 = t1[i].i1; + t1[i].i1 = t3[i].i1 = 7; + } + + for (int i = 0; i < 2; i++) { + if (t1[i].i3 != 1) { + printf("failed %s. t1[%d].i3 (%d) != t1[%d].i1 (%d)\n", T::name, i, + t1[i].i3, i, t1[i].i1); + return 1; + } + if (t3[i].i3 != 1) { + printf("failed %s. t3[%d].i3 (%d) != t3[%d].i1 (%d)\n", T::name, i, + t3[i].i3, i, t3[i].i1); + return 1; + } + } + + int pt0 = omp_target_is_present(&t1[0], omp_get_default_device()); + int pt1 = omp_target_is_present(&t2[1], omp_get_default_device()); + int pt2 = omp_target_is_present(&t3[2], omp_get_default_device()); + + // CHECK: present check for TY1: t1 0, t2 0, t3 0, expected 3x 0 + // CHECK: present check for TY2: t1 0, t2 0, t3 0, expected 3x 0 + // CHECK: present check for TY3: t1 0, t2 0, t3 0, expected 3x 0 + // CHECK: present check for TY4: t1 0, t2 0, t3 0, expected 3x 0 + printf("present check for %s: t1 %i, t2 %i, t3 %i, expected 3x 0\n", T::name, + pt0, pt1, pt2); + return pt0 + pt1 + pt2; +} + +template <typename T> int testTypeNestedPtr(T t1[2], T t2[3], T t3[4]) { + for (int i = 0; i < 2; i++) + t1[i].n.i1 = t3[i].n.i1 = 1; + +#pragma omp target map(tofrom : t1[0:2], t2[0:3], t3[0:4]) + for (int i = 0; i < 2; i++) { + t1[i].n.i3 = t3[i].n.i3 = t1[i].n.i1; + t1[i].n.i1 = t3[i].n.i1 = 7; + } + + for (int i = 0; i < 2; i++) { + if (t1[i].n.i3 != t1[i].n.i1) { + printf("failed %s-ptr. t1[%d].i3 (%d) != t1[%d].i1 (%d)\n", T::name, i, + t1[i].n.i3, i, t1[i].n.i1); + return 1; + } + if (t3[i].n.i3 != t3[i].n.i1) { + printf("failed %s-ptr. t3[%d].i3 (%d) != t3[%d].i1 (%d)\n", T::name, i, + t3[i].n.i3, i, t3[i].n.i1); + return 1; + } + } + + int pt0 = omp_target_is_present(&t1[0], omp_get_default_device()); + int pt1 = omp_target_is_present(&t2[1], omp_get_default_device()); + int pt2 = omp_target_is_present(&t3[2], omp_get_default_device()); + + // CHECK: present check for TY3-ptr: t1 0, t2 0, t3 0, expected 3x 0 + // CHECK: present check for TY4-ptr: t1 0, t2 0, t3 0, expected 3x 0 + // CHECK: present check for TY5-ptr: t1 0, t2 0, t3 0, expected 3x 0 + printf("present check for %s-ptr: t1 %i, t2 %i, t3 %i, expected 3x 0\n", + T::name, pt0, pt1, pt2); + return pt0 + pt1 + pt2; +} + +template <typename T> int testTypeNested() { + T t1[2], t2[3], t3[4]; + testTypeNestedPtr(t1, t2, t3); + for (int i = 0; i < 2; i++) + t1[i].n.i1 = t3[i].n.i1 = 1; + +#pragma omp target map(tofrom : t1, t2, t3) + for (int i = 0; i < 2; i++) { + t1[i].n.i3 = t3[i].n.i3 = t1[i].n.i1; + t1[i].n.i1 = t3[i].n.i1 = 7; + } + + for (int i = 0; i < 2; i++) { + if (t1[i].n.i3 != t1[i].n.i1) { + printf("failed %s. t1[%d].i3 (%d) != t1[%d].i1 (%d)\n", T::name, i, + t1[i].n.i3, i, t1[i].n.i1); + return 1; + } + if (t3[i].n.i3 != t3[i].n.i1) { + printf("failed %s. t3[%d].i3 (%d) != t3[%d].i1 (%d)\n", T::name, i, + t3[i].n.i3, i, t3[i].n.i1); + return 1; + } + } + + int pt0 = omp_target_is_present(&t1[0], omp_get_default_device()); + int pt1 = omp_target_is_present(&t2[1], omp_get_default_device()); + int pt2 = omp_target_is_present(&t3[2], omp_get_default_device()); + + // CHECK: present check for TY1: t1 0, t2 0, t3 0, expected 3x 0 + // CHECK: present check for TY2: t1 0, t2 0, t3 0, expected 3x 0 + printf("present check for %s: t1 %i, t2 %i, t3 %i, expected 3x 0\n", T::name, + pt0, pt1, pt2); + return pt0 + pt1 + pt2; +} + +int main(int argc, char **argv) { + int r = 0; + r += testType<TY1>(); + r += testType<TY2>(); + r += testTypeNested<TY3>(); + r += testTypeNested<TY4>(); + { + TY2 a[9]; + TY5 t1[2], t2[3], t3[4]; + int i = 0; + for (int j = 0; j < 2; j++) + t1[j].n = a[i++]; + for (int j = 0; j < 3; j++) + t2[j].n = a[i++]; + for (int j = 0; j < 4; j++) + t3[j].n = a[i++]; + r += testTypeNestedPtr<TY5>(t1, t2, t3); + } + return r; +} >From 1315cbe6d4fb9ca756b9dc3d282c59038d82e6c0 Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Tue, 20 Jan 2026 15:57:53 -0800 Subject: [PATCH 2/4] Update clang codegen tests. --- clang/test/OpenMP/declare_mapper_codegen.cpp | 25 +- ..._of_structs_with_nested_mapper_codegen.cpp | 220 +++++++------- ..._of_structs_with_nested_mapper_codegen.cpp | 286 +++++++++--------- 3 files changed, 252 insertions(+), 279 deletions(-) diff --git a/clang/test/OpenMP/declare_mapper_codegen.cpp b/clang/test/OpenMP/declare_mapper_codegen.cpp index 7dc32d0ae12ff..69f9dae7d4988 100644 --- a/clang/test/OpenMP/declare_mapper_codegen.cpp +++ b/clang/test/OpenMP/declare_mapper_codegen.cpp @@ -92,10 +92,7 @@ class C { // CK0-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1 // CK0-DAG: [[PTREND:%.+]] = getelementptr %class.C, ptr [[BEGIN]], i64 [[SIZE]] // CK0-DAG: [[PTRSNE:%.+]] = icmp ne ptr [[BPTR]], [[BEGIN]] -// CK0-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16 -// CK0-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0 -// CK0-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]] -// CK0-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]] +// CK0-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]] // CK0-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 // CK0-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 // CK0-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]] @@ -592,10 +589,7 @@ class C { // CK1-DAG: [[PTREND:%.+]] = getelementptr %class.C, ptr [[BEGIN]], i64 [[SIZE]] // CK1-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1 // CK1-DAG: [[PTRSNE:%.+]] = icmp ne ptr [[BPTR]], [[BEGIN]] -// CK1-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16 -// CK1-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0 -// CK1-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]] -// CK1-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]] +// CK1-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]] // CK1-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 // CK1-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 // CK1-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]] @@ -702,10 +696,7 @@ class C { // CK2-DAG: [[PTREND:%.+]] = getelementptr %class.C, ptr [[BEGIN]], i64 [[SIZE]] // CK2-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1 // CK2-DAG: [[PTRSNE:%.+]] = icmp ne ptr [[BPTR]], [[BEGIN]] -// CK2-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16 -// CK2-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0 -// CK2-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]] -// CK2-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]] +// CK2-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]] // CK2-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 // CK2-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 // CK2-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]] @@ -897,10 +888,7 @@ class C { // CK4-DAG: [[PTREND:%.+]] = getelementptr %class.C, ptr [[BEGIN]], i64 [[SIZE]] // CK4-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1 // CK4-DAG: [[PTRSNE:%.+]] = icmp ne ptr [[BPTR]], [[BEGIN]] -// CK4-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16 -// CK4-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0 -// CK4-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]] -// CK4-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]] +// CK4-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]] // CK4-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 // CK4-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 // CK4-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]] @@ -1087,10 +1075,7 @@ void foo(){ // CK5-DAG: [[PTREND:%.+]] = getelementptr %struct.myvec, ptr [[BEGIN]], i64 [[SIZE]] // CK5-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1 // CK5-DAG: [[PTRSNE:%.+]] = icmp ne ptr [[BPTR]], [[BEGIN]] -// CK5-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16 -// CK5-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0 -// CK5-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]] -// CK5-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]] +// CK5-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]] // CK5-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 // CK5-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 // CK5-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]] 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 index 5df1e958ad55a..054e8f22633cb 100644 --- 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 @@ -37,15 +37,15 @@ void foo() { // CHECK-LABEL: define {{[^@]+}}@_Z3foov // CHECK-SAME: () #[[ATTR0:[0-9]+]] { // CHECK-NEXT: entry: -// CHECK-NEXT: [[SA:%.*]] = alloca [10 x %struct.D], align 4 +// 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: [[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: [[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 @@ -99,10 +99,10 @@ void foo() { // 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: [[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: [[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 @@ -117,18 +117,15 @@ void foo() { // 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: [[TMP10:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP9]] // 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-NEXT: [[TMP11:%.*]] = and i1 [[TMP10]], [[DOTOMP_ARRAY__INIT__DELETE]] +// CHECK-NEXT: br i1 [[TMP11]], 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: [[TMP12:%.*]] = mul nuw i64 [[TMP6]], 12 +// CHECK-NEXT: [[TMP13:%.*]] = and i64 [[TMP4]], -4 +// CHECK-NEXT: [[TMP14:%.*]] = or i64 [[TMP13]], 512 +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP12]], i64 [[TMP14]], ptr [[TMP5]]) // CHECK-NEXT: br label [[OMP_ARRAYMAP_HEAD]] // CHECK: omp.arraymap.head: // CHECK-NEXT: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP7]] @@ -138,115 +135,115 @@ void foo() { // 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-NEXT: [[TMP15:%.*]] = getelementptr i32, ptr [[H]], i32 1 +// CHECK-NEXT: [[TMP16:%.*]] = ptrtoint ptr [[TMP15]] to i64 +// CHECK-NEXT: [[TMP17:%.*]] = ptrtoint ptr [[E]] to i64 +// CHECK-NEXT: [[TMP18:%.*]] = sub i64 [[TMP16]], [[TMP17]] +// CHECK-NEXT: [[TMP19:%.*]] = sdiv exact i64 [[TMP18]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64) +// CHECK-NEXT: [[TMP20:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]]) +// CHECK-NEXT: [[TMP21:%.*]] = shl i64 [[TMP20]], 48 +// CHECK-NEXT: [[TMP22:%.*]] = add nuw i64 0, [[TMP21]] +// CHECK-NEXT: [[TMP23:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP24:%.*]] = icmp eq i64 [[TMP23]], 0 +// CHECK-NEXT: br i1 [[TMP24]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]] // CHECK: omp.type.alloc: -// CHECK-NEXT: [[TMP28:%.*]] = and i64 [[TMP25]], -4 +// CHECK-NEXT: [[TMP25:%.*]] = and i64 [[TMP22]], -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-NEXT: [[TMP26:%.*]] = icmp eq i64 [[TMP23]], 1 +// CHECK-NEXT: br i1 [[TMP26]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]] // CHECK: omp.type.to: -// CHECK-NEXT: [[TMP30:%.*]] = and i64 [[TMP25]], -3 +// CHECK-NEXT: [[TMP27:%.*]] = and i64 [[TMP22]], -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-NEXT: [[TMP28:%.*]] = icmp eq i64 [[TMP23]], 2 +// CHECK-NEXT: br i1 [[TMP28]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]] // CHECK: omp.type.from: -// CHECK-NEXT: [[TMP32:%.*]] = and i64 [[TMP25]], -2 +// CHECK-NEXT: [[TMP29:%.*]] = and i64 [[TMP22]], -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-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP25]], [[OMP_TYPE_ALLOC]] ], [ [[TMP27]], [[OMP_TYPE_TO]] ], [ [[TMP29]], [[OMP_TYPE_FROM]] ], [ [[TMP22]], [[OMP_TYPE_TO_ELSE]] ] +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP19]], i64 [[OMP_MAPTYPE]], ptr null) +// CHECK-NEXT: [[TMP30:%.*]] = add nuw i64 281474976711171, [[TMP21]] +// CHECK-NEXT: [[TMP31:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP32:%.*]] = icmp eq i64 [[TMP31]], 0 +// CHECK-NEXT: br i1 [[TMP32]], label [[OMP_TYPE_ALLOC1:%.*]], label [[OMP_TYPE_ALLOC_ELSE2:%.*]] // CHECK: omp.type.alloc1: -// CHECK-NEXT: [[TMP36:%.*]] = and i64 [[TMP33]], -4 +// CHECK-NEXT: [[TMP33:%.*]] = and i64 [[TMP30]], -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-NEXT: [[TMP34:%.*]] = icmp eq i64 [[TMP31]], 1 +// CHECK-NEXT: br i1 [[TMP34]], label [[OMP_TYPE_TO3:%.*]], label [[OMP_TYPE_TO_ELSE4:%.*]] // CHECK: omp.type.to3: -// CHECK-NEXT: [[TMP38:%.*]] = and i64 [[TMP33]], -3 +// CHECK-NEXT: [[TMP35:%.*]] = and i64 [[TMP30]], -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-NEXT: [[TMP36:%.*]] = icmp eq i64 [[TMP31]], 2 +// CHECK-NEXT: br i1 [[TMP36]], label [[OMP_TYPE_FROM5:%.*]], label [[OMP_TYPE_END6]] // CHECK: omp.type.from5: -// CHECK-NEXT: [[TMP40:%.*]] = and i64 [[TMP33]], -2 +// CHECK-NEXT: [[TMP37:%.*]] = and i64 [[TMP30]], -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: [[OMP_MAPTYPE7:%.*]] = phi i64 [ [[TMP33]], [[OMP_TYPE_ALLOC1]] ], [ [[TMP35]], [[OMP_TYPE_TO3]] ], [ [[TMP37]], [[OMP_TYPE_FROM5]] ], [ [[TMP30]], [[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-NEXT: [[TMP38:%.*]] = add nuw i64 281474976711171, [[TMP21]] +// CHECK-NEXT: [[TMP39:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP40:%.*]] = icmp eq i64 [[TMP39]], 0 +// CHECK-NEXT: br i1 [[TMP40]], label [[OMP_TYPE_ALLOC8:%.*]], label [[OMP_TYPE_ALLOC_ELSE9:%.*]] // CHECK: omp.type.alloc8: -// CHECK-NEXT: [[TMP44:%.*]] = and i64 [[TMP41]], -4 +// CHECK-NEXT: [[TMP41:%.*]] = and i64 [[TMP38]], -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-NEXT: [[TMP42:%.*]] = icmp eq i64 [[TMP39]], 1 +// CHECK-NEXT: br i1 [[TMP42]], label [[OMP_TYPE_TO10:%.*]], label [[OMP_TYPE_TO_ELSE11:%.*]] // CHECK: omp.type.to10: -// CHECK-NEXT: [[TMP46:%.*]] = and i64 [[TMP41]], -3 +// CHECK-NEXT: [[TMP43:%.*]] = and i64 [[TMP38]], -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-NEXT: [[TMP44:%.*]] = icmp eq i64 [[TMP39]], 2 +// CHECK-NEXT: br i1 [[TMP44]], label [[OMP_TYPE_FROM12:%.*]], label [[OMP_TYPE_END13]] // CHECK: omp.type.from12: -// CHECK-NEXT: [[TMP48:%.*]] = and i64 [[TMP41]], -2 +// CHECK-NEXT: [[TMP45:%.*]] = and i64 [[TMP38]], -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: [[OMP_MAPTYPE14:%.*]] = phi i64 [ [[TMP41]], [[OMP_TYPE_ALLOC8]] ], [ [[TMP43]], [[OMP_TYPE_TO10]] ], [ [[TMP45]], [[OMP_TYPE_FROM12]] ], [ [[TMP38]], [[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-NEXT: [[TMP46:%.*]] = add nuw i64 281474976711171, [[TMP21]] +// CHECK-NEXT: [[TMP47:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP48:%.*]] = icmp eq i64 [[TMP47]], 0 +// CHECK-NEXT: br i1 [[TMP48]], label [[OMP_TYPE_ALLOC15:%.*]], label [[OMP_TYPE_ALLOC_ELSE16:%.*]] // CHECK: omp.type.alloc15: -// CHECK-NEXT: [[TMP52:%.*]] = and i64 [[TMP49]], -4 +// CHECK-NEXT: [[TMP49:%.*]] = and i64 [[TMP46]], -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-NEXT: [[TMP50:%.*]] = icmp eq i64 [[TMP47]], 1 +// CHECK-NEXT: br i1 [[TMP50]], label [[OMP_TYPE_TO17:%.*]], label [[OMP_TYPE_TO_ELSE18:%.*]] // CHECK: omp.type.to17: -// CHECK-NEXT: [[TMP54:%.*]] = and i64 [[TMP49]], -3 +// CHECK-NEXT: [[TMP51:%.*]] = and i64 [[TMP46]], -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-NEXT: [[TMP52:%.*]] = icmp eq i64 [[TMP47]], 2 +// CHECK-NEXT: br i1 [[TMP52]], label [[OMP_TYPE_FROM19:%.*]], label [[OMP_TYPE_END20]] // CHECK: omp.type.from19: -// CHECK-NEXT: [[TMP56:%.*]] = and i64 [[TMP49]], -2 +// CHECK-NEXT: [[TMP53:%.*]] = and i64 [[TMP46]], -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: [[OMP_MAPTYPE21:%.*]] = phi i64 [ [[TMP49]], [[OMP_TYPE_ALLOC15]] ], [ [[TMP51]], [[OMP_TYPE_TO17]] ], [ [[TMP53]], [[OMP_TYPE_FROM19]] ], [ [[TMP46]], [[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-NEXT: [[TMP54:%.*]] = and i64 [[TMP4]], 8 +// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP54]], 0 +// CHECK-NEXT: [[TMP55:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY22]], [[DOTOMP_ARRAY__DEL__DELETE]] +// CHECK-NEXT: br i1 [[TMP55]], 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: [[TMP56:%.*]] = mul nuw i64 [[TMP6]], 12 +// CHECK-NEXT: [[TMP57:%.*]] = and i64 [[TMP4]], -4 +// CHECK-NEXT: [[TMP58:%.*]] = or i64 [[TMP57]], 512 +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP56]], i64 [[TMP58]], ptr [[TMP5]]) // CHECK-NEXT: br label [[OMP_DONE]] // CHECK: omp.done: // CHECK-NEXT: ret void @@ -260,18 +257,15 @@ void foo() { // 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: [[TMP10:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP9]] // 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-NEXT: [[TMP11:%.*]] = and i1 [[TMP10]], [[DOTOMP_ARRAY__INIT__DELETE]] +// CHECK-NEXT: br i1 [[TMP11]], 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: [[TMP12:%.*]] = mul nuw i64 [[TMP6]], 4 +// CHECK-NEXT: [[TMP13:%.*]] = and i64 [[TMP4]], -4 +// CHECK-NEXT: [[TMP14:%.*]] = or i64 [[TMP13]], 512 +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP12]], i64 [[TMP14]], ptr [[TMP5]]) // CHECK-NEXT: br label [[OMP_ARRAYMAP_HEAD]] // CHECK: omp.arraymap.head: // CHECK-NEXT: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP7]] @@ -279,44 +273,44 @@ void foo() { // 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-NEXT: [[TMP15:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]]) +// CHECK-NEXT: [[TMP16:%.*]] = shl i64 [[TMP15]], 48 +// CHECK-NEXT: [[TMP17:%.*]] = add nuw i64 1, [[TMP16]] +// CHECK-NEXT: [[TMP18:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP19:%.*]] = icmp eq i64 [[TMP18]], 0 +// CHECK-NEXT: br i1 [[TMP19]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]] // CHECK: omp.type.alloc: -// CHECK-NEXT: [[TMP23:%.*]] = and i64 [[TMP20]], -4 +// CHECK-NEXT: [[TMP20:%.*]] = and i64 [[TMP17]], -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-NEXT: [[TMP21:%.*]] = icmp eq i64 [[TMP18]], 1 +// CHECK-NEXT: br i1 [[TMP21]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]] // CHECK: omp.type.to: -// CHECK-NEXT: [[TMP25:%.*]] = and i64 [[TMP20]], -3 +// CHECK-NEXT: [[TMP22:%.*]] = and i64 [[TMP17]], -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-NEXT: [[TMP23:%.*]] = icmp eq i64 [[TMP18]], 2 +// CHECK-NEXT: br i1 [[TMP23]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]] // CHECK: omp.type.from: -// CHECK-NEXT: [[TMP27:%.*]] = and i64 [[TMP20]], -2 +// CHECK-NEXT: [[TMP24:%.*]] = and i64 [[TMP17]], -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: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP20]], [[OMP_TYPE_ALLOC]] ], [ [[TMP22]], [[OMP_TYPE_TO]] ], [ [[TMP24]], [[OMP_TYPE_FROM]] ], [ [[TMP17]], [[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-NEXT: [[TMP25:%.*]] = and i64 [[TMP4]], 8 +// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP25]], 0 +// CHECK-NEXT: [[TMP26:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY1]], [[DOTOMP_ARRAY__DEL__DELETE]] +// CHECK-NEXT: br i1 [[TMP26]], 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: [[TMP27:%.*]] = mul nuw i64 [[TMP6]], 4 +// CHECK-NEXT: [[TMP28:%.*]] = and i64 [[TMP4]], -4 +// CHECK-NEXT: [[TMP29:%.*]] = or i64 [[TMP28]], 512 +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP27]], i64 [[TMP29]], ptr [[TMP5]]) // CHECK-NEXT: br label [[OMP_DONE]] // CHECK: omp.done: // CHECK-NEXT: ret void diff --git a/clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_codegen.cpp b/clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_codegen.cpp index 0fc6de0e4279a..a017fb098ced9 100644 --- a/clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_codegen.cpp +++ b/clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_codegen.cpp @@ -33,19 +33,19 @@ void foo() { // CHECK-LABEL: define {{[^@]+}}@_Z3foov // CHECK-SAME: () #[[ATTR0:[0-9]+]] { // CHECK-NEXT: entry: -// CHECK-NEXT: [[SA:%.*]] = alloca [10 x %struct.D], align 4 +// 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: [[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: [[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: [[ARRAYIDX2:%.*]] = getelementptr inbounds nuw [10 x %struct.D], ptr [[SA]], i64 0, i64 0 +// CHECK-NEXT: [[ARRAYIDX2:%.*]] = getelementptr inbounds nuw [10 x [[STRUCT_D]]], ptr [[SA]], i64 0, i64 0 // CHECK-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 // CHECK-NEXT: store ptr [[SA]], ptr [[TMP0]], align 8 // CHECK-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 @@ -95,11 +95,11 @@ void foo() { // CHECK-NEXT: entry: // CHECK-NEXT: [[SA_ADDR:%.*]] = alloca ptr, align 8 // CHECK-NEXT: store ptr [[SA]], ptr [[SA_ADDR]], align 8 -// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SA_ADDR]], align 8 -// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x %struct.D], ptr [[TMP0]], i64 0, i64 1 -// CHECK-NEXT: [[E:%.*]] = getelementptr inbounds nuw [[STRUCT_D:%.*]], ptr [[ARRAYIDX]], i32 0, i32 0 +// 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: [[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 @@ -109,141 +109,138 @@ void foo() { // 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: [[TMP10:%.*]] = udiv exact i64 [[TMP3]], 12 -// CHECK-NEXT: [[TMP11:%.*]] = getelementptr [[STRUCT_D:%.*]], ptr [[TMP2]], i64 [[TMP10]] -// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP10]], 1 -// CHECK-NEXT: [[TMP14:%.*]] = and i64 [[TMP4]], 8 -// CHECK-NEXT: [[TMP15:%.*]] = icmp ne ptr [[TMP1]], [[TMP2]] -// CHECK-NEXT: [[TMP16:%.*]] = and i64 [[TMP4]], 16 -// CHECK-NEXT: [[TMP17:%.*]] = icmp ne i64 [[TMP16]], 0 -// CHECK-NEXT: [[TMP18:%.*]] = and i1 [[TMP15]], [[TMP17]] -// CHECK-NEXT: [[TMP19:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP18]] -// CHECK-NEXT: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP14]], 0 -// CHECK-NEXT: [[TMP20:%.*]] = and i1 [[TMP19]], [[DOTOMP_ARRAY__INIT__DELETE]] -// CHECK-NEXT: br i1 [[TMP20]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]] +// CHECK-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:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP9]] +// CHECK-NEXT: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP8]], 0 +// CHECK-NEXT: [[TMP11:%.*]] = and i1 [[TMP10]], [[DOTOMP_ARRAY__INIT__DELETE]] +// CHECK-NEXT: br i1 [[TMP11]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]] // CHECK: .omp.array..init: -// CHECK-NEXT: [[TMP21:%.*]] = mul nuw i64 [[TMP10]], 12 -// CHECK-NEXT: [[TMP22:%.*]] = and i64 [[TMP4]], -4 -// CHECK-NEXT: [[TMP23:%.*]] = or i64 [[TMP22]], 512 -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP21]], i64 [[TMP23]], ptr [[TMP5]]) +// CHECK-NEXT: [[TMP12:%.*]] = mul nuw i64 [[TMP6]], 12 +// CHECK-NEXT: [[TMP13:%.*]] = and i64 [[TMP4]], -4 +// CHECK-NEXT: [[TMP14:%.*]] = or i64 [[TMP13]], 512 +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP12]], i64 [[TMP14]], ptr [[TMP5]]) // CHECK-NEXT: br label [[OMP_ARRAYMAP_HEAD]] // CHECK: omp.arraymap.head: -// CHECK-NEXT: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP11]] +// 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_END25:%.*]] ] +// 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: [[TMP24:%.*]] = getelementptr i32, ptr [[H]], i32 1 -// CHECK-NEXT: [[TMP25:%.*]] = ptrtoint ptr [[TMP24]] to i64 -// CHECK-NEXT: [[TMP26:%.*]] = ptrtoint ptr [[E]] to i64 -// CHECK-NEXT: [[TMP27:%.*]] = sub i64 [[TMP25]], [[TMP26]] -// CHECK-NEXT: [[TMP28:%.*]] = sdiv exact i64 [[TMP27]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64) -// CHECK-NEXT: [[TMP29:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]]) -// CHECK-NEXT: [[TMP30:%.*]] = shl i64 [[TMP29]], 48 -// CHECK-NEXT: [[TMP31:%.*]] = add nuw i64 0, [[TMP30]] -// CHECK-NEXT: [[TMP32:%.*]] = and i64 [[TMP4]], 3 -// CHECK-NEXT: [[TMP33:%.*]] = icmp eq i64 [[TMP32]], 0 -// CHECK-NEXT: br i1 [[TMP33]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]] +// CHECK-NEXT: [[TMP15:%.*]] = getelementptr i32, ptr [[H]], i32 1 +// CHECK-NEXT: [[TMP16:%.*]] = ptrtoint ptr [[TMP15]] to i64 +// CHECK-NEXT: [[TMP17:%.*]] = ptrtoint ptr [[E]] to i64 +// CHECK-NEXT: [[TMP18:%.*]] = sub i64 [[TMP16]], [[TMP17]] +// CHECK-NEXT: [[TMP19:%.*]] = sdiv exact i64 [[TMP18]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64) +// CHECK-NEXT: [[TMP20:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]]) +// CHECK-NEXT: [[TMP21:%.*]] = shl i64 [[TMP20]], 48 +// CHECK-NEXT: [[TMP22:%.*]] = add nuw i64 0, [[TMP21]] +// CHECK-NEXT: [[TMP23:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP24:%.*]] = icmp eq i64 [[TMP23]], 0 +// CHECK-NEXT: br i1 [[TMP24]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]] // CHECK: omp.type.alloc: -// CHECK-NEXT: [[TMP34:%.*]] = and i64 [[TMP31]], -4 +// CHECK-NEXT: [[TMP25:%.*]] = and i64 [[TMP22]], -4 // CHECK-NEXT: br label [[OMP_TYPE_END:%.*]] // CHECK: omp.type.alloc.else: -// CHECK-NEXT: [[TMP35:%.*]] = icmp eq i64 [[TMP32]], 1 -// CHECK-NEXT: br i1 [[TMP35]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]] +// CHECK-NEXT: [[TMP26:%.*]] = icmp eq i64 [[TMP23]], 1 +// CHECK-NEXT: br i1 [[TMP26]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]] // CHECK: omp.type.to: -// CHECK-NEXT: [[TMP36:%.*]] = and i64 [[TMP31]], -3 +// CHECK-NEXT: [[TMP27:%.*]] = and i64 [[TMP22]], -3 // CHECK-NEXT: br label [[OMP_TYPE_END]] // CHECK: omp.type.to.else: -// CHECK-NEXT: [[TMP37:%.*]] = icmp eq i64 [[TMP32]], 2 -// CHECK-NEXT: br i1 [[TMP37]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]] +// CHECK-NEXT: [[TMP28:%.*]] = icmp eq i64 [[TMP23]], 2 +// CHECK-NEXT: br i1 [[TMP28]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]] // CHECK: omp.type.from: -// CHECK-NEXT: [[TMP38:%.*]] = and i64 [[TMP31]], -2 +// CHECK-NEXT: [[TMP29:%.*]] = and i64 [[TMP22]], -2 // CHECK-NEXT: br label [[OMP_TYPE_END]] // CHECK: omp.type.end: -// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP34]], [[OMP_TYPE_ALLOC]] ], [ [[TMP36]], [[OMP_TYPE_TO]] ], [ [[TMP38]], [[OMP_TYPE_FROM]] ], [ [[TMP31]], [[OMP_TYPE_TO_ELSE]] ] -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP28]], i64 [[OMP_MAPTYPE]], ptr null) -// CHECK-NEXT: [[TMP39:%.*]] = add nuw i64 281474976711171, [[TMP30]] -// CHECK-NEXT: [[TMP40:%.*]] = and i64 [[TMP4]], 3 -// CHECK-NEXT: [[TMP41:%.*]] = icmp eq i64 [[TMP40]], 0 -// CHECK-NEXT: br i1 [[TMP41]], label [[OMP_TYPE_ALLOC6:%.*]], label [[OMP_TYPE_ALLOC_ELSE7:%.*]] +// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP25]], [[OMP_TYPE_ALLOC]] ], [ [[TMP27]], [[OMP_TYPE_TO]] ], [ [[TMP29]], [[OMP_TYPE_FROM]] ], [ [[TMP22]], [[OMP_TYPE_TO_ELSE]] ] +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP19]], i64 [[OMP_MAPTYPE]], ptr null) +// CHECK-NEXT: [[TMP30:%.*]] = add nuw i64 281474976711171, [[TMP21]] +// CHECK-NEXT: [[TMP31:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP32:%.*]] = icmp eq i64 [[TMP31]], 0 +// CHECK-NEXT: br i1 [[TMP32]], label [[OMP_TYPE_ALLOC1:%.*]], label [[OMP_TYPE_ALLOC_ELSE2:%.*]] // CHECK: omp.type.alloc1: -// CHECK-NEXT: [[TMP42:%.*]] = and i64 [[TMP39]], -4 -// CHECK-NEXT: br label [[OMP_TYPE_END11:%.*]] +// CHECK-NEXT: [[TMP33:%.*]] = and i64 [[TMP30]], -4 +// CHECK-NEXT: br label [[OMP_TYPE_END6:%.*]] // CHECK: omp.type.alloc.else2: -// CHECK-NEXT: [[TMP43:%.*]] = icmp eq i64 [[TMP40]], 1 -// CHECK-NEXT: br i1 [[TMP43]], label [[OMP_TYPE_TO8:%.*]], label [[OMP_TYPE_TO_ELSE9:%.*]] +// CHECK-NEXT: [[TMP34:%.*]] = icmp eq i64 [[TMP31]], 1 +// CHECK-NEXT: br i1 [[TMP34]], label [[OMP_TYPE_TO3:%.*]], label [[OMP_TYPE_TO_ELSE4:%.*]] // CHECK: omp.type.to3: -// CHECK-NEXT: [[TMP44:%.*]] = and i64 [[TMP39]], -3 -// CHECK-NEXT: br label [[OMP_TYPE_END11]] +// CHECK-NEXT: [[TMP35:%.*]] = and i64 [[TMP30]], -3 +// CHECK-NEXT: br label [[OMP_TYPE_END6]] // CHECK: omp.type.to.else4: -// CHECK-NEXT: [[TMP45:%.*]] = icmp eq i64 [[TMP40]], 2 -// CHECK-NEXT: br i1 [[TMP45]], label [[OMP_TYPE_FROM10:%.*]], label [[OMP_TYPE_END11]] +// CHECK-NEXT: [[TMP36:%.*]] = icmp eq i64 [[TMP31]], 2 +// CHECK-NEXT: br i1 [[TMP36]], label [[OMP_TYPE_FROM5:%.*]], label [[OMP_TYPE_END6]] // CHECK: omp.type.from5: -// CHECK-NEXT: [[TMP46:%.*]] = and i64 [[TMP39]], -2 -// CHECK-NEXT: br label [[OMP_TYPE_END11]] +// CHECK-NEXT: [[TMP37:%.*]] = and i64 [[TMP30]], -2 +// CHECK-NEXT: br label [[OMP_TYPE_END6]] // CHECK: omp.type.end6: -// CHECK-NEXT: [[OMP_MAPTYPE12:%.*]] = phi i64 [ [[TMP42]], [[OMP_TYPE_ALLOC6]] ], [ [[TMP44]], [[OMP_TYPE_TO8]] ], [ [[TMP46]], [[OMP_TYPE_FROM10]] ], [ [[TMP39]], [[OMP_TYPE_TO_ELSE9]] ] -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 4, i64 [[OMP_MAPTYPE12]], ptr null) -// CHECK-NEXT: [[TMP47:%.*]] = add nuw i64 281474976711171, [[TMP30]] -// CHECK-NEXT: [[TMP48:%.*]] = and i64 [[TMP4]], 3 -// CHECK-NEXT: [[TMP49:%.*]] = icmp eq i64 [[TMP48]], 0 -// CHECK-NEXT: br i1 [[TMP49]], label [[OMP_TYPE_ALLOC13:%.*]], label [[OMP_TYPE_ALLOC_ELSE14:%.*]] +// CHECK-NEXT: [[OMP_MAPTYPE7:%.*]] = phi i64 [ [[TMP33]], [[OMP_TYPE_ALLOC1]] ], [ [[TMP35]], [[OMP_TYPE_TO3]] ], [ [[TMP37]], [[OMP_TYPE_FROM5]] ], [ [[TMP30]], [[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: [[TMP38:%.*]] = add nuw i64 281474976711171, [[TMP21]] +// CHECK-NEXT: [[TMP39:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP40:%.*]] = icmp eq i64 [[TMP39]], 0 +// CHECK-NEXT: br i1 [[TMP40]], label [[OMP_TYPE_ALLOC8:%.*]], label [[OMP_TYPE_ALLOC_ELSE9:%.*]] // CHECK: omp.type.alloc8: -// CHECK-NEXT: [[TMP50:%.*]] = and i64 [[TMP47]], -4 -// CHECK-NEXT: br label [[OMP_TYPE_END18:%.*]] +// CHECK-NEXT: [[TMP41:%.*]] = and i64 [[TMP38]], -4 +// CHECK-NEXT: br label [[OMP_TYPE_END13:%.*]] // CHECK: omp.type.alloc.else9: -// CHECK-NEXT: [[TMP51:%.*]] = icmp eq i64 [[TMP48]], 1 -// CHECK-NEXT: br i1 [[TMP51]], label [[OMP_TYPE_TO15:%.*]], label [[OMP_TYPE_TO_ELSE16:%.*]] +// CHECK-NEXT: [[TMP42:%.*]] = icmp eq i64 [[TMP39]], 1 +// CHECK-NEXT: br i1 [[TMP42]], label [[OMP_TYPE_TO10:%.*]], label [[OMP_TYPE_TO_ELSE11:%.*]] // CHECK: omp.type.to10: -// CHECK-NEXT: [[TMP52:%.*]] = and i64 [[TMP47]], -3 -// CHECK-NEXT: br label [[OMP_TYPE_END18]] +// CHECK-NEXT: [[TMP43:%.*]] = and i64 [[TMP38]], -3 +// CHECK-NEXT: br label [[OMP_TYPE_END13]] // CHECK: omp.type.to.else11: -// CHECK-NEXT: [[TMP53:%.*]] = icmp eq i64 [[TMP48]], 2 -// CHECK-NEXT: br i1 [[TMP53]], label [[OMP_TYPE_FROM17:%.*]], label [[OMP_TYPE_END18]] +// CHECK-NEXT: [[TMP44:%.*]] = icmp eq i64 [[TMP39]], 2 +// CHECK-NEXT: br i1 [[TMP44]], label [[OMP_TYPE_FROM12:%.*]], label [[OMP_TYPE_END13]] // CHECK: omp.type.from12: -// CHECK-NEXT: [[TMP54:%.*]] = and i64 [[TMP47]], -2 -// CHECK-NEXT: br label [[OMP_TYPE_END18]] +// CHECK-NEXT: [[TMP45:%.*]] = and i64 [[TMP38]], -2 +// CHECK-NEXT: br label [[OMP_TYPE_END13]] // CHECK: omp.type.end13: -// CHECK-NEXT: [[OMP_MAPTYPE19:%.*]] = phi i64 [ [[TMP50]], [[OMP_TYPE_ALLOC13]] ], [ [[TMP52]], [[OMP_TYPE_TO15]] ], [ [[TMP54]], [[OMP_TYPE_FROM17]] ], [ [[TMP47]], [[OMP_TYPE_TO_ELSE16]] ] -// CHECK-NEXT: call void @.omp_mapper._ZTS1C.default(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[F]], i64 4, i64 [[OMP_MAPTYPE19]], ptr null) #[[ATTR3]] -// CHECK-NEXT: [[TMP55:%.*]] = add nuw i64 281474976711171, [[TMP30]] -// CHECK-NEXT: [[TMP56:%.*]] = and i64 [[TMP4]], 3 -// CHECK-NEXT: [[TMP57:%.*]] = icmp eq i64 [[TMP56]], 0 -// CHECK-NEXT: br i1 [[TMP57]], label [[OMP_TYPE_ALLOC20:%.*]], label [[OMP_TYPE_ALLOC_ELSE21:%.*]] +// CHECK-NEXT: [[OMP_MAPTYPE14:%.*]] = phi i64 [ [[TMP41]], [[OMP_TYPE_ALLOC8]] ], [ [[TMP43]], [[OMP_TYPE_TO10]] ], [ [[TMP45]], [[OMP_TYPE_FROM12]] ], [ [[TMP38]], [[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: [[TMP46:%.*]] = add nuw i64 281474976711171, [[TMP21]] +// CHECK-NEXT: [[TMP47:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP48:%.*]] = icmp eq i64 [[TMP47]], 0 +// CHECK-NEXT: br i1 [[TMP48]], label [[OMP_TYPE_ALLOC15:%.*]], label [[OMP_TYPE_ALLOC_ELSE16:%.*]] // CHECK: omp.type.alloc15: -// CHECK-NEXT: [[TMP58:%.*]] = and i64 [[TMP55]], -4 -// CHECK-NEXT: br label [[OMP_TYPE_END25]] +// CHECK-NEXT: [[TMP49:%.*]] = and i64 [[TMP46]], -4 +// CHECK-NEXT: br label [[OMP_TYPE_END20]] // CHECK: omp.type.alloc.else16: -// CHECK-NEXT: [[TMP59:%.*]] = icmp eq i64 [[TMP56]], 1 -// CHECK-NEXT: br i1 [[TMP59]], label [[OMP_TYPE_TO22:%.*]], label [[OMP_TYPE_TO_ELSE23:%.*]] +// CHECK-NEXT: [[TMP50:%.*]] = icmp eq i64 [[TMP47]], 1 +// CHECK-NEXT: br i1 [[TMP50]], label [[OMP_TYPE_TO17:%.*]], label [[OMP_TYPE_TO_ELSE18:%.*]] // CHECK: omp.type.to17: -// CHECK-NEXT: [[TMP60:%.*]] = and i64 [[TMP55]], -3 -// CHECK-NEXT: br label [[OMP_TYPE_END25]] +// CHECK-NEXT: [[TMP51:%.*]] = and i64 [[TMP46]], -3 +// CHECK-NEXT: br label [[OMP_TYPE_END20]] // CHECK: omp.type.to.else18: -// CHECK-NEXT: [[TMP61:%.*]] = icmp eq i64 [[TMP56]], 2 -// CHECK-NEXT: br i1 [[TMP61]], label [[OMP_TYPE_FROM24:%.*]], label [[OMP_TYPE_END25]] +// CHECK-NEXT: [[TMP52:%.*]] = icmp eq i64 [[TMP47]], 2 +// CHECK-NEXT: br i1 [[TMP52]], label [[OMP_TYPE_FROM19:%.*]], label [[OMP_TYPE_END20]] // CHECK: omp.type.from19: -// CHECK-NEXT: [[TMP62:%.*]] = and i64 [[TMP55]], -2 -// CHECK-NEXT: br label [[OMP_TYPE_END25]] +// CHECK-NEXT: [[TMP53:%.*]] = and i64 [[TMP46]], -2 +// CHECK-NEXT: br label [[OMP_TYPE_END20]] // CHECK: omp.type.end20: -// CHECK-NEXT: [[OMP_MAPTYPE26:%.*]] = phi i64 [ [[TMP58]], [[OMP_TYPE_ALLOC20]] ], [ [[TMP60]], [[OMP_TYPE_TO22]] ], [ [[TMP62]], [[OMP_TYPE_FROM24]] ], [ [[TMP55]], [[OMP_TYPE_TO_ELSE23]] ] -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[H]], i64 4, i64 [[OMP_MAPTYPE26]], ptr null) +// CHECK-NEXT: [[OMP_MAPTYPE21:%.*]] = phi i64 [ [[TMP49]], [[OMP_TYPE_ALLOC15]] ], [ [[TMP51]], [[OMP_TYPE_TO17]] ], [ [[TMP53]], [[OMP_TYPE_FROM19]] ], [ [[TMP46]], [[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]], [[TMP11]] +// 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_ISARRAY27:%.*]] = icmp sgt i64 [[TMP10]], 1 -// CHECK-NEXT: [[TMP63:%.*]] = and i64 [[TMP4]], 8 -// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP63]], 0 -// CHECK-NEXT: [[TMP64:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY27]], [[DOTOMP_ARRAY__DEL__DELETE]] -// CHECK-NEXT: br i1 [[TMP64]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] +// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY22:%.*]] = icmp sgt i64 [[TMP6]], 1 +// CHECK-NEXT: [[TMP54:%.*]] = and i64 [[TMP4]], 8 +// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP54]], 0 +// CHECK-NEXT: [[TMP55:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY22]], [[DOTOMP_ARRAY__DEL__DELETE]] +// CHECK-NEXT: br i1 [[TMP55]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] // CHECK: .omp.array..del: -// CHECK-NEXT: [[TMP65:%.*]] = mul nuw i64 [[TMP10]], 12 -// CHECK-NEXT: [[TMP66:%.*]] = and i64 [[TMP4]], -4 -// CHECK-NEXT: [[TMP67:%.*]] = or i64 [[TMP66]], 512 -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP65]], i64 [[TMP67]], ptr [[TMP5]]) +// CHECK-NEXT: [[TMP56:%.*]] = mul nuw i64 [[TMP6]], 12 +// CHECK-NEXT: [[TMP57:%.*]] = and i64 [[TMP4]], -4 +// CHECK-NEXT: [[TMP58:%.*]] = or i64 [[TMP57]], 512 +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP56]], i64 [[TMP58]], ptr [[TMP5]]) // CHECK-NEXT: br label [[OMP_DONE]] // CHECK: omp.done: // CHECK-NEXT: ret void @@ -252,68 +249,65 @@ void foo() { // 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: [[TMP10:%.*]] = udiv exact i64 [[TMP3]], 4 -// CHECK-NEXT: [[TMP11:%.*]] = getelementptr [[STRUCT_C:%.*]], ptr [[TMP2]], i64 [[TMP10]] -// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP10]], 1 -// CHECK-NEXT: [[TMP14:%.*]] = and i64 [[TMP4]], 8 -// CHECK-NEXT: [[TMP15:%.*]] = icmp ne ptr [[TMP1]], [[TMP2]] -// CHECK-NEXT: [[TMP16:%.*]] = and i64 [[TMP4]], 16 -// CHECK-NEXT: [[TMP17:%.*]] = icmp ne i64 [[TMP16]], 0 -// CHECK-NEXT: [[TMP18:%.*]] = and i1 [[TMP15]], [[TMP17]] -// CHECK-NEXT: [[TMP19:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP18]] -// CHECK-NEXT: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP14]], 0 -// CHECK-NEXT: [[TMP20:%.*]] = and i1 [[TMP19]], [[DOTOMP_ARRAY__INIT__DELETE]] -// CHECK-NEXT: br i1 [[TMP20]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]] +// CHECK-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:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP9]] +// CHECK-NEXT: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP8]], 0 +// CHECK-NEXT: [[TMP11:%.*]] = and i1 [[TMP10]], [[DOTOMP_ARRAY__INIT__DELETE]] +// CHECK-NEXT: br i1 [[TMP11]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]] // CHECK: .omp.array..init: -// CHECK-NEXT: [[TMP21:%.*]] = mul nuw i64 [[TMP10]], 4 -// CHECK-NEXT: [[TMP22:%.*]] = and i64 [[TMP4]], -4 -// CHECK-NEXT: [[TMP23:%.*]] = or i64 [[TMP22]], 512 -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP21]], i64 [[TMP23]], ptr [[TMP5]]) +// CHECK-NEXT: [[TMP12:%.*]] = mul nuw i64 [[TMP6]], 4 +// CHECK-NEXT: [[TMP13:%.*]] = and i64 [[TMP4]], -4 +// CHECK-NEXT: [[TMP14:%.*]] = or i64 [[TMP13]], 512 +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP12]], i64 [[TMP14]], ptr [[TMP5]]) // CHECK-NEXT: br label [[OMP_ARRAYMAP_HEAD]] // CHECK: omp.arraymap.head: -// CHECK-NEXT: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP11]] +// 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: [[TMP24:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]]) -// CHECK-NEXT: [[TMP25:%.*]] = shl i64 [[TMP24]], 48 -// CHECK-NEXT: [[TMP26:%.*]] = add nuw i64 1, [[TMP25]] -// CHECK-NEXT: [[TMP27:%.*]] = and i64 [[TMP4]], 3 -// CHECK-NEXT: [[TMP28:%.*]] = icmp eq i64 [[TMP27]], 0 -// CHECK-NEXT: br i1 [[TMP28]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]] +// CHECK-NEXT: [[TMP15:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]]) +// CHECK-NEXT: [[TMP16:%.*]] = shl i64 [[TMP15]], 48 +// CHECK-NEXT: [[TMP17:%.*]] = add nuw i64 1, [[TMP16]] +// CHECK-NEXT: [[TMP18:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP19:%.*]] = icmp eq i64 [[TMP18]], 0 +// CHECK-NEXT: br i1 [[TMP19]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]] // CHECK: omp.type.alloc: -// CHECK-NEXT: [[TMP29:%.*]] = and i64 [[TMP26]], -4 +// CHECK-NEXT: [[TMP20:%.*]] = and i64 [[TMP17]], -4 // CHECK-NEXT: br label [[OMP_TYPE_END]] // CHECK: omp.type.alloc.else: -// CHECK-NEXT: [[TMP30:%.*]] = icmp eq i64 [[TMP27]], 1 -// CHECK-NEXT: br i1 [[TMP30]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]] +// CHECK-NEXT: [[TMP21:%.*]] = icmp eq i64 [[TMP18]], 1 +// CHECK-NEXT: br i1 [[TMP21]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]] // CHECK: omp.type.to: -// CHECK-NEXT: [[TMP31:%.*]] = and i64 [[TMP26]], -3 +// CHECK-NEXT: [[TMP22:%.*]] = and i64 [[TMP17]], -3 // CHECK-NEXT: br label [[OMP_TYPE_END]] // CHECK: omp.type.to.else: -// CHECK-NEXT: [[TMP32:%.*]] = icmp eq i64 [[TMP27]], 2 -// CHECK-NEXT: br i1 [[TMP32]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]] +// CHECK-NEXT: [[TMP23:%.*]] = icmp eq i64 [[TMP18]], 2 +// CHECK-NEXT: br i1 [[TMP23]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]] // CHECK: omp.type.from: -// CHECK-NEXT: [[TMP33:%.*]] = and i64 [[TMP26]], -2 +// CHECK-NEXT: [[TMP24:%.*]] = and i64 [[TMP17]], -2 // CHECK-NEXT: br label [[OMP_TYPE_END]] // CHECK: omp.type.end: -// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP29]], [[OMP_TYPE_ALLOC]] ], [ [[TMP31]], [[OMP_TYPE_TO]] ], [ [[TMP33]], [[OMP_TYPE_FROM]] ], [ [[TMP26]], [[OMP_TYPE_TO_ELSE]] ] +// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP20]], [[OMP_TYPE_ALLOC]] ], [ [[TMP22]], [[OMP_TYPE_TO]] ], [ [[TMP24]], [[OMP_TYPE_FROM]] ], [ [[TMP17]], [[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]], [[TMP11]] +// 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_ISARRAY6:%.*]] = icmp sgt i64 [[TMP10]], 1 -// CHECK-NEXT: [[TMP34:%.*]] = and i64 [[TMP4]], 8 -// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP34]], 0 -// CHECK-NEXT: [[TMP35:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY6]], [[DOTOMP_ARRAY__DEL__DELETE]] -// CHECK-NEXT: br i1 [[TMP35]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] +// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY1:%.*]] = icmp sgt i64 [[TMP6]], 1 +// CHECK-NEXT: [[TMP25:%.*]] = and i64 [[TMP4]], 8 +// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP25]], 0 +// CHECK-NEXT: [[TMP26:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY1]], [[DOTOMP_ARRAY__DEL__DELETE]] +// CHECK-NEXT: br i1 [[TMP26]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] // CHECK: .omp.array..del: -// CHECK-NEXT: [[TMP36:%.*]] = mul nuw i64 [[TMP10]], 4 -// CHECK-NEXT: [[TMP37:%.*]] = and i64 [[TMP4]], -4 -// CHECK-NEXT: [[TMP38:%.*]] = or i64 [[TMP37]], 512 -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP36]], i64 [[TMP38]], ptr [[TMP5]]) +// CHECK-NEXT: [[TMP27:%.*]] = mul nuw i64 [[TMP6]], 4 +// CHECK-NEXT: [[TMP28:%.*]] = and i64 [[TMP4]], -4 +// CHECK-NEXT: [[TMP29:%.*]] = or i64 [[TMP28]], 512 +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP27]], i64 [[TMP29]], ptr [[TMP5]]) // CHECK-NEXT: br label [[OMP_DONE]] // CHECK: omp.done: // CHECK-NEXT: ret void >From 854a548f432ba9a728142eacc8be2fe0755282d2 Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Tue, 20 Jan 2026 16:48:13 -0800 Subject: [PATCH 3/4] Clang-format fixes. --- offload/libomptarget/omptarget.cpp | 2 +- offload/test/mapping/declare_mapper_target_checks.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp index 4ae018b1b27eb..bd99edee5e1b3 100644 --- a/offload/libomptarget/omptarget.cpp +++ b/offload/libomptarget/omptarget.cpp @@ -1110,7 +1110,7 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum, int64_t DataSize = ArgSizes[I]; bool IsImplicit = ArgTypes[I] & OMP_TGT_MAPTYPE_IMPLICIT; bool UpdateRef = !(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) || - (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ); + (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ); bool ForceDelete = ArgTypes[I] & OMP_TGT_MAPTYPE_DELETE; bool HasPresentModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_PRESENT; bool HasHoldModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_OMPX_HOLD; diff --git a/offload/test/mapping/declare_mapper_target_checks.cpp b/offload/test/mapping/declare_mapper_target_checks.cpp index 3136138f9b030..1c5edae99a7c8 100644 --- a/offload/test/mapping/declare_mapper_target_checks.cpp +++ b/offload/test/mapping/declare_mapper_target_checks.cpp @@ -75,7 +75,7 @@ template <typename T> int testTypeNestedPtr(T t1[2], T t2[3], T t3[4]) { for (int i = 0; i < 2; i++) t1[i].n.i1 = t3[i].n.i1 = 1; -#pragma omp target map(tofrom : t1[0:2], t2[0:3], t3[0:4]) +#pragma omp target map(tofrom : t1[0 : 2], t2[0 : 3], t3[0 : 4]) for (int i = 0; i < 2; i++) { t1[i].n.i3 = t3[i].n.i3 = t1[i].n.i1; t1[i].n.i1 = t3[i].n.i1 = 7; >From e2486cb7425852e2db602e1eace88fe1b6591043 Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Tue, 20 Jan 2026 16:50:42 -0800 Subject: [PATCH 4/4] Update MLIR test. --- mlir/test/Target/LLVMIR/omptarget-llvm.mlir | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/mlir/test/Target/LLVMIR/omptarget-llvm.mlir b/mlir/test/Target/LLVMIR/omptarget-llvm.mlir index 0b4d63125f82f..0548c6a178d0c 100644 --- a/mlir/test/Target/LLVMIR/omptarget-llvm.mlir +++ b/mlir/test/Target/LLVMIR/omptarget-llvm.mlir @@ -564,10 +564,7 @@ module attributes {omp.target_triples = ["amdgcn-amd-amdhsa"]} { // CHECK: %[[VAL_20:.*]] = icmp sgt i64 %[[VAL_15]], 1 // CHECK: %[[VAL_21:.*]] = and i64 %[[VAL_22:.*]], 8 // CHECK: %[[VAL_23:.*]] = icmp ne ptr %[[VAL_24:.*]], %[[VAL_19]] -// CHECK: %[[VAL_25:.*]] = and i64 %[[VAL_22]], 16 -// CHECK: %[[VAL_26:.*]] = icmp ne i64 %[[VAL_25]], 0 -// CHECK: %[[VAL_27:.*]] = and i1 %[[VAL_23]], %[[VAL_26]] -// CHECK: %[[VAL_28:.*]] = or i1 %[[VAL_20]], %[[VAL_27]] +// CHECK: %[[VAL_28:.*]] = or i1 %[[VAL_20]], %[[VAL_23]] // CHECK: %[[VAL_29:.*]] = icmp eq i64 %[[VAL_21]], 0 // CHECK: %[[VAL_30:.*]] = and i1 %[[VAL_28]], %[[VAL_29]] // CHECK: br i1 %[[VAL_30]], label %[[VAL_31:.*]], label %[[VAL_32:.*]] _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
