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
49void 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, format: "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, format: "error: arr[%d]=%d\n", i, arr[i]);
74 }
75}
76
77void 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, format: "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, format: "error: arr[%d]=%d\n", i, arr[i]);
104 } else if (arr[i] != 99) {
105 fprintf(stderr, format: "error: arr[%d]=%d\n", i, arr[i]);
106 }
107 }
108}
109
110int main() {
111 check_not_present();
112 check_is_present();
113 return 0;
114}
115

source code of offload/test/mapping/target_data_array_extension_at_exit.c