Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 3 commits into from
Aug 14, 2025
Merged

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
21123df
[NFC][Offload] Add missing maps to OpenMP offloading tests.
abhinavgaba Aug 6, 2025
6ebf7a2
[NFC][OpenMP] Fix tests mapping pointees, but expecting pointers to b…
abhinavgaba Aug 11, 2025
8a94ab8
Apply clang-format.
abhinavgaba Aug 11, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 2 additions & 1 deletion offload/test/mapping/data_member_ref.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]) \
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

map(from : D1, D2)
{
// CHECK: Device 123456.
D1 = *Baz.VRef.Data;
Expand Down
4 changes: 2 additions & 2 deletions offload/test/mapping/declare_mapper_nested_default_mappers.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
4 changes: 2 additions & 2 deletions offload/test/mapping/declare_mapper_nested_mappers.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 1 addition & 1 deletion offload/test/mapping/ptr_and_obj_motion.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
20 changes: 11 additions & 9 deletions offload/test/mapping/target_derefence_array_pointrs.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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));
Expand All @@ -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
Expand All @@ -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)));
Expand All @@ -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
Expand Down
5 changes: 3 additions & 2 deletions offload/test/mapping/target_has_device_addr.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]) \
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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 identifier.

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).

use_device_addr(xpp[1])
#pragma omp target has_device_addr(xpp[1])
{
xpp[1][1] = 222;
// CHECK: 222
Expand Down
Loading