xref: /llvm-project/offload/test/mapping/target_has_device_addr.c (revision 330d8983d25d08580fc1642fea48b2473f47a9da)
1 // RUN: %libomptarget-compile-generic -fopenmp-version=51
2 // RUN: %libomptarget-run-generic 2>&1 \
3 // RUN: | %fcheck-generic
4 
5 // UNSUPPORTED: amdgcn-amd-amdhsa
6 
7 #include <omp.h>
8 #include <stdio.h>
9 #include <stdlib.h>
10 
11 #define N 1024
12 #define FROM 64
13 #define LENGTH 128
14 
foo()15 void foo() {
16   const int device_id = omp_get_default_device();
17   float *A;
18   A = (float *)omp_target_alloc((FROM + LENGTH) * sizeof(float), device_id);
19 
20   float *A_dev = NULL;
21 #pragma omp target has_device_addr(A[FROM : LENGTH]) map(A_dev)
22   { A_dev = A; }
23   // CHECK: Success
24   if (A_dev == NULL || A_dev != A)
25     fprintf(stderr, "Failure %p %p \n", A_dev, A);
26   else
27     fprintf(stderr, "Success\n");
28 }
29 
bar()30 void bar() {
31   short x[10];
32   short *xp = &x[0];
33 
34   x[1] = 111;
35 #pragma omp target data map(tofrom : xp[0 : 2]) use_device_addr(xp[0 : 2])
36 #pragma omp target has_device_addr(xp[0 : 2])
37   {
38     xp[1] = 222;
39     // CHECK: 222
40     printf("%d %p\n", xp[1], &xp[1]);
41   }
42   // CHECK: 222
43   printf("%d %p\n", xp[1], &xp[1]);
44 }
45 
moo()46 void moo() {
47   short *b = malloc(sizeof(short));
48   b = b - 1;
49 
50   b[1] = 111;
51 #pragma omp target data map(tofrom : b[1]) use_device_addr(b[1])
52 #pragma omp target has_device_addr(b[1])
53   {
54     b[1] = 222;
55     // CHECK: 222
56     printf("%hd %p %p %p\n", b[1], b, &b[1], &b);
57   }
58   // CHECK: 222
59   printf("%hd %p %p %p\n", b[1], b, &b[1], &b);
60 }
61 
zoo()62 void zoo() {
63   short x[10];
64   short *(xp[10]);
65   xp[1] = &x[0];
66   short **xpp = &xp[0];
67 
68   x[1] = 111;
69 #pragma omp target data map(tofrom : xpp[1][1]) use_device_addr(xpp[1][1])
70 #pragma omp target has_device_addr(xpp[1][1])
71   {
72     xpp[1][1] = 222;
73     // CHECK: 222
74     printf("%d %p %p\n", xpp[1][1], xpp[1], &xpp[1][1]);
75   }
76   // CHECK: 222
77   printf("%d %p %p\n", xpp[1][1], xpp[1], &xpp[1][1]);
78 }
xoo()79 void xoo() {
80   short a[10], b[10];
81   a[1] = 111;
82   b[1] = 111;
83 #pragma omp target data map(to : a[0 : 2], b[0 : 2]) use_device_addr(a, b)
84 #pragma omp target has_device_addr(a) has_device_addr(b[0])
85   {
86     a[1] = 222;
87     b[1] = 222;
88     // CHECK: 222 222
89     printf("%hd %hd %p %p %p\n", a[1], b[1], &a, b, &b);
90   }
91   // CHECK:111
92   printf("%hd %hd %p %p %p\n", a[1], b[1], &a, b, &b); // 111 111 p1d p2d p3d
93 }
main()94 int main() {
95   foo();
96   bar();
97   moo();
98   zoo();
99   xoo();
100   return 0;
101 }
102