1 // RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu 2 // RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu 3 // RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu 4 // RUN: %libomptarget-compile-run-and-check-x86_64-unknown-linux-gnu 5 // RUN: %libomptarget-compile-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