Author: Abhinav Gaba Date: 2026-08-02T23:07:32-07:00 New Revision: 187c39ba03cfed96c8aab9d0c951d0bea6f27b82
URL: https://github.com/llvm/llvm-project/commit/187c39ba03cfed96c8aab9d0c951d0bea6f27b82 DIFF: https://github.com/llvm/llvm-project/commit/187c39ba03cfed96c8aab9d0c951d0bea6f27b82.diff LOG: [OpenMP] Propagate ALWAYS/DELETE/CLOSE map-type modifiers to mapper entries (#210210) Per OpenMP, when a map/motion clause uses a mapper, any map-type-modifying modifier on that clause applies to each map the declared mapper specifies. This change propagates the `ALWAYS`, `DELETE`, and `CLOSE` bits from the outer clause's map type into every entry emitted by emitUserDefinedMapper, except `ATTACH` entries (`ATTACH`|`ALWAYS` is reserved for `attach(always)`, and the other bits have no meaning for an `ATTACH` entry). `PRESENT` is intentionally NOT propagated here: it requires distinguishing pointee entries from the struct's own storage and is handled in a follow-up. Co-authored-by: Claude Opus 4.8 <[email protected]> Added: Modified: clang/docs/ReleaseNotes.md clang/test/OpenMP/declare_mapper_codegen.cpp clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_codegen.cpp clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp mlir/test/Target/LLVMIR/omptarget-llvm.mlir offload/test/mapping/mapper_map_always_from.c offload/test/mapping/mapper_map_always_to_enter_data.c offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c offload/test/mapping/mapper_map_ptee_only_always_array.c offload/test/offloading/fortran/mapper-map-always-to-enter-data.f90 Removed: ################################################################################ diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md index a38b99ff8e075..7108392abbaa1 100644 --- a/clang/docs/ReleaseNotes.md +++ b/clang/docs/ReleaseNotes.md @@ -546,6 +546,8 @@ features cannot lower the translation-unit ABI level; - Added parsing and semantic support for `dims` modifier in `num_teams` and `thread_limit` clauses for OpenMP 6.1 or later. +- Map-type-modifying modifiers applied to a list item with a user-defined mapper + are now propagated onto the maps the mapper expands to. ### SYCL Support diff --git a/clang/test/OpenMP/declare_mapper_codegen.cpp b/clang/test/OpenMP/declare_mapper_codegen.cpp index ff891396fe72e..eb90f218f5ca8 100644 --- a/clang/test/OpenMP/declare_mapper_codegen.cpp +++ b/clang/test/OpenMP/declare_mapper_codegen.cpp @@ -148,7 +148,9 @@ class C { // CK0-DAG: br label %[[TYEND]] // CK0-DAG: [[TYEND]] // CK0-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] -// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 [[CUSIZE]], i64 [[PHITYPE0]], {{.*}}) +// CK0-DAG: [[MODMASK0:%.+]] = and i64 [[TYPE]], 1036 +// CK0-DAG: [[PHITYPE0_MOD:%.+]] = or i64 [[PHITYPE0]], [[MODMASK0]] +// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 [[CUSIZE]], i64 [[PHITYPE0_MOD]], {{.*}}) // 281474976710659 == 0x1,000,000,003 // CK0-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]] // CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 @@ -171,7 +173,9 @@ class C { // CK0-DAG: br label %[[TYEND]] // CK0-DAG: [[TYEND]] // CK0-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] -// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[TYPE1]], {{.*}}) +// CK0-DAG: [[MODMASK1:%.+]] = and i64 [[TYPE]], 1036 +// CK0-DAG: [[TYPE1_MOD:%.+]] = or i64 [[TYPE1]], [[MODMASK1]] +// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[TYPE1_MOD]], {{.*}}) // 281474976710675 == 0x1,000,000,013 // CK0-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]] // CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 @@ -194,7 +198,9 @@ class C { // CK0-DAG: br label %[[TYEND]] // CK0-DAG: [[TYEND]] // CK0-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] -// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BBEGIN]], ptr [[BARRBEGINGEP]], i64 16, i64 [[TYPE2]], {{.*}}) +// CK0-DAG: [[MODMASK2:%.+]] = and i64 [[TYPE]], 1036 +// CK0-DAG: [[TYPE2_MOD:%.+]] = or i64 [[TYPE2]], [[MODMASK2]] +// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BBEGIN]], ptr [[BARRBEGINGEP]], i64 16, i64 [[TYPE2_MOD]], {{.*}}) // CK0: [[PTRNEXT]] = getelementptr %class.C, ptr [[PTR]], i32 1 // CK0: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]] // CK0: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]] @@ -636,7 +642,9 @@ class C { // CK1-DAG: br label %[[TYEND]] // CK1-DAG: [[TYEND]] // CK1-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] -// CK1: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[TYPE1]], {{.*}}) +// CK1-DAG: [[MODMASK:%.+]] = and i64 [[TYPE]], 1036 +// CK1-DAG: [[TYPE1_MOD:%.+]] = or i64 [[TYPE1]], [[MODMASK]] +// CK1: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[TYPE1_MOD]], {{.*}}) // CK1: [[PTRNEXT]] = getelementptr %class.C, ptr [[PTR]], i32 1 // CK1: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]] // CK1: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]] @@ -743,7 +751,9 @@ class C { // CK2-DAG: br label %[[TYEND]] // CK2-DAG: [[TYEND]] // CK2-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] -// CK2: call void [[BMPRFUNC]](ptr [[HANDLE]], ptr [[PTR]], ptr [[BBEGIN]], i64 8, i64 [[TYPE1]], {{.*}}) +// CK2-DAG: [[MODMASK:%.+]] = and i64 [[TYPE]], 1036 +// CK2-DAG: [[TYPE1_MOD:%.+]] = or i64 [[TYPE1]], [[MODMASK]] +// CK2: call void [[BMPRFUNC]](ptr [[HANDLE]], ptr [[PTR]], ptr [[BBEGIN]], i64 8, i64 [[TYPE1_MOD]], {{.*}}) // CK2: [[PTRNEXT]] = getelementptr %class.C, ptr [[PTR]], i32 1 // CK2: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]] // CK2: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]] @@ -948,7 +958,9 @@ class C { // CK4-DAG: br label %[[TYEND]] // CK4-DAG: [[TYEND]] // CK4-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] -// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 [[CUSIZE]], i64 [[PHITYPE0]], {{.*}}) +// CK4-DAG: [[MODMASK0:%.+]] = and i64 [[TYPE]], 1036 +// CK4-DAG: [[PHITYPE0_MOD:%.+]] = or i64 [[PHITYPE0]], [[MODMASK0]] +// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 [[CUSIZE]], i64 [[PHITYPE0_MOD]], {{.*}}) // 281474976710659 == 0x1,000,000,003 // CK4-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]] // CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 @@ -971,7 +983,9 @@ class C { // CK4-DAG: br label %[[TYEND]] // CK4-DAG: [[TYEND]] // CK4-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] -// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[TYPE1]], {{.*}}) +// CK4-DAG: [[MODMASK1:%.+]] = and i64 [[TYPE]], 1036 +// CK4-DAG: [[TYPE1_MOD:%.+]] = or i64 [[TYPE1]], [[MODMASK1]] +// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[TYPE1_MOD]], {{.*}}) // 281474976710675 == 0x1,000,000,013 // CK4-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]] // CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 @@ -994,7 +1008,9 @@ class C { // CK4-DAG: br label %[[TYEND]] // CK4-DAG: [[TYEND]] // CK4-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] -// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BBEGIN]], ptr [[BARRBEGINGEP]], i64 16, i64 [[TYPE2]], {{.*}}) +// CK4-DAG: [[MODMASK2:%.+]] = and i64 [[TYPE]], 1036 +// CK4-DAG: [[TYPE2_MOD:%.+]] = or i64 [[TYPE2]], [[MODMASK2]] +// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BBEGIN]], ptr [[BARRBEGINGEP]], i64 16, i64 [[TYPE2_MOD]], {{.*}}) // CK4: [[PTRNEXT]] = getelementptr %class.C, ptr [[PTR]], i32 1 // CK4: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]] // CK4: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]] @@ -1126,7 +1142,9 @@ void foo(){ // CK5-DAG: br label %[[TYEND]] // CK5-DAG: [[TYEND]] // CK5-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] -// CK5: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 {{.*}}, i64 [[TYPE1]], {{.*}}) +// CK5-DAG: [[MODMASK:%.+]] = and i64 [[TYPE]], 1036 +// CK5-DAG: [[TYPE1_MOD:%.+]] = or i64 [[TYPE1]], [[MODMASK]] +// CK5: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 {{.*}}, i64 [[TYPE1_MOD]], {{.*}}) // CK5: [[PTRNEXT]] = getelementptr %struct.myvec, ptr [[PTR]], i32 1 // CK5: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]] // CK5: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]] 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 c9bc3a54c38f2..d3270669da3da 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 @@ -139,7 +139,7 @@ void foo() { // 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: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END22:%.*]] ] // 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 @@ -170,87 +170,95 @@ void foo() { // CHECK-NEXT: br label [[OMP_TYPE_END]] // CHECK: omp.type.end: // CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP24]], [[OMP_TYPE_ALLOC]] ], [ [[TMP26]], [[OMP_TYPE_TO]] ], [ [[TMP28]], [[OMP_TYPE_FROM]] ], [ [[TMP21]], [[OMP_TYPE_TO_ELSE]] ] -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP18]], i64 [[OMP_MAPTYPE]], ptr null) -// CHECK-NEXT: [[TMP29:%.*]] = add nuw i64 281474976711171, [[TMP20]] -// CHECK-NEXT: [[TMP30:%.*]] = and i64 [[TMP4]], 3 -// CHECK-NEXT: [[TMP31:%.*]] = icmp eq i64 [[TMP30]], 0 -// CHECK-NEXT: br i1 [[TMP31]], label [[OMP_TYPE_ALLOC1:%.*]], label [[OMP_TYPE_ALLOC_ELSE2:%.*]] +// CHECK-NEXT: [[TMP29:%.*]] = and i64 [[TMP4]], 1036 +// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS:%.*]] = or i64 [[OMP_MAPTYPE]], [[TMP29]] +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP18]], i64 [[OMP_MAPTYPE_WITH_MODIFIERS]], ptr null) +// CHECK-NEXT: [[TMP30:%.*]] = add nuw i64 281474976711171, [[TMP20]] +// 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: [[TMP32:%.*]] = and i64 [[TMP29]], -4 +// CHECK-NEXT: [[TMP33:%.*]] = and i64 [[TMP30]], -4 // CHECK-NEXT: br label [[OMP_TYPE_END6:%.*]] // CHECK: omp.type.alloc.else2: -// CHECK-NEXT: [[TMP33:%.*]] = icmp eq i64 [[TMP30]], 1 -// CHECK-NEXT: br i1 [[TMP33]], 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: [[TMP34:%.*]] = and i64 [[TMP29]], -3 +// CHECK-NEXT: [[TMP35:%.*]] = and i64 [[TMP30]], -3 // CHECK-NEXT: br label [[OMP_TYPE_END6]] // CHECK: omp.type.to.else4: -// CHECK-NEXT: [[TMP35:%.*]] = icmp eq i64 [[TMP30]], 2 -// CHECK-NEXT: br i1 [[TMP35]], 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: [[TMP36:%.*]] = and i64 [[TMP29]], -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 [ [[TMP32]], [[OMP_TYPE_ALLOC1]] ], [ [[TMP34]], [[OMP_TYPE_TO3]] ], [ [[TMP36]], [[OMP_TYPE_FROM5]] ], [ [[TMP29]], [[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: [[TMP37:%.*]] = add nuw i64 281474976711171, [[TMP20]] -// CHECK-NEXT: [[TMP38:%.*]] = and i64 [[TMP4]], 3 -// CHECK-NEXT: [[TMP39:%.*]] = icmp eq i64 [[TMP38]], 0 -// CHECK-NEXT: br i1 [[TMP39]], label [[OMP_TYPE_ALLOC8:%.*]], label [[OMP_TYPE_ALLOC_ELSE9:%.*]] -// CHECK: omp.type.alloc8: -// CHECK-NEXT: [[TMP40:%.*]] = and i64 [[TMP37]], -4 -// CHECK-NEXT: br label [[OMP_TYPE_END13:%.*]] -// CHECK: omp.type.alloc.else9: -// CHECK-NEXT: [[TMP41:%.*]] = icmp eq i64 [[TMP38]], 1 -// CHECK-NEXT: br i1 [[TMP41]], label [[OMP_TYPE_TO10:%.*]], label [[OMP_TYPE_TO_ELSE11:%.*]] -// CHECK: omp.type.to10: -// CHECK-NEXT: [[TMP42:%.*]] = and i64 [[TMP37]], -3 -// CHECK-NEXT: br label [[OMP_TYPE_END13]] -// CHECK: omp.type.to.else11: -// CHECK-NEXT: [[TMP43:%.*]] = icmp eq i64 [[TMP38]], 2 -// CHECK-NEXT: br i1 [[TMP43]], label [[OMP_TYPE_FROM12:%.*]], label [[OMP_TYPE_END13]] -// CHECK: omp.type.from12: -// CHECK-NEXT: [[TMP44:%.*]] = and i64 [[TMP37]], -2 -// CHECK-NEXT: br label [[OMP_TYPE_END13]] -// CHECK: omp.type.end13: -// CHECK-NEXT: [[OMP_MAPTYPE14:%.*]] = phi i64 [ [[TMP40]], [[OMP_TYPE_ALLOC8]] ], [ [[TMP42]], [[OMP_TYPE_TO10]] ], [ [[TMP44]], [[OMP_TYPE_FROM12]] ], [ [[TMP37]], [[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: [[TMP45:%.*]] = add nuw i64 281474976711171, [[TMP20]] -// CHECK-NEXT: [[TMP46:%.*]] = and i64 [[TMP4]], 3 -// CHECK-NEXT: [[TMP47:%.*]] = icmp eq i64 [[TMP46]], 0 -// CHECK-NEXT: br i1 [[TMP47]], label [[OMP_TYPE_ALLOC15:%.*]], label [[OMP_TYPE_ALLOC_ELSE16:%.*]] -// CHECK: omp.type.alloc15: -// CHECK-NEXT: [[TMP48:%.*]] = and i64 [[TMP45]], -4 -// CHECK-NEXT: br label [[OMP_TYPE_END20]] -// CHECK: omp.type.alloc.else16: -// CHECK-NEXT: [[TMP49:%.*]] = icmp eq i64 [[TMP46]], 1 -// CHECK-NEXT: br i1 [[TMP49]], label [[OMP_TYPE_TO17:%.*]], label [[OMP_TYPE_TO_ELSE18:%.*]] -// CHECK: omp.type.to17: -// CHECK-NEXT: [[TMP50:%.*]] = and i64 [[TMP45]], -3 -// CHECK-NEXT: br label [[OMP_TYPE_END20]] -// CHECK: omp.type.to.else18: -// CHECK-NEXT: [[TMP51:%.*]] = icmp eq i64 [[TMP46]], 2 -// CHECK-NEXT: br i1 [[TMP51]], label [[OMP_TYPE_FROM19:%.*]], label [[OMP_TYPE_END20]] -// CHECK: omp.type.from19: -// CHECK-NEXT: [[TMP52:%.*]] = and i64 [[TMP45]], -2 -// CHECK-NEXT: br label [[OMP_TYPE_END20]] -// CHECK: omp.type.end20: -// CHECK-NEXT: [[OMP_MAPTYPE21:%.*]] = phi i64 [ [[TMP48]], [[OMP_TYPE_ALLOC15]] ], [ [[TMP50]], [[OMP_TYPE_TO17]] ], [ [[TMP52]], [[OMP_TYPE_FROM19]] ], [ [[TMP45]], [[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_MAPTYPE7:%.*]] = phi i64 [ [[TMP33]], [[OMP_TYPE_ALLOC1]] ], [ [[TMP35]], [[OMP_TYPE_TO3]] ], [ [[TMP37]], [[OMP_TYPE_FROM5]] ], [ [[TMP30]], [[OMP_TYPE_TO_ELSE4]] ] +// CHECK-NEXT: [[TMP38:%.*]] = and i64 [[TMP4]], 1036 +// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS8:%.*]] = or i64 [[OMP_MAPTYPE7]], [[TMP38]] +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS8]], ptr null) +// CHECK-NEXT: [[TMP39:%.*]] = add nuw i64 281474976711171, [[TMP20]] +// CHECK-NEXT: [[TMP40:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP41:%.*]] = icmp eq i64 [[TMP40]], 0 +// CHECK-NEXT: br i1 [[TMP41]], label [[OMP_TYPE_ALLOC9:%.*]], label [[OMP_TYPE_ALLOC_ELSE10:%.*]] +// CHECK: omp.type.alloc9: +// CHECK-NEXT: [[TMP42:%.*]] = and i64 [[TMP39]], -4 +// CHECK-NEXT: br label [[OMP_TYPE_END14:%.*]] +// CHECK: omp.type.alloc.else10: +// CHECK-NEXT: [[TMP43:%.*]] = icmp eq i64 [[TMP40]], 1 +// CHECK-NEXT: br i1 [[TMP43]], label [[OMP_TYPE_TO11:%.*]], label [[OMP_TYPE_TO_ELSE12:%.*]] +// CHECK: omp.type.to11: +// CHECK-NEXT: [[TMP44:%.*]] = and i64 [[TMP39]], -3 +// CHECK-NEXT: br label [[OMP_TYPE_END14]] +// CHECK: omp.type.to.else12: +// CHECK-NEXT: [[TMP45:%.*]] = icmp eq i64 [[TMP40]], 2 +// CHECK-NEXT: br i1 [[TMP45]], label [[OMP_TYPE_FROM13:%.*]], label [[OMP_TYPE_END14]] +// CHECK: omp.type.from13: +// CHECK-NEXT: [[TMP46:%.*]] = and i64 [[TMP39]], -2 +// CHECK-NEXT: br label [[OMP_TYPE_END14]] +// CHECK: omp.type.end14: +// CHECK-NEXT: [[OMP_MAPTYPE15:%.*]] = phi i64 [ [[TMP42]], [[OMP_TYPE_ALLOC9]] ], [ [[TMP44]], [[OMP_TYPE_TO11]] ], [ [[TMP46]], [[OMP_TYPE_FROM13]] ], [ [[TMP39]], [[OMP_TYPE_TO_ELSE12]] ] +// CHECK-NEXT: [[TMP47:%.*]] = and i64 [[TMP4]], 1036 +// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS16:%.*]] = or i64 [[OMP_MAPTYPE15]], [[TMP47]] +// CHECK-NEXT: call void @.omp_mapper._ZTS1C.default(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[F]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS16]], ptr null) #[[ATTR3]] +// CHECK-NEXT: [[TMP48:%.*]] = add nuw i64 281474976711171, [[TMP20]] +// CHECK-NEXT: [[TMP49:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP50:%.*]] = icmp eq i64 [[TMP49]], 0 +// CHECK-NEXT: br i1 [[TMP50]], label [[OMP_TYPE_ALLOC17:%.*]], label [[OMP_TYPE_ALLOC_ELSE18:%.*]] +// CHECK: omp.type.alloc17: +// CHECK-NEXT: [[TMP51:%.*]] = and i64 [[TMP48]], -4 +// CHECK-NEXT: br label [[OMP_TYPE_END22]] +// CHECK: omp.type.alloc.else18: +// CHECK-NEXT: [[TMP52:%.*]] = icmp eq i64 [[TMP49]], 1 +// CHECK-NEXT: br i1 [[TMP52]], label [[OMP_TYPE_TO19:%.*]], label [[OMP_TYPE_TO_ELSE20:%.*]] +// CHECK: omp.type.to19: +// CHECK-NEXT: [[TMP53:%.*]] = and i64 [[TMP48]], -3 +// CHECK-NEXT: br label [[OMP_TYPE_END22]] +// CHECK: omp.type.to.else20: +// CHECK-NEXT: [[TMP54:%.*]] = icmp eq i64 [[TMP49]], 2 +// CHECK-NEXT: br i1 [[TMP54]], label [[OMP_TYPE_FROM21:%.*]], label [[OMP_TYPE_END22]] +// CHECK: omp.type.from21: +// CHECK-NEXT: [[TMP55:%.*]] = and i64 [[TMP48]], -2 +// CHECK-NEXT: br label [[OMP_TYPE_END22]] +// CHECK: omp.type.end22: +// CHECK-NEXT: [[OMP_MAPTYPE23:%.*]] = phi i64 [ [[TMP51]], [[OMP_TYPE_ALLOC17]] ], [ [[TMP53]], [[OMP_TYPE_TO19]] ], [ [[TMP55]], [[OMP_TYPE_FROM21]] ], [ [[TMP48]], [[OMP_TYPE_TO_ELSE20]] ] +// CHECK-NEXT: [[TMP56:%.*]] = and i64 [[TMP4]], 1036 +// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS24:%.*]] = or i64 [[OMP_MAPTYPE23]], [[TMP56]] +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[H]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS24]], 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: [[TMP53:%.*]] = and i64 [[TMP4]], 8 -// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP53]], 0 -// CHECK-NEXT: [[TMP54:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY22]], [[DOTOMP_ARRAY__DEL__DELETE]] -// CHECK-NEXT: br i1 [[TMP54]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] +// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY25:%.*]] = 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_ISARRAY25]], [[DOTOMP_ARRAY__DEL__DELETE]] +// CHECK-NEXT: br i1 [[TMP58]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] // CHECK: .omp.array..del: -// CHECK-NEXT: [[TMP55:%.*]] = mul nuw i64 [[TMP6]], 12 -// CHECK-NEXT: [[TMP56:%.*]] = and i64 [[TMP4]], -4 -// CHECK-NEXT: [[TMP57:%.*]] = or i64 [[TMP56]], 512 -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP55]], i64 [[TMP57]], ptr [[TMP5]]) +// 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 @@ -303,21 +311,23 @@ void foo() { // CHECK-NEXT: br label [[OMP_TYPE_END]] // CHECK: omp.type.end: // 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: [[TMP25:%.*]] = and i64 [[TMP4]], 1036 +// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS:%.*]] = or i64 [[OMP_MAPTYPE]], [[TMP25]] +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[A]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS]], 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: [[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-NEXT: [[TMP26:%.*]] = and i64 [[TMP4]], 8 +// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP26]], 0 +// CHECK-NEXT: [[TMP27:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY1]], [[DOTOMP_ARRAY__DEL__DELETE]] +// CHECK-NEXT: br i1 [[TMP27]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] // CHECK: .omp.array..del: -// 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: [[TMP28:%.*]] = mul nuw i64 [[TMP6]], 4 +// CHECK-NEXT: [[TMP29:%.*]] = and i64 [[TMP4]], -4 +// CHECK-NEXT: [[TMP30:%.*]] = or i64 [[TMP29]], 512 +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP28]], i64 [[TMP30]], 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 f530c92c8f1fe..8723ad765c5fb 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 @@ -136,7 +136,7 @@ void foo() { // 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: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END22:%.*]] ] // 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 @@ -167,87 +167,95 @@ void foo() { // CHECK-NEXT: br label [[OMP_TYPE_END]] // CHECK: omp.type.end: // CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP24]], [[OMP_TYPE_ALLOC]] ], [ [[TMP26]], [[OMP_TYPE_TO]] ], [ [[TMP28]], [[OMP_TYPE_FROM]] ], [ [[TMP21]], [[OMP_TYPE_TO_ELSE]] ] -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP18]], i64 [[OMP_MAPTYPE]], ptr null) -// CHECK-NEXT: [[TMP29:%.*]] = add nuw i64 281474976711171, [[TMP20]] -// CHECK-NEXT: [[TMP30:%.*]] = and i64 [[TMP4]], 3 -// CHECK-NEXT: [[TMP31:%.*]] = icmp eq i64 [[TMP30]], 0 -// CHECK-NEXT: br i1 [[TMP31]], label [[OMP_TYPE_ALLOC1:%.*]], label [[OMP_TYPE_ALLOC_ELSE2:%.*]] +// CHECK-NEXT: [[TMP29:%.*]] = and i64 [[TMP4]], 1036 +// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS:%.*]] = or i64 [[OMP_MAPTYPE]], [[TMP29]] +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP18]], i64 [[OMP_MAPTYPE_WITH_MODIFIERS]], ptr null) +// CHECK-NEXT: [[TMP30:%.*]] = add nuw i64 281474976711171, [[TMP20]] +// 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: [[TMP32:%.*]] = and i64 [[TMP29]], -4 +// CHECK-NEXT: [[TMP33:%.*]] = and i64 [[TMP30]], -4 // CHECK-NEXT: br label [[OMP_TYPE_END6:%.*]] // CHECK: omp.type.alloc.else2: -// CHECK-NEXT: [[TMP33:%.*]] = icmp eq i64 [[TMP30]], 1 -// CHECK-NEXT: br i1 [[TMP33]], 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: [[TMP34:%.*]] = and i64 [[TMP29]], -3 +// CHECK-NEXT: [[TMP35:%.*]] = and i64 [[TMP30]], -3 // CHECK-NEXT: br label [[OMP_TYPE_END6]] // CHECK: omp.type.to.else4: -// CHECK-NEXT: [[TMP35:%.*]] = icmp eq i64 [[TMP30]], 2 -// CHECK-NEXT: br i1 [[TMP35]], 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: [[TMP36:%.*]] = and i64 [[TMP29]], -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 [ [[TMP32]], [[OMP_TYPE_ALLOC1]] ], [ [[TMP34]], [[OMP_TYPE_TO3]] ], [ [[TMP36]], [[OMP_TYPE_FROM5]] ], [ [[TMP29]], [[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: [[TMP37:%.*]] = add nuw i64 281474976711171, [[TMP20]] -// CHECK-NEXT: [[TMP38:%.*]] = and i64 [[TMP4]], 3 -// CHECK-NEXT: [[TMP39:%.*]] = icmp eq i64 [[TMP38]], 0 -// CHECK-NEXT: br i1 [[TMP39]], label [[OMP_TYPE_ALLOC8:%.*]], label [[OMP_TYPE_ALLOC_ELSE9:%.*]] -// CHECK: omp.type.alloc8: -// CHECK-NEXT: [[TMP40:%.*]] = and i64 [[TMP37]], -4 -// CHECK-NEXT: br label [[OMP_TYPE_END13:%.*]] -// CHECK: omp.type.alloc.else9: -// CHECK-NEXT: [[TMP41:%.*]] = icmp eq i64 [[TMP38]], 1 -// CHECK-NEXT: br i1 [[TMP41]], label [[OMP_TYPE_TO10:%.*]], label [[OMP_TYPE_TO_ELSE11:%.*]] -// CHECK: omp.type.to10: -// CHECK-NEXT: [[TMP42:%.*]] = and i64 [[TMP37]], -3 -// CHECK-NEXT: br label [[OMP_TYPE_END13]] -// CHECK: omp.type.to.else11: -// CHECK-NEXT: [[TMP43:%.*]] = icmp eq i64 [[TMP38]], 2 -// CHECK-NEXT: br i1 [[TMP43]], label [[OMP_TYPE_FROM12:%.*]], label [[OMP_TYPE_END13]] -// CHECK: omp.type.from12: -// CHECK-NEXT: [[TMP44:%.*]] = and i64 [[TMP37]], -2 -// CHECK-NEXT: br label [[OMP_TYPE_END13]] -// CHECK: omp.type.end13: -// CHECK-NEXT: [[OMP_MAPTYPE14:%.*]] = phi i64 [ [[TMP40]], [[OMP_TYPE_ALLOC8]] ], [ [[TMP42]], [[OMP_TYPE_TO10]] ], [ [[TMP44]], [[OMP_TYPE_FROM12]] ], [ [[TMP37]], [[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: [[TMP45:%.*]] = add nuw i64 281474976711171, [[TMP20]] -// CHECK-NEXT: [[TMP46:%.*]] = and i64 [[TMP4]], 3 -// CHECK-NEXT: [[TMP47:%.*]] = icmp eq i64 [[TMP46]], 0 -// CHECK-NEXT: br i1 [[TMP47]], label [[OMP_TYPE_ALLOC15:%.*]], label [[OMP_TYPE_ALLOC_ELSE16:%.*]] -// CHECK: omp.type.alloc15: -// CHECK-NEXT: [[TMP48:%.*]] = and i64 [[TMP45]], -4 -// CHECK-NEXT: br label [[OMP_TYPE_END20]] -// CHECK: omp.type.alloc.else16: -// CHECK-NEXT: [[TMP49:%.*]] = icmp eq i64 [[TMP46]], 1 -// CHECK-NEXT: br i1 [[TMP49]], label [[OMP_TYPE_TO17:%.*]], label [[OMP_TYPE_TO_ELSE18:%.*]] -// CHECK: omp.type.to17: -// CHECK-NEXT: [[TMP50:%.*]] = and i64 [[TMP45]], -3 -// CHECK-NEXT: br label [[OMP_TYPE_END20]] -// CHECK: omp.type.to.else18: -// CHECK-NEXT: [[TMP51:%.*]] = icmp eq i64 [[TMP46]], 2 -// CHECK-NEXT: br i1 [[TMP51]], label [[OMP_TYPE_FROM19:%.*]], label [[OMP_TYPE_END20]] -// CHECK: omp.type.from19: -// CHECK-NEXT: [[TMP52:%.*]] = and i64 [[TMP45]], -2 -// CHECK-NEXT: br label [[OMP_TYPE_END20]] -// CHECK: omp.type.end20: -// CHECK-NEXT: [[OMP_MAPTYPE21:%.*]] = phi i64 [ [[TMP48]], [[OMP_TYPE_ALLOC15]] ], [ [[TMP50]], [[OMP_TYPE_TO17]] ], [ [[TMP52]], [[OMP_TYPE_FROM19]] ], [ [[TMP45]], [[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_MAPTYPE7:%.*]] = phi i64 [ [[TMP33]], [[OMP_TYPE_ALLOC1]] ], [ [[TMP35]], [[OMP_TYPE_TO3]] ], [ [[TMP37]], [[OMP_TYPE_FROM5]] ], [ [[TMP30]], [[OMP_TYPE_TO_ELSE4]] ] +// CHECK-NEXT: [[TMP38:%.*]] = and i64 [[TMP4]], 1036 +// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS8:%.*]] = or i64 [[OMP_MAPTYPE7]], [[TMP38]] +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS8]], ptr null) +// CHECK-NEXT: [[TMP39:%.*]] = add nuw i64 281474976711171, [[TMP20]] +// CHECK-NEXT: [[TMP40:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP41:%.*]] = icmp eq i64 [[TMP40]], 0 +// CHECK-NEXT: br i1 [[TMP41]], label [[OMP_TYPE_ALLOC9:%.*]], label [[OMP_TYPE_ALLOC_ELSE10:%.*]] +// CHECK: omp.type.alloc9: +// CHECK-NEXT: [[TMP42:%.*]] = and i64 [[TMP39]], -4 +// CHECK-NEXT: br label [[OMP_TYPE_END14:%.*]] +// CHECK: omp.type.alloc.else10: +// CHECK-NEXT: [[TMP43:%.*]] = icmp eq i64 [[TMP40]], 1 +// CHECK-NEXT: br i1 [[TMP43]], label [[OMP_TYPE_TO11:%.*]], label [[OMP_TYPE_TO_ELSE12:%.*]] +// CHECK: omp.type.to11: +// CHECK-NEXT: [[TMP44:%.*]] = and i64 [[TMP39]], -3 +// CHECK-NEXT: br label [[OMP_TYPE_END14]] +// CHECK: omp.type.to.else12: +// CHECK-NEXT: [[TMP45:%.*]] = icmp eq i64 [[TMP40]], 2 +// CHECK-NEXT: br i1 [[TMP45]], label [[OMP_TYPE_FROM13:%.*]], label [[OMP_TYPE_END14]] +// CHECK: omp.type.from13: +// CHECK-NEXT: [[TMP46:%.*]] = and i64 [[TMP39]], -2 +// CHECK-NEXT: br label [[OMP_TYPE_END14]] +// CHECK: omp.type.end14: +// CHECK-NEXT: [[OMP_MAPTYPE15:%.*]] = phi i64 [ [[TMP42]], [[OMP_TYPE_ALLOC9]] ], [ [[TMP44]], [[OMP_TYPE_TO11]] ], [ [[TMP46]], [[OMP_TYPE_FROM13]] ], [ [[TMP39]], [[OMP_TYPE_TO_ELSE12]] ] +// CHECK-NEXT: [[TMP47:%.*]] = and i64 [[TMP4]], 1036 +// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS16:%.*]] = or i64 [[OMP_MAPTYPE15]], [[TMP47]] +// CHECK-NEXT: call void @.omp_mapper._ZTS1C.default(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[F]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS16]], ptr null) #[[ATTR3]] +// CHECK-NEXT: [[TMP48:%.*]] = add nuw i64 281474976711171, [[TMP20]] +// CHECK-NEXT: [[TMP49:%.*]] = and i64 [[TMP4]], 3 +// CHECK-NEXT: [[TMP50:%.*]] = icmp eq i64 [[TMP49]], 0 +// CHECK-NEXT: br i1 [[TMP50]], label [[OMP_TYPE_ALLOC17:%.*]], label [[OMP_TYPE_ALLOC_ELSE18:%.*]] +// CHECK: omp.type.alloc17: +// CHECK-NEXT: [[TMP51:%.*]] = and i64 [[TMP48]], -4 +// CHECK-NEXT: br label [[OMP_TYPE_END22]] +// CHECK: omp.type.alloc.else18: +// CHECK-NEXT: [[TMP52:%.*]] = icmp eq i64 [[TMP49]], 1 +// CHECK-NEXT: br i1 [[TMP52]], label [[OMP_TYPE_TO19:%.*]], label [[OMP_TYPE_TO_ELSE20:%.*]] +// CHECK: omp.type.to19: +// CHECK-NEXT: [[TMP53:%.*]] = and i64 [[TMP48]], -3 +// CHECK-NEXT: br label [[OMP_TYPE_END22]] +// CHECK: omp.type.to.else20: +// CHECK-NEXT: [[TMP54:%.*]] = icmp eq i64 [[TMP49]], 2 +// CHECK-NEXT: br i1 [[TMP54]], label [[OMP_TYPE_FROM21:%.*]], label [[OMP_TYPE_END22]] +// CHECK: omp.type.from21: +// CHECK-NEXT: [[TMP55:%.*]] = and i64 [[TMP48]], -2 +// CHECK-NEXT: br label [[OMP_TYPE_END22]] +// CHECK: omp.type.end22: +// CHECK-NEXT: [[OMP_MAPTYPE23:%.*]] = phi i64 [ [[TMP51]], [[OMP_TYPE_ALLOC17]] ], [ [[TMP53]], [[OMP_TYPE_TO19]] ], [ [[TMP55]], [[OMP_TYPE_FROM21]] ], [ [[TMP48]], [[OMP_TYPE_TO_ELSE20]] ] +// CHECK-NEXT: [[TMP56:%.*]] = and i64 [[TMP4]], 1036 +// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS24:%.*]] = or i64 [[OMP_MAPTYPE23]], [[TMP56]] +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[H]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS24]], 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: [[TMP53:%.*]] = and i64 [[TMP4]], 8 -// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP53]], 0 -// CHECK-NEXT: [[TMP54:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY22]], [[DOTOMP_ARRAY__DEL__DELETE]] -// CHECK-NEXT: br i1 [[TMP54]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] +// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY25:%.*]] = 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_ISARRAY25]], [[DOTOMP_ARRAY__DEL__DELETE]] +// CHECK-NEXT: br i1 [[TMP58]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] // CHECK: .omp.array..del: -// CHECK-NEXT: [[TMP55:%.*]] = mul nuw i64 [[TMP6]], 12 -// CHECK-NEXT: [[TMP56:%.*]] = and i64 [[TMP4]], -4 -// CHECK-NEXT: [[TMP57:%.*]] = or i64 [[TMP56]], 512 -// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP55]], i64 [[TMP57]], ptr [[TMP5]]) +// 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 @@ -300,21 +308,23 @@ void foo() { // CHECK-NEXT: br label [[OMP_TYPE_END]] // CHECK: omp.type.end: // 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: [[TMP25:%.*]] = and i64 [[TMP4]], 1036 +// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS:%.*]] = or i64 [[OMP_MAPTYPE]], [[TMP25]] +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[A]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS]], 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: [[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-NEXT: [[TMP26:%.*]] = and i64 [[TMP4]], 8 +// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP26]], 0 +// CHECK-NEXT: [[TMP27:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY1]], [[DOTOMP_ARRAY__DEL__DELETE]] +// CHECK-NEXT: br i1 [[TMP27]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] // CHECK: .omp.array..del: -// 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: [[TMP28:%.*]] = mul nuw i64 [[TMP6]], 4 +// CHECK-NEXT: [[TMP29:%.*]] = and i64 [[TMP4]], -4 +// CHECK-NEXT: [[TMP30:%.*]] = or i64 [[TMP29]], 512 +// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP28]], i64 [[TMP30]], ptr [[TMP5]]) // CHECK-NEXT: br label [[OMP_DONE]] // CHECK: omp.done: // CHECK-NEXT: ret void diff --git a/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp b/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp index cfdc9b861090c..bd9189dce7c9f 100644 --- a/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp +++ b/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp @@ -98,7 +98,7 @@ void foo(S2 *arr) { // CHECK: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP7]] // CHECK: br i1 [[OMP_ARRAYMAP_ISEMPTY]], label [[OMP_DONE:%.*]], label [[OMP_ARRAYMAP_BODY:%.*]] // CHECK: omp.arraymap.body: -// CHECK: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END23:%.*]] ] +// CHECK: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END25:%.*]] ] // CHECK: [[Z:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 1 // CHECK: [[S1P:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0 // CHECK: [[S1P1:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0 @@ -135,87 +135,95 @@ void foo(S2 *arr) { // CHECK: br label [[OMP_TYPE_END]] // CHECK: omp.type.end: // CHECK: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP26]], [[OMP_TYPE_ALLOC]] ], [ [[TMP28]], [[OMP_TYPE_TO]] ], [ [[TMP30]], [[OMP_TYPE_FROM]] ], [ [[TMP23]], [[OMP_TYPE_TO_ELSE]] ] -// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[S1P]], i64 [[TMP20]], i64 [[OMP_MAPTYPE]], ptr null) -// CHECK: [[TMP31:%.*]] = add nuw i64 281474976710659, [[TMP22]] -// CHECK: [[TMP32:%.*]] = and i64 [[TMP4]], 3 -// CHECK: [[TMP33:%.*]] = icmp eq i64 [[TMP32]], 0 -// CHECK: br i1 [[TMP33]], label [[OMP_TYPE_ALLOC4:%.*]], label [[OMP_TYPE_ALLOC_ELSE5:%.*]] +// CHECK: [[TMP31:%.*]] = and i64 [[TMP4]], 1036 +// CHECK: [[OMP_MAPTYPE_WITH_MODIFIERS:%.*]] = or i64 [[OMP_MAPTYPE]], [[TMP31]] +// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[S1P]], i64 [[TMP20]], i64 [[OMP_MAPTYPE_WITH_MODIFIERS]], ptr null) +// CHECK: [[TMP32:%.*]] = add nuw i64 281474976710659, [[TMP22]] +// CHECK: [[TMP33:%.*]] = and i64 [[TMP4]], 3 +// CHECK: [[TMP34:%.*]] = icmp eq i64 [[TMP33]], 0 +// CHECK: br i1 [[TMP34]], label [[OMP_TYPE_ALLOC4:%.*]], label [[OMP_TYPE_ALLOC_ELSE5:%.*]] // CHECK: omp.type.alloc4: -// CHECK: [[TMP34:%.*]] = and i64 [[TMP31]], -4 +// CHECK: [[TMP35:%.*]] = and i64 [[TMP32]], -4 // CHECK: br label [[OMP_TYPE_END9:%.*]] // CHECK: omp.type.alloc.else5: -// CHECK: [[TMP35:%.*]] = icmp eq i64 [[TMP32]], 1 -// CHECK: br i1 [[TMP35]], label [[OMP_TYPE_TO6:%.*]], label [[OMP_TYPE_TO_ELSE7:%.*]] +// CHECK: [[TMP36:%.*]] = icmp eq i64 [[TMP33]], 1 +// CHECK: br i1 [[TMP36]], label [[OMP_TYPE_TO6:%.*]], label [[OMP_TYPE_TO_ELSE7:%.*]] // CHECK: omp.type.to6: -// CHECK: [[TMP36:%.*]] = and i64 [[TMP31]], -3 +// CHECK: [[TMP37:%.*]] = and i64 [[TMP32]], -3 // CHECK: br label [[OMP_TYPE_END9]] // CHECK: omp.type.to.else7: -// CHECK: [[TMP37:%.*]] = icmp eq i64 [[TMP32]], 2 -// CHECK: br i1 [[TMP37]], label [[OMP_TYPE_FROM8:%.*]], label [[OMP_TYPE_END9]] +// CHECK: [[TMP38:%.*]] = icmp eq i64 [[TMP33]], 2 +// CHECK: br i1 [[TMP38]], label [[OMP_TYPE_FROM8:%.*]], label [[OMP_TYPE_END9]] // CHECK: omp.type.from8: -// CHECK: [[TMP38:%.*]] = and i64 [[TMP31]], -2 +// CHECK: [[TMP39:%.*]] = and i64 [[TMP32]], -2 // CHECK: br label [[OMP_TYPE_END9]] // CHECK: omp.type.end9: -// CHECK: [[OMP_MAPTYPE10:%.*]] = phi i64 [ [[TMP34]], [[OMP_TYPE_ALLOC4]] ], [ [[TMP36]], [[OMP_TYPE_TO6]] ], [ [[TMP38]], [[OMP_TYPE_FROM8]] ], [ [[TMP31]], [[OMP_TYPE_TO_ELSE7]] ] -// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[Z]], i64 4, i64 [[OMP_MAPTYPE10]], ptr null) -// CHECK: [[TMP39:%.*]] = add nuw i64 281474976710675, [[TMP22]] -// CHECK: [[TMP40:%.*]] = and i64 [[TMP4]], 3 -// CHECK: [[TMP41:%.*]] = icmp eq i64 [[TMP40]], 0 -// CHECK: br i1 [[TMP41]], label [[OMP_TYPE_ALLOC11:%.*]], label [[OMP_TYPE_ALLOC_ELSE12:%.*]] -// CHECK: omp.type.alloc11: -// CHECK: [[TMP42:%.*]] = and i64 [[TMP39]], -4 -// CHECK: br label [[OMP_TYPE_END16:%.*]] -// CHECK: omp.type.alloc.else12: -// CHECK: [[TMP43:%.*]] = icmp eq i64 [[TMP40]], 1 -// CHECK: br i1 [[TMP43]], label [[OMP_TYPE_TO13:%.*]], label [[OMP_TYPE_TO_ELSE14:%.*]] -// CHECK: omp.type.to13: -// CHECK: [[TMP44:%.*]] = and i64 [[TMP39]], -3 -// CHECK: br label [[OMP_TYPE_END16]] -// CHECK: omp.type.to.else14: -// CHECK: [[TMP45:%.*]] = icmp eq i64 [[TMP40]], 2 -// CHECK: br i1 [[TMP45]], label [[OMP_TYPE_FROM15:%.*]], label [[OMP_TYPE_END16]] -// CHECK: omp.type.from15: -// CHECK: [[TMP46:%.*]] = and i64 [[TMP39]], -2 -// CHECK: br label [[OMP_TYPE_END16]] -// CHECK: omp.type.end16: -// CHECK: [[OMP_MAPTYPE17:%.*]] = phi i64 [ [[TMP42]], [[OMP_TYPE_ALLOC11]] ], [ [[TMP44]], [[OMP_TYPE_TO13]] ], [ [[TMP46]], [[OMP_TYPE_FROM15]] ], [ [[TMP39]], [[OMP_TYPE_TO_ELSE14]] ] -// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P]], ptr [[X]], i64 4, i64 [[OMP_MAPTYPE17]], ptr null) -// CHECK: [[TMP47:%.*]] = add nuw i64 281474976710675, [[TMP22]] -// CHECK: [[TMP48:%.*]] = and i64 [[TMP4]], 3 -// CHECK: [[TMP49:%.*]] = icmp eq i64 [[TMP48]], 0 -// CHECK: br i1 [[TMP49]], label [[OMP_TYPE_ALLOC18:%.*]], label [[OMP_TYPE_ALLOC_ELSE19:%.*]] -// CHECK: omp.type.alloc18: -// CHECK: [[TMP50:%.*]] = and i64 [[TMP47]], -4 -// CHECK: br label [[OMP_TYPE_END23]] -// CHECK: omp.type.alloc.else19: -// CHECK: [[TMP51:%.*]] = icmp eq i64 [[TMP48]], 1 -// CHECK: br i1 [[TMP51]], label [[OMP_TYPE_TO20:%.*]], label [[OMP_TYPE_TO_ELSE21:%.*]] -// CHECK: omp.type.to20: -// CHECK: [[TMP52:%.*]] = and i64 [[TMP47]], -3 -// CHECK: br label [[OMP_TYPE_END23]] -// CHECK: omp.type.to.else21: -// CHECK: [[TMP53:%.*]] = icmp eq i64 [[TMP48]], 2 -// CHECK: br i1 [[TMP53]], label [[OMP_TYPE_FROM22:%.*]], label [[OMP_TYPE_END23]] -// CHECK: omp.type.from22: -// CHECK: [[TMP54:%.*]] = and i64 [[TMP47]], -2 -// CHECK: br label [[OMP_TYPE_END23]] -// CHECK: omp.type.end23: -// CHECK: [[OMP_MAPTYPE24:%.*]] = phi i64 [ [[TMP50]], [[OMP_TYPE_ALLOC18]] ], [ [[TMP52]], [[OMP_TYPE_TO20]] ], [ [[TMP54]], [[OMP_TYPE_FROM22]] ], [ [[TMP47]], [[OMP_TYPE_TO_ELSE21]] ] -// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P2]], ptr [[Y]], i64 4, i64 [[OMP_MAPTYPE24]], ptr null) +// CHECK: [[OMP_MAPTYPE10:%.*]] = phi i64 [ [[TMP35]], [[OMP_TYPE_ALLOC4]] ], [ [[TMP37]], [[OMP_TYPE_TO6]] ], [ [[TMP39]], [[OMP_TYPE_FROM8]] ], [ [[TMP32]], [[OMP_TYPE_TO_ELSE7]] ] +// CHECK: [[TMP40:%.*]] = and i64 [[TMP4]], 1036 +// CHECK: [[OMP_MAPTYPE_WITH_MODIFIERS11:%.*]] = or i64 [[OMP_MAPTYPE10]], [[TMP40]] +// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[Z]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS11]], ptr null) +// CHECK: [[TMP41:%.*]] = add nuw i64 281474976710675, [[TMP22]] +// CHECK: [[TMP42:%.*]] = and i64 [[TMP4]], 3 +// CHECK: [[TMP43:%.*]] = icmp eq i64 [[TMP42]], 0 +// CHECK: br i1 [[TMP43]], label [[OMP_TYPE_ALLOC12:%.*]], label [[OMP_TYPE_ALLOC_ELSE13:%.*]] +// CHECK: omp.type.alloc12: +// CHECK: [[TMP44:%.*]] = and i64 [[TMP41]], -4 +// CHECK: br label [[OMP_TYPE_END17:%.*]] +// CHECK: omp.type.alloc.else13: +// CHECK: [[TMP45:%.*]] = icmp eq i64 [[TMP42]], 1 +// CHECK: br i1 [[TMP45]], label [[OMP_TYPE_TO14:%.*]], label [[OMP_TYPE_TO_ELSE15:%.*]] +// CHECK: omp.type.to14: +// CHECK: [[TMP46:%.*]] = and i64 [[TMP41]], -3 +// CHECK: br label [[OMP_TYPE_END17]] +// CHECK: omp.type.to.else15: +// CHECK: [[TMP47:%.*]] = icmp eq i64 [[TMP42]], 2 +// CHECK: br i1 [[TMP47]], label [[OMP_TYPE_FROM16:%.*]], label [[OMP_TYPE_END17]] +// CHECK: omp.type.from16: +// CHECK: [[TMP48:%.*]] = and i64 [[TMP41]], -2 +// CHECK: br label [[OMP_TYPE_END17]] +// CHECK: omp.type.end17: +// CHECK: [[OMP_MAPTYPE18:%.*]] = phi i64 [ [[TMP44]], [[OMP_TYPE_ALLOC12]] ], [ [[TMP46]], [[OMP_TYPE_TO14]] ], [ [[TMP48]], [[OMP_TYPE_FROM16]] ], [ [[TMP41]], [[OMP_TYPE_TO_ELSE15]] ] +// CHECK: [[TMP49:%.*]] = and i64 [[TMP4]], 1036 +// CHECK: [[OMP_MAPTYPE_WITH_MODIFIERS19:%.*]] = or i64 [[OMP_MAPTYPE18]], [[TMP49]] +// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P]], ptr [[X]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS19]], ptr null) +// CHECK: [[TMP50:%.*]] = add nuw i64 281474976710675, [[TMP22]] +// CHECK: [[TMP51:%.*]] = and i64 [[TMP4]], 3 +// CHECK: [[TMP52:%.*]] = icmp eq i64 [[TMP51]], 0 +// CHECK: br i1 [[TMP52]], label [[OMP_TYPE_ALLOC20:%.*]], label [[OMP_TYPE_ALLOC_ELSE21:%.*]] +// CHECK: omp.type.alloc20: +// CHECK: [[TMP53:%.*]] = and i64 [[TMP50]], -4 +// CHECK: br label [[OMP_TYPE_END25]] +// CHECK: omp.type.alloc.else21: +// CHECK: [[TMP54:%.*]] = icmp eq i64 [[TMP51]], 1 +// CHECK: br i1 [[TMP54]], label [[OMP_TYPE_TO22:%.*]], label [[OMP_TYPE_TO_ELSE23:%.*]] +// CHECK: omp.type.to22: +// CHECK: [[TMP55:%.*]] = and i64 [[TMP50]], -3 +// CHECK: br label [[OMP_TYPE_END25]] +// CHECK: omp.type.to.else23: +// CHECK: [[TMP56:%.*]] = icmp eq i64 [[TMP51]], 2 +// CHECK: br i1 [[TMP56]], label [[OMP_TYPE_FROM24:%.*]], label [[OMP_TYPE_END25]] +// CHECK: omp.type.from24: +// CHECK: [[TMP57:%.*]] = and i64 [[TMP50]], -2 +// CHECK: br label [[OMP_TYPE_END25]] +// CHECK: omp.type.end25: +// CHECK: [[OMP_MAPTYPE26:%.*]] = phi i64 [ [[TMP53]], [[OMP_TYPE_ALLOC20]] ], [ [[TMP55]], [[OMP_TYPE_TO22]] ], [ [[TMP57]], [[OMP_TYPE_FROM24]] ], [ [[TMP50]], [[OMP_TYPE_TO_ELSE23]] ] +// CHECK: [[TMP58:%.*]] = and i64 [[TMP4]], 1036 +// CHECK: [[OMP_MAPTYPE_WITH_MODIFIERS27:%.*]] = or i64 [[OMP_MAPTYPE26]], [[TMP58]] +// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P2]], ptr [[Y]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS27]], ptr null) // CHECK: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1 // CHECK: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP7]] // CHECK: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]] // CHECK: omp.arraymap.exit: -// CHECK: [[OMP_ARRAYINIT_ISARRAY25:%.*]] = icmp sgt i64 [[TMP6]], 1 -// CHECK: [[TMP55:%.*]] = and i64 [[TMP4]], 8 -// CHECK: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP55]], 0 -// CHECK: [[TMP56:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY25]], [[DOTOMP_ARRAY__DEL__DELETE]] -// CHECK: br i1 [[TMP56]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] +// CHECK: [[OMP_ARRAYINIT_ISARRAY28:%.*]] = icmp sgt i64 [[TMP6]], 1 +// CHECK: [[TMP59:%.*]] = and i64 [[TMP4]], 8 +// CHECK: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP59]], 0 +// CHECK: [[TMP60:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY28]], [[DOTOMP_ARRAY__DEL__DELETE]] +// CHECK: br i1 [[TMP60]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] // CHECK: .omp.array..del: -// CHECK: [[TMP57:%.*]] = mul nuw i64 [[TMP6]], 16 -// CHECK: [[TMP58:%.*]] = and i64 [[TMP4]], -4 -// CHECK: [[TMP59:%.*]] = or i64 [[TMP58]], 512 -// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP57]], i64 [[TMP59]], ptr [[TMP5]]) +// CHECK: [[TMP61:%.*]] = mul nuw i64 [[TMP6]], 16 +// CHECK: [[TMP62:%.*]] = and i64 [[TMP4]], -4 +// CHECK: [[TMP63:%.*]] = or i64 [[TMP62]], 512 +// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP61]], i64 [[TMP63]], ptr [[TMP5]]) // CHECK: br label [[OMP_DONE]] // CHECK: omp.done: // CHECK: ret void @@ -268,7 +276,7 @@ void foo(S2 *arr) { // CHECK-60: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP7]] // CHECK-60: br i1 [[OMP_ARRAYMAP_ISEMPTY]], label [[OMP_DONE:%.*]], label [[OMP_ARRAYMAP_BODY:%.*]] // CHECK-60: omp.arraymap.body: -// CHECK-60: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END23:%.*]] ] +// CHECK-60: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END25:%.*]] ] // CHECK-60: [[Z:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 1 // CHECK-60: [[S1P:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0 // CHECK-60: [[S1P1:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0 @@ -305,87 +313,95 @@ void foo(S2 *arr) { // CHECK-60: br label [[OMP_TYPE_END]] // CHECK-60: omp.type.end: // CHECK-60: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP26]], [[OMP_TYPE_ALLOC]] ], [ [[TMP28]], [[OMP_TYPE_TO]] ], [ [[TMP30]], [[OMP_TYPE_FROM]] ], [ [[TMP23]], [[OMP_TYPE_TO_ELSE]] ] -// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[S1P]], i64 [[TMP20]], i64 [[OMP_MAPTYPE]], ptr null) -// CHECK-60: [[TMP31:%.*]] = add nuw i64 281474976710659, [[TMP22]] -// CHECK-60: [[TMP32:%.*]] = and i64 [[TMP4]], 3 -// CHECK-60: [[TMP33:%.*]] = icmp eq i64 [[TMP32]], 0 -// CHECK-60: br i1 [[TMP33]], label [[OMP_TYPE_ALLOC4:%.*]], label [[OMP_TYPE_ALLOC_ELSE5:%.*]] +// CHECK-60: [[TMP31:%.*]] = and i64 [[TMP4]], 1036 +// CHECK-60: [[OMP_MAPTYPE_WITH_MODIFIERS:%.*]] = or i64 [[OMP_MAPTYPE]], [[TMP31]] +// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[S1P]], i64 [[TMP20]], i64 [[OMP_MAPTYPE_WITH_MODIFIERS]], ptr null) +// CHECK-60: [[TMP32:%.*]] = add nuw i64 281474976710659, [[TMP22]] +// CHECK-60: [[TMP33:%.*]] = and i64 [[TMP4]], 3 +// CHECK-60: [[TMP34:%.*]] = icmp eq i64 [[TMP33]], 0 +// CHECK-60: br i1 [[TMP34]], label [[OMP_TYPE_ALLOC4:%.*]], label [[OMP_TYPE_ALLOC_ELSE5:%.*]] // CHECK-60: omp.type.alloc4: -// CHECK-60: [[TMP34:%.*]] = and i64 [[TMP31]], -4 +// CHECK-60: [[TMP35:%.*]] = and i64 [[TMP32]], -4 // CHECK-60: br label [[OMP_TYPE_END9:%.*]] // CHECK-60: omp.type.alloc.else5: -// CHECK-60: [[TMP35:%.*]] = icmp eq i64 [[TMP32]], 1 -// CHECK-60: br i1 [[TMP35]], label [[OMP_TYPE_TO6:%.*]], label [[OMP_TYPE_TO_ELSE7:%.*]] +// CHECK-60: [[TMP36:%.*]] = icmp eq i64 [[TMP33]], 1 +// CHECK-60: br i1 [[TMP36]], label [[OMP_TYPE_TO6:%.*]], label [[OMP_TYPE_TO_ELSE7:%.*]] // CHECK-60: omp.type.to6: -// CHECK-60: [[TMP36:%.*]] = and i64 [[TMP31]], -3 +// CHECK-60: [[TMP37:%.*]] = and i64 [[TMP32]], -3 // CHECK-60: br label [[OMP_TYPE_END9]] // CHECK-60: omp.type.to.else7: -// CHECK-60: [[TMP37:%.*]] = icmp eq i64 [[TMP32]], 2 -// CHECK-60: br i1 [[TMP37]], label [[OMP_TYPE_FROM8:%.*]], label [[OMP_TYPE_END9]] +// CHECK-60: [[TMP38:%.*]] = icmp eq i64 [[TMP33]], 2 +// CHECK-60: br i1 [[TMP38]], label [[OMP_TYPE_FROM8:%.*]], label [[OMP_TYPE_END9]] // CHECK-60: omp.type.from8: -// CHECK-60: [[TMP38:%.*]] = and i64 [[TMP31]], -2 +// CHECK-60: [[TMP39:%.*]] = and i64 [[TMP32]], -2 // CHECK-60: br label [[OMP_TYPE_END9]] // CHECK-60: omp.type.end9: -// CHECK-60: [[OMP_MAPTYPE10:%.*]] = phi i64 [ [[TMP34]], [[OMP_TYPE_ALLOC4]] ], [ [[TMP36]], [[OMP_TYPE_TO6]] ], [ [[TMP38]], [[OMP_TYPE_FROM8]] ], [ [[TMP31]], [[OMP_TYPE_TO_ELSE7]] ] -// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[Z]], i64 4, i64 [[OMP_MAPTYPE10]], ptr null) -// CHECK-60: [[TMP39:%.*]] = add nuw i64 281474976710675, [[TMP22]] -// CHECK-60: [[TMP40:%.*]] = and i64 [[TMP4]], 3 -// CHECK-60: [[TMP41:%.*]] = icmp eq i64 [[TMP40]], 0 -// CHECK-60: br i1 [[TMP41]], label [[OMP_TYPE_ALLOC11:%.*]], label [[OMP_TYPE_ALLOC_ELSE12:%.*]] -// CHECK-60: omp.type.alloc11: -// CHECK-60: [[TMP42:%.*]] = and i64 [[TMP39]], -4 -// CHECK-60: br label [[OMP_TYPE_END16:%.*]] -// CHECK-60: omp.type.alloc.else12: -// CHECK-60: [[TMP43:%.*]] = icmp eq i64 [[TMP40]], 1 -// CHECK-60: br i1 [[TMP43]], label [[OMP_TYPE_TO13:%.*]], label [[OMP_TYPE_TO_ELSE14:%.*]] -// CHECK-60: omp.type.to13: -// CHECK-60: [[TMP44:%.*]] = and i64 [[TMP39]], -3 -// CHECK-60: br label [[OMP_TYPE_END16]] -// CHECK-60: omp.type.to.else14: -// CHECK-60: [[TMP45:%.*]] = icmp eq i64 [[TMP40]], 2 -// CHECK-60: br i1 [[TMP45]], label [[OMP_TYPE_FROM15:%.*]], label [[OMP_TYPE_END16]] -// CHECK-60: omp.type.from15: -// CHECK-60: [[TMP46:%.*]] = and i64 [[TMP39]], -2 -// CHECK-60: br label [[OMP_TYPE_END16]] -// CHECK-60: omp.type.end16: -// CHECK-60: [[OMP_MAPTYPE17:%.*]] = phi i64 [ [[TMP42]], [[OMP_TYPE_ALLOC11]] ], [ [[TMP44]], [[OMP_TYPE_TO13]] ], [ [[TMP46]], [[OMP_TYPE_FROM15]] ], [ [[TMP39]], [[OMP_TYPE_TO_ELSE14]] ] -// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P]], ptr [[X]], i64 4, i64 [[OMP_MAPTYPE17]], ptr null) -// CHECK-60: [[TMP47:%.*]] = add nuw i64 281474976710675, [[TMP22]] -// CHECK-60: [[TMP48:%.*]] = and i64 [[TMP4]], 3 -// CHECK-60: [[TMP49:%.*]] = icmp eq i64 [[TMP48]], 0 -// CHECK-60: br i1 [[TMP49]], label [[OMP_TYPE_ALLOC18:%.*]], label [[OMP_TYPE_ALLOC_ELSE19:%.*]] -// CHECK-60: omp.type.alloc18: -// CHECK-60: [[TMP50:%.*]] = and i64 [[TMP47]], -4 -// CHECK-60: br label [[OMP_TYPE_END23]] -// CHECK-60: omp.type.alloc.else19: -// CHECK-60: [[TMP51:%.*]] = icmp eq i64 [[TMP48]], 1 -// CHECK-60: br i1 [[TMP51]], label [[OMP_TYPE_TO20:%.*]], label [[OMP_TYPE_TO_ELSE21:%.*]] -// CHECK-60: omp.type.to20: -// CHECK-60: [[TMP52:%.*]] = and i64 [[TMP47]], -3 -// CHECK-60: br label [[OMP_TYPE_END23]] -// CHECK-60: omp.type.to.else21: -// CHECK-60: [[TMP53:%.*]] = icmp eq i64 [[TMP48]], 2 -// CHECK-60: br i1 [[TMP53]], label [[OMP_TYPE_FROM22:%.*]], label [[OMP_TYPE_END23]] -// CHECK-60: omp.type.from22: -// CHECK-60: [[TMP54:%.*]] = and i64 [[TMP47]], -2 -// CHECK-60: br label [[OMP_TYPE_END23]] -// CHECK-60: omp.type.end23: -// CHECK-60: [[OMP_MAPTYPE24:%.*]] = phi i64 [ [[TMP50]], [[OMP_TYPE_ALLOC18]] ], [ [[TMP52]], [[OMP_TYPE_TO20]] ], [ [[TMP54]], [[OMP_TYPE_FROM22]] ], [ [[TMP47]], [[OMP_TYPE_TO_ELSE21]] ] -// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P2]], ptr [[Y]], i64 4, i64 [[OMP_MAPTYPE24]], ptr null) +// CHECK-60: [[OMP_MAPTYPE10:%.*]] = phi i64 [ [[TMP35]], [[OMP_TYPE_ALLOC4]] ], [ [[TMP37]], [[OMP_TYPE_TO6]] ], [ [[TMP39]], [[OMP_TYPE_FROM8]] ], [ [[TMP32]], [[OMP_TYPE_TO_ELSE7]] ] +// CHECK-60: [[TMP40:%.*]] = and i64 [[TMP4]], 1036 +// CHECK-60: [[OMP_MAPTYPE_WITH_MODIFIERS11:%.*]] = or i64 [[OMP_MAPTYPE10]], [[TMP40]] +// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[Z]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS11]], ptr null) +// CHECK-60: [[TMP41:%.*]] = add nuw i64 281474976710675, [[TMP22]] +// CHECK-60: [[TMP42:%.*]] = and i64 [[TMP4]], 3 +// CHECK-60: [[TMP43:%.*]] = icmp eq i64 [[TMP42]], 0 +// CHECK-60: br i1 [[TMP43]], label [[OMP_TYPE_ALLOC12:%.*]], label [[OMP_TYPE_ALLOC_ELSE13:%.*]] +// CHECK-60: omp.type.alloc12: +// CHECK-60: [[TMP44:%.*]] = and i64 [[TMP41]], -4 +// CHECK-60: br label [[OMP_TYPE_END17:%.*]] +// CHECK-60: omp.type.alloc.else13: +// CHECK-60: [[TMP45:%.*]] = icmp eq i64 [[TMP42]], 1 +// CHECK-60: br i1 [[TMP45]], label [[OMP_TYPE_TO14:%.*]], label [[OMP_TYPE_TO_ELSE15:%.*]] +// CHECK-60: omp.type.to14: +// CHECK-60: [[TMP46:%.*]] = and i64 [[TMP41]], -3 +// CHECK-60: br label [[OMP_TYPE_END17]] +// CHECK-60: omp.type.to.else15: +// CHECK-60: [[TMP47:%.*]] = icmp eq i64 [[TMP42]], 2 +// CHECK-60: br i1 [[TMP47]], label [[OMP_TYPE_FROM16:%.*]], label [[OMP_TYPE_END17]] +// CHECK-60: omp.type.from16: +// CHECK-60: [[TMP48:%.*]] = and i64 [[TMP41]], -2 +// CHECK-60: br label [[OMP_TYPE_END17]] +// CHECK-60: omp.type.end17: +// CHECK-60: [[OMP_MAPTYPE18:%.*]] = phi i64 [ [[TMP44]], [[OMP_TYPE_ALLOC12]] ], [ [[TMP46]], [[OMP_TYPE_TO14]] ], [ [[TMP48]], [[OMP_TYPE_FROM16]] ], [ [[TMP41]], [[OMP_TYPE_TO_ELSE15]] ] +// CHECK-60: [[TMP49:%.*]] = and i64 [[TMP4]], 1036 +// CHECK-60: [[OMP_MAPTYPE_WITH_MODIFIERS19:%.*]] = or i64 [[OMP_MAPTYPE18]], [[TMP49]] +// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P]], ptr [[X]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS19]], ptr null) +// CHECK-60: [[TMP50:%.*]] = add nuw i64 281474976710675, [[TMP22]] +// CHECK-60: [[TMP51:%.*]] = and i64 [[TMP4]], 3 +// CHECK-60: [[TMP52:%.*]] = icmp eq i64 [[TMP51]], 0 +// CHECK-60: br i1 [[TMP52]], label [[OMP_TYPE_ALLOC20:%.*]], label [[OMP_TYPE_ALLOC_ELSE21:%.*]] +// CHECK-60: omp.type.alloc20: +// CHECK-60: [[TMP53:%.*]] = and i64 [[TMP50]], -4 +// CHECK-60: br label [[OMP_TYPE_END25]] +// CHECK-60: omp.type.alloc.else21: +// CHECK-60: [[TMP54:%.*]] = icmp eq i64 [[TMP51]], 1 +// CHECK-60: br i1 [[TMP54]], label [[OMP_TYPE_TO22:%.*]], label [[OMP_TYPE_TO_ELSE23:%.*]] +// CHECK-60: omp.type.to22: +// CHECK-60: [[TMP55:%.*]] = and i64 [[TMP50]], -3 +// CHECK-60: br label [[OMP_TYPE_END25]] +// CHECK-60: omp.type.to.else23: +// CHECK-60: [[TMP56:%.*]] = icmp eq i64 [[TMP51]], 2 +// CHECK-60: br i1 [[TMP56]], label [[OMP_TYPE_FROM24:%.*]], label [[OMP_TYPE_END25]] +// CHECK-60: omp.type.from24: +// CHECK-60: [[TMP57:%.*]] = and i64 [[TMP50]], -2 +// CHECK-60: br label [[OMP_TYPE_END25]] +// CHECK-60: omp.type.end25: +// CHECK-60: [[OMP_MAPTYPE26:%.*]] = phi i64 [ [[TMP53]], [[OMP_TYPE_ALLOC20]] ], [ [[TMP55]], [[OMP_TYPE_TO22]] ], [ [[TMP57]], [[OMP_TYPE_FROM24]] ], [ [[TMP50]], [[OMP_TYPE_TO_ELSE23]] ] +// CHECK-60: [[TMP58:%.*]] = and i64 [[TMP4]], 1036 +// CHECK-60: [[OMP_MAPTYPE_WITH_MODIFIERS27:%.*]] = or i64 [[OMP_MAPTYPE26]], [[TMP58]] +// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P2]], ptr [[Y]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS27]], ptr null) // CHECK-60: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1 // CHECK-60: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP7]] // CHECK-60: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]] // CHECK-60: omp.arraymap.exit: -// CHECK-60: [[OMP_ARRAYINIT_ISARRAY25:%.*]] = icmp sgt i64 [[TMP6]], 1 -// CHECK-60: [[TMP55:%.*]] = and i64 [[TMP4]], 8 -// CHECK-60: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP55]], 0 -// CHECK-60: [[TMP56:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY25]], [[DOTOMP_ARRAY__DEL__DELETE]] -// CHECK-60: br i1 [[TMP56]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] +// CHECK-60: [[OMP_ARRAYINIT_ISARRAY28:%.*]] = icmp sgt i64 [[TMP6]], 1 +// CHECK-60: [[TMP59:%.*]] = and i64 [[TMP4]], 8 +// CHECK-60: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP59]], 0 +// CHECK-60: [[TMP60:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY28]], [[DOTOMP_ARRAY__DEL__DELETE]] +// CHECK-60: br i1 [[TMP60]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] // CHECK-60: .omp.array..del: -// CHECK-60: [[TMP57:%.*]] = mul nuw i64 [[TMP6]], 16 -// CHECK-60: [[TMP58:%.*]] = and i64 [[TMP4]], -4 -// CHECK-60: [[TMP59:%.*]] = or i64 [[TMP58]], 512 -// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP57]], i64 [[TMP59]], ptr [[TMP5]]) +// CHECK-60: [[TMP61:%.*]] = mul nuw i64 [[TMP6]], 16 +// CHECK-60: [[TMP62:%.*]] = and i64 [[TMP4]], -4 +// CHECK-60: [[TMP63:%.*]] = or i64 [[TMP62]], 512 +// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP61]], i64 [[TMP63]], ptr [[TMP5]]) // CHECK-60: br label [[OMP_DONE]] // CHECK-60: omp.done: // CHECK-60: ret void diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h index 05307414c94d0..d4798f5a92dc6 100644 --- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h +++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h @@ -3627,14 +3627,20 @@ class OpenMPIRBuilder { /// for (unsigned i = 0; i < size; i++) { /// // For each component specified by this mapper: /// for (auto c : begin[i]->all_components) { + /// // Map-type-modifying bits (ALWAYS, DELETE, CLOSE) from the outer + /// // map clause are propagated to each component, except ATTACH + /// // entries (ATTACH|ALWAYS is reserved for attach(always), and other + /// // modifier bits have no meaning for ATTACH). + /// imported_modifier_bits = type & (ALWAYS | DELETE | CLOSE); + /// effective_type = c.isAttach() ? c.arg_type + /// : c.arg_type | imported_modifier_bits; /// if (c.hasMapper()) /// (*c.Mapper())(rt_mapper_handle, c.arg_base, c.arg_begin, - /// c.arg_size, - /// c.arg_type, c.arg_name); + /// c.arg_size, effective_type, c.arg_name); /// else /// __tgt_push_mapper_component(rt_mapper_handle, c.arg_base, - /// c.arg_begin, c.arg_size, c.arg_type, - /// c.arg_name); + /// c.arg_begin, c.arg_size, + /// effective_type, c.arg_name); /// } /// } /// // Delete the array section. diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index a227397b06f4e..9704f68df4ec7 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -10615,8 +10615,44 @@ Expected<Function *> OpenMPIRBuilder::emitUserDefinedMapper( CurMapType->addIncoming(FromMapType, FromBB); CurMapType->addIncoming(MemberMapType, ToElseBB); - Value *OffloadingArgs[] = {MapperHandle, CurBaseArg, CurBeginArg, - CurSizeArg, CurMapType, CurNameArg}; + // Propagate map-type-modifying bits from the outer map clause to each map + // inserted by the mapper. + // + // OpenMP 6.0:281:34: The effect of the mapper modifier is to remove the + // list item from the map clause and to apply the clauses specified in the + // declared mapper to the construct on which the map clause appears... + // If any modifier with the map-type-modifying property appears in the map + // clause then the effect is as if that modifier appears in each map clause + // specified in the declared mapper. + // + // Map-type-modifying bits: ALWAYS, DELETE, CLOSE, PRESENT. + // TODO: PRESENT is not propagated here yet. Doing so requires + // distinguishing pointee entries from the struct's own storage; it is + // handled in a follow-on. + Value *ImportedModifierBits = Builder.CreateAnd( + MapType, + Builder.getInt64( + static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>( + OpenMPOffloadMappingFlags::OMP_MAP_ALWAYS | + OpenMPOffloadMappingFlags::OMP_MAP_DELETE | + OpenMPOffloadMappingFlags::OMP_MAP_CLOSE))); + Value *CurMapTypeWithModifiers = Builder.CreateOr( + CurMapType, ImportedModifierBits, "omp.maptype.with.modifiers"); + + // ATTACH entries must not receive map-type-modifying bits: ATTACH|ALWAYS is + // reserved for the attach(always) map-type modifier, and other modifier + // bits (DELETE, CLOSE) have no meaning for an ATTACH entry. + auto RawType = + static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>( + Info->Types[I]); + constexpr uint64_t AttachBit = + static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>( + OpenMPOffloadMappingFlags::OMP_MAP_ATTACH); + Value *FinalMapType = + (RawType & AttachBit) ? CurMapType : CurMapTypeWithModifiers; + + Value *OffloadingArgs[] = {MapperHandle, CurBaseArg, CurBeginArg, + CurSizeArg, FinalMapType, CurNameArg}; auto ChildMapperFn = CustomMapperCB(I); if (!ChildMapperFn) diff --git a/mlir/test/Target/LLVMIR/omptarget-llvm.mlir b/mlir/test/Target/LLVMIR/omptarget-llvm.mlir index bb6085840cf07..baf3404e468ab 100644 --- a/mlir/test/Target/LLVMIR/omptarget-llvm.mlir +++ b/mlir/test/Target/LLVMIR/omptarget-llvm.mlir @@ -599,7 +599,11 @@ module attributes {omp.target_triples = ["amdgcn-amd-amdhsa"]} { // CHECK: br label %[[VAL_42]] // CHECK: omp.type.end: ; preds = %[[VAL_59]], %[[VAL_56]], %[[VAL_55]], %[[VAL_51]] // CHECK: %[[VAL_61:.*]] = phi i64 [ 0, %[[VAL_51]] ], [ 1, %[[VAL_55]] ], [ 2, %[[VAL_59]] ], [ 3, %[[VAL_56]] ] -// CHECK: call void @__tgt_push_mapper_component(ptr %[[VAL_37]], ptr %[[VAL_45]], ptr %[[VAL_45]], i64 4, i64 %[[VAL_61]], ptr @2) +// Map-type modifiers (ALWAYS|DELETE|CLOSE = 1036) from the outer clause are +// propagated to each entry the mapper pushes. +// CHECK: %[[VAL_MODMASK:.*]] = and i64 %{{.*}}, 1036 +// CHECK: %[[VAL_MTYPEMOD:.*]] = or i64 %[[VAL_61]], %[[VAL_MODMASK]] +// CHECK: call void @__tgt_push_mapper_component(ptr %[[VAL_37]], ptr %[[VAL_45]], ptr %[[VAL_45]], i64 4, i64 %[[VAL_MTYPEMOD]], ptr @2) // CHECK: %[[VAL_44]] = getelementptr %[[VAL_18]], ptr %[[VAL_43]], i32 1 // CHECK: %[[VAL_62:.*]] = icmp eq ptr %[[VAL_44]], %[[VAL_17]] // CHECK: br i1 %[[VAL_62]], label %[[VAL_63:.*]], label %[[VAL_41]] diff --git a/offload/test/mapping/mapper_map_always_from.c b/offload/test/mapping/mapper_map_always_from.c index 848bdbb448458..1048b32e990ef 100644 --- a/offload/test/mapping/mapper_map_always_from.c +++ b/offload/test/mapping/mapper_map_always_from.c @@ -1,17 +1,12 @@ -// Show that the ALWAYS map-type modifier on the outer map clause should be -// propagated to the entries pushed by a user-defined mapper. +// Show that the ALWAYS map-type modifier on the outer map clause is propagated +// to the entries pushed by a user-defined mapper. // // The mapper transfers s.y. We pre-map s.y so that on the target region below // it already has a device copy with a nonzero reference count. Without ALWAYS, -// the `from` at the end of the target region is suppressed (a present, -// ref-counted entry is not copied back), so the write of 111 is lost. ALWAYS -// must force the copy-back -- but only if ALWAYS is propagated from the outer -// clause to the mapper's s.y entry. -// -// FIXME: ALWAYS is not propagated to the mapper's entries yet, so the copy-back -// is currently suppressed and s.y reads back as 0 (its pre-map value). Once -// ALWAYS is propagated: -// EXPECTED: s.y = 111 +// the `from` at the end of the target region would be suppressed (a present, +// ref-counted entry is not copied back), so the write of 111 would be lost. +// ALWAYS forces the copy-back, which only happens if ALWAYS is propagated from +// the outer clause to the mapper's s.y entry. // RUN: %libomptarget-compile-run-and-check-generic @@ -37,10 +32,8 @@ int main() { s.y = 111; } - // ALWAYS should force s.y back even though it is still mapped (ref count > - // 0), but the modifier is not propagated yet, so the copy-back does not - // happen. - printf("s.y = %d\n", s.y); // CHECK: s.y = 0 + // ALWAYS forces s.y back even though it is still mapped (ref count > 0). + printf("s.y = %d\n", s.y); // CHECK: s.y = 111 #pragma omp target exit data map(delete : s.y) } diff --git a/offload/test/mapping/mapper_map_always_to_enter_data.c b/offload/test/mapping/mapper_map_always_to_enter_data.c index 185cc76a75246..b1742d9d14ff5 100644 --- a/offload/test/mapping/mapper_map_always_to_enter_data.c +++ b/offload/test/mapping/mapper_map_always_to_enter_data.c @@ -1,20 +1,15 @@ -// Show that the ALWAYS map-type modifier on the outer map clause should be -// propagated to the entries pushed by a user-defined mapper, using no target -// construct at all: the device data is inspected with omp_get_mapped_ptr() and +// Show that the ALWAYS map-type modifier on the outer map clause is propagated +// to the entries pushed by a user-defined mapper, using no target construct at +// all: the device data is inspected with omp_get_mapped_ptr() and // omp_target_memcpy(), so what is checked is purely the data-motion done by // `target enter data`. // // The mapper transfers s.y. We pre-map s.y so that it already has a device copy // with a nonzero reference count, and set that copy to a known value. Without -// ALWAYS, the `to` of the second `enter data` is suppressed (a present, -// ref-counted entry is not copied), so the host's 111 does not reach the -// device. ALWAYS must force the copy -- but only if ALWAYS is propagated from -// the outer clause to the mapper's s.y entry. -// -// FIXME: ALWAYS is not propagated to the mapper's entries yet, so the transfer -// is currently suppressed and the device copy still reads as 0 (its pre-set -// value). Once ALWAYS is propagated: -// EXPECTED: device s.y = 111 +// ALWAYS, the `to` of the second `enter data` would be suppressed (a present, +// ref-counted entry is not copied), so the host's 111 would not reach the +// device. ALWAYS forces the copy, which only happens if ALWAYS is propagated +// from the outer clause to the mapper's s.y entry. // RUN: %libomptarget-compile-run-and-check-generic @@ -48,13 +43,12 @@ int main() { #pragma omp target enter data map(always, to : s) - // ALWAYS should force s.y to the device even though it is already mapped - // (ref count > 0), but the modifier is not propagated yet, so the transfer - // does not happen. + // ALWAYS forces s.y to the device even though it is already mapped + // (ref count > 0). int dev_y_val = -1; omp_target_memcpy(&dev_y_val, dev_y, sizeof(int), 0, 0, host, dev); - printf("device s.y = %d\n", dev_y_val); // CHECK: device s.y = 0 + printf("device s.y = %d\n", dev_y_val); // CHECK: device s.y = 111 #pragma omp target exit data map(delete : s.y) } diff --git a/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c b/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c index fb03accec0083..1b91d87071cbd 100644 --- a/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c +++ b/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c @@ -54,8 +54,6 @@ int main() { print_status(&s2.s1p->y, "y"); // CHECK: y is not present print_status(&s2.z, "z"); // CHECK: z is not present print_status(&s2.s1p->dummy, "dummy"); // CHECK: dummy is not present - print_status(&s2.s1p->p, "p"); // CHECK: p is present - // FIXME: the DELETE modifier is not propagated to mapper entries yet. - // EXPECTED: p is not present - print_status(&s2.s1p->p[0], "p[0]"); // CHECK: p[0] is not present + print_status(&s2.s1p->p, "p"); // CHECK: p is not present + print_status(&s2.s1p->p[0], "p[0]"); // CHECK: p[0] is not present } diff --git a/offload/test/mapping/mapper_map_ptee_only_always_array.c b/offload/test/mapping/mapper_map_ptee_only_always_array.c index 64aaabd80dfe2..48e2a492131d3 100644 --- a/offload/test/mapping/mapper_map_ptee_only_always_array.c +++ b/offload/test/mapping/mapper_map_ptee_only_always_array.c @@ -49,18 +49,10 @@ int main() { printf("\n"); printf("After map(always,from)\n"); - printf("s[0].x = %d\n", s1[0].x); // CHECK: s[0].x = 111 - // FIXME: modifier-bits are not propagated to mapper entries yet. - // EXPECTED: s[0].x = 222 - printf("s[1].x = %d\n", s1[1].x); // CHECK: s[1].x = 111 - // FIXME: modifier-bits are not propagated to mapper entries yet. - // EXPECTED: s[1].x = 222 - printf("s[0].p[0] = %d\n", s1[0].p[0]); // CHECK: s[0].p[0] = 111 - // FIXME: modifier-bits are not propagated to mapper entries yet. - // EXPECTED: s[0].p[0] = 222 - printf("s[1].p[0] = %d\n", s1[1].p[0]); // CHECK: s[1].p[0] = 0 - // FIXME: modifier-bits are not propagated to mapper entries yet. - // EXPECTED: s[1].p[0] = 222 + printf("s[0].x = %d\n", s1[0].x); // CHECK: s[0].x = 222 + printf("s[1].x = %d\n", s1[1].x); // CHECK: s[1].x = 222 + printf("s[0].p[0] = %d\n", s1[0].p[0]); // CHECK: s[0].p[0] = 222 + printf("s[1].p[0] = %d\n", s1[1].p[0]); // CHECK: s[1].p[0] = 222 printf("\n"); #pragma omp target exit data map(delete : s1) diff --git a/offload/test/offloading/fortran/mapper-map-always-to-enter-data.f90 b/offload/test/offloading/fortran/mapper-map-always-to-enter-data.f90 index 52e32b1d6ccad..1c2a67ad92d1a 100644 --- a/offload/test/offloading/fortran/mapper-map-always-to-enter-data.f90 +++ b/offload/test/offloading/fortran/mapper-map-always-to-enter-data.f90 @@ -1,20 +1,15 @@ -! Show that the ALWAYS map-type modifier on the outer map clause should be -! propagated to the entries pushed by a user-defined mapper. The test uses no -! target region at all: the device data is inspected with omp_get_mapped_ptr() -! and omp_target_memcpy(), so what is checked is purely the data-motion done by +! Show that the ALWAYS map-type modifier on the outer map clause is propagated +! to the entries pushed by a user-defined mapper. The test uses no target +! region at all: the device data is inspected with omp_get_mapped_ptr() and +! omp_target_memcpy(), so what is checked is purely the data-motion done by ! `target enter data`. ! ! The mapper transfers s%y. We pre-map s%y so that it already has a device copy ! with a nonzero reference count, and set that copy to a known value. Without -! ALWAYS, the `to` of the second `enter data` is suppressed (a present, -! ref-counted entry is not copied), so the host's 111 does not reach the device. -! ALWAYS must force the copy -- but only if ALWAYS is propagated from the outer -! clause to the mapper's s%y entry. -! -! FIXME: ALWAYS is not propagated to the mapper's entries yet, so the transfer -! is currently suppressed and the device copy still reads as 0 (its pre-set -! value). Once ALWAYS is propagated: -! EXPECTED: device s%y = 111 +! ALWAYS, the `to` of the second `enter data` would be suppressed (a present, +! ref-counted entry is not copied), so the host's 111 would not reach the +! device. ALWAYS forces the copy, which only happens if ALWAYS is propagated +! from the outer clause to the mapper's s%y entry. ! REQUIRES: flang @@ -56,16 +51,15 @@ program main !$omp target enter data map(always, to: s) - ! ALWAYS should force s%y to the device even though it is already mapped - ! (ref count > 0), but the modifier is not propagated yet, so the transfer - ! does not happen. + ! ALWAYS forces s%y to the device even though it is already mapped + ! (ref count > 0). dev_y_val = -1 rc = omp_target_memcpy(c_loc(dev_y_val), dev_y, & int(c_sizeof(dev_y_val), c_size_t), & 0_c_size_t, 0_c_size_t, host, dev) print *, "device s%y =", dev_y_val - ! CHECK: device s%y = 0 + ! CHECK: device s%y = 111 !$omp target exit data map(delete: s%y) end program main _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
