llvm-project

1
// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
2
// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
3
// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
4
// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
5
// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
6

7
// REQUIRES: unified_shared_memory
8
// UNSUPPORTED: amdgcn-amd-amdhsa
9

10
#pragma omp declare target
11
int *ptr1;
12
#pragma omp end declare target
13
int a[10];
14

15
#include <stdio.h>
16
#include <stdlib.h>
17
int main() {
18
  ptr1 = (int *)malloc(sizeof(int) * 100);
19
  int *ptr2;
20
  ptr2 = (int *)malloc(sizeof(int) * 100);
21
#pragma omp target map(ptr1, ptr1[ : 100])
22
  { ptr1[1] = 6; }
23
  // CHECK: 6
24
  printf(" %d \n", ptr1[1]);
25
#pragma omp target data map(ptr1[ : 5])
26
  {
27
#pragma omp target map(ptr1[2], ptr1, ptr1[3]) map(ptr2, ptr2[2])
28
    {
29
      ptr1[2] = 7;
30
      ptr1[3] = 9;
31
      ptr2[2] = 7;
32
    }
33
  }
34
  // CHECK: 7 7 9
35
  printf(" %d %d %d \n", ptr2[2], ptr1[2], ptr1[3]);
36
  free(ptr1);
37
#pragma omp target map(ptr2, ptr2[ : 100])
38
  { ptr2[1] = 6; }
39
  // CHECK: 6
40
  printf(" %d \n", ptr2[1]);
41
  free(ptr2);
42

43
  a[1] = 111;
44
  int *p = &a[0];
45
  // CHECK: 111
46
  printf("%d %p %p\n", p[1], p, &p); // 111 hst_p1 hst_p2
47
#pragma omp target data map(to : p[1 : 3]) map(p)
48
#pragma omp target data use_device_addr(p)
49
  {
50
#pragma omp target has_device_addr(p)
51
    {
52
      // CHECK: 111
53
      printf("%d %p %p\n", p[1], p, &p); // 111 dev_p1 dev_p2
54
      p[1] = 222;
55
      // CHECK: 222
56
      printf("%d %p %p\n", p[1], p, &p); // 222 dev_p1 dev_p2
57
    }
58
  }
59
  // CHECK: 111
60
  printf("%d %p %p\n", p[1], p, &p); // 111 hst_p1 hst_p2
61
  return 0;
62
}
63