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 
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 
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 
107 int main() {
108   check_not_present();
109   check_is_present();
110   return 0;
111 }
112