https://github.com/TIFitis created
https://github.com/llvm/llvm-project/pull/159680
When a parent is mapped as `alloc`, declare mapper’s child `to/from` flags are
dropped, so device updates aren’t copied back to host.
Fix - In OpenMPIRBuilder’s `emitUserDefinedMapper` alloc branch, clear
`TO/FROM` then re-OR child’s `TO/FROM` bits from the original map-type to
preserve copy semantics.
This fixes #156466.
>From d7cb68c2296535e8e2318bc45f78ae3c6ba60b95 Mon Sep 17 00:00:00 2001
From: Akash Banerjee <akash.baner...@amd.com>
Date: Fri, 19 Sep 2025 01:07:48 +0100
Subject: [PATCH] [OpenMP][MLIR] Preserve to/from under alloc in declare mapper
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit
When a parent is mapped as alloc, declare mapper’s child to/from flags are
dropped, so device updates aren’t copied back to host.
Fix - In OpenMPIRBuilder’s emitUserDefinedMapper alloc branch, clear TO/FROM
then re-OR child’s to/from bits from the original map-type to preserve copy
semantics.
This fixes #156466.
---
clang/test/OpenMP/declare_mapper_codegen.cpp | 31 +++++++-----
..._of_structs_with_nested_mapper_codegen.cpp | 23 +++++----
..._of_structs_with_nested_mapper_codegen.cpp | 21 +++++---
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 10 +++-
mlir/test/Target/LLVMIR/omptarget-llvm.mlir | 3 +-
...mapper_alloc_parent_tofrom_propagation.cpp | 49 +++++++++++++++++++
.../target-declare-mapper-allocatable.f90 | 48 ++++++++++++++++++
7 files changed, 155 insertions(+), 30 deletions(-)
create mode 100644
offload/test/offloading/declare_mapper_alloc_parent_tofrom_propagation.cpp
create mode 100644
offload/test/offloading/fortran/target-declare-mapper-allocatable.f90
diff --git a/clang/test/OpenMP/declare_mapper_codegen.cpp
b/clang/test/OpenMP/declare_mapper_codegen.cpp
index 7dc32d0ae12ff..179de14c0c3c2 100644
--- a/clang/test/OpenMP/declare_mapper_codegen.cpp
+++ b/clang/test/OpenMP/declare_mapper_codegen.cpp
@@ -133,6 +133,7 @@ class C {
// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label
%[[ALLOCELSE:[^,]+]]
// CK0-DAG: [[ALLOC]]
// CK0-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
+// CK0-DAG: [[ALLOCTYPE_TF0:%.+]] = or i64 [[ALLOCTYPE]], {{[-]?[0-9]+}}
// CK0-DAG: br label %[[TYEND:[^,]+]]
// CK0-DAG: [[ALLOCELSE]]
// CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
@@ -147,7 +148,7 @@ class C {
// CK0-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
// CK0-DAG: br label %[[TYEND]]
// CK0-DAG: [[TYEND]]
-// CK0-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [
[[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]],
%[[TOELSE]] ]
+// CK0-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE_TF0]], %[[ALLOC]] ], [
[[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]],
%[[TOELSE]] ]
// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]],
ptr [[ABEGIN]], i64 [[CUSIZE]], i64 [[PHITYPE0]], {{.*}})
// 281474976710659 == 0x1,000,000,003
// CK0-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]]
@@ -156,6 +157,7 @@ class C {
// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label
%[[ALLOCELSE:[^,]+]]
// CK0-DAG: [[ALLOC]]
// CK0-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
+// CK0-DAG: [[ALLOCTYPE_TF1:%.+]] = or i64 [[ALLOCTYPE]], {{[-]?[0-9]+}}
// CK0-DAG: br label %[[TYEND:[^,]+]]
// CK0-DAG: [[ALLOCELSE]]
// CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
@@ -170,8 +172,8 @@ class C {
// CK0-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
// 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: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE_TF1]], %[[ALLOC]] ], [
[[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]],
%[[TOELSE]] ]
+// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]],
ptr [[ABEGIN]], i64 4, i64 %{{.*}}, {{.*}})
// 281474976710675 == 0x1,000,000,013
// CK0-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]]
// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
@@ -179,6 +181,7 @@ class C {
// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label
%[[ALLOCELSE:[^,]+]]
// CK0-DAG: [[ALLOC]]
// CK0-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
+// CK0-DAG: [[ALLOCTYPE_TF2:%.+]] = or i64 [[ALLOCTYPE]], {{[-]?[0-9]+}}
// CK0-DAG: br label %[[TYEND:[^,]+]]
// CK0-DAG: [[ALLOCELSE]]
// CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
@@ -193,8 +196,8 @@ class C {
// CK0-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
// 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: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE_TF2]], %[[ALLOC]] ], [
[[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]],
%[[TOELSE]] ]
+// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BBEGIN]],
ptr [[BARRBEGINGEP]], i64 16, i64 %{{.*}}, {{.*}})
// CK0: [[PTRNEXT]] = getelementptr %class.C, ptr [[PTR]], i32 1
// CK0: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]]
// CK0: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
@@ -624,6 +627,7 @@ class C {
// CK1-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label
%[[ALLOCELSE:[^,]+]]
// CK1-DAG: [[ALLOC]]
// CK1-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
+// CK1-DAG: [[ALLOCTYPE_TF:%.+]] = or i64 [[ALLOCTYPE]], {{[-]?[0-9]+}}
// CK1-DAG: br label %[[TYEND:[^,]+]]
// CK1-DAG: [[ALLOCELSE]]
// CK1-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
@@ -638,7 +642,7 @@ class C {
// CK1-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
// CK1-DAG: br label %[[TYEND]]
// CK1-DAG: [[TYEND]]
-// CK1-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [
[[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]],
%[[TOELSE]] ]
+// CK1-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE_TF]], %[[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: [[PTRNEXT]] = getelementptr %class.C, ptr [[PTR]], i32 1
// CK1: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]]
@@ -734,6 +738,7 @@ class C {
// CK2-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label
%[[ALLOCELSE:[^,]+]]
// CK2-DAG: [[ALLOC]]
// CK2-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
+// CK2-DAG: [[ALLOCTYPE_TF:%.+]] = or i64 [[ALLOCTYPE]], {{[-]?[0-9]+}}
// CK2-DAG: br label %[[TYEND:[^,]+]]
// CK2-DAG: [[ALLOCELSE]]
// CK2-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
@@ -748,7 +753,7 @@ class C {
// CK2-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
// CK2-DAG: br label %[[TYEND]]
// CK2-DAG: [[TYEND]]
-// CK2-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [
[[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]],
%[[TOELSE]] ]
+// CK2-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE_TF]], %[[ALLOC]] ], [
[[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]],
%[[TOELSE]] ]
// CK2: call void [[BMPRFUNC]](ptr [[HANDLE]], ptr [[PTR]], ptr [[BBEGIN]],
i64 8, i64 [[TYPE1]], {{.*}})
// CK2: [[PTRNEXT]] = getelementptr %class.C, ptr [[PTR]], i32 1
// CK2: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]]
@@ -939,6 +944,7 @@ class C {
// CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label
%[[ALLOCELSE:[^,]+]]
// CK4-DAG: [[ALLOC]]
// CK4-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
+// CK4-DAG: [[ALLOCTYPE_TF0:%.+]] = or i64 [[ALLOCTYPE]], {{[-]?[0-9]+}}
// CK4-DAG: br label %[[TYEND:[^,]+]]
// CK4-DAG: [[ALLOCELSE]]
// CK4-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
@@ -953,7 +959,7 @@ class C {
// CK4-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
// CK4-DAG: br label %[[TYEND]]
// CK4-DAG: [[TYEND]]
-// CK4-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [
[[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]],
%[[TOELSE]] ]
+// CK4-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE_TF0]], %[[ALLOC]] ], [
[[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]],
%[[TOELSE]] ]
// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]],
ptr [[ABEGIN]], i64 [[CUSIZE]], i64 [[PHITYPE0]], {{.*}})
// 281474976710659 == 0x1,000,000,003
// CK4-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]]
@@ -962,6 +968,7 @@ class C {
// CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label
%[[ALLOCELSE:[^,]+]]
// CK4-DAG: [[ALLOC]]
// CK4-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
+// CK4-DAG: [[ALLOCTYPE_TF1:%.+]] = or i64 [[ALLOCTYPE]], {{[-]?[0-9]+}}
// CK4-DAG: br label %[[TYEND:[^,]+]]
// CK4-DAG: [[ALLOCELSE]]
// CK4-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
@@ -976,7 +983,7 @@ class C {
// CK4-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
// CK4-DAG: br label %[[TYEND]]
// CK4-DAG: [[TYEND]]
-// CK4-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [
[[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]],
%[[TOELSE]] ]
+// CK4-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE_TF1]], %[[ALLOC]] ], [
[[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]],
%[[TOELSE]] ]
// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]],
ptr [[ABEGIN]], i64 4, i64 [[TYPE1]], {{.*}})
// 281474976710675 == 0x1,000,000,013
// CK4-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]]
@@ -985,6 +992,7 @@ class C {
// CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label
%[[ALLOCELSE:[^,]+]]
// CK4-DAG: [[ALLOC]]
// CK4-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
+// CK4-DAG: [[ALLOCTYPE_TF2:%.+]] = or i64 [[ALLOCTYPE]], {{[-]?[0-9]+}}
// CK4-DAG: br label %[[TYEND:[^,]+]]
// CK4-DAG: [[ALLOCELSE]]
// CK4-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
@@ -999,7 +1007,7 @@ class C {
// CK4-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
// CK4-DAG: br label %[[TYEND]]
// CK4-DAG: [[TYEND]]
-// CK4-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [
[[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]],
%[[TOELSE]] ]
+// CK4-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE_TF2]], %[[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: [[PTRNEXT]] = getelementptr %class.C, ptr [[PTR]], i32 1
// CK4: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]]
@@ -1120,6 +1128,7 @@ void foo(){
// CK5-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label
%[[ALLOCELSE:[^,]+]]
// CK5-DAG: [[ALLOC]]
// CK5-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
+// CK5-DAG: [[ALLOCTYPE_TF:%.+]] = or i64 [[ALLOCTYPE]], {{[-]?[0-9]+}}
// CK5-DAG: br label %[[TYEND:[^,]+]]
// CK5-DAG: [[ALLOCELSE]]
// CK5-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
@@ -1134,7 +1143,7 @@ void foo(){
// CK5-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
// CK5-DAG: br label %[[TYEND]]
// CK5-DAG: [[TYEND]]
-// CK5-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [
[[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]],
%[[TOELSE]] ]
+// CK5-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE_TF]], %[[ALLOC]] ], [
[[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]],
%[[TOELSE]] ]
// CK5: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]],
ptr [[ABEGIN]], i64 {{.*}}, i64 [[TYPE1]], {{.*}})
// CK5: [[PTRNEXT]] = getelementptr %struct.myvec, ptr [[PTR]], i32 1
// CK5: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]]
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..457c60e30a156 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
@@ -150,7 +150,8 @@ void foo() {
// CHECK-NEXT: [[TMP27:%.*]] = icmp eq i64 [[TMP26]], 0
// CHECK-NEXT: br i1 [[TMP27]], label [[OMP_TYPE_ALLOC:%.*]], label
[[OMP_TYPE_ALLOC_ELSE:%.*]]
// CHECK: omp.type.alloc:
-// CHECK-NEXT: [[TMP28:%.*]] = and i64 [[TMP25]], -4
+// CHECK-NEXT: [[ALLOCTYPE0:%.*]] = and i64 [[TMP25]], -4
+// CHECK-NEXT: [[ALLOCTYPE0_TF:%.*]] = or i64 [[ALLOCTYPE0]],
[[TFIMM0:[0-9]+]]
// CHECK-NEXT: br label [[OMP_TYPE_END:%.*]]
// CHECK: omp.type.alloc.else:
// CHECK-NEXT: [[TMP29:%.*]] = icmp eq i64 [[TMP26]], 1
@@ -165,14 +166,15 @@ void foo() {
// CHECK-NEXT: [[TMP32:%.*]] = and i64 [[TMP25]], -2
// CHECK-NEXT: br label [[OMP_TYPE_END]]
// CHECK: omp.type.end:
-// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP28]],
[[OMP_TYPE_ALLOC]] ], [ [[TMP30]], [[OMP_TYPE_TO]] ], [ [[TMP32]],
[[OMP_TYPE_FROM]] ], [ [[TMP25]], [[OMP_TYPE_TO_ELSE]] ]
+// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[ALLOCTYPE0_TF]],
[[OMP_TYPE_ALLOC]] ], [ [[TMP30]], [[OMP_TYPE_TO]] ], [ [[TMP32]],
[[OMP_TYPE_FROM]] ], [ [[TMP25]], [[OMP_TYPE_TO_ELSE]] ]
// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr
[[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP22]], i64 [[OMP_MAPTYPE]], ptr
null)
// CHECK-NEXT: [[TMP33:%.*]] = add nuw i64 281474976711171, [[TMP24]]
// CHECK-NEXT: [[TMP34:%.*]] = and i64 [[TMP4]], 3
// CHECK-NEXT: [[TMP35:%.*]] = icmp eq i64 [[TMP34]], 0
// CHECK-NEXT: br i1 [[TMP35]], label [[OMP_TYPE_ALLOC1:%.*]], label
[[OMP_TYPE_ALLOC_ELSE2:%.*]]
// CHECK: omp.type.alloc1:
-// CHECK-NEXT: [[TMP36:%.*]] = and i64 [[TMP33]], -4
+// CHECK-NEXT: [[ALLOCTYPE1:%.*]] = and i64 [[TMP33]], -4
+// CHECK-NEXT: [[ALLOCTYPE1_TF:%.*]] = or i64 [[ALLOCTYPE1]],
[[TFIMM1:[0-9]+]]
// CHECK-NEXT: br label [[OMP_TYPE_END6:%.*]]
// CHECK: omp.type.alloc.else2:
// CHECK-NEXT: [[TMP37:%.*]] = icmp eq i64 [[TMP34]], 1
@@ -187,14 +189,15 @@ void foo() {
// CHECK-NEXT: [[TMP40:%.*]] = and i64 [[TMP33]], -2
// CHECK-NEXT: br label [[OMP_TYPE_END6]]
// CHECK: omp.type.end6:
-// CHECK-NEXT: [[OMP_MAPTYPE7:%.*]] = phi i64 [ [[TMP36]],
[[OMP_TYPE_ALLOC1]] ], [ [[TMP38]], [[OMP_TYPE_TO3]] ], [ [[TMP40]],
[[OMP_TYPE_FROM5]] ], [ [[TMP33]], [[OMP_TYPE_TO_ELSE4]] ]
+// CHECK-NEXT: [[OMP_MAPTYPE7:%.*]] = phi i64 [ [[ALLOCTYPE1_TF]],
[[OMP_TYPE_ALLOC1]] ], [ [[TMP38]], [[OMP_TYPE_TO3]] ], [ [[TMP40]],
[[OMP_TYPE_FROM5]] ], [ [[TMP33]], [[OMP_TYPE_TO_ELSE4]] ]
// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr
[[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 4, i64 [[OMP_MAPTYPE7]], ptr null)
// CHECK-NEXT: [[TMP41:%.*]] = add nuw i64 281474976711171, [[TMP24]]
// CHECK-NEXT: [[TMP42:%.*]] = and i64 [[TMP4]], 3
// CHECK-NEXT: [[TMP43:%.*]] = icmp eq i64 [[TMP42]], 0
// CHECK-NEXT: br i1 [[TMP43]], label [[OMP_TYPE_ALLOC8:%.*]], label
[[OMP_TYPE_ALLOC_ELSE9:%.*]]
// CHECK: omp.type.alloc8:
-// CHECK-NEXT: [[TMP44:%.*]] = and i64 [[TMP41]], -4
+// CHECK-NEXT: [[ALLOCTYPE2:%.*]] = and i64 [[TMP41]], -4
+// CHECK-NEXT: [[ALLOCTYPE2_TF:%.*]] = or i64 [[ALLOCTYPE2]],
[[TFIMM2:[0-9]+]]
// CHECK-NEXT: br label [[OMP_TYPE_END13:%.*]]
// CHECK: omp.type.alloc.else9:
// CHECK-NEXT: [[TMP45:%.*]] = icmp eq i64 [[TMP42]], 1
@@ -209,14 +212,15 @@ void foo() {
// CHECK-NEXT: [[TMP48:%.*]] = and i64 [[TMP41]], -2
// CHECK-NEXT: br label [[OMP_TYPE_END13]]
// CHECK: omp.type.end13:
-// CHECK-NEXT: [[OMP_MAPTYPE14:%.*]] = phi i64 [ [[TMP44]],
[[OMP_TYPE_ALLOC8]] ], [ [[TMP46]], [[OMP_TYPE_TO10]] ], [ [[TMP48]],
[[OMP_TYPE_FROM12]] ], [ [[TMP41]], [[OMP_TYPE_TO_ELSE11]] ]
+// CHECK-NEXT: [[OMP_MAPTYPE14:%.*]] = phi i64 [ [[ALLOCTYPE2_TF]],
[[OMP_TYPE_ALLOC8]] ], [ [[TMP46]], [[OMP_TYPE_TO10]] ], [ [[TMP48]],
[[OMP_TYPE_FROM12]] ], [ [[TMP41]], [[OMP_TYPE_TO_ELSE11]] ]
// CHECK-NEXT: call void @.omp_mapper._ZTS1C.default(ptr [[TMP0]], ptr
[[OMP_ARRAYMAP_PTRCURRENT]], ptr [[F]], i64 4, i64 [[OMP_MAPTYPE14]], ptr null)
#[[ATTR3]]
// CHECK-NEXT: [[TMP49:%.*]] = add nuw i64 281474976711171, [[TMP24]]
// CHECK-NEXT: [[TMP50:%.*]] = and i64 [[TMP4]], 3
// CHECK-NEXT: [[TMP51:%.*]] = icmp eq i64 [[TMP50]], 0
// CHECK-NEXT: br i1 [[TMP51]], label [[OMP_TYPE_ALLOC15:%.*]], label
[[OMP_TYPE_ALLOC_ELSE16:%.*]]
// CHECK: omp.type.alloc15:
-// CHECK-NEXT: [[TMP52:%.*]] = and i64 [[TMP49]], -4
+// CHECK-NEXT: [[ALLOCTYPE3:%.*]] = and i64 [[TMP49]], -4
+// CHECK-NEXT: [[ALLOCTYPE3_TF:%.*]] = or i64 [[ALLOCTYPE3]],
[[TFIMM3:[0-9]+]]
// CHECK-NEXT: br label [[OMP_TYPE_END20]]
// CHECK: omp.type.alloc.else16:
// CHECK-NEXT: [[TMP53:%.*]] = icmp eq i64 [[TMP50]], 1
@@ -231,7 +235,7 @@ void foo() {
// CHECK-NEXT: [[TMP56:%.*]] = and i64 [[TMP49]], -2
// CHECK-NEXT: br label [[OMP_TYPE_END20]]
// CHECK: omp.type.end20:
-// CHECK-NEXT: [[OMP_MAPTYPE21:%.*]] = phi i64 [ [[TMP52]],
[[OMP_TYPE_ALLOC15]] ], [ [[TMP54]], [[OMP_TYPE_TO17]] ], [ [[TMP56]],
[[OMP_TYPE_FROM19]] ], [ [[TMP49]], [[OMP_TYPE_TO_ELSE18]] ]
+// CHECK-NEXT: [[OMP_MAPTYPE21:%.*]] = phi i64 [ [[ALLOCTYPE3_TF]],
[[OMP_TYPE_ALLOC15]] ], [ [[TMP54]], [[OMP_TYPE_TO17]] ], [ [[TMP56]],
[[OMP_TYPE_FROM19]] ], [ [[TMP49]], [[OMP_TYPE_TO_ELSE18]] ]
// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr
[[OMP_ARRAYMAP_PTRCURRENT]], ptr [[H]], i64 4, i64 [[OMP_MAPTYPE21]], ptr null)
// CHECK-NEXT: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_D]], ptr
[[OMP_ARRAYMAP_PTRCURRENT]], i32 1
// CHECK-NEXT: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr
[[OMP_ARRAYMAP_NEXT]], [[TMP7]]
@@ -287,6 +291,7 @@ void foo() {
// CHECK-NEXT: br i1 [[TMP22]], label [[OMP_TYPE_ALLOC:%.*]], label
[[OMP_TYPE_ALLOC_ELSE:%.*]]
// CHECK: omp.type.alloc:
// CHECK-NEXT: [[TMP23:%.*]] = and i64 [[TMP20]], -4
+// CHECK-NEXT: [[TMP23_OR:%.*]] = or i64 [[TMP23]], [[TFIMM0:[0-9]+]]
// CHECK-NEXT: br label [[OMP_TYPE_END]]
// CHECK: omp.type.alloc.else:
// CHECK-NEXT: [[TMP24:%.*]] = icmp eq i64 [[TMP21]], 1
@@ -301,7 +306,7 @@ void foo() {
// CHECK-NEXT: [[TMP27:%.*]] = and i64 [[TMP20]], -2
// CHECK-NEXT: br label [[OMP_TYPE_END]]
// CHECK: omp.type.end:
-// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP23]],
[[OMP_TYPE_ALLOC]] ], [ [[TMP25]], [[OMP_TYPE_TO]] ], [ [[TMP27]],
[[OMP_TYPE_FROM]] ], [ [[TMP20]], [[OMP_TYPE_TO_ELSE]] ]
+// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP23_OR]],
[[OMP_TYPE_ALLOC]] ], [ [[TMP25]], [[OMP_TYPE_TO]] ], [ [[TMP27]],
[[OMP_TYPE_FROM]] ], [ [[TMP20]], [[OMP_TYPE_TO_ELSE]] ]
// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr
[[OMP_ARRAYMAP_PTRCURRENT]], ptr [[A]], i64 4, i64 [[OMP_MAPTYPE]], ptr null)
// CHECK-NEXT: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_C]], ptr
[[OMP_ARRAYMAP_PTRCURRENT]], i32 1
// CHECK-NEXT: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr
[[OMP_ARRAYMAP_NEXT]], [[TMP7]]
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..13c43cebc9401 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
@@ -147,7 +147,8 @@ void foo() {
// CHECK-NEXT: [[TMP33:%.*]] = icmp eq i64 [[TMP32]], 0
// CHECK-NEXT: br i1 [[TMP33]], label [[OMP_TYPE_ALLOC:%.*]], label
[[OMP_TYPE_ALLOC_ELSE:%.*]]
// CHECK: omp.type.alloc:
-// CHECK-NEXT: [[TMP34:%.*]] = and i64 [[TMP31]], -4
+// CHECK-NEXT: [[ALLOCTYPE0:%.*]] = and i64 [[TMP31]], -4
+// CHECK-NEXT: [[ALLOCTYPE0_TF:%.*]] = or i64 [[ALLOCTYPE0]],
[[TFIMM0:[0-9]+]]
// CHECK-NEXT: br label [[OMP_TYPE_END:%.*]]
// CHECK: omp.type.alloc.else:
// CHECK-NEXT: [[TMP35:%.*]] = icmp eq i64 [[TMP32]], 1
@@ -162,14 +163,15 @@ void foo() {
// CHECK-NEXT: [[TMP38:%.*]] = and i64 [[TMP31]], -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: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[ALLOCTYPE0_TF]],
[[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: omp.type.alloc1:
-// CHECK-NEXT: [[TMP42:%.*]] = and i64 [[TMP39]], -4
+// CHECK-NEXT: [[ALLOCTYPE1:%.*]] = and i64 [[TMP39]], -4
+// CHECK-NEXT: [[ALLOCTYPE1_TF:%.*]] = or i64 [[ALLOCTYPE1]],
[[TFIMM1:[0-9]+]]
// CHECK-NEXT: br label [[OMP_TYPE_END11:%.*]]
// CHECK: omp.type.alloc.else2:
// CHECK-NEXT: [[TMP43:%.*]] = icmp eq i64 [[TMP40]], 1
@@ -184,14 +186,15 @@ void foo() {
// CHECK-NEXT: [[TMP46:%.*]] = and i64 [[TMP39]], -2
// CHECK-NEXT: br label [[OMP_TYPE_END11]]
// 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: [[OMP_MAPTYPE12:%.*]] = phi i64 [ [[ALLOCTYPE1_TF]],
[[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: omp.type.alloc8:
-// CHECK-NEXT: [[TMP50:%.*]] = and i64 [[TMP47]], -4
+// CHECK-NEXT: [[ALLOCTYPE2:%.*]] = and i64 [[TMP47]], -4
+// CHECK-NEXT: [[ALLOCTYPE2_TF:%.*]] = or i64 [[ALLOCTYPE2]],
[[TFIMM2:[0-9]+]]
// CHECK-NEXT: br label [[OMP_TYPE_END18:%.*]]
// CHECK: omp.type.alloc.else9:
// CHECK-NEXT: [[TMP51:%.*]] = icmp eq i64 [[TMP48]], 1
@@ -206,7 +209,7 @@ void foo() {
// CHECK-NEXT: [[TMP54:%.*]] = and i64 [[TMP47]], -2
// CHECK-NEXT: br label [[OMP_TYPE_END18]]
// 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: [[OMP_MAPTYPE19:%.*]] = phi i64 [ [[ALLOCTYPE2_TF]],
[[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
@@ -214,6 +217,7 @@ void foo() {
// CHECK-NEXT: br i1 [[TMP57]], label [[OMP_TYPE_ALLOC20:%.*]], label
[[OMP_TYPE_ALLOC_ELSE21:%.*]]
// CHECK: omp.type.alloc15:
// CHECK-NEXT: [[TMP58:%.*]] = and i64 [[TMP55]], -4
+// CHECK-NEXT: [[TMP58_OR:%.*]] = or i64 [[TMP58]], [[TFIMM3:[0-9]+]]
// CHECK-NEXT: br label [[OMP_TYPE_END25]]
// CHECK: omp.type.alloc.else16:
// CHECK-NEXT: [[TMP59:%.*]] = icmp eq i64 [[TMP56]], 1
@@ -228,7 +232,7 @@ void foo() {
// CHECK-NEXT: [[TMP62:%.*]] = and i64 [[TMP55]], -2
// CHECK-NEXT: br label [[OMP_TYPE_END25]]
// 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: [[OMP_MAPTYPE26:%.*]] = phi i64 [ [[TMP58_OR]],
[[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_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_D]], ptr
[[OMP_ARRAYMAP_PTRCURRENT]], i32 1
// CHECK-NEXT: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr
[[OMP_ARRAYMAP_NEXT]], [[TMP11]]
@@ -284,6 +288,7 @@ void foo() {
// CHECK-NEXT: br i1 [[TMP28]], label [[OMP_TYPE_ALLOC:%.*]], label
[[OMP_TYPE_ALLOC_ELSE:%.*]]
// CHECK: omp.type.alloc:
// CHECK-NEXT: [[TMP29:%.*]] = and i64 [[TMP26]], -4
+// CHECK-NEXT: [[TMP29_OR:%.*]] = or i64 [[TMP29]], [[TFIMM4:[0-9]+]]
// CHECK-NEXT: br label [[OMP_TYPE_END]]
// CHECK: omp.type.alloc.else:
// CHECK-NEXT: [[TMP30:%.*]] = icmp eq i64 [[TMP27]], 1
@@ -298,7 +303,7 @@ void foo() {
// CHECK-NEXT: [[TMP33:%.*]] = and i64 [[TMP26]], -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 [ [[TMP29_OR]],
[[OMP_TYPE_ALLOC]] ], [ [[TMP31]], [[OMP_TYPE_TO]] ], [ [[TMP33]],
[[OMP_TYPE_FROM]] ], [ [[TMP26]], [[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]]
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 220eee3cb8b08..04c9c66d63757 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -8887,7 +8887,9 @@ Expected<Function *>
OpenMPIRBuilder::emitUserDefinedMapper(
BasicBlock *EndBB = BasicBlock::Create(M.getContext(), "omp.type.end");
Value *IsAlloc = Builder.CreateIsNull(LeftToFrom);
Builder.CreateCondBr(IsAlloc, AllocBB, AllocElseBB);
- // In case of alloc, clear OMP_MAP_TO and OMP_MAP_FROM.
+ // In case of alloc, clear OMP_MAP_TO and OMP_MAP_FROM, then re-OR any
+ // explicit child TO/FROM intent from the mapper-declared type to avoid
+ // losing copy semantics when the parent map-type is alloc (partial maps).
emitBlock(AllocBB, MapperFn);
Value *AllocMapType = Builder.CreateAnd(
MemberMapType,
@@ -8895,6 +8897,12 @@ Expected<Function *>
OpenMPIRBuilder::emitUserDefinedMapper(
~static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
OpenMPOffloadMappingFlags::OMP_MAP_TO |
OpenMPOffloadMappingFlags::OMP_MAP_FROM)));
+ Value *TFMaskAlloc = Builder.getInt64(
+ static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
+ OpenMPOffloadMappingFlags::OMP_MAP_TO |
+ OpenMPOffloadMappingFlags::OMP_MAP_FROM));
+ Value *ChildTFAlloc = Builder.CreateAnd(OriMapType, TFMaskAlloc);
+ AllocMapType = Builder.CreateOr(AllocMapType, ChildTFAlloc);
Builder.CreateBr(EndBB);
emitBlock(AllocElseBB, MapperFn);
Value *IsTo = Builder.CreateICmpEQ(
diff --git a/mlir/test/Target/LLVMIR/omptarget-llvm.mlir
b/mlir/test/Target/LLVMIR/omptarget-llvm.mlir
index e6ea3aaeec656..0b8d5e82eeb7e 100644
--- a/mlir/test/Target/LLVMIR/omptarget-llvm.mlir
+++ b/mlir/test/Target/LLVMIR/omptarget-llvm.mlir
@@ -589,6 +589,7 @@ module attributes {omp.target_triples =
["amdgcn-amd-amdhsa"]} {
// CHECK: br i1 %[[VAL_50]], label %[[VAL_51:.*]], label %[[VAL_52:.*]]
// CHECK: omp.type.alloc: ; preds =
%[[VAL_41]]
// CHECK: %[[VAL_53:.*]] = and i64 %[[VAL_48]], -4
+// CHECK: %[[VAL_53_OR:.*]] = or i64 %[[VAL_53]], {{[0-9]+}}
// CHECK: br label %[[VAL_42]]
// CHECK: omp.type.alloc.else: ; preds =
%[[VAL_41]]
// CHECK: %[[VAL_54:.*]] = icmp eq i64 %[[VAL_49]], 1
@@ -603,7 +604,7 @@ module attributes {omp.target_triples =
["amdgcn-amd-amdhsa"]} {
// CHECK: %[[VAL_60:.*]] = and i64 %[[VAL_48]], -2
// CHECK: br label %[[VAL_42]]
// CHECK: omp.type.end: ; preds =
%[[VAL_59]], %[[VAL_56]], %[[VAL_55]], %[[VAL_51]]
-// CHECK: %[[VAL_61:.*]] = phi i64 [ %[[VAL_53]], %[[VAL_51]] ], [
%[[VAL_57]], %[[VAL_55]] ], [ %[[VAL_60]], %[[VAL_59]] ], [ %[[VAL_48]],
%[[VAL_56]] ]
+// CHECK: %[[VAL_61:.*]] = phi i64 [ %[[VAL_53_OR]], %[[VAL_51]] ], [
%[[VAL_57]], %[[VAL_55]] ], [ %[[VAL_60]], %[[VAL_59]] ], [ %[[VAL_48]],
%[[VAL_56]] ]
// CHECK: call void @__tgt_push_mapper_component(ptr %[[VAL_37]], ptr
%[[VAL_43]], ptr %[[VAL_45]], i64 4, i64 %[[VAL_61]], ptr @2)
// CHECK: %[[VAL_44]] = getelementptr %[[VAL_18]], ptr %[[VAL_43]],
i32 1
// CHECK: %[[VAL_62:.*]] = icmp eq ptr %[[VAL_44]], %[[VAL_17]]
diff --git
a/offload/test/offloading/declare_mapper_alloc_parent_tofrom_propagation.cpp
b/offload/test/offloading/declare_mapper_alloc_parent_tofrom_propagation.cpp
new file mode 100644
index 0000000000000..4f224b74be945
--- /dev/null
+++ b/offload/test/offloading/declare_mapper_alloc_parent_tofrom_propagation.cpp
@@ -0,0 +1,49 @@
+// REQUIRES: amdgpu
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <cstdio>
+#include <cstdlib>
+
+struct vec {
+ int len;
+ int *data;
+};
+
+// Map the dynamic payload with tofrom semantics via a user-defined mapper.
+#pragma omp declare mapper(default : vec v) map(tofrom : v.data [0:v.len])
+
+int main() {
+ vec s{};
+ s.len = 16;
+ s.data = (int *)malloc(sizeof(int) * s.len);
+ for (int i = 0; i < s.len; ++i)
+ s.data[i] = 1;
+
+ // Offload with the mapper and update payload on device. Avoid reading s.len
+ // on device; use a firstprivate copy of the length.
+ int n = s.len;
+ // Intentionally map the struct itself with 'alloc'. The mapper specifies
+ // tofrom semantics for the payload. Without the fix that propagates mapper
+ // to/from into ALLOC branches for components, the device writes would not
+ // be copied back and this test would fail.
+#pragma omp target map(mapper(default), alloc : s) firstprivate(n)
+ {
+ for (int i = 0; i < n; ++i)
+ s.data[i] = 7;
+ }
+
+ long sum = 0;
+ for (int i = 0; i < s.len; ++i)
+ sum += s.data[i];
+
+ if (sum == 7L * s.len) {
+ std::printf("Test passed!\n");
+ } else {
+ std::printf("Test failed! sum=%ld\n", sum);
+ }
+
+ free(s.data);
+ return 0;
+}
+
+// CHECK: Test passed!
diff --git
a/offload/test/offloading/fortran/target-declare-mapper-allocatable.f90
b/offload/test/offloading/fortran/target-declare-mapper-allocatable.f90
new file mode 100644
index 0000000000000..d8d5e1b5631a5
--- /dev/null
+++ b/offload/test/offloading/fortran/target-declare-mapper-allocatable.f90
@@ -0,0 +1,48 @@
+! This test validates that declare mapper for a derived type with an
+! allocatable component preserves TO/FROM semantics for the component,
+! ensuring the payload is copied back to the host on target exit.
+
+! REQUIRES: flang, amdgpu
+
+! RUN: %libomptarget-compile-fortran-run-and-check-generic
+
+program target_declare_mapper_allocatable
+ implicit none
+
+ type :: real_t
+ real, allocatable :: real_arr(:)
+ end type real_t
+
+ ! Map the allocatable array payload via a named mapper.
+ !$omp declare mapper (xyz : real_t :: t) map(tofrom: t%real_arr)
+
+ type(real_t) :: r
+ integer :: i
+ logical :: ok
+
+ allocate(r%real_arr(10))
+ r%real_arr = 1.0
+
+ !$omp target map(mapper(xyz), tofrom: r)
+ do i = 1, size(r%real_arr)
+ r%real_arr(i) = 3.0
+ end do
+ !$omp end target
+
+ ok = .true.
+ do i = 1, size(r%real_arr)
+ if (r%real_arr(i) /= 3.0) ok = .false.
+ end do
+ if (ok) then
+ print *, "Test passed!"
+ else
+ print *, "Test failed!"
+ do i = 1, size(r%real_arr)
+ print *, r%real_arr(i)
+ end do
+ end if
+
+ deallocate(r%real_arr)
+end program target_declare_mapper_allocatable
+
+! CHECK: Test passed!
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
