llvm-project

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

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

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

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

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
}
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
}
94
int main() {
95
  foo();
96
  bar();
97
  moo();
98
  zoo();
99
  xoo();
100
  return 0;
101
}
102