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