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