1 // RUN: %libomptarget-compile-aarch64-unknown-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-aarch64-unknown-linux-gnu | %fcheck-aarch64-unknown-linux-gnu
2 // RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-powerpc64-ibm-linux-gnu | %fcheck-powerpc64-ibm-linux-gnu
3 // RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-powerpc64le-ibm-linux-gnu | %fcheck-powerpc64le-ibm-linux-gnu
4 // RUN: %libomptarget-compile-x86_64-pc-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-x86_64-pc-linux-gnu | %fcheck-x86_64-pc-linux-gnu -allow-empty
5 // RUN: %libomptarget-compile-nvptx64-nvidia-cuda && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-nvptx64-nvidia-cuda | %fcheck-nvptx64-nvidia-cuda -allow-empty
6
7 #include <assert.h>
8 #include <omp.h>
9 #include <stdio.h>
10 #include <stdlib.h>
11
12 const int magic_num = 7;
13
main(int argc,char * argv[])14 int main(int argc, char *argv[]) {
15 const int N = 128;
16 const int num_devices = omp_get_num_devices();
17
18 // No target device, just return
19 if (num_devices == 0) {
20 printf("PASS\n");
21 return 0;
22 }
23
24 const int src_device = 0;
25 int dst_device = num_devices - 1;
26
27 int length = N * sizeof(int);
28 int *src_ptr = omp_target_alloc(length, src_device);
29 int *dst_ptr = omp_target_alloc(length, dst_device);
30
31 assert(src_ptr && "src_ptr is NULL");
32 assert(dst_ptr && "dst_ptr is NULL");
33
34 #pragma omp target teams distribute parallel for device(src_device) \
35 is_device_ptr(src_ptr)
36 for (int i = 0; i < N; ++i) {
37 src_ptr[i] = magic_num;
38 }
39
40 int rc =
41 omp_target_memcpy(dst_ptr, src_ptr, length, 0, 0, dst_device, src_device);
42
43 assert(rc == 0 && "error in omp_target_memcpy");
44
45 int *buffer = malloc(length);
46
47 assert(buffer && "failed to allocate host buffer");
48
49 #pragma omp target teams distribute parallel for device(dst_device) \
50 map(from: buffer[0:N]) is_device_ptr(dst_ptr)
51 for (int i = 0; i < N; ++i) {
52 buffer[i] = dst_ptr[i] + magic_num;
53 }
54
55 for (int i = 0; i < N; ++i)
56 assert(buffer[i] == 2 * magic_num);
57
58 printf("PASS\n");
59
60 // Free host and device memory
61 free(buffer);
62 omp_target_free(src_ptr, src_device);
63 omp_target_free(dst_ptr, dst_device);
64
65 return 0;
66 }
67
68 // CHECK: PASS
69