[NFC][Offload] Add missing maps to OpenMP offloading tests. #153103

abhinavgaba · 2025-08-11T23:35:19Z

A few tests were only mapping a pointee, like: map(pp[0][0]), on an int**pp, but expecting the pointers, like pp, pp[0] to also be mapped, which is incorrect.

This change fixes six such tests.

We had cases were we were mapping expressions requiring 2/3 level pointer dereference, and expecting the full chain to be mapped, like: ```c // assuming no previous mapping #pragma omp target map(p->q->r) printf("%d\n", p->q->r); // May lead to a seg-fault ``` That does not conform to the OpenMP spec. The proper mapping that would guarantee the usage in the target region to succeed would be something like: ```c #pragma omp target map(p->q->r) map(alloc:p->q) ```

…e mapped. A few tests were only mapping a pointee, like: `map(spp[0][0])`, but expecting the pointers, like `spp`, `spp[0]` to also be mapped, which is incorrect. This change fixes four such tests.

github-actions · 2025-08-11T23:37:41Z

✅ With the latest revision this PR passed the C/C++ code formatter.

abhinavgaba · 2025-08-12T00:01:29Z

offload/test/mapping/target_has_device_addr.c

@@ -66,8 +66,9 @@ void zoo() {
  short **xpp = &xp[0];

  x[1] = 111;
-#pragma omp target data map(tofrom : xpp[1][1]) use_device_addr(xpp[1][1])
-#pragma omp target has_device_addr(xpp[1][1])
+#pragma omp target data map(tofrom : xpp[1][1]) map(xpp[1])                    \


use_device_addr(xpp[1][1:1]) is non-conforming from OpenMP's perspective, as it violates this restriction from OpenMP 5.1 2.14.2:

If a list item in a use_device_addr clause is an array section, the base expression must be a base language identiﬁer.

The same is present in OpenMP 6.0 as well.

It is supposed to apply to array-elements/pointer-dereferences etc. as well. That is to be clarified in a future version of OpenMP).

abhinavgaba · 2025-08-12T00:04:14Z

offload/test/mapping/data_member_ref.cpp

@@ -60,7 +60,8 @@ int main() {
  printf("Host %d %d.\n", Bar.VRef.Data, V.Data);
  // CHECK: Host 123456.
  printf("Host %d.\n", *Baz.VRef.Data);
-#pragma omp target map(*Baz.VRef.Data) map(from : D1, D2)
+#pragma omp target map(Baz.VRef.Data) map(*Baz.VRef.Data) map(V1.Data[0 : 0])  \


For pointer-attachment on V1.Data to trigger, we need to map(V1.Data[0:0]). Otherwise, the implicit map(V1) will cause V1.Data to retain the host pointer in the target region.

llvmbot · 2025-08-12T00:05:39Z

@llvm/pr-subscribers-offload

Author: Abhinav Gaba (abhinavgaba)

Changes

A few tests were only mapping a pointee, like: map(pp[0][0]), on an int**pp, but expecting the pointers, like pp, pp[0] to also be mapped, which is incorrect.

This change fixes six such tests.

Full diff: https://github.com/llvm/llvm-project/pull/153103.diff

6 Files Affected:

(modified) offload/test/mapping/data_member_ref.cpp (+2-1)
(modified) offload/test/mapping/declare_mapper_nested_default_mappers.cpp (+2-2)
(modified) offload/test/mapping/declare_mapper_nested_mappers.cpp (+2-2)
(modified) offload/test/mapping/ptr_and_obj_motion.c (+1-1)
(modified) offload/test/mapping/target_derefence_array_pointrs.cpp (+11-9)
(modified) offload/test/mapping/target_has_device_addr.c (+3-2)

diff --git a/offload/test/mapping/data_member_ref.cpp b/offload/test/mapping/data_member_ref.cpp
index fdb8abcaa6506..7947a62c169f4 100644
--- a/offload/test/mapping/data_member_ref.cpp
+++ b/offload/test/mapping/data_member_ref.cpp
@@ -60,7 +60,8 @@ int main() {
   printf("Host %d %d.\n", Bar.VRef.Data, V.Data);
   // CHECK: Host 123456.
   printf("Host %d.\n", *Baz.VRef.Data);
-#pragma omp target map(*Baz.VRef.Data) map(from : D1, D2)
+#pragma omp target map(Baz.VRef.Data) map(*Baz.VRef.Data) map(V1.Data[0 : 0])  \
+    map(from : D1, D2)
   {
     // CHECK: Device 123456.
     D1 = *Baz.VRef.Data;
diff --git a/offload/test/mapping/declare_mapper_nested_default_mappers.cpp b/offload/test/mapping/declare_mapper_nested_default_mappers.cpp
index c6c5657ae6166..45fd042aedb0e 100644
--- a/offload/test/mapping/declare_mapper_nested_default_mappers.cpp
+++ b/offload/test/mapping/declare_mapper_nested_default_mappers.cpp
@@ -44,8 +44,8 @@ int main() {
 
   int spp00fa = -1, spp00fca = -1, spp00fb_r = -1;
   __intptr_t p = reinterpret_cast<__intptr_t>(&x[0]);
-#pragma omp target map(tofrom: spp[0][0]) firstprivate(p)                           \
-                   map(from: spp00fa, spp00fca, spp00fb_r)
+#pragma omp target map(tofrom : spp[0][0]) map(alloc : spp[0]) firstprivate(p) \
+    map(from : spp00fa, spp00fca, spp00fb_r)
   {
     spp00fa = spp[0][0].f.a;
     spp00fca = spp[0][0].f.c.a;
diff --git a/offload/test/mapping/declare_mapper_nested_mappers.cpp b/offload/test/mapping/declare_mapper_nested_mappers.cpp
index a9e3f05e0f5fd..a59ed6980ec4c 100644
--- a/offload/test/mapping/declare_mapper_nested_mappers.cpp
+++ b/offload/test/mapping/declare_mapper_nested_mappers.cpp
@@ -42,8 +42,8 @@ int main() {
   int spp00fa = -1, spp00fb_r = -1, spp00fg1 = -1, spp00fg_r = -1;
   __intptr_t p = reinterpret_cast<__intptr_t>(&x[0]),
              p1 = reinterpret_cast<__intptr_t>(&y[0]);
-#pragma omp target map(tofrom : spp[0][0]) firstprivate(p, p1)                  \
-                   map(from: spp00fa, spp00fb_r, spp00fg1, spp00fg_r)
+#pragma omp target map(tofrom : spp[0][0]) map(alloc : spp[0])                 \
+    firstprivate(p, p1) map(from : spp00fa, spp00fb_r, spp00fg1, spp00fg_r)
   {
     spp00fa = spp[0][0].f.a;
     spp00fb_r = spp[0][0].f.b == reinterpret_cast<void *>(p) ? 1 : 0;
diff --git a/offload/test/mapping/ptr_and_obj_motion.c b/offload/test/mapping/ptr_and_obj_motion.c
index 8fa2c9865b4ac..a94c07aadc1bc 100644
--- a/offload/test/mapping/ptr_and_obj_motion.c
+++ b/offload/test/mapping/ptr_and_obj_motion.c
@@ -17,7 +17,7 @@ void init(double vertexx[]) {
 }
 
 void change(DV *dvptr) {
-#pragma omp target map(dvptr->dataptr[0 : 100])
+#pragma omp target map(dvptr->dataptr[0 : 100]) map(alloc : dvptr -> dataptr)
   {
     printf("In change: %lf, expected 77.0\n", dvptr->dataptr[77]);
     dvptr->dataptr[77] += 1.0;
diff --git a/offload/test/mapping/target_derefence_array_pointrs.cpp b/offload/test/mapping/target_derefence_array_pointrs.cpp
index a6dd4069a8f58..d213c87443634 100644
--- a/offload/test/mapping/target_derefence_array_pointrs.cpp
+++ b/offload/test/mapping/target_derefence_array_pointrs.cpp
@@ -18,23 +18,24 @@ void foo(int **t1d) {
 
   for (j = 0; j < 3; j++)
     (*t1d)[j] = 0;
-#pragma omp target map(tofrom : (*t1d)[0 : 3])
+#pragma omp target map(tofrom : (*t1d)[0 : 3]) map(alloc : *t1d)
   { (*t1d)[1] = 1; }
   // CHECK: 1
   printf("%d\n", (*t1d)[1]);
-#pragma omp target map(tofrom : (**t2d)[0 : 3])
+#pragma omp target map(tofrom : (**t2d)[0 : 3]) map(alloc : **t2d, *t2d)
   { (**t2d)[1] = 2; }
   // CHECK: 2
   printf("%d\n", (**t2d)[1]);
-#pragma omp target map(tofrom : (***t3d)[0 : 3])
+#pragma omp target map(tofrom : (***t3d)[0 : 3])                               \
+    map(alloc : ***t3d, **t3d, *t3d)
   { (***t3d)[1] = 3; }
   // CHECK: 3
   printf("%d\n", (***t3d)[1]);
-#pragma omp target map(tofrom : (**t1d))
+#pragma omp target map(tofrom : (**t1d)) map(alloc : *t1d)
   { (*t1d)[0] = 4; }
   // CHECK: 4
   printf("%d\n", (*t1d)[0]);
-#pragma omp target map(tofrom : (*(*(t1d + a) + b)))
+#pragma omp target map(tofrom : (*(*(t1d + a) + b))) map(to : *(t1d + a))
   { *(*(t1d + a) + b) = 5; }
   // CHECK: 5
   printf("%d\n", *(*(t1d + a) + b));
@@ -49,7 +50,7 @@ void bar() {
   for (int i = 0; i < 3; i++) {
     (**a)[1] = i;
   }
-#pragma omp target map((**a)[ : 3])
+#pragma omp target map((**a)[ : 3]) map(alloc : **a, *a)
   {
     (**a)[1] = 6;
     // CHECK: 6
@@ -73,7 +74,8 @@ void zoo(int **f, SSA *sa) {
   *(f + sa->i + 1) = t;
   *(sa->sa->i + *(f + sa->i + 1)) = 4;
   printf("%d\n", *(sa->sa->i + *(1 + sa->i + f)));
-#pragma omp target map(sa, *(sa->sa->i + *(1 + sa->i + f)))
+#pragma omp target map(*(sa->sa->i + *(1 + sa->i + f))) map(alloc : sa->sa)    \
+    map(to : sa->i) map(to : sa->sa->i) map(to : *(1 + sa->i + f))
   { *(sa->sa->i + *(1 + sa->i + f)) = 7; }
   // CHECK: 7
   printf("%d\n", *(sa->sa->i + *(1 + sa->i + f)));
@@ -87,13 +89,13 @@ void xoo() {
 
 void yoo(int **x) {
   *x = (int *)malloc(2 * sizeof(int));
-#pragma omp target map(**x)
+#pragma omp target map(**x) map(alloc : *x)
   {
     **x = 8;
     // CHECK: 8
     printf("%d\n", **x);
   }
-#pragma omp target map(*(*x + 1))
+#pragma omp target map(*(*x + 1)) map(alloc : *x)
   {
     *(*x + 1) = 9;
     // CHECK: 9
diff --git a/offload/test/mapping/target_has_device_addr.c b/offload/test/mapping/target_has_device_addr.c
index e8bfff868c7ed..f238832c44054 100644
--- a/offload/test/mapping/target_has_device_addr.c
+++ b/offload/test/mapping/target_has_device_addr.c
@@ -66,8 +66,9 @@ void zoo() {
   short **xpp = &xp[0];
 
   x[1] = 111;
-#pragma omp target data map(tofrom : xpp[1][1]) use_device_addr(xpp[1][1])
-#pragma omp target has_device_addr(xpp[1][1])
+#pragma omp target data map(tofrom : xpp[1][1]) map(xpp[1])                    \
+    use_device_addr(xpp[1])
+#pragma omp target has_device_addr(xpp[1])
   {
     xpp[1][1] = 222;
     // CHECK: 222

abhinavgaba · 2025-08-12T00:07:03Z

@dreachem, please check if the updates look ok to you.

dreachem

Test changes look good to me.

abhinavgaba added 2 commits August 11, 2025 16:13

[NFC][OpenMP] Fix tests mapping pointees, but expecting pointers to b…

6ebf7a2

…e mapped. A few tests were only mapping a pointee, like: `map(spp[0][0])`, but expecting the pointers, like `spp`, `spp[0]` to also be mapped, which is incorrect. This change fixes four such tests.

Apply clang-format.

8a94ab8

abhinavgaba commented Aug 12, 2025

View reviewed changes

abhinavgaba marked this pull request as ready for review August 12, 2025 00:05

llvmbot added the offload label Aug 12, 2025

abhinavgaba requested a review from alexey-bataev August 12, 2025 00:05

dreachem reviewed Aug 12, 2025

View reviewed changes

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

[NFC][Offload] Add missing maps to OpenMP offloading tests. #153103

[NFC][Offload] Add missing maps to OpenMP offloading tests. #153103

Uh oh!

abhinavgaba commented Aug 11, 2025 •

edited

Loading

Uh oh!

github-actions bot commented Aug 11, 2025 •

edited

Loading

Uh oh!

abhinavgaba Aug 12, 2025

Uh oh!

abhinavgaba Aug 12, 2025

Uh oh!

llvmbot commented Aug 12, 2025

Uh oh!

abhinavgaba commented Aug 12, 2025

Uh oh!

dreachem left a comment

Uh oh!

Uh oh!

[NFC][Offload] Add missing maps to OpenMP offloading tests. #153103

Are you sure you want to change the base?

[NFC][Offload] Add missing maps to OpenMP offloading tests. #153103

Uh oh!

Conversation

abhinavgaba commented Aug 11, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

github-actions bot commented Aug 11, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

abhinavgaba Aug 12, 2025

Choose a reason for hiding this comment

Uh oh!

abhinavgaba Aug 12, 2025

Choose a reason for hiding this comment

Uh oh!

llvmbot commented Aug 12, 2025

Uh oh!

abhinavgaba commented Aug 12, 2025

Uh oh!

dreachem left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

abhinavgaba commented Aug 11, 2025 •

edited

Loading

github-actions bot commented Aug 11, 2025 •

edited

Loading