1 // --------------------------------------------------
2 // Check extends before
3 // --------------------------------------------------
4
5 // RUN: %libomptarget-compile-generic \
6 // RUN: -fopenmp-version=51 -DEXTENDS=BEFORE
7 // RUN: %libomptarget-run-generic 2>&1 \
8 // RUN: | %fcheck-generic
9
10 // --------------------------------------------------
11 // Check extends after
12 // --------------------------------------------------
13
14 // RUN: %libomptarget-compile-generic \
15 // RUN: -fopenmp-version=51 -DEXTENDS=AFTER
16 // RUN: %libomptarget-run-generic 2>&1 \
17 // RUN: | %fcheck-generic
18
19
20 // END.
21
22 #include <stdio.h>
23
24 #define BEFORE 0
25 #define AFTER 1
26
27 #define SIZE 100
28
29 #if EXTENDS == BEFORE
30 # define SMALL_BEG (SIZE-2)
31 # define SMALL_END SIZE
32 # define LARGE_BEG 0
33 # define LARGE_END SIZE
34 #elif EXTENDS == AFTER
35 # define SMALL_BEG 0
36 # define SMALL_END 2
37 # define LARGE_BEG 0
38 # define LARGE_END SIZE
39 #else
40 # error EXTENDS undefined
41 #endif
42
43 #define SMALL SMALL_BEG:(SMALL_END-SMALL_BEG)
44 #define LARGE LARGE_BEG:(LARGE_END-LARGE_BEG)
45
check_not_present()46 void check_not_present() {
47 int arr[SIZE];
48
49 for (int i = 0; i < SIZE; ++i)
50 arr[i] = 99;
51
52 // CHECK-LABEL: checking not present
53 fprintf(stderr, "checking not present\n");
54
55 // arr[LARGE] isn't (fully) present at the end of the target data region, so
56 // the device-to-host transfer should not be performed, or it might fail.
57 #pragma omp target data map(tofrom: arr[LARGE])
58 {
59 #pragma omp target exit data map(delete: arr[LARGE])
60 #pragma omp target enter data map(alloc: arr[SMALL])
61 #pragma omp target map(alloc: arr[SMALL])
62 for (int i = SMALL_BEG; i < SMALL_END; ++i)
63 arr[i] = 88;
64 }
65
66 // CHECK-NOT: Libomptarget
67 // CHECK-NOT: error
68 for (int i = 0; i < SIZE; ++i) {
69 if (arr[i] != 99)
70 fprintf(stderr, "error: arr[%d]=%d\n", i, arr[i]);
71 }
72 }
73
check_is_present()74 void check_is_present() {
75 int arr[SIZE];
76
77 for (int i = 0; i < SIZE; ++i)
78 arr[i] = 99;
79
80 // CHECK-LABEL: checking is present
81 fprintf(stderr, "checking is present\n");
82
83 // arr[SMALL] is (fully) present at the end of the target data region, and the
84 // device-to-host transfer should be performed only for it even though more
85 // of the array is then present.
86 #pragma omp target data map(tofrom: arr[SMALL])
87 {
88 #pragma omp target exit data map(delete: arr[SMALL])
89 #pragma omp target enter data map(alloc: arr[LARGE])
90 #pragma omp target map(alloc: arr[LARGE])
91 for (int i = LARGE_BEG; i < LARGE_END; ++i)
92 arr[i] = 88;
93 }
94
95 // CHECK-NOT: Libomptarget
96 // CHECK-NOT: error
97 for (int i = 0; i < SIZE; ++i) {
98 if (SMALL_BEG <= i && i < SMALL_END) {
99 if (arr[i] != 88)
100 fprintf(stderr, "error: arr[%d]=%d\n", i, arr[i]);
101 } else if (arr[i] != 99) {
102 fprintf(stderr, "error: arr[%d]=%d\n", i, arr[i]);
103 }
104 }
105 }
106
main()107 int main() {
108 check_not_present();
109 check_is_present();
110 return 0;
111 }
112