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
15void 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, format: "Failure %p %p \n", A_dev, A);
26 else
27 fprintf(stderr, format: "Success\n");
28}
29
30void 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(format: "%d %p\n", xp[1], &xp[1]);
41 }
42 // CHECK: 222
43 printf(format: "%d %p\n", xp[1], &xp[1]);
44}
45
46void moo() {
47 short *b = malloc(size: 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(format: "%hd %p %p %p\n", b[1], b, &b[1], &b);
57 }
58 // CHECK: 222
59 printf(format: "%hd %p %p %p\n", b[1], b, &b[1], &b);
60}
61
62void 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(format: "%d %p %p\n", xpp[1][1], xpp[1], &xpp[1][1]);
75 }
76 // CHECK: 222
77 printf(format: "%d %p %p\n", xpp[1][1], xpp[1], &xpp[1][1]);
78}
79void 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(format: "%hd %hd %p %p %p\n", a[1], b[1], &a, b, &b);
90 }
91 // CHECK:111
92 printf(format: "%hd %hd %p %p %p\n", a[1], b[1], &a, b, &b); // 111 111 p1d p2d p3d
93}
94int main() {
95 foo();
96 bar();
97 moo();
98 zoo();
99 xoo();
100 return 0;
101}
102

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