Skip to content

Commit 6716366

Browse files
abhinavgabaabhinavgaba
committed
[NFC][OpenMP][Offload] Add tests for use_device_ptr(fb_preserve/nullify).
The fallback modifiers are currently part of OpenMP 6.1. The tests mostly fail for now. The associated libomptarget and clang parsing/sema changes are in #169438, #169603 and #170578, with clang codegen to follow.
1 parent 7ba7101 commit 6716366

12 files changed

+317
-1
lines changed

‎offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback.cpp‎

Lines changed: 34 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,34 @@
1+
// RUN: %libomptarget-compilexx-generic
2+
// RUN: %libomptarget-run-generic 2>&1 \
3+
// RUN: | %fcheck-generic
4+
// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
5+
// RUN: | %fcheck-generic
6+
7+
// Test that when a use_device_ptr lookup fails, the
8+
// privatized pointer retains its original value by
9+
// default.
10+
//
11+
// This is necessary because we must assume that the
12+
// pointee is device-accessible, even if it was not
13+
// previously mapped.
14+
15+
// XFAIL: *
16+
17+
#include <stdio.h>
18+
19+
int x = 0;
20+
21+
struct ST {
22+
int *a = &x;
23+
24+
void f1() {
25+
printf("%p\n", a); // CHECK: 0x[[#%x,ADDR:]]
26+
#pragma omp target data use_device_ptr(a)
27+
printf("%p\n", a); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
28+
}
29+
};
30+
31+
int main() {
32+
ST s;
33+
s.f1();
34+
}

‎offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp‎

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,30 @@
1+
// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
2+
// RUN: %libomptarget-run-generic 2>&1 \
3+
// RUN: | %fcheck-generic -check-prefixes=CHECK,OFFLOAD
4+
// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
5+
// RUN: | %fcheck-generic -check-prefixes=CHECK,NOOFFLOAD
6+
7+
// Test that when a use_device_ptr lookup fails, the
8+
// privatized pointer is set to null because of fb_nullify.
9+
10+
// XFAIL: *
11+
12+
#include <stdio.h>
13+
14+
int x = 0;
15+
16+
struct ST {
17+
int *a = &x;
18+
19+
void f1() {
20+
printf("%p\n", a); // CHECK: 0x[[#%x,ADDR:]]
21+
#pragma omp target data use_device_ptr(fb_nullify : a)
22+
printf("%p\n", a); // OFFLOAD-NEXT: (nil)
23+
// NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
24+
}
25+
};
26+
27+
int main() {
28+
ST s;
29+
s.f1();
30+
}

‎offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_preserve.cpp‎

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,30 @@
1+
// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
2+
// RUN: %libomptarget-run-generic 2>&1 \
3+
// RUN: | %fcheck-generic
4+
// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
5+
// RUN: | %fcheck-generic
6+
7+
// Test that when a use_device_ptr lookup fails, the
8+
// privatized pointer retains its original value
9+
// because of fb_preserve.
10+
11+
// XFAIL: *
12+
13+
#include <stdio.h>
14+
15+
int x = 0;
16+
17+
struct ST {
18+
int *a = &x;
19+
20+
void f1() {
21+
printf("%p\n", a); // CHECK: 0x[[#%x,ADDR:]]
22+
#pragma omp target data use_device_ptr(fb_preserve : a)
23+
printf("%p\n", a); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
24+
}
25+
};
26+
27+
int main() {
28+
ST s;
29+
s.f1();
30+
}

‎offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback.cpp‎

Lines changed: 35 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,35 @@
1+
// RUN: %libomptarget-compilexx-generic
2+
// RUN: %libomptarget-run-generic 2>&1 \
3+
// RUN: | %fcheck-generic
4+
// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
5+
// RUN: | %fcheck-generic
6+
7+
// Test that when a use_device_ptr lookup fails, the
8+
// privatized pointer retains its original value by
9+
// default.
10+
//
11+
// This is necessary because we must assume that the
12+
// pointee is device-accessible, even if it was not
13+
// previously mapped.
14+
15+
// XFAIL: *
16+
17+
#include <stdio.h>
18+
19+
int x = 0;
20+
int *y = &x;
21+
22+
struct ST {
23+
int *&b = y;
24+
25+
void f2() {
26+
printf("%p\n", b); // CHECK: 0x[[#%x,ADDR:]]
27+
#pragma omp target data use_device_ptr(b)
28+
printf("%p\n", b); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
29+
}
30+
};
31+
32+
int main() {
33+
ST s;
34+
s.f2();
35+
}

‎offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp‎

Lines changed: 31 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,31 @@
1+
// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
2+
// RUN: %libomptarget-run-generic 2>&1 \
3+
// RUN: | %fcheck-generic -check-prefixes=CHECK,OFFLOAD
4+
// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
5+
// RUN: | %fcheck-generic -check-prefixes=CHECK,NOOFFLOAD
6+
7+
// Test that when a use_device_ptr lookup fails, the
8+
// privatized pointer is set to null because of fb_nullify.
9+
10+
// XFAIL: *
11+
12+
#include <stdio.h>
13+
14+
int x = 0;
15+
int *y = &x;
16+
17+
struct ST {
18+
int *&b = y;
19+
20+
void f2() {
21+
printf("%p\n", b); // CHECK: 0x[[#%x,ADDR:]]
22+
#pragma omp target data use_device_ptr(fb_nullify : b)
23+
printf("%p\n", b); // OFFLOAD-NEXT: (nil)
24+
// NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
25+
}
26+
};
27+
28+
int main() {
29+
ST s;
30+
s.f2();
31+
}

‎offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_preserve.cpp‎

Lines changed: 31 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,31 @@
1+
// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
2+
// RUN: %libomptarget-run-generic 2>&1 \
3+
// RUN: | %fcheck-generic
4+
// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
5+
// RUN: | %fcheck-generic
6+
7+
// Test that when a use_device_ptr lookup fails, the
8+
// privatized pointer retains its original value
9+
// because of fb_preserve.
10+
11+
// XFAIL: *
12+
13+
#include <stdio.h>
14+
15+
int x = 0;
16+
int *y = &x;
17+
18+
struct ST {
19+
int *&b = y;
20+
21+
void f2() {
22+
printf("%p\n", b); // CHECK: 0x[[#%x,ADDR:]]
23+
#pragma omp target data use_device_ptr(fb_preserve : b)
24+
printf("%p\n", b); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
25+
}
26+
};
27+
28+
int main() {
29+
ST s;
30+
s.f2();
31+
}

offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c renamed to offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.cpp

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,8 @@
1-
// RUN: %libomptarget-compilexx-run-and-check-generic
1+
// RUN: %libomptarget-compilexx-generic
2+
// RUN: %libomptarget-run-generic 2>&1 \
3+
// RUN: | %fcheck-generic
4+
// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
5+
// RUN: | %fcheck-generic
26

37
// Test that when a use_device_ptr lookup fails, the
48
// privatized pointer retains its original value by

‎offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_nullify.cpp‎

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,23 @@
1+
// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
2+
// RUN: %libomptarget-run-generic 2>&1 \
3+
// RUN: | %fcheck-generic -check-prefixes=CHECK,OFFLOAD
4+
// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
5+
// RUN: | %fcheck-generic -check-prefixes=CHECK,NOOFFLOAD
6+
7+
// Test that when a use_device_ptr lookup fails, the
8+
// privatized pointer is set to null because of fb_nullify.
9+
10+
// XFAIL: *
11+
12+
#include <stdio.h>
13+
int x;
14+
int *xp = &x;
15+
16+
void f1() {
17+
printf("%p\n", xp); // CHECK: 0x[[#%x,ADDR:]]
18+
#pragma omp target data use_device_ptr(fb_nullify : xp)
19+
printf("%p\n", xp); // OFFLOAD-NEXT: (nil)
20+
// NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
21+
}
22+
23+
int main() { f1(); }

‎offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_preserve.cpp‎

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,23 @@
1+
// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
2+
// RUN: %libomptarget-run-generic 2>&1 \
3+
// RUN: | %fcheck-generic
4+
// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
5+
// RUN: | %fcheck-generic
6+
7+
// Test that when a use_device_ptr lookup fails, the
8+
// privatized pointer retains its original value
9+
// because of fb_preserve.
10+
11+
// XFAIL: *
12+
13+
#include <stdio.h>
14+
int x;
15+
int *xp = &x;
16+
17+
void f1() {
18+
printf("%p\n", xp); // CHECK: 0x[[#%x,ADDR:]]
19+
#pragma omp target data use_device_ptr(fb_preserve : xp)
20+
printf("%p\n", xp); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
21+
}
22+
23+
int main() { f1(); }

‎offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback.cpp‎

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,26 @@
1+
// RUN: %libomptarget-compilexx-generic
2+
// RUN: %libomptarget-run-generic 2>&1 \
3+
// RUN: | %fcheck-generic
4+
// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
5+
// RUN: | %fcheck-generic
6+
7+
// Test that when a use_device_ptr lookup fails, the
8+
// privatized pointer retains its original value by
9+
// default.
10+
//
11+
// This is necessary because we must assume that the
12+
// pointee is device-accessible, even if it was not
13+
// previously mapped.
14+
15+
#include <stdio.h>
16+
int x;
17+
int *xp = &x;
18+
int *&xpr = xp;
19+
20+
void f2() {
21+
printf("%p\n", xpr); // CHECK: 0x[[#%x,ADDR:]]
22+
#pragma omp target data use_device_ptr(xpr)
23+
printf("%p\n", xpr); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
24+
}
25+
26+
int main() { f2(); }

0 commit comments

Comments
 (0)