| 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(size: sizeof(int) * 100); |
|---|
| 19 | int *ptr2; |
|---|
| 20 | ptr2 = (int *)malloc(size: sizeof(int) * 100); |
|---|
| 21 | #pragma omp target map(ptr1, ptr1[ : 100]) |
|---|
| 22 | { ptr1[1] = 6; } |
|---|
| 23 | // CHECK: 6 |
|---|
| 24 | printf(format: " %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(format: " %d %d %d \n" , ptr2[2], ptr1[2], ptr1[3]); |
|---|
| 36 | free(ptr: ptr1); |
|---|
| 37 | #pragma omp target map(ptr2, ptr2[ : 100]) |
|---|
| 38 | { ptr2[1] = 6; } |
|---|
| 39 | // CHECK: 6 |
|---|
| 40 | printf(format: " %d \n" , ptr2[1]); |
|---|
| 41 | free(ptr: ptr2); |
|---|
| 42 | |
|---|
| 43 | a[1] = 111; |
|---|
| 44 | int *p = &a[0]; |
|---|
| 45 | // CHECK: 111 |
|---|
| 46 | printf(format: "%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(format: "%d %p %p\n" , p[1], p, &p); // 111 dev_p1 dev_p2 |
|---|
| 54 | p[1] = 222; |
|---|
| 55 | // CHECK: 222 |
|---|
| 56 | printf(format: "%d %p %p\n" , p[1], p, &p); // 222 dev_p1 dev_p2 |
|---|
| 57 | } |
|---|
| 58 | } |
|---|
| 59 | // CHECK: 111 |
|---|
| 60 | printf(format: "%d %p %p\n" , p[1], p, &p); // 111 hst_p1 hst_p2 |
|---|
| 61 | return 0; |
|---|
| 62 | } |
|---|
| 63 | |
|---|
