[clang] [llvm] [NFC][OpenMP] Add mapper-specific tests exercising pointee section mapping. (PR #204269)
Abhinav Gaba via cfe-commits
cfe-commits at lists.llvm.org
Mon Jul 27 17:46:56 PDT 2026
https://github.com/abhinavgaba updated https://github.com/llvm/llvm-project/pull/204269
>From 48084622e5343e8dae1a38486bde92473cddcfa1 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Tue, 16 Jun 2026 16:21:12 -0700
Subject: [PATCH 1/6] [NFC][OpenMP] Add mapper-specific tests exercising
pointee seciton mapping.
Also add a couple of tests that require correct propagation of map-type-modifier
bits into the mapper.
---
...t_map_nested_ptr_member_mapper_codegen.cpp | 211 ++++++++++++++++++
.../mapping/declare_mapper_nested_mappers.cpp | 7 +-
.../test/mapping/declare_mapper_target.cpp | 2 +-
.../mapping/declare_mapper_target_data.cpp | 2 +-
.../declare_mapper_target_data_enter_exit.cpp | 2 +-
.../mapping/declare_mapper_target_update.cpp | 2 +-
offload/test/mapping/mapper_map_ptee_only.c | 49 ++++
.../mapper_map_ptee_only_2_ptr_indirections.c | 61 +++++
...r_map_ptee_only_2_ptr_indirections_array.c | 93 ++++++++
.../mapping/mapper_map_ptee_only_2ndlevel.c | 52 +++++
.../mapper_map_ptee_only_2ndlevel_array.c | 82 +++++++
.../mapper_map_ptee_only_always_array.c | 78 +++++++
.../test/mapping/mapper_map_ptee_only_array.c | 59 +++++
.../mapper_map_ptr_ptee_nomapper_del_ptee.c | 43 ++++
.../mapper_map_ptr_ptee_nomapper_del_ptr.c | 43 ++++
.../mapper_target_update_present_ptee.c | 52 +++++
.../multiple_deletes_within_one_struct.c | 44 ++++
17 files changed, 875 insertions(+), 7 deletions(-)
create mode 100644 clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp
create mode 100644 offload/test/mapping/mapper_map_ptee_only.c
create mode 100644 offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c
create mode 100644 offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections_array.c
create mode 100644 offload/test/mapping/mapper_map_ptee_only_2ndlevel.c
create mode 100644 offload/test/mapping/mapper_map_ptee_only_2ndlevel_array.c
create mode 100644 offload/test/mapping/mapper_map_ptee_only_always_array.c
create mode 100644 offload/test/mapping/mapper_map_ptee_only_array.c
create mode 100644 offload/test/mapping/mapper_map_ptr_ptee_nomapper_del_ptee.c
create mode 100644 offload/test/mapping/mapper_map_ptr_ptee_nomapper_del_ptr.c
create mode 100644 offload/test/mapping/mapper_target_update_present_ptee.c
create mode 100644 offload/test/mapping/multiple_deletes_within_one_struct.c
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
new file mode 100644
index 0000000000000..d445c8014c656
--- /dev/null
+++ b/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp
@@ -0,0 +1,211 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --functions ".*mapper.*" --function-signature --check-globals --filter-out-after "getelem.*kernel" --filter-out "= alloca.*" --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _ --global-value-regex "\.offload_.*" --global-hex-value-regex ".offload_maptypes.*"
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+// S2 mapper for: map(to: arr[0:2])
+// Per-element entries (i = array element index; N = __tgt_mapper_num_components()):
+// &arr[i], &arr[i].s1p, sizeof(s1p..z), MEMBER_OF(N) | ALLOC
+// &arr[i], &arr[i].z, sizeof(int), MEMBER_OF(N+1) | TO | FROM
+// &arr[i].s1p, &arr[i].s1p->x, sizeof(int), MEMBER_OF(N+1) | TO | FROM | PTR_AND_OBJ
+// &arr[i].s1p, &arr[i].s1p->y, sizeof(int), MEMBER_OF(N+1) | TO | FROM | PTR_AND_OBJ
+// FIXME: should use attach-style codegen for s1p->x/y instead of PTR_AND_OBJ,
+// which unnecessarily also maps s1p.
+
+typedef struct {
+ int x;
+ int y;
+} S1;
+
+typedef struct {
+ S1 *s1p;
+ int z;
+} S2;
+
+#pragma omp declare mapper(default : S2 s2) map(s2.z, s2.s1p->x, s2.s1p->y)
+
+void foo(S2 *arr) {
+ // &arr, &arr[0], 2*sizeof(S2), TARGET_PARAM | TO
+ // (mapper handles individual members)
+#pragma omp target enter data map(to: arr[0:2])
+ {}
+}
+
+#endif
+//.
+// CHECK: @.offload_sizes = private unnamed_addr constant [2 x i64] [i64 32, i64 8]
+// CHECK: @.offload_maptypes = private unnamed_addr constant [2 x i64] [i64 [[#0x1]], i64 [[#0x4000]]]
+//.
+// CHECK-LABEL: define {{[^@]+}}@_Z3fooP2S2
+// CHECK-SAME: (ptr noundef [[ARR:%.*]]) #[[ATTR0:[0-9]+]] {
+// CHECK: entry:
+// CHECK: store ptr [[ARR]], ptr [[ARR_ADDR:%.*]], align 8
+// CHECK: [[TMP0:%.*]] = load ptr, ptr [[ARR_ADDR]], align 8
+// CHECK: [[TMP1:%.*]] = load ptr, ptr [[ARR_ADDR]], align 8
+// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw [[STRUCT_S2:%.*]], ptr [[TMP1]], i64 0
+// CHECK: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
+// CHECK: store ptr [[TMP0]], ptr [[TMP2]], align 8
+// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
+// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8
+// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
+// CHECK: store ptr @.omp_mapper._ZTS2S2.default, ptr [[TMP4]], align 8
+// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
+// CHECK: store ptr [[ARR_ADDR]], ptr [[TMP5]], align 8
+// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1
+// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP6]], align 8
+// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
+// CHECK: store ptr null, ptr [[TMP7]], align 8
+// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 2, ptr [[TMP8]], ptr [[TMP9]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr [[DOTOFFLOAD_MAPPERS]])
+// CHECK: ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@.omp_mapper._ZTS2S2.default
+// CHECK-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]], ptr noundef [[TMP2:%.*]], i64 noundef [[TMP3:%.*]], i64 noundef [[TMP4:%.*]], ptr noundef [[TMP5:%.*]]) #[[ATTR1:[0-9]+]] {
+// CHECK: entry:
+// CHECK: [[TMP6:%.*]] = udiv exact i64 [[TMP3]], 16
+// CHECK: [[TMP7:%.*]] = getelementptr [[STRUCT_S2:%.*]], ptr [[TMP2]], i64 [[TMP6]]
+// CHECK: [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP6]], 1
+// CHECK: [[TMP8:%.*]] = and i64 [[TMP4]], 8
+// CHECK: [[TMP9:%.*]] = icmp ne ptr [[TMP1]], [[TMP2]]
+// CHECK: [[TMP10:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP9]]
+// CHECK: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP8]], 0
+// CHECK: [[TMP11:%.*]] = and i1 [[TMP10]], [[DOTOMP_ARRAY__INIT__DELETE]]
+// CHECK: br i1 [[TMP11]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]]
+// CHECK: .omp.array..init:
+// CHECK: [[TMP12:%.*]] = mul nuw i64 [[TMP6]], 16
+// CHECK: [[TMP13:%.*]] = and i64 [[TMP4]], -4
+// CHECK: [[TMP14:%.*]] = or i64 [[TMP13]], 512
+// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP12]], i64 [[TMP14]], ptr [[TMP5]])
+// CHECK: br label [[OMP_ARRAYMAP_HEAD]]
+// CHECK: omp.arraymap.head:
+// 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: [[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
+// CHECK: [[TMP15:%.*]] = load ptr, ptr [[S1P1]], align 8
+// CHECK: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP15]], i32 0, i32 0
+// CHECK: [[S1P2:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
+// CHECK: [[S1P3:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
+// CHECK: [[TMP16:%.*]] = load ptr, ptr [[S1P3]], align 8
+// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[TMP16]], i32 0, i32 1
+// CHECK: [[TMP17:%.*]] = getelementptr i32, ptr [[Z]], i32 1
+// CHECK: [[TMP18:%.*]] = ptrtoaddr ptr [[TMP17]] to i64
+// CHECK: [[TMP19:%.*]] = ptrtoaddr ptr [[S1P]] to i64
+// CHECK: [[TMP20:%.*]] = sub i64 [[TMP18]], [[TMP19]]
+// CHECK: [[TMP21:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]])
+// CHECK: [[TMP22:%.*]] = shl i64 [[TMP21]], 48
+// CHECK: [[TMP23:%.*]] = add nuw i64 0, [[TMP22]]
+// CHECK: [[TMP24:%.*]] = and i64 [[TMP4]], 3
+// CHECK: [[TMP25:%.*]] = icmp eq i64 [[TMP24]], 0
+// CHECK: br i1 [[TMP25]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]]
+// CHECK: omp.type.alloc:
+// CHECK: [[TMP26:%.*]] = and i64 [[TMP23]], -4
+// CHECK: br label [[OMP_TYPE_END:%.*]]
+// CHECK: omp.type.alloc.else:
+// CHECK: [[TMP27:%.*]] = icmp eq i64 [[TMP24]], 1
+// CHECK: br i1 [[TMP27]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]]
+// CHECK: omp.type.to:
+// CHECK: [[TMP28:%.*]] = and i64 [[TMP23]], -3
+// CHECK: br label [[OMP_TYPE_END]]
+// CHECK: omp.type.to.else:
+// CHECK: [[TMP29:%.*]] = icmp eq i64 [[TMP24]], 2
+// CHECK: br i1 [[TMP29]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]]
+// CHECK: omp.type.from:
+// CHECK: [[TMP30:%.*]] = and i64 [[TMP23]], -2
+// 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: omp.type.alloc4:
+// CHECK: [[TMP34:%.*]] = and i64 [[TMP31]], -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: omp.type.to6:
+// CHECK: [[TMP36:%.*]] = and i64 [[TMP31]], -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: omp.type.from8:
+// CHECK: [[TMP38:%.*]] = and i64 [[TMP31]], -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_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.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: br label [[OMP_DONE]]
+// CHECK: omp.done:
+// CHECK: ret void
+//
diff --git a/offload/test/mapping/declare_mapper_nested_mappers.cpp b/offload/test/mapping/declare_mapper_nested_mappers.cpp
index a59ed6980ec4c..5966bb6127af8 100644
--- a/offload/test/mapping/declare_mapper_nested_mappers.cpp
+++ b/offload/test/mapping/declare_mapper_nested_mappers.cpp
@@ -7,7 +7,8 @@ typedef struct {
int a;
double *b;
} C;
-#pragma omp declare mapper(id1 : C s) map(to : s.a) map(from : s.b[0 : 2])
+#pragma omp declare mapper(id1 : C s) map(to : s.a) map(alloc : s.b) \
+ map(from : s.b[0 : 2])
typedef struct {
int e;
@@ -16,7 +17,7 @@ typedef struct {
short *g;
} D;
#pragma omp declare mapper(default : D r) map(from : r.e) \
- map(mapper(id1), tofrom : r.f) map(tofrom : r.g[0 : r.h])
+ map(mapper(id1), tofrom : r.f) map(alloc : r.g) map(tofrom : r.g[0 : r.h])
int main() {
constexpr int N = 10;
@@ -56,7 +57,7 @@ int main() {
spp[0][0].f.b[1] = 40;
spp[0][0].g[1] = 50;
}
- printf("%d %d %d %d\n", spp00fa, spp00fb_r, spp00fg1, spp00fg_r);
+ printf("%d %d %d %d\n", spp00fa, spp00fb_r, spp00fg1, spp00fg_r);
// CHECK: 222 0 30 0
printf("%d %d %4.5f %d %d %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1],
spp[0][0].f.b == &x[0] ? 1 : 0, spp[0][0].g[1],
diff --git a/offload/test/mapping/declare_mapper_target.cpp b/offload/test/mapping/declare_mapper_target.cpp
index 4d7237e94657a..c642949b2a367 100644
--- a/offload/test/mapping/declare_mapper_target.cpp
+++ b/offload/test/mapping/declare_mapper_target.cpp
@@ -10,7 +10,7 @@ class C {
int *a;
};
-#pragma omp declare mapper(id : C s) map(s.a[0 : NUM])
+#pragma omp declare mapper(id : C s) map(alloc : s.a) map(s.a[0 : NUM])
int main() {
C c;
diff --git a/offload/test/mapping/declare_mapper_target_data.cpp b/offload/test/mapping/declare_mapper_target_data.cpp
index 7f09844014005..fd2c8f4886401 100644
--- a/offload/test/mapping/declare_mapper_target_data.cpp
+++ b/offload/test/mapping/declare_mapper_target_data.cpp
@@ -10,7 +10,7 @@ class C {
int *a;
};
-#pragma omp declare mapper(id : C s) map(s.a[0 : NUM])
+#pragma omp declare mapper(id : C s) map(alloc : s.a) map(s.a[0 : NUM])
int main() {
C c;
diff --git a/offload/test/mapping/declare_mapper_target_data_enter_exit.cpp b/offload/test/mapping/declare_mapper_target_data_enter_exit.cpp
index f5fad8b8fe332..245462eab5785 100644
--- a/offload/test/mapping/declare_mapper_target_data_enter_exit.cpp
+++ b/offload/test/mapping/declare_mapper_target_data_enter_exit.cpp
@@ -10,7 +10,7 @@ class C {
int *a;
};
-#pragma omp declare mapper(id : C s) map(s.a[0 : NUM])
+#pragma omp declare mapper(id : C s) map(alloc : s.a) map(s.a[0 : NUM])
int main() {
C c;
diff --git a/offload/test/mapping/declare_mapper_target_update.cpp b/offload/test/mapping/declare_mapper_target_update.cpp
index fe4597b76908f..4e053e3e56139 100644
--- a/offload/test/mapping/declare_mapper_target_update.cpp
+++ b/offload/test/mapping/declare_mapper_target_update.cpp
@@ -10,7 +10,7 @@ class C {
int *a;
};
-#pragma omp declare mapper(id : C s) map(s.a[0 : NUM])
+#pragma omp declare mapper(id : C s) map(alloc : s.a) map(s.a[0 : NUM])
int main() {
C c;
diff --git a/offload/test/mapping/mapper_map_ptee_only.c b/offload/test/mapping/mapper_map_ptee_only.c
new file mode 100644
index 0000000000000..6e9904bb1917e
--- /dev/null
+++ b/offload/test/mapping/mapper_map_ptee_only.c
@@ -0,0 +1,49 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// This test ensures that when mapping `s.x` and `s.p[0:10] via a mapper, the
+// storage of `s.p` and hence `s.dummy` is not unnecessarily allocated.
+
+int x[10];
+
+typedef struct {
+ int x;
+ int dummy[10000];
+ int *p;
+} S;
+
+#pragma omp declare mapper(default : S s) map(s.x, s.p[0 : 10])
+
+S s1;
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+
+int main() {
+ s1.p = (int *)&x;
+
+#pragma omp target enter data map(alloc : s1)
+ printf("After mapping\n");
+ print_status(&s1.x, "x"); // CHECK: x is present
+ // FIXME: mapper should not map s.dummy or s.p; will be fixed when mapper
+ // emits attach-style maps for pointer members.
+ print_status(&s1.dummy, "dummy"); // CHECK: dummy is present
+ // EXPECTED: dummy is not present
+ // FIXME: mapper should not map s.dummy or s.p; will be fixed when mapper
+ // emits attach-style maps for pointer members.
+ print_status(&s1.p, "p"); // CHECK: p is present
+ // EXPECTED: p is not present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is present
+ printf("\n");
+
+#pragma omp target exit data map(delete : s1)
+ printf("After deleting\n");
+ print_status(&s1.x, "x"); // CHECK: x is not present
+ print_status(&s1.dummy, "dummy"); // CHECK: dummy is not present
+ print_status(&s1.p, "p"); // CHECK: p is not present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is not present
+}
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
new file mode 100644
index 0000000000000..5263a7da6a32b
--- /dev/null
+++ b/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c
@@ -0,0 +1,61 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// The mapper maps s2.z, s2.s1p->x, s2.s1p->y, and s2.s1p->p[0:10].
+// Check that s2.s1p->dummy and s2.s1p->p itself are not mapped, and that all
+// mapped fields are correctly removed on delete.
+
+int x[10];
+
+typedef struct {
+ int x;
+ int y;
+ int dummy[10000];
+ int *p;
+} S1;
+
+typedef struct {
+ S1 *s1p;
+ int z;
+} S2;
+
+#pragma omp declare mapper(default : S2 s2) \
+ map(s2.z, s2.s1p -> x, s2.s1p->y, s2.s1p->p[0 : 10])
+
+S1 s1;
+S2 s2;
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+
+int main() {
+ s2.s1p = &s1;
+ s2.s1p->p = (int *)&x;
+
+#pragma omp target enter data map(alloc : s2)
+ printf("After mapping\n");
+ print_status(&s2.s1p->x, "x"); // CHECK: x is present
+ print_status(&s2.s1p->y, "y"); // CHECK: y is present
+ print_status(&s2.z, "z"); // CHECK: z is present
+ print_status(&s2.s1p->dummy, "dummy"); // CHECK: dummy is not present
+ print_status(&s2.s1p->p, "p"); // CHECK: p is present
+ // FIXME: mapper should emit attach-style maps for pointer members.
+ // EXPECTED: p is not present
+ print_status(&s2.s1p->p[0], "p[0]"); // CHECK: p[0] is present
+ printf("\n");
+
+#pragma omp target exit data map(delete : s2)
+ printf("After deleting\n");
+ print_status(&s2.s1p->x, "x"); // CHECK: x is not present
+ 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: mapper should emit attach-style maps for pointer members.
+ // EXPECTED: 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_2_ptr_indirections_array.c b/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections_array.c
new file mode 100644
index 0000000000000..09eed525d4340
--- /dev/null
+++ b/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections_array.c
@@ -0,0 +1,93 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// Array variant of mapper_map_ptee_only_2_ptr_indirections.c.
+// The mapper maps s2.z, s2.s1p->x, s2.s1p->y, and s2.s1p->p[0:10].
+// s2.s1p->dummy and s2.s1p->p itself are not mapped.
+
+int x[2][10];
+
+typedef struct {
+ int x;
+ int y;
+ int dummy[10000];
+ int *p;
+} S1;
+
+typedef struct {
+ S1 *s1p;
+ int z;
+} S2;
+
+#pragma omp declare mapper(default : S2 s2) \
+ map(s2.z, s2.s1p -> x, s2.s1p->y, s2.s1p->p[0 : 10])
+
+S1 s1arr[2];
+S2 s2arr[2];
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+
+int main() {
+ s2arr[0].s1p = &s1arr[0];
+ s2arr[1].s1p = &s1arr[1];
+ s2arr[0].s1p->p = (int *)&x[0];
+ s2arr[1].s1p->p = (int *)&x[1];
+
+#pragma omp target enter data map(alloc : s2arr)
+ printf("After mapping\n");
+ print_status(&s2arr[0].s1p->x, "s2arr[0].x"); // CHECK: s2arr[0].x is present
+ print_status(&s2arr[0].s1p->y, "s2arr[0].y"); // CHECK: s2arr[0].y is present
+ print_status(&s2arr[0].z, "s2arr[0].z"); // CHECK: s2arr[0].z is present
+ print_status(&s2arr[0].s1p->dummy,
+ "s2arr[0].dummy"); // CHECK: s2arr[0].dummy is not present
+ // FIXME: mapper should not map s1p->p; will be fixed when mapper emits
+ // attach-style maps for pointer members.
+ print_status(&s2arr[0].s1p->p,
+ "s2arr[0].p"); // CHECK: s2arr[0].p is present
+ // EXPECTED: s2arr[0].p is not present
+ print_status(&s2arr[0].s1p->p[0],
+ "s2arr[0].p[0]"); // CHECK: s2arr[0].p[0] is present
+ print_status(&s2arr[1].s1p->x, "s2arr[1].x"); // CHECK: s2arr[1].x is present
+ print_status(&s2arr[1].s1p->y, "s2arr[1].y"); // CHECK: s2arr[1].y is present
+ print_status(&s2arr[1].z, "s2arr[1].z"); // CHECK: s2arr[1].z is present
+ print_status(&s2arr[1].s1p->dummy,
+ "s2arr[1].dummy"); // CHECK: s2arr[1].dummy is not present
+ // FIXME: mapper should not map s1p->p; will be fixed when mapper emits
+ // attach-style maps for pointer members.
+ print_status(&s2arr[1].s1p->p,
+ "s2arr[1].p"); // CHECK: s2arr[1].p is present
+ // EXPECTED: s2arr[1].p is not present
+ print_status(&s2arr[1].s1p->p[0],
+ "s2arr[1].p[0]"); // CHECK: s2arr[1].p[0] is present
+ printf("\n");
+
+#pragma omp target exit data map(delete : s2arr)
+ printf("After deleting\n");
+ print_status(&s2arr[0].s1p->x,
+ "s2arr[0].x"); // CHECK: s2arr[0].x is not present
+ print_status(&s2arr[0].s1p->y,
+ "s2arr[0].y"); // CHECK: s2arr[0].y is not present
+ print_status(&s2arr[0].z, "s2arr[0].z"); // CHECK: s2arr[0].z is not present
+ print_status(&s2arr[0].s1p->dummy,
+ "s2arr[0].dummy"); // CHECK: s2arr[0].dummy is not present
+ print_status(&s2arr[0].s1p->p,
+ "s2arr[0].p"); // CHECK: s2arr[0].p is not present
+ print_status(&s2arr[0].s1p->p[0],
+ "s2arr[0].p[0]"); // CHECK: s2arr[0].p[0] is not present
+ print_status(&s2arr[1].s1p->x,
+ "s2arr[1].x"); // CHECK: s2arr[1].x is not present
+ print_status(&s2arr[1].s1p->y,
+ "s2arr[1].y"); // CHECK: s2arr[1].y is not present
+ print_status(&s2arr[1].z, "s2arr[1].z"); // CHECK: s2arr[1].z is not present
+ print_status(&s2arr[1].s1p->dummy,
+ "s2arr[1].dummy"); // CHECK: s2arr[1].dummy is not present
+ print_status(&s2arr[1].s1p->p,
+ "s2arr[1].p"); // CHECK: s2arr[1].p is not present
+ print_status(&s2arr[1].s1p->p[0],
+ "s2arr[1].p[0]"); // CHECK: s2arr[1].p[0] is not present
+}
diff --git a/offload/test/mapping/mapper_map_ptee_only_2ndlevel.c b/offload/test/mapping/mapper_map_ptee_only_2ndlevel.c
new file mode 100644
index 0000000000000..d7803b4e24d9a
--- /dev/null
+++ b/offload/test/mapping/mapper_map_ptee_only_2ndlevel.c
@@ -0,0 +1,52 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// This test ensures that when mapping `s2.s1.x` and `s2.s1.p[0:10] via a
+// mapper, the storage of `s2.s1.p` and hence `s2.s1.dummy` is not unnecessarily
+// allocated.
+
+int x[10];
+
+typedef struct {
+ int x;
+ int dummy[10000];
+ int *p;
+} S1;
+
+typedef struct {
+ S1 s1;
+} S2;
+
+#pragma omp declare mapper(default : S2 s2) map(s2.s1.x, s2.s1.p[0 : 10])
+
+S2 s2;
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+
+int main() {
+ s2.s1.p = (int *)&x;
+
+#pragma omp target enter data map(alloc : s2)
+ printf("After mapping\n");
+ print_status(&s2.s1.x, "x"); // CHECK: x is present
+ print_status(&s2.s1.dummy, "dummy"); // CHECK: dummy is not present
+ print_status(&s2.s1.p, "p"); // CHECK: p is present
+ // FIXME: mapper should emit attach-style maps for pointer members.
+ // EXPECTED: p is not present
+ print_status(&s2.s1.p[0], "p[0]"); // CHECK: p[0] is present
+ printf("\n");
+
+#pragma omp target exit data map(delete : s2)
+ printf("After deleting\n");
+ print_status(&s2.s1.x, "x"); // CHECK: x is not present
+ print_status(&s2.s1.dummy, "dummy"); // CHECK: dummy is not present
+ print_status(&s2.s1.p, "p"); // CHECK: p is present
+ // FIXME: mapper should emit attach-style maps for pointer members.
+ // EXPECTED: p is not present
+ print_status(&s2.s1.p[0], "p[0]"); // CHECK: p[0] is not present
+}
diff --git a/offload/test/mapping/mapper_map_ptee_only_2ndlevel_array.c b/offload/test/mapping/mapper_map_ptee_only_2ndlevel_array.c
new file mode 100644
index 0000000000000..c745779ab4b97
--- /dev/null
+++ b/offload/test/mapping/mapper_map_ptee_only_2ndlevel_array.c
@@ -0,0 +1,82 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// Test that a mapper on a nested struct correctly maps and deletes data when
+// applied to an array of structs. s2arr[i].s1.x and s2arr[i].s1.p[0:10] are
+// mapped; after delete they are gone.
+// It's not ideal that s2arr[i].s1.dummy and s2arr[i].s1.p are also present
+// after enter-data (the runtime allocates the full contiguous storage of the
+// array, otherwise it's tricky to map an array with "holes"), but everything is
+// correctly removed after delete.
+
+int x[2][10];
+
+typedef struct {
+ int x;
+ int dummy[10000];
+ int *p;
+} S1;
+
+typedef struct {
+ S1 s1;
+} S2;
+
+#pragma omp declare mapper(default : S2 s2) map(s2.s1.x, s2.s1.p[0 : 10])
+
+S2 s2arr[2];
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+
+int main() {
+ s2arr[0].s1.x = 111;
+ s2arr[1].s1.x = 222;
+ s2arr[0].s1.p = (int *)&x[0];
+ s2arr[1].s1.p = (int *)&x[1];
+
+#pragma omp target enter data map(to : s2arr)
+ printf("After mapping\n");
+ print_status(&s2arr[0].s1.x,
+ "s2arr[0].s1.x"); // CHECK: s2arr[0].s1.x is present
+ // dummy/p being present is not ideal, but that's what we get with the
+ // current implementation because we need to map the full contiguous
+ // chunk for the array first before invoking the mapper.
+ print_status(&s2arr[0].s1.dummy,
+ "s2arr[0].s1.dummy"); // CHECK: s2arr[0].s1.dummy is present
+ print_status(&s2arr[0].s1.p,
+ "s2arr[0].s1.p"); // CHECK: s2arr[0].s1.p is present
+ print_status(&s2arr[0].s1.p[0],
+ "s2arr[0].s1.p[0]"); // CHECK: s2arr[0].s1.p[0] is present
+ print_status(&s2arr[1].s1.x,
+ "s2arr[1].s1.x"); // CHECK: s2arr[1].s1.x is present
+ print_status(&s2arr[1].s1.dummy,
+ "s2arr[1].s1.dummy"); // CHECK: s2arr[1].s1.dummy is present
+ print_status(&s2arr[1].s1.p,
+ "s2arr[1].s1.p"); // CHECK: s2arr[1].s1.p is present
+ print_status(&s2arr[1].s1.p[0],
+ "s2arr[1].s1.p[0]"); // CHECK: s2arr[1].s1.p[0] is present
+
+ printf("\n");
+#pragma omp target exit data map(delete : s2arr)
+ printf("After deleting\n");
+ print_status(&s2arr[0].s1.x,
+ "s2arr[0].s1.x"); // CHECK: s2arr[0].s1.x is not present
+ print_status(&s2arr[0].s1.dummy,
+ "s2arr[0].s1.dummy"); // CHECK: s2arr[0].s1.dummy is not present
+ print_status(&s2arr[0].s1.p,
+ "s2arr[0].s1.p"); // CHECK: s2arr[0].s1.p is not present
+ print_status(&s2arr[0].s1.p[0],
+ "s2arr[0].s1.p[0]"); // CHECK: s2arr[0].s1.p[0] is not present
+ print_status(&s2arr[1].s1.x,
+ "s2arr[1].s1.x"); // CHECK: s2arr[1].s1.x is not present
+ print_status(&s2arr[1].s1.dummy,
+ "s2arr[1].s1.dummy"); // CHECK: s2arr[1].s1.dummy is not present
+ print_status(&s2arr[1].s1.p,
+ "s2arr[1].s1.p"); // CHECK: s2arr[1].s1.p is not present
+ print_status(&s2arr[1].s1.p[0],
+ "s2arr[1].s1.p[0]"); // CHECK: s2arr[1].s1.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
new file mode 100644
index 0000000000000..64aaabd80dfe2
--- /dev/null
+++ b/offload/test/mapping/mapper_map_ptee_only_always_array.c
@@ -0,0 +1,78 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// Test that a mapper that maps a var and a pointee, correctly propagates the
+// always and present bits into the individual maps "pushed" by it.
+
+int x[2][10];
+
+typedef struct {
+ int x;
+ int dummy[10000];
+ int *p;
+} S;
+
+#pragma omp declare mapper(default : S s) map(s.x, s.p[0 : 10])
+
+S s1[2];
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+
+int main() {
+ x[0][0] = x[0][1] = s1[0].x = s1[1].x = 111;
+ s1[0].p = (int *)&x[0];
+ s1[1].p = (int *)&x[1];
+
+#pragma omp target enter data map(alloc : s1)
+ printf("After mapping\n");
+ print_status(&s1[0].x, "s1[0].x"); // CHECK: s1[0].x is present
+ // dummy/p being present is not ideal, but that's what we get with the
+ // current implementation because we need to map the full contiguous
+ // chunk for the array first before invoking the mapper.
+ print_status(&s1[0].dummy, "s1[0].dummy"); // CHECK: s1[0].dummy is present
+ print_status(&s1[0].p, "s1[0].p"); // CHECK: s1[0].p is present
+ print_status(&s1[0].p[0], "s1[0].p[0]"); // CHECK: s1[0].p[0] is present
+ print_status(&s1[1].x, "s1[1].x"); // CHECK: s1[1].x is present
+ print_status(&s1[1].dummy, "s1[1].dummy"); // CHECK: s1[1].dummy is present
+ print_status(&s1[1].p, "s1[1].p"); // CHECK: s1[1].p is present
+ print_status(&s1[1].p[0], "s1[1].p[0]"); // CHECK: s1[1].p[0] is present
+
+#pragma omp target map(always, present, from : s1)
+ {
+ s1[0].p[0] = s1[1].p[0] = s1[0].x = s1[1].x = 222;
+ }
+
+ 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("\n");
+
+#pragma omp target exit data map(delete : s1)
+ printf("After deleting\n");
+ print_status(&s1[0].x, "s1[0].x"); // CHECK: s1[0].x is not present
+ print_status(&s1[0].dummy,
+ "s1[0].dummy"); // CHECK: s1[0].dummy is not present
+ print_status(&s1[0].p, "s1[0].p"); // CHECK: s1[0].p is not present
+ print_status(&s1[0].p[0], "s1[0].p[0]"); // CHECK: s1[0].p[0] is not present
+ print_status(&s1[1].x, "s1[1].x"); // CHECK: s1[1].x is not present
+ print_status(&s1[1].dummy,
+ "s1[1].dummy"); // CHECK: s1[1].dummy is not present
+ print_status(&s1[1].p, "s1[1].p"); // CHECK: s1[1].p is not present
+ print_status(&s1[1].p[0], "s1[1].p[0]"); // CHECK: s1[1].p[0] is not present
+}
diff --git a/offload/test/mapping/mapper_map_ptee_only_array.c b/offload/test/mapping/mapper_map_ptee_only_array.c
new file mode 100644
index 0000000000000..c5f194d4b6c21
--- /dev/null
+++ b/offload/test/mapping/mapper_map_ptee_only_array.c
@@ -0,0 +1,59 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// Test that a mapper that maps a var and a pointee, works on an array of
+// structs in terms of properly allocating/deleting the data.
+
+int x[2][10];
+
+typedef struct {
+ int x;
+ int dummy[10000];
+ int *p;
+} S;
+
+#pragma omp declare mapper(default : S s) map(s.x, s.p[0 : 10])
+
+S s1[2];
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+
+int main() {
+ s1[0].x = 111;
+ s1[1].x = 222;
+ s1[0].p = (int *)&x[0];
+ s1[1].p = (int *)&x[1];
+
+#pragma omp target enter data map(to : s1)
+ printf("After mapping\n");
+ print_status(&s1[0].x, "s1[0].x"); // CHECK: s1[0].x is present
+ // dummy/p being present is not ideal, but that's what we get with the
+ // current implementation because we need to map the full contiguous
+ // chunk for the array first before invoking the mapper.
+ print_status(&s1[0].dummy, "s1[0].dummy"); // CHECK: s1[0].dummy is present
+ print_status(&s1[0].p, "s1[0].p"); // CHECK: s1[0].p is present
+ print_status(&s1[0].p[0], "s1[0].p[0]"); // CHECK: s1[0].p[0] is present
+ print_status(&s1[1].x, "s1[1].x"); // CHECK: s1[1].x is present
+ print_status(&s1[1].dummy, "s1[1].dummy"); // CHECK: s1[1].dummy is present
+ print_status(&s1[1].p, "s1[1].p"); // CHECK: s1[1].p is present
+ print_status(&s1[1].p[0], "s1[1].p[0]"); // CHECK: s1[1].p[0] is present
+
+ printf("\n");
+#pragma omp target exit data map(delete : s1)
+ printf("After deleting\n");
+ print_status(&s1[0].x, "s1[0].x"); // CHECK: s1[0].x is not present
+ print_status(&s1[0].dummy,
+ "s1[0].dummy"); // CHECK: s1[0].dummy is not present
+ print_status(&s1[0].p, "s1[0].p"); // CHECK: s1[0].p is not present
+ print_status(&s1[0].p[0], "s1[0].p[0]"); // CHECK: s1[0].p[0] is not present
+ print_status(&s1[1].x, "s1[1].x"); // CHECK: s1[1].x is not present
+ print_status(&s1[1].dummy,
+ "s1[1].dummy"); // CHECK: s1[1].dummy is not present
+ print_status(&s1[1].p, "s1[1].p"); // CHECK: s1[1].p is not present
+ print_status(&s1[1].p[0], "s1[1].p[0]"); // CHECK: s1[1].p[0] is not present
+}
diff --git a/offload/test/mapping/mapper_map_ptr_ptee_nomapper_del_ptee.c b/offload/test/mapping/mapper_map_ptr_ptee_nomapper_del_ptee.c
new file mode 100644
index 0000000000000..26a7098016207
--- /dev/null
+++ b/offload/test/mapping/mapper_map_ptr_ptee_nomapper_del_ptee.c
@@ -0,0 +1,43 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// This test ensures that when mapping `s.p` and `s.p[0:10] via a mapper, they
+// occupy different storage blocks in memory, so deleting `s.p[0:10]` does not
+// delete `s.p`.
+
+int x[10];
+
+typedef struct {
+ int *p;
+} S;
+
+#pragma omp declare mapper(default : S s) map(s.p, s.p[0 : 10])
+
+S s1;
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+int main() {
+ s1.p = (int *)&x;
+
+#pragma omp target enter data map(alloc : s1)
+ printf("After mapping ptr and ptee\n");
+ print_status(&s1.p, "p"); // CHECK: p is present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is present
+ printf("\n");
+
+#pragma omp target exit data map(delete : s1.p[0 : 10])
+ printf("After deleting ptee\n");
+ print_status(&s1.p, "p"); // CHECK: p is present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is not present
+ printf("\n");
+
+#pragma omp target exit data map(delete : s1.p)
+ printf("After deleting ptr\n");
+ print_status(&s1.p, "p"); // CHECK: p is not present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is not present
+}
diff --git a/offload/test/mapping/mapper_map_ptr_ptee_nomapper_del_ptr.c b/offload/test/mapping/mapper_map_ptr_ptee_nomapper_del_ptr.c
new file mode 100644
index 0000000000000..7457a988e5933
--- /dev/null
+++ b/offload/test/mapping/mapper_map_ptr_ptee_nomapper_del_ptr.c
@@ -0,0 +1,43 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// This test ensures that when mapping `s.p` and `s.p[0:10] via a mapper, they
+// occupy different storage blocks in memory, so deleting `s.p` does not delete
+// `s.p[0:10]`.
+
+int x[10];
+
+typedef struct {
+ int *p;
+} S;
+
+#pragma omp declare mapper(default : S s) map(s.p, s.p[0 : 10])
+
+S s1;
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+int main() {
+ s1.p = (int *)&x;
+
+#pragma omp target enter data map(alloc : s1)
+ printf("After mapping ptr and ptee\n");
+ print_status(&s1.p, "p"); // CHECK: p is present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is present
+ printf("\n");
+
+#pragma omp target exit data map(delete : s1.p)
+ printf("After deleting ptr\n");
+ print_status(&s1.p, "p"); // CHECK: p is not present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is present
+ printf("\n");
+
+#pragma omp target exit data map(delete : s1.p[0 : 10])
+ printf("After deleting ptee\n");
+ print_status(&s1.p, "p"); // CHECK: p is not present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is not present
+}
diff --git a/offload/test/mapping/mapper_target_update_present_ptee.c b/offload/test/mapping/mapper_target_update_present_ptee.c
new file mode 100644
index 0000000000000..06d75e343ddd2
--- /dev/null
+++ b/offload/test/mapping/mapper_target_update_present_ptee.c
@@ -0,0 +1,52 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// RUN: %libomptarget-compile-generic -DOUT_OF_BOUNDS
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic --check-prefix=CHECK-OOB
+
+#include <stdio.h>
+
+int x[10];
+typedef struct {
+ int y;
+ int *p;
+} S;
+
+#ifdef OUT_OF_BOUNDS
+// s.p[0:20] extends beyond the mapped region x[0:10]; present check should
+// fail.
+// NOTE: While OpenMP 6.0:296 implies that the present "motion" modifier should
+// not be propagated to s.p[0:20], it is under discussion whether that was
+// intentional. Not propagating it would require treating the present modifier
+// differently for data-motion clauses (to/from) vs. map clauses.
+#pragma omp declare mapper(S s) map(s.y, s.p[0 : 20])
+#else
+#pragma omp declare mapper(S s) map(s.y, s.p[0 : 2])
+#endif
+S s;
+
+void f1() {
+#pragma omp target update to(present : s)
+
+#pragma omp target data use_device_addr(s, x)
+#pragma omp target has_device_addr(s, x)
+ {
+ s.y = s.y + 222;
+ x[0] = x[0] + 222;
+ }
+}
+
+int main() {
+ x[0] = 111;
+ s.y = 111;
+ s.p = &x[0];
+
+ // CHECK-OOB: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]]
+ fprintf(stderr, "addr=%p, size=%zu\n", &s.p[0], 20 * sizeof(s.p[0]));
+
+#pragma omp target data map(from : s.y, x)
+ {
+ f1();
+ }
+
+ printf("%d %d\n", x[0], s.y); // CHECK: 333 333
+}
diff --git a/offload/test/mapping/multiple_deletes_within_one_struct.c b/offload/test/mapping/multiple_deletes_within_one_struct.c
new file mode 100644
index 0000000000000..d5292a2acc0f7
--- /dev/null
+++ b/offload/test/mapping/multiple_deletes_within_one_struct.c
@@ -0,0 +1,44 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// This test ensures that when mapping `s.x` and `s.p[0:10], the
+// storage of `s.p` and `s.dummy` is not unnecessarily allocated.
+
+// Secondly, after deleting s1.x and s1.p[0], they are no longer
+// present on the device.
+
+int g[10];
+
+typedef struct {
+ int x;
+ int dummy[10000];
+ int *p;
+} S;
+
+S s1;
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+
+int main() {
+ s1.p = (int *)&g;
+
+#pragma omp target enter data map(alloc : s1.x, s1.p[0 : 10])
+ printf("After mapping\n");
+ print_status(&s1.x, "x"); // CHECK: x is present
+ print_status(&s1.dummy, "dummy"); // CHECK: dummy is not present
+ print_status(&s1.p, "p"); // CHECK: p is not present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is present
+ printf("\n");
+
+#pragma omp target exit data map(delete : s1.x) map(delete : s1.p[0 : 10])
+ printf("After deleting\n");
+ print_status(&s1.x, "x"); // CHECK: x is not present
+ print_status(&s1.dummy, "dummy"); // CHECK: dummy is not present
+ print_status(&s1.p, "p"); // CHECK: p is not present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is not present
+}
>From 428b5e5eec3f66f348e1426b652cef9d1257ee8a Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Thu, 16 Jul 2026 17:46:30 -0700
Subject: [PATCH 2/6] [OpenMP][test] Add mapper present-check test for absent
struct member
A user-defined mapper maps a member (s.x) and a pointee (s.p[0:10]). Pre-map
only the pointee, leaving s.x absent, then map(present: s) via the mapper.
This currently PASSES with no present-check failure even though s.x is absent:
the mapper emits the s.x member component without the PRESENT bit, so its
absence is a silent no-op. The test documents this gap with a FIXME; the
EXPECTED behavior is a present-check failure on s.x. Propagating PRESENT to
member entries is blocked while pointer members use PTR_AND_OBJ.
Co-Authored-By: Claude Opus 4.8 <noreply at anthropic.com>
---
.../mapper_map_ptee_then_present_absent_mbr.c | 60 +++++++++++++++++++
1 file changed, 60 insertions(+)
create mode 100644 offload/test/mapping/mapper_map_ptee_then_present_absent_mbr.c
diff --git a/offload/test/mapping/mapper_map_ptee_then_present_absent_mbr.c b/offload/test/mapping/mapper_map_ptee_then_present_absent_mbr.c
new file mode 100644
index 0000000000000..6293c049152db
--- /dev/null
+++ b/offload/test/mapping/mapper_map_ptee_then_present_absent_mbr.c
@@ -0,0 +1,60 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// A user-defined mapper maps a struct member (s.x) and a pointee (s.p[0:10]).
+// We pre-map ONLY the pointee, leaving s.x absent on the device, then do a
+// map(present: s) that invokes the mapper.
+//
+// FIXME: This currently PASSES (no present-check failure) even though s.x is
+// absent. The mapper emits the s.x member component WITHOUT the PRESENT bit, so
+// its absence is a silent no-op; PRESENT is only propagated to the pointee entry
+// (s.p[0:10]), which is present here. We eventually need to propagate PRESENT to
+// non-pointee (member) data as well, so that an absent member triggers the
+// present check.
+//
+// EXPECTED: the present check should FAIL (assert / abort) because s.x is not
+// present on the device. Propagating PRESENT to member entries is blocked while
+// pointer members use PTR_AND_OBJ: a single combined entry allocates the whole
+// struct (including the pointer's own storage), so propagating PRESENT there
+// would wrongly require the pointer's pointee to be present. Enable this once
+// Clang emits attach-style maps throughout instead of PTR_AND_OBJ.
+
+int x[10];
+
+typedef struct {
+ int x;
+ int dummy[10000];
+ int *p;
+} S;
+
+#pragma omp declare mapper(default : S s) map(s.x, s.p[0 : 10])
+
+S s1;
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+
+int main() {
+ s1.p = (int *)&x;
+
+ // Pre-map ONLY the pointee s1.p[0:10], NOT s1.x.
+#pragma omp target enter data map(alloc : s1.p[0 : 10])
+ printf("After mapping\n");
+ print_status(&s1.x, "x"); // CHECK: x is not present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is present
+ printf("\n");
+
+ // s1.x is NOT present on the device.
+ // FIXME: this should trigger a present-check failure on s1.x, but currently
+ // passes because PRESENT is not propagated to the s.x member entry.
+ // EXPECTED: run-fail with a "present" motion-modifier error for s1.x.
+#pragma omp target enter data map(present, alloc : s1)
+ printf("present check passed\n"); // CHECK: present check passed
+
+#pragma omp target exit data map(delete : s1)
+ return 0;
+}
>From 70f6b86b315c49c646dc1e21ed268a2475616453 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Mon, 20 Jul 2026 17:36:29 -0700
Subject: [PATCH 3/6] [OpenMP][test] Add version-based present-check RUN lines
and always-map variant
Address review on the mapper present tests:
- mapper_target_update_present_ptee.c: add RUN lines for all four
OpenMP-version x bounds combinations (5.2/6.0 x inbounds/out-of-bounds).
The present motion vs map-type modifier divergence was unintentional and is
fixed as an OpenMP 6.0 erratum, so for <= 5.2 present is not applied to the
pointee; the checks reflect the current (pre-propagation) state with FIXMEs
for the eventual 6.0 present-check failure on the out-of-bounds pointee.
- mapper_enter_data_always_present_ptee.c: new variant driving the mapper via
'target enter data map(always, present, to : s)' instead of target update.
Documents the current run-fail state (an explicit-extension error on the
mapper's combined member entry) with a FIXME for the attach-style fix.
- mapper_map_ptee_only_2_ptr_indirections.c: clarify the after-delete FIXME --
the pointer's storage stays present because the mapper drops DELETE, not
because of attach-style mapping.
Co-Authored-By: Claude Opus 4.8 <noreply at anthropic.com>
---
.../mapper_enter_data_always_present_ptee.c | 76 +++++++++++++++++++
.../mapper_map_ptee_only_2_ptr_indirections.c | 2 +-
.../mapper_target_update_present_ptee.c | 36 ++++++---
3 files changed, 102 insertions(+), 12 deletions(-)
create mode 100644 offload/test/mapping/mapper_enter_data_always_present_ptee.c
diff --git a/offload/test/mapping/mapper_enter_data_always_present_ptee.c b/offload/test/mapping/mapper_enter_data_always_present_ptee.c
new file mode 100644
index 0000000000000..9fd3a6bf0d794
--- /dev/null
+++ b/offload/test/mapping/mapper_enter_data_always_present_ptee.c
@@ -0,0 +1,76 @@
+// Same as mapper_target_update_present_ptee.c, but drives the mapper with a
+// "target enter data map(always, present, to : s)" clause instead of a
+// "target update to(present : s)" motion clause. Both invoke the mapper; this
+// checks present propagation to the pointee is consistent across the two paths.
+//
+// FIXME: this test currently run-fails at every version/bounds combination.
+// The mapper maps the struct member (s.y) with a combined entry whose size does
+// not match the member's own storage, so the map clause aborts with an
+// "explicit extension not allowed" error before the present modifier is ever
+// considered. This is fixed once the mapper emits attach-style maps for pointer
+// members (so the member and pointee occupy separate, correctly-sized entries).
+//
+// EXPECTED final state:
+// inbounds, 5.2 and 6.0: run succeeds, prints "333 333".
+// out-of-bounds (s.p[0:20] over the mapped x[0:10]):
+// 5.2: succeeds (present is not applied to the pointee before 6.0).
+// 6.0: run-fails; the failure should be the 'present' map-type-modifier
+// check on s.p[0:20] (an accompanying "explicit extension" message is
+// incidental -- for a map clause it is user error to map 20 elements
+// when only 10 are present).
+
+// RUN: %libomptarget-compile-generic -fopenmp-version=52
+// RUN: %libomptarget-run-fail-generic 2>&1 | %fcheck-generic
+// RUN: %libomptarget-compile-generic -fopenmp-version=60
+// RUN: %libomptarget-run-fail-generic 2>&1 | %fcheck-generic
+// RUN: %libomptarget-compile-generic -fopenmp-version=52 -DOUT_OF_BOUNDS
+// RUN: %libomptarget-run-fail-generic 2>&1 | %fcheck-generic
+// RUN: %libomptarget-compile-generic -fopenmp-version=60 -DOUT_OF_BOUNDS
+// RUN: %libomptarget-run-fail-generic 2>&1 | %fcheck-generic
+
+#include <stdio.h>
+
+int x[10];
+typedef struct {
+ int y;
+ int *p;
+} S;
+
+#ifdef OUT_OF_BOUNDS
+#pragma omp declare mapper(S s) map(s.y, s.p[0 : 20])
+#else
+#pragma omp declare mapper(S s) map(s.y, s.p[0 : 2])
+#endif
+S s;
+
+void f1() {
+#pragma omp target enter data map(always, present, to : s)
+
+#pragma omp target data use_device_addr(s, x)
+#pragma omp target has_device_addr(s, x)
+ {
+ s.y = s.y + 222;
+ x[0] = x[0] + 222;
+ }
+
+#pragma omp target exit data map(release : s)
+}
+
+int main() {
+ x[0] = 111;
+ s.y = 111;
+ s.p = &x[0];
+
+ fprintf(stderr, "addr=%p, size=%zu\n", &s.p[0], 20 * sizeof(s.p[0]));
+
+ // FIXME: the map clause aborts here with an "explicit extension not allowed"
+ // error on the mapper's combined member entry, at every version. Fixed once
+ // the mapper uses attach-style maps for pointer members.
+ // CHECK: explicit extension not allowed
+#pragma omp target data map(from : s.y, x)
+ {
+ f1();
+ }
+
+ printf("%d %d\n", x[0], 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 5263a7da6a32b..fb03accec0083 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
@@ -55,7 +55,7 @@ int main() {
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: mapper should emit attach-style maps for pointer members.
+ // 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
}
diff --git a/offload/test/mapping/mapper_target_update_present_ptee.c b/offload/test/mapping/mapper_target_update_present_ptee.c
index 06d75e343ddd2..86c16d7d139d9 100644
--- a/offload/test/mapping/mapper_target_update_present_ptee.c
+++ b/offload/test/mapping/mapper_target_update_present_ptee.c
@@ -1,7 +1,23 @@
-// RUN: %libomptarget-compile-run-and-check-generic
-// RUN: %libomptarget-compile-generic -DOUT_OF_BOUNDS
+// Inbounds: the pointee region is fully present; the present check passes at
+// every OpenMP version.
+// RUN: %libomptarget-compile-generic -fopenmp-version=52
+// RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefix=CHECK-52
+// RUN: %libomptarget-compile-generic -fopenmp-version=60
+// RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefix=CHECK-60
+
+// Out-of-bounds: s.p[0:20] extends beyond the mapped region x[0:10].
+// FIXME: at OpenMP 6.0 the present modifier should be propagated to the pointee
+// s.p[0:20], so this should run-fail with a present error. It currently does
+// NOT, because the mapper drops the present modifier for the pointee;
+// propagating it is done in a follow-on. For OpenMP <= 5.2 present is
+// (correctly) not applied to the pointee, so the run must succeed there.
+// EXPECTED (6.0): run-fail with a present-modifier error for s.p[0:20].
+// RUN: %libomptarget-compile-generic -fopenmp-version=52 -DOUT_OF_BOUNDS
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic --check-prefix=CHECK-52-OOB
+// RUN: %libomptarget-compile-generic -fopenmp-version=60 -DOUT_OF_BOUNDS
// RUN: %libomptarget-run-generic 2>&1 \
-// RUN: | %fcheck-generic --check-prefix=CHECK-OOB
+// RUN: | %fcheck-generic --check-prefix=CHECK-60-OOB
#include <stdio.h>
@@ -12,12 +28,6 @@ typedef struct {
} S;
#ifdef OUT_OF_BOUNDS
-// s.p[0:20] extends beyond the mapped region x[0:10]; present check should
-// fail.
-// NOTE: While OpenMP 6.0:296 implies that the present "motion" modifier should
-// not be propagated to s.p[0:20], it is under discussion whether that was
-// intentional. Not propagating it would require treating the present modifier
-// differently for data-motion clauses (to/from) vs. map clauses.
#pragma omp declare mapper(S s) map(s.y, s.p[0 : 20])
#else
#pragma omp declare mapper(S s) map(s.y, s.p[0 : 2])
@@ -40,7 +50,8 @@ int main() {
s.y = 111;
s.p = &x[0];
- // CHECK-OOB: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]]
+ // CHECK-52-OOB: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]]
+ // CHECK-60-OOB: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]]
fprintf(stderr, "addr=%p, size=%zu\n", &s.p[0], 20 * sizeof(s.p[0]));
#pragma omp target data map(from : s.y, x)
@@ -48,5 +59,8 @@ int main() {
f1();
}
- printf("%d %d\n", x[0], s.y); // CHECK: 333 333
+ // Inbounds: present check passes at both versions.
+ // CHECK-52: 333 333
+ // CHECK-60: 333 333
+ printf("%d %d\n", x[0], s.y);
}
>From 96e9d8129ec84802a4023040c98cb8671a9e1488 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Thu, 23 Jul 2026 19:18:42 -0700
Subject: [PATCH 4/6] [OpenMP][test] Address review: restructure present-check
mapper tests
Move the present-check pointee tests into this PR (PR1) with RUN lines and
FIXME/EXPECTED comments matching the current pre-feature behavior, per review:
- mapper_map_present_ptee.c
- mapper_map_mbr_ptee_then_present_mbr_ptee.c
- mapper_map_mbr_then_present_mbr_ptee.c
Later PRs update these checks as the behavior becomes correct. Also add
mapper_map_always_from.c showing the current (pre-propagation) ALWAYS behavior
(s.y = 0), which the ALWAYS/DELETE/CLOSE propagation PR flips to s.y = 111.
Also:
- mapper_map_ptee_then_present_absent_mbr.c: reword the FIXME to just say
PRESENT should be propagated to the s.x member entry (drop the PTR_AND_OBJ
wording, which belongs in the propagation PR).
- mapper_target_update_present_ptee.c: add OOB CHECK/EXPECTED lines and a
subset-transfer FIXME; use fprintf consistently.
- target_map_nested_ptr_member_mapper_codegen.cpp: add the OpenMP 6.0 RUN
line now (CHECK-60 currently identical to CHECK, with a FIXME) so the
propagation PR only needs to flip the pointee map-type values.
Co-Authored-By: Claude Opus 4.8 <noreply at anthropic.com>
---
...t_map_nested_ptr_member_mapper_codegen.cpp | 181 ++++++++++++++++++
offload/test/mapping/mapper_map_always_from.c | 46 +++++
...apper_map_mbr_ptee_then_present_mbr_ptee.c | 58 ++++++
.../mapper_map_mbr_then_present_mbr_ptee.c | 62 ++++++
.../test/mapping/mapper_map_present_ptee.c | 65 +++++++
.../mapper_map_ptee_then_present_absent_mbr.c | 13 +-
.../mapper_target_update_present_ptee.c | 32 +++-
7 files changed, 441 insertions(+), 16 deletions(-)
create mode 100644 offload/test/mapping/mapper_map_always_from.c
create mode 100644 offload/test/mapping/mapper_map_mbr_ptee_then_present_mbr_ptee.c
create mode 100644 offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c
create mode 100644 offload/test/mapping/mapper_map_present_ptee.c
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 d445c8014c656..cfdc9b861090c 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
@@ -3,6 +3,14 @@
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+// PRESENT is propagated to pointee (attach-ptr) entries only at OpenMP >= 6.0.
+// FIXME: that propagation is done in a follow-on; until then the CHECK-60
+// output below is identical to CHECK (the pointee entries carry map-type mask
+// 1036 = ALWAYS|DELETE|CLOSE, without PRESENT). Once PRESENT is propagated, the
+// attach-ptr pointee entries should use mask 5132 = ALWAYS|DELETE|CLOSE|PRESENT
+// at 6.0.
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK-60
+
// expected-no-diagnostics
#ifndef HEADER
#define HEADER
@@ -40,6 +48,9 @@ void foo(S2 *arr) {
// CHECK: @.offload_sizes = private unnamed_addr constant [2 x i64] [i64 32, i64 8]
// CHECK: @.offload_maptypes = private unnamed_addr constant [2 x i64] [i64 [[#0x1]], i64 [[#0x4000]]]
//.
+// CHECK-60: @.offload_sizes = private unnamed_addr constant [2 x i64] [i64 32, i64 8]
+// CHECK-60: @.offload_maptypes = private unnamed_addr constant [2 x i64] [i64 [[#0x1]], i64 [[#0x4000]]]
+//.
// CHECK-LABEL: define {{[^@]+}}@_Z3fooP2S2
// CHECK-SAME: (ptr noundef [[ARR:%.*]]) #[[ATTR0:[0-9]+]] {
// CHECK: entry:
@@ -209,3 +220,173 @@ void foo(S2 *arr) {
// CHECK: omp.done:
// CHECK: ret void
//
+//
+// CHECK-60-LABEL: define {{[^@]+}}@_Z3fooP2S2
+// CHECK-60-SAME: (ptr noundef [[ARR:%.*]]) #[[ATTR0:[0-9]+]] {
+// CHECK-60: entry:
+// CHECK-60: store ptr [[ARR]], ptr [[ARR_ADDR:%.*]], align 8
+// CHECK-60: [[TMP0:%.*]] = load ptr, ptr [[ARR_ADDR]], align 8
+// CHECK-60: [[TMP1:%.*]] = load ptr, ptr [[ARR_ADDR]], align 8
+// CHECK-60: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw [[STRUCT_S2:%.*]], ptr [[TMP1]], i64 0
+// CHECK-60: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
+// CHECK-60: store ptr [[TMP0]], ptr [[TMP2]], align 8
+// CHECK-60: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
+// CHECK-60: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8
+// CHECK-60: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
+// CHECK-60: store ptr @.omp_mapper._ZTS2S2.default, ptr [[TMP4]], align 8
+// CHECK-60: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
+// CHECK-60: store ptr [[ARR_ADDR]], ptr [[TMP5]], align 8
+// CHECK-60: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1
+// CHECK-60: store ptr [[ARRAYIDX]], ptr [[TMP6]], align 8
+// CHECK-60: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
+// CHECK-60: store ptr null, ptr [[TMP7]], align 8
+// CHECK-60: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK-60: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK-60: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 2, ptr [[TMP8]], ptr [[TMP9]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr [[DOTOFFLOAD_MAPPERS]])
+// CHECK-60: ret void
+//
+//
+// CHECK-60-LABEL: define {{[^@]+}}@.omp_mapper._ZTS2S2.default
+// CHECK-60-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]], ptr noundef [[TMP2:%.*]], i64 noundef [[TMP3:%.*]], i64 noundef [[TMP4:%.*]], ptr noundef [[TMP5:%.*]]) #[[ATTR1:[0-9]+]] {
+// CHECK-60: entry:
+// CHECK-60: [[TMP6:%.*]] = udiv exact i64 [[TMP3]], 16
+// CHECK-60: [[TMP7:%.*]] = getelementptr [[STRUCT_S2:%.*]], ptr [[TMP2]], i64 [[TMP6]]
+// CHECK-60: [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP6]], 1
+// CHECK-60: [[TMP8:%.*]] = and i64 [[TMP4]], 8
+// CHECK-60: [[TMP9:%.*]] = icmp ne ptr [[TMP1]], [[TMP2]]
+// CHECK-60: [[TMP10:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP9]]
+// CHECK-60: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP8]], 0
+// CHECK-60: [[TMP11:%.*]] = and i1 [[TMP10]], [[DOTOMP_ARRAY__INIT__DELETE]]
+// CHECK-60: br i1 [[TMP11]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]]
+// CHECK-60: .omp.array..init:
+// CHECK-60: [[TMP12:%.*]] = mul nuw i64 [[TMP6]], 16
+// CHECK-60: [[TMP13:%.*]] = and i64 [[TMP4]], -4
+// CHECK-60: [[TMP14:%.*]] = or i64 [[TMP13]], 512
+// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP12]], i64 [[TMP14]], ptr [[TMP5]])
+// CHECK-60: br label [[OMP_ARRAYMAP_HEAD]]
+// CHECK-60: omp.arraymap.head:
+// 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: [[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
+// CHECK-60: [[TMP15:%.*]] = load ptr, ptr [[S1P1]], align 8
+// CHECK-60: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP15]], i32 0, i32 0
+// CHECK-60: [[S1P2:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
+// CHECK-60: [[S1P3:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
+// CHECK-60: [[TMP16:%.*]] = load ptr, ptr [[S1P3]], align 8
+// CHECK-60: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[TMP16]], i32 0, i32 1
+// CHECK-60: [[TMP17:%.*]] = getelementptr i32, ptr [[Z]], i32 1
+// CHECK-60: [[TMP18:%.*]] = ptrtoaddr ptr [[TMP17]] to i64
+// CHECK-60: [[TMP19:%.*]] = ptrtoaddr ptr [[S1P]] to i64
+// CHECK-60: [[TMP20:%.*]] = sub i64 [[TMP18]], [[TMP19]]
+// CHECK-60: [[TMP21:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]])
+// CHECK-60: [[TMP22:%.*]] = shl i64 [[TMP21]], 48
+// CHECK-60: [[TMP23:%.*]] = add nuw i64 0, [[TMP22]]
+// CHECK-60: [[TMP24:%.*]] = and i64 [[TMP4]], 3
+// CHECK-60: [[TMP25:%.*]] = icmp eq i64 [[TMP24]], 0
+// CHECK-60: br i1 [[TMP25]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]]
+// CHECK-60: omp.type.alloc:
+// CHECK-60: [[TMP26:%.*]] = and i64 [[TMP23]], -4
+// CHECK-60: br label [[OMP_TYPE_END:%.*]]
+// CHECK-60: omp.type.alloc.else:
+// CHECK-60: [[TMP27:%.*]] = icmp eq i64 [[TMP24]], 1
+// CHECK-60: br i1 [[TMP27]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]]
+// CHECK-60: omp.type.to:
+// CHECK-60: [[TMP28:%.*]] = and i64 [[TMP23]], -3
+// CHECK-60: br label [[OMP_TYPE_END]]
+// CHECK-60: omp.type.to.else:
+// CHECK-60: [[TMP29:%.*]] = icmp eq i64 [[TMP24]], 2
+// CHECK-60: br i1 [[TMP29]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]]
+// CHECK-60: omp.type.from:
+// CHECK-60: [[TMP30:%.*]] = and i64 [[TMP23]], -2
+// 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: omp.type.alloc4:
+// CHECK-60: [[TMP34:%.*]] = and i64 [[TMP31]], -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: omp.type.to6:
+// CHECK-60: [[TMP36:%.*]] = and i64 [[TMP31]], -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: omp.type.from8:
+// CHECK-60: [[TMP38:%.*]] = and i64 [[TMP31]], -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_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.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: br label [[OMP_DONE]]
+// CHECK-60: omp.done:
+// CHECK-60: ret void
+//
diff --git a/offload/test/mapping/mapper_map_always_from.c b/offload/test/mapping/mapper_map_always_from.c
new file mode 100644
index 0000000000000..848bdbb448458
--- /dev/null
+++ b/offload/test/mapping/mapper_map_always_from.c
@@ -0,0 +1,46 @@
+// 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 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
+
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <stdio.h>
+
+typedef struct {
+ int x;
+ int y;
+ int z;
+} S;
+
+#pragma omp declare mapper(default : S s) map(tofrom : s.y)
+
+S s;
+
+int main() {
+ s.y = 0;
+
+#pragma omp target enter data map(alloc : s.y)
+
+#pragma omp target map(always, from : s)
+ {
+ 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
+
+#pragma omp target exit data map(delete : s.y)
+}
diff --git a/offload/test/mapping/mapper_map_mbr_ptee_then_present_mbr_ptee.c b/offload/test/mapping/mapper_map_mbr_ptee_then_present_mbr_ptee.c
new file mode 100644
index 0000000000000..6dea1acfdd9e3
--- /dev/null
+++ b/offload/test/mapping/mapper_map_mbr_ptee_then_present_mbr_ptee.c
@@ -0,0 +1,58 @@
+// Check that it's ok to first map a member of a struct and its pointee, and
+// then do a map(present) on a mapper that maps them internally.
+//
+// FIXME: This currently run-fails, because the mapper does not yet emit
+// attach-style maps for the pointee: the combined entry over the whole struct
+// s1 (40016 bytes) triggers an "explicit extension not allowed" error against
+// the 4-byte device allocation of s1.x. Once attach-style maps are emitted for
+// the pointee, the present check should pass and the run should complete the
+// present/delete sequence below.
+// RUN: %libomptarget-compile-generic
+// RUN: %libomptarget-run-fail-generic 2>&1 \
+// RUN: | %fcheck-generic --check-prefix=CHECK
+
+#include <omp.h>
+#include <stdio.h>
+
+int x[10];
+
+typedef struct {
+ int x;
+ int dummy[10000];
+ int *p;
+} S;
+
+#pragma omp declare mapper(default : S s) map(s.x, s.p[0 : 10])
+
+S s1;
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+
+int main() {
+ s1.p = (int *)&x;
+
+#pragma omp target enter data map(alloc : s1.x, s1.p[0 : 10])
+ // EXPECTED: After mapping
+ print_status(&s1.x, "x"); // EXPECTED: x is present
+ print_status(&s1.dummy, "dummy"); // EXPECTED: dummy is not present
+ print_status(&s1.p, "p"); // EXPECTED: p is not present
+ print_status(&s1.p[0], "p[0]"); // EXPECTED: p[0] is present
+
+ // This present check currently fails (explicit extension); once attach-style
+ // maps are emitted for the pointee, it should pass.
+ // clang-format off
+ // CHECK: omptarget message: explicit extension not allowed
+ // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory
+ // clang-format on
+#pragma omp target enter data map(present, alloc : s1)
+
+#pragma omp target exit data map(delete : s1)
+ // EXPECTED: After deleting
+ print_status(&s1.x, "x"); // EXPECTED: x is not present
+ print_status(&s1.dummy, "dummy"); // EXPECTED: dummy is not present
+ print_status(&s1.p, "p"); // EXPECTED: p is not present
+ print_status(&s1.p[0], "p[0]"); // EXPECTED: p[0] is not present
+}
diff --git a/offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c b/offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c
new file mode 100644
index 0000000000000..422bfa484c2dc
--- /dev/null
+++ b/offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c
@@ -0,0 +1,62 @@
+// The mapper maps a struct member (s.x) and a pointee (s.p[0:10]). We pre-map
+// only s.x, then do map(present) on the mapper. The pointee s.p[0:10] is not
+// present, so once PRESENT is propagated to the pointee (a follow-on, at OpenMP
+// >= 6.0) the check must fail; at <= 5.2 present is not propagated, so it would
+// pass.
+//
+// FIXME: This currently run-fails at BOTH versions, because the mapper does not
+// yet emit attach-style maps for the pointee: the combined entry over the whole
+// struct s1 (40016 bytes) triggers an "explicit extension not allowed" error
+// against the 4-byte device allocation of s1.x. Once attach-style maps are
+// emitted for the pointee:
+// EXPECTED (5.2): the run completes ("done").
+// EXPECTED (6.0): the present check fails for the absent pointee s1.p[0:10].
+// RUN: %libomptarget-compile-generic -fopenmp-version=52
+// RUN: %libomptarget-run-fail-generic 2>&1 \
+// RUN: | %fcheck-generic --check-prefixes=CHECK
+// RUN: %libomptarget-compile-generic -fopenmp-version=60
+// RUN: %libomptarget-run-fail-generic 2>&1 \
+// RUN: | %fcheck-generic --check-prefixes=CHECK
+
+#include <omp.h>
+#include <stdio.h>
+
+int x[10];
+
+typedef struct {
+ int x;
+ int dummy[10000];
+ int *p;
+} S;
+
+#pragma omp declare mapper(default : S s) map(s.x, s.p[0 : 10])
+
+S s1;
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ fprintf(stderr, "%s is %spresent\n", name, present ? "" : "not ");
+}
+
+int main() {
+ s1.p = (int *)&x;
+
+ fprintf(stderr, "addr=%p, size=%ld\n", &s1.p[0], 10 * sizeof(s1.p[0]));
+
+#pragma omp target enter data map(alloc : s1.x)
+ print_status(&s1.x, "x"); // EXPECTED: x is present
+ print_status(&s1.dummy, "dummy"); // EXPECTED: dummy is not present
+ print_status(&s1.p, "p"); // EXPECTED: p is not present
+ print_status(&s1.p[0], "p[0]"); // EXPECTED: p[0] is not present
+
+#pragma omp target enter data map(present, alloc : s1)
+ // Once attach-style maps are emitted for the pointee, at 5.2 the run
+ // completes past this point; at 6.0 the present check on the absent pointee
+ // s1.p[0:10] fails here.
+ // clang-format off
+ // CHECK: omptarget message: explicit extension not allowed
+ // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory
+ // clang-format on
+
+ fprintf(stderr, "done\n");
+}
diff --git a/offload/test/mapping/mapper_map_present_ptee.c b/offload/test/mapping/mapper_map_present_ptee.c
new file mode 100644
index 0000000000000..065bbcb62e28f
--- /dev/null
+++ b/offload/test/mapping/mapper_map_present_ptee.c
@@ -0,0 +1,65 @@
+// A user-defined mapper maps s.y and, with present written directly in the
+// mapper's own clause, the pointee s.p. present in the mapper clause applies at
+// every OpenMP version -- this is not the outer-clause propagation that a
+// follow-on gates on the version.
+//
+// FIXME: This currently run-fails at BOTH the inbounds and out-of-bounds cases,
+// because the mapper does not yet emit attach-style maps for the pointee: the
+// present check runs against the whole struct s (16 bytes), which is not
+// present. Once attach-style maps are emitted for the pointee:
+// EXPECTED (inbounds): the present check passes and the run prints "333 333".
+// EXPECTED (out-of-bounds): the present check fails only for the pointee
+// region s.p[0:20] that is not present.
+// RUN: %libomptarget-compile-generic
+// RUN: %libomptarget-run-fail-generic 2>&1 \
+// RUN: | %fcheck-generic --check-prefix=CHECK
+// RUN: %libomptarget-compile-generic -DOUT_OF_BOUNDS
+// RUN: %libomptarget-run-fail-generic 2>&1 \
+// RUN: | %fcheck-generic --check-prefix=CHECK
+
+#include <stdio.h>
+
+int x[10];
+typedef struct {
+ int y;
+ int *p;
+} S;
+
+#ifdef OUT_OF_BOUNDS
+#pragma omp declare mapper(S s) map(s.y) map(present, tofrom : s.p[0 : 20])
+#else
+#pragma omp declare mapper(S s) map(s.y) map(present, tofrom : s.p[0 : 2])
+#endif
+S s;
+
+void f1() {
+ // The mapper runs here; the present check currently fails here at both cases.
+#pragma omp target update to(s)
+
+#pragma omp target data use_device_addr(s, x)
+#pragma omp target has_device_addr(s, x)
+ {
+ s.y = s.y + 222;
+ x[0] = x[0] + 222;
+ }
+}
+
+int main() {
+ x[0] = 111;
+ s.y = 111;
+ s.p = &x[0];
+
+ fprintf(stderr, "addr=%p, size=%zu\n", &s.p[0], 20 * sizeof(s.p[0]));
+
+#pragma omp target data map(from : s.y, x)
+ {
+ f1();
+ }
+
+ // EXPECTED (inbounds): 333 333
+ fprintf(stderr, "%d %d\n", x[0], s.y);
+ // clang-format off
+ // CHECK: omptarget message: device mapping required by 'present' motion modifier does not exist for host address
+ // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory
+ // clang-format on
+}
diff --git a/offload/test/mapping/mapper_map_ptee_then_present_absent_mbr.c b/offload/test/mapping/mapper_map_ptee_then_present_absent_mbr.c
index 6293c049152db..3ad7b3d97a6f1 100644
--- a/offload/test/mapping/mapper_map_ptee_then_present_absent_mbr.c
+++ b/offload/test/mapping/mapper_map_ptee_then_present_absent_mbr.c
@@ -9,17 +9,12 @@
//
// FIXME: This currently PASSES (no present-check failure) even though s.x is
// absent. The mapper emits the s.x member component WITHOUT the PRESENT bit, so
-// its absence is a silent no-op; PRESENT is only propagated to the pointee entry
-// (s.p[0:10]), which is present here. We eventually need to propagate PRESENT to
-// non-pointee (member) data as well, so that an absent member triggers the
-// present check.
+// its absence is a silent no-op. PRESENT should be propagated to the s.x member
+// entry pushed by the mapper, so that an absent member triggers the present
+// check, but it currently is not.
//
// EXPECTED: the present check should FAIL (assert / abort) because s.x is not
-// present on the device. Propagating PRESENT to member entries is blocked while
-// pointer members use PTR_AND_OBJ: a single combined entry allocates the whole
-// struct (including the pointer's own storage), so propagating PRESENT there
-// would wrongly require the pointer's pointee to be present. Enable this once
-// Clang emits attach-style maps throughout instead of PTR_AND_OBJ.
+// present on the device.
int x[10];
diff --git a/offload/test/mapping/mapper_target_update_present_ptee.c b/offload/test/mapping/mapper_target_update_present_ptee.c
index 86c16d7d139d9..10f11392ed76c 100644
--- a/offload/test/mapping/mapper_target_update_present_ptee.c
+++ b/offload/test/mapping/mapper_target_update_present_ptee.c
@@ -5,12 +5,12 @@
// RUN: %libomptarget-compile-generic -fopenmp-version=60
// RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefix=CHECK-60
-// Out-of-bounds: s.p[0:20] extends beyond the mapped region x[0:10].
-// FIXME: at OpenMP 6.0 the present modifier should be propagated to the pointee
-// s.p[0:20], so this should run-fail with a present error. It currently does
-// NOT, because the mapper drops the present modifier for the pointee;
-// propagating it is done in a follow-on. For OpenMP <= 5.2 present is
-// (correctly) not applied to the pointee, so the run must succeed there.
+// Out-of-bounds: s.p[0:20] extends beyond the mapped region x[0:10]. At OpenMP
+// 6.0 the present modifier should be propagated to the pointee s.p[0:20], so
+// this should run-fail with a present error; at <= 5.2 present is (correctly)
+// not applied to the pointee, so the run succeeds.
+// FIXME: the present modifier is not yet propagated to the pointee, so the OOB
+// run currently succeeds at 6.0 too; propagating it is done in a follow-on.
// EXPECTED (6.0): run-fail with a present-modifier error for s.p[0:20].
// RUN: %libomptarget-compile-generic -fopenmp-version=52 -DOUT_OF_BOUNDS
// RUN: %libomptarget-run-generic 2>&1 \
@@ -62,5 +62,23 @@ int main() {
// Inbounds: present check passes at both versions.
// CHECK-52: 333 333
// CHECK-60: 333 333
- printf("%d %d\n", x[0], s.y);
+ fprintf(stderr, "%d %d\n", x[0], s.y);
+
+ // Out-of-bounds: present is not yet propagated to the pointee, so both
+ // versions currently complete past the update.
+ //
+ // 5.2: present is (correctly) never applied to the pointee, so the update
+ // completes.
+ // FIXME: even at 5.2, the update should still have happened for the subset of
+ // the pointee that is present (x[0:10]), instead of being silently ignored
+ // (x[0] currently reads back as garbage). Once that is fixed:
+ // EXPECTED-52-OOB: 333 333
+ // CHECK-52-OOB: done
+ //
+ // 6.0: present should be propagated to the pointee, so the update should fail
+ // the present check for s.p[0:20]. That propagation is done in a follow-on;
+ // until then the run completes.
+ // EXPECTED-60-OOB: run-fail with a present-modifier error for s.p[0:20].
+ // CHECK-60-OOB: done
+ fprintf(stderr, "done\n");
}
>From 636995a0a03689987c87f8aaea8c49c7922717e7 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Mon, 27 Jul 2026 02:16:44 -0700
Subject: [PATCH 5/6] [OpenMP][test] Mark failing declare_mapper tests
UNSUPPORTED on intelgpu
These five declare_mapper mapping tests currently fail on Intel GPUs
(spirv64-intel): four produce a wrong result at runtime and
declare_mapper_nested_mappers.cpp hits a module-build failure. Mark them
UNSUPPORTED: intelgpu for now so the suite is green on Intel GPUs;
investigating and fixing the underlying issues is tracked as a separate
effort. The tests still run and pass on the host and other targets.
Co-Authored-By: Claude Opus 4.8 <noreply at anthropic.com>
---
offload/test/mapping/declare_mapper_nested_mappers.cpp | 2 ++
offload/test/mapping/declare_mapper_target.cpp | 2 ++
offload/test/mapping/declare_mapper_target_data.cpp | 2 ++
offload/test/mapping/declare_mapper_target_data_enter_exit.cpp | 2 ++
offload/test/mapping/declare_mapper_target_update.cpp | 2 ++
5 files changed, 10 insertions(+)
diff --git a/offload/test/mapping/declare_mapper_nested_mappers.cpp b/offload/test/mapping/declare_mapper_nested_mappers.cpp
index 5966bb6127af8..0785b1fe6f167 100644
--- a/offload/test/mapping/declare_mapper_nested_mappers.cpp
+++ b/offload/test/mapping/declare_mapper_nested_mappers.cpp
@@ -1,5 +1,7 @@
// RUN: %libomptarget-compilexx-run-and-check-generic
+// UNSUPPORTED: intelgpu
+
#include <cstdio>
#include <cstdlib>
diff --git a/offload/test/mapping/declare_mapper_target.cpp b/offload/test/mapping/declare_mapper_target.cpp
index c642949b2a367..d5feec40bb656 100644
--- a/offload/test/mapping/declare_mapper_target.cpp
+++ b/offload/test/mapping/declare_mapper_target.cpp
@@ -1,5 +1,7 @@
// RUN: %libomptarget-compile-run-and-check-generic
+// UNSUPPORTED: intelgpu
+
#include <cstdio>
#include <cstdlib>
diff --git a/offload/test/mapping/declare_mapper_target_data.cpp b/offload/test/mapping/declare_mapper_target_data.cpp
index fd2c8f4886401..e0fc332c4ac5c 100644
--- a/offload/test/mapping/declare_mapper_target_data.cpp
+++ b/offload/test/mapping/declare_mapper_target_data.cpp
@@ -1,5 +1,7 @@
// RUN: %libomptarget-compile-run-and-check-generic
+// UNSUPPORTED: intelgpu
+
#include <cstdio>
#include <cstdlib>
diff --git a/offload/test/mapping/declare_mapper_target_data_enter_exit.cpp b/offload/test/mapping/declare_mapper_target_data_enter_exit.cpp
index 245462eab5785..38e2b0ef9fd44 100644
--- a/offload/test/mapping/declare_mapper_target_data_enter_exit.cpp
+++ b/offload/test/mapping/declare_mapper_target_data_enter_exit.cpp
@@ -1,5 +1,7 @@
// RUN: %libomptarget-compile-run-and-check-generic
+// UNSUPPORTED: intelgpu
+
#include <cstdio>
#include <cstdlib>
diff --git a/offload/test/mapping/declare_mapper_target_update.cpp b/offload/test/mapping/declare_mapper_target_update.cpp
index 4e053e3e56139..978eaf6f01bd6 100644
--- a/offload/test/mapping/declare_mapper_target_update.cpp
+++ b/offload/test/mapping/declare_mapper_target_update.cpp
@@ -1,5 +1,7 @@
// RUN: %libomptarget-compile-run-and-check-generic
+// UNSUPPORTED: intelgpu
+
#include <cstdio>
#include <cstdlib>
>From ece9735c44d13dc1c22f2b2957f1328a1a488351 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Mon, 27 Jul 2026 15:24:13 -0700
Subject: [PATCH 6/6] [OpenMP][test] Add mapper ALWAYS-propagation tests with
no target construct
Companions to mapper_map_always_from.c that exercise ALWAYS propagation into a
user-defined mapper's entries without entering any target region. The device
copy is inspected directly via omp_get_mapped_ptr() and omp_target_memcpy(), so
the check covers only the data-motion performed by `target enter data`.
The Fortran variant is the first Fortran offload test that runs on the host
(x86_64) target: it needs no device kernel, so it does not depend on a device
runtime being available.
Both tests currently CHECK the pre-propagation behavior (the transfer is
suppressed for the already-mapped entry) and record the expected value in a
FIXME.
---
.../mapping/mapper_map_always_to_enter_data.c | 60 ++++++++++++++++
.../mapper-map-always-to-enter-data.f90 | 71 +++++++++++++++++++
2 files changed, 131 insertions(+)
create mode 100644 offload/test/mapping/mapper_map_always_to_enter_data.c
create mode 100644 offload/test/offloading/fortran/mapper-map-always-to-enter-data.f90
diff --git a/offload/test/mapping/mapper_map_always_to_enter_data.c b/offload/test/mapping/mapper_map_always_to_enter_data.c
new file mode 100644
index 0000000000000..185cc76a75246
--- /dev/null
+++ b/offload/test/mapping/mapper_map_always_to_enter_data.c
@@ -0,0 +1,60 @@
+// 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
+// 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
+
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+typedef struct {
+ int x;
+ int y;
+ int z;
+} S;
+
+#pragma omp declare mapper(default : S s) map(tofrom : s.y)
+
+S s;
+
+int main() {
+ int dev = omp_get_default_device();
+ int host = omp_get_initial_device();
+
+ s.y = 0;
+
+#pragma omp target enter data map(alloc : s.y)
+
+ // The device copy of s.y is uninitialized after `alloc`; set it to 0 so that
+ // a missing transfer below is distinguishable from garbage.
+ int *dev_y = (int *)omp_get_mapped_ptr(&s.y, dev);
+ omp_target_memcpy(dev_y, &s.y, sizeof(int), 0, 0, dev, host);
+
+ s.y = 111;
+
+#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.
+ 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
+
+#pragma omp target exit data map(delete : s.y)
+}
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
new file mode 100644
index 0000000000000..52e32b1d6ccad
--- /dev/null
+++ b/offload/test/offloading/fortran/mapper-map-always-to-enter-data.f90
@@ -0,0 +1,71 @@
+! 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
+! `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
+
+! REQUIRES: flang
+
+! RUN: %libomptarget-compile-fortran-generic
+! RUN: %libomptarget-run-generic | %fcheck-generic
+
+program main
+ use omp_lib
+ use iso_c_binding
+ implicit none
+
+ type :: s_type
+ integer(c_int) :: x
+ integer(c_int) :: y
+ integer(c_int) :: z
+ end type s_type
+
+ !$omp declare mapper(s_type :: s) map(tofrom: s%y)
+
+ type(s_type), target :: s
+ integer(c_int), target :: dev_y_val
+ integer :: dev, host, rc
+ type(c_ptr) :: dev_y
+
+ dev = omp_get_default_device()
+ host = omp_get_initial_device()
+
+ s%y = 0
+
+ !$omp target enter data map(alloc: s%y)
+
+ ! The device copy of s%y is uninitialized after `alloc`; set it to 0 so that
+ ! a missing transfer below is distinguishable from garbage.
+ dev_y = omp_get_mapped_ptr(c_loc(s%y), dev)
+ rc = omp_target_memcpy(dev_y, c_loc(s%y), int(c_sizeof(s%y), c_size_t), &
+ 0_c_size_t, 0_c_size_t, dev, host)
+
+ s%y = 111
+
+ !$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.
+ 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
+
+ !$omp target exit data map(delete: s%y)
+end program main
More information about the cfe-commits
mailing list